From f11f0d4aed0e904928f9ddb2af319561f307d16c Mon Sep 17 00:00:00 2001 From: Not Zed Date: Tue, 28 Jan 2020 21:59:06 +1030 Subject: [PATCH] Implement the few extensions previously implemented. Not entirely pleased, still needs work. --- Makefile | 1 - src/notzed.zcl/classes/api/Native.java | 2 +- .../classes/au/notzed/zcl/CLCommandQueue.java | 64 +++--- .../classes/au/notzed/zcl/CLContext.java | 85 +++++--- .../classes/au/notzed/zcl/CLDevice.java | 5 + .../classes/au/notzed/zcl/CLEventList.java | 6 +- .../classes/au/notzed/zcl/CLExtendable.java | 13 +- .../classes/au/notzed/zcl/CLExtension.java | 19 +- .../classes/au/notzed/zcl/CLMemory.java | 46 ++++- .../classes/au/notzed/zcl/CLObject.java | 3 +- .../classes/au/notzed/zcl/CLPlatform.java | 79 +++++--- .../au/notzed/zcl/internal/EventInfo.java | 47 +++++ .../classes/au/notzed/zcl/khr/GLEvent.java | 42 +++- .../classes/au/notzed/zcl/khr/GLSharing.java | 184 +++++++++++++++--- src/notzed.zcl/classes/module-info.java | 2 +- src/notzed.zcl/gen/gen.make | 12 +- src/notzed.zcl/gen/generate-api | 168 +++++++--------- src/notzed.zcl/gen/opencl-ext.txt | 11 ++ src/notzed.zcl/gen/opencl.pm | 127 +++++++++++- 19 files changed, 648 insertions(+), 268 deletions(-) create mode 100644 src/notzed.zcl/classes/au/notzed/zcl/internal/EventInfo.java create mode 100644 src/notzed.zcl/gen/opencl-ext.txt diff --git a/Makefile b/Makefile index 0145faf..f21e8aa 100644 --- a/Makefile +++ b/Makefile @@ -22,7 +22,6 @@ include java.make # Work in progress idea for java.make extension to create execution templates test_demos := $(subst /,.,$(subst .java,,$(shell find src/notzed.zcl.demo/classes -name 'Test*.java' -printf '%P\n'))) -$(info $(test_demos)) notzed.zcl.demo_DEMOS=au.notzed.zcl.tools.clinfo \ $(test_demos) diff --git a/src/notzed.zcl/classes/api/Native.java b/src/notzed.zcl/classes/api/Native.java index db0dfbe..0516199 100644 --- a/src/notzed.zcl/classes/api/Native.java +++ b/src/notzed.zcl/classes/api/Native.java @@ -786,7 +786,7 @@ public class Native { */ private static void cleaner() { if (dolog) - log().log(Level.INFO, "Native finaliser started"); + log().log(Level.DEBUG, "Native finaliser started"); try { while (true) { CHandle stale = (CHandle) references.remove(); diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLCommandQueue.java b/src/notzed.zcl/classes/au/notzed/zcl/CLCommandQueue.java index 063b480..73e7ae3 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLCommandQueue.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLCommandQueue.java @@ -18,6 +18,7 @@ package au.notzed.zcl; import static au.notzed.zcl.CL.*; import static au.notzed.zcl.CLLib.*; +import au.notzed.zcl.internal.EventInfo; import jdk.incubator.foreign.*; import api.Native; import api.Allocator; @@ -30,7 +31,7 @@ import java.nio.ByteBuffer; import java.nio.ByteOrder; import java.nio.BufferOverflowException; import java.nio.BufferUnderflowException; -//import au.notzed.zcl.khr.GLSharing; +import au.notzed.zcl.khr.*; import java.lang.invoke.MethodHandle; import java.util.function.Function; @@ -1476,30 +1477,6 @@ public class CLCommandQueue extends CLExtendable { } } - /** - * Simplify wait/event handling. - * - * To use, create an EventInfo from the passed in arguments. - * In the enqueue command pass in .wait and .event from this - * structure. If the command succeeds, then call post(). - */ - static private class EventInfo { - final int nwait; - final MemoryAddress wait; - final MemoryAddress event; - - EventInfo(Allocator 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; - } - - void post(CLEventList events) { - if (events != null) - events.incrementSlot(); - } - } - /** * Call clEnqueueNDRangeKernel. * @@ -1919,26 +1896,31 @@ public class CLCommandQueue extends CLExtendable { return getDevice().platform; } - //protected GLSharing getGLSharing() { - // return getExtension(GLSharing.class, CLPlatform.cl_khr_gl_sharing); - //} + protected GLSharing getGLSharing() { + return getExtension(CLPlatform.cl_khr_gl_sharing, (p) -> { + if (getDevice().hasDeviceExtension(GLSharing.NAME)) + return new GLSharing(p); + else + throw new UnsupportedOperationException(); + }); + } /* Experimental: Alternative interface to extensions. */ - //public void enqueueAcquireGLObjects( - // CLMemory[] mem_objects, - // CLEventList waiters, - // CLEventList events) { - // getGLSharing().enqueueAcquireGLObjects(this, mem_objects, waiters, events); - //} - - //public void enqueueReleaseGLObjects( - // CLMemory[] mem_objects, - // CLEventList waiters, - // CLEventList events) { - // getGLSharing().enqueueReleaseGLObjects(this, mem_objects, waiters, events); - //} + public void enqueueAcquireGLObjects( + CLMemory[] mem_objects, + CLEventList waiters, + CLEventList events) { + getGLSharing().enqueueAcquireGLObjects(this, mem_objects, waiters, events); + } + + public void enqueueReleaseGLObjects( + CLMemory[] mem_objects, + CLEventList waiters, + CLEventList events) { + getGLSharing().enqueueReleaseGLObjects(this, mem_objects, waiters, events); + } /** * Invoke task.queue for this queue with no event lists. diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLContext.java b/src/notzed.zcl/classes/au/notzed/zcl/CLContext.java index 564964c..bd5bc02 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLContext.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLContext.java @@ -30,6 +30,7 @@ import java.io.ByteArrayOutputStream; import java.io.IOException; import java.io.InputStream; import java.nio.charset.Charset; +import au.notzed.zcl.khr.*; /** * Interface for cl_context @@ -990,39 +991,61 @@ public class CLContext extends CLExtendable { return getDevices()[0].platform; } - //protected GLSharing getGLSharing() { - // return getExtension(GLSharing.class, CLPlatform.cl_khr_gl_sharing); - //} - - //protected GLEvent getGLEvent() { - // return getExtension(GLEvent.class, CLPlatform.cl_khr_gl_event); - //} - /* Experimental: extension interface mechanism */ - // public CLBuffer createFromGLBuffer( - // long flags, - // int bufobj) { - // return getGLSharing().createFromGLBuffer(this, flags, bufobj); - // } - - // public CLImage createFromGLTexture( - // long flags /* flags */, - // int target /* target */, - // int miplevel /* miplevel */, - // int texture /* texture */) { - // return getGLSharing().createFromGLTexture(this, flags, target, miplevel, texture); - // } - - // public CLImage createFromGLRenderbuffer( - // long flags /* flags */, - // int renderbuffer /* renderbuffer */) { - // return getGLSharing().createFromGLRenderbuffer(this, flags, renderbuffer); - // } - - // public CLEvent createEventFromGLsync(long glsync) { - // return getGLEvent().clCreateEventFromGLsync(this, glsync); - // } + protected GLSharing getGLSharing() { + return getExtension(CLPlatform.cl_khr_gl_sharing, (p) -> { + if (getDevices()[0].hasDeviceExtension(GLSharing.NAME)) + return new GLSharing(p); + else + throw new UnsupportedOperationException(); + }); + } + + protected GLEvent getGLEvent() { + return getExtension(CLPlatform.cl_khr_gl_sharing, (p) -> { + if (getDevices()[0].hasDeviceExtension(GLEvent.NAME)) + return new GLEvent(p); + else + throw new UnsupportedOperationException(); + }); + } + + /** + * @since cl_khr_gl_sharing extension + */ + public CLBuffer createFromGLBuffer( + long flags, + int bufobj) { + return getGLSharing().createFromGLBuffer(this, flags, bufobj); + } + + /** + * @since cl_khr_gl_sharing extension + */ + public CLImage createFromGLTexture( + long flags /* flags */, + int target /* target */, + int miplevel /* miplevel */, + int texture /* texture */) { + return getGLSharing().createFromGLTexture(this, flags, target, miplevel, texture); + } + + /** + * @since cl_khr_gl_sharing extension + */ + public CLImage createFromGLRenderbuffer( + long flags /* flags */, + int renderbuffer /* renderbuffer */) { + return getGLSharing().createFromGLRenderbuffer(this, flags, renderbuffer); + } + + /** + * @since cl_khr_gl_event extension + */ + public CLEvent createEventFromGLsync(MemoryAddress glsync) { + return getGLEvent().clCreateEventFromGLsync(this, glsync); + } } diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLDevice.java b/src/notzed.zcl/classes/au/notzed/zcl/CLDevice.java index c97d51c..c8fbd00 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLDevice.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLDevice.java @@ -20,6 +20,7 @@ import java.lang.invoke.MethodHandle; import jdk.incubator.foreign.*; import static au.notzed.zcl.CL.*; import static au.notzed.zcl.CLLib.*; +import java.util.stream.Stream; /** * Interface for cl_device_id. @@ -383,6 +384,10 @@ public class CLDevice extends CLExtendable { return getInfoString(CL_DEVICE_EXTENSIONS); } + public boolean hasDeviceExtension(String name) { + return Stream.of(getDeviceExtensions().split(" ")).anyMatch(name::equals); + } + public long getPrintfBufferSize() { return getInfoSizeT(CL_DEVICE_PRINTF_BUFFER_SIZE); } diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLEventList.java b/src/notzed.zcl/classes/au/notzed/zcl/CLEventList.java index 21a30f8..a633bc1 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLEventList.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLEventList.java @@ -112,7 +112,7 @@ public final class CLEventList implements AutoCloseable { * * This is used internally by CLCommandQueue.EventInfo to write directly to the event list. */ - MemoryAddress slots() { + public MemoryAddress slots() { return cevents; } @@ -124,7 +124,7 @@ public final class CLEventList implements AutoCloseable { * @throws IllegalStateException if the CLEventList has been released. * @throws ArrayIndexOutOfBoundsException if the CLEventList is full. */ - MemoryAddress currentSlot() { + public MemoryAddress currentSlot() { if (index < jevents.length) { MemoryAddress addr = cevents.addOffset(index * 8); @@ -140,7 +140,7 @@ public final class CLEventList implements AutoCloseable { * * This is used internally by CLCommandQueue.EventInfo to write directly to the event list. */ - void incrementSlot() { + public void incrementSlot() { index++; } diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLExtendable.java b/src/notzed.zcl/classes/au/notzed/zcl/CLExtendable.java index 44837df..213f5fe 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLExtendable.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLExtendable.java @@ -17,6 +17,7 @@ package au.notzed.zcl; import jdk.incubator.foreign.MemoryAddress; +import java.util.function.Function; /** * Extendable object. These keep track of the platform and api revision to be @@ -41,6 +42,9 @@ public abstract class CLExtendable extends CLObject { * Retrieve the platform. This should not cache the lookup. It cannot return * null by definition. This should not be called by any implementing class. * + * TODO: this is expensive to call every object instance the time, find a better way. + * Probably pass platform or parent in, with ability to look it up (on demand) if null. + * TODO: move the whole class to CLObject? * @return */ protected abstract CLPlatform initPlatform(); @@ -66,12 +70,13 @@ public abstract class CLExtendable extends CLObject { * CLExtenable. * * @param - * @param klass * @param id The extension id code on CLPlatform. + * @param create constructor method. This should perform + * extension availability checks. It is only called once per + * platform. * @return */ - protected T getExtension(Class klass, int id) { - //return platform.getExtension(klass, id); - return null; + protected T getExtension(int id, Function create) { + return platform.getExtension(id, create); } } diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLExtension.java b/src/notzed.zcl/classes/au/notzed/zcl/CLExtension.java index e4e165e..f881a38 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLExtension.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLExtension.java @@ -16,14 +16,10 @@ */ package au.notzed.zcl; -import jdk.incubator.foreign.MemoryAddress; -import java.lang.invoke.MethodHandle; - /** * Experimental code for extension support. *

- * Extensions need to be per-platform and are backed by a C structure - * which holds the function pointers. + * Extensions need to be per-platform. *

* Because some extensions effectively expand the method or property sets * of the basic objects it may make sense for the extension methods to appear @@ -31,17 +27,8 @@ import java.lang.invoke.MethodHandle; *

* Actually the above is always true otherwise it becomes one huge fuckup to use. */ -public abstract class CLExtension extends CLObject { - - protected CLExtension(MemoryAddress p) { - super(p); - } - - public abstract String getName(); +public interface CLExtension { - @Override - MethodHandle getInfoFunc() { - throw new UnsupportedOperationException(); - } + public String getName(); } diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLMemory.java b/src/notzed.zcl/classes/au/notzed/zcl/CLMemory.java index 0625cbf..f1151f0 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLMemory.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLMemory.java @@ -18,6 +18,7 @@ package au.notzed.zcl; import static au.notzed.zcl.CL.*; import static au.notzed.zcl.CLLib.*; +import au.notzed.zcl.khr.*; import jdk.incubator.foreign.*; import api.Native; import api.Callback; @@ -63,7 +64,7 @@ import java.nio.ByteOrder; * performance penalty over simply passing the byte offset as with the array * methods. It may change (again) in the future? */ -public abstract class CLMemory extends CLObject { +public abstract class CLMemory extends CLExtendable { /** * If use USE_HOST_PTR was used then this keeps track of the @@ -120,13 +121,17 @@ public abstract class CLMemory extends CLObject { static void release(MemoryAddress p) { // note: no way to free the hostSegment, even if we could - System.out.println("*** release clmemory"); try { clReleaseMemObject(p); } catch (Throwable t) { } } + @Override + protected CLPlatform initPlatform() { + return getContext().getDevices()[0].getPlatform(); + } + /** * If CL_MEM_USE_HOST_PTR was used at creation then this must * be invoked to avoid a memory leak. It also must be invoked @@ -237,6 +242,43 @@ public abstract class CLMemory extends CLObject { return getInfoSizeT(CL_MEM_OFFSET); } + protected GLSharing getGLSharing() { + return getExtension(CLPlatform.cl_khr_gl_sharing, (p) -> { + if (getContext().getDevices()[0].hasDeviceExtension(GLSharing.NAME)) + return new GLSharing(p); + else + throw new UnsupportedOperationException(); + }); + } + + /** + * @since cl_khr_gl_sharing + */ + public GLSharing.GLObjectInfo getGLObjectInfo() { + return getGLSharing().getGLObjectInfo(this); + } + + /** + * @since cl_khr_gl_sharing + */ + public int getGLTextureTarget() { + return getGLSharing().getGLTextureTarget(this); + } + + /** + * @since cl_khr_gl_sharing + */ + public int getGLMIPMapLevelTextureTarget() { + return getGLSharing().getGLMIPMapLevel(this); + } + + /** + * @since cl_khr_gl_sharing + */ + public int getGLNumSamples() { + return getGLSharing().getGLNumSamples(this); + } + /** * Allocates a buffer suitable for opencl use - sets the byte order. * diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLObject.java b/src/notzed.zcl/classes/au/notzed/zcl/CLObject.java index ba54975..6619c11 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLObject.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLObject.java @@ -62,7 +62,8 @@ public abstract class CLObject extends Native { // new 5-param version // this one is static so it can be accessed at creation time - protected static MemoryAddress getInfo(MemoryAddress self, int id, MethodHandle getInfo, Allocator frame, long size) throws CLRuntimeException { + // public so extensions can see it, move to internal package + public static MemoryAddress getInfo(MemoryAddress self, int id, MethodHandle getInfo, Allocator frame, long size) throws CLRuntimeException { try { MemoryAddress addr = frame.alloca(size); int res; diff --git a/src/notzed.zcl/classes/au/notzed/zcl/CLPlatform.java b/src/notzed.zcl/classes/au/notzed/zcl/CLPlatform.java index 9b11165..64e6971 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/CLPlatform.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/CLPlatform.java @@ -19,6 +19,7 @@ package au.notzed.zcl; import static au.notzed.zcl.CL.*; import static au.notzed.zcl.CLLib.*; import java.util.function.ToDoubleFunction; +import java.util.function.Function; import jdk.incubator.foreign.*; import api.Native; import api.Memory; @@ -195,6 +196,28 @@ public class CLPlatform extends CLObject { } } + /** + * Calls clGetExtensionFunctionAddressForPlatform.If not available then + * it falls back to clGetExtensionFunctionAddress. + * @param name extension function name + * @return MemoryAddress of function entry point, or MemoryAddress.NULL. + */ + public MemoryAddress clGetExtensionFunctionAddressForPlatform(String name) { + try (Allocator frame = Memory.stack()) { + MemoryAddress cname = toByteV(frame, name); + + if (apiVersion >= VERSION_1_2) { + return CLLib.clGetExtensionFunctionAddressForPlatform(addr(), cname); + } else { + return clGetExtensionFunctionAddress(cname); + } + } catch (RuntimeException | Error t) { + throw t; + } catch (Throwable t) { + throw new RuntimeException(t); + } + } + /** * Get the platform api versiom. * @@ -266,34 +289,38 @@ public class CLPlatform extends CLObject { *

* Extensions are bound to platforms. */ - //final CLExtension[] extensions = new CLExtension[2]; - - //native CLExtension createExtension(int extension); // throws something - - //public T getExtension(Class klass, int id) { - // synchronized (extensions) { - // if (extensions[id] == null) { - // extensions[id] = createExtension(id); - // } - // return klass.cast(extensions[id]); - // } - //} + private final CLExtension[] extension = new CLExtension[2]; /** - * Retrieve an extension by name. + * Retrieve an extension by extension id. + * + * The extension id must be one of the supported extensions. + * This is called internally by the class/method to which the + * extension applies. It should check the published list of + * extensions for validity and throw + * UnsupportedOperationException if it isn't. * - * @param - * @param klass Required type of extension - * @param name - * @return + * TODO: pass the class, constructor.invoke? id = getfield()? */ - //public T getExtension(Class klass, String name) { - // switch (name) { - // case GLSharing.NAME: - // return getExtension(klass, cl_khr_gl_sharing); - // case au.notzed.zcl.khr.GLEvent.NAME: - // return getExtension(klass, cl_khr_gl_event); - // } - // return null; - //} + @SuppressWarnings("unchecked") + T getExtension(int id, Function create) { + synchronized (extension) { + T x = (T)extension[id]; + + if (x == null) { + // or constructor.invoke + switch (id) { + case cl_khr_gl_sharing: + x = create.apply(this); + break; + case cl_khr_gl_event: + x = create.apply(this); + break; + } + extension[id] = x; + } + + return x; + } + } } diff --git a/src/notzed.zcl/classes/au/notzed/zcl/internal/EventInfo.java b/src/notzed.zcl/classes/au/notzed/zcl/internal/EventInfo.java new file mode 100644 index 0000000..53236bb --- /dev/null +++ b/src/notzed.zcl/classes/au/notzed/zcl/internal/EventInfo.java @@ -0,0 +1,47 @@ +/* + * 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 au.notzed.zcl.internal; + +import jdk.incubator.foreign.*; +import au.notzed.zcl.CLEventList; +import api.Allocator; + +/** + * Simplify wait/event handling. + * + * To use, create an EventInfo from the passed in arguments. + * In the enqueue command pass in .wait and .event from this + * structure. If the command succeeds, then call post(). + * + */ +public class EventInfo { + public final int nwait; + public final MemoryAddress wait; + public final MemoryAddress event; + + public EventInfo(Allocator 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; + } + + public void post(CLEventList events) { + if (events != null) + events.incrementSlot(); + } +} diff --git a/src/notzed.zcl/classes/au/notzed/zcl/khr/GLEvent.java b/src/notzed.zcl/classes/au/notzed/zcl/khr/GLEvent.java index 449da72..3fa07ce 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/khr/GLEvent.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/khr/GLEvent.java @@ -19,29 +19,57 @@ package au.notzed.zcl.khr; import au.notzed.zcl.CLContext; import au.notzed.zcl.CLEvent; import au.notzed.zcl.CLExtension; +import au.notzed.zcl.CLPlatform; +import au.notzed.zcl.CLRuntimeException; +import au.notzed.zcl.CLext; import jdk.incubator.foreign.MemoryAddress; +import java.lang.invoke.MethodHandle; +import java.util.function.Function; +import api.*; /** - * cl_khr_gl_sharing extension interface. + * cl_khr_gl_event extension interface. */ -public class GLEvent extends CLExtension { +public class GLEvent implements CLExtension { - public GLEvent(MemoryAddress p) { - super(p); - } + private final CLPlatform plat; + private final MethodHandle clCreateEventFromGLsyncKHR; + + public GLEvent(CLPlatform plat) { + Function find = plat::clGetExtensionFunctionAddressForPlatform; - private native static void release(long p); + this.clCreateEventFromGLsyncKHR = CLext.clCreateEventFromGLsyncKHR(find); + this.plat = plat; + } @Override public String getName() { return NAME; } + public final static int ID = CLPlatform.cl_khr_gl_event; public final static String NAME = "cl_khr_gl_event"; public static final int CL_COMMAND_GL_FENCE_SYNC_OBJECT_KHR = 0x200D; - public native CLEvent clCreateEventFromGLsync(CLContext ctx, long glsync); + public CLEvent clCreateEventFromGLsync(CLContext ctx, MemoryAddress glsync) { + try (Allocator frame = Memory.stack()) { + MemoryAddress cret = frame.alloca(8); + MemoryAddress ce; + int res; + + ce = (MemoryAddress)clCreateEventFromGLsyncKHR.invokeExact(ctx.addr(), glsync, cret); + res = Native.getInt(cret); + if (res != 0) + throw new CLRuntimeException(res); + + return Native.resolve(ce, CLEvent::new); + } catch (RuntimeException | Error t) { + throw t; + } catch (Throwable t) { + throw new RuntimeException(t); + } + } } diff --git a/src/notzed.zcl/classes/au/notzed/zcl/khr/GLSharing.java b/src/notzed.zcl/classes/au/notzed/zcl/khr/GLSharing.java index 245c907..4f98358 100644 --- a/src/notzed.zcl/classes/au/notzed/zcl/khr/GLSharing.java +++ b/src/notzed.zcl/classes/au/notzed/zcl/khr/GLSharing.java @@ -16,6 +16,7 @@ */ package au.notzed.zcl.khr; +import static au.notzed.zcl.CLext.*; import au.notzed.zcl.CLBuffer; import au.notzed.zcl.CLCommandQueue; import au.notzed.zcl.CLContext; @@ -24,26 +25,51 @@ import au.notzed.zcl.CLEventList; import au.notzed.zcl.CLExtension; import au.notzed.zcl.CLImage; import au.notzed.zcl.CLMemory; +import au.notzed.zcl.CLPlatform; +import au.notzed.zcl.CLObject; import au.notzed.zcl.CLRuntimeException; +import au.notzed.zcl.internal.EventInfo; import jdk.incubator.foreign.MemoryAddress; +import java.lang.invoke.MethodHandle; +import java.util.function.Function; +import api.*; /** * cl_khr_gl_sharing extension interface. */ -public class GLSharing extends CLExtension { - - public GLSharing(MemoryAddress p) { - super(p); +public class GLSharing implements CLExtension { + private final CLPlatform plat; + + private final MethodHandle clCreateFromGLBuffer; + private final MethodHandle clCreateFromGLTexture; + private final MethodHandle clCreateFromGLRenderbuffer; + private final MethodHandle clEnqueueAcquireGLObjects; + private final MethodHandle clEnqueueReleaseGLObjects; + private final MethodHandle clGetGLContextInfoKHR; + private final MethodHandle clGetGLObjectInfo; + private final MethodHandle clGetGLTextureInfo; + + public GLSharing(CLPlatform plat) { + Function find = plat::clGetExtensionFunctionAddressForPlatform; + + this.clCreateFromGLBuffer = clCreateFromGLBuffer(find); + this.clCreateFromGLTexture = clCreateFromGLTexture(find); + this.clCreateFromGLRenderbuffer = clCreateFromGLRenderbuffer(find); + this.clEnqueueAcquireGLObjects = clEnqueueAcquireGLObjects(find); + this.clEnqueueReleaseGLObjects = clEnqueueReleaseGLObjects(find); + this.clGetGLContextInfoKHR = clGetGLContextInfoKHR(find); + this.clGetGLObjectInfo = clGetGLObjectInfo(find); + this.clGetGLTextureInfo = clGetGLTextureInfo(find); + this.plat = plat; } - private native static void release(long p); - @Override public String getName() { return NAME; } + public final static int ID = CLPlatform.cl_khr_gl_sharing; public final static String NAME = "cl_khr_gl_sharing"; /* cl_gl_object_type = 0x2000 - 0x200F enum values are currently taken */ @@ -95,19 +121,67 @@ public class GLSharing extends CLExtension { return new CLContextProperty.TagValue(CL_CGL_SHAREGROUP_KHR, id); } - public native CLBuffer createFromGLBuffer(CLContext ctx, + public CLBuffer createFromGLBuffer(CLContext ctx, long flags, - int bufobj); + int bufobj) { + try (Allocator frame = Memory.stack()) { + MemoryAddress cres = frame.alloca(8); + MemoryAddress ce; + int res; + + ce = (MemoryAddress)clCreateFromGLBuffer.invokeExact(ctx.addr(), flags, bufobj, cres); + res = Native.getInt(cres); + if (res != 0) + throw new CLRuntimeException(res); + return Native.resolve(ce, CLBuffer::new); + } catch (RuntimeException | Error t) { + throw t; + } catch (Throwable t) { + throw new RuntimeException(t); + } + } - public native CLImage createFromGLTexture(CLContext ctx, + public CLImage createFromGLTexture(CLContext ctx, long flags /* flags */, int target /* target */, int miplevel /* miplevel */, - int texture /* texture */); + int texture /* texture */) { + try (Allocator frame = Memory.stack()) { + MemoryAddress cres = frame.alloca(8); + MemoryAddress ce; + int res; + + ce = (MemoryAddress)clCreateFromGLTexture.invokeExact(ctx.addr(), flags, target, miplevel, texture, cres); + res = Native.getInt(cres); + if (res != 0) + throw new CLRuntimeException(res); + return Native.resolve(ce, CLImage::new); + } catch (RuntimeException | Error t) { + throw t; + } catch (Throwable t) { + throw new RuntimeException(t); + } + } - public native CLImage createFromGLRenderbuffer(CLContext cl_context /* context */, + public CLImage createFromGLRenderbuffer(CLContext ctx /* context */, long flags /* flags */, - int renderbuffer /* renderbuffer */); + int renderbuffer /* renderbuffer */) { + try (Allocator frame = Memory.stack()) { + MemoryAddress cres = frame.alloca(8); + MemoryAddress ce; + int res; + + ce = (MemoryAddress)clCreateFromGLRenderbuffer.invokeExact(ctx.addr(), flags, renderbuffer, cres); + res = Native.getInt(cres); + if (res != 0) + throw new CLRuntimeException(res); + return Native.resolve(ce, CLImage::new); + } catch (RuntimeException | Error t) { + throw t; + } catch (Throwable t) { + throw new RuntimeException(t); + } + } public static class GLObjectInfo { @@ -121,43 +195,101 @@ public class GLSharing extends CLExtension { this.gl_object_type = gl_object_type; this.gl_object_name = gl_object_name; } - } - public native GLObjectInfo getGLObjectInfo(CLMemory mem); + public GLObjectInfo getGLObjectInfo(CLMemory mem) { + try (Allocator frame = Memory.stack()) { + MemoryAddress ctype = frame.alloca(8); + MemoryAddress cname = frame.alloca(8); + MemoryAddress ce; + int res; + + res = (int)clGetGLObjectInfo.invokeExact(mem.addr(), ctype, cname); + if (res != 0) + throw new CLRuntimeException(res); + return new GLObjectInfo(Native.getInt(ctype), Native.getInt(cname)); + } catch (RuntimeException | Error t) { + throw t; + } catch (Throwable t) { + throw new RuntimeException(t); + } + } - native int getGLTextureInfoInt(CLMemory mem, int param); + public int getGLTextureInfoInt(CLMemory mem, int param) { + try (Allocator frame = Memory.stack()) { + return Native.getInt(CLObject.getInfo(mem.addr(), CL_GL_TEXTURE_TARGET, clGetGLTextureInfo, frame, 4)); + } + } - public int getGLTExtureTarget(CLMemory mem) { + public int getGLTextureTarget(CLMemory mem) { return getGLTextureInfoInt(mem, CL_GL_TEXTURE_TARGET); } - public int getGLMIPMAPLevel(CLMemory mem) { + public int getGLMIPMapLevel(CLMemory mem) { return getGLTextureInfoInt(mem, CL_GL_MIPMAP_LEVEL); } - public native void enqueueAcquireGLObjects( + public int getGLNumSamples(CLMemory mem) { + return getGLTextureInfoInt(mem, CL_GL_NUM_SAMPLES); + } + + public void enqueueAcquireGLObjects( CLCommandQueue queue /* command_queue */, CLMemory[] mem_objects /* mem_objects */, - CLEventList waiters, - CLEventList events); + CLEventList wait, + CLEventList event) { + try (Allocator frame = Memory.stack()) { + EventInfo info = new EventInfo(frame, wait, event); + MemoryAddress cmem_objects = Native.toAddrV(frame, mem_objects); + int res; + + res = (int)clEnqueueAcquireGLObjects.invokeExact(queue.addr(), mem_objects.length, cmem_objects, + info.nwait, info.wait, info.event); + if (res != 0) + throw new CLRuntimeException(res); + + info.post(event); + } catch (RuntimeException | Error t) { + throw t; + } catch (Throwable t) { + throw new RuntimeException(t); + } + } - public native void enqueueReleaseGLObjects( + public void enqueueReleaseGLObjects( CLCommandQueue queue /* command_queue */, CLMemory[] mem_objects /* mem_objects */, - CLEventList waiters, - CLEventList events); + CLEventList wait, + CLEventList event) { + try (Allocator frame = Memory.stack()) { + EventInfo info = new EventInfo(frame, wait, event); + MemoryAddress cmem_objects = Native.toAddrV(frame, mem_objects); + int res; + + res = (int)clEnqueueReleaseGLObjects.invokeExact(queue.addr(), mem_objects.length, cmem_objects, + info.nwait, info.wait, info.event); + if (res != 0) + throw new CLRuntimeException(res); + + info.post(event); + } catch (RuntimeException | Error t) { + throw t; + } catch (Throwable t) { + throw new RuntimeException(t); + } + } + /* native T getGLContextInfoKHRAny( - CLContextProperty[] properties /* properties */, + CLContextProperty[] properties int ctype, int param_name) throws CLRuntimeException; native T getGLContextInfoKHRAnyV( - CLContextProperty[] properties /* properties */, + CLContextProperty[] properties int ctype, int param_name) throws CLRuntimeException; - + */ /* public CLDevice getCurrendDeviceForGLConextKHR(CLContextProperty[] properties) throws CLRuntimeException { return getGLContextInfoKHRAny(properties, CLObject.CTYPE_DEVICE, CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR); diff --git a/src/notzed.zcl/classes/module-info.java b/src/notzed.zcl/classes/module-info.java index 2f7089f..4f375e6 100644 --- a/src/notzed.zcl/classes/module-info.java +++ b/src/notzed.zcl/classes/module-info.java @@ -24,7 +24,7 @@ module notzed.zcl { requires transitive jdk.incubator.foreign; exports au.notzed.zcl; - //exports au.notzed.zcl.khr; + exports au.notzed.zcl.khr; exports api to notzed.zcl.demo; } diff --git a/src/notzed.zcl/gen/gen.make b/src/notzed.zcl/gen/gen.make index 2be9c5f..8a62902 100644 --- a/src/notzed.zcl/gen/gen.make +++ b/src/notzed.zcl/gen/gen.make @@ -16,7 +16,8 @@ notzed.zcl_generated = \ au/notzed/zcl/Call_pv_v.java notzed.zcl_JAVA_GENERATED = \ - $(notzed.zcl_generated) + $(notzed.zcl_generated) \ + au/notzed/zcl/CLext.java $(notzed.zcl_genjavadir)/au/notzed/zcl/CL.java: src/notzed.zcl/include/CL/cl.h $(export_defines) @install -d $(@D) @@ -40,3 +41,12 @@ $(addprefix $(notzed.zcl_genjavadir)/,$(notzed.zcl_generated)): $(generate_api) --raw-calls \ -c CLLib -lOpenCL --func-file src/notzed.zcl/gen/opencl.txt \ ./$(opencl_pm) + +$(addprefix $(notzed.zcl_genjavadir)/,au/notzed/zcl/CLext.java): $(generate_api) $(opencl_pm) + perl $(generate_api) \ + -d $(notzed.zcl_genjavadir) \ + -t au.notzed.zcl \ + --raw-calls \ + -c CLext --func-file src/notzed.zcl/gen/opencl-ext.txt \ + --no-types \ + ./$(opencl_pm) diff --git a/src/notzed.zcl/gen/generate-api b/src/notzed.zcl/gen/generate-api index 7e2dee0..07e397f 100755 --- a/src/notzed.zcl/gen/generate-api +++ b/src/notzed.zcl/gen/generate-api @@ -7,6 +7,8 @@ # replace a datatype with another, do not generate any code for it # -r name=new +# if no libraries (-l) it outputs method generator factories + @matchStruct = (); $meta = ""; # @classes = ( { name => 'class', match => [ func-pattern, ... ], match_file => [ file, ... ], enum => [ enum-pattern, ... ], enum_file => [ file, ...] } ) @@ -22,6 +24,8 @@ $package = ""; $rawCalls = 0; # calls visited by all output types %usedCalls = (); +# don't output any types +$noTypes = 0; while (@ARGV) { my $cmd = shift(@ARGV); @@ -73,6 +77,8 @@ while (@ARGV) { $output = shift(@ARGV); } elsif ($cmd eq "--enclosing-type") { $enclosingType = shift(@ARGV); + } elsif ($cmd eq "--no-types") { + $noTypes = 1; } else { $meta = $cmd; } @@ -592,8 +598,8 @@ $all = join ('|', keys %roots); if ($all) { push @matchStruct, qr/^($all)$/; } -#print "structures:\n"; -#print Dumper(@matchStruct); +print "structures:\n"; +print Dumper(@matchStruct); # make a map for all callbacks (call: type) to generated names for $c (grep { $_ =~ m/^call:/n } keys %data) { @@ -679,6 +685,9 @@ for $k (findStructs(\%data, @matchStruct)) { my $signature = structSignature(\%struct, ($struct{type} eq "union")); my $name = StudlyCaps($struct{name}); + # yuck: it's just easier here + last if ($noTypes); + if (!$enclosingType) { my $classname = $packagePrefix.$name; @@ -861,7 +870,7 @@ END print $dst "\tpublic static MemoryLayout layout() { return Native.parseUnion(\"$signature\"); }\n"; } else { print $dst "\tpublic static MemoryLayout layout() { return Native.parseStruct(\"$signature\"); }\n"; - } + } print $dst "}\n"; @@ -896,40 +905,46 @@ for $c (@classes) { import jdk.incubator.foreign.*; import java.lang.invoke.MethodHandle; import api.Native; +import java.util.function.Function; $importPoineter END } - print $dst "class $class{name} {\n"; + print $dst "public class $class{name} {\n"; - print $dst "\tstatic final String[] libraries = {"; - print $dst join(",", map { "\"$_\"" } @libs); - print $dst "};\n"; + if (@libs) { + print $dst "\tstatic final String[] libraries = {"; + print $dst join(",", map { "\"$_\"" } @libs); + print $dst "};\n"; + } # enums to ints # TODO: interfaces? # TODO: static lib class? # typedef enums might appear twice in the data, so ignore duplicates # also, some api's have multiple definitions (?) - my %visited = (); - my @match_enum = @{$class{enum}}; - for $k (sort(findDefinition(\%data, 'enum', @match_enum))) { - my %enum = %{$data{$k}}; - my @values = @{$enum{values}}; - my $type = "int"; - - if ($enum{value_type} =~ m/^[ui](\d+)/) { - $type = "long" if ($1 > 32) - } + if (!$noTypes) { + my %visited = (); + my @match_enum = @{$class{enum}}; - print $dst "\n\t// enum $enum{name}\n"; - for $vi (@values) { - my %value = %{$vi}; + for $k (sort(findDefinition(\%data, 'enum', @match_enum))) { + my %enum = %{$data{$k}}; + my @values = @{$enum{values}}; + my $type = "int"; - if (!$visited{$value{label}}) { - #print $dst "\tpublic static final $type $value{label} = ($type)$value{value};\n"; - print $dst "\tpublic static final $type $value{label} = $value{value};\n"; - $visited{$value{label}} = 1; + if ($enum{value_type} =~ m/^[ui](\d+)/) { + $type = "long" if ($1 > 32) + } + + print $dst "\n\t// enum $enum{name}\n"; + for $vi (@values) { + my %value = %{$vi}; + + if (!$visited{$value{label}}) { + #print $dst "\tpublic static final $type $value{label} = ($type)$value{value};\n"; + print $dst "\tpublic static final $type $value{label} = $value{value};\n"; + $visited{$value{label}} = 1; + } } } } @@ -952,34 +967,32 @@ END } } - # function handles - #print "class $class{name} -> match:\n".Dumper(\@match); - - for $k (sort(findDefinition(\%data, 'func', @match))) { - my %func = %{$data{$k}}; - my @params = @{$func{arguments}}; - my $signature = funcSignature(\%func); - my $name = ($func{name}); + if (@libs) { + # function handles + for $k (sort(findDefinition(\%data, 'func', @match))) { + my %func = %{$data{$k}}; + my @params = @{$func{arguments}}; + my $signature = funcSignature(\%func); + my $name = ($func{name}); - print $dst "\tfinal static MethodHandle $name;\n"; - } + print $dst "\tpublic final static MethodHandle $name;\n"; + } - # function handle init - print $dst "\tstatic {\n"; - print $dst "\t\tLibraryLookup[] libs = Native.loadLibraries(libraries);\n"; + # function handle init + print $dst "\tstatic {\n"; + print $dst "\t\tLibraryLookup[] libs = Native.loadLibraries(libraries);\n"; - for $k (sort(findDefinition(\%data, 'func', @match))) { - my %func = %{$data{$k}}; - my @params = @{$func{arguments}}; - my $signature = funcSignature(\%func); - my $name = ($func{name}); + for $k (sort(findDefinition(\%data, 'func', @match))) { + my %func = %{$data{$k}}; + my @params = @{$func{arguments}}; + my $signature = funcSignature(\%func); + my $name = ($func{name}); - print $dst "\t\t$name = Native.downcallHandle(libs, \"$name\", \"$signature\");\n"; - } - print $dst "\t}\n"; + print $dst "\t\t$name = Native.downcallHandle(libs, \"$name\", \"$signature\");\n"; + } + print $dst "\t}\n"; - # function handle invocation - if ($rawCalls) { + # function handle invocation for $k (sort(findDefinition(\%data, 'func', @match))) { my %func = %{$data{$k}}; my @params = @{$func{arguments}}; @@ -1012,71 +1025,20 @@ END print $dst ");\n"; print $dst "\t}\n\n"; } - print $dst "}\n"; } else { + # function handle factories for $k (sort(findDefinition(\%data, 'func', @match))) { my %func = %{$data{$k}}; my @params = @{$func{arguments}}; my $signature = funcSignature(\%func); my $name = ($func{name}); - my %res = %{$func{result}}; - my $result = typeToJava(\%{$func{result}}); - - print $dst "\tpublic static $result $name("; - - for $pi (@params) { - my %param = %{$pi}; - my $type = typeToJava($pi); - - $type =~ s/Callback/Pointer/; - # HACK - $type =~ s/Pointer/Pointer/; - - print $dst "$type $param{name}"; - print $dst ", " if ($pi != $params[$#params]); - } - - print $dst ") {\n"; - # see also call below - print $dst "\t\ttry {\n"; - print $dst "\t\t\t"; - if ($res{type} =~ m/(struct|union)/n) { - if ($res{deref}) { - print $dst "MemoryAddress add = (MemoryAddress)"; - } else { - print $dst "MemorySegment seg = (MemorySegment)"; - } - } elsif ($result ne "void") { - print $dst "return ($result)"; - } - print $dst "$name.invokeExact("; - for $pi (@params) { - my %param = %{$pi}; - - print $dst "$param{name}"; - if ($param{deref}) { - print $dst ".addr()"; - } elsif ($param{type} =~ m/^struct|union/) { - print $dst ".addr().segment()"; - } - print $dst ", " if ($pi != $params[$#params]); - } - print $dst ");\n"; - if ($res{type} =~ m/(struct|union)/n) { - if ($res{deref}) { - print $dst "\t\t\treturn $result.create(add);\n"; - } else { - print $dst "\t\t\treturn $result.create(seg.baseAddress());\n"; - } - } - print $dst "\t\t}\n"; - print $dst "\t\tcatch (Throwable t) { throw new RuntimeException(t); }\n"; - print $dst "\t}\n\n"; + print $dst "\tpublic static MethodHandle $name(Function addr) {\n"; + print $dst "\t\treturn Native.downcallHandle(addr.apply(\"$name\"), \"$signature\");\n"; + print $dst "\t}\n"; } - - print $dst "}\n"; } + print $dst "}\n"; if (!$enclosingType) { close($dst); diff --git a/src/notzed.zcl/gen/opencl-ext.txt b/src/notzed.zcl/gen/opencl-ext.txt new file mode 100644 index 0000000..f4d0c41 --- /dev/null +++ b/src/notzed.zcl/gen/opencl-ext.txt @@ -0,0 +1,11 @@ +clCreateFromGLBuffer +clCreateFromGLTexture +clCreateFromGLRenderbuffer +clGetGLObjectInfo +clGetGLTextureInfo +clEnqueueAcquireGLObjects +clEnqueueReleaseGLObjects +clCreateFromGLTexture2D +clCreateFromGLTexture3D +clGetGLContextInfoKHR +clCreateEventFromGLsyncKHR diff --git a/src/notzed.zcl/gen/opencl.pm b/src/notzed.zcl/gen/opencl.pm index 2908010..7bd45eb 100644 --- a/src/notzed.zcl/gen/opencl.pm +++ b/src/notzed.zcl/gen/opencl.pm @@ -1,7 +1,3 @@ - -# Note that this was generated by a gcc plugin -# it grabs a lot of junk that just gets ignored by the generator - %data = ( 'func:__ctype_get_mb_cur_max' => { name => '__ctype_get_mb_cur_max', type => 'func', result => { ctype => 'long unsigned int', type => 'u64', }, @@ -1680,6 +1676,113 @@ { size => 64, name => 'arg_3', deref => 'u64:u64:${_cl_event}', type => 'struct:_cl_event',}, { size => 64, name => 'arg_4', deref => 'u64:u64:${_cl_event}', type => 'struct:_cl_event',}, ]}, +'func:clCreateFromGLBuffer' => { name => 'clCreateFromGLBuffer', type => 'func', + result => { deref => 'u64:${_cl_mem}', type => 'struct:_cl_mem', }, + arguments => [ + { size => 64, name => 'arg_0', deref => 'u64:${_cl_context}', type => 'struct:_cl_context',}, + { size => 64, name => 'arg_1', ctype => 'long unsigned int', type => 'u64',}, + { size => 32, name => 'arg_2', ctype => 'unsigned int', type => 'u32',}, + { size => 64, name => 'arg_3', deref => 'u64:i32', ctype => 'int', type => 'i32',}, +]}, +'func:clCreateFromGLTexture' => { name => 'clCreateFromGLTexture', type => 'func', + result => { deref => 'u64:${_cl_mem}', type => 'struct:_cl_mem', }, + arguments => [ + { size => 64, name => 'arg_0', deref => 'u64:${_cl_context}', type => 'struct:_cl_context',}, + { size => 64, name => 'arg_1', ctype => 'long unsigned int', type => 'u64',}, + { size => 32, name => 'arg_2', ctype => 'unsigned int', type => 'u32',}, + { size => 32, name => 'arg_3', ctype => 'int', type => 'i32',}, + { size => 32, name => 'arg_4', ctype => 'unsigned int', type => 'u32',}, + { size => 64, name => 'arg_5', deref => 'u64:i32', ctype => 'int', type => 'i32',}, +]}, +'func:clCreateFromGLRenderbuffer' => { name => 'clCreateFromGLRenderbuffer', type => 'func', + result => { deref => 'u64:${_cl_mem}', type => 'struct:_cl_mem', }, + arguments => [ + { size => 64, name => 'arg_0', deref => 'u64:${_cl_context}', type => 'struct:_cl_context',}, + { size => 64, name => 'arg_1', ctype => 'long unsigned int', type => 'u64',}, + { size => 32, name => 'arg_2', ctype => 'unsigned int', type => 'u32',}, + { size => 64, name => 'arg_3', deref => 'u64:i32', ctype => 'int', type => 'i32',}, +]}, +'func:clCreateEventFromGLsyncKHR' => { name => 'clCreateEventFromGLsyncKHR', type => 'func', + result => { deref => 'u64:${_cl_event}', type => 'struct:_cl_event', }, + arguments => [ + { size => 64, name => 'arg_0', deref => 'u64:${_cl_context}', type => 'struct:_cl_context',}, + { size => 64, name => 'arg_1', deref => 'u64:${__GLsync}', type => 'struct:__GLsync',}, + { size => 64, name => 'arg_2', deref => 'u64:i32', ctype => 'int', type => 'i32',}, +]}, +'func:clGetGLObjectInfo' => { name => 'clGetGLObjectInfo', type => 'func', + result => { ctype => 'int', type => 'i32', }, + arguments => [ + { size => 64, name => 'arg_0', deref => 'u64:${_cl_mem}', type => 'struct:_cl_mem',}, + { size => 64, name => 'arg_1', deref => 'u64:u32', ctype => 'unsigned int', type => 'u32',}, + { size => 64, name => 'arg_2', deref => 'u64:u32', ctype => 'unsigned int', type => 'u32',}, +]}, +'func:clGetGLTextureInfo' => { name => 'clGetGLTextureInfo', type => 'func', + result => { ctype => 'int', type => 'i32', }, + arguments => [ + { size => 64, name => 'arg_0', deref => 'u64:${_cl_mem}', type => 'struct:_cl_mem',}, + { size => 32, name => 'arg_1', ctype => 'unsigned int', type => 'u32',}, + { size => 64, name => 'arg_2', ctype => 'long unsigned int', type => 'u64',}, + { size => 64, name => 'arg_3', deref => 'u64:v', type => 'void',}, + { size => 64, name => 'arg_4', deref => 'u64:u64', ctype => 'long unsigned int', type => 'u64',}, +]}, +'func:clEnqueueAcquireGLObjects' => { name => 'clEnqueueAcquireGLObjects', type => 'func', + result => { ctype => 'int', type => 'i32', }, + arguments => [ + { size => 64, name => 'arg_0', deref => 'u64:${_cl_command_queue}', type => 'struct:_cl_command_queue',}, + { size => 32, name => 'arg_1', ctype => 'unsigned int', type => 'u32',}, + { size => 64, name => 'arg_2', deref => 'u64:u64:${_cl_mem}', type => 'struct:_cl_mem',}, + { size => 32, name => 'arg_3', ctype => 'unsigned int', type => 'u32',}, + { size => 64, name => 'arg_4', deref => 'u64:u64:${_cl_event}', type => 'struct:_cl_event',}, + { size => 64, name => 'arg_5', deref => 'u64:u64:${_cl_event}', type => 'struct:_cl_event',}, +]}, +'func:clEnqueueReleaseGLObjects' => { name => 'clEnqueueReleaseGLObjects', type => 'func', + result => { ctype => 'int', type => 'i32', }, + arguments => [ + { size => 64, name => 'arg_0', deref => 'u64:${_cl_command_queue}', type => 'struct:_cl_command_queue',}, + { size => 32, name => 'arg_1', ctype => 'unsigned int', type => 'u32',}, + { size => 64, name => 'arg_2', deref => 'u64:u64:${_cl_mem}', type => 'struct:_cl_mem',}, + { size => 32, name => 'arg_3', ctype => 'unsigned int', type => 'u32',}, + { size => 64, name => 'arg_4', deref => 'u64:u64:${_cl_event}', type => 'struct:_cl_event',}, + { size => 64, name => 'arg_5', deref => 'u64:u64:${_cl_event}', type => 'struct:_cl_event',}, +]}, +'func:clCreateFromGLTexture2D' => { name => 'clCreateFromGLTexture2D', type => 'func', + result => { deref => 'u64:${_cl_mem}', type => 'struct:_cl_mem', }, + arguments => [ + { size => 64, name => 'arg_0', deref => 'u64:${_cl_context}', type => 'struct:_cl_context',}, + { size => 64, name => 'arg_1', ctype => 'long unsigned int', type => 'u64',}, + { size => 32, name => 'arg_2', ctype => 'unsigned int', type => 'u32',}, + { size => 32, name => 'arg_3', ctype => 'int', type => 'i32',}, + { size => 32, name => 'arg_4', ctype => 'unsigned int', type => 'u32',}, + { size => 64, name => 'arg_5', deref => 'u64:i32', ctype => 'int', type => 'i32',}, +]}, +'func:clCreateFromGLTexture3D' => { name => 'clCreateFromGLTexture3D', type => 'func', + result => { deref => 'u64:${_cl_mem}', type => 'struct:_cl_mem', }, + arguments => [ + { size => 64, name => 'arg_0', deref => 'u64:${_cl_context}', type => 'struct:_cl_context',}, + { size => 64, name => 'arg_1', ctype => 'long unsigned int', type => 'u64',}, + { size => 32, name => 'arg_2', ctype => 'unsigned int', type => 'u32',}, + { size => 32, name => 'arg_3', ctype => 'int', type => 'i32',}, + { size => 32, name => 'arg_4', ctype => 'unsigned int', type => 'u32',}, + { size => 64, name => 'arg_5', deref => 'u64:i32', ctype => 'int', type => 'i32',}, +]}, +'func:clGetGLContextInfoKHR' => { name => 'clGetGLContextInfoKHR', type => 'func', + result => { ctype => 'int', type => 'i32', }, + arguments => [ + { size => 64, name => 'arg_0', deref => 'u64:i64', ctype => 'long int', type => 'i64',}, + { size => 32, name => 'arg_1', ctype => 'unsigned int', type => 'u32',}, + { size => 64, name => 'arg_2', ctype => 'long unsigned int', type => 'u64',}, + { size => 64, name => 'arg_3', deref => 'u64:v', type => 'void',}, + { size => 64, name => 'arg_4', deref => 'u64:u64', ctype => 'long unsigned int', type => 'u64',}, +]}, +'call:clGetGLContextInfoKHR_fn' => { name => 'clGetGLContextInfoKHR_fn', deref => 'u64:(u64:i64u32u64u64:vu64:u64)i32', + result => { ctype => 'int', type => 'i32', }, + arguments => [ + { size => 64, name => 'arg_0', deref => 'u64:i64', ctype => 'long int', type => 'i64',}, + { size => 32, name => 'arg_1', ctype => 'unsigned int', type => 'u32',}, + { size => 64, name => 'arg_2', ctype => 'long unsigned int', type => 'u64',}, + { size => 64, name => 'arg_3', deref => 'u64:v', type => 'void',}, + { size => 64, name => 'arg_4', deref => 'u64:u64', ctype => 'long unsigned int', type => 'u64',}, +]}, 'struct:__pthread_cond_s___wseq32' => { name => '__pthread_cond_s___wseq32', type => 'struct', size => 64, fields => [ { name => '__low', size => 32, offset => 0, ctype => 'unsigned int', type => 'u32',}, { name => '__high', size => 32, offset => 32, ctype => 'unsigned int', type => 'u32',}, @@ -2287,6 +2390,22 @@ # clCreateCommandQueue # clCreateSampler # clEnqueueTask +# cl_gl_object_type +# cl_gl_texture_info +# cl_gl_platform_info +# cl_GLsync +# clCreateFromGLBuffer +# clCreateFromGLTexture +# clCreateFromGLRenderbuffer +# clGetGLObjectInfo +# clGetGLTextureInfo +# clEnqueueAcquireGLObjects +# clEnqueueReleaseGLObjects +# clCreateFromGLTexture2D +# clCreateFromGLTexture3D +# cl_gl_context_info +# clGetGLContextInfoKHR +# clGetGLContextInfoKHR_fn # __pthread_cond_s___wseq32 # __pthread_cond_s___g1_start32 # ()v -- 2.39.5