From f9187e0f1d68f76faef79ba862f22fa216152588 Mon Sep 17 00:00:00 2001 From: Not Zed Date: Thu, 9 Dec 2021 10:43:27 +1030 Subject: [PATCH] Move to using api.Frame for all marshalling operations Cleanup dead code --- Makefile | 3 +- README | 6 + .../au/notzed/zcl/test/TestAllocate.java | 84 ----- .../au/notzed/zcl/test/TestCopies.java | 176 +---------- src/notzed.zcl/classes/api/Allocator.java | 42 --- src/notzed.zcl/classes/api/Callback.java | 1 - src/notzed.zcl/classes/api/Frame.java | 108 ++++++- src/notzed.zcl/classes/api/Memory.java | 298 ++---------------- src/notzed.zcl/classes/api/Native.java | 99 +----- .../classes/au/notzed/zcl/CLBuffer.java | 6 +- .../classes/au/notzed/zcl/CLBufferInfo.java | 9 +- .../classes/au/notzed/zcl/CLCommandQueue.java | 133 ++++---- .../classes/au/notzed/zcl/CLContext.java | 217 +++++++------ .../au/notzed/zcl/CLContextNotify.java | 9 +- .../classes/au/notzed/zcl/CLEvent.java | 10 +- .../classes/au/notzed/zcl/CLEventList.java | 9 +- .../classes/au/notzed/zcl/CLImage.java | 2 +- .../classes/au/notzed/zcl/CLImageDesc.java | 24 +- .../classes/au/notzed/zcl/CLImageFormat.java | 6 +- .../classes/au/notzed/zcl/CLKernel.java | 59 ++-- .../classes/au/notzed/zcl/CLMemory.java | 23 +- .../classes/au/notzed/zcl/CLObject.java | 112 +++---- .../classes/au/notzed/zcl/CLPlatform.java | 24 +- .../classes/au/notzed/zcl/CLProgram.java | 51 +-- .../classes/au/notzed/zcl/CLProperty.java | 8 +- .../classes/au/notzed/zcl/EventInfo.java | 3 +- .../tests/au/notzed/zcl/CLEventTest.java | 8 +- 27 files changed, 513 insertions(+), 1017 deletions(-) delete mode 100644 src/notzed.zcl/classes/api/Allocator.java diff --git a/Makefile b/Makefile index 4047315..796d6d6 100644 --- a/Makefile +++ b/Makefile @@ -28,8 +28,7 @@ notzed.zcl.demo_DEMOS=au.notzed.zcl.tools.clinfo \ notzed.zcl.fxdemo_DEMOS=fxdemo.fract.Mandelbrot fxdemo.fract.Test -DEMOFLAGS=--add-exports jdk.incubator.foreign/jdk.incubator.foreign.unsafe=notzed.zcl \ - -Dforeign.restricted=permit +DEMOFLAGS=--enable-native-access=notzed.zcl # module class basename(class) define java_demo= diff --git a/README b/README index 673c5b7..4c00a31 100644 --- a/README +++ b/README @@ -15,6 +15,12 @@ make run-clinfo make run-Mandelbrot make run-Mandelbrot ARGV="--gamma=15 --rotate=0" +foreign-abi branch TODO + + - api.Frame should probably be a SegmentAllocator + - where pointers are passed in to native calls, MemoryAddress is used for the mid-level generated api - these + should probably be Addressable to avoid the clutter of .address() invocations in the callee. + - some apis take or use bytebuffer, maybe should just be MemorySegment INTRODUCTION ------------ diff --git a/src/notzed.zcl.demo/classes/au/notzed/zcl/test/TestAllocate.java b/src/notzed.zcl.demo/classes/au/notzed/zcl/test/TestAllocate.java index 9789955..89b1d7f 100644 --- a/src/notzed.zcl.demo/classes/au/notzed/zcl/test/TestAllocate.java +++ b/src/notzed.zcl.demo/classes/au/notzed/zcl/test/TestAllocate.java @@ -71,90 +71,6 @@ public class TestAllocate { System.out.printf(" %12.9f 2xC_POINTER\n", m * 1E-9); results.compute(String.format("2xC_POINTER"), (k, v) -> v == null ? m : min(v, m)); } - if (true) { - - now = System.nanoTime(); - for (int l = 0; l < X; l++) { - try (Allocator a = api.Memory.stack()) { - MemorySegment arg0 = a.allocs(8); - MemorySegment arg1 = a.allocs(8); - MemorySegment arg2 = a.allocs(8); - MemorySegment arg3 = a.allocs(8); - } - } - long s = System.nanoTime() - now; - System.out.printf(" %12.9f 4x8 stack\n", s * 1E-9); - results.compute(String.format("4x8 stack"), (k, v) -> v == null ? s : min(v, s)); - } - if (false) { - - now = System.nanoTime(); - for (int l = 0; l < X; l++) { - try (Allocator a = api.Memory.stack()) { - MemorySegment arg0 = a.allocs(8); - MemorySegment arg1 = a.allocs(8); - } - } - long h = System.nanoTime() - now; - System.out.printf(" %12.9f 2x8 stack\n", h * 1E-9); - results.compute(String.format("2x8 stack"), (k, v) -> v == null ? h : min(v, h)); - } - if (true) { - - now = System.nanoTime(); - for (int l = 0; l < X; l++) { - try (Allocator a = api.Memory.stack2()) { - MemorySegment arg0 = a.allocs(8); - MemorySegment arg1 = a.allocs(8); - MemorySegment arg2 = a.allocs(8); - MemorySegment arg3 = a.allocs(8); - } - } - long r = System.nanoTime() - now; - System.out.printf(" %12.9f 4x8 stack2\n", r * 1E-9); - results.compute(String.format("4x8 stack2"), (k, v) -> v == null ? r : min(v, r)); - } - if (false) { - - now = System.nanoTime(); - for (int l = 0; l < X; l++) { - try (Allocator a = api.Memory.stack2()) { - MemorySegment arg0 = a.allocs(8); - MemorySegment arg1 = a.allocs(8); - } - } - long p = System.nanoTime() - now; - System.out.printf(" %12.9f 2x8 stack2\n", p * 1E-9); - results.compute(String.format("2x8 stack2"), (k, v) -> v == null ? p : min(v, p)); - } - if (false) { - - now = System.nanoTime(); - for (int l = 0; l < X; l++) { - try (Allocator a = api.Memory.stack3()) { - MemorySegment arg0 = a.allocs(8); - MemorySegment arg1 = a.allocs(8); - } - } - long g = System.nanoTime() - now; - System.out.printf(" %12.9f 2x8 stack3\n", g * 1E-9); - results.compute(String.format("2x8 stack3"), (k, v) -> v == null ? g : min(v, g)); - } - if (false) { - - now = System.nanoTime(); - for (int l = 0; l < X; l++) { - try (Allocator a = api.Memory.stack3()) { - MemorySegment arg0 = a.allocs(8); - MemorySegment arg1 = a.allocs(8); - MemorySegment arg2 = a.allocs(8); - MemorySegment arg3 = a.allocs(8); - } - } - long q = System.nanoTime() - now; - System.out.printf(" %12.9f 4x8 stack3\n", q * 1E-9); - results.compute(String.format("4x8 stack3"), (k, v) -> v == null ? q : min(v, q)); - } if (false) { now = System.nanoTime(); diff --git a/src/notzed.zcl.demo/classes/au/notzed/zcl/test/TestCopies.java b/src/notzed.zcl.demo/classes/au/notzed/zcl/test/TestCopies.java index 7ce89e9..1d7481a 100644 --- a/src/notzed.zcl.demo/classes/au/notzed/zcl/test/TestCopies.java +++ b/src/notzed.zcl.demo/classes/au/notzed/zcl/test/TestCopies.java @@ -54,137 +54,11 @@ public class TestCopies { int X = 1024 * 1024 * 10; for (int c = 0; c < 5; c++) { - { - long now; - if (true) { - now = System.nanoTime(); - for (int l = 0; l < X; l++) { - try (ResourceScope scope = ResourceScope.newConfinedScope()) { - MemorySegment arg0 = MemorySegment.allocateNative(8, 8, scope); - MemorySegment arg1 = MemorySegment.allocateNative(8, 8, scope); - MemorySegment arg2 = MemorySegment.allocateNative(8, 8, scope); - MemorySegment arg3 = MemorySegment.allocateNative(8, 8, scope); - } - } - long t = System.nanoTime() - now; - System.out.printf(" %12.9f 4x8 scope\n", t * 1E-9); - results.compute(String.format("4x8 scope"), (k, v) -> v == null ? t : min(v, t)); - - now = System.nanoTime(); - for (int l = 0; l < X; l++) { - try (ResourceScope scope = ResourceScope.newConfinedScope()) { - MemorySegment arg0 = MemorySegment.allocateNative(8, 8, scope); - MemorySegment arg1 = MemorySegment.allocateNative(8, 8, scope); - } - } - long n = System.nanoTime() - now; - System.out.printf(" %12.9f 2x8 scope\n", n * 1E-9); - results.compute(String.format("2x8 scope"), (k, v) -> v == null ? n : min(v, n)); - - now = System.nanoTime(); - for (int l = 0; l < X; l++) { - try (ResourceScope scope = ResourceScope.newConfinedScope()) { - MemorySegment arg0 = MemorySegment.allocateNative(CLinker.C_POINTER, scope); - MemorySegment arg1 = MemorySegment.allocateNative(CLinker.C_POINTER, scope); - MemorySegment arg2 = MemorySegment.allocateNative(CLinker.C_POINTER, scope); - MemorySegment arg3 = MemorySegment.allocateNative(CLinker.C_POINTER, scope); - } - } - long u = System.nanoTime() - now; - System.out.printf(" %12.9f 4xC_POINTER\n", u * 1E-9); - results.compute(String.format("4xC_POINTER"), (k, v) -> v == null ? u : min(v, u)); - - now = System.nanoTime(); - for (int l = 0; l < X; l++) { - try (ResourceScope scope = ResourceScope.newConfinedScope()) { - MemorySegment arg0 = MemorySegment.allocateNative(CLinker.C_POINTER, scope); - MemorySegment arg1 = MemorySegment.allocateNative(CLinker.C_POINTER, scope); - } - } - long m = System.nanoTime() - now; - System.out.printf(" %12.9f 2xC_POINTER\n", m * 1E-9); - results.compute(String.format("2xC_POINTER"), (k, v) -> v == null ? m : min(v, m)); - } - - now = System.nanoTime(); - for (int l = 0; l < X; l++) { - try (Allocator a = api.Memory.stack()) { - MemorySegment arg0 = a.allocs(8); - MemorySegment arg1 = a.allocs(8); - MemorySegment arg2 = a.allocs(8); - MemorySegment arg3 = a.allocs(8); - } - } - long s = System.nanoTime() - now; - System.out.printf(" %12.9f 4x8 stack\n", s * 1E-9); - results.compute(String.format("4x8 stack"), (k, v) -> v == null ? s : min(v, s)); - - now = System.nanoTime(); - for (int l = 0; l < X; l++) { - try (Allocator a = api.Memory.stack()) { - MemorySegment arg0 = a.allocs(8); - MemorySegment arg1 = a.allocs(8); - } - } - long h = System.nanoTime() - now; - System.out.printf(" %12.9f 2x8 stack\n", h * 1E-9); - results.compute(String.format("2x8 stack"), (k, v) -> v == null ? h : min(v, h)); - - now = System.nanoTime(); - for (int l = 0; l < X; l++) { - try (Allocator a = api.Memory.stack2()) { - MemorySegment arg0 = a.allocs(8); - MemorySegment arg1 = a.allocs(8); - MemorySegment arg2 = a.allocs(8); - MemorySegment arg3 = a.allocs(8); - } - } - long r = System.nanoTime() - now; - System.out.printf(" %12.9f 4x8 stack2\n", r * 1E-9); - results.compute(String.format("4x8 stack2"), (k, v) -> v == null ? r : min(v, r)); - - now = System.nanoTime(); - for (int l = 0; l < X; l++) { - try (Allocator a = api.Memory.stack2()) { - MemorySegment arg0 = a.allocs(8); - MemorySegment arg1 = a.allocs(8); - } - } - long p = System.nanoTime() - now; - System.out.printf(" %12.9f 2x8 stack2\n", p * 1E-9); - results.compute(String.format("2x8 stack2"), (k, v) -> v == null ? p : min(v, p)); - - now = System.nanoTime(); - for (int l = 0; l < X; l++) { - try (Allocator a = api.Memory.stack3()) { - MemorySegment arg0 = a.allocs(8); - MemorySegment arg1 = a.allocs(8); - } - } - long g = System.nanoTime() - now; - System.out.printf(" %12.9f 2x8 stack3\n", g * 1E-9); - results.compute(String.format("2x8 stack3"), (k, v) -> v == null ? g : min(v, g)); - - now = System.nanoTime(); - for (int l = 0; l < X; l++) { - try (Allocator a = api.Memory.stack3()) { - MemorySegment arg0 = a.allocs(8); - MemorySegment arg1 = a.allocs(8); - MemorySegment arg2 = a.allocs(8); - MemorySegment arg3 = a.allocs(8); - } - } - long q = System.nanoTime() - now; - System.out.printf(" %12.9f 4x8 stack3\n", q * 1E-9); - results.compute(String.format("4x8 stack3"), (k, v) -> v == null ? q : min(v, q)); - - } - for (int len: lengths) { long[] data = new long[len]; long now; - if (false) { + if (true) { now = System.nanoTime(); for (int l = 0; l < X; l++) { try (ResourceScope scope = ResourceScope.newConfinedScope()) { @@ -197,12 +71,12 @@ public class TestCopies { results.compute(String.format("%3d copyLoop", len), (k, v) -> v == null ? t : min(v, t)); } - if (false) { + if (true) { now = System.nanoTime(); for (int l = 0; l < X; l++) { try (ResourceScope scope = ResourceScope.newConfinedScope()) { MemorySegment seg = MemorySegment.allocateNative(data.length * 8, scope); - //copyBulk(data, seg); + copyBulk(data, seg); } } long s = System.nanoTime() - now; @@ -211,44 +85,16 @@ public class TestCopies { } } - if (false) - for (int len: lengths) { - long[] data = new long[len]; - long now; - - now = System.nanoTime(); - try (ResourceScope scope = ResourceScope.newConfinedScope()) { - MemorySegment seg = MemorySegment.allocateNative(data.length * 8, scope); - for (int l = 0; l < X; l++) { - copyLoop(data, seg); - } - } - long t = System.nanoTime() - now; - System.out.printf(" %12.9f %3d copyLoop pre-alloc\n", t * 1E-9, len); - results.compute(String.format("%3d copyLoop pre-alloc", len), (k, v) -> v == null ? t : min(v, t)); - - now = System.nanoTime(); - try (ResourceScope scope = ResourceScope.newConfinedScope()) { - MemorySegment seg = MemorySegment.allocateNative(data.length * 8, scope); - for (int l = 0; l < X; l++) { - copyBulk(data, seg); - } - } - long s = System.nanoTime() - now; - System.out.printf(" %12.9f %3d copyBulk pre-alloc\n", s * 1E-9, len); - results.compute(String.format("%3d copyBulk pre-alloc", len), (k, v) -> v == null ? s : min(v, s)); - } - - // if have stack allocator: + // custom allocator for (int len: lengths) { long[] data = new long[len]; long now; - if (false) { + if (true) { now = System.nanoTime(); for (int l = 0; l < X; l++) { - try (Allocator a = api.Memory.stack()) { - MemorySegment seg = a.allocs(data.length * 8); + try (Frame frame = api.Memory.createFrame()) { + MemorySegment seg = frame.allocate(CLinker.C_POINTER, data.length); copyLoop(data, seg); } } @@ -257,12 +103,12 @@ public class TestCopies { results.compute(String.format("%3d copyLoop stack", len), (k, v) -> v == null ? t : min(v, t)); } - if (false) { + if (true) { now = System.nanoTime(); for (int l = 0; l < X; l++) { - try (Allocator a = api.Memory.stack()) { - MemorySegment seg = a.allocs(data.length * 8); - //copyBulk(data, seg); + try (Frame frame = api.Memory.createFrame()) { + MemorySegment seg = frame.allocate(CLinker.C_POINTER, data.length); + copyBulk(data, seg); } } long s = System.nanoTime() - now; diff --git a/src/notzed.zcl/classes/api/Allocator.java b/src/notzed.zcl/classes/api/Allocator.java deleted file mode 100644 index 87f96f9..0000000 --- a/src/notzed.zcl/classes/api/Allocator.java +++ /dev/null @@ -1,42 +0,0 @@ -/* - * Copyright (C) 2020 Michael Zucchi - * - * This program is free software: you can redistribute it and/or modify - * it under the terms of the GNU General Public License as published by - * the Free Software Foundation, either version 3 of the License, or - * (at your option) any later version. - * - * This program is distributed in the hope that it will be useful, - * but WITHOUT ANY WARRANTY; without even the implied warranty of - * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the - * GNU General Public License for more details. - * - * You should have received a copy of the GNU General Public License - * along with this program. If not, see . - */ - -package api; - -import jdk.incubator.foreign.MemoryAddress; -import jdk.incubator.foreign.MemorySegment; - -/** - * An interface to a 'pool' type allocator. - * - * That is, one where individual items cannot necessarily be freed. - */ -public interface Allocator extends AutoCloseable { - @Override - public void close(); - - /** - * allocate memory - */ - public MemoryAddress alloca(long size); - - /** - Allocate bounded memory. - Whether you can close the segment depends on the implementation. - */ - public MemorySegment allocs(long size); -} diff --git a/src/notzed.zcl/classes/api/Callback.java b/src/notzed.zcl/classes/api/Callback.java index 7e28435..c6c2208 100644 --- a/src/notzed.zcl/classes/api/Callback.java +++ b/src/notzed.zcl/classes/api/Callback.java @@ -55,7 +55,6 @@ public class Callback extends Native implements AutoCloseable { }; private static void release(MemoryAddress p) { - freeUpcallStub(p); } @Override diff --git a/src/notzed.zcl/classes/api/Frame.java b/src/notzed.zcl/classes/api/Frame.java index b5ca8c9..bb89a5c 100644 --- a/src/notzed.zcl/classes/api/Frame.java +++ b/src/notzed.zcl/classes/api/Frame.java @@ -16,13 +16,25 @@ */ package api; +import jdk.incubator.foreign.CLinker; import jdk.incubator.foreign.MemoryAccess; +import jdk.incubator.foreign.MemoryAddress; +import jdk.incubator.foreign.MemoryLayout; import jdk.incubator.foreign.MemorySegment; +// TODO: after writing this i discovered SegmentAllocator, perhaps should be one of those. public interface Frame extends AutoCloseable { MemorySegment allocate(long size, long alignment); + /** + * Copy string to native memory. + * + * @param value may be null + * @return + */ + MemorySegment copy(String value); + @Override public void close(); @@ -30,12 +42,50 @@ public interface Frame extends AutoCloseable { return allocate(size, 1); } + default MemorySegment allocate(MemoryLayout layout) { + return allocate(layout.byteSize(), layout.byteAlignment()); + } + + default MemorySegment allocate(MemoryLayout layout, long count) { + return allocate(layout.byteSize() * count, layout.byteAlignment()); + } + + default MemorySegment allocateInt() { + return allocate(CLinker.C_INT); + } + + default MemorySegment allocateInt(int count) { + return allocate(CLinker.C_INT, count); + } + + default MemorySegment allocateLong() { + return allocate(CLinker.C_LONG); + } + + default MemorySegment allocateLong(int count) { + return allocate(CLinker.C_LONG, count); + } + + default MemorySegment allocateAddress() { + return allocate(CLinker.C_POINTER); + } + + default MemorySegment allocateAddress(int count) { + return allocate(CLinker.C_POINTER, count); + } + default MemorySegment copy(byte value) { MemorySegment mem = allocate(1); MemoryAccess.setByte(mem, value); return mem; } + default MemorySegment copy(short value) { + MemorySegment mem = allocate(2, 2); + MemoryAccess.setShort(mem, value); + return mem; + } + default MemorySegment copy(int value) { MemorySegment mem = allocate(4, 4); MemoryAccess.setInt(mem, value); @@ -48,22 +98,66 @@ public interface Frame extends AutoCloseable { return mem; } - default MemorySegment copy(byte[] value, int alignment) { - MemorySegment mem = allocate(value.length, alignment); + default MemorySegment copy(float value) { + MemorySegment mem = allocate(4, 4); + MemoryAccess.setFloat(mem, value); + return mem; + } + + default MemorySegment copy(double value) { + MemorySegment mem = allocate(8, 8); + MemoryAccess.setDouble(mem, value); + return mem; + } + + default MemorySegment copy(byte[] value) { + MemorySegment mem = allocate(value.length, 1); mem.copyFrom(MemorySegment.ofArray(value)); return mem; } - default MemorySegment copy(int[] value, int alignment) { - MemorySegment mem = allocate(value.length * 4, alignment); + default MemorySegment copy(int[] value) { + MemorySegment mem = allocate(value.length * 4, 4); mem.copyFrom(MemorySegment.ofArray(value)); return mem; } - default MemorySegment copy(long[] value, int alignment) { - MemorySegment mem = allocate(value.length * 8, alignment); + default MemorySegment copy(long[] value) { + MemorySegment mem = allocate(value.length * 8, 8); mem.copyFrom(MemorySegment.ofArray(value)); return mem; } - + + default MemorySegment copy(float[] value) { + MemorySegment mem = allocate(value.length * 4, 4); + mem.copyFrom(MemorySegment.ofArray(value)); + return mem; + } + + default MemorySegment copy(T[] array) { + MemorySegment mem = allocateAddress(array.length); + for (int i = 0; i < array.length; i++) + MemoryAccess.setAddressAtIndex(mem, i, array[i].address()); + return mem; + } + + default MemorySegment copy(T value) { + return copy(value.address()); + } + + default MemorySegment copy(MemoryAddress value) { + MemorySegment mem = allocateAddress(); + MemoryAccess.setAddress(mem, value); + return mem; + } + + // create an array pointing to strings + default MemorySegment copy(String[] array) { + MemorySegment list = allocateAddress(array.length); + for (int i = 0; i < array.length; i++) { + MemoryAccess.setAddressAtIndex(list, i, copy(array[i])); + } + return list; + } + } diff --git a/src/notzed.zcl/classes/api/Memory.java b/src/notzed.zcl/classes/api/Memory.java index 86dbe17..4b2765d 100644 --- a/src/notzed.zcl/classes/api/Memory.java +++ b/src/notzed.zcl/classes/api/Memory.java @@ -17,292 +17,37 @@ package api; import jdk.incubator.foreign.*; -import java.lang.invoke.MethodHandle; -import java.lang.invoke.MethodType; import java.lang.ref.Cleaner; /** - * A utility library for memory operations including a stack allocator. + * A utility for memory operations including a stack allocator. *

* The stack allocator works like this - * try (Memory.Frame f = Memory.frame()) { - * MemoryAddress a = Memory.alloca(size); + * 

+ * try (Frame f = Memory.createFrame()) {
+ *		MemorySegment a = f.allocate(size);
  * }
+ *
* Any memory allocated is freed when the frame is closed. *

- * Note that unlike C's alloca() the frame scope is not per-function - * but application defined. + * This is MUCH faster than using MemorySegment.allocate(). */ public class Memory { - private static final ThreadLocal stacks = ThreadLocal.withInitial(() -> Native.resolve(malloc(4096), (p) -> new Stack(p, 4096))); - private static final ThreadLocal stackAllocators = new ThreadLocal<>(); - - /* older idea */ - // eh, this isn't really necessary, it's just a safety check/debug? - //private static final ThreadLocal frames = new ThreadLocal<>(); - /** - * Alloc memory in the current frame. The size is not enforced. - */ - /* - public static MemoryAddress alloca(long size) { - if (frames.get() == null) - throw new UnsupportedOperationException("Must have a frame"); - return stack().alloca(size); - }*/ - /** - * Allocate a segment in the current frame. The segment must not be closed. - */ - /* - public static MemorySegment slicea(long size) { - if (frames.get() == null) - throw new UnsupportedOperationException("Must have a frame"); - return stack().slicea(size); - } - - public static Frame frame() { - Stack stack = stack(); - long fp = stack.sp; - Frame old = frames.get(); - Frame gnu = () -> { - frames.set(old); - stack.sp = fp; - }; - frames.set(gnu); - - return gnu; - } - - public interface Frame extends AutoCloseable { - @Override - public void close(); - } - */ - static class Stack extends Native { - - final MemorySegment stack; - final MemoryAddress base; - long sp; - Thread origin; - - Stack(MemoryAddress p, long size) { - super(p); - - // TODO: use cleanup stuff, perhaps owner thread - stack = MemorySegment.globalNativeSegment().asSlice(addr(), size); - base = stack.address(); - sp = size; - - origin = Thread.currentThread(); - } - - public String toString() { - return "api.Memory$Stack thread=" + origin.toString(); - } - - private static void release(MemoryAddress p) { - System.err.printf("** release stack " + p); - free(p); - } - - public MemoryAddress alloca(long size) { - sp -= (size + 7) & ~7; - return base.addOffset(sp); - } - - public MemorySegment allocs(long size) { - sp -= (size + 7) & ~7; - return stack.asSlice(sp, size); - } - - public MemorySegment slicea(long size) { - sp -= (size + 7) & ~7; - return stack.asSlice(sp, size); - } - } - - /* - Alternatative version. - This allows different allocators to be used for functions that use one. - Individual memory cannot be freed. - */ - /** - * Create a stack allocator. - * The stack allocator uses thread-specific backing buffer. - * This should only be used for small allocations. - */ - public static Allocator stack() { - Stack stack = stacks.get(); - long fp = stack.sp; - Allocator old = stackAllocators.get(); - Thread me = Thread.currentThread(); - - Allocator gnu = new Allocator() { - public void close() { - stack.sp = fp; - stackAllocators.set(old); - } - - public MemoryAddress alloca(long size) { - if (stackAllocators.get() != this || me != Thread.currentThread()) - throw new IllegalStateException(); - return stack.alloca(size); - } - - public MemorySegment allocs(long size) { - if (stackAllocators.get() != this || me != Thread.currentThread()) - throw new IllegalStateException(); - return stack.allocs(size); - } - }; - - stackAllocators.set(gnu); - return gnu; - } - - static final MethodHandle malloc; - static final MethodHandle free; - - static { - // So for whatever reason it's been decided that MemorySegment can't be freed normally on another thread - CLinker abi = CLinker.getInstance(); - SymbolLookup libc = CLinker.systemLookup(); - - try { - malloc = abi.downcallHandle(libc.lookup("malloc").get(), - MethodType.methodType(MemoryAddress.class, - long.class), - FunctionDescriptor.of(CLinker.C_POINTER, - CLinker.C_LONG)); - free = abi.downcallHandle(libc.lookup("free").get(), - MethodType.methodType(void.class, - MemoryAddress.class), - FunctionDescriptor.ofVoid(CLinker.C_POINTER)); - } catch (Exception x) { - throw new RuntimeException(x); - } - } - - /** - * Create a sized memory segment from a segment allocated with malloc. - * Closing this segment has no effect on the original, and the original must be used for that. - */ - public static MemorySegment ofNative(MemoryAddress addr, long size) { - return MemorySegment.globalNativeSegment().asSlice(addr, size); - } - - /** - * Get the physical address. I mean how is this an "offset"? - */ - public static long toLong(MemoryAddress addr) { - return addr.toRawLongValue(); - } - - /** - * Allocate C memory. - * This is not usable by java directly, see memsize() - */ - public static MemoryAddress malloc(long size) { - try { - MemoryAddress addr = (MemoryAddress)malloc.invokeExact(size); - - return addr; - } catch (Throwable t) { - throw new RuntimeException(t); - } - } - - public static void free(MemoryAddress ptr) { - try { - free.invokeExact(ptr); - } catch (Throwable t) { - throw new RuntimeException(t); - } - } - - interface ThreadAllocator { - - Allocator create(); - } - // segment based stack without using malloc, locked to single thread - private static final ThreadLocal stacks2 = ThreadLocal.withInitial(() -> new ThreadAllocator() { - private final ResourceScope scope = ResourceScope.newConfinedScope(); - private final MemorySegment root = MemorySegment.allocateNative(4096, 4096, scope); - private MemorySegment base = root; - - //{ - // System.out.printf("0x%016x init stack\n", root.address().toRawLongValue()); - //} - @Override - public Allocator create() { - MemorySegment here = base; - - //System.out.printf("0x%016x new frame\n", here.address().toRawLongValue()); - return new Allocator() { - @Override - public void close() { - base = here; - //System.out.printf("0x%016x close frame\n", here.address().toRawLongValue()); - } - - @Override - public MemoryAddress alloca(long size) { - return allocs(size).address(); - } - - @Override - public MemorySegment allocs(long size) { - long alloc = (size + 7) & ~7; - MemorySegment seg = base.asSlice(0, size); - - //System.out.printf("0x%016x alloc %d\n", base.address().toRawLongValue(), size); - base = base.asSlice(alloc); - return seg; - } - }; - } - }); - - public static Allocator stack2() { - return stacks2.get().create(); - } - - public static Allocator stack3() { - return new Allocator() { - private final ResourceScope scope = ResourceScope.newConfinedScope(); - private final MemorySegment root = MemorySegment.allocateNative(4096, 4096, scope); - private MemorySegment base = root; - - @Override - public void close() { - scope.close(); - } - - @Override - public MemoryAddress alloca(long size) { - return allocs(size).address(); - } - - @Override - public MemorySegment allocs(long size) { - long alloc = (size + 7) & ~7; - MemorySegment seg = base.asSlice(0, size); - - //System.out.printf("0x%016x alloc %d\n", base.address().toRawLongValue(), size); - base = base.asSlice(alloc); - return seg; - } - }; + static final ResourceScope scope = ResourceScope.newSharedScope(Cleaner.create()); + private static final ThreadLocal stacks = ThreadLocal.withInitial(() -> new Stack(scope)); + public static Frame createFrame() { + return stacks.get().createFrame(); } - static class Stack4 { + static class Stack { private final MemorySegment stack; private long sp; private Thread thread = Thread.currentThread(); - Stack4(ResourceScope scope) { + Stack(ResourceScope scope) { stack = MemorySegment.allocateNative(4096, 4096, scope); sp = 4096; } @@ -330,6 +75,17 @@ public class Memory { } } + @Override + public MemorySegment copy(String value) { + if (value != null) { + if (scope == null) + scope = ResourceScope.newConfinedScope(); + return CLinker.toCString(value, scope); + } else { + return MemorySegment.globalNativeSegment().asSlice(0, 0); + } + } + @Override public void close() { sp = tos; @@ -343,12 +99,4 @@ public class Memory { } } - - static final ResourceScope scope4 = ResourceScope.newSharedScope(Cleaner.create()); - private static final ThreadLocal stacks4 = ThreadLocal.withInitial(() -> new Stack4(scope4)); - - public static Frame createFrame() { - return stacks4.get().createFrame(); - } - } diff --git a/src/notzed.zcl/classes/api/Native.java b/src/notzed.zcl/classes/api/Native.java index 68dde60..550aeb5 100644 --- a/src/notzed.zcl/classes/api/Native.java +++ b/src/notzed.zcl/classes/api/Native.java @@ -19,7 +19,6 @@ package api; import java.io.StringReader; import java.lang.invoke.*; import java.lang.reflect.Method; -import java.nio.ByteOrder; import java.util.ArrayList; import java.util.Arrays; import java.util.List; @@ -46,37 +45,12 @@ import java.lang.System.Logger.Level; *

* FIXME: there are MemorySegment based accessors for primitive types now, use those */ -public class Native { +public class Native implements Addressable { private final MemoryAddress p; private final static boolean dolog = true; - public final static VarHandle byteHandle = MemoryHandles.varHandle(byte.class, ByteOrder.nativeOrder()); - public final static VarHandle shortHandle = MemoryHandles.varHandle(short.class, ByteOrder.nativeOrder()); - //final static VarHandle intHandle = MemoryHandles.varHandle(int.class, ByteOrder.nativeOrder()); - public final static VarHandle intHandle = CLinker.C_INT.varHandle(int.class); - public final static VarHandle longHandle = MemoryHandles.varHandle(long.class, ByteOrder.nativeOrder()); - public final static VarHandle floatHandle = MemoryHandles.varHandle(float.class, ByteOrder.nativeOrder()); - public final static VarHandle doubleHandle = MemoryHandles.varHandle(double.class, ByteOrder.nativeOrder()); - public final static VarHandle addrHandle = MemoryHandles.asAddressVarHandle(MemoryHandles.varHandle(long.class, ByteOrder.nativeOrder())); - - final static SequenceLayout byteVLayout = MemoryLayout.sequenceLayout(CLinker.C_CHAR); - final static SequenceLayout shortVLayout = MemoryLayout.sequenceLayout(CLinker.C_SHORT); - final static SequenceLayout intVLayout = MemoryLayout.sequenceLayout(CLinker.C_INT); - final static SequenceLayout longVLayout = MemoryLayout.sequenceLayout(CLinker.C_LONG); - final static SequenceLayout floatVLayout = MemoryLayout.sequenceLayout(CLinker.C_FLOAT); - final static SequenceLayout doubleVLayout = MemoryLayout.sequenceLayout(CLinker.C_DOUBLE); - final static SequenceLayout addrVLayout = MemoryLayout.sequenceLayout(CLinker.C_POINTER); - - public final static VarHandle byteVHandle = byteVLayout.varHandle(byte.class, MemoryLayout.PathElement.sequenceElement()); - public final static VarHandle shortVHandle = shortVLayout.varHandle(short.class, MemoryLayout.PathElement.sequenceElement()); - public final static VarHandle intVHandle = intVLayout.varHandle(int.class, MemoryLayout.PathElement.sequenceElement()); - public final static VarHandle longVHandle = longVLayout.varHandle(long.class, MemoryLayout.PathElement.sequenceElement()); - public final static VarHandle floatVHandle = floatVLayout.varHandle(float.class, MemoryLayout.PathElement.sequenceElement()); - public final static VarHandle doubleVHandle = doubleVLayout.varHandle(double.class, MemoryLayout.PathElement.sequenceElement()); - public final static VarHandle addrVHandle = addrVLayout.varHandle(long.class, MemoryLayout.PathElement.sequenceElement()); - protected Native(MemoryAddress p) { this.p = p; } @@ -85,56 +59,17 @@ public class Native { return System.getLogger("notzed.native"); } - public MemoryAddress addr() { + public MemoryAddress address() { return p; } public static MemoryAddress addr(Native o) { - return o != null ? o.addr() : MemoryAddress.NULL; + return o != null ? o.address() : MemoryAddress.NULL; } public static MemoryAddress addr(MemorySegment o) { return o != null ? o.address() : MemoryAddress.NULL; } - - /* helpers - java to native */ - public static MemorySegment allocAddrV(ResourceScope frame, int len) { - return MemorySegment.allocateNative(CLinker.C_POINTER.byteSize() * len, CLinker.C_POINTER.byteAlignment(), frame); - } - - public static MemorySegment allocV(ResourceScope frame, ValueLayout type, int len) { - return MemorySegment.allocateNative(type.byteSize() * len, type.byteAlignment(), frame); - } - - public static MemorySegment toAddrV(ResourceScope frame, T[] array, int len) { - MemorySegment list = MemorySegment.allocateNative(CLinker.C_POINTER.byteSize() * len, CLinker.C_POINTER.byteAlignment(), frame); - - for (int i = 0; i < len; i++) - MemoryAccess.setAddressAtIndex(list, i, array[i].addr()); - - return list; - } - - public static MemorySegment toAddrV(ResourceScope frame, T[] array) { - return toAddrV(frame, array, array.length); - } - - public static MemorySegment toLongV(ResourceScope frame, long[] array) { - MemorySegment list = MemorySegment.allocateNative(CLinker.C_LONG.byteSize() * array.length, CLinker.C_LONG.byteAlignment(), frame); - - for (int i = 0; i < array.length; i++) - MemoryAccess.setLongAtIndex(list, i, array[i]); - - return list; - } - - public static MemorySegment toCStringV(String[]array, ResourceScope frame) { - MemorySegment list = MemorySegment.allocateNative(CLinker.C_POINTER.byteSize() * array.length, CLinker.C_POINTER.byteAlignment(), frame); - for (int i=0;i T[] toObjectV(MemorySegment list, T[] array, Function create) { @@ -147,18 +82,10 @@ public class Native { return toObjectV(list, createArray.apply((int)(list.byteSize() >>> 3)), create); } - public static String toString(MemoryAddress cstr) { - return CLinker.toJavaString(cstr); - } - public static long[] toLongV(MemorySegment valp) { - int len = (int)(valp.byteSize() >>> 3); - long[] list = new long[len]; - - for (int i = 0; i < list.length; i++) - list[i] = MemoryAccess.getLongAtIndex(valp, i); - - return list; + long[] val = new long[(int)(valp.byteSize() >> 3)]; + MemorySegment.ofArray(val).copyFrom(valp); + return val; } /* abi stuff */ @@ -244,10 +171,6 @@ public class Native { } } - public static void freeUpcallStub(MemoryAddress addr) { - // NOOP? - } - public static void loadLibraries(String... libraries) { for (int i = 0; i < libraries.length; i++) System.loadLibrary(libraries[i]); @@ -549,7 +472,7 @@ public class Native { //if (dolog) // log().log(Level.DEBUG, () -> String.format(" resolv $%016x %s", Memory.toLong(p), create)); - if (Memory.toLong(p) == 0) + if (p.toRawLongValue() == 0) return null; // Instantiation needs to be synchronized for obvious reasons. @@ -571,7 +494,7 @@ public class Native { } { T x = o; - log().log(Level.DEBUG, () -> String.format(fmt, Memory.toLong(p), x.getClass().getName())); + log().log(Level.DEBUG, () -> String.format(fmt, p.toRawLongValue(), x.getClass().getName())); } } @@ -610,7 +533,7 @@ public class Native { if (ref != null) { if (dolog) - log().log(Level.DEBUG, () -> String.format(" force $%016x %s", Memory.toLong(p), getClass().getName())); + log().log(Level.DEBUG, () -> String.format(" force $%016x %s", p.toRawLongValue(), getClass().getName())); ref.enqueue(); } @@ -689,7 +612,7 @@ public class Native { try { if (p != null) { if (dolog) - log().log(Level.DEBUG, () -> String.format(" releas $%016x %s", Memory.toLong(p), jtype.getName())); + log().log(Level.DEBUG, () -> String.format(" releas $%016x %s", p.toRawLongValue(), jtype.getName())); Method mm = jtype.getDeclaredMethod("release", MemoryAddress.class); mm.setAccessible(true); @@ -763,7 +686,7 @@ public class Native { while (h != null) { Native o = h.get(); System.out.printf(" $%016x: %s %-40s %s\n", - Memory.toLong(h.p), + h.p.toRawLongValue(), o == null ? "dead" : "live", h.jtype.getName(), o); diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLBuffer.java b/src/notzed.zcl/classes/au/notzed/zcl/CLBuffer.java index f2c339f..d27aeeb 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLBuffer.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLBuffer.java @@ -65,13 +65,13 @@ public class CLBuffer extends CLMemory { public CLBuffer createSubBuffer(long flags, CLBufferInfo info) throws CLException { requireAPIVersion(CLPlatform.VERSION_1_1); - try (Allocator frame = Memory.stack()) { - MemorySegment pres = frame.allocs(8); + try (Frame frame = Memory.createFrame()) { + MemorySegment pres = frame.allocateInt(); MemorySegment pinfo = info.toNative(frame); MemoryAddress b; int res; - b = clCreateSubBuffer(addr(), flags, CL_BUFFER_CREATE_TYPE_REGION, pinfo.address(), pres.address()); + b = clCreateSubBuffer(address(), flags, CL_BUFFER_CREATE_TYPE_REGION, pinfo.address(), pres.address()); if ((res = MemoryAccess.getInt(pres)) != 0) throw new CLRuntimeException(res); diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLBufferInfo.java b/src/notzed.zcl/classes/au/notzed/zcl/CLBufferInfo.java index 2fa4bf8..4a2241b 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLBufferInfo.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLBufferInfo.java @@ -16,7 +16,8 @@ */ package au.notzed.zcl; -import api.Allocator; +import api.Frame; +import jdk.incubator.foreign.CLinker; import jdk.incubator.foreign.MemoryAccess; import jdk.incubator.foreign.MemorySegment; @@ -38,8 +39,8 @@ public abstract class CLBufferInfo { this.size = size; } - MemorySegment toNative(Allocator frame) { - MemorySegment addr = frame.allocs(2*8); // FIXME: size_t + MemorySegment toNative(Frame frame) { + MemorySegment addr = frame.allocate(CLinker.C_LONG, 2); // FIXME: size_t MemoryAccess.setLong(addr, origin); MemoryAccess.setLongAtIndex(addr, 1, size); @@ -48,5 +49,5 @@ public abstract class CLBufferInfo { } } - abstract MemorySegment toNative(Allocator frame); + abstract MemorySegment toNative(Frame frame); } diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLCommandQueue.java b/src/notzed.zcl/classes/au/notzed/zcl/CLCommandQueue.java index a6ea924..592a145 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLCommandQueue.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLCommandQueue.java @@ -22,6 +22,7 @@ import jdk.incubator.foreign.*; import api.Native; import api.Memory; import api.Callback; +import api.Frame; import java.util.ArrayList; @@ -97,7 +98,7 @@ public class CLCommandQueue extends CLObject { */ public void flush() throws CLException { try { - int res = clFlush(addr()); + int res = clFlush(address()); if (res != 0) throw new CLException(res); } catch (CLException | RuntimeException | Error t) { @@ -114,7 +115,7 @@ public class CLCommandQueue extends CLObject { */ public void finish() throws CLException { try { - int res = clFinish(addr()); + int res = clFinish(address()); if (res != 0) throw new CLException(res); } catch (CLException | RuntimeException | Error t) { @@ -144,11 +145,11 @@ public class CLCommandQueue extends CLObject { if (size > buffer.byteSize()) throw new BufferOverflowException(); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { EventInfo info = new EventInfo(frame, wait, event); int res; - res = clEnqueueReadBuffer(addr(), mem.addr(), blocking ? 1 : 0, + res = clEnqueueReadBuffer(address(), mem.address(), blocking ? 1 : 0, mem_offset, size, buffer.address(), info.nwait, info.wait, info.event); @@ -323,15 +324,15 @@ public class CLCommandQueue extends CLObject { buffer_row_pitch, buffer_slice_pitch, host_row_pitch, host_slice_pitch, buffer.byteSize()); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { EventInfo info = new EventInfo(frame, wait, event); - MemorySegment cbuffer_origin = toLongV(frame, buffer_origin); - MemorySegment chost_origin = toLongV(frame, host_origin); - MemorySegment cregion = toLongV(frame, region); + MemorySegment cbuffer_origin = frame.copy(buffer_origin); + MemorySegment chost_origin = frame.copy(host_origin); + MemorySegment cregion = frame.copy(region); int res; res = clEnqueueReadBufferRect( - addr(), mem.addr(), blocking ? 1 : 0, + address(), mem.address(), blocking ? 1 : 0, cbuffer_origin.address(), chost_origin.address(), cregion.address(), buffer_row_pitch, buffer_slice_pitch, @@ -435,11 +436,11 @@ public class CLCommandQueue extends CLObject { if (size > buffer.byteSize()) throw new BufferUnderflowException(); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { EventInfo info = new EventInfo(frame, wait, event); int res; - res = clEnqueueWriteBuffer(addr(), mem.addr(), blocking ? 1 : 0, + res = clEnqueueWriteBuffer(address(), mem.address(), blocking ? 1 : 0, mem_offset, size, buffer.address(), info.nwait, info.wait, info.event); @@ -569,15 +570,15 @@ public class CLCommandQueue extends CLObject { buffer_row_pitch, buffer_slice_pitch, host_row_pitch, host_slice_pitch, buffer.byteSize()); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { EventInfo info = new EventInfo(frame, wait, event); - MemorySegment cbuffer_origin = toLongV(frame, buffer_origin); - MemorySegment chost_origin = toLongV(frame, host_origin); - MemorySegment cregion = toLongV(frame, region); + MemorySegment cbuffer_origin = frame.copy(buffer_origin); + MemorySegment chost_origin = frame.copy(host_origin); + MemorySegment cregion = frame.copy(region); int res; res = clEnqueueWriteBufferRect( - addr(), mem.addr(), blocking ? 1 : 0, + address(), mem.address(), blocking ? 1 : 0, cbuffer_origin.address(), chost_origin.address(), cregion.address(), buffer_row_pitch, buffer_slice_pitch, @@ -668,12 +669,12 @@ public class CLCommandQueue extends CLObject { CLEventList wait, CLEventList event) throws CLException { requireAPIVersion(CLPlatform.VERSION_1_2); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { EventInfo info = new EventInfo(frame, wait, event); int res; res = clEnqueueFillBuffer( - addr(), buffer.addr(), + address(), buffer.address(), pattern.address(), pattern.byteSize(), offset * pattern.byteSize(), size * pattern.byteSize(), @@ -856,12 +857,12 @@ public class CLCommandQueue extends CLObject { CLEventList wait, CLEventList event) throws CLException { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { EventInfo info = new EventInfo(frame, wait, event); int res; res = clEnqueueCopyBuffer( - addr(), srcmem.addr(), dstmem.addr(), + address(), srcmem.address(), dstmem.address(), srcoffset, dstoffset, size, info.nwait, info.wait, info.event); @@ -927,14 +928,14 @@ public class CLCommandQueue extends CLObject { if (buffer.byteSize() < slice * region[2]) throw new BufferOverflowException(); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { EventInfo info = new EventInfo(frame, wait, event); - MemorySegment corigin = toLongV(frame, origin); - MemorySegment cregion = toLongV(frame, region); + MemorySegment corigin = frame.copy(origin); + MemorySegment cregion = frame.copy(region); int res; res = (int)op.invokeExact( - addr(), image.addr(), blocking ? 1 : 0, + address(), image.address(), blocking ? 1 : 0, corigin.address(), cregion.address(), row_pitch, slice_pitch, buffer.address(), @@ -1352,20 +1353,20 @@ public class CLCommandQueue extends CLObject { long size, CLEventList wait, CLEventList event) throws CLException { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { EventInfo info = new EventInfo(frame, wait, event); - MemorySegment pres = MemorySegment.allocateNative(CLinker.C_INT, frame); + MemorySegment pres = frame.allocateInt(); MemoryAddress cmap; int res; - cmap = clEnqueueMapBuffer(addr(), buffer.addr(), blocking ? 1 : 0, + cmap = clEnqueueMapBuffer(address(), buffer.address(), blocking ? 1 : 0, flags, offset, size, info.nwait, info.wait, info.event, pres.address()); if ((res = MemoryAccess.getInt(pres)) != 0) throw new CLException(res); - return Memory.ofNative(cmap, size).asByteBuffer(); + return cmap.asSegment(size, ResourceScope.globalScope()).asByteBuffer(); } catch (CLException | RuntimeException | Error t) { throw t; } catch (Throwable t) { @@ -1388,19 +1389,19 @@ public class CLCommandQueue extends CLObject { throw new IllegalArgumentException("image_row_pitch must contain 1 element"); // image_slice_pitch may be null for 2D images, checking this isn't worth it - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { EventInfo info = new EventInfo(frame, wait, event); - MemorySegment corigin = toLongV(frame, origin); - MemorySegment cregion = toLongV(frame, region); - MemorySegment cstride = MemorySegment.allocateNative(CLinker.C_LONG, frame); - MemorySegment cslice = MemorySegment.allocateNative(CLinker.C_LONG, frame); - MemorySegment pres = MemorySegment.allocateNative(CLinker.C_LONG, frame); + MemorySegment corigin = frame.copy(origin); + MemorySegment cregion = frame.copy(region); + MemorySegment cstride = frame.allocateLong(); + MemorySegment cslice = frame.allocateLong(); + MemorySegment pres = frame.allocateLong(); MemoryAddress cmap; long stride, slice; int res; cmap = clEnqueueMapImage( - addr(), image.addr(), blocking ? 1 : 0, + address(), image.address(), blocking ? 1 : 0, flags, corigin.address(), cregion.address(), cstride.address(), cslice.address(), @@ -1421,7 +1422,7 @@ public class CLCommandQueue extends CLObject { ? stride * region[1] + region[0] // 2D : slice * region[2] + stride * region[1] + region[0]; // 3D - return Memory.ofNative(cmap, size).asByteBuffer(); + return cmap.asSegment(size, ResourceScope.globalScope()).asByteBuffer(); } catch (CLException | RuntimeException | Error t) { throw t; } catch (Throwable t) { @@ -1434,12 +1435,12 @@ public class CLCommandQueue extends CLObject { CLEventList event) throws CLException { MemoryAddress cmap = MemorySegment.ofByteBuffer(mapped).address(); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { EventInfo info = new EventInfo(frame, wait, event); int res; res = clEnqueueUnmapMemObject( - addr(), mem.addr(), + address(), mem.address(), cmap, info.nwait, info.wait, info.event); @@ -1467,11 +1468,11 @@ public class CLCommandQueue extends CLObject { CLEventList waiters, CLEventList events) throws CLException; - private MemorySegment wsToLongV(ResourceScope frame, long[] ws, int dim) { + private MemorySegment wsToLongV(Frame frame, long[] ws, int dim) { if (ws != null) { if (ws.length < dim) throw new IllegalArgumentException(); - return toLongV(frame, ws); + return frame.copy(ws); } else { // or null?? return MemorySegment.globalNativeSegment(); @@ -1499,14 +1500,14 @@ public class CLCommandQueue extends CLObject { long[] local_size, CLEventList waiters, CLEventList events) throws CLException { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { MemorySegment gwo = wsToLongV(frame, global_offset, work_dim); MemorySegment gws = wsToLongV(frame, global_size, work_dim); MemorySegment lws = wsToLongV(frame, local_size, work_dim); EventInfo info = new EventInfo(frame, waiters, events); int res; - res = clEnqueueNDRangeKernel(addr(), kernel.addr(), + res = clEnqueueNDRangeKernel(address(), kernel.address(), work_dim, gwo.address(), gws.address(), lws.address(), info.nwait, info.wait, info.event); if (res != 0) @@ -1531,11 +1532,11 @@ public class CLCommandQueue extends CLObject { public void enqueueTask(CLKernel kernel, CLEventList waiters, CLEventList events) throws CLException { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { EventInfo info = new EventInfo(frame, waiters, events); int res; - res = clEnqueueTask(addr(), kernel.addr(), + res = clEnqueueTask(address(), kernel.address(), info.nwait, info.wait, info.event); if (res != 0) throw new CLException(res); @@ -1565,10 +1566,10 @@ public class CLCommandQueue extends CLObject { Object... args) throws CLException { // This basically just passes the memory objects to opencl, the rest are handled by the lambda. // This means args and mem_list are the same value - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { EventInfo info = new EventInfo(frame, waiters, events); - MemorySegment memstage = Native.allocAddrV(frame, args.length); - MemorySegment memptrs = Native.allocAddrV(frame, args.length); + MemorySegment memstage = frame.allocate(CLinker.C_POINTER, args.length); + MemorySegment memptrs = frame.allocate(CLinker.C_POINTER, args.length); int nmem = 0; int res; @@ -1576,7 +1577,7 @@ public class CLCommandQueue extends CLObject { for (Object a: args) { if (a instanceof CLMemory) { - MemoryAccess.setAddressAtIndex(memstage, nmem, ((CLMemory)a).addr()); + MemoryAccess.setAddressAtIndex(memstage, nmem, ((CLMemory)a).address()); MemoryAccess.setAddressAtIndex(memptrs, nmem, memstage.asSlice(nmem * CLinker.C_POINTER.byteSize())); nmem++; } @@ -1592,7 +1593,7 @@ public class CLCommandQueue extends CLObject { if (args[i] instanceof CLMemory) tmem++; } - + MemorySegment seg = memargs.asSegment(tmem * CLinker.C_POINTER.byteSize(), ResourceScope.globalScope()); for (int i = 0; i < args.length; i++) { @@ -1600,7 +1601,7 @@ public class CLCommandQueue extends CLObject { MemoryAddress mem = MemoryAccess.getAddressAtIndex(seg, xmem); long size = ((CLMemory)args[i]).getSize(); - save[i] = Memory.ofNative(mem, size).asByteBuffer().order(ByteOrder.nativeOrder()); + save[i] = mem.asSegment(size, ResourceScope.globalScope()).asByteBuffer().order(ByteOrder.nativeOrder()); xmem++; } } @@ -1609,7 +1610,7 @@ public class CLCommandQueue extends CLObject { }), (p) -> new Callback<>(p, kernel)); - res = clEnqueueNativeKernel(addr(), call.addr(), memstage.address(), nmem * CLinker.C_POINTER.byteSize(), nmem, memstage.address(), memptrs.address(), info.nwait, info.wait, info.event); + res = clEnqueueNativeKernel(address(), call.address(), memstage.address(), nmem * CLinker.C_POINTER.byteSize(), nmem, memstage.address(), memptrs.address(), info.nwait, info.wait, info.event); if (res != 0) throw new CLException(res); @@ -1636,13 +1637,13 @@ public class CLCommandQueue extends CLObject { public void enqueueMarkerWithWaitList( CLEventList waiters, CLEventList events) throws CLException { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { EventInfo info = new EventInfo(frame, waiters, events); if (haveAPIVersion(CLPlatform.VERSION_1_2)) { - clEnqueueMarkerWithWaitList(addr(), info.nwait, info.wait, info.event); + clEnqueueMarkerWithWaitList(address(), info.nwait, info.wait, info.event); } else { - clEnqueueWaitForEvents(addr(), info.nwait, info.wait); - clEnqueueMarker(addr(), info.event); + clEnqueueWaitForEvents(address(), info.nwait, info.wait); + clEnqueueMarker(address(), info.event); } info.post(events); } catch (CLException | RuntimeException | Error t) { @@ -1666,14 +1667,14 @@ public class CLCommandQueue extends CLObject { public void enqueueBarrierWithWaitList( CLEventList waiters, CLEventList events) throws CLException { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { EventInfo info = new EventInfo(frame, waiters, events); if (haveAPIVersion(CLPlatform.VERSION_1_2)) { - clEnqueueBarrierWithWaitList(addr(), info.nwait, info.wait, info.event); + clEnqueueBarrierWithWaitList(address(), info.nwait, info.wait, info.event); } else { - clEnqueueWaitForEvents(addr(), info.nwait, info.wait); - clEnqueueBarrier(addr()); - clEnqueueMarker(addr(), info.event); + clEnqueueWaitForEvents(address(), info.nwait, info.wait); + clEnqueueBarrier(address()); + clEnqueueMarker(address(), info.event); } info.post(events); } catch (CLException | RuntimeException | Error t) { @@ -1915,12 +1916,12 @@ public class CLCommandQueue extends CLObject { CLEventList wait, CLEventList event) { GLext gl = getGLext(); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { EventInfo info = new EventInfo(frame, wait, event); - MemorySegment cmem_objects = Native.toAddrV(frame, mem_objects); + MemorySegment cmem_objects = frame.copy(mem_objects); int res; - res = gl.clEnqueueAcquireGLObjects(addr(), mem_objects.length, cmem_objects.address(), + res = gl.clEnqueueAcquireGLObjects(address(), mem_objects.length, cmem_objects.address(), info.nwait, info.wait, info.event); if (res != 0) throw new CLRuntimeException(res); @@ -1938,12 +1939,12 @@ public class CLCommandQueue extends CLObject { CLEventList wait, CLEventList event) { GLext gl = getGLext(); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { EventInfo info = new EventInfo(frame, wait, event); - MemorySegment cmem_objects = Native.toAddrV(frame, mem_objects); + MemorySegment cmem_objects = frame.copy(mem_objects); int res; - res = gl.clEnqueueReleaseGLObjects(addr(), mem_objects.length, cmem_objects.address(), + res = gl.clEnqueueReleaseGLObjects(address(), mem_objects.length, cmem_objects.address(), info.nwait, info.wait, info.event); if (res != 0) throw new CLRuntimeException(res); diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLContext.java b/src/notzed.zcl/classes/au/notzed/zcl/CLContext.java index 3035dcb..a23d145 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLContext.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLContext.java @@ -20,9 +20,9 @@ import static au.notzed.zcl.CL.*; import static au.notzed.zcl.CLLib.*; import jdk.incubator.foreign.*; import api.Memory; -import api.Allocator; import api.Native; import api.Callback; +import api.Frame; import java.lang.invoke.MethodHandle; import java.nio.ByteBuffer; import java.nio.ByteOrder; @@ -89,7 +89,7 @@ public class CLContext extends CLObject { * @return new property */ public static CLContextProperty PLATFORM(CLPlatform platform) { - return new CLContextProperty.TagValue(CL.CL_CONTEXT_PLATFORM, Memory.toLong(platform.addr())); + return new CLContextProperty.TagValue(CL.CL_CONTEXT_PLATFORM, platform.address().toRawLongValue()); } /** @@ -113,18 +113,18 @@ public class CLContext extends CLObject { * @throws CLRuntimeException */ public static CLContext createContext(CLContextProperty[] properties, CLDevice[] devices, CLContextNotify notify) throws CLRuntimeException { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { MemorySegment pprops = CLProperty.toNative(frame, properties); - MemorySegment pdevs = MemorySegment.allocateNative(devices.length * CLinker.C_POINTER.byteSize(), CLinker.C_POINTER.byteAlignment(), frame); - MemorySegment pres = MemorySegment.allocateNative(CLinker.C_INT, frame); + MemorySegment pdevs = frame.allocateAddress(devices.length); + MemorySegment pres = frame.allocateInt(); Callback call = CLContextNotify.call(notify); MemoryAddress cl; int res; for (int i = 0; i < devices.length; i++) - MemoryAccess.setAddressAtIndex(pdevs, i, devices[i].addr()); + MemoryAccess.setAddressAtIndex(pdevs, i, devices[i].address()); - cl = clCreateContext(pprops.address(), devices.length, pdevs.address(), call.addr(), MemoryAddress.NULL, pres.address()); + cl = clCreateContext(pprops.address(), devices.length, pdevs.address(), call.address(), MemoryAddress.NULL, pres.address()); res = MemoryAccess.getInt(pres); if (res != 0) @@ -148,10 +148,10 @@ public class CLContext extends CLObject { * Directly lookup platform for a context where we have no other context to find it */ static CLPlatform findPlatform(MemoryAddress ccl) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - MemorySegment devices = getInfoAny(ccl, CL_CONTEXT_DEVICES, clGetContextInfo, a); + try (Frame frame = Memory.createFrame()) { + MemorySegment devices = getInfoAny(ccl, CL_CONTEXT_DEVICES, clGetContextInfo, frame); MemoryAddress dev0 = MemoryAccess.getAddress(devices); - MemoryAddress plat = MemoryAccess.getAddress(getInfo(dev0, CL_DEVICE_PLATFORM, clGetDeviceInfo, a, 8)); + MemoryAddress plat = MemoryAccess.getAddress(getInfo(dev0, CL_DEVICE_PLATFORM, clGetDeviceInfo, frame, 8)); return Native.resolve(plat, CLPlatform::new); } catch (RuntimeException | Error t) { @@ -172,14 +172,14 @@ public class CLContext extends CLObject { * @throws CLRuntimeException */ public static CLContext createContextFromType(CLContextProperty[] properties, long device_type, CLContextNotify notify) throws CLRuntimeException { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { MemorySegment pprops = CLProperty.toNative(frame, properties); - MemorySegment pres = MemorySegment.allocateNative(CLinker.C_INT, frame); + MemorySegment pres = frame.allocateInt(); MemoryAddress cl; Callback call = CLContextNotify.call(notify); int res; - cl = clCreateContextFromType(pprops.address(), device_type, call.addr(), MemoryAddress.NULL, pres.address()); + cl = clCreateContextFromType(pprops.address(), device_type, call.address(), MemoryAddress.NULL, pres.address()); res = MemoryAccess.getInt(pres); if (res != 0) @@ -204,12 +204,12 @@ public class CLContext extends CLObject { * deprecated as of OpenCL 2.0 */ public CLCommandQueue createCommandQueue(CLDevice dev, long properties) throws CLRuntimeException { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment pres = MemorySegment.allocateNative(CLinker.C_INT, frame); + try (Frame frame = Memory.createFrame()) { + MemorySegment pres = frame.allocateInt(); MemoryAddress q; int res; - q = clCreateCommandQueue(addr(), dev.addr(), properties, pres.address()); + q = clCreateCommandQueue(address(), dev.address(), properties, pres.address()); res = MemoryAccess.getInt(pres); if (res != 0) @@ -236,13 +236,13 @@ public class CLContext extends CLObject { // Fallback if opencl2 not supported? requireAPIVersion(CLPlatform.VERSION_2_0); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { MemorySegment pprops = CLProperty.toNative(frame, properties); - MemorySegment pres = MemorySegment.allocateNative(CLinker.C_INT, frame); + MemorySegment pres = frame.allocateInt(); MemoryAddress q; int res; - q = clCreateCommandQueueWithProperties(addr(), dev.addr(), pprops.address(), pres.address()); + q = clCreateCommandQueueWithProperties(address(), dev.address(), pprops.address(), pres.address()); res = MemoryAccess.getInt(pres); if (res != 0) @@ -269,7 +269,7 @@ public class CLContext extends CLObject { try { int res; - res = clSetDefaultDeviceCommandQueue(addr(), dev.addr(), q.addr()); + res = clSetDefaultDeviceCommandQueue(address(), dev.address(), q.address()); if (res != 0) throw new CLException(res); } catch (CLException | RuntimeException | Error t) { @@ -308,12 +308,12 @@ public class CLContext extends CLObject { if (hostseg != null && hostseg.byteSize() < size) throw new CLRuntimeException(CL_INVALID_HOST_PTR); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment pres = MemorySegment.allocateNative(CLinker.C_INT, frame); + try (Frame frame = Memory.createFrame()) { + MemorySegment pres = frame.allocateInt(); MemoryAddress pbuffer; int res; - pbuffer = clCreateBuffer(addr(), flags, size, addr(hostseg), pres.address()); + pbuffer = clCreateBuffer(address(), flags, size, addr(hostseg), pres.address()); if ((res = MemoryAccess.getInt(pres)) != 0) throw new CLRuntimeException(res); @@ -403,9 +403,8 @@ public class CLContext extends CLObject { throw new NullPointerException(); // Must copy. - try (ResourceScope scope = ResourceScope.newConfinedScope()) { - MemorySegment mem = MemorySegment.allocateNative(hostp.length, 16, scope); - mem.asByteBuffer().order(ByteOrder.nativeOrder()).put(hostp); + try (Frame frame = Memory.createFrame()) { + MemorySegment mem = frame.copy(hostp); return createBuffer(flags, hostp.length, mem); } } @@ -420,9 +419,8 @@ public class CLContext extends CLObject { throw new NullPointerException(); // Must copy. - try (ResourceScope scope = ResourceScope.newConfinedScope()) { - MemorySegment mem = MemorySegment.allocateNative(hostp.length, 16, scope); - mem.asByteBuffer().order(ByteOrder.nativeOrder()).asFloatBuffer().put(hostp); + try (Frame frame = Memory.createFrame()) { + MemorySegment mem = frame.copy(hostp); return createBuffer(flags, hostp.length * 4, mem); } } @@ -448,9 +446,9 @@ public class CLContext extends CLObject { * @throws UnsupportedOperationException */ public CLImage createImage(long flags, CLImageFormat fmt, CLImageDesc desc, MemorySegment hostseg) throws CLRuntimeException, UnsupportedOperationException { - try (Allocator frame = Memory.stack()) { + try (Frame frame = Memory.createFrame()) { MemorySegment cfmt = CLImageFormat.toNative(frame, fmt); - MemorySegment cres = frame.allocs(8); + MemorySegment cres = frame.allocateInt(); MemoryAddress ci; int res; @@ -458,14 +456,14 @@ public class CLContext extends CLObject { if (haveAPIVersion(CLPlatform.VERSION_1_2)) { MemorySegment cdesc = CLImageDesc.toNative(frame, desc); - ci = clCreateImage(addr(), flags, cfmt.address(), cdesc.address(), addr(hostseg), cres.address()); + ci = clCreateImage(address(), flags, cfmt.address(), cdesc.address(), addr(hostseg), cres.address()); } else { switch (desc.imageType) { case CL_MEM_OBJECT_IMAGE2D: - ci = clCreateImage2D(addr(), flags, cfmt.address(), desc.imageWidth, desc.imageHeight, desc.imageRowPitch, addr(hostseg), cres.address()); + ci = clCreateImage2D(address(), flags, cfmt.address(), desc.imageWidth, desc.imageHeight, desc.imageRowPitch, addr(hostseg), cres.address()); break; case CL_MEM_OBJECT_IMAGE3D: - ci = clCreateImage3D(addr(), flags, cfmt.address(), desc.imageWidth, desc.imageHeight, desc.imageDepth, + ci = clCreateImage3D(address(), flags, cfmt.address(), desc.imageWidth, desc.imageHeight, desc.imageDepth, desc.imageRowPitch, desc.imageSlicePitch, addr(hostseg), cres.address()); break; default: @@ -513,9 +511,8 @@ public class CLContext extends CLObject { throw new NullPointerException(); // Must copy. - try (ResourceScope scope = ResourceScope.newConfinedScope()) { - MemorySegment mem = MemorySegment.allocateNative(hostp.length, 16, scope); - mem.copyFrom(MemorySegment.ofArray(hostp)); + try (Frame frame = Memory.createFrame()) { + MemorySegment mem = frame.copy(hostp); return createImage(flags, fmt, desc, mem); } } @@ -530,9 +527,8 @@ public class CLContext extends CLObject { throw new NullPointerException(); // Must copy. - try (ResourceScope scope = ResourceScope.newConfinedScope()) { - MemorySegment mem = MemorySegment.allocateNative(hostp.length * 4, 16, scope); - mem.copyFrom(MemorySegment.ofArray(hostp)); + try (Frame frame = Memory.createFrame()) { + MemorySegment mem = frame.copy(hostp); return createImage(flags, fmt, desc, mem); } } @@ -552,13 +548,13 @@ public class CLContext extends CLObject { public CLPipe createPipe(long flags, int packetSize, int maxPackets, CLPipeProperty[] properties) throws CLRuntimeException, UnsupportedOperationException { requireAPIVersion(CLPlatform.VERSION_2_0); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { MemorySegment pprops = CLProperty.toNative(frame, properties); - MemorySegment pres = MemorySegment.allocateNative(CLinker.C_INT, frame); + MemorySegment pres = frame.allocateInt(); int res; MemoryAddress cp; - cp = clCreatePipe(addr(), flags, packetSize, maxPackets, pprops.address(), pres.address()); + cp = clCreatePipe(address(), flags, packetSize, maxPackets, pprops.address(), pres.address()); if ((res = MemoryAccess.getInt(pres)) != 0) throw new CLRuntimeException(res); @@ -579,20 +575,20 @@ public class CLContext extends CLObject { * @throws CLRuntimeException */ public CLImageFormat[] getSupportedImageFormats(long flags, int type) throws CLRuntimeException { - try (Allocator frame = Memory.stack()) { - MemorySegment cnum = frame.allocs(8); + try (Frame frame = Memory.createFrame()) { + MemorySegment cnum = frame.allocateInt(); MemorySegment list; int num; int res; - res = clGetSupportedImageFormats(addr(), flags, type, 0, MemoryAddress.NULL, cnum.address()); + res = clGetSupportedImageFormats(address(), flags, type, 0, MemoryAddress.NULL, cnum.address()); if (res != 0) throw new CLRuntimeException(res); num = MemoryAccess.getInt(cnum); - list = frame.allocs(num * 8); + list = frame.allocateAddress(8); - res = clGetSupportedImageFormats(addr(), flags, type, num, list.address(), cnum.address()); + res = clGetSupportedImageFormats(address(), flags, type, num, list.address(), cnum.address()); CLImageFormat[] out = new CLImageFormat[num]; for (int i = 0; i < out.length; i++) @@ -648,12 +644,12 @@ public class CLContext extends CLObject { */ //@Deprecated public CLSampler createSampler(boolean norm, int addr_mode, int filter_mode) throws CLRuntimeException { - try (Allocator frame = Memory.stack()) { - MemorySegment cres = frame.allocs(8); + try (Frame frame = Memory.createFrame()) { + MemorySegment cres = frame.allocateInt(); int res; MemoryAddress cs; - cs = clCreateSampler(addr(), norm ? 1 : 0, addr_mode, filter_mode, cres.address()); + cs = clCreateSampler(address(), norm ? 1 : 0, addr_mode, filter_mode, cres.address()); if ((res = MemoryAccess.getInt(cres)) != 0) throw new CLRuntimeException(res); @@ -677,13 +673,13 @@ public class CLContext extends CLObject { // Fallback if opencl2 not supported? requireAPIVersion(CLPlatform.VERSION_2_0); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { + try (Frame frame = Memory.createFrame()) { MemorySegment pprops = CLProperty.toNative(frame, properties); - MemorySegment pres = MemorySegment.allocateNative(CLinker.C_INT, frame); + MemorySegment pres = frame.allocateInt(); int res; MemoryAddress cs; - cs = clCreateSamplerWithProperties(addr(), pprops.address(), pres.address()); + cs = clCreateSamplerWithProperties(address(), pprops.address(), pres.address()); if ((res = MemoryAccess.getInt(pres)) != 0) throw new CLRuntimeException(res); @@ -703,15 +699,20 @@ public class CLContext extends CLObject { return len; } - static void copy(MemorySegment addr, byte[][] list) { + static void copxy(MemorySegment addr, byte[][] list) { for (int i = 0, k = 0; i < list.length; i++) { addr.asSlice(k).copyFrom(MemorySegment.ofArray(list[i])); k += list[i].length; } } - static void copy(MemorySegment addr, byte[] row) { - addr.copyFrom(MemorySegment.ofArray(row)); + static MemorySegment concat(Frame frame, byte[]... list) { + MemorySegment buffer = frame.allocate(length(list)); + for (int i = 0, k = 0; i < list.length; i++) { + buffer.asSlice(k).copyFrom(MemorySegment.ofArray(list[i])); + k += list[i].length; + } + return buffer; } /** @@ -725,19 +726,15 @@ public class CLContext extends CLObject { public CLProgram createProgramWithSource(byte[]... strings) throws CLException { long size = length(strings); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment buffer = MemorySegment.allocateNative(size, frame); - MemorySegment cstring = MemorySegment.allocateNative(CLinker.C_POINTER, frame); - MemorySegment clength = MemorySegment.allocateNative(CLinker.C_LONG, frame); - MemorySegment pres = MemorySegment.allocateNative(CLinker.C_INT, frame); + try (Frame frame = Memory.createFrame()) { + MemorySegment buffer = concat(frame, strings); + MemorySegment cstring = frame.copy(buffer.address()); + MemorySegment clength = frame.copy(size); + MemorySegment pres = frame.allocateInt(); MemoryAddress cp; int res; - copy(buffer, strings); - MemoryAccess.setAddress(cstring, buffer.address()); - MemoryAccess.setLong(clength, size); - - cp = clCreateProgramWithSource(addr(), 1, cstring.address(), clength.address(), pres.address()); + cp = clCreateProgramWithSource(address(), 1, cstring.address(), clength.address(), pres.address()); if ((res = MemoryAccess.getInt(pres)) != 0) throw new CLException(res); @@ -781,15 +778,13 @@ public class CLContext extends CLObject { public CLProgram createProgramWithIL(byte[] il) throws CLException { requireAPIVersion(CLPlatform.VERSION_2_1); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment buffer = MemorySegment.allocateNative(il.length, 1, frame); - MemorySegment pres = MemorySegment.allocateNative(CLinker.C_INT, frame); + try (Frame frame = Memory.createFrame()) { + MemorySegment buffer = frame.copy(il); + MemorySegment pres = frame.allocateInt(); MemoryAddress cp; int res; - copy(buffer, il); - - cp = clCreateProgramWithIL(addr(), buffer.address(), il.length, pres.address()); + cp = clCreateProgramWithIL(address(), buffer.address(), il.length, pres.address()); if ((res = MemoryAccess.getInt(pres)) != 0) throw new CLException(res); @@ -813,27 +808,27 @@ public class CLContext extends CLObject { */ public CLProgram createProgramWithBinary(CLDevice[] devices, byte[][] binaries, int[] status) throws CLException { long size = length(binaries); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment buffer = MemorySegment.allocateNative(size, 1, frame); - MemorySegment cdevs = toAddrV(frame, devices); - MemorySegment barray = Native.allocAddrV(frame, binaries.length); - MemorySegment larray = Native.allocV(frame, CLinker.C_LONG, binaries.length); - MemorySegment cstatus = status != null ? MemorySegment.allocateNative(4 * status.length, 4, frame) : MemorySegment.globalNativeSegment(); - MemorySegment pres = MemorySegment.allocateNative(CLinker.C_INT, frame); + try (Frame frame = Memory.createFrame()) { + MemorySegment buffer = frame.allocate(size); + MemorySegment cdevs = frame.copy(devices); + MemorySegment barray = frame.allocateAddress(binaries.length); + MemorySegment larray = frame.allocateLong(binaries.length); + MemorySegment cstatus = status != null ? frame.allocateInt(status.length) : MemorySegment.globalNativeSegment(); + MemorySegment pres = frame.allocateInt(); MemoryAddress cp; int res; for (int i = 0, o = 0; i < binaries.length; i++) { MemorySegment addr = buffer.asSlice(o); - copy(addr, binaries[i]); + addr.copyFrom(MemorySegment.ofArray(binaries[i])); MemoryAccess.setAddressAtIndex(barray, i, addr); MemoryAccess.setLongAtIndex(larray, i, binaries[i].length); o += binaries[i].length; } - cp = clCreateProgramWithBinary(addr(), devices.length, cdevs.address(), larray.address(), barray.address(), cstatus.address(), pres.address()); + cp = clCreateProgramWithBinary(address(), devices.length, cdevs.address(), larray.address(), barray.address(), cstatus.address(), pres.address()); if ((res = MemoryAccess.getInt(pres)) != 0) throw new CLException(res); @@ -863,14 +858,14 @@ public class CLContext extends CLObject { public CLProgram createProgramWithBuiltInKernels(CLDevice[] devices, String names) throws CLException, UnsupportedOperationException { requireAPIVersion(CLPlatform.VERSION_1_2); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment cdevs = toAddrV(frame, devices); - MemorySegment cnames = CLinker.toCString(names, frame); - MemorySegment pres = MemorySegment.allocateNative(CLinker.C_INT, frame); + try (Frame frame = Memory.createFrame()) { + MemorySegment cdevs = frame.copy(devices); + MemorySegment cnames = frame.copy(names); + MemorySegment pres = frame.allocateInt(); MemoryAddress cp; int res; - cp = clCreateProgramWithBuiltInKernels(addr(), devices.length, cdevs.address(), cnames.address(), pres.address()); + cp = clCreateProgramWithBuiltInKernels(address(), devices.length, cdevs.address(), cnames.address(), pres.address()); if ((res = MemoryAccess.getInt(pres)) != 0) throw new CLException(res); @@ -898,16 +893,16 @@ public class CLContext extends CLObject { public CLProgram linkProgram(CLDevice[] devices, String options, CLProgram[] programs, CLNotify notify) throws CLException, UnsupportedOperationException { requireAPIVersion(CLPlatform.VERSION_1_2); - try (ResourceScope frame = ResourceScope.newConfinedScope(); + try (Frame frame = Memory.createFrame(); Callback> cnotify = CLNotify.call(notify, (x) -> new CLProgram(x, getObjectPlatform()))) { - MemorySegment cdevs = toAddrV(frame, devices); - MemorySegment coptions = options != null ? CLinker.toCString(options, frame) : MemorySegment.globalNativeSegment(); - MemorySegment cprogs = toAddrV(frame, programs); - MemorySegment pres = MemorySegment.allocateNative(CLinker.C_INT, frame); + MemorySegment cdevs = frame.copy(devices); + MemorySegment coptions = frame.copy(options); + MemorySegment cprogs = frame.copy(programs); + MemorySegment pres = frame.allocateInt(); MemoryAddress cp; int res; - cp = clLinkProgram(addr(), devices.length, cdevs.address(), coptions.address(), programs.length, cprogs.address(), cnotify.addr(), MemoryAddress.NULL, pres.address()); + cp = clLinkProgram(address(), devices.length, cdevs.address(), coptions.address(), programs.length, cprogs.address(), cnotify.address(), MemoryAddress.NULL, pres.address()); if ((res = MemoryAccess.getInt(pres)) != 0) throw new CLException(res); @@ -930,12 +925,12 @@ public class CLContext extends CLObject { public CLEvent createUserEvent() throws CLException { requireAPIVersion(CLPlatform.VERSION_1_1); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment pres = MemorySegment.allocateNative(CLinker.C_INT, frame); + try (Frame frame = Memory.createFrame()) { + MemorySegment pres = frame.allocateInt(); MemoryAddress ce; int res; - ce = clCreateUserEvent(addr(), pres.address()); + ce = clCreateUserEvent(address(), pres.address()); if ((res = MemoryAccess.getInt(pres)) != 0) throw new CLException(res); @@ -1011,12 +1006,12 @@ public class CLContext extends CLObject { long flags, int bufobj) { GLext gl = getGLext(); - try (Allocator frame = Memory.stack()) { - MemorySegment cres = frame.allocs(8); + try (Frame frame = Memory.createFrame()) { + MemorySegment cres = frame.allocateInt(); MemoryAddress ce; int res; - ce = gl.clCreateFromGLBuffer(addr(), flags, bufobj, cres.address()); + ce = gl.clCreateFromGLBuffer(address(), flags, bufobj, cres.address()); if ((res = MemoryAccess.getInt(cres)) != 0) throw new CLRuntimeException(res); return Native.resolve(ce, (b) -> new CLBuffer(b, getObjectPlatform())); @@ -1036,13 +1031,13 @@ public class CLContext extends CLObject { int miplevel, int texture) { GLext gl = getGLext(); - try (Allocator frame = Memory.stack()) { - MemorySegment cres = frame.allocs(8); + try (Frame frame = Memory.createFrame()) { + MemorySegment pres = frame.allocateInt(); MemoryAddress ce; int res; - ce = gl.clCreateFromGLTexture(addr(), flags, target, miplevel, texture, cres.address()); - if ((res = MemoryAccess.getInt(cres)) != 0) + ce = gl.clCreateFromGLTexture(address(), flags, target, miplevel, texture, pres.address()); + if ((res = MemoryAccess.getInt(pres)) != 0) throw new CLRuntimeException(res); return Native.resolve(ce, (x) -> new CLImage(x, getObjectPlatform())); } catch (RuntimeException | Error t) { @@ -1059,13 +1054,13 @@ public class CLContext extends CLObject { long flags, int renderbuffer) { GLext gl = getGLext(); - try (Allocator frame = Memory.stack()) { - MemorySegment cres = frame.allocs(8); + try (Frame frame = Memory.createFrame()) { + MemorySegment pres = frame.allocateInt(); MemoryAddress ce; int res; - ce = gl.clCreateFromGLRenderbuffer(addr(), flags, renderbuffer, cres.address()); - if ((res = MemoryAccess.getInt(cres)) != 0) + ce = gl.clCreateFromGLRenderbuffer(address(), flags, renderbuffer, pres.address()); + if ((res = MemoryAccess.getInt(pres)) != 0) throw new CLRuntimeException(res); return Native.resolve(ce, (x) -> new CLImage(x, getObjectPlatform())); } catch (RuntimeException | Error t) { @@ -1102,13 +1097,13 @@ public class CLContext extends CLObject { */ public CLEvent clCreateEventFromGLsyncKHR(MemoryAddress glsync) { GLext gl = getGLext(); - try (Allocator frame = Memory.stack()) { - MemorySegment cret = frame.allocs(8); + try (Frame frame = Memory.createFrame()) { + MemorySegment pres = frame.allocateInt(); MemoryAddress ce; int res; - ce = gl.clCreateEventFromGLsyncKHR(addr(), glsync, cret.address()); - if ((res = MemoryAccess.getInt(cret)) != 0) + ce = gl.clCreateEventFromGLsyncKHR(address(), glsync, pres.address()); + if ((res = MemoryAccess.getInt(pres)) != 0) throw new CLRuntimeException(res); return Native.resolve(ce, CLEvent::new); diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLContextNotify.java b/src/notzed.zcl/classes/au/notzed/zcl/CLContextNotify.java index c005742..a18255d 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLContextNotify.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLContextNotify.java @@ -20,7 +20,8 @@ import java.nio.ByteBuffer; import jdk.incubator.foreign.MemorySegment; import api.Native; import api.Callback; -import api.Memory; +import jdk.incubator.foreign.CLinker; +import jdk.incubator.foreign.ResourceScope; /** * Callback for CLContext.createContext*() @@ -29,13 +30,13 @@ public interface CLContextNotify { public void notify(String what, ByteBuffer error_info); - @SuppressWarnings("unchecked") + //@SuppressWarnings("unchecked") static Callback call(CLContextNotify notify) { if (notify != null) { return Native.resolve( Call_pBpvJpv_v.stub((cwhat, cinfo, cinfolen, dummy) -> { - MemorySegment seg = Memory.ofNative(cinfo, cinfolen); - notify.notify(Native.toString(cwhat), seg.asByteBuffer()); + MemorySegment seg = cinfo.asSegment(cinfolen, ResourceScope.globalScope()); + notify.notify(CLinker.toJavaString(cwhat), seg.asByteBuffer()); }), (p) -> new Callback<>(p, notify)); } else { diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLEvent.java b/src/notzed.zcl/classes/au/notzed/zcl/CLEvent.java index fd2f283..1351968 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLEvent.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLEvent.java @@ -66,7 +66,7 @@ public class CLEvent extends CLObject { @Override CLPlatform getObjectPlatform() { if (platform == null) - platform = findPlatform(addr()); + platform = findPlatform(address()); return platform; } @@ -75,8 +75,8 @@ public class CLEvent extends CLObject { * This goes via the context - every one must have a context? */ static CLPlatform findPlatform(MemoryAddress cevent) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - MemoryAddress ccl = MemoryAccess.getAddress(getInfo(cevent, CL_EVENT_CONTEXT, clGetEventInfo, a, 8)); + try (Frame frame = Memory.createFrame()) { + MemoryAddress ccl = MemoryAccess.getAddress(getInfo(cevent, CL_EVENT_CONTEXT, clGetEventInfo, frame, 8)); return CLContext.findPlatform(ccl); } catch (RuntimeException | Error t) { @@ -98,7 +98,7 @@ public class CLEvent extends CLObject { public void setUserEventStatus(int status) throws CLRuntimeException { requireAPIVersion(CLPlatform.VERSION_1_1); try { - int res = clSetUserEventStatus(addr(), status); + int res = clSetUserEventStatus(address(), status); if (res != 0) throw new CLRuntimeException(res); @@ -125,7 +125,7 @@ public class CLEvent extends CLObject { Callback callback = CLEventNotify.call(notify); try { - int res = clSetEventCallback(addr(), type, callback.addr(), MemoryAddress.NULL); + int res = clSetEventCallback(address(), type, callback.address(), MemoryAddress.NULL); if (res != 0) throw new CLRuntimeException(res); diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLEventList.java b/src/notzed.zcl/classes/au/notzed/zcl/CLEventList.java index 811ade2..cf4eace 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLEventList.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLEventList.java @@ -20,6 +20,7 @@ import jdk.incubator.foreign.MemoryAddress; import jdk.incubator.foreign.MemorySegment; import api.Native; import static au.notzed.zcl.CLLib.*; +import jdk.incubator.foreign.CLinker; import jdk.incubator.foreign.MemoryAccess; import jdk.incubator.foreign.ResourceScope; @@ -44,7 +45,9 @@ import jdk.incubator.foreign.ResourceScope; */ public final class CLEventList implements AutoCloseable { - // FIXME: pass scope in instead of being autocloseable + /** + * Scope is required for native allocation + */ final ResourceScope scope = ResourceScope.newConfinedScope(); /** * Raw event values. @@ -67,7 +70,7 @@ public final class CLEventList implements AutoCloseable { */ public CLEventList(int capacity) { this.jevents = new CLEvent[capacity]; - this.cevents = MemorySegment.allocateNative(8 * capacity, 8, scope); + this.cevents = MemorySegment.allocateNative(CLinker.C_POINTER.byteSize() * capacity, CLinker.C_POINTER.byteAlignment(), scope); } /** @@ -168,7 +171,7 @@ public final class CLEventList implements AutoCloseable { * @param event */ public void add(CLEvent event) { - MemoryAccess.setAddressAtIndex(cevents, index, event.addr()); + MemoryAccess.setAddressAtIndex(cevents, index, event.address()); jevents[index++] = event; } diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLImage.java b/src/notzed.zcl/classes/au/notzed/zcl/CLImage.java index 6fcf662..bd12b63 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLImage.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLImage.java @@ -93,7 +93,7 @@ public class CLImage extends CLMemory { @Override public String toString() { - return String.format("[%s: %dx%dx%d 0x%x]", getClass().getSimpleName(), getWidth(), getHeight(), getDepth(), Memory.toLong(addr())); + return String.format("[%s: %dx%dx%d 0x%x]", getClass().getSimpleName(), getWidth(), getHeight(), getDepth(), address().toRawLongValue()); } /** diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLImageDesc.java b/src/notzed.zcl/classes/au/notzed/zcl/CLImageDesc.java index 55b02cc..8024c3d 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLImageDesc.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLImageDesc.java @@ -18,16 +18,16 @@ package au.notzed.zcl; import jdk.incubator.foreign.*; import api.Native; -import api.Allocator; +import api.Frame; /** * Holder for cl_image_desc equivalent. *

panama notes

- To maintain compatability with the - * previous api this remains as a simple pojo and marshalling is - * done as required. -

-This also means it doesn't have to deal with allocation and deallocation and so on. + * To maintain compatability with the + * previous api this remains as a simple pojo and marshalling is + * done as required. + *

+ * This also means it doesn't have to deal with allocation and deallocation and so on. */ public class CLImageDesc { @@ -86,8 +86,8 @@ public class CLImageDesc { * This is just hand-rolled for now. I'm not really sure how to approach it * since these are just going to be used temporarily */ - public static MemorySegment toNative(Allocator frame, CLImageDesc d) { - MemorySegment addr = frame.allocs(sizeof); + public static MemorySegment toNative(Frame frame, CLImageDesc d) { + MemorySegment addr = frame.allocate(LAYOUT); MemoryAccess.setInt(addr, d.imageType); MemoryAccess.setLongAtIndex(addr, 1, d.imageWidth); @@ -98,7 +98,7 @@ public class CLImageDesc { MemoryAccess.setLongAtIndex(addr, 6, d.imageSlicePitch); MemoryAccess.setLongAtIndex(addr, 7, d.numMipLevels); MemoryAccess.setLongAtIndex(addr, 8, d.numSamples); - MemoryAccess.setAddressAtIndex(addr, 9, d.memObject.addr()); + MemoryAccess.setAddressAtIndex(addr, 9, d.memObject.address()); return addr; } @@ -107,8 +107,10 @@ public class CLImageDesc { return new CLImageFormat(MemoryAccess.getInt(addr), MemoryAccess.getIntAtIndex(addr, 1)); } - public static final long sizeof = 72; + static final MemoryLayout LAYOUT = Native.parseStruct("[u32(image_type)x32u64(image_width)u64(image_height)u64(image_depth)u64(image_array_size)u64(image_row_pitch)u64(image_slice_pitch)u32(num_mip_levels)u32(num_samples)u64(buffer):${_cl_mem}u64(mem_object):${_cl_mem}]"); - public static MemoryLayout layout() { return Native.parseStruct("[u32(image_type)x32u64(image_width)u64(image_height)u64(image_depth)u64(image_array_size)u64(image_row_pitch)u64(image_slice_pitch)u32(num_mip_levels)u32(num_samples)u64(buffer):${_cl_mem}u64(mem_object):${_cl_mem}]"); } + static MemoryLayout layout() { + return LAYOUT; + } } diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLImageFormat.java b/src/notzed.zcl/classes/au/notzed/zcl/CLImageFormat.java index ea47a6c..9911097 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLImageFormat.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLImageFormat.java @@ -18,7 +18,7 @@ package au.notzed.zcl; import jdk.incubator.foreign.*; import api.Native; -import api.Allocator; +import api.Frame; /** * Holder for cl_image_format equivalent. @@ -96,8 +96,8 @@ public class CLImageFormat { * This is just hand-rolled for now. I'm not really sure how to approach it * since these are just going to be used temporarily */ - static MemorySegment toNative(Allocator frame, CLImageFormat fmt) { - MemorySegment addr = frame.allocs(2 * 4); + static MemorySegment toNative(Frame frame, CLImageFormat fmt) { + MemorySegment addr = frame.allocate(CLinker.C_INT, 2); MemoryAccess.setInt(addr, fmt.channelOrder); MemoryAccess.setIntAtIndex(addr, 1, fmt.channelDataType); diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLKernel.java b/src/notzed.zcl/classes/au/notzed/zcl/CLKernel.java index bfdd96b..2c12aaf 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLKernel.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLKernel.java @@ -20,7 +20,7 @@ import static au.notzed.zcl.CL.*; import static au.notzed.zcl.CLLib.*; import jdk.incubator.foreign.*; import api.Native; -import api.Allocator; +import api.Frame; import api.Memory; import java.nio.ByteBuffer; import java.lang.invoke.MethodHandle; @@ -56,12 +56,12 @@ public class CLKernel extends CLObject { public CLKernel cloneKernel() throws CLRuntimeException { requireAPIVersion(CLPlatform.VERSION_2_1); - try (Allocator a = Memory.stack()) { - MemorySegment cres = a.allocs(8); - MemoryAddress ck = clCloneKernel(addr(), cres.address()); + try (Frame frame = Memory.createFrame()) { + MemorySegment pres = frame.allocateInt(); + MemoryAddress ck = clCloneKernel(address(), pres.address()); int res; - if ((res = MemoryAccess.getInt(cres)) != 0) + if ((res = MemoryAccess.getInt(pres)) != 0) throw new CLRuntimeException(res); return Native.resolve(ck, (x) -> new CLKernel(x, getObjectPlatform())); @@ -80,7 +80,7 @@ public class CLKernel extends CLObject { try { int res; - res = clSetKernelArg(addr(), index, size, pval); + res = clSetKernelArg(address(), index, size, pval); if (res != 0) throw new CLRuntimeException(res); @@ -118,10 +118,9 @@ public class CLKernel extends CLObject { * @param o */ public void setArg(int index, CLObject o) { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment pval = MemorySegment.allocateNative(CLinker.C_POINTER, frame); + try (Frame frame = Memory.createFrame()) { + MemorySegment pval = frame.copy(o); - MemoryAccess.setAddress(pval, o.addr()); setKernelArg(index, CLinker.C_POINTER.byteSize(), pval.address()); } } @@ -134,10 +133,9 @@ public class CLKernel extends CLObject { * @param val */ public void setArg(int index, byte val) { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment pval = MemorySegment.allocateNative(1, frame); + try (Frame frame = Memory.createFrame()) { + MemorySegment pval = frame.copy(val); - MemoryAccess.setByte(pval, val); setKernelArg(index, 1, pval.address()); } } @@ -150,10 +148,9 @@ public class CLKernel extends CLObject { * @param val */ public void setArg(int index, short val) { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment pval = MemorySegment.allocateNative(CLinker.C_SHORT, frame); + try (Frame frame = Memory.createFrame()) { + MemorySegment pval = frame.copy(val); - MemoryAccess.setShort(pval, val); setKernelArg(index, CLinker.C_SHORT.byteSize(), pval.address()); } } @@ -166,10 +163,9 @@ public class CLKernel extends CLObject { * @param val */ public void setArg(int index, int val) { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment pval = MemorySegment.allocateNative(CLinker.C_INT, frame); + try (Frame frame = Memory.createFrame()) { + MemorySegment pval = frame.copy(val); - MemoryAccess.setInt(pval, val); setKernelArg(index, CLinker.C_INT.byteSize(), pval.address()); } } @@ -182,10 +178,9 @@ public class CLKernel extends CLObject { * @param val */ public void setArg(int index, long val) { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment pval = MemorySegment.allocateNative(CLinker.C_LONG, frame); + try (Frame frame = Memory.createFrame()) { + MemorySegment pval = frame.copy(val); - MemoryAccess.setLong(pval, val); setKernelArg(index, CLinker.C_LONG.byteSize(), pval.address()); } } @@ -197,10 +192,9 @@ public class CLKernel extends CLObject { * @param val */ public void setArg(int index, float val) { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment pval = MemorySegment.allocateNative(CLinker.C_FLOAT, frame); + try (Frame frame = Memory.createFrame()) { + MemorySegment pval = frame.copy(val); - MemoryAccess.setFloat(pval, val); setKernelArg(index, CLinker.C_FLOAT.byteSize(), pval.address()); } } @@ -212,10 +206,9 @@ public class CLKernel extends CLObject { * @param val */ public void setArg(int index, double val) { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment pval = MemorySegment.allocateNative(CLinker.C_DOUBLE, frame); + try (Frame frame = Memory.createFrame()) { + MemorySegment pval = frame.copy(val); - MemoryAccess.setDouble(pval, val); setKernelArg(index, CLinker.C_DOUBLE.byteSize(), pval.address()); } } @@ -230,10 +223,16 @@ public class CLKernel extends CLObject { throw new UnsupportedOperationException("not yet"); } - public void setArg(int index, MemorySegment val) { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment pval = MemorySegment.allocateNative(val.byteSize(), 16, frame); + /** + * Set a non-native memory segment argument + * @param index + * @param val + */ + private void setArg(int index, MemorySegment val) { + try (Frame frame = Memory.createFrame()) { + MemorySegment pval = frame.allocate(val.byteSize(), 16); + pval.copyFrom(val); setKernelArg(index, pval.byteSize(), pval.address()); } } diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLMemory.java b/src/notzed.zcl/classes/au/notzed/zcl/CLMemory.java index e84c15d..2c5d8ea 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLMemory.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLMemory.java @@ -22,7 +22,7 @@ import jdk.incubator.foreign.*; import api.Native; import api.Callback; import api.Memory; -import api.Allocator; +import api.Frame; import java.lang.invoke.MethodHandle; import java.nio.ByteBuffer; @@ -94,12 +94,12 @@ public abstract class CLMemory extends CLObject { */ @Deprecated private static CLMemory create(MemoryAddress p, CLPlatform plat) { - if (Memory.toLong(p) == 0) + if (p.toRawLongValue() == 0) return null; - try (ResourceScope a = ResourceScope.newConfinedScope()) { - MemorySegment addr = getInfo(p, CL_MEM_TYPE, clGetMemObjectInfo, a, 4); + try (Frame frame = Memory.createFrame()) { + MemorySegment addr = getInfo(p, CL_MEM_TYPE, clGetMemObjectInfo, frame, 4); int type = MemoryAccess.getInt(addr); switch (type) { @@ -159,7 +159,7 @@ public abstract class CLMemory extends CLObject { destroyCallback = CLNotify.call(notify, (x) -> this); try { - int res = clSetMemObjectDestructorCallback(addr(), destroyCallback.addr(), MemoryAddress.NULL); + int res = clSetMemObjectDestructorCallback(address(), destroyCallback.address(), MemoryAddress.NULL); if (res != 0) throw new CLException(res); } catch (RuntimeException | Error t) { @@ -264,13 +264,12 @@ public abstract class CLMemory extends CLObject { * @since cl_khr_gl_sharing */ public GLObjectInfo getGLObjectInfo() { - try (Allocator frame = Memory.stack()) { - MemorySegment ctype = frame.allocs(8); - MemorySegment cname = frame.allocs(8); - MemoryAddress ce; + try (Frame frame = Memory.createFrame()) { + MemorySegment ctype = frame.allocateInt(); + MemorySegment cname = frame.allocateInt(); int res; - res = getGLext().clGetGLObjectInfo(addr(), ctype.address(), cname.address()); + res = getGLext().clGetGLObjectInfo(address(), ctype.address(), cname.address()); if (res != 0) throw new CLRuntimeException(res); return new GLObjectInfo(MemoryAccess.getInt(ctype), MemoryAccess.getInt(cname)); @@ -285,8 +284,8 @@ public abstract class CLMemory extends CLObject { * @since cl_khr_gl_sharing */ public int getGLTextureInfoInt(int param) { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { - return MemoryAccess.getInt(getInfo(addr(), CL_GL_TEXTURE_TARGET, getGLext().clGetGLTextureInfo, frame, 4)); + try (Frame frame = Memory.createFrame()) { + return MemoryAccess.getInt(getInfo(address(), CL_GL_TEXTURE_TARGET, getGLext().clGetGLTextureInfo, frame, 4)); } } diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLObject.java b/src/notzed.zcl/classes/au/notzed/zcl/CLObject.java index c7512fa..0cecffe 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLObject.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLObject.java @@ -16,6 +16,8 @@ */ package au.notzed.zcl; +import api.Frame; +import api.Memory; import java.lang.invoke.MethodHandle; import java.util.function.Function; import java.util.function.IntFunction; @@ -99,9 +101,9 @@ public abstract class CLObject extends Native { // some are static for access before object instantiation // new 5-param version - static MemorySegment getInfo(MemoryAddress self, int id, MethodHandle getInfo, ResourceScope frame, long size) throws CLRuntimeException { + static MemorySegment getInfo(MemoryAddress self, int id, MethodHandle getInfo, Frame frame, long size) throws CLRuntimeException { try { - MemorySegment addr = MemorySegment.allocateNative(size, 8, frame); + MemorySegment addr = frame.allocate(size, 8); int res; res = (int)getInfo.invokeExact(self, id, size, addr.address(), MemoryAddress.NULL); @@ -117,14 +119,14 @@ public abstract class CLObject extends Native { } protected int getInfoInt(int id, MethodHandle getInfo) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - return MemoryAccess.getInt(getInfo(addr(), id, getInfo, a, 4)); + try (Frame frame = Memory.createFrame()) { + return MemoryAccess.getInt(getInfo(address(), id, getInfo, frame, 4)); } } protected long getInfoLong(int id, MethodHandle getInfo) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - return MemoryAccess.getLong(getInfo(addr(), id, getInfo, a, 8)); + try (Frame frame = Memory.createFrame()) { + return MemoryAccess.getLong(getInfo(address(), id, getInfo, frame, 8)); } } @@ -133,21 +135,21 @@ public abstract class CLObject extends Native { } protected T getInfoAny(int id, MethodHandle getInfo, Function create) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - return Native.resolve(MemoryAccess.getAddress(getInfo(addr(), id, getInfo, a, 8)), create); + try (Frame frame = Memory.createFrame()) { + return Native.resolve(MemoryAccess.getAddress(getInfo(address(), id, getInfo, frame, 8)), create); } } protected T getInfoJava(int id, MethodHandle getInfo, Function create) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - return create.apply(getInfo(addr(), id, getInfo, a, 8)); + try (Frame frame = Memory.createFrame()) { + return create.apply(getInfo(address(), id, getInfo, frame, 8)); } } // new 5-param version for get any - static MemorySegment getInfoAny(MemoryAddress addr, int id, MethodHandle getInfo, ResourceScope a) throws CLRuntimeException { + static MemorySegment getInfoAny(MemoryAddress addr, int id, MethodHandle getInfo, Frame frame) throws CLRuntimeException { try { - MemorySegment sizep = MemorySegment.allocateNative(8, 8, a); + MemorySegment sizep = frame.allocateLong(); MemorySegment valp; long size; int res; @@ -155,7 +157,7 @@ public abstract class CLObject extends Native { res = (int)getInfo.invokeExact(addr, id, 0L, MemoryAddress.NULL, sizep.address()); size = MemoryAccess.getLong(sizep); - valp = MemorySegment.allocateNative(size, a); + valp = frame.allocate(size); res = (int)getInfo.invokeExact(addr, id, size, valp.address(), sizep.address()); @@ -171,32 +173,32 @@ public abstract class CLObject extends Native { } protected byte[] getInfoByteV(int id, MethodHandle getInfo) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - return getInfoAny(addr(), id, getInfo, a).toByteArray(); + try (Frame frame = Memory.createFrame()) { + return getInfoAny(address(), id, getInfo, frame).toByteArray(); } } protected String getInfoString(int id, MethodHandle getInfo) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - return infoToString(getInfoAny(addr(), id, getInfo, a)); + try (Frame frame = Memory.createFrame()) { + return infoToString(getInfoAny(address(), id, getInfo, frame)); } } protected T[] getInfoPropertyV(int id, MethodHandle getInfo, BiFunction create, IntFunction createArray) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - return CLProperty.fromNative(getInfoAny(addr(), id, getInfo, a), create, createArray); + try (Frame frame = Memory.createFrame()) { + return CLProperty.fromNative(getInfoAny(address(), id, getInfo, frame), create, createArray); } } protected long[] getInfoLongV(int id, MethodHandle getInfo) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - return Native.toLongV(getInfoAny(addr(), id, getInfo, a)); + try (Frame frame = Memory.createFrame()) { + return Native.toLongV(getInfoAny(address(), id, getInfo, frame)); } } protected T[] getInfoAnyV(int id, MethodHandle getInfo, Function create, IntFunction createArray) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - return Native.toObjectV(getInfoAny(addr(), id, getInfo, a), create, createArray); + try (Frame frame = Memory.createFrame()) { + return Native.toObjectV(getInfoAny(address(), id, getInfo, frame), create, createArray); } catch (RuntimeException | Error t) { throw t; } catch (Throwable t) { @@ -205,13 +207,13 @@ public abstract class CLObject extends Native { } // new 6-param version - protected MemorySegment getInfo(T ctx, int id, MethodHandle getInfo, ResourceScope a, long size) throws CLRuntimeException { + protected MemorySegment getInfo(T ctx, int id, MethodHandle getInfo, Frame frame, long size) throws CLRuntimeException { try { // TODO: use ... ValueLayout? - MemorySegment addr = MemorySegment.allocateNative(size, 16, a); + MemorySegment addr = frame.allocate(size, 16); int res; - res = (int)getInfo.invokeExact(addr(), ctx.addr(), id, size, addr.address(), MemoryAddress.NULL); + res = (int)getInfo.invokeExact(address(), ctx.address(), id, size, addr.address(), MemoryAddress.NULL); if (res != 0) throw new CLRuntimeException(res); @@ -224,14 +226,14 @@ public abstract class CLObject extends Native { } protected int getInfoInt(T ctx, int id, MethodHandle getInfo) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - return MemoryAccess.getInt(getInfo(ctx, id, getInfo, a, 4)); + try (Frame frame = Memory.createFrame()) { + return MemoryAccess.getInt(getInfo(ctx, id, getInfo, frame, 4)); } } protected long getInfoLong(T ctx, int id, MethodHandle getInfo) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - return MemoryAccess.getLong(getInfo(ctx, id, getInfo, a, 8)); + try (Frame frame = Memory.createFrame()) { + return MemoryAccess.getLong(getInfo(ctx, id, getInfo, frame, 8)); } } @@ -241,19 +243,19 @@ public abstract class CLObject extends Native { // new 6-param get-any // TODO: pass type descriptor for alignment - protected MemorySegment getInfoAny(T ctx, int id, MethodHandle getInfo, ResourceScope a) throws CLRuntimeException { + protected MemorySegment getInfoAny(T ctx, int id, MethodHandle getInfo, Frame frame) throws CLRuntimeException { try { - MemorySegment sizep = MemorySegment.allocateNative(CLinker.C_LONG, a); + MemorySegment sizep = frame.allocateLong(); MemorySegment valp; long size; int res; - res = (int)getInfo.invokeExact(addr(), ctx.addr(), id, 0L, MemoryAddress.NULL, sizep.address()); + res = (int)getInfo.invokeExact(address(), ctx.address(), id, 0L, MemoryAddress.NULL, sizep.address()); size = MemoryAccess.getLong(sizep); - valp = MemorySegment.allocateNative(size, 16, a); + valp = frame.allocate(size, 16); - res = (int)getInfo.invokeExact(addr(), ctx.addr(), id, size, valp.address(), sizep.address()); + res = (int)getInfo.invokeExact(address(), ctx.address(), id, size, valp.address(), sizep.address()); if (res != 0) throw new CLRuntimeException(res); @@ -267,8 +269,8 @@ public abstract class CLObject extends Native { } protected byte[] getInfoByteV(T ctx, int id, MethodHandle getInfo) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - MemorySegment seg = getInfoAny(ctx, id, getInfo, a); + try (Frame frame = Memory.createFrame()) { + MemorySegment seg = getInfoAny(ctx, id, getInfo, frame); return seg.toByteArray(); } } @@ -279,15 +281,15 @@ public abstract class CLObject extends Native { } protected String getInfoString(T ctx, int id, MethodHandle getInfo) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - MemorySegment seg = getInfoAny(ctx, id, getInfo, a); + try (Frame frame = Memory.createFrame()) { + MemorySegment seg = getInfoAny(ctx, id, getInfo, frame); return infoToString(seg); } } protected long[] getInfoLongV(T ctx, int id, MethodHandle getInfo) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - MemorySegment valp = getInfoAny(ctx, id, getInfo, a); + try (Frame frame = Memory.createFrame()) { + MemorySegment valp = getInfoAny(ctx, id, getInfo, frame); int len = (int)(valp.byteSize() >>> 3); long[] list = new long[len]; @@ -302,11 +304,11 @@ public abstract class CLObject extends Native { } // indexed version - protected MemorySegment getInfo(int index, int id, MethodHandle getInfo, ResourceScope a, long size) throws CLRuntimeException { + protected MemorySegment getInfo(int index, int id, MethodHandle getInfo, Frame frame, long size) throws CLRuntimeException { try { // TODO: type for alignment - MemorySegment addr = MemorySegment.allocateNative(size, 16, a); - int res = (int)getInfo.invokeExact(addr(), index, id, size, addr, MemoryAddress.NULL); + MemorySegment addr = frame.allocate(size, 16); + int res = (int)getInfo.invokeExact(address(), index, id, size, addr, MemoryAddress.NULL); if (res != 0) throw new IllegalArgumentException(); @@ -320,30 +322,30 @@ public abstract class CLObject extends Native { } protected int getInfoInt(int index, int id, MethodHandle getInfo) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - return MemoryAccess.getInt(getInfo(index, id, getInfo, a, 4)); + try (Frame frame = Memory.createFrame()) { + return MemoryAccess.getInt(getInfo(index, id, getInfo, frame, 4)); } } protected long getInfoLong(int index, int id, MethodHandle getInfo) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - return MemoryAccess.getLong(getInfo(index, id, getInfo, a, 8)); + try (Frame frame = Memory.createFrame()) { + return MemoryAccess.getLong(getInfo(index, id, getInfo, frame, 8)); } } - protected MemorySegment getInfoAny(int index, int id, MethodHandle getInfo, ResourceScope a) throws CLRuntimeException { + protected MemorySegment getInfoAny(int index, int id, MethodHandle getInfo, Frame frame) throws CLRuntimeException { try { - MemorySegment sizep = MemorySegment.allocateNative(CLinker.C_LONG, a); + MemorySegment sizep = frame.allocateInt(); MemorySegment valp; long size; int res; - res = (int)getInfo.invokeExact(addr(), index, id, 0L, MemoryAddress.NULL, sizep.address()); + res = (int)getInfo.invokeExact(address(), index, id, 0L, MemoryAddress.NULL, sizep.address()); size = MemoryAccess.getLong(sizep); - valp = MemorySegment.allocateNative(size, a); + valp = frame.allocate(size, 16); - res = (int)getInfo.invokeExact(addr(), index, id, size, valp.address(), sizep.address()); + res = (int)getInfo.invokeExact(address(), index, id, size, valp.address(), sizep.address()); if (res != 0) throw new CLRuntimeException(res); @@ -357,8 +359,8 @@ public abstract class CLObject extends Native { } protected String getInfoString(int index, int id, MethodHandle getInfo) { - try (ResourceScope a = ResourceScope.newConfinedScope()) { - MemorySegment valp = getInfoAny(index, id, getInfo, a); + try (Frame frame = Memory.createFrame()) { + MemorySegment valp = getInfoAny(index, id, getInfo, frame); return infoToString(valp); } } diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLPlatform.java b/src/notzed.zcl/classes/au/notzed/zcl/CLPlatform.java index d2614ef..078ee08 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLPlatform.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLPlatform.java @@ -16,6 +16,8 @@ */ package au.notzed.zcl; +import api.Frame; +import api.Memory; import static au.notzed.zcl.CL.*; import static au.notzed.zcl.CLLib.*; import java.util.function.ToDoubleFunction; @@ -87,8 +89,8 @@ public class CLPlatform extends CLObject { * @throws CLRuntimeException */ public static CLPlatform[] getPlatforms() /*throws CLRuntimeException*/ { - try (ResourceScope scope = ResourceScope.newConfinedScope()) { - MemorySegment lenp = MemorySegment.allocateNative(CLinker.C_INT, scope); + try (Frame frame = Memory.createFrame()) { + MemorySegment lenp = frame.allocateInt(); MemorySegment list; int len; int res; @@ -98,7 +100,7 @@ public class CLPlatform extends CLObject { throw new CLRuntimeException(res); len = MemoryAccess.getInt(lenp); - list = MemorySegment.allocateNative(MemoryLayout.sequenceLayout(len, CLinker.C_POINTER), scope); + list = frame.allocate(MemoryLayout.sequenceLayout(len, CLinker.C_POINTER)); res = (int)clGetPlatformIDs.invokeExact(len, list.address(), lenp.address()); if (res != 0) @@ -122,12 +124,12 @@ public class CLPlatform extends CLObject { * @throws CLRuntimeException */ public CLDevice[] getDevices(long type) /*throws CLRuntimeException*/ { - try (ResourceScope scope = ResourceScope.newConfinedScope()) { - MemorySegment lenp = MemorySegment.allocateNative(CLinker.C_INT, scope); + try (Frame frame = Memory.createFrame()) { + MemorySegment lenp = frame.allocateInt();; MemorySegment list; int res, len; - res = (int)clGetDeviceIDs.invokeExact(addr(), type, 0, MemoryAddress.NULL, lenp.address()); + res = (int)clGetDeviceIDs.invokeExact(address(), type, 0, MemoryAddress.NULL, lenp.address()); if (res == CL_DEVICE_NOT_FOUND) return new CLDevice[0]; @@ -135,9 +137,9 @@ public class CLPlatform extends CLObject { throw new CLRuntimeException(res); len = MemoryAccess.getInt(lenp); - list = MemorySegment.allocateNative(MemoryLayout.sequenceLayout(len, CLinker.C_POINTER), scope); + list = frame.allocate(MemoryLayout.sequenceLayout(len, CLinker.C_POINTER)); - res = (int)clGetDeviceIDs.invokeExact(addr(), type, len, list.address(), lenp.address()); + res = (int)clGetDeviceIDs.invokeExact(address(), type, len, list.address(), lenp.address()); if (res != 0) throw new CLRuntimeException(res); @@ -213,11 +215,11 @@ public class CLPlatform extends CLObject { * @return MemoryAddress of function entry point, or MemoryAddress.NULL. */ public MemoryAddress clGetExtensionFunctionAddressForPlatform(String name) { - try (ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment cname = CLinker.toCString(name, frame); + try (Frame frame = Memory.createFrame()) { + MemorySegment cname = frame.copy(name); if (apiVersion >= VERSION_1_2) { - return CLLib.clGetExtensionFunctionAddressForPlatform(addr(), cname.address()); + return CLLib.clGetExtensionFunctionAddressForPlatform(address(), cname.address()); } else { return clGetExtensionFunctionAddress(cname.address()); } diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLProgram.java b/src/notzed.zcl/classes/au/notzed/zcl/CLProgram.java index b8104cd..bafe559 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLProgram.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLProgram.java @@ -22,6 +22,7 @@ import jdk.incubator.foreign.*; import api.Native; import api.Memory; import api.Callback; +import api.Frame; import java.lang.invoke.MethodHandle; import java.util.stream.LongStream; @@ -56,13 +57,13 @@ public class CLProgram extends CLObject { * @throws CLException */ public void buildProgram(CLDevice[] devices, String options, CLNotify notify) throws CLException { - try(ResourceScope frame = ResourceScope.newConfinedScope(); + try (Frame frame = Memory.createFrame(); Callback> call = CLNotify.call(notify, (x) -> new CLProgram(x, getObjectPlatform()))) { - MemorySegment pdevs = toAddrV(frame, devices); - MemorySegment poptions = options != null ? CLinker.toCString(options, frame) : MemorySegment.globalNativeSegment(); + MemorySegment pdevs = frame.copy(devices); + MemorySegment poptions = frame.copy(options); int res; - res = clBuildProgram(addr(), devices.length, pdevs.address(), poptions.address(), call.addr(), MemoryAddress.NULL); + res = clBuildProgram(address(), devices.length, pdevs.address(), poptions.address(), call.address(), MemoryAddress.NULL); if (res != 0) throw new CLException(res); } catch (CLException | RuntimeException | Error t) { @@ -117,15 +118,15 @@ public class CLProgram extends CLObject { throw new IllegalArgumentException(); } - try(ResourceScope frame = ResourceScope.newConfinedScope(); + try (Frame frame = Memory.createFrame(); Callback> call = CLNotify.call(notify, (x) -> new CLProgram(x, getObjectPlatform()))) { - MemorySegment cdevs = toAddrV(frame, devices); - MemorySegment poptions = CLinker.toCString(options, frame); - MemorySegment cheaders = toAddrV(frame, headers); - MemorySegment cnames = Native.toCStringV(header_names, frame); + MemorySegment cdevs = frame.copy(devices); + MemorySegment poptions = frame.copy(options); + MemorySegment cheaders = frame.copy(headers); + MemorySegment cnames = frame.copy(header_names); int res; - res = clCompileProgram(addr(), devices.length, cdevs.address(), poptions.address(), nheaders, cheaders.address(), cnames.address(), call.addr(), MemoryAddress.NULL); + res = clCompileProgram(address(), devices.length, cdevs.address(), poptions.address(), nheaders, cheaders.address(), cnames.address(), call.address(), MemoryAddress.NULL); if (res != 0) throw new CLException(res); } catch (CLException | RuntimeException | Error t) { @@ -144,13 +145,13 @@ public class CLProgram extends CLObject { * @throws CLRuntimeException */ public CLKernel createKernel(String name) throws CLException { - try(ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment pres = MemorySegment.allocateNative(CLinker.C_INT, frame); - MemorySegment pname = CLinker.toCString(name, frame); + try (Frame frame = Memory.createFrame()) { + MemorySegment pres = frame.allocateInt(); + MemorySegment pname = frame.copy(name); int res; MemoryAddress ck; - ck = clCreateKernel(addr(), pname.address(), pres.address()); + ck = clCreateKernel(address(), pname.address(), pres.address()); if ((res = MemoryAccess.getInt(pres)) != 0) throw new CLException(res); return resolve(ck, (x) -> new CLKernel(x, getObjectPlatform())); @@ -168,19 +169,19 @@ public class CLProgram extends CLObject { * @throws CLRuntimeException */ public CLKernel[] createKernelsInProgram() throws CLRuntimeException { - int size = getNumKernels(); + int count = getNumKernels(); - try(ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment csize = MemorySegment.allocateNative(CLinker.C_INT, frame); - MemorySegment ckern = Native.allocAddrV(frame, size); + try (Frame frame = Memory.createFrame()) { + MemorySegment csize = frame.allocateInt(); + MemorySegment ckern = frame.allocateAddress(count); int res; - res = clCreateKernelsInProgram(addr(), size, ckern.address(), csize.address()); + res = clCreateKernelsInProgram(address(), count, ckern.address(), csize.address()); if (res != 0) throw new CLRuntimeException(); - size = MemoryAccess.getInt(csize); - return Native.toObjectV(ckern, new CLKernel[size], (x) -> new CLKernel(x, getObjectPlatform())); + count = MemoryAccess.getInt(csize); + return Native.toObjectV(ckern, new CLKernel[count], (x) -> new CLKernel(x, getObjectPlatform())); } catch (RuntimeException | Error t) { throw t; } catch (Throwable t) { @@ -211,10 +212,10 @@ public class CLProgram extends CLObject { long[] sizes = getInfoSizeTA(CL_PROGRAM_BINARY_SIZES); long size = LongStream.of(sizes).sum(); - try (ResourceScope frame = ResourceScope.newConfinedScope()) { - MemorySegment seg = MemorySegment.allocateNative(size, frame); + try (Frame frame = Memory.createFrame()) { + MemorySegment seg = frame.allocate(size); MemoryAddress data = seg.address(); - MemorySegment cptrs = Native.allocAddrV(frame, sizes.length); + MemorySegment cptrs = frame.allocateAddress(sizes.length); long off = 0; int res; @@ -222,7 +223,7 @@ public class CLProgram extends CLObject { MemoryAccess.setAddressAtIndex(cptrs, i, data.addOffset(off)); off += sizes[i]; } - res = clGetProgramInfo(addr(), CL_PROGRAM_BINARIES, sizes.length * CLinker.C_POINTER.byteSize(), cptrs.address(), MemoryAddress.NULL); + res = clGetProgramInfo(address(), CL_PROGRAM_BINARIES, sizes.length * CLinker.C_POINTER.byteSize(), cptrs.address(), MemoryAddress.NULL); if (res != 0) throw new CLRuntimeException(res); diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLProperty.java b/src/notzed.zcl/classes/au/notzed/zcl/CLProperty.java index a77cf8f..71ffbd1 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLProperty.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLProperty.java @@ -17,6 +17,7 @@ package au.notzed.zcl; +import api.Frame; import static au.notzed.zcl.CL.*; import jdk.incubator.foreign.MemorySegment; @@ -24,7 +25,6 @@ import java.util.function.BiFunction; import java.util.function.IntFunction; import java.util.ArrayList; import jdk.incubator.foreign.MemoryAccess; -import jdk.incubator.foreign.ResourceScope; /** * All property arrays implement this interface. @@ -104,9 +104,9 @@ public interface CLProperty { return size + 1; } - static MemorySegment toNative(ResourceScope frame, T[] properties) { + static MemorySegment toNative(Frame frame, T[] properties) { if (properties != null && properties.length > 0) { - MemorySegment addr = MemorySegment.allocateNative(getSize(properties) * 8, 8, frame); + MemorySegment addr = frame.allocateLong(getSize(properties) * 8); int i = 0; for (CLProperty p: properties) @@ -114,7 +114,7 @@ public interface CLProperty { MemoryAccess.setLongAtIndex(addr, i, 0L); return addr; } else { - MemorySegment addr = MemorySegment.allocateNative(8, 8, frame); + MemorySegment addr = frame.allocateLong(); MemoryAccess.setLong(addr, 0L); return addr; } diff --git a/src/notzed.zcl/classes/au/notzed/zcl/EventInfo.java b/src/notzed.zcl/classes/au/notzed/zcl/EventInfo.java index 4d59123..e117294 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/EventInfo.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/EventInfo.java @@ -17,6 +17,7 @@ package au.notzed.zcl; +import api.Frame; import jdk.incubator.foreign.*; /** @@ -33,7 +34,7 @@ class EventInfo { public final MemoryAddress event; // TODO: frame isn't used, maybe it should be removed? if so check callers where scope isn't doing anyting - public EventInfo(ResourceScope frame, CLEventList waiters, CLEventList events) { + public EventInfo(Frame frame, CLEventList waiters, CLEventList events) { nwait = waiters != null ? waiters.size() : 0; wait = nwait > 0 ? waiters.slots() : MemoryAddress.NULL; event = events != null ? events.currentSlot() : MemoryAddress.NULL; diff --git a/src/notzed.zcl/tests/au/notzed/zcl/CLEventTest.java b/src/notzed.zcl/tests/au/notzed/zcl/CLEventTest.java index 2790c14..8146156 100644 --- a/src/notzed.zcl/tests/au/notzed/zcl/CLEventTest.java +++ b/src/notzed.zcl/tests/au/notzed/zcl/CLEventTest.java @@ -143,10 +143,10 @@ public class CLEventTest { int countEvent(MemoryAddress x) { int res = -1; - try (Allocator a = Memory.stack()) { - MemorySegment rc = a.allocs(8); - clGetEventInfo(x, CL_EVENT_REFERENCE_COUNT, 4, rc.address(), MemoryAddress.NULL); - res = MemoryAccess.getInt(rc); + try (Frame frame = Memory.createFrame()) { + MemorySegment pres = frame.allocateInt(); + clGetEventInfo(x, CL_EVENT_REFERENCE_COUNT, 4, pres.address(), MemoryAddress.NULL); + res = MemoryAccess.getInt(pres); } catch (Throwable T) { } return res; -- 2.39.5