Started a panama (no jni) based implementation.
authorNot Zed <notzed@gmail.com>
Fri, 24 Jan 2020 02:48:05 +0000 (13:18 +1030)
committerNot Zed <notzed@gmail.com>
Fri, 24 Jan 2020 02:48:05 +0000 (13:18 +1030)
63 files changed:
README
config.make.in
java.make
nbproject/build-impl.xml
nbproject/genfiles.properties
nbproject/project.properties
nbproject/project.xml
src/notzed.zcl/classes/api/Allocator.java [moved from src/notzed.zcl/jni/zcl-init-dll.c with 54% similarity]
src/notzed.zcl/classes/api/Callback.java [new file with mode: 0644]
src/notzed.zcl/classes/api/Memory.java [new file with mode: 0644]
src/notzed.zcl/classes/api/Native.java [new file with mode: 0644]
src/notzed.zcl/classes/au/notzed/zcl/CLBuffer.java
src/notzed.zcl/classes/au/notzed/zcl/CLBufferInfo.java
src/notzed.zcl/classes/au/notzed/zcl/CLCommandQueue.java
src/notzed.zcl/classes/au/notzed/zcl/CLContext.java
src/notzed.zcl/classes/au/notzed/zcl/CLContextNotify.java
src/notzed.zcl/classes/au/notzed/zcl/CLDevice.java
src/notzed.zcl/classes/au/notzed/zcl/CLDeviceProperty.java
src/notzed.zcl/classes/au/notzed/zcl/CLEvent.java
src/notzed.zcl/classes/au/notzed/zcl/CLEventList.java
src/notzed.zcl/classes/au/notzed/zcl/CLEventNotify.java
src/notzed.zcl/classes/au/notzed/zcl/CLException.java
src/notzed.zcl/classes/au/notzed/zcl/CLExtendable.java
src/notzed.zcl/classes/au/notzed/zcl/CLExtension.java
src/notzed.zcl/classes/au/notzed/zcl/CLImage.java
src/notzed.zcl/classes/au/notzed/zcl/CLImageDesc.java
src/notzed.zcl/classes/au/notzed/zcl/CLImageFormat.java
src/notzed.zcl/classes/au/notzed/zcl/CLKernel.java
src/notzed.zcl/classes/au/notzed/zcl/CLMemory.java
src/notzed.zcl/classes/au/notzed/zcl/CLNotify.java
src/notzed.zcl/classes/au/notzed/zcl/CLObject.java
src/notzed.zcl/classes/au/notzed/zcl/CLPipe.java
src/notzed.zcl/classes/au/notzed/zcl/CLPlatform.java
src/notzed.zcl/classes/au/notzed/zcl/CLProgram.java
src/notzed.zcl/classes/au/notzed/zcl/CLProperty.java
src/notzed.zcl/classes/au/notzed/zcl/CLRuntimeException.java
src/notzed.zcl/classes/au/notzed/zcl/CLSampler.java
src/notzed.zcl/classes/au/notzed/zcl/khr/GLEvent.java
src/notzed.zcl/classes/au/notzed/zcl/khr/GLSharing.java
src/notzed.zcl/classes/module-info.java
src/notzed.zcl/gen/export-defines [new file with mode: 0755]
src/notzed.zcl/gen/gen.make [new file with mode: 0644]
src/notzed.zcl/gen/generate-api [new file with mode: 0755]
src/notzed.zcl/gen/opencl.pm [new file with mode: 0644]
src/notzed.zcl/gen/opencl.txt [new file with mode: 0644]
src/notzed.zcl/include/CL/cl.h [moved from src/notzed.zcl/jni/include/CL/cl.h with 100% similarity]
src/notzed.zcl/include/CL/cl_d3d10.h [moved from src/notzed.zcl/jni/include/CL/cl_d3d10.h with 100% similarity]
src/notzed.zcl/include/CL/cl_d3d11.h [moved from src/notzed.zcl/jni/include/CL/cl_d3d11.h with 100% similarity]
src/notzed.zcl/include/CL/cl_dx9_media_sharing.h [moved from src/notzed.zcl/jni/include/CL/cl_dx9_media_sharing.h with 100% similarity]
src/notzed.zcl/include/CL/cl_egl.h [moved from src/notzed.zcl/jni/include/CL/cl_egl.h with 100% similarity]
src/notzed.zcl/include/CL/cl_ext.h [moved from src/notzed.zcl/jni/include/CL/cl_ext.h with 100% similarity]
src/notzed.zcl/include/CL/cl_gl.h [moved from src/notzed.zcl/jni/include/CL/cl_gl.h with 100% similarity]
src/notzed.zcl/include/CL/cl_gl_ext.h [moved from src/notzed.zcl/jni/include/CL/cl_gl_ext.h with 100% similarity]
src/notzed.zcl/include/CL/cl_platform.h [moved from src/notzed.zcl/jni/include/CL/cl_platform.h with 100% similarity]
src/notzed.zcl/include/CL/opencl.h [moved from src/notzed.zcl/jni/include/CL/opencl.h with 100% similarity]
src/notzed.zcl/jni/jni.make [deleted file]
src/notzed.zcl/jni/zcl-extension.h [deleted file]
src/notzed.zcl/jni/zcl-generate [deleted file]
src/notzed.zcl/jni/zcl-init-so.c [deleted file]
src/notzed.zcl/jni/zcl-jni.c [deleted file]
src/notzed.zcl/jni/zcl-jni.def [deleted file]
src/notzed.zcl/jni/zcl-khr-gl-event.c [deleted file]
src/notzed.zcl/jni/zcl-khr-gl-sharing.c [deleted file]

diff --git a/README b/README
index a51ca98..591c71f 100644 (file)
--- a/README
+++ b/README
@@ -1,4 +1,14 @@
 
+foreign-abi branch note
+-----------------------
+
+This is still work in progress, expect breakage, out of date doco,
+other issues.
+
+Various native support code is in the temporary package api.*,
+eventually to be moved to a new nativez.
+
+
 INTRODUCTION
 ------------
 
index f0186c5..dba9824 100644 (file)
@@ -3,15 +3,12 @@ TARGET ?= linux-amd64
 
 JAVA_HOME ?= /usr/local/jdk-13+33
 
-# See also JAVACFLAGS --module-path
-NATIVEZ_HOME=../nativez/bin/$(TARGET)
-
-JAVAMODPATH = $(NATIVEZ_HOME)/lib
 JAVACFLAGS += -source 13
+JAVACFLAGS += --add-exports jdk.incubator.foreign/jdk.incubator.foreign.unsafe=notzed.zcl
 
-JAVAC ?= javac
-JAR ?= jar
-JMOD ?= jmod
+JAVAC ?= $(JAVA_HOME)/bin/javac
+JAR ?= $(JAVA_HOME)/bin/jar
+JMOD ?= $(JAVA_HOME)/bin/jmod
 
 # Linux options
 linux-amd64_CPPFLAGS = \
index ee876b1..01df66d 100644 (file)
--- a/java.make
+++ b/java.make
@@ -347,4 +347,3 @@ dist:
         --transform=s,^,$(dist_NAME)-$(dist_VERSION)/, \
         config.make java.make Makefile src             \
         $(dist_EXTRA)
-
index 489097b..9cc58e9 100644 (file)
@@ -69,7 +69,43 @@ is divided into following sections:
         <property name="dist.jlink.output" value="${dist.jlink.dir}/${application.title}"/>
     </target>
     <target depends="-pre-init,-init-private,-init-user,-init-project,-init-macrodef-property" name="-do-init">
-        <property name="platform.java" value="${java.home}/bin/java"/>
+        <j2semodularproject1:property xmlns:j2semodularproject1="http://www.netbeans.org/ns/j2se-modular-project/1" name="platform.home" value="platforms.${platform.active}.home"/>
+        <j2semodularproject1:property xmlns:j2semodularproject1="http://www.netbeans.org/ns/j2se-modular-project/1" name="platform.bootcp" value="platforms.${platform.active}.bootclasspath"/>
+        <j2semodularproject1:property xmlns:j2semodularproject1="http://www.netbeans.org/ns/j2se-modular-project/1" name="platform.compiler" value="platforms.${platform.active}.compile"/>
+        <j2semodularproject1:property xmlns:j2semodularproject1="http://www.netbeans.org/ns/j2se-modular-project/1" name="platform.javac.tmp" value="platforms.${platform.active}.javac"/>
+        <condition property="platform.javac" value="${platform.home}/bin/javac">
+            <equals arg1="${platform.javac.tmp}" arg2="$${platforms.${platform.active}.javac}"/>
+        </condition>
+        <property name="platform.javac" value="${platform.javac.tmp}"/>
+        <j2semodularproject1:property xmlns:j2semodularproject1="http://www.netbeans.org/ns/j2se-modular-project/1" name="platform.java.tmp" value="platforms.${platform.active}.java"/>
+        <condition property="platform.java" value="${platform.home}/bin/java">
+            <equals arg1="${platform.java.tmp}" arg2="$${platforms.${platform.active}.java}"/>
+        </condition>
+        <property name="platform.java" value="${platform.java.tmp}"/>
+        <j2semodularproject1:property xmlns:j2semodularproject1="http://www.netbeans.org/ns/j2se-modular-project/1" name="platform.javadoc.tmp" value="platforms.${platform.active}.javadoc"/>
+        <condition property="platform.javadoc" value="${platform.home}/bin/javadoc">
+            <equals arg1="${platform.javadoc.tmp}" arg2="$${platforms.${platform.active}.javadoc}"/>
+        </condition>
+        <property name="platform.javadoc" value="${platform.javadoc.tmp}"/>
+        <condition property="platform.invalid" value="true">
+            <or>
+                <contains string="${platform.javac}" substring="$${platforms."/>
+                <contains string="${platform.java}" substring="$${platforms."/>
+                <contains string="${platform.javadoc}" substring="$${platforms."/>
+            </or>
+        </condition>
+        <fail unless="platform.home">Must set platform.home</fail>
+        <fail unless="platform.bootcp">Must set platform.bootcp</fail>
+        <fail unless="platform.java">Must set platform.java</fail>
+        <fail unless="platform.javac">Must set platform.javac</fail>
+        <fail if="platform.invalid">
+ The J2SE Platform is not correctly set up.
+ Your active platform is: ${platform.active}, but the corresponding property "platforms.${platform.active}.home" is not found in the project's properties files. 
+ Either open the project in the IDE and setup the Platform with the same name or add it manually.
+ For example like this:
+     ant -Duser.properties.file=&lt;path_to_property_file&gt; jar (where you put the property "platforms.${platform.active}.home" in a .properties file)
+  or ant -Dplatforms.${platform.active}.home=&lt;path_to_JDK_home&gt; jar (where no properties file is used) 
+  </fail>
         <j2semodularproject1:modsource_regexp xmlns:j2semodularproject1="http://www.netbeans.org/ns/j2se-modular-project/1" modsource="${test.src.dir.path}" property="have.tests.test.src.dir.regexp"/>
         <dirset dir="${basedir}/${test.src.dir}" id="have.tests.test.src.dir.set" includes="*/*">
             <filename regex="${have.tests.test.src.dir.regexp}"/>
@@ -94,11 +130,16 @@ is divided into following sections:
         <union id="have.tests.patchset">
             <dirset refid="have.tests.test.src.dir.patchset"/>
         </union>
+        <j2semodularproject1:modsource_regexp xmlns:j2semodularproject1="http://www.netbeans.org/ns/j2se-modular-project/1" modsource="${src.gen2.dir.path}" property="have.sources.src.gen2.dir.regexp"/>
+        <dirset dir="${basedir}/${src.gen2.dir}" id="have.sources.src.gen2.dir.set" includes="*/*">
+            <filename regex="${have.sources.src.gen2.dir.regexp}"/>
+        </dirset>
         <j2semodularproject1:modsource_regexp xmlns:j2semodularproject1="http://www.netbeans.org/ns/j2se-modular-project/1" modsource="${src.dir.path}" property="have.sources.src.dir.regexp"/>
         <dirset dir="${basedir}/${src.dir}" id="have.sources.src.dir.set" includes="*/*">
             <filename regex="${have.sources.src.dir.regexp}"/>
         </dirset>
         <union id="have.sources.set">
+            <dirset refid="have.sources.src.gen2.dir.set"/>
             <dirset refid="have.sources.src.dir.set"/>
         </union>
         <condition property="have.sources">
@@ -106,6 +147,9 @@ is divided into following sections:
                 <resourcecount count="0" when="greater">
                     <union refid="have.sources.set"/>
                 </resourcecount>
+                <resourcecount count="0" when="greater">
+                    <union refid="have.sources.set"/>
+                </resourcecount>
             </or>
         </condition>
         <condition property="main.class.available">
@@ -177,20 +221,6 @@ is divided into following sections:
         <condition else="" property="javac.profile.cmd.line.arg" value="-profile ${javac.profile}">
             <isset property="profile.available"/>
         </condition>
-        <condition else="false" property="jdkBug6558476">
-            <and>
-                <matches pattern="1\.[56]" string="${java.specification.version}"/>
-                <not>
-                    <os family="unix"/>
-                </not>
-            </and>
-        </condition>
-        <condition else="false" property="javac.fork">
-            <or>
-                <istrue value="${jdkBug6558476}"/>
-                <istrue value="${javac.external.vm}"/>
-            </or>
-        </condition>
         <condition property="main.class.check.available">
             <and>
                 <isset property="libs.CopyLibs.classpath"/>
@@ -262,6 +292,7 @@ is divided into following sections:
         <!-- You can override this target in the ../build.xml file. -->
     </target>
     <target depends="-pre-init,-init-private,-init-user,-init-project,-do-init" name="-init-check">
+        <fail unless="src.gen2.dir">Must set src.gen2.dir</fail>
         <fail unless="src.dir">Must set src.dir</fail>
         <fail unless="test.src.dir">Must set test.src.dir</fail>
         <fail unless="build.dir">Must set build.dir</fail>
@@ -293,7 +324,7 @@ is divided into following sections:
             <attribute default="${build.modules.dir}" name="destdir"/>
             <attribute default="${javac.classpath}" name="classpath"/>
             <attribute default="${javac.modulepath}" name="modulepath"/>
-            <attribute default="${src.dir}/*/${src.dir.path}" name="modulesourcepath"/>
+            <attribute default="${src.gen2.dir}/*/${src.gen2.dir.path}:${src.dir}/*/${src.dir.path}" name="modulesourcepath"/>
             <attribute default="${javac.upgrademodulepath}" name="upgrademodulepath"/>
             <attribute default="${javac.processorpath}" name="processorpath"/>
             <attribute default="${javac.processormodulepath}" name="processormodulepath"/>
@@ -314,7 +345,7 @@ is divided into following sections:
                         </path>
                     </resourcecount>
                 </condition>
-                <javac debug="@{debug}" deprecation="${javac.deprecation}" destdir="@{destdir}" encoding="${source.encoding}" excludes="@{excludes}" fork="${javac.fork}" includeantruntime="false" includes="@{includes}" source="${javac.source}" target="${javac.target}" tempdir="${java.io.tmpdir}">
+                <javac debug="@{debug}" deprecation="${javac.deprecation}" destdir="@{destdir}" encoding="${source.encoding}" excludes="@{excludes}" executable="${platform.javac}" fork="yes" includeantruntime="false" includes="@{includes}" source="${javac.source}" target="${javac.target}" tempdir="${java.io.tmpdir}">
                     <classpath>
                         <path path="@{classpath}"/>
                     </classpath>
@@ -346,7 +377,7 @@ is divided into following sections:
     </target>
     <target depends="-init-macrodef-javac" name="-init-macrodef-javac-depend">
         <macrodef name="depend" uri="http://www.netbeans.org/ns/j2se-modular-project/1">
-            <attribute default="${src.dir}" name="srcdir"/>
+            <attribute default="${src.gen2.dir}:${src.dir}" name="srcdir"/>
             <attribute default="${build.classes.dir}" name="destdir"/>
             <attribute default="${javac.classpath}" name="classpath"/>
             <sequential>
@@ -407,7 +438,7 @@ is divided into following sections:
             <sequential>
                 <property location="${build.dir}/empty" name="empty.dir"/>
                 <property name="junit.forkmode" value="perTest"/>
-                <junit dir="${work.dir}" errorproperty="tests.failed" failureproperty="tests.failed" fork="true" forkmode="${junit.forkmode}" showoutput="true" tempdir="${build.dir}">
+                <junit dir="${work.dir}" errorproperty="tests.failed" failureproperty="tests.failed" fork="true" forkmode="${junit.forkmode}" jvm="${platform.java}" showoutput="true" tempdir="${build.dir}">
                     <syspropertyset>
                         <propertyref prefix="test-sys-prop."/>
                         <mapper from="test-sys-prop.*" to="*" type="glob"/>
@@ -492,7 +523,7 @@ is divided into following sections:
                     </fileset>
                 </union>
                 <taskdef classname="org.testng.TestNGAntTask" classpath="${run.test.classpath}" name="testng"/>
-                <testng classfilesetref="test.set" failureProperty="tests.failed" listeners="org.testng.reporters.VerboseReporter" methods="${testng.methods.arg}" mode="${testng.mode}" outputdir="${build.test.results.dir}" suitename="notzed.zcl" testname="TestNG tests" workingDir="${work.dir}">
+                <testng classfilesetref="test.set" failureProperty="tests.failed" jvm="${platform.java}" listeners="org.testng.reporters.VerboseReporter" methods="${testng.methods.arg}" mode="${testng.mode}" outputdir="${build.test.results.dir}" suitename="notzed.zcl" testname="TestNG tests" workingDir="${work.dir}">
                     <xmlfileset dir="${build.test.classes.dir}" includes="@{testincludes}"/>
                     <propertyset>
                         <propertyref prefix="test-sys-prop."/>
@@ -723,6 +754,9 @@ is divided into following sections:
                     <classpath>
                         <path path="@{classpath}"/>
                     </classpath>
+                    <bootclasspath>
+                        <path path="${platform.bootcp}"/>
+                    </bootclasspath>
                 </nbjpdastart>
             </sequential>
         </macrodef>
@@ -772,7 +806,7 @@ is divided into following sections:
             <attribute default="jvm" name="jvm"/>
             <element name="customize" optional="true"/>
             <sequential>
-                <java classname="@{classname}" dir="${work.dir}" failonerror="${java.failonerror}" fork="true" module="@{modulename}">
+                <java classname="@{classname}" dir="${work.dir}" failonerror="${java.failonerror}" fork="true" jvm="${platform.java}" module="@{modulename}">
                     <classpath>
                         <path path="@{classpath}"/>
                     </classpath>
@@ -842,15 +876,6 @@ is divided into following sections:
         <propertyfile file="${built-jar.properties}">
             <entry key="${basedir}" value=""/>
         </propertyfile>
-        <antcall target="-maybe-call-dep">
-            <param name="call.built.properties" value="${built-jar.properties}"/>
-            <param location="${project.notzed_nativez}" name="call.subproject"/>
-            <param location="${project.notzed_nativez}/build.xml" name="call.script"/>
-            <param name="call.target" value="jar"/>
-            <param name="transfer.built-jar.properties" value="${built-jar.properties}"/>
-            <param name="transfer.not.archive.disabled" value="true"/>
-            <param name="transfer.do.jlink" value="false"/>
-        </antcall>
     </target>
     <target depends="init,-check-automatic-build,-clean-after-automatic-build" name="-verify-automatic-build"/>
     <target depends="init" name="-check-automatic-build">
@@ -999,10 +1024,16 @@ is divided into following sections:
                 <include name="*"/>
             </dirset>
         </pathconvert>
-        <j2semodularproject1:depend xmlns:j2semodularproject1="http://www.netbeans.org/ns/j2se-modular-project/1" srcdir="${src.dir}:${build.generated.subdirs}"/>
+        <j2semodularproject1:depend xmlns:j2semodularproject1="http://www.netbeans.org/ns/j2se-modular-project/1" srcdir="${src.gen2.dir}:${src.dir}:${build.generated.subdirs}"/>
     </target>
     <target depends="init,deps-jar,-pre-pre-compile,-pre-compile,-compile-depend" if="have.sources" name="-do-compile">
         <j2semodularproject1:javac xmlns:j2semodularproject1="http://www.netbeans.org/ns/j2se-modular-project/1" gensrcdir="${build.generated.sources.dir}"/>
+        <j2semodularproject1:modsource_regexp xmlns:j2semodularproject1="http://www.netbeans.org/ns/j2se-modular-project/1" filePattern="(.*$)" modsource="${src.gen2.dir.path}" property="src.gen2.dir.path.regexp"/>
+        <echo message="Copying resources from ${src.gen2.dir}"/>
+        <copy todir="${build.modules.dir}">
+            <fileset dir="${src.gen2.dir}" excludes="${build.classes.excludes},${excludes}" includes="${includes}"/>
+            <regexpmapper from="${src.gen2.dir.path.regexp}" to="\1/\3"/>
+        </copy>
         <j2semodularproject1:modsource_regexp xmlns:j2semodularproject1="http://www.netbeans.org/ns/j2se-modular-project/1" filePattern="(.*$)" modsource="${src.dir.path}" property="src.dir.path.regexp"/>
         <echo message="Copying resources from ${src.dir}"/>
         <copy todir="${build.modules.dir}">
@@ -1170,8 +1201,8 @@ is divided into following sections:
                 <isset property="main.class.available"/>
             </and>
         </condition>
-        <property name="platform.jlink" value="${jdk.home}/bin/jlink"/>
-        <property name="jlink.systemmodules.internal" value="${jdk.home}/jmods"/>
+        <property name="platform.jlink" value="${platform.home}/bin/jlink"/>
+        <property name="jlink.systemmodules.internal" value="${platform.home}/jmods"/>
         <exec executable="${platform.jlink}">
             <arg value="--module-path"/>
             <arg path="${jlink.systemmodules.internal}:${run.modulepath}:${dist.dir}"/>
@@ -1414,19 +1445,25 @@ is divided into following sections:
                 </not>
             </and>
         </condition>
+        <exec executable="${platform.java}" failonerror="false" outputproperty="platform.version.output">
+            <arg value="-version"/>
+        </exec>
         <condition else="" property="bug5101868workaround" value="*.java">
-            <matches pattern="1\.[56](\..*)?" string="${java.version}"/>
+            <matches multiline="true" pattern="1\.[56](\..*)?" string="${platform.version.output}"/>
         </condition>
         <condition else="" property="javadoc.html5.cmd.line.arg" value="-html5">
             <and>
                 <isset property="javadoc.html5"/>
-                <available file="${jdk.home}${file.separator}lib${file.separator}jrt-fs.jar"/>
+                <available file="${platform.home}${file.separator}lib${file.separator}jrt-fs.jar"/>
             </and>
         </condition>
-        <javadoc additionalparam="-J-Dfile.encoding=${file.encoding} ${javadoc.additionalparam}" author="${javadoc.author}" charset="UTF-8" destdir="${dist.javadoc.dir}" docencoding="UTF-8" encoding="${javadoc.encoding.used}" failonerror="true" noindex="${javadoc.noindex}" nonavbar="${javadoc.nonavbar}" notree="${javadoc.notree}" private="${javadoc.private}" source="${javac.source}" splitindex="${javadoc.splitindex}" use="${javadoc.use}" useexternalfile="true" version="${javadoc.version}" windowtitle="${javadoc.windowtitle}">
+        <javadoc additionalparam="-J-Dfile.encoding=${file.encoding} ${javadoc.additionalparam}" author="${javadoc.author}" charset="UTF-8" destdir="${dist.javadoc.dir}" docencoding="UTF-8" encoding="${javadoc.encoding.used}" executable="${platform.javadoc}" failonerror="true" noindex="${javadoc.noindex}" nonavbar="${javadoc.nonavbar}" notree="${javadoc.notree}" private="${javadoc.private}" source="${javac.source}" splitindex="${javadoc.splitindex}" use="${javadoc.use}" useexternalfile="true" version="${javadoc.version}" windowtitle="${javadoc.windowtitle}">
             <classpath>
                 <path path="${javac.classpath}"/>
             </classpath>
+            <fileset dir="${src.gen2.dir}" excludes="${bug5101868workaround},${excludes}" includes="${includes}">
+                <filename name="**/*.java"/>
+            </fileset>
             <fileset dir="${src.dir}" excludes="${bug5101868workaround},${excludes}" includes="${includes}">
                 <filename name="**/*.java"/>
             </fileset>
@@ -1438,6 +1475,9 @@ is divided into following sections:
             <arg line="${javadoc.html5.cmd.line.arg}"/>
         </javadoc>
         <copy todir="${dist.javadoc.dir}">
+            <fileset dir="${src.gen2.dir}" excludes="${excludes}" includes="${includes}">
+                <filename name="**/doc-files/**"/>
+            </fileset>
             <fileset dir="${src.dir}" excludes="${excludes}" includes="${includes}">
                 <filename name="**/doc-files/**"/>
             </fileset>
@@ -1695,15 +1735,6 @@ is divided into following sections:
         <propertyfile file="${built-clean.properties}">
             <entry key="${basedir}" value=""/>
         </propertyfile>
-        <antcall target="-maybe-call-dep">
-            <param name="call.built.properties" value="${built-clean.properties}"/>
-            <param location="${project.notzed_nativez}" name="call.subproject"/>
-            <param location="${project.notzed_nativez}/build.xml" name="call.script"/>
-            <param name="call.target" value="clean"/>
-            <param name="transfer.built-clean.properties" value="${built-clean.properties}"/>
-            <param name="transfer.not.archive.disabled" value="true"/>
-            <param name="transfer.do.jlink" value="false"/>
-        </antcall>
     </target>
     <target depends="init" name="-do-clean">
         <delete dir="${build.dir}"/>
index c2d597c..29eae41 100644 (file)
@@ -3,6 +3,6 @@ build.xml.script.CRC32=b55362bc
 build.xml.stylesheet.CRC32=32069288@1.6.1
 # This file is used by a NetBeans-based IDE to track changes in generated files such as build-impl.xml.
 # Do not edit this file. You may delete it but then the IDE will never regenerate such files for you.
-nbproject/build-impl.xml.data.CRC32=e7caf820
-nbproject/build-impl.xml.script.CRC32=2c69bf13
-nbproject/build-impl.xml.stylesheet.CRC32=0f0529df@1.6.1
+nbproject/build-impl.xml.data.CRC32=2c0455af
+nbproject/build-impl.xml.script.CRC32=68c8182d
+nbproject/build-impl.xml.stylesheet.CRC32=0f0529df@1.7
index f1aaa57..20b0bb1 100644 (file)
@@ -40,11 +40,10 @@ includes=**
 jar.compress=false
 javac.classpath=
 # Space-separated list of extra javac options
-javac.compilerargs=-Xlint:unchecked
+javac.compilerargs=-Xlint:unchecked --add-exports jdk.incubator.foreign/jdk.incubator.foreign.unsafe=notzed.zcl
 javac.deprecation=false
 javac.external.vm=false
-javac.modulepath=\
-    ${reference.notzed_nativez.notzed_nativez_jar}
+javac.modulepath=
 javac.processormodulepath=
 javac.processorpath=\
     ${javac.classpath}
@@ -75,9 +74,7 @@ jlink.additionalmodules=
 jlink.additionalparam=
 jlink.launcher=true
 jlink.launcher.name=notzed.zcl
-platform.active=default_platform
-project.notzed_nativez=../nativez
-reference.notzed_nativez.notzed_nativez_jar=${project.notzed_nativez}/dist/notzed.nativez.jar
+platform.active=JDK_15
 run.classpath=
 # Space-separated list of JVM arguments used when running the project.
 # You may also define separate properties like run-sys-prop.name=value instead of -Dname=value.
@@ -94,5 +91,9 @@ run.test.modulepath=\
 source.encoding=UTF-8
 src.dir=src
 src.dir.path=classes
+src.gen.dir=bin/gen
+src.gen.dir.path=classes
+src.gen2.dir=bin/gen
+src.gen2.dir.path=classes
 test.src.dir=src
 test.src.dir.path=tests
index 15787d7..575ebab 100644 (file)
@@ -4,22 +4,15 @@
     <configuration>
         <data xmlns="http://www.netbeans.org/ns/j2se-modular-project/1">
             <name>notzed.zcl</name>
+            <explicit-platform explicit-source-supported="true"/>
             <source-roots>
+                <root id="src.gen2.dir" pathref="src.gen2.dir.path"/>
                 <root id="src.dir" pathref="src.dir.path"/>
             </source-roots>
             <test-roots>
                 <root id="test.src.dir" pathref="test.src.dir.path"/>
             </test-roots>
         </data>
-        <references xmlns="http://www.netbeans.org/ns/ant-project-references/1">
-            <reference>
-                <foreign-project>notzed_nativez</foreign-project>
-                <artifact-type>jar</artifact-type>
-                <script>build.xml</script>
-                <target>jar</target>
-                <clean-target>clean</clean-target>
-                <id>notzed.nativez.jar</id>
-            </reference>
-        </references>
+        <references xmlns="http://www.netbeans.org/ns/ant-project-references/1"/>
     </configuration>
 </project>
similarity index 54%
rename from src/notzed.zcl/jni/zcl-init-dll.c
rename to src/notzed.zcl/classes/api/Allocator.java
index 1a9a10e..87f96f9 100644 (file)
@@ -1,5 +1,5 @@
 /*
- * Copyright (C) 2015 Michael Zucchi
+ * 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
  * along with this program.  If not, see <http://www.gnu.org/licenses/>.
  */
 
-/*
-  shared library init for mingwin
- */
-
-#include <windows.h>
-
-#include <CL/cl.h>
-#include <jni.h>
-
-#include "zcl-extension.h"
+package api;
 
-#define ZCL_DL_TABLE 1
-#include "zcl-functions.h"
+import jdk.incubator.foreign.MemoryAddress;
+import jdk.incubator.foreign.MemorySegment;
 
-static void *lib_opencl;
-
-int init_dynamic(JNIEnv *env) {
-       lib_opencl = LoadLibrary("OpenCL.dll");
-
-       if (lib_opencl) {
-               for (int i=0;i<zcl_fn_SIZEOF;i++)
-                       zcl_fp[i] = GetProcAddress(lib_opencl, zcl_fn[i]);
-               return 0;
-       }
-
-       return -1;
+/**
+ * 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
new file mode 100644 (file)
index 0000000..936ae55
--- /dev/null
@@ -0,0 +1,61 @@
+/*
+ * 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 <http://www.gnu.org/licenses/>.
+ */
+
+/*
+  Not sure on this.
+
+  Something needs to handle gc of callback stubs.
+  This uses Native to do it.
+ */
+
+package api;
+
+import jdk.incubator.foreign.*;
+
+public class Callback<T> extends Native implements AutoCloseable {
+       T func;
+
+       public Callback(MemoryAddress addr, T func) {
+               super(addr);
+
+               this.func = func;
+       }
+
+       /*
+        * A callback that resolve to MemoryAddress.NULL.
+        * This can be released safely any number of times.
+        * This must never be registerd with ref queue.
+        */
+       public static final Callback NULL = new Callback<>(MemoryAddress.NULL, null) {
+                       @Override
+                       public void release() {
+                       }
+                       @Override
+                       public void close() {
+                       }
+               };
+
+       private static void release(MemoryAddress p) {
+               System.err.printf("** release upcall stub %016x\n", Memory.toLong(p));
+               freeUpcallStub(p);
+       }
+
+       @Override
+       public void close() {
+               release();
+       }
+}
diff --git a/src/notzed.zcl/classes/api/Memory.java b/src/notzed.zcl/classes/api/Memory.java
new file mode 100644 (file)
index 0000000..8107d6c
--- /dev/null
@@ -0,0 +1,221 @@
+/*
+ * 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 <http://www.gnu.org/licenses/>.
+ */
+
+package api;
+
+import jdk.incubator.foreign.*;
+import jdk.incubator.foreign.unsafe.ForeignUnsafe;
+import java.lang.invoke.MethodHandles;
+import java.lang.invoke.MethodHandle;
+import java.lang.invoke.MethodType;
+
+/**
+ * A utility library for memory operations including a stack allocator.
+ *
+ * The stack allocator works like this
+ *  try (Memory.Frame f = Memory.frame()) {
+ *    MemoryAddress a = Memory.alloca(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.
+ */
+public class Memory {
+
+       private static final ThreadLocal<Stack> stacks = ThreadLocal.withInitial(() -> Native.resolve(malloc(4096), (p) -> new Stack(p, 4096)));
+       private static final ThreadLocal<Allocator> stackAllocators = new ThreadLocal<>();
+
+       /* older idea */
+
+       // eh, this isn't really necessary, it's just a safety check/debug?
+       //private static final ThreadLocal<Frame> 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;
+
+               Stack(MemoryAddress p, long size) {
+                       super(p);
+
+                       stack = ForeignUnsafe.ofNativeUnchecked(addr(), size);
+                       base = stack.baseAddress();
+                       sp = size;
+               }
+
+               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
+               SystemABI abi = SystemABI.getInstance();
+               MethodHandles.Lookup lookup = MethodHandles.lookup();
+               LibraryLookup libc = LibraryLookup.ofDefault();
+
+               try {
+                       malloc = abi.downcallHandle(libc.lookup("malloc"),
+                                                   MethodType.methodType(MemoryAddress.class,
+                                                                         long.class),
+                                                   FunctionDescriptor.of(MemoryLayouts.SysV.C_POINTER,
+                                                                         false,
+                                                                         MemoryLayouts.SysV.C_ULONG));
+                       free = abi.downcallHandle(libc.lookup("free"),
+                                                   MethodType.methodType(void.class,
+                                                                         MemoryAddress.class),
+                                                   FunctionDescriptor.ofVoid(false,
+                                                                             MemoryLayouts.SysV.C_POINTER));
+               } catch (NoSuchMethodException 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 ForeignUnsafe.ofNativeUnchecked(addr, size);
+       }
+
+       /**
+        * Get the physical address.  I mean how is this an "offset"?
+        */
+       public static long toLong(MemoryAddress addr) {
+               return ForeignUnsafe.getUnsafeOffset(addr);
+       }
+
+       /**
+        * 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);
+               }
+       }
+}
diff --git a/src/notzed.zcl/classes/api/Native.java b/src/notzed.zcl/classes/api/Native.java
new file mode 100644 (file)
index 0000000..ef599df
--- /dev/null
@@ -0,0 +1,931 @@
+/*
+ * 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 <http://www.gnu.org/licenses/>.
+ */
+
+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;
+import java.util.function.Function;
+import java.util.function.IntFunction;
+import jdk.incubator.foreign.*;
+import jdk.incubator.foreign.unsafe.ForeignUnsafe;
+import java.lang.ref.ReferenceQueue;
+import java.lang.ref.WeakReference;
+import java.lang.reflect.Constructor;
+import java.lang.reflect.InvocationTargetException;
+import java.lang.System.Logger.Level;
+
+public class Native {
+
+       private final MemoryAddress p;
+
+       private final static boolean dolog = true;
+
+       final static VarHandle byteHandle = MemoryHandles.varHandle(byte.class, ByteOrder.nativeOrder());
+       final static VarHandle shortHandle = MemoryHandles.varHandle(short.class, ByteOrder.nativeOrder());
+       final static VarHandle intHandle = MemoryHandles.varHandle(int.class, ByteOrder.nativeOrder());
+       final static VarHandle longHandle = MemoryHandles.varHandle(long.class, ByteOrder.nativeOrder());
+       final static VarHandle floatHandle = MemoryHandles.varHandle(float.class, ByteOrder.nativeOrder());
+       final static VarHandle doubleHandle = MemoryHandles.varHandle(double.class, ByteOrder.nativeOrder());
+       final static VarHandle addrHandle = MemoryHandles.varHandle(MemoryAddress.class, ByteOrder.nativeOrder());
+
+       final static VarHandle byteVHandle = MemoryHandles.withStride(byteHandle, 1);
+       final static VarHandle shortVHandle = MemoryHandles.withStride(shortHandle, 2);
+       final static VarHandle intVHandle = MemoryHandles.withStride(intHandle, 4);
+       final static VarHandle longVHandle = MemoryHandles.withStride(longHandle, 8);
+       final static VarHandle floatVHandle = MemoryHandles.withStride(floatHandle, 4);
+       final static VarHandle doubleVHandle = MemoryHandles.withStride(doubleHandle, 8);
+       final static VarHandle addrVHandle = MemoryHandles.withStride(addrHandle, 8);
+
+       protected Native(MemoryAddress p) {
+               this.p = p;
+       }
+
+       static System.Logger log() {
+               return System.getLogger("notzed.native");
+       }
+
+       public MemoryAddress addr() {
+               return p;
+       }
+
+       public static MemoryAddress addr(Native o) {
+               return o != null ? o.addr() : MemoryAddress.NULL;
+       }
+
+       public static MemoryAddress addr(MemorySegment o) {
+               return o != null ? o.baseAddress() : MemoryAddress.NULL;
+       }
+
+       public static byte getByte(MemoryAddress p) {
+               return (byte)byteHandle.get(p);
+       }
+
+       public static byte getByte(MemoryAddress p, long i) {
+               return (byte)byteVHandle.get(p, i);
+       }
+
+       public static void setByte(MemoryAddress p, byte v) {
+               byteHandle.set(p, v);
+       }
+
+       public static void setByte(MemoryAddress p, int i, byte v) {
+               byteVHandle.set(p, i, v);
+       }
+
+       public static int getInt(MemoryAddress p) {
+               return (int)intHandle.get(p);
+       }
+
+       public static int getInt(MemoryAddress p, long i) {
+               return (int)intVHandle.get(p, i);
+       }
+
+       public static void setInt(MemoryAddress p, int v) {
+               intHandle.set(p, v);
+       }
+
+       public static void setInt(MemoryAddress p, long i, int v) {
+               intVHandle.set(p, i, v);
+       }
+
+       public static short getShort(MemoryAddress p) {
+               return (short)shortHandle.get(p);
+       }
+
+       public static void setShort(MemoryAddress p, short v) {
+               shortHandle.set(p, v);
+       }
+
+       public static void setShort(MemoryAddress p, long i, short v) {
+               shortVHandle.set(p, i, v);
+       }
+
+       public static long getLong(MemoryAddress p) {
+               return (long)longHandle.get(p);
+       }
+
+       public static long getLong(MemoryAddress p, long i) {
+               return (long)longVHandle.get(p, i);
+       }
+
+       public static void setLong(MemoryAddress p, long v) {
+               longHandle.set(p, v);
+       }
+
+       public static void setLong(MemoryAddress p, long i, long v) {
+               longVHandle.set(p, i, v);
+       }
+
+       public static float getFloat(MemoryAddress p) {
+               return (float)floatHandle.get(p);
+       }
+
+       public static float getFloat(MemoryAddress p, long i) {
+               return (float)floatVHandle.get(p, i);
+       }
+
+       public static void setFloat(MemoryAddress p, float v) {
+               floatHandle.set(p, v);
+       }
+
+       public static void setFloat(MemoryAddress p, long i, float v) {
+               floatVHandle.set(p, i, v);
+       }
+
+       public static double getDouble(MemoryAddress p) {
+               return (int)doubleHandle.get(p);
+       }
+
+       public static double getDouble(MemoryAddress p, long i) {
+               return (int)doubleVHandle.get(p, i);
+       }
+
+       public static void setDouble(MemoryAddress p, double v) {
+               doubleHandle.set(p, v);
+       }
+
+       public static void setDouble(MemoryAddress p, long i, double v) {
+               doubleVHandle.set(p, i, v);
+       }
+
+       public static MemoryAddress getAddr(MemoryAddress p) {
+               return (MemoryAddress)addrHandle.get(p);
+       }
+
+       public static MemoryAddress getAddr(MemoryAddress p, long i) {
+               return (MemoryAddress)addrVHandle.get(p, i);
+       }
+
+       public static void setAddr(MemoryAddress p, MemoryAddress v) {
+               addrHandle.set(p, v);
+       }
+
+       public static void setAddr(MemoryAddress p, long i, MemoryAddress v) {
+               addrVHandle.set(p, i, v);
+       }
+
+       /* helpers - java to native */
+       public static <T extends Native> MemoryAddress toAddrV(Allocator frame, T[]array) {
+               MemoryAddress list = frame.alloca(8 *  array.length);
+
+               for (int i=0;i<array.length;i++)
+                       setAddr(list, i, array[i].addr());
+
+               return list;
+       }
+
+       public static <T extends Native> MemoryAddress toAddrV(Allocator frame, String[]array) {
+                if (array != null) {
+                        MemoryAddress list = frame.alloca(8 *  array.length);
+
+                        for (int i=0;i<array.length;i++)
+                                setAddr(list, i, toByteV(frame, array[i]));
+
+                        return list;
+                } else {
+                        return MemoryAddress.NULL;
+                }
+        }
+
+       public static <T extends Native> MemoryAddress toLongV(Allocator frame, long[]array) {
+               MemoryAddress list = frame.alloca(8 *  array.length);
+
+               for (int i=0;i<array.length;i++)
+                       setLong(list, i, array[i]);
+
+               return list;
+       }
+
+       public static <T extends Native> MemoryAddress toByteV(Allocator frame, byte[] data) {
+               if (data != null) {
+                       MemoryAddress list = frame.alloca(data.length);
+
+                       for (int i=0;i<data.length;i++)
+                               setByte(list, i, data[i]);
+
+                       return list;
+               } else {
+                       return MemoryAddress.NULL;
+               }
+       }
+
+       public static <T extends Native> MemoryAddress toByteV(Allocator frame, String string) {
+               if (string != null) {
+                       byte[] data = string.getBytes();
+                       MemoryAddress list = frame.alloca(data.length + 1);
+
+                       for (int i=0;i<data.length;i++)
+                               setByte(list, i, data[i]);
+                       setByte(list, data.length, (byte)0);
+
+                       return list;
+               } else {
+                       return MemoryAddress.NULL;
+               }
+       }
+
+       /* helpers - native to java */
+       public static <T extends Native> T[] toObjectV(MemoryAddress list, T[] array, Function<MemoryAddress,T> create) {
+               for (int i=0;i<array.length;i++)
+                       array[i] = Native.resolve(getAddr(list, i), create);
+               return array;
+       }
+
+       public static <T extends Native> T[] toObjectV(MemorySegment list, Function<MemoryAddress,T> create, IntFunction<T[]> createArray) {
+               return toObjectV(list.baseAddress(), createArray.apply((int)(list.byteSize() >>> 3)), create);
+       }
+
+       public static String toString(MemoryAddress cstr) {
+               MemorySegment seg = ForeignUnsafe.ofNativeUnchecked(cstr, Integer.MAX_VALUE);
+               MemoryAddress add = seg.baseAddress();
+               byte[] data;
+               int len = 0;
+
+               while (getByte(add, len) != 0)
+                       len++;
+
+               data = new byte[len];
+               for (int i=0;i<data.length;i++)
+                       data[i] = getByte(add, i);
+
+               return new String(data);
+       }
+
+       public static long[] toLongV(MemorySegment valp) {
+               MemoryAddress val = valp.baseAddress();
+               int len = (int)(valp.byteSize() >>> 3);
+               long[] list = new long[len];
+
+               for (int i=0;i<list.length;i++)
+                       list[i] = getLong(val, i);
+
+               return list;
+       }
+
+       /* abi stuff */
+
+       public static MethodHandle downcallHandle(LibraryLookup[] libs, String name, String signature) {
+               Signature sig = Signature.parse(signature);
+               int n = sig.classes.length;
+               SystemABI abi = SystemABI.getInstance();
+
+               if (sig.classes.length != sig.layouts.length)
+                       throw new RuntimeException("layout class mismatch");
+               Class<?> resClass = sig.classes[n-1];
+               MemoryLayout resLayout = sig.layouts[n-1];
+
+               for (LibraryLookup lib: libs) {
+                       try {
+                               return abi.downcallHandle(lib.lookup(name),
+                                                         MethodType.methodType(resClass, Arrays.copyOf(sig.classes, n-1)),
+                                                         FunctionDescriptor.of(resLayout, false, Arrays.copyOf(sig.layouts, n-1)));
+                       } catch (NoSuchMethodException x) {
+                       }
+               }
+               // or use some 'unsupportedoperation' one?
+               System.err.println("not found: "+ name);
+
+               return null;
+       }
+
+       public static MethodHandle downcallHandle(MemoryAddress addr, String signature) {
+               Signature sig = Signature.parse(signature);
+               int n = sig.classes.length;
+               SystemABI abi = SystemABI.getInstance();
+
+               if (sig.classes.length != sig.layouts.length)
+                       throw new RuntimeException("layout class mismatch");
+               Class<?> resClass = sig.classes[n-1];
+               MemoryLayout resLayout = sig.layouts[n-1];
+
+               return abi.downcallHandle(addr,
+                                         MethodType.methodType(resClass, Arrays.copyOf(sig.classes, n-1)),
+                                         FunctionDescriptor.of(resLayout, false, Arrays.copyOf(sig.layouts, n-1)));
+       }
+
+       // instance must be of a functional interface
+       public static MemoryAddress upcallStub(MethodHandles.Lookup lookup, Object instance, String signature) {
+               Signature sig = Signature.parse(signature);
+               int n = sig.classes.length;
+               SystemABI abi = SystemABI.getInstance();
+
+               if (sig.classes.length != sig.layouts.length)
+                       throw new RuntimeException("layout class mismatch");
+               Class<?> resClass = sig.classes[n-1];
+               MemoryLayout resLayout = sig.layouts[n-1];
+
+
+               Method m = instance.getClass().getMethods()[0];
+               MethodType mt = MethodType.methodType(m.getReturnType(), m.getParameterTypes());
+
+               //System.out.printf("instance %s\n", instance);
+               //System.out.printf("declaring class %s\n", m.getDeclaringClass());
+
+               try {
+                       return abi.upcallStub(lookup.findVirtual(
+                                                     m.getDeclaringClass(),
+                                                     m.getName(),
+                                                     mt)
+                                             .bindTo(instance),
+                                             FunctionDescriptor.of(resLayout, false, Arrays.copyOf(sig.layouts, n-1)));
+               } catch (NoSuchMethodException | IllegalAccessException x) {
+                       throw new RuntimeException(x);
+               }
+       }
+
+       public static void freeUpcallStub(MemoryAddress addr) {
+               SystemABI.getInstance().freeUpcallStub(addr);
+       }
+
+       public static LibraryLookup[] loadLibraries(String... libraries) {
+               LibraryLookup[] libs = new LibraryLookup[libraries.length];
+               MethodHandles.Lookup lookup = MethodHandles.lookup();
+               for (int i=0;i<libraries.length;i++)
+                       libs[i] = LibraryLookup.ofLibrary(lookup, libraries[i]);
+               return libs;
+       }
+
+       public static MemoryLayout parseStruct(String layout) {
+               Signature sig = Signature.parse(layout);
+
+               return MemoryLayout.ofStruct(sig.layouts);
+       }
+
+       public static MemoryLayout parseUnion(String layout) {
+               Signature sig = Signature.parse(layout);
+
+               return MemoryLayout.ofUnion(sig.layouts);
+       }
+
+       public static void dumpSignature(String signature) {
+               Signature sig = Signature.parse(signature);
+               for (int i=0;i<sig.classes.length;i++) {
+                       System.out.printf(" %-40s %s\n", sig.classes[i].getName(), sig.layouts[i]);
+               }
+       }
+
+       static class Signature {
+               Class classes[];
+               MemoryLayout layouts[];
+
+               public Signature(Class[] classes, MemoryLayout[] layouts) {
+                       this.classes = classes;
+                       this.layouts = layouts;
+               }
+
+               public static Signature parse(String s) {
+                       List<Class> argClass = new ArrayList<>();
+                       List<MemoryLayout> argLayout = new ArrayList<>();
+                       StringReader r = new StringReader(s);
+                       int c;
+                       int size = 0;
+                       StringBuilder sb = new StringBuilder();
+                       int pointerDepth = 0;
+                       int type = 0; // func/struct/union
+
+                       //System.out.printf("parse: '%s'\n", s);
+                       try {
+                               c = r.read();
+                               switch (c) {
+                               case '(':
+                                       type = 0;
+                                       break;
+                               case '[':
+                                       type = 1;
+                                       break;
+                               default:
+                                       throw new RuntimeException("Unknown type: " + (char)c);
+                               }
+                               c = r.read();
+
+                               while ( c != -1) {
+                                       switch (c) {
+                                       case 'u':
+                                       case 'i':
+                                       case 'f':
+                                       case 'x': {
+                                               int d;
+                                               size = 0;
+                                               while ((d = r.read()) != -1 && d >= '0' && d <= '9')
+                                                       size = size * 10 + (d - '0');
+
+                                               // named field
+                                               if (d == '(') {
+                                                       while ((d = r.read()) != -1 && d != ')')
+                                                               ;
+                                                       d = r.read();
+                                               }
+
+                                               if (d == ':') {
+                                                       //System.out.printf(" pointer: size=%d\n", size);
+                                                       pointerDepth++;
+                                               } else {
+                                                       if (d == '=') {
+                                                               // actually it's bitfield, but ignore internal details
+                                                               d = r.read();
+                                                               if (d != '[')
+                                                                       throw new UnsupportedOperationException("expecting [");
+                                                               while ((d = r.read()) != -1 && d != ']')
+                                                                       ;
+                                                               d = r.read();
+                                                       } else {
+                                                               //System.out.printf(" prim: *=%d size=%d\n", pointerDepth, size);
+                                                       }
+
+                                                       if (pointerDepth > 0) {
+                                                               argClass.add(MemoryAddress.class);
+                                                               argLayout.add(MemoryLayouts.SysV.C_POINTER);
+                                                       } else {
+                                                               switch (c) {
+                                                               case 'u':
+                                                                       switch (size) {
+                                                                       case 8:
+                                                                               argClass.add(byte.class);
+                                                                               argLayout.add(MemoryLayouts.SysV.C_UCHAR);
+                                                                               break;
+                                                                       case 16:
+                                                                               // char.class?
+                                                                               argClass.add(short.class);
+                                                                               argLayout.add(MemoryLayouts.SysV.C_USHORT);
+                                                                               break;
+                                                                       case 32:
+                                                                               argClass.add(int.class);
+                                                                               argLayout.add(MemoryLayouts.SysV.C_UINT);
+                                                                               break;
+                                                                       case 64:
+                                                                               argClass.add(long.class);
+                                                                               argLayout.add(MemoryLayouts.SysV.C_ULONG);
+                                                                               break;
+                                                                       }
+                                                                       break;
+                                                               case 'i':
+                                                                       switch (size) {
+                                                                       case 8:
+                                                                               argClass.add(byte.class);
+                                                                               argLayout.add(MemoryLayouts.SysV.C_SCHAR);
+                                                                               break;
+                                                                       case 16:
+                                                                               argClass.add(short.class);
+                                                                               argLayout.add(MemoryLayouts.SysV.C_SHORT);
+                                                                               break;
+                                                                       case 32:
+                                                                               argClass.add(int.class);
+                                                                               argLayout.add(MemoryLayouts.SysV.C_INT);
+                                                                               break;
+                                                                       case 64:
+                                                                               argClass.add(long.class);
+                                                                               argLayout.add(MemoryLayouts.SysV.C_LONG);
+                                                                               break;
+                                                                       }
+                                                                       break;
+                                                               case 'f':
+                                                                       switch (size) {
+                                                                       case 32:
+                                                                               argClass.add(float.class);
+                                                                               argLayout.add(MemoryLayouts.SysV.C_FLOAT);
+                                                                               break;
+                                                                       case 64:
+                                                                               argClass.add(double.class);
+                                                                               argLayout.add(MemoryLayouts.SysV.C_DOUBLE);
+                                                                               break;
+                                                                       }
+                                                                       break;
+                                                               case 'x':
+                                                                       switch (size) {
+                                                                       case 8:
+                                                                               argLayout.add(MemoryLayouts.PAD_8);
+                                                                               break;
+                                                                       case 16:
+                                                                               argLayout.add(MemoryLayouts.PAD_16);
+                                                                               break;
+                                                                       case 32:
+                                                                               argLayout.add(MemoryLayouts.PAD_32);
+                                                                               break;
+                                                                       case 64:
+                                                                               argLayout.add(MemoryLayouts.PAD_64);
+                                                                               break;
+                                                                       }
+                                                                       break;
+                                                               }
+                                                       }
+                                                       c = d;
+                                                       pointerDepth = 0;
+                                                       continue;
+                                               }
+                                               break;
+                                       }
+                                       case 'v':
+                                               //System.out.printf("void *=%d\n", pointerDepth);
+                                               if (pointerDepth > 0) {
+                                                       argClass.add(MemoryAddress.class);
+                                                       argLayout.add(MemoryLayouts.SysV.C_POINTER);
+                                               } else {
+                                                       // can only be return value
+                                                       argClass.add(void.class);
+                                                       argLayout.add(null);
+                                               }
+                                               pointerDepth = 0;
+                                               break;
+                                       case '$':
+                                               c = r.read();
+                                               if (c != '{')
+                                                       throw new RuntimeException();
+                                               sb.setLength(0);
+                                               boolean cap = true;
+                                               while ((c = r.read()) != -1 && c != '}') {
+                                                       if (c == '_') {
+                                                               cap = true;
+                                                               continue;
+                                                       }
+                                                       if (cap) {
+                                                               c = Character.toUpperCase(c);
+                                                               cap = false;
+                                                       }
+                                                       sb.append((char)c);
+                                               }
+                                               //System.out.printf(" type: *=%d  %s\n", pointerDepth, sb);
+                                               if (pointerDepth > 0) {
+                                                       argClass.add(MemoryAddress.class);
+                                                       argLayout.add(MemoryLayouts.SysV.C_POINTER);
+                                               } else {
+                                                       argClass.add(MemorySegment.class);
+                                                       try {
+                                                               Class<?> stype = Class.forName("api." + sb.toString());
+                                                               argLayout.add((MemoryLayout)stype.getDeclaredMethod("layout").invoke(null));
+                                                       } catch (Exception x) {
+                                                               throw new RuntimeException(x);
+                                                       }
+                                               }
+                                               pointerDepth = 0;
+                                               break;
+                                       case '(': // named field, ignore it
+                                               while ((c = r.read()) != -1 && c != ')')
+                                                       ;
+                                               break;
+                                       case '|':
+                                               // wrong!  could be inline union sub-part
+                                               if (type == 1 || type == 2)
+                                                       type = 2;
+                                               else
+                                                       throw new UnsupportedOperationException("union in thing?");
+                                               break;
+                                       case ')':
+                                               break;
+                                       case ']':
+                                               break;
+                                       default:
+                                               throw new UnsupportedOperationException("Unknown token: " + (char)c);
+                                       }
+                                       c = r.read();
+                               }
+                       } catch (java.io.IOException x) {
+                               throw new RuntimeException(x);
+                       }
+
+                       return new Signature(argClass.stream().toArray(Class[]::new),
+                                            argLayout.stream().toArray(MemoryLayout[]::new));
+               }
+       }
+
+       /* ********************************************************************** */
+       /* GC handling */
+       /* ********************************************************************** */
+
+       /**
+        * Resource index.
+        */
+       static private final PointerTable map = new PointerTable();
+
+       /**
+        * Reference queue for stale objects.
+        */
+       static private final ReferenceQueue<Native> references = new ReferenceQueue<>();
+
+       private static <T extends Native> T createInstance(Class<T> jtype, MemoryAddress p) {
+               cleanerStep();
+               try {
+                       Class[] params = {MemoryAddress.class};
+                       Constructor<T> cc = jtype.getDeclaredConstructor(params);
+
+                       cc.setAccessible(true);
+
+                       return cc.newInstance(p);
+               } catch (NoSuchMethodException | SecurityException | InstantiationException | IllegalAccessException | IllegalArgumentException | InvocationTargetException ex) {
+                       log().log(Level.ERROR, "createInstance", ex);
+                       throw new RuntimeException(ex);
+               }
+       }
+
+       /*
+       public static <T extends Native> T resolve(Class<T> jtype, MemoryAddress p) {
+               T o;
+
+               //if (dolog)
+               log().log(Level.DEBUG, () -> String.format("  resolve $%016x %s", p.offset(), jtype.getName()));
+
+               // Instantiation needs to be synchronized for obvious reasons.
+               synchronized (map) {
+                       CHandle h = (CHandle) map.get(p);
+
+                       if (h == null || (o = jtype.cast(h.get())) == null) {
+                               o = createInstance(jtype, p);
+                               h = new CHandle(o, references, p);
+                               map.putAlways(h);
+                       }
+               }
+               return o;
+               }*/
+
+       public static <T extends Native> T resolve(MemoryAddress p, Function<MemoryAddress,T> create) {
+               T o;
+               boolean step = false;
+
+               if (dolog)
+                       log().log(Level.DEBUG, () -> String.format("  resolv $%016x %s", Memory.toLong(p), create));
+
+               // ??? who the fuck knows if this will work
+               if (p.offset() == 0)
+                       return null;
+
+               // Instantiation needs to be synchronized for obvious reasons.
+               synchronized (map) {
+                       CHandle h = (CHandle) map.get(p);
+
+                       if (h == null || (o = (T)(h.get())) == null) {
+                               o = create.apply(p);
+                               h = new CHandle(o, references, p);
+                               map.put(h);
+                               step = true;
+                       }
+               }
+
+               if (step)
+                       cleanerStep();
+
+               return o;
+       }
+
+       /*
+       public static <T extends Native> void register(T o) {
+               T o;
+               boolean step = false;
+
+               if (dolog)
+                       log().log(Level.DEBUG, () -> String.format("  regist $%016x %s", o.addr().offset(), o.getClass().getName()));
+
+               CHandle h = new CHandle(o, references, o.addr());
+
+               synchronized (map) {
+                       map.put(h);
+                       step = true;
+               }
+
+               if (step)
+                       cleanerStep();
+
+               return o;
+               }*/
+
+       public void release() {
+               WeakReference<? extends Native> ref;
+
+               synchronized (map) {
+                       ref = map.remove(p);
+               }
+
+               if (ref != null) {
+                       if (dolog)
+                               log().log(Level.DEBUG, () -> String.format("  force  $%016x %s", Memory.toLong(p), getClass().getName()));
+
+                       ref.enqueue();
+               }
+       }
+
+       public static <T extends Native> void release(T a) {
+               if (a != null)
+                       a.release();
+       }
+
+       public static void release(Native... list) {
+               for (Native o : list)
+                       release(o);
+       }
+
+       static {
+               Thread cleanup = new Thread(Native::cleaner, "Native cleaner");
+               cleanup.setPriority(Thread.MAX_PRIORITY);
+               cleanup.setDaemon(true);
+               cleanup.start();
+       }
+
+       private static void cleanerStep() {
+               try {
+                       CHandle stale = (CHandle) references.poll();
+                       if (stale != null) {
+                               synchronized (map) {
+                                       map.remove(stale.p);
+                               }
+                               stale.release();
+                       }
+               } catch (Throwable ex) {
+               }
+       }
+
+       /**
+        * Cleaner thread.
+        * <p>
+        * This polls the reference queue and releases objects via
+        * their static release method.
+        */
+       private static void cleaner() {
+               if (dolog)
+                       log().log(Level.INFO, "Native finaliser started");
+               try {
+                       while (true) {
+                               CHandle stale = (CHandle) references.remove();
+                               do {
+                                       try {
+                                               synchronized (map) {
+                                                       map.remove(stale.p);
+                                               }
+                                               stale.release();
+                                       } catch (Throwable ex) {
+                                       }
+                                       stale = (CHandle) references.poll();
+                               } while (stale != null);
+                       }
+               } catch (InterruptedException ex) {
+               }
+       }
+
+       private static class CHandle extends WeakReference<Native> {
+               protected MemoryAddress p;
+               final Class<? extends Native> jtype;
+               CHandle next;
+
+               CHandle(Native referent, ReferenceQueue<Native> references, MemoryAddress p) {
+                       super(referent, references);
+                       this.p = p;
+                       this.jtype = referent.getClass();
+               }
+
+               void release() {
+                       try {
+                               if (p != null) {
+                                       if (dolog)
+                                               log().log(Level.DEBUG, () -> String.format("  releas $%016x %s", Memory.toLong(p), jtype.getName()));
+
+                                       Method mm = jtype.getDeclaredMethod("release", MemoryAddress.class);
+                                       mm.setAccessible(true);
+                                       mm.invoke(null, p);
+                               }
+                       } catch (NoSuchMethodException | SecurityException | IllegalAccessException | IllegalArgumentException | InvocationTargetException ex) {
+                               log().log(Level.ERROR, jtype.getName(), ex);
+                       } finally {
+                               p = null;
+                       }
+               }
+
+               @Override
+               public boolean equals(Object obj) {
+                       return (obj instanceof CHandle) && ((CHandle) obj).p == p;
+               }
+
+               @Override
+               public int hashCode() {
+                       //return p.hashCode();
+                       return hashCode(p);
+               }
+
+               /**
+                * Simple hashcode for native pointers.
+                * <p>
+                * This simply strips the bottom 4 bits from the pointer as
+                * on a 64-bit system the low 3 bits are typically zero and the 4th
+                * isn't very well distributed.
+                *
+                * @param p
+                * @return
+                */
+               public static final int hashCode(long p) {
+                       return (int) p >>> 4;
+               }
+
+               /**
+                * Sigh, memoryaddress has a miserable hashCode(), it's even worse than Long.hashCode()
+                */
+               public static final int hashCode(MemoryAddress p) {
+                       return p.hashCode() >>> 5;
+               }
+       }
+
+       /**
+        * Lightweight pointer hashtable.
+        * <p>
+        * This serves two purposes:
+        * <ol>
+        * <li>Track and resolve unique objects based on memory address;
+        * <li>Hold hard references to the WeakReference as required by the gc system.
+        * </ol>
+        * <p>
+        * CHandle's are chained directly from the index table, the p field
+        * is used as a key directly, and hash values are not cached. This combines
+        * to save significant memory per node.
+        */
+       private static class PointerTable {
+
+               int mask = 63;
+               int size = 0;
+               CHandle[] table = new CHandle[64];
+
+               private void resize(int length) {
+                       CHandle[] ntable = new CHandle[length];
+                       int nmask = length - 1;
+
+                       for (int i = 0; i < table.length; i++) {
+                               CHandle h = table[i];
+
+                               while (h != null) {
+                                       CHandle n = h.next;
+                                       int k = h.hashCode() & nmask;
+
+                                       h.next = ntable[k];
+                                       ntable[k] = h;
+
+                                       h = n;
+                               }
+                       }
+
+                       table = ntable;
+                       mask = nmask;
+               }
+
+               public CHandle put(CHandle h) {
+                       CHandle o = remove(h.p);
+
+                       putAlways(h);
+
+                       return o;
+               }
+
+               public void putAlways(CHandle h) {
+                       if (size > table.length * 2)
+                               resize(table.length * 2);
+
+                       int i = h.hashCode() & mask;
+
+                       h.next = table[i];
+                       table[i] = h;
+                       size += 1;
+               }
+
+               public CHandle get(MemoryAddress p) {
+                       int i = CHandle.hashCode(p) & mask;
+                       CHandle h = table[i];
+
+                       while (h != null && !h.p.equals(p))
+                               h = h.next;
+                       return h;
+               }
+
+               public CHandle remove(MemoryAddress p) {
+                       int i = CHandle.hashCode(p) & mask;
+                       CHandle h = table[i];
+                       CHandle a = null;
+
+                       while (h != null && !h.p.equals(p)) {
+                               a = h;
+                               h = h.next;
+                       }
+                       if (h != null) {
+                               if (a != null)
+                                       a.next = h.next;
+                               else
+                                       table[i] = h.next;
+                               size -= 1;
+                       }
+
+                       return h;
+               }
+       }
+}
index 731823a..d0f59af 100644 (file)
  */
 package au.notzed.zcl;
 
+import jdk.incubator.foreign.*;
+
 import java.nio.ByteBuffer;
+import java.lang.invoke.MethodHandle;
 
 /**
  * Interface for memory buffers.
@@ -25,42 +28,23 @@ import java.nio.ByteBuffer;
 public class CLBuffer extends CLMemory {
 
        /**
-        * Use USE_HOST_PTR was used then this keeps track of the host ptr reference
-        * to avoid java freeing it.
-        */
-       private final ByteBuffer hostPtr;
-
-       /**
         * Create an interface for a native pointer of type cl_mem that refers to a
         * buffer object.
         *
         * @param p Native pointer.
         */
-       public CLBuffer(long p) {
-               super(p);
-               hostPtr = null;
+       public CLBuffer(MemoryAddress p) {
+               this(p, null);
        }
 
-       public CLBuffer(long p, ByteBuffer hostPtr) {
-               super(p);
-
-               this.hostPtr = hostPtr;
+       public CLBuffer(MemoryAddress p, MemorySegment seg) {
+               super(p, seg);
        }
 
-       static void release(long p) {
+       static void release(MemoryAddress p) {
                CLMemory.release(p);
        }
 
-       @Override
-       int getInfoType() {
-               return TYPE_MEM_OBJECT;
-       }
-
-       @Override
-       public ByteBuffer getHostPtr() {
-               return hostPtr;
-       }
-
        /**
         * Parameter to define a region for createSubBuffer().
         *
index 94ef9ec..fa97c8a 100644 (file)
@@ -20,7 +20,7 @@ package au.notzed.zcl;
  * Parameters for Buffer.createSubBuffer()
  * <p>
  * Specific types are defined on CLBuffer.
- * 
+ *
  */
 public abstract class CLBufferInfo {
 
index ed2539b..c73cfa7 100644 (file)
  */
 package au.notzed.zcl;
 
-import java.nio.ByteBuffer;
 import static au.notzed.zcl.CL.*;
-import au.notzed.zcl.khr.GLSharing;
+import static au.notzed.zcl.CLLib.*;
+import jdk.incubator.foreign.*;
+import api.Native;
+import api.Allocator;
+import api.Memory;
+import api.Callback;
+
+import java.util.ArrayList;
+
+import java.nio.ByteBuffer;
+import java.nio.ByteOrder;
+import java.nio.BufferOverflowException;
+import java.nio.BufferUnderflowException;
+//import static au.au.notzed.zcl.CL.*;
+//import au.notzed.zcl.khr.GLSharing;
+
+import java.lang.invoke.MethodHandle;
+import java.util.function.Function;
 
 /**
  * Interface for cl_command_queue.
@@ -44,15 +60,24 @@ public class CLCommandQueue extends CLExtendable {
         *
         * @param p Native pointer.
         */
-       public CLCommandQueue(long p) {
+       CLCommandQueue(MemoryAddress p) {
                super(p);
        }
 
-       private native static void release(long p);
+       public static CLCommandQueue create(MemoryAddress p) {
+               return Native.resolve(p, CLCommandQueue::new);
+       }
+
+       private static void release(MemoryAddress p) {
+               try {
+                       clReleaseCommandQueue(p);
+               } catch (Throwable t) {
+               }
+       }
 
        @Override
-       int getInfoType() {
-               return TYPE_COMMAND_QUEUE;
+       MethodHandle getInfoFunc() {
+               return clGetCommandQueueInfo;
        }
 
        public static CLQueueProperty PROPERTIES(long flags) {
@@ -82,36 +107,92 @@ public class CLCommandQueue extends CLExtendable {
         *
         * @throws CLException
         */
-       public native void flush() throws CLException;
+       public void flush() throws CLException {
+               try {
+                       int res = clFlush(addr());
+                       if (res != 0)
+                               throw new CLException(res);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         * Calls clFinish.
         *
         * @throws CLException
         */
-       public native void finish() throws CLException;
+       public void finish() throws CLException {
+               try {
+                       int res = clFinish(addr());
+                       if (res != 0)
+                               throw new CLException(res);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
-        * Calls clEnqueueReadBuffer for a direct Buffer.
+        * Calls clEnqueueReadBuffer for a MemorySegment.
         *
         * @param mem
         * @param blocking
         * @param mem_offset source memory offset in bytes.
         * @param size memory transfer size in bytes.
-        * @param buffer destination buffer. Must be a direct buffer.
+        * @param buffer destination buffer.  It must be a native segment.
         * @param wait
         * @param event
         * @throws CLException
         */
-       native public void enqueueReadBuffer(CLBuffer mem, boolean blocking,
+       public void enqueueReadBuffer(CLBuffer mem, boolean blocking,
+               long mem_offset, long size,
+               MemorySegment buffer,
+               CLEventList wait,
+               CLEventList event) throws CLException {
+               if (size > buffer.byteSize())
+                       throw new BufferOverflowException();
+
+               try (Allocator frame = Memory.stack()) {
+                       EventInfo info = new EventInfo(frame, wait, event);
+                       int res;
+
+                       res = clEnqueueReadBuffer(addr(), mem.addr(), blocking ? 1 : 0,
+                               mem_offset, size,
+                               buffer.baseAddress(),
+                               info.nwait, info.wait, info.event);
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       info.post(event);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
+
+       public void enqueueReadBuffer(CLBuffer mem, boolean blocking,
                long mem_offset, long size,
                ByteBuffer buffer,
                CLEventList wait,
-               CLEventList event) throws CLException;
+               CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.ofByteBuffer(buffer)) {
+                       enqueueReadBuffer(mem, blocking,
+                               mem_offset, size,
+                               seg,
+                               wait, event);
+               }
+       }
 
        /**
         * Calls clEnqueueReadBuffer for a byte array.
         *
+        * panama note: the buffer must always be copied.
+        *
         * @param mem
         * @param blocking ignored, array reads are always blocking.
         * @param mem_offset source memory offset in bytes.
@@ -122,15 +203,22 @@ public class CLCommandQueue extends CLExtendable {
         * @param event
         * @throws CLException
         */
-       native public void enqueueReadBuffer(CLBuffer mem, boolean blocking,
-               long mem_offset, long size,
-               byte[] buffer, long buf_offset,
+       public void enqueueReadBuffer(CLBuffer mem, boolean blocking,
+               long mem_offset, int size,
+               byte[] buffer, int buf_offset,
                CLEventList wait,
-               CLEventList event) throws CLException;
+               CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.allocateNative(size)) {
+                       enqueueReadBuffer(mem, true, mem_offset, size, seg, wait, event);
+                       seg.asByteBuffer().order(ByteOrder.nativeOrder()).get(buffer, buf_offset, size);
+               }
+       }
 
        /**
         * Calls clEnqueueReadBuffer for a short array.
         *
+        * Note that mem_offset and size are in terms of the buffer element size.
+        *
         * @param mem
         * @param blocking ignored, array reads are always blocking.
         * @param mem_offset source memory offset in shorts.
@@ -141,35 +229,78 @@ public class CLCommandQueue extends CLExtendable {
         * @param event
         * @throws CLException
         */
-       native public void enqueueReadBuffer(CLBuffer mem, boolean blocking,
-               long mem_offset, long size,
-               short[] buffer, long buf_offset,
+       public void enqueueReadBuffer(CLBuffer mem, boolean blocking,
+               long mem_offset, int size,
+               short[] buffer, int buf_offset,
                CLEventList wait,
-               CLEventList event) throws CLException;
+               CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.allocateNative(size * 2)) {
+                       enqueueReadBuffer(mem, true, mem_offset * 2, size * 2, seg, wait, event);
+                       seg.asByteBuffer().order(ByteOrder.nativeOrder()).asShortBuffer().get(buffer, buf_offset, size);
+               }
+       }
 
-       native public void enqueueReadBuffer(CLBuffer mem, boolean blocking,
-               long mem_offset, long size,
-               int[] buffer, long buf_offset,
+       public void enqueueReadBuffer(CLBuffer mem, boolean blocking,
+               long mem_offset, int size,
+               int[] buffer, int buf_offset,
                CLEventList wait,
-               CLEventList event) throws CLException;
+               CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.allocateNative(size * 4)) {
+                       enqueueReadBuffer(mem, true, mem_offset * 4, size * 4, seg, wait, event);
+                       seg.asByteBuffer().order(ByteOrder.nativeOrder()).asIntBuffer().get(buffer, buf_offset, size);
+               }
+       }
 
-       native public void enqueueReadBuffer(CLBuffer mem, boolean blocking,
-               long mem_offset, long size,
-               long[] buffer, long buf_offset,
+       public void enqueueReadBuffer(CLBuffer mem, boolean blocking,
+               long mem_offset, int size,
+               long[] buffer, int buf_offset,
                CLEventList wait,
-               CLEventList event) throws CLException;
+               CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.allocateNative(size * 8)) {
+                       enqueueReadBuffer(mem, true, mem_offset * 8, size * 8, seg, wait, event);
+                       seg.asByteBuffer().order(ByteOrder.nativeOrder()).asLongBuffer().get(buffer, buf_offset, size);
+               }
+       }
 
-       native public void enqueueReadBuffer(CLBuffer mem, boolean blocking,
-               long mem_offset, long size,
-               float[] buffer, long buf_offset,
+       public void enqueueReadBuffer(CLBuffer mem, boolean blocking,
+               long mem_offset, int size,
+               float[] buffer, int buf_offset,
                CLEventList wait,
-               CLEventList event) throws CLException;
+               CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.allocateNative(size * 4)) {
+                       enqueueReadBuffer(mem, true, mem_offset * 4, size * 4, seg, wait, event);
+                       seg.asByteBuffer().order(ByteOrder.nativeOrder()).asFloatBuffer().get(buffer, buf_offset, size);
+               }
+       }
 
-       native public void enqueueReadBuffer(CLBuffer mem, boolean blocking,
-               long mem_offset, long size,
-               double[] buffer, long buf_offset,
+       public void enqueueReadBuffer(CLBuffer mem, boolean blocking,
+               long mem_offset, int size,
+               double[] buffer, int buf_offset,
                CLEventList wait,
-               CLEventList event) throws CLException;
+               CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.allocateNative(size * 8)) {
+                       enqueueReadBuffer(mem, true, mem_offset * 8, size * 8, seg, wait, event);
+                       seg.asByteBuffer().order(ByteOrder.nativeOrder()).asDoubleBuffer().get(buffer, buf_offset, size);
+               }
+       }
+
+       private void checkBufferRectRange(
+               long[] buffer_origin, long[] host_origin, long[] region,
+               long buffer_row_pitch, long buffer_slice_pitch, long host_row_pitch, long host_slice_pitch,
+               long limit) {
+               if (buffer_origin.length != 3
+                   || host_origin.length != 3
+                   || region.length != 3)
+                       throw new IllegalArgumentException("origin and region must be 3-dimensional");
+
+               long stride = host_row_pitch == 0 ? region[0] : host_row_pitch;
+               long slice = host_slice_pitch == 0 ? region[1] * stride : host_slice_pitch;
+
+               if (host_origin[0] + host_origin[1] * stride + host_origin[2] * slice
+                   + region[0] + (region[1]-1) * stride + (region[2]-1) * slice
+                   > limit)
+                       throw new BufferOverflowException();
+       }
 
        /**
         *
@@ -189,17 +320,89 @@ public class CLCommandQueue extends CLExtendable {
         * @throws UnsupportedOperationException
         * @since OpenCL 1.1
         */
-       public native void enqueueReadBufferRect(CLBuffer mem, boolean blocking,
+       public void enqueueReadBufferRect(CLBuffer mem, boolean blocking,
+               long[] buffer_origin, long[] host_origin, long[] region,
+               long buffer_row_pitch, long buffer_slice_pitch, long host_row_pitch, long host_slice_pitch,
+               MemorySegment buffer,
+               CLEventList wait, CLEventList event) throws CLException, UnsupportedOperationException {
+               checkBufferRectRange(
+                       buffer_origin, host_origin, region,
+                       buffer_row_pitch, buffer_slice_pitch, host_row_pitch, host_slice_pitch,
+                       buffer.byteSize());
+
+               try (Allocator frame = Memory.stack()) {
+                       EventInfo info = new EventInfo(frame, wait, event);
+                       MemoryAddress cbuffer_origin = toLongV(frame, buffer_origin);
+                       MemoryAddress chost_origin = toLongV(frame, host_origin);
+                       MemoryAddress cregion = toLongV(frame, region);
+                       int res;
+
+                       res = clEnqueueReadBufferRect(
+                               addr(), mem.addr(), blocking ? 1 : 0,
+                               cbuffer_origin, chost_origin, cregion,
+                               buffer_row_pitch,
+                               buffer_slice_pitch,
+                               host_row_pitch,
+                               host_slice_pitch,
+                               buffer.baseAddress(),
+                               info.nwait, info.wait, info.event);
+
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       info.post(event);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
+
+       public void enqueueReadBufferRect(CLBuffer mem, boolean blocking,
                long[] buffer_origin, long[] host_origin, long[] region,
                long buffer_row_pitch, long buffer_slice_pitch, long host_row_pitch, long host_slice_pitch,
                ByteBuffer buffer,
-               CLEventList waiters, CLEventList events) throws CLException, UnsupportedOperationException;
+               CLEventList waiters, CLEventList events) throws CLException, UnsupportedOperationException {
+
+               try (MemorySegment seg = MemorySegment.ofByteBuffer(buffer)) {
+                       enqueueReadBufferRect(mem, blocking,
+                               buffer_origin, host_origin, region,
+                               buffer_row_pitch, buffer_slice_pitch, host_row_pitch, host_slice_pitch,
+                               seg,
+                               waiters, events);
+               }
+       }
 
-       public native void enqueueReadBufferRect(CLBuffer mem, boolean blocking,
+       public void enqueueReadBufferRect(CLBuffer mem, boolean blocking,
                long[] buffer_origin, long[] host_origin, long[] region,
                long buffer_row_pitch, long buffer_slice_pitch, long host_row_pitch, long host_slice_pitch,
                byte[] buffer,
-               CLEventList waiters, CLEventList events) throws CLException, UnsupportedOperationException;
+               CLEventList waiters, CLEventList events) throws CLException, UnsupportedOperationException {
+               // This allocates only what is required and then copies row by row to the target host_origin.
+               // yikes this isn't very nice.
+               long stride = region[0];
+               long slice = region[1] * stride;
+               long size = slice * region[2];
+
+               try (MemorySegment seg = MemorySegment.allocateNative(size)) {
+                       enqueueReadBufferRect(mem, blocking,
+                               buffer_origin, OFFSET_0x0x0, region,
+                               buffer_row_pitch, buffer_slice_pitch, stride, slice,
+                               seg,
+                               waiters, events);
+                       long host_stride = host_row_pitch == 0 ? region[0] : host_row_pitch;
+                       long host_slice = host_slice_pitch == 0 ? region[1] * host_stride : host_slice_pitch;
+                       long host_offset = host_origin[0] + host_origin[1] * host_stride + host_origin[2] * host_slice;
+                       ByteBuffer bb = seg.asByteBuffer().order(ByteOrder.nativeOrder());
+
+                       for (long z = 0; z < region[2]; z++) {
+                               for (long y = 0; y < region[1]; y++) {
+                                       long o = host_offset + y * host_stride + z * host_slice;
+                                       bb.get(buffer, (int)o, (int)region[0]);
+                               }
+                       }
+               }
+       }
 
        public native void enqueueReadBufferRect(CLBuffer mem, boolean blocking,
                long[] buffer_origin, long[] host_origin, long[] region,
@@ -231,47 +434,111 @@ public class CLCommandQueue extends CLExtendable {
                double[] buffer,
                CLEventList waiters, CLEventList events) throws CLException, UnsupportedOperationException;
 
-       native public void enqueueWriteBuffer(CLBuffer mem, boolean blocking,
+       public void enqueueWriteBuffer(CLBuffer mem, boolean blocking,
                long mem_offset, long size,
-               ByteBuffer buffer,
+               MemorySegment buffer,
                CLEventList wait,
-               CLEventList event) throws CLException;
+               CLEventList event) throws CLException {
+               if (size > buffer.byteSize())
+                       throw new BufferUnderflowException();
+
+               try (Allocator frame = Memory.stack()) {
+                       EventInfo info = new EventInfo(frame, wait, event);
+                       int res;
+
+                       res = clEnqueueWriteBuffer(addr(), mem.addr(), blocking ? 1 : 0,
+                               mem_offset, size,
+                               buffer.baseAddress(),
+                               info.nwait, info.wait, info.event);
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       info.post(event);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
-       native public void enqueueWriteBuffer(CLBuffer mem, boolean blocking,
+       public void enqueueWriteBuffer(CLBuffer mem, boolean blocking,
                long mem_offset, long size,
-               byte[] buffer, long buf_offset,
+               ByteBuffer buffer,
                CLEventList wait,
-               CLEventList event) throws CLException;
+               CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.ofByteBuffer(buffer)) {
+                       enqueueWriteBuffer(mem, blocking,
+                               mem_offset, size,
+                               seg,
+                               wait, event);
+               }
+       }
 
-       native public void enqueueWriteBuffer(CLBuffer mem, boolean blocking,
-               long mem_offset, long size,
-               short[] buffer, long buf_offset,
+       public void enqueueWriteBuffer(CLBuffer mem, boolean blocking,
+               long mem_offset, int size,
+               byte[] buffer, int buf_offset,
                CLEventList wait,
-               CLEventList event) throws CLException;
+               CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.allocateNative(size)) {
+                       seg.asByteBuffer().order(ByteOrder.nativeOrder()).put(buffer, buf_offset, size);
+                       enqueueWriteBuffer(mem, true, mem_offset, size, seg, wait, event);
+               }
+       }
 
-       native public void enqueueWriteBuffer(CLBuffer mem, boolean blocking,
-               long mem_offset, long size,
-               int[] buffer, long buf_offset,
+       public void enqueueWriteBuffer(CLBuffer mem, boolean blocking,
+               long mem_offset, int size,
+               short[] buffer, int buf_offset,
                CLEventList wait,
-               CLEventList event) throws CLException;
+               CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.allocateNative(size * 2L)) {
+                       seg.asByteBuffer().order(ByteOrder.nativeOrder()).asShortBuffer().put(buffer, buf_offset, size);
+                       enqueueWriteBuffer(mem, true, mem_offset * 2, size * 2L, seg, wait, event);
+               }
+       }
 
-       native public void enqueueWriteBuffer(CLBuffer mem, boolean blocking,
-               long mem_offset, long size,
-               long[] buffer, long buf_offset,
+       public void enqueueWriteBuffer(CLBuffer mem, boolean blocking,
+               long mem_offset, int size,
+               int[] buffer, int buf_offset,
                CLEventList wait,
-               CLEventList event) throws CLException;
+               CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.allocateNative(size * 4L)) {
+                       seg.asByteBuffer().order(ByteOrder.nativeOrder()).asIntBuffer().put(buffer, buf_offset, size);
+                       enqueueWriteBuffer(mem, true, mem_offset * 4, size * 4L, seg, wait, event);
+               }
+       }
 
-       native public void enqueueWriteBuffer(CLBuffer mem, boolean blocking,
-               long mem_offset, long size,
-               float[] buffer, long buf_offset,
+       public void enqueueWriteBuffer(CLBuffer mem, boolean blocking,
+               long mem_offset, int size,
+               long[] buffer, int buf_offset,
                CLEventList wait,
-               CLEventList event) throws CLException;
+               CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.allocateNative(size * 8L)) {
+                       seg.asByteBuffer().order(ByteOrder.nativeOrder()).asLongBuffer().put(buffer, buf_offset, size);
+                       enqueueWriteBuffer(mem, true, mem_offset * 8, size * 8L, seg, wait, event);
+               }
+       }
 
-       native public void enqueueWriteBuffer(CLBuffer mem, boolean blocking,
-               long mem_offset, long size,
-               double[] buffer, long buf_offset,
+       public void enqueueWriteBuffer(CLBuffer mem, boolean blocking,
+               long mem_offset, int size,
+               float[] buffer, int buf_offset,
                CLEventList wait,
-               CLEventList event) throws CLException;
+               CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.allocateNative(size * 4L)) {
+                       seg.asByteBuffer().order(ByteOrder.nativeOrder()).asFloatBuffer().put(buffer, buf_offset, size);
+                       enqueueWriteBuffer(mem, true, mem_offset * 4, size * 4L, seg, wait, event);
+               }
+       }
+
+       public void enqueueWriteBuffer(CLBuffer mem, boolean blocking,
+               long mem_offset, int size,
+               double[] buffer, int buf_offset,
+               CLEventList wait,
+               CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.allocateNative(size * 8L)) {
+                       seg.asByteBuffer().order(ByteOrder.nativeOrder()).asDoubleBuffer().put(buffer, buf_offset, size);
+                       enqueueWriteBuffer(mem, true, mem_offset * 8, size * 8L, seg, wait, event);
+               }
+       }
 
        /**
         *
@@ -291,11 +558,60 @@ public class CLCommandQueue extends CLExtendable {
         * @throws UnsupportedOperationException
         * @since OpenCL 1.1
         */
-       public native void enqueueWriteBufferRect(CLBuffer mem, boolean blocking,
+       public void enqueueWriteBufferRect(CLBuffer mem, boolean blocking,
+               long[] buffer_origin, long[] host_origin, long[] region,
+               long buffer_row_pitch, long buffer_slice_pitch, long host_row_pitch, long host_slice_pitch,
+               MemorySegment buffer,
+               CLEventList wait, CLEventList event) throws CLException, UnsupportedOperationException {
+
+               requireAPIVersion(CLPlatform.VERSION_1_1);
+
+               checkBufferRectRange(
+                       buffer_origin, host_origin, region,
+                       buffer_row_pitch, buffer_slice_pitch, host_row_pitch, host_slice_pitch,
+                       buffer.byteSize());
+
+               try (Allocator frame = Memory.stack()) {
+                       EventInfo info = new EventInfo(frame, wait, event);
+                       MemoryAddress cbuffer_origin = toLongV(frame, buffer_origin);
+                       MemoryAddress chost_origin = toLongV(frame, host_origin);
+                       MemoryAddress cregion = toLongV(frame, region);
+                       int res;
+
+                       res = clEnqueueWriteBufferRect(
+                               addr(), mem.addr(), blocking ? 1 : 0,
+                               cbuffer_origin, chost_origin, cregion,
+                               buffer_row_pitch,
+                               buffer_slice_pitch,
+                               host_row_pitch,
+                               host_slice_pitch,
+                               buffer.baseAddress(),
+                               info.nwait, info.wait, info.event);
+
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       info.post(event);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
+
+       public void enqueueWriteBufferRect(CLBuffer mem, boolean blocking,
                long[] buffer_origin, long[] host_origin, long[] region,
                long buffer_row_pitch, long buffer_slice_pitch, long host_row_pitch, long host_slice_pitch,
                ByteBuffer buffer,
-               CLEventList waiters, CLEventList events) throws CLException, UnsupportedOperationException;
+               CLEventList waiters, CLEventList events) throws CLException, UnsupportedOperationException {
+               try (MemorySegment seg = MemorySegment.ofByteBuffer(buffer)) {
+                       enqueueWriteBufferRect(mem, blocking,
+                               buffer_origin, host_origin, region,
+                               buffer_row_pitch, buffer_slice_pitch, host_row_pitch, host_slice_pitch,
+                               seg,
+                               waiters, events);
+               }
+       }
 
        public native void enqueueWriteBufferRect(CLBuffer mem, boolean blocking,
                long[] buffer_origin, long[] host_origin, long[] region,
@@ -334,7 +650,10 @@ public class CLCommandQueue extends CLExtendable {
                CLEventList waiters, CLEventList events) throws CLException, UnsupportedOperationException;
 
        /**
-        * Calls clEnqueueFillBuffer for byte types.
+        * Calls clEnqueueFillBuffer.
+        *
+        * This fills the buffer with a pattern.  The offset and size
+        * is in multiples of the pattern size.
         *
         * @param buffer
         * @param pattern pattern to fill
@@ -345,12 +664,75 @@ public class CLCommandQueue extends CLExtendable {
         * @since OpenCL 1.2
         * @throws CLException
         */
-       public native void enqueueFillBuffer(CLBuffer buffer,
-               byte[] pattern,
+       public void enqueueFillBuffer(CLBuffer buffer,
+                                     MemorySegment pattern,
+                                     long offset,
+                                     long size,
+                                     CLEventList wait,
+                                     CLEventList event) throws CLException {
+               requireAPIVersion(CLPlatform.VERSION_1_2);
+               try (Allocator frame = Memory.stack()) {
+                       EventInfo info = new EventInfo(frame, wait, event);
+                       int res;
+
+                       res = clEnqueueFillBuffer(
+                               addr(), buffer.addr(),
+                               pattern.baseAddress(), pattern.byteSize(),
+                               offset * pattern.byteSize(),
+                               size * pattern.byteSize(),
+                               info.nwait, info.wait, info.event);
+
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       info.post(event);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
+       public void enqueueFillBuffer(CLBuffer buffer,
+               ByteBuffer pattern,
                long offset,
                long size,
-               CLEventList waiters,
-               CLEventList events) throws CLException;
+               CLEventList wait,
+               CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.ofByteBuffer(pattern)) {
+                       enqueueFillBuffer(buffer,
+                               pattern,
+                               offset, size,
+                               wait, event);
+               }
+       }
+
+       /**
+        * Calls clEnqueueFillBuffer for byte types.
+        *
+        * @param buffer
+        * @param pattern pattern to fill
+        * @param offset offset in multiples of the pattern size.
+        * @param size number of elements in multiples of the pattern size.
+        * @param waiters
+        * @param events
+        * @since OpenCL 1.2
+        * @throws CLException
+        */
+       public void enqueueFillBuffer(CLBuffer buffer,
+                                     byte[] pattern,
+                                     long offset,
+                                     long size,
+                                     CLEventList wait,
+                                     CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.allocateNative(pattern.length)) {
+                       MemoryAddress add = seg.baseAddress();
+
+                       for (int i=0;i<pattern.length;i++)
+                               setByte(add, i, pattern[i]);
+
+                       enqueueFillBuffer(buffer, seg, offset, size, wait, event);
+               }
+       }
 
        /**
         * Calls clEnqueueFillBuffer for short types.
@@ -364,12 +746,21 @@ public class CLCommandQueue extends CLExtendable {
         * @since OpenCL 1.2
         * @throws CLException
         */
-       public native void enqueueFillBuffer(CLBuffer buffer,
-               short[] pattern,
-               long offset,
-               long size,
-               CLEventList waiters,
-               CLEventList events) throws CLException;
+       public void enqueueFillBuffer(CLBuffer buffer,
+                                            short[] pattern,
+                                            long offset,
+                                            long size,
+                                            CLEventList wait,
+                                            CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.allocateNative(pattern.length * 2)) {
+                       MemoryAddress add = seg.baseAddress();
+
+                       for (int i=0;i<pattern.length;i++)
+                               setShort(add, i, pattern[i]);
+
+                       enqueueFillBuffer(buffer, seg, offset * 2, size * 2, wait, event);
+               }
+       }
 
        /**
         * Calls clEnqueueFillBuffer for integer types.
@@ -383,12 +774,21 @@ public class CLCommandQueue extends CLExtendable {
         * @since OpenCL 1.2
         * @throws CLException
         */
-       public native void enqueueFillBuffer(CLBuffer buffer,
-               int[] pattern,
-               long offset,
-               long size,
-               CLEventList waiters,
-               CLEventList events) throws CLException;
+       public void enqueueFillBuffer(CLBuffer buffer,
+                                            int[] pattern,
+                                            long offset,
+                                            long size,
+                                            CLEventList wait,
+                                            CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.allocateNative(pattern.length * 4)) {
+                       MemoryAddress add = seg.baseAddress();
+
+                       for (int i=0;i<pattern.length;i++)
+                               setInt(add, i, pattern[i]);
+
+                       enqueueFillBuffer(buffer, seg, offset * 4, size * 4, wait, event);
+               }
+       }
 
        /**
         * Calls clEnqueueFillBuffer for long types.
@@ -414,19 +814,28 @@ public class CLCommandQueue extends CLExtendable {
         *
         * @param buffer
         * @param pattern pattern to fill
-        * @param offset offset in multiples of the pattern size.
-        * @param size number of elements in multiples of the pattern size.
+        * @param offset offset in multiples of the pattern size in floats.
+        * @param size number of elements in multiples of the pattern size in floats.
         * @param waiters
         * @param events
         * @since OpenCL 1.2
         * @throws CLException
         */
-       public native void enqueueFillBuffer(CLBuffer buffer,
-               float[] pattern,
-               long offset,
-               long size,
-               CLEventList waiters,
-               CLEventList events) throws CLException;
+       public void enqueueFillBuffer(CLBuffer buffer,
+                                     float[] pattern,
+                                     long offset,
+                                     long size,
+                                     CLEventList wait,
+                                     CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.allocateNative(pattern.length * 4)) {
+                       MemoryAddress add = seg.baseAddress();
+
+                       for (int i=0;i<pattern.length;i++)
+                               setFloat(add, i, pattern[i]);
+
+                       enqueueFillBuffer(buffer, seg, offset * 4, size * 4, wait, event);
+               }
+       }
 
        /**
         * Calls clEnqueueFillBuffer for double types.
@@ -447,9 +856,29 @@ public class CLCommandQueue extends CLExtendable {
                CLEventList waiters,
                CLEventList events) throws CLException;
 
-       public native void enqueueCopyBuffer(CLBuffer srcmem, CLBuffer dstmem, long srcoffset, long dstoffset, long size,
-               CLEventList waiters,
-               CLEventList events) throws CLException;
+       public void enqueueCopyBuffer(CLBuffer srcmem, CLBuffer dstmem, long srcoffset, long dstoffset, long size,
+               CLEventList wait,
+               CLEventList event) throws CLException {
+
+               try (Allocator frame = Memory.stack()) {
+                       EventInfo info = new EventInfo(frame, wait, event);
+                       int res;
+
+                       res = clEnqueueCopyBuffer(
+                               addr(), srcmem.addr(), dstmem.addr(),
+                               srcoffset, dstoffset, size,
+                               info.nwait, info.wait, info.event);
+
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       info.post(event);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         * @param srcmem
@@ -467,7 +896,7 @@ public class CLCommandQueue extends CLExtendable {
         * @throws UnsupportedOperationException
         * @since OpenCL 1.1
         */
-       public native void enqueueCopyBufferRect(CLBuffer srcmem, CLBuffer dstmem,
+       public void enqueueCopyBufferRect(CLBuffer srcmem, CLBuffer dstmem,
                long[] src_origin,
                long[] dst_origin,
                long[] region,
@@ -476,7 +905,115 @@ public class CLCommandQueue extends CLExtendable {
                long dst_row_pitch,
                long dst_slice_pitch,
                CLEventList waiters,
-               CLEventList events) throws CLException, UnsupportedOperationException;
+               CLEventList events) throws CLException, UnsupportedOperationException {
+       }
+
+       /*
+        * Parameterised image<>buffer copy function, for read or write
+        */
+       private void doImageBuffer(CLImage image, boolean blocking,
+               long[] origin,
+               long[] region,
+               long row_pitch,
+               long slice_pitch,
+               MemorySegment buffer,
+               CLEventList wait,
+               CLEventList event,
+               MethodHandle op) throws CLException {
+
+               if (origin.length != 3
+                   || region.length != 3)
+                       throw new IllegalArgumentException("origin and region must be 3-dimensional");
+
+               long stride = row_pitch == 0 ? region[0] * image.getElementSize() : row_pitch;
+               long slice = slice_pitch == 0 ? region[1] * stride : slice_pitch;
+
+               if (buffer.byteSize() < slice * region[2])
+                       throw new BufferOverflowException();
+
+               try (Allocator frame = Memory.stack()) {
+                       EventInfo info = new EventInfo(frame, wait, event);
+                       MemoryAddress corigin = toLongV(frame, origin);
+                       MemoryAddress cregion = toLongV(frame, region);
+                       int res;
+
+                       res = (int)op.invokeExact(
+                               addr(), image.addr(), blocking ? 1 : 0,
+                               corigin, cregion,
+                               row_pitch, slice_pitch,
+                               buffer.baseAddress(),
+                               info.nwait, info.wait, info.event);
+
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       info.post(event);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
+
+       private interface CopyAppend {
+               void apply(int doff, int size);
+
+               default void copy(int w, int h, int d, int stride, int slice) {
+                       for (int z=0;z<d;z++) {
+                               for (int y=0;y<h;y++) {
+                                       apply(y * stride + z * slice, w);
+                               }
+                       }
+               }
+       }
+
+       /*
+        * Parameterised image<>buffer copy function, for read or write
+        * of primitive types.
+        * This first copies to the exact size required
+        * and then copies to the target.
+        */
+       private void doImageBufferPrimitive(
+               CLImage image, boolean blocking,
+               long[] origin,
+               long[] region,
+               long row_pitch,
+               long slice_pitch,
+               boolean readMode,
+               Function<MemorySegment, CopyAppend> writer,
+               long write_size,
+               long dshift,
+               MethodHandle op,
+               CLEventList wait,
+               CLEventList event) throws CLException {
+               //  Just transfer the minimum size, then copy to target
+               long elsize = image.getElementSize();
+
+               // these sizes are in array elmeents
+               long stride = row_pitch == 0 ? (region[0] * elsize) >> dshift : row_pitch;
+               long slice = slice_pitch == 0 ? (region[1] * stride) >> dshift : slice_pitch;
+
+               // Test target size
+               if (slice * region[2] > write_size)
+                       throw new BufferOverflowException();
+
+               // these are in bytes
+               long alloc = (region[0] * region[1] * region[2]) << dshift;
+               long xstride = region[0] << dshift;
+               long xslice = xstride * region[1];
+
+               try (MemorySegment seg = MemorySegment.allocateNative(alloc)) {
+                       if (!readMode)
+                               writer.apply(seg).copy((int)region[0], (int)region[1], (int)region[2],(int)stride, (int)slice);
+                       doImageBuffer(image, blocking, origin, region, xstride, 0, seg, wait, event, op);
+                       if (readMode)
+                               writer.apply(seg).copy((int)region[0], (int)region[1], (int)region[2],(int)stride, (int)slice);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         *
@@ -491,14 +1028,36 @@ public class CLCommandQueue extends CLExtendable {
         * @param events
         * @throws CLException
         */
-       public native void enqueueReadImage(CLImage image, boolean blocking,
+       public void enqueueReadImage(CLImage image, boolean blocking,
+               long[] origin,
+               long[] region,
+               long row_pitch,
+               long slice_pitch,
+               MemorySegment buffer,
+               CLEventList wait,
+               CLEventList event) throws CLException {
+               doImageBuffer(image, blocking, origin, region, row_pitch, slice_pitch, buffer, wait, event, clEnqueueReadImage);
+       }
+
+       /**
+        * @see #enqueueReadImage(CLImage,boolean,long[],long[],long,long,MemorySegment,CLEventList,CLEventList)
+        */
+       public void enqueueReadImage(CLImage image, boolean blocking,
                long[] origin,
                long[] region,
                long row_pitch,
                long slice_pitch,
                ByteBuffer buffer,
-               CLEventList waiters,
-               CLEventList events) throws CLException;
+               CLEventList wait,
+               CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.ofByteBuffer(buffer)) {
+                       enqueueReadImage(image, blocking,
+                               origin, region,
+                               row_pitch, slice_pitch,
+                               seg,
+                               wait, event);
+               }
+       }
 
        /**
         *
@@ -514,23 +1073,50 @@ public class CLCommandQueue extends CLExtendable {
         * @param events
         * @throws CLException
         */
-       public native void enqueueReadImage(CLImage image, boolean blocking,
+       public void enqueueReadImage(CLImage image, boolean blocking,
                long[] origin,
                long[] region,
                long row_pitch,
                long slice_pitch,
-               byte[] buffer, long buff_offset,
-               CLEventList waiters,
-               CLEventList events) throws CLException;
+               byte[] buffer, int buff_offset,
+               CLEventList wait,
+               CLEventList event) throws CLException {
+
+               doImageBufferPrimitive(image, blocking, origin, region, row_pitch, slice_pitch,
+                       true,
+                       (seg) -> {
+                               var sb = seg.asByteBuffer();
+                               return (off, size) -> sb.get(buffer, buff_offset + off, size);
+                       },
+                       buff_offset - buffer.length,
+                       0,
+                       clEnqueueReadImage,
+                       wait, event);
+       }
 
-       public native void enqueueReadImage(CLImage image, boolean blocking,
+       /**
+        * @see #enqueueReadImage(CLImage,boolean,long[],long[],long,long,byte[],int,CLEventList,CLEventList)
+        */
+       public void enqueueReadImage(CLImage image, boolean blocking,
                long[] origin,
                long[] region,
                long row_pitch,
                long slice_pitch,
-               short[] buffer, long buff_offset,
-               CLEventList waiters,
-               CLEventList events) throws CLException;
+               short[] buffer, int buff_offset,
+               CLEventList wait,
+               CLEventList event) throws CLException {
+
+               doImageBufferPrimitive(image, blocking, origin, region, row_pitch, slice_pitch,
+                       true,
+                       (seg) -> {
+                               var sb = seg.asByteBuffer().order(ByteOrder.nativeOrder()).asShortBuffer();
+                               return (off, size) -> sb.get(buffer, buff_offset + off, size);
+                       },
+                       buff_offset - buffer.length,
+                       1,
+                       clEnqueueReadImage,
+                       wait, event);
+       }
 
        public native void enqueueReadImage(CLImage image, boolean blocking,
                long[] origin,
@@ -550,14 +1136,25 @@ public class CLCommandQueue extends CLExtendable {
                CLEventList waiters,
                CLEventList events) throws CLException;
 
-       public native void enqueueReadImage(CLImage image, boolean blocking,
+       public void enqueueReadImage(CLImage image, boolean blocking,
                long[] origin,
                long[] region,
                long row_pitch,
                long slice_pitch,
-               float[] buffer, long buff_offset,
-               CLEventList waiters,
-               CLEventList events) throws CLException;
+               float[] buffer, int buff_offset,
+               CLEventList wait,
+               CLEventList event) throws CLException {
+               doImageBufferPrimitive(image, blocking, origin, region, row_pitch, slice_pitch,
+                       true,
+                       (seg) -> {
+                               var sb = seg.asByteBuffer().order(ByteOrder.nativeOrder()).asFloatBuffer();
+                               return (off, size) -> sb.get(buffer, buff_offset + off, size);
+                       },
+                       buff_offset - buffer.length,
+                       2,
+                       clEnqueueReadImage,
+                       wait, event);
+       }
 
        public native void enqueueReadImage(CLImage image, boolean blocking,
                long[] origin,
@@ -568,23 +1165,54 @@ public class CLCommandQueue extends CLExtendable {
                CLEventList waiters,
                CLEventList events) throws CLException;
 
-       public native void enqueueWriteImage(CLImage image, boolean blocking,
+       public void enqueueWriteImage(CLImage image, boolean blocking,
+               long[] origin,
+               long[] region,
+               long row_pitch,
+               long slice_pitch,
+               MemorySegment buffer,
+               CLEventList wait,
+               CLEventList event) throws CLException {
+               doImageBuffer(image, blocking, origin, region, row_pitch, slice_pitch, buffer, wait, event, clEnqueueWriteImage);
+       }
+
+       public void enqueueWriteImage(CLImage image, boolean blocking,
                long[] origin,
                long[] region,
                long row_pitch,
                long slice_pitch,
                ByteBuffer buffer,
-               CLEventList waiters,
-               CLEventList events) throws CLException;
+               CLEventList wait,
+               CLEventList event) throws CLException {
+               try (MemorySegment seg = MemorySegment.ofByteBuffer(buffer)) {
+                       enqueueWriteImage(image, blocking,
+                               origin, region,
+                               row_pitch, slice_pitch,
+                               seg,
+                               wait, event);
+               }
+       }
 
-       public native void enqueueWriteImage(CLImage image, boolean blocking,
+       public void enqueueWriteImage(CLImage image, boolean blocking,
                long[] origin,
                long[] region,
                long row_pitch,
                long slice_pitch,
-               byte[] buffer, long buff_offset,
-               CLEventList waiters,
-               CLEventList events) throws CLException;
+               byte[] buffer, int buff_offset,
+               CLEventList wait,
+               CLEventList event) throws CLException {
+
+               doImageBufferPrimitive(image, blocking, origin, region, row_pitch, slice_pitch,
+                       false,
+                       (seg) -> {
+                               var sb = seg.asByteBuffer();
+                               return (off, size) -> sb.put(buffer, buff_offset + off, size);
+                       },
+                       buff_offset - buffer.length,
+                       0,
+                       clEnqueueWriteImage,
+                       wait, event);
+       }
 
        public native void enqueueWriteImage(CLImage image, boolean blocking,
                long[] origin,
@@ -613,14 +1241,25 @@ public class CLCommandQueue extends CLExtendable {
                CLEventList waiters,
                CLEventList events) throws CLException;
 
-       public native void enqueueWriteImage(CLImage image, boolean blocking,
+       public void enqueueWriteImage(CLImage image, boolean blocking,
                long[] origin,
                long[] region,
                long row_pitch,
                long slice_pitch,
-               float[] buffer, long buff_offset,
-               CLEventList waiters,
-               CLEventList events) throws CLException;
+               float[] buffer, int buff_offset,
+               CLEventList wait,
+               CLEventList event) throws CLException {
+               doImageBufferPrimitive(image, blocking, origin, region, row_pitch, slice_pitch,
+                       false,
+                       (seg) -> {
+                               var sb = seg.asByteBuffer().order(ByteOrder.nativeOrder()).asFloatBuffer();
+                               return (off, size) -> sb.put(buffer, buff_offset + off, size);
+                       },
+                       buff_offset - buffer.length,
+                       2,
+                       clEnqueueWriteImage,
+                       wait, event);
+       }
 
        public native void enqueueWriteImage(CLImage image, boolean blocking,
                long[] origin,
@@ -692,25 +1331,163 @@ public class CLCommandQueue extends CLExtendable {
                CLEventList waiters,
                CLEventList events) throws CLException;
 
-       public native ByteBuffer enqueueMapBuffer(CLBuffer buffer, boolean blocking,
-               long flags,
-               long offset,
-               long size,
-               CLEventList waiters,
-               CLEventList events) throws CLException;
+       static class MapData {
+               MemoryAddress raw;
+               ByteBuffer buffer;
 
-       public native ByteBuffer enqueueMapImage(CLImage image, boolean blocking,
-               long flags,
-               long[] origin,
-               long[] region,
-               long[] image_row_pitch,
-               long[] image_slice_pitch,
-               CLEventList waiters,
-               CLEventList events) throws CLException;
+               public MapData(MemoryAddress p, ByteBuffer buffer) {
+                       this.raw = p;
+                       this.buffer = buffer;
+               }
+       }
 
-       public native void enqueueUnmapMemObject(CLMemory mem, ByteBuffer mapped,
-               CLEventList waiters,
-               CLEventList events) throws CLException;
+       private final ArrayList<MapData> maps = new ArrayList<>();
+
+       private void addMap(MemoryAddress raw, ByteBuffer bb) {
+               synchronized (maps) {
+                       maps.add(new MapData(raw, bb));
+               }
+       }
+
+       private MemoryAddress getMap(ByteBuffer bb) {
+               synchronized (maps) {
+                       // Note this can't use a hashtable as ByteBuffer
+                       // hashes on state like position.
+                       for (int i=0;i<maps.size();i++) {
+                               MapData d = maps.get(i);
+
+                               if (d.buffer == bb) {
+                                       maps.remove(i);
+                                       return d.raw;
+                               }
+                       }
+               }
+               throw new IllegalArgumentException();
+       }
+
+       public ByteBuffer enqueueMapBuffer(CLBuffer buffer, boolean blocking,
+                                          long flags,
+                                          long offset,
+                                          long size,
+                                          CLEventList wait,
+                                          CLEventList event) throws CLException {
+               try (Allocator frame = Memory.stack()) {
+                       EventInfo info = new EventInfo(frame, wait, event);
+                       MemoryAddress cres = frame.alloca(8);
+                       MemoryAddress cmap;
+                       int res;
+
+                       cmap = clEnqueueMapBuffer(addr(), buffer.addr(), blocking ? 1 : 0,
+                                                 flags, offset, size,
+                                                 info.nwait, info.wait, info.event,
+                                                 cres);
+                       res = getInt(cres);
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       // Need to map the segment and track it separately
+                       ByteBuffer bb = Memory.ofNative(cmap, size).asByteBuffer();
+
+                       addMap(cmap, bb);
+
+                       info.post(event);
+
+                       return bb;
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
+
+       /* There is some scope for improvement in this shittastic interface */
+       public ByteBuffer enqueueMapImage(CLImage image, boolean blocking,
+                                                long flags,
+                                                long[] origin,
+                                                long[] region,
+                                                long[] image_row_pitch,
+                                                long[] image_slice_pitch,
+                                                CLEventList wait,
+                                                CLEventList event) throws CLException {
+               if (origin.length != 3 || region.length != 3)
+                       throw new IllegalArgumentException("origin and region must be 3-dimensional");
+               if (image_row_pitch.length < 1)
+                       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 (Allocator frame = Memory.stack()) {
+                       EventInfo info = new EventInfo(frame, wait, event);
+                       MemoryAddress corigin = toLongV(frame, origin);
+                       MemoryAddress cregion = toLongV(frame, region);
+                       MemoryAddress cstride = frame.alloca(8);
+                       MemoryAddress cslice = frame.alloca(8);
+                       MemoryAddress cres = frame.alloca(8);
+                       MemoryAddress cmap;
+                       long stride, slice;
+                       int res;
+
+                       cmap = clEnqueueMapImage(
+                               addr(), image.addr(), blocking ? 1 : 0,
+                               flags,
+                               corigin, cregion,
+                               cstride, cslice,
+                               info.nwait, info.wait, info.event,
+                               cres);
+                       res = getInt(cres);
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       stride = getLong(cstride);
+                       slice = getLong(cslice);
+
+                       image_row_pitch[0] = stride;
+                       if (image_slice_pitch != null)
+                               image_slice_pitch[0] = slice;
+
+                       // These look weird but that's what the docs say
+                       long size = slice == 0
+                                   ?                     stride * region[1] + region[0] // 2D
+                                   : slice * region[2] + stride * region[1] + region[0]; // 3D
+
+                       // Need to map the segment and track it separately
+                       ByteBuffer bb = Memory.ofNative(cmap, size).asByteBuffer();
+
+                       addMap(cmap, bb);
+
+                       info.post(event);
+
+                       return bb;
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
+
+       public void enqueueUnmapMemObject(CLMemory mem, ByteBuffer mapped,
+                                         CLEventList wait,
+                                         CLEventList event) throws CLException {
+               MemoryAddress cmap = getMap(mapped);
+
+               try (Allocator frame = Memory.stack()) {
+                       EventInfo info = new EventInfo(frame, wait, event);
+                       int res;
+
+                       res = clEnqueueUnmapMemObject(
+                               addr(), mem.addr(),
+                               cmap,
+                               info.nwait, info.wait, info.event);
+
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       info.post(event);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         *
@@ -725,6 +1502,39 @@ public class CLCommandQueue extends CLExtendable {
                CLEventList waiters,
                CLEventList events) throws CLException;
 
+       private MemoryAddress wsToLongV(Allocator frame, long[] ws, int dim) {
+               if (ws != null) {
+                       if (ws.length < dim)
+                               throw new IllegalArgumentException();
+                       return toLongV(frame, ws);
+               } else {
+                       return MemoryAddress.NULL;
+               }
+       }
+
+       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;
+                       if (nwait > 0) {
+                               wait = frame.alloca(8 * nwait);
+                               for (int i=0;i<nwait;i++)
+                                       Native.setAddr(wait, i, waiters.get(i).addr());
+                       } else {
+                               wait = MemoryAddress.NULL;
+                       }
+                       event = events != null ? frame.alloca(8) : MemoryAddress.NULL;
+               }
+
+               void post(CLEventList events) {
+                       if (events != null)
+                               events.add(resolve(getAddr(event), CLEvent::new));
+               }
+       }
+
        /**
         * Call clEnqueueNDRangeKernel.
         *
@@ -739,12 +1549,33 @@ public class CLCommandQueue extends CLExtendable {
         * @param events
         * @throws CLException
         */
-       public native void enqueueNDRangeKernel(CLKernel kernel, int work_dim,
-               long[] global_offset,
-               long[] global_size,
-               long[] local_size,
-               CLEventList waiters,
-               CLEventList events) throws CLException;
+       public void enqueueNDRangeKernel(CLKernel kernel,
+                                        int work_dim,
+                                        long[] global_offset,
+                                        long[] global_size,
+                                        long[] local_size,
+                                        CLEventList waiters,
+                                        CLEventList events) throws CLException {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress gwo = wsToLongV(frame, global_offset, work_dim);
+                       MemoryAddress gws = wsToLongV(frame, global_size, work_dim);
+                       MemoryAddress lws = wsToLongV(frame, local_size, work_dim);
+                       EventInfo info = new EventInfo(frame, waiters, events);
+                       int res;
+
+                       res = clEnqueueNDRangeKernel(addr(), kernel.addr(),
+                                                    work_dim, gwo, gws, lws,
+                                                    info.nwait, info.wait, info.event);
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       info.post(events);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         *
@@ -752,12 +1583,28 @@ public class CLCommandQueue extends CLExtendable {
         * @param waiters
         * @param events
         * @throws CLException
-        * @deprecated as of OpenCL 2.0
+        * x-deprecated as of OpenCL 2.0
         */
-       @Deprecated
-       public native void enqueueTask(CLKernel kernel,
-               CLEventList waiters,
-               CLEventList events) throws CLException;
+       //@Deprecated
+       public void enqueueTask(CLKernel kernel,
+                               CLEventList waiters,
+                               CLEventList events) throws CLException {
+               try (Allocator frame = Memory.stack()) {
+                       EventInfo info = new EventInfo(frame, waiters, events);
+                       int res;
+
+                       res = clEnqueueTask(addr(), kernel.addr(),
+                                           info.nwait, info.wait, info.event);
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       info.post(events);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         * Invoke a native kernel.
@@ -769,11 +1616,61 @@ public class CLCommandQueue extends CLExtendable {
         * be replaced by a ByteBuffer in the kernel.
         * @throws CLException
         */
-       public native void enqueueNativeKernel(
+       public void enqueueNativeKernel(
                CLNativeKernel kernel,
                CLEventList waiters,
                CLEventList events,
-               Object... args) throws CLException;
+               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 (Allocator frame = Memory.stack()) {
+                       EventInfo info = new EventInfo(frame, waiters, events);
+                       MemoryAddress memstage = frame.alloca(8 * args.length);
+                       MemoryAddress memptrs = frame.alloca(8 * args.length);
+                       int nmem = 0;
+                       int res;
+
+                       Object[] save = args.clone();
+
+                       for (Object a: args) {
+                               if (a instanceof CLMemory) {
+                                       setAddr(memstage, nmem, ((CLMemory)a).addr());
+                                       setAddr(memptrs, nmem, memstage.addOffset(nmem * 8));
+                                       nmem++;
+                               }
+                       }
+
+                       Callback<CLNativeKernel> call = Native.resolve(
+                               Call_pv_v.stub((MemoryAddress memargs) -> {
+                                               int xmem = 0;
+
+                                               for (int i=0;i<args.length;i++) {
+                                                       if (args[i] instanceof CLMemory) {
+                                                               MemoryAddress mem = getAddr(memargs, xmem);
+                                                               long size = ((CLMemory)args[i]).getSize();
+
+                                                               save[i] = Memory.ofNative(mem, size).asByteBuffer().order(ByteOrder.nativeOrder());
+                                                               xmem++;
+                                                       }
+                                               }
+
+                                               kernel.invoke(save);
+                                       }),
+                               (p) -> new Callback<>(p, kernel));
+
+                       res = clEnqueueNativeKernel(addr(), call.addr(), memstage, nmem * 8, nmem, memstage, memptrs, info.nwait, info.wait, info.event);
+
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       info.post(events);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+
+       }
 
        /**
         * Enqueues a marker point.
@@ -785,9 +1682,24 @@ public class CLCommandQueue extends CLExtendable {
         * @param events
         * @throws CLException
         */
-       public native void enqueueMarkerWithWaitList(
+       public void enqueueMarkerWithWaitList(
                CLEventList waiters,
-               CLEventList events) throws CLException;
+               CLEventList events) throws CLException {
+               try (Allocator frame = Memory.stack()) {
+                       EventInfo info = new EventInfo(frame, waiters, events);
+                       if (haveAPIVersion(CLPlatform.VERSION_1_2)) {
+                               clEnqueueMarkerWithWaitList(addr(), info.nwait, info.wait, info.event);
+                       } else {
+                               clEnqueueWaitForEvents(addr(), info.nwait, info.wait);
+                               clEnqueueMarker(addr(), info.event);
+                       }
+                       info.post(events);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         * Enqueues a barrier.
@@ -800,9 +1712,25 @@ public class CLCommandQueue extends CLExtendable {
         * @param events
         * @throws CLException
         */
-       public native void enqueueBarrierWithWaitList(
+       public  void enqueueBarrierWithWaitList(
                CLEventList waiters,
-               CLEventList events) throws CLException;
+               CLEventList events) throws CLException {
+               try (Allocator frame = Memory.stack()) {
+                       EventInfo info = new EventInfo(frame, waiters, events);
+                       if (haveAPIVersion(CLPlatform.VERSION_1_2)) {
+                               clEnqueueBarrierWithWaitList(addr(), info.nwait, info.wait, info.event);
+                       } else {
+                               clEnqueueWaitForEvents(addr(), info.nwait, info.wait);
+                               clEnqueueBarrier(addr());
+                               clEnqueueMarker(addr(), info.event);
+                       }
+                       info.post(events);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         * Call clEnqueueSVMFree().
@@ -1002,11 +1930,11 @@ public class CLCommandQueue extends CLExtendable {
        }
 
        public CLContext getContext() {
-               return getInfoAny(CTYPE_CONTEXT, CL_QUEUE_CONTEXT);
+               return getInfoAny(CL_QUEUE_CONTEXT, clGetCommandQueueInfo, CLContext::new);
        }
 
        public CLDevice getDevice() {
-               return getInfoAny(CTYPE_DEVICE, CL_QUEUE_DEVICE);
+               return getInfoAny(CL_QUEUE_DEVICE, clGetCommandQueueInfo, CLDevice::new);
        }
 
        public long getProperties() {
@@ -1026,26 +1954,26 @@ 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(GLSharing.class, CLPlatform.cl_khr_gl_sharing);
+       //}
 
        /*
         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.
index a2d51c3..34d89d0 100644 (file)
  */
 package au.notzed.zcl;
 
-import java.nio.ByteBuffer;
 import static au.notzed.zcl.CL.*;
-import au.notzed.zcl.khr.GLEvent;
-import au.notzed.zcl.khr.GLSharing;
+import static au.notzed.zcl.CLLib.*;
+import jdk.incubator.foreign.*;
+import api.Memory;
+import api.Allocator;
+import api.Native;
+import api.Callback;
+import java.util.stream.Stream;
+
+import java.lang.invoke.MethodHandle;
+
+import java.nio.ByteBuffer;
+import java.nio.ByteOrder;
 import java.io.ByteArrayOutputStream;
 import java.io.IOException;
 import java.io.InputStream;
 import java.nio.charset.Charset;
 
+import java.util.function.Function;
+import java.util.function.IntFunction;
+
 /**
  * Interface for cl_context
  */
@@ -33,24 +45,34 @@ public class CLContext extends CLExtendable {
        /**
         * If a notify callback is supplied to createContext() then this is used to track the reference.
         */
-       final CLContextNotify notify;
+       final Callback<CLContextNotify> notify;
 
        /**
         * Creates an interface for a native pointer of type cl_context.
         *
         * @param p Native pointer.
         */
-       public CLContext(long p) {
+       public CLContext(MemoryAddress p) {
                super(p);
                this.notify = null;
        }
 
-       CLContext(long p, CLContextNotify notify) {
+       CLContext(MemoryAddress p, Callback<CLContextNotify> notify) {
                super(p);
                this.notify = notify;
        }
 
-       private native static void release(long p);
+       static CLContext create(MemoryAddress p) {
+               return Native.resolve(p, CLContext::new);
+       }
+
+       private static void release(MemoryAddress p) {
+               try {
+                       clReleaseContext(p);
+               } catch (Throwable t) {
+                       t.printStackTrace();
+               }
+       }
 
        @Override
        public String toString() {
@@ -65,8 +87,8 @@ public class CLContext extends CLExtendable {
        }
 
        @Override
-       int getInfoType() {
-               return TYPE_CONTEXT;
+       MethodHandle getInfoFunc() {
+               return clGetContextInfo;
        }
 
        /**
@@ -76,7 +98,7 @@ public class CLContext extends CLExtendable {
         * @return new property
         */
        public static CLContextProperty PLATFORM(CLPlatform platform) {
-               return new CLContextProperty.TagValue(CL.CL_CONTEXT_PLATFORM, platform.getP());
+               return new CLContextProperty.TagValue(CL.CL_CONTEXT_PLATFORM, platform.addr().offset());
        }
 
        /**
@@ -99,7 +121,31 @@ public class CLContext extends CLExtendable {
         * release() should be called when it is no longer needed.
         * @throws CLRuntimeException
         */
-       public static native CLContext createContext(CLContextProperty[] properties, CLDevice[] devices, CLContextNotify notify) throws CLRuntimeException;
+       public static CLContext createContext(CLContextProperty[] properties, CLDevice[] devices, CLContextNotify notify) throws CLRuntimeException {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress pprops = CLProperty.toNative(frame, properties);
+                       MemoryAddress pdevs = frame.alloca(devices.length * 8);
+                       MemoryAddress pres = frame.alloca(8);
+                       Callback<CLContextNotify> call = CLContextNotify.call(notify);
+                       MemoryAddress cl;
+                       int res;
+
+                       for (int i=0;i<devices.length;i++)
+                               setAddr(pdevs, i, devices[i].addr());
+
+                       cl = clCreateContext(pprops, devices.length, pdevs, call.addr(), MemoryAddress.NULL, pres);
+                       res = getInt(pres);
+
+                       if (res != 0)
+                               throw new CLRuntimeException(res);
+
+                       return Native.resolve(cl, CLContext::new);
+               } catch (RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        public static CLContext createContext(CLContextProperty[] properties, CLDevice[] devices) throws CLRuntimeException {
                return createContext(properties, devices, (String what, ByteBuffer error_info) -> {
@@ -117,7 +163,28 @@ public class CLContext extends CLExtendable {
         * release() should be called when it is no longer needed.
         * @throws CLRuntimeException
         */
-       public static native CLContext createContextFromType(CLContextProperty[] properties, long device_type, CLContextNotify notify) throws CLRuntimeException;
+       public static CLContext createContextFromType(CLContextProperty[] properties, long device_type, CLContextNotify notify) throws CLRuntimeException {
+
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress pprops = CLProperty.toNative(frame, properties);
+                       MemoryAddress pres = frame.alloca(8);
+                       MemoryAddress cl;
+                       Callback<CLContextNotify> call = CLContextNotify.call(notify);
+                       int res;
+
+                       cl = clCreateContextFromType(pprops, device_type, call.addr(), MemoryAddress.NULL, pres);
+                       res = getInt(pres);
+
+                       if (res != 0)
+                               throw new CLRuntimeException(res);
+
+                       return Native.resolve(cl, (p) -> new CLContext(p, call));
+               } catch (RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         * Calls clCreateCommandQueue.
@@ -127,10 +194,27 @@ public class CLContext extends CLExtendable {
         * @return Newly created command queue.<p>
         * release() should be called when it is no longer needed.
         * @throws CLRuntimeException
-        * @deprecated as of OpenCL 2.0
+        * deprecated as of OpenCL 2.0
         */
-       @Deprecated
-       native public CLCommandQueue createCommandQueue(CLDevice dev, long properties) throws CLRuntimeException;
+       public CLCommandQueue createCommandQueue(CLDevice dev, long properties) throws CLRuntimeException {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress pres = frame.alloca(8);
+                       MemoryAddress q;
+                       int res;
+
+                       q = clCreateCommandQueue(addr(), dev.addr(), properties, pres);
+                       res = getInt(pres);
+
+                       if (res != 0)
+                               throw new CLRuntimeException(res);
+
+                       return Native.resolve(q, CLCommandQueue::new);
+               } catch (RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         * Calls clCreateCommandQueueWithProperties.
@@ -141,7 +225,29 @@ public class CLContext extends CLExtendable {
         * @throws CLRuntimeException
         * @since OpenCL 2.0
         */
-       native public CLCommandQueue createCommandQueue(CLDevice dev, CLQueueProperty[] properties) throws CLRuntimeException;
+       public CLCommandQueue createCommandQueue(CLDevice dev, CLQueueProperty[] properties) throws CLRuntimeException {
+               // Fallback if opencl2 not supported?
+               requireAPIVersion(CLPlatform.VERSION_2_0);
+
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress pprops = CLProperty.toNative(frame, properties);
+                       MemoryAddress pres = frame.alloca(8);
+                       MemoryAddress q;
+                       int res;
+
+                       q = clCreateCommandQueueWithProperties(addr(), dev.addr(), pprops, pres);
+                       res = getInt(pres);
+
+                       if (res != 0)
+                               throw new CLRuntimeException(res);
+
+                       return Native.resolve(q, CLCommandQueue::new);
+               } catch (RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         *
@@ -150,23 +256,105 @@ public class CLContext extends CLExtendable {
         * @throws CLException
         * @since OpenCL 2.1
         */
-       native public void setDefaultDeviceCommandQueue(CLDevice dev, CLCommandQueue q) throws CLException;
+       public void setDefaultDeviceCommandQueue(CLDevice dev, CLCommandQueue q) throws CLException {
+               requireAPIVersion(CLPlatform.VERSION_2_1);
 
+               try {
+                       int res;
+
+                       res = clSetDefaultDeviceCommandQueue(addr(), dev.addr(), q.addr());
+                       if (res != 0)
+                               throw new CLException(res);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
+
+       /* ********************************************************************** */
+
+       /**
+        * Calls clCreateBuffer with an empty host pointer.
+        *
+        * @param flags CL_MEM_* flags.
+        * @param size Size in bytes.
+        * @throws CLRuntimeException
+        */
        public CLBuffer createBuffer(long flags, long size) throws CLRuntimeException {
-               return createBuffer(flags, size, null);
+               return createBuffer(flags, size, (MemorySegment)null);
        }
 
        /**
         * Calls clCreateBuffer.
         *
+        * If flags contains CL_MEM_USE_HOST_PTR then hostp must not
+        * be null.  Currently in this case the CLBuffer must be
+        * explicitly released() on the caller thread.
+        *
         * @param flags CL_MEM_* flags.
         * @param size Size in bytes.
         * @param hostp Optional host memory pointer.
-        * @return Newly allocated buffer.<p>
-        * release() should be called when it is no longer needed.
+        * @return Newly allocated buffer.
         * @throws CLRuntimeException
         */
-       native public CLBuffer createBuffer(long flags, long size, ByteBuffer hostp) throws CLRuntimeException;
+       public CLBuffer createBuffer(long flags, long size, MemorySegment hostseg) throws CLRuntimeException {
+               if (hostseg != null && hostseg.byteSize() < size)
+                       throw new CLRuntimeException(CL_INVALID_HOST_PTR);
+
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress pres = frame.alloca(8);
+                       MemoryAddress pbuffer;
+                       int res;
+
+                       pbuffer = clCreateBuffer(addr(), flags, size, addr(hostseg), pres);
+                       res = getInt(pres);
+
+                       if (res != 0)
+                               throw new CLRuntimeException(res);
+
+                       if (hostseg != null && (flags & CL_MEM_USE_HOST_PTR) != 0)
+                               return resolve(pbuffer, (x) -> new CLBuffer(x, hostseg));
+                       else
+                               return resolve(pbuffer, CLBuffer::new);
+               } catch (RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
+
+       /**
+        * Calls clCreateBuffer.
+        *
+        * If flags contains CL_MEM_USE_HOST_PTR then hostp must not
+        * be null.  Currently the CLBuffer must be explicitly
+        * released() on the caller thread.
+        *
+        * @param flags CL_MEM_* flags.
+        * @param size Size in bytes.
+        * @param hostp Optional host memory pointer.
+        * @return Newly allocated buffer.
+        * @throws CLRuntimeException
+        */
+       public CLBuffer createBuffer(long flags, long size, ByteBuffer hostp) throws CLRuntimeException {
+               if (hostp != null) {
+                       MemorySegment hostseg = MemorySegment.ofByteBuffer(hostp);
+                       try {
+                               CLBuffer buffer = createBuffer(flags, size, hostseg);
+
+                               if ((flags & CL_MEM_USE_HOST_PTR) != 0)
+                                       hostseg = null;
+
+                               return buffer;
+                       } finally {
+                               if (hostseg != null)
+                                       hostseg.close();
+                       }
+               } else {
+                       return createBuffer(flags, size, (MemorySegment)null);
+               }
+       }
 
        /**
         * Wraps a full buffer.
@@ -192,41 +380,150 @@ public class CLContext extends CLExtendable {
         * @throws CLRuntimeException if the clCreateBuffer fails, or with CL_INVALID_VALUE if CL_MEM_USE_HOST_PTR is specified in flags.
         * @throws NullPointerException if hostp is null.
         */
-       native public CLBuffer createBuffer(long flags, byte[] hostp) throws CLRuntimeException;
+       public CLBuffer createBuffer(long flags, byte[] hostp) throws CLRuntimeException {
+               if ((flags & CL_MEM_USE_HOST_PTR) != 0)
+                       throw new CLRuntimeException(CL_INVALID_VALUE);
+
+               if (hostp == null)
+                       throw new NullPointerException();
 
-       native public CLBuffer createBuffer(long flags, short[] hostp) throws CLRuntimeException;
+               // Must copy.
+               try (MemorySegment mem = MemorySegment.allocateNative(hostp.length, 16)) {
+                       mem.asByteBuffer().order(ByteOrder.nativeOrder()).put(hostp);
+                       return createBuffer(flags, hostp.length, mem);
+               }
+       }
 
-       native public CLBuffer createBuffer(long flags, int[] hostp) throws CLRuntimeException;
+       //native public CLBuffer createBuffer(long flags, short[] hostp) throws CLRuntimeException;
 
-       native public CLBuffer createBuffer(long flags, float[] hostp) throws CLRuntimeException;
+       //native public CLBuffer createBuffer(long flags, int[] hostp) throws CLRuntimeException;
 
-       native public CLBuffer createBuffer(long flags, double[] hostp) throws CLRuntimeException;
+       public CLBuffer createBuffer(long flags, float[] hostp) throws CLRuntimeException {
+               if ((flags & CL_MEM_USE_HOST_PTR) != 0)
+                       throw new CLRuntimeException(CL_INVALID_VALUE);
+
+               if (hostp == null)
+                       throw new NullPointerException();
+
+               // Must copy.
+               try (MemorySegment mem = MemorySegment.allocateNative(hostp.length, 16)) {
+                       mem.asByteBuffer().order(ByteOrder.nativeOrder()).asFloatBuffer().put(hostp);
+                       return createBuffer(flags, hostp.length * 4, mem);
+               }
+       }
+
+       //native public CLBuffer createBuffer(long flags, double[] hostp) throws CLRuntimeException;
+
+       /* ********************************************************************** */
 
        public CLImage createImage(long flags, CLImageFormat fmt, CLImageDesc desc) throws CLRuntimeException, UnsupportedOperationException {
-               return createImage(flags, fmt, desc, (ByteBuffer) null);
+               return createImage(flags, fmt, desc, (MemorySegment) null);
        }
 
        /**
-        * Calls clCreateImage.
+        * Calls clCreateImage, or clCreateImageXD on OpenCL 1.1 or less.
         *
         * @param flags CL_MEM_* flags.
         * @param fmt Image format.
-        * @param desc Image descriptor. For OpenCL 1.2 only CL_MEM_OBJECT_IMAGE2D and CL_MEM_OBJECT_IMAGE3D are allowed.
+        * @param desc Image descriptor. For OpenCL 1.2, only CL_MEM_OBJECT_IMAGE2D and CL_MEM_OBJECT_IMAGE3D are allowed.
         * @param hostp
         * @return Newly allocated image.<p>
         * release() should be called when it is no longer needed.
         * @throws CLRuntimeException
         * @throws UnsupportedOperationException
         */
-       native public CLImage createImage(long flags, CLImageFormat fmt, CLImageDesc desc, ByteBuffer hostp) throws CLRuntimeException, UnsupportedOperationException;
+       public CLImage createImage(long flags, CLImageFormat fmt, CLImageDesc desc, MemorySegment hostseg) throws CLRuntimeException, UnsupportedOperationException {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress cfmt = CLImageFormat.toNative(frame, fmt);
+                       MemoryAddress cres = frame.alloca(8);
+                       MemoryAddress ci;
+                       int res;
+
+                       // FIXME: perform range checks
+
+                       if (haveAPIVersion(CLPlatform.VERSION_1_2)) {
+                               MemoryAddress cdesc = CLImageDesc.toNative(frame, desc);
+
+                               ci = clCreateImage(addr(), flags, cfmt, cdesc, addr(hostseg), cres);
+                       } else {
+                               switch (desc.imageType) {
+                               case CL_MEM_OBJECT_IMAGE2D:
+                                       ci = clCreateImage2D(addr(), flags, cfmt, desc.imageWidth, desc.imageHeight, desc.imageRowPitch, addr(hostseg), cres);
+                                       break;
+                               case CL_MEM_OBJECT_IMAGE3D:
+                                       ci = clCreateImage3D(addr(), flags, cfmt, desc.imageWidth, desc.imageHeight, desc.imageDepth,
+                                                            desc.imageRowPitch, desc.imageSlicePitch, addr(hostseg), cres);
+                                       break;
+                               default:
+                                       throw new UnsupportedOperationException("Requires OpenCL 1.2");
+                               }
+                       }
+
+                       res = getInt(cres);
+                       if (res != 0)
+                               throw new CLRuntimeException(res);
 
-       native public CLImage createImage(long flags, CLImageFormat fmt, CLImageDesc desc, byte[] hostp) throws CLRuntimeException, UnsupportedOperationException;
+                       if (hostseg != null && (flags & CL_MEM_USE_HOST_PTR) != 0)
+                               return resolve(ci, (x) -> new CLImage(x, hostseg));
+                       else
+                               return resolve(ci, CLImage::new);
+               } catch (RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
+
+       public CLImage createImage(long flags, CLImageFormat fmt, CLImageDesc desc, ByteBuffer hostp) throws CLRuntimeException, UnsupportedOperationException {
+               if (hostp != null) {
+                       MemorySegment hostseg = MemorySegment.ofByteBuffer(hostp);
+                       try {
+                               CLImage image = createImage(flags, fmt, desc, hostseg);
 
-       native public CLImage createImage(long flags, CLImageFormat fmt, CLImageDesc desc, short[] hostp) throws CLRuntimeException, UnsupportedOperationException;
+                               if ((flags & CL_MEM_USE_HOST_PTR) != 0)
+                                       hostseg = null;
 
-       native public CLImage createImage(long flags, CLImageFormat fmt, CLImageDesc desc, int[] hostp) throws CLRuntimeException, UnsupportedOperationException;
+                               return image;
+                       } finally {
+                               if (hostseg != null)
+                                       hostseg.close();
+                       }
+               } else {
+                       return createImage(flags, fmt, desc, (MemorySegment)null);
+               }
+       }
+
+       public CLImage createImage(long flags, CLImageFormat fmt, CLImageDesc desc, byte[] hostp) throws CLRuntimeException, UnsupportedOperationException {
+               if ((flags & CL_MEM_USE_HOST_PTR) != 0)
+                       throw new CLRuntimeException(CL_INVALID_VALUE);
+
+               if (hostp == null)
+                       throw new NullPointerException();
+
+               // Must copy.
+               try (MemorySegment mem = MemorySegment.allocateNative(hostp.length, 16)) {
+                       mem.asByteBuffer().order(ByteOrder.nativeOrder()).put(hostp);
+                       return createImage(flags, fmt, desc, mem);
+               }
+       }
 
-       native public CLImage createImage(long flags, CLImageFormat fmt, CLImageDesc desc, float[] hostp) throws CLRuntimeException, UnsupportedOperationException;
+       //native public CLImage createImage(long flags, CLImageFormat fmt, CLImageDesc desc, short[] hostp) throws CLRuntimeException, UnsupportedOperationException;
+
+       //native public CLImage createImage(long flags, CLImageFormat fmt, CLImageDesc desc, int[] hostp) throws CLRuntimeException, UnsupportedOperationException;
+
+       public CLImage createImage(long flags, CLImageFormat fmt, CLImageDesc desc, float[] hostp) throws CLRuntimeException, UnsupportedOperationException {
+               if ((flags & CL_MEM_USE_HOST_PTR) != 0)
+                       throw new CLRuntimeException(CL_INVALID_VALUE);
+
+               if (hostp == null)
+                       throw new NullPointerException();
+
+               // Must copy.
+               try (MemorySegment mem = MemorySegment.allocateNative(hostp.length * 4, 16)) {
+                       mem.asByteBuffer().order(ByteOrder.nativeOrder()).asFloatBuffer().put(hostp);
+                       return createImage(flags, fmt, desc, mem);
+               }
+       }
 
        /**
         * Calls clCreatePipe
@@ -240,7 +537,27 @@ public class CLContext extends CLExtendable {
         * @throws UnsupportedOperationException
         * @since OpenCL 2.0
         */
-       native public CLPipe createPipe(long flags, int packetSize, int maxPackets, CLPipeProperty[] properties) throws CLRuntimeException, UnsupportedOperationException;
+       public CLPipe createPipe(long flags, int packetSize, int maxPackets, CLPipeProperty[] properties) throws CLRuntimeException, UnsupportedOperationException {
+               requireAPIVersion(CLPlatform.VERSION_2_0);
+
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress cprops = CLProperty.toNative(frame, properties);
+                       MemoryAddress cres = frame.alloca(8);
+                       int res;
+                       MemoryAddress cp;
+
+                       cp = clCreatePipe(addr(), flags, packetSize, maxPackets, cprops, cres);
+                       res = getInt(cres);
+                       if (res != 0)
+                               throw new CLRuntimeException(res);
+
+                       return resolve(cp, CLPipe::new);
+               } catch (RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         * Calls clGetSupportedImageFormats.
@@ -250,7 +567,33 @@ public class CLContext extends CLExtendable {
         * @return List of supported image formats.
         * @throws CLRuntimeException
         */
-       native public CLImageFormat[] getSupportedImageFormats(long flags, int type) throws CLRuntimeException;
+       public CLImageFormat[] getSupportedImageFormats(long flags, int type) throws CLRuntimeException {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress cnum = frame.alloca(8);
+                       MemoryAddress list;
+                       int num;
+                       int res;
+
+                       res = clGetSupportedImageFormats(addr(), flags, type, 0, MemoryAddress.NULL, cnum);
+                       if (res != 0)
+                               throw new CLRuntimeException(res);
+
+                       num = getInt(cnum);
+                       list = frame.alloca(num * 8);
+
+                       res = clGetSupportedImageFormats(addr(), flags, type, num, list, cnum);
+
+                       CLImageFormat[] out = new CLImageFormat[num];
+                       for (int i=0;i<out.length;i++)
+                               out[i] = CLImageFormat.fromNative(getAddr(list, i));
+
+                       return out;
+               } catch (RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         * Allocate shared virtual memory.
@@ -264,7 +607,9 @@ public class CLContext extends CLExtendable {
         * @throws CLRuntimeException
         * @since OpenCL 2.0
         */
-       native public ByteBuffer SVMAlloc(long flags, long size, int alignment) throws CLRuntimeException;
+       public ByteBuffer SVMAlloc(long flags, long size, int alignment) throws CLRuntimeException {
+               throw new UnsupportedOperationException();
+       }
 
        /**
         * Free memory allocated by SVMAlloc. It is up to the caller to ensure the memory is not in use.
@@ -273,7 +618,9 @@ public class CLContext extends CLExtendable {
         * @throws CLRuntimeException
         * @since OpenCL 2.0
         */
-       native public void SVMFree(ByteBuffer mem) throws CLRuntimeException;
+       public void SVMFree(ByteBuffer mem) throws CLRuntimeException {
+               throw new UnsupportedOperationException();
+       }
 
        /**
         * Calls clCreateSampler.
@@ -286,10 +633,27 @@ public class CLContext extends CLExtendable {
         * @return Newly created sampler.<p>
         * release() should be called when it is no longer needed.
         * @throws CLRuntimeException
-        * @deprecated as of OpenCL 2.0
+        * xdeprecated as of OpenCL 2.0
         */
-       @Deprecated
-       native public CLSampler createSampler(boolean norm, int addr_mode, int filter_mode) throws CLRuntimeException;
+       //@Deprecated
+       public CLSampler createSampler(boolean norm, int addr_mode, int filter_mode) throws CLRuntimeException {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress cres = frame.alloca(8);
+                       int res;
+                       MemoryAddress cs;
+
+                       cs = clCreateSampler(addr(), norm ? 1 : 0, addr_mode, filter_mode, cres);
+                       res = getInt(cres);
+                       if (res != 0)
+                               throw new CLRuntimeException(res);
+
+                       return resolve(cs, CLSampler::new);
+               } catch (RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         * Calls clCreateSamplerWithProperties
@@ -299,7 +663,50 @@ public class CLContext extends CLExtendable {
         * @throws CLRuntimeException
         * @since OpenCL 2.0
         */
-       native public CLSampler createSampler(CLSamplerProperty[] props) throws CLRuntimeException;
+       public CLSampler createSampler(CLSamplerProperty[] props) throws CLRuntimeException {
+               // Fallback if opencl2 not supported?
+               requireAPIVersion(CLPlatform.VERSION_2_0);
+
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress cres = frame.alloca(8);
+                       MemoryAddress cprops = CLProperty.toNative(frame, props);
+                       int res;
+                       MemoryAddress cs;
+
+                       cs = clCreateSamplerWithProperties(addr(), cprops, cres);
+                       res = getInt(cres);
+                       if (res != 0)
+                               throw new CLRuntimeException(res);
+
+                       return resolve(cs, CLSampler::new);
+               } catch (RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+
+       }
+
+       static long length(byte[][] strings) {
+               long len = 0;
+               for (int i=0;i<strings.length;i++)
+                       len += strings[i].length;
+               return len;
+       }
+
+       static void copy(MemoryAddress addr, byte[][]list) {
+               for (int i=0,k=0;i<list.length;i++) {
+                       byte[] row = list[i];
+
+                       for (int j=0;j<row.length;j++,k++)
+                               setByte(addr, k, row[j]);
+               }
+       }
+
+       static void copy(MemoryAddress addr, byte[]row) {
+               for (int j=0;j<row.length;j++)
+                       setByte(addr, j, row[j]);
+       }
 
        /**
         * Calls clCreateProgramWithSource.
@@ -309,13 +716,40 @@ public class CLContext extends CLExtendable {
         * release() should be called when it is no longer needed or allow for GC cleanup.
         * @throws CLException
         */
-       native public CLProgram createProgramWithSource(byte[]... strings) throws CLException;
+       public CLProgram createProgramWithSource(byte[]... strings) throws CLException {
+               long size = length(strings);
+               try (Allocator frame = Memory.stack();
+                    MemorySegment buffer = MemorySegment.allocateNative(size)) {
+                       MemoryAddress cstring = frame.alloca(8);
+                       MemoryAddress clength = frame.alloca(8);
+                       MemoryAddress cret = frame.alloca(8);
+                       MemoryAddress base = buffer.baseAddress();
+                       MemoryAddress cp;
+                       int res;
+
+                       copy(base, strings);
+                       setAddr(cstring, base);
+                       setLong(clength, size);
+
+                       cp = clCreateProgramWithSource(addr(), 1, cstring, clength, cret);
+
+                       res = getInt(cret);
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       return resolve(cp, CLProgram::new);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        public CLProgram createProgramWithSource(String... strings) throws CLException {
                byte[][] bytes = new byte[strings.length][];
-               for (int i = 0; i < strings.length; i++) {
+
+               for (int i = 0; i < strings.length; i++)
                        bytes[i] = strings[i].getBytes(Charset.forName("UTF-8"));
-               }
 
                return createProgramWithSource(bytes);
        }
@@ -340,7 +774,31 @@ public class CLContext extends CLExtendable {
         * @return
         * @throws CLException
         */
-       native public CLProgram createProgramWithIL(byte[] il) throws CLException;
+       public CLProgram createProgramWithIL(byte[] il) throws CLException {
+               requireAPIVersion(CLPlatform.VERSION_2_1);
+
+               try (Allocator frame = Memory.stack();
+                    MemorySegment buffer = MemorySegment.allocateNative(il.length, 1)) {
+                       MemoryAddress cret = frame.alloca(8);
+                       MemoryAddress base = buffer.baseAddress();
+                       MemoryAddress cp;
+                       int res;
+
+                       copy(base, il);
+
+                       cp = clCreateProgramWithIL(addr(), base, il.length, cret);
+
+                       res = getInt(cret);
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       return resolve(cp, CLProgram::new);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         * Calls clCreateProgramWithBinary.
@@ -352,7 +810,47 @@ public class CLContext extends CLExtendable {
         * release() should be called when it is no longer needed.
         * @throws CLException
         */
-       native public CLProgram createProgramWithBinary(CLDevice[] devices, byte[][] binaries, int[] status) throws CLException;
+       public CLProgram createProgramWithBinary(CLDevice[] devices, byte[][] binaries, int[] status) throws CLException {
+               long size = length(binaries);
+               try (Allocator frame = Memory.stack();
+                    MemorySegment buffer = MemorySegment.allocateNative(size, 1)) {
+                       MemoryAddress cdevs = toAddrV(frame, devices);
+                       MemoryAddress barray = frame.alloca(8 * binaries.length);
+                       MemoryAddress larray = frame.alloca(8 * binaries.length);
+                       MemoryAddress cstatus = status != null ? frame.alloca(4 * devices.length) : MemoryAddress.NULL;
+                       MemoryAddress cret = frame.alloca(8);
+                       MemoryAddress base = buffer.baseAddress();
+                       MemoryAddress cp;
+                       int res;
+
+                       for (int i=0, o=0;i<binaries.length;i++) {
+                               MemoryAddress addr = base.addOffset(o);
+
+                               copy(addr, binaries[i]);
+                               setAddr(barray, i, addr);
+
+                               setLong(larray,  i, binaries[i].length);
+                               o += binaries[i].length;
+                       }
+
+                       cp = clCreateProgramWithBinary(addr(), devices.length, cdevs, larray, barray, cstatus, cret);
+
+                       res = getInt(cret);
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       if (status != null) {
+                               for (int i=0;i<status.length;i++)
+                                       status[i] = getInt(cstatus, i);
+                       }
+
+                       return resolve(cp, CLProgram::new);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         * Calls clCreateProgramWithBuiltInKernels.
@@ -364,7 +862,29 @@ public class CLContext extends CLExtendable {
         * @throws UnsupportedOperationException
         * @since OpenCL 1.2
         */
-       native public CLProgram createProgramWithBuiltInKernels(CLDevice[] devices, String names) throws CLException, UnsupportedOperationException;
+       public CLProgram createProgramWithBuiltInKernels(CLDevice[] devices, String names) throws CLException, UnsupportedOperationException {
+               requireAPIVersion(CLPlatform.VERSION_1_2);
+
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress cdevs = toAddrV(frame, devices);
+                       MemoryAddress cnames = toByteV(frame, names);
+                       MemoryAddress cret = frame.alloca(8);
+                       MemoryAddress cp;
+                       int res;
+
+                       cp = clCreateProgramWithBuiltInKernels(addr(), devices.length, cdevs, cnames, cret);
+
+                       res = getInt(cret);
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       return resolve(cp, CLProgram::new);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         * Calls clLinkProgram.
@@ -379,7 +899,31 @@ public class CLContext extends CLExtendable {
         * @throws UnsupportedOperationException
         * @since OpenCL 1.2
         */
-       native public CLProgram linkProgram(CLDevice[] devices, String options, CLProgram[] programs, CLNotify<CLProgram> notify) throws CLException, UnsupportedOperationException;
+       public CLProgram linkProgram(CLDevice[] devices, String options, CLProgram[] programs, CLNotify<CLProgram> notify) throws CLException, UnsupportedOperationException {
+               requireAPIVersion(CLPlatform.VERSION_1_2);
+
+               try (Allocator frame = Memory.stack();
+                    Callback<CLNotify<CLProgram>> cnotify = CLNotify.call(notify, CLProgram::new)) {
+                       MemoryAddress cdevs = toAddrV(frame, devices);
+                       MemoryAddress coptions = toByteV(frame, options);
+                       MemoryAddress cprogs = toAddrV(frame, programs);
+                       MemoryAddress cret = frame.alloca(8);
+                       MemoryAddress cp;
+                       int res;
+
+                       cp = clLinkProgram(addr(), devices.length, cdevs, coptions, programs.length, cprogs, cnotify.addr(), MemoryAddress.NULL, cret);
+
+                       res = getInt(cret);
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       return resolve(cp, CLProgram::new);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         * Calls clCreateUserEvent.
@@ -389,7 +933,27 @@ public class CLContext extends CLExtendable {
         * @throws CLRuntimeException
         * @since OpenCL 1.1
         */
-       native public CLEvent createUserEvent() throws CLRuntimeException;
+       public CLEvent createUserEvent() throws CLException {
+               requireAPIVersion(CLPlatform.VERSION_1_1);
+
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress cret = frame.alloca(8);
+                       MemoryAddress ce;
+                       int res;
+
+                       ce = clCreateUserEvent(addr(), cret);
+
+                       res = getInt(cret);
+                       if (res != 0)
+                               throw new CLException(res);
+
+                       return resolve(ce, CLEvent::new);
+               } catch (CLException | RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         * gets CL_CONTEXT_NUM_DEVICES.
@@ -406,7 +970,7 @@ public class CLContext extends CLExtendable {
         * @return List of devices.
         */
        public CLDevice[] getDevices() {
-               return getInfoAnyV(CTYPE_DEVICE, CL_CONTEXT_DEVICES);
+               return getInfoAnyV(CL_CONTEXT_DEVICES, clGetContextInfo, CLDevice::new, CLDevice[]::new);
        }
 
        /**
@@ -414,45 +978,48 @@ public class CLContext extends CLExtendable {
         *
         * @return List of properties used at create time.
         */
-       public native CLContextProperty[] getProperties();
+       public CLContextProperty[] getProperties() {
+               return getInfoPropertyV(CL_CONTEXT_PROPERTIES, clGetContextInfo, CLContextProperty.TagValue::new, CLContextProperty[]::new);
+       }
 
        @Override
        protected CLPlatform initPlatform() {
                return getDevices()[0].platform;
        }
 
-       protected GLSharing getGLSharing() {
-               return getExtension(GLSharing.class, CLPlatform.cl_khr_gl_sharing);
-       }
+       //protected GLSharing getGLSharing() {
+       //      return getExtension(GLSharing.class, CLPlatform.cl_khr_gl_sharing);
+       //}
 
-       protected GLEvent getGLEvent() {
-               return getExtension(GLEvent.class, CLPlatform.cl_khr_gl_event);
-       }
+       //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 CLBuffer createFromGLBuffer(
+       //      long flags,
+       //      int bufobj) {
+       //      return getGLSharing().createFromGLBuffer(this, flags, bufobj);
+       // }
 
-       public CLImage createFromGLRenderbuffer(
-               long flags /* flags */,
-               int renderbuffer /* renderbuffer */) {
-               return getGLSharing().createFromGLRenderbuffer(this, flags, renderbuffer);
-       }
+       // public CLImage createFromGLTexture(
+       //      long flags /* flags */,
+       //      int target /* target */,
+       //      int miplevel /* miplevel */,
+       //      int texture /* texture */) {
+       //      return getGLSharing().createFromGLTexture(this, flags, target, miplevel, texture);
+       // }
 
-       public CLEvent createEventFromGLsync(long glsync) {
-               return getGLEvent().clCreateEventFromGLsync(this, glsync);
-       }
+       // 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);
+       // }
 }
index 8c3a304..e8f18f7 100644 (file)
 package au.notzed.zcl;
 
 import java.nio.ByteBuffer;
+import java.util.function.Function;
+import jdk.incubator.foreign.MemoryAddress;
+import jdk.incubator.foreign.MemorySegment;
+import api.Native;
+import api.Callback;
+import api.Memory;
 
 /**
  * Callback for CLContext.createContext*()
@@ -25,4 +31,19 @@ public interface CLContextNotify {
 
        public void notify(String what, ByteBuffer error_info);
 
+       @SuppressWarnings("unchecked")
+       static Callback<CLContextNotify> 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());
+                                       }),
+                               (p) -> new Callback<>(p, notify));
+               } else {
+                       return Callback.NULL;
+               }
+       }
+
+
 }
index 37dc9fb..f6e474c 100644 (file)
  */
 package au.notzed.zcl;
 
+import api.Native;
+import java.lang.invoke.MethodHandle;
+import java.util.function.Function;
+import jdk.incubator.foreign.*;
 import static au.notzed.zcl.CL.*;
+import static au.notzed.zcl.CLLib.*;
 
 /**
  * Interface for cl_device_id.
@@ -28,11 +33,18 @@ public class CLDevice extends CLExtendable {
         *
         * @param p Native pointer.
         */
-       public CLDevice(long p) {
+       public CLDevice(MemoryAddress p) {
                super(p);
        }
 
-       private native static void release(long p);
+       private static void release(MemoryAddress p) {
+               try {
+                       clReleaseDevice(p);
+               } catch (Throwable t) {
+               }
+       }
+
+       //private native static void release(long p);
 
        @Override
        public String toString() {
@@ -40,8 +52,8 @@ public class CLDevice extends CLExtendable {
        }
 
        @Override
-       int getInfoType() {
-               return TYPE_DEVICE;
+       MethodHandle getInfoFunc() {
+               return clGetDeviceInfo;
        }
 
        @Override
@@ -55,7 +67,7 @@ public class CLDevice extends CLExtendable {
         * @param times { device_timestamp, host_timestamp }
         * @since OpenCL 2.1
         */
-       public native void getDeviceAndHostTimer(long[] times);
+       //public native void getDeviceAndHostTimer(long[] times);
 
        /**
         * Gets device view of host clock.
@@ -63,7 +75,7 @@ public class CLDevice extends CLExtendable {
         * @return host_timestamp.
         * @since OpenCN 2.1
         */
-       public native long getHostTimer();
+       //public native long getHostTimer();
 
        /**
         * Create a property representing CL_DEVICE_PARTITION_EQUALLY.
@@ -107,7 +119,7 @@ public class CLDevice extends CLExtendable {
         * @throws CLException
         * @since OpenCL 1.2
         */
-       public native CLDevice[] createSubDevices(CLDeviceProperty[] properties, int num_devices) throws CLException, UnsupportedOperationException;
+       //public native CLDevice[] createSubDevices(CLDeviceProperty[] properties, int num_devices) throws CLException, UnsupportedOperationException;
 
        /**
         * get CL_DEVICE_TYPE
@@ -349,7 +361,7 @@ public class CLDevice extends CLExtendable {
        }
 
        public CLPlatform getPlatform() {
-               return getInfoAny(CTYPE_PLATFORM, CL_DEVICE_PLATFORM);
+               return getInfoAny(CL_DEVICE_PLATFORM, clGetDeviceInfo, CLPlatform::new);
        }
 
        public String getName() {
@@ -389,22 +401,22 @@ public class CLDevice extends CLExtendable {
        }
 
        public CLDevice getParentDevice() {
-               return getInfoAny(CTYPE_DEVICE, CL_DEVICE_PARENT_DEVICE);
+               return getInfoAny(CL_DEVICE_PARENT_DEVICE, clGetDeviceInfo, CLDevice::new);
        }
 
        public int getPartitionMaxSubDevices() {
                return getInfoInt(CL_DEVICE_PARTITION_MAX_SUB_DEVICES);
        }
 
-       public native CLDeviceProperty[] getPartitionProperties();
+       //public native CLDeviceProperty[] getPartitionProperties();
 
        public long getPartitionAffinityDomain() {
                return getInfoLong(CL_DEVICE_PARTITION_AFFINITY_DOMAIN);
        }
 
-       public native CLDeviceProperty[] getPartitionType();
+       //public native CLDeviceProperty[] getPartitionType();
 
-       public static CLDevice[] newArray(int n) {
-               return new CLDevice[n];
-       }
+       //public static CLDevice[] newArray(int n) {
+       //      return new CLDevice[n];
+       //}
 }
index 4d1d0b4..f3f1afb 100644 (file)
@@ -16,6 +16,9 @@
  */
 package au.notzed.zcl;
 
+import api.Native;
+import jdk.incubator.foreign.MemoryAddress;
+
 /**
  * Properties for CLDevice sub-device creation.
  *
@@ -62,6 +65,14 @@ public interface CLDeviceProperty extends CLProperty {
                        dst[o++] = CL.CL_DEVICE_PARTITION_BY_COUNTS_LIST_END;
                        return o;
                }
+
+               public int toLong(MemoryAddress dst, int o) {
+                       Native.setLong(dst, o++, CL.CL_DEVICE_PARTITION_BY_COUNTS);
+                       for (int c : counts)
+                               Native.setLong(dst, o++, c);
+                       Native.setLong(dst, o++, CL.CL_DEVICE_PARTITION_BY_COUNTS_LIST_END);
+                       return o;
+               }
        }
 
 }
index f8dcb60..6e63126 100644 (file)
  */
 package au.notzed.zcl;
 
+import jdk.incubator.foreign.*;
 import static au.notzed.zcl.CL.*;
+import static au.notzed.zcl.CLLib.*;
+import api.Native;
+import api.Callback;
+import java.lang.invoke.MethodHandle;
 
 /**
  * Interface for cl_event.
@@ -24,19 +29,40 @@ import static au.notzed.zcl.CL.*;
 public class CLEvent extends CLObject {
 
        /**
+        * This is used to retain a reference for any callback set
+        */
+       Callback<CLEventNotify> callback;
+
+       /**
         * Create an interface for a native pointer of type cl_event.
         *
         * @param p Native pointer.
         */
-       public CLEvent(long p) {
+       public CLEvent(MemoryAddress p) {
                super(p);
        }
 
-       private native static void release(long p);
+       public static CLEvent create(MemoryAddress p) {
+               return Native.resolve(p, CLEvent::new);
+       }
+
+       private static void release(MemoryAddress p) {
+               try {
+                       System.err.printf("** release event %016x\n", api.Memory.toLong(p));
+                       clReleaseEvent(p);
+               } catch (Throwable t) {
+               }
+       }
+
+       public void release() {
+               Native.release(callback);
+               callback = null;
+               super.release();
+       }
 
        @Override
-       int getInfoType() {
-               return TYPE_EVENT;
+       MethodHandle getInfoFunc() {
+               return clGetEventInfo;
        }
 
        /**
@@ -48,7 +74,20 @@ public class CLEvent extends CLObject {
         * @throws CLRuntimeException
         * @since OpenCL 1.1
         */
-       public native void setUserEventStatus(int status) throws CLRuntimeException;
+       public void setUserEventStatus(int status) throws CLRuntimeException {
+               // avoid platform lookup for costs
+               //requiredAPIVersion(CLPlatform.VERSION_1_1);
+               try {
+                       int res = clSetUserEventStatus(addr(), status);
+
+                       if (res != 0)
+                               throw new CLRuntimeException(res);
+               } catch (RuntimeException | Error t) {
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         * Call clSetEventCallback(type, notify).
@@ -58,7 +97,28 @@ public class CLEvent extends CLObject {
         * @throws CLRuntimeException
         * @since OpenCL 1.1
         */
-       public native void setEventCallback(int type, CLEventNotify notify) throws CLRuntimeException;
+       public void setEventCallback(int type, CLEventNotify notify) throws CLRuntimeException {
+               //requiredAPIVersion(CLPlatform.VERSION_1_1);
+
+               Native.release(callback);
+
+               if (notify != null) {
+                       callback = CLEventNotify.call(notify);
+
+                       try {
+                               int res = clSetEventCallback(addr(), type, callback.addr(), MemoryAddress.NULL);
+
+                               if (res != 0)
+                                       throw new CLRuntimeException(res);
+                       } catch (RuntimeException | Error t) {
+                               throw t;
+                       } catch (Throwable t) {
+                               throw new RuntimeException(t);
+                       }
+               } else {
+                       callback = null;
+               }
+       }
 
        /**
         * Get CL_EVENT_COMMAND_QUEUE.
@@ -68,7 +128,7 @@ public class CLEvent extends CLObject {
         * @return
         */
        public CLCommandQueue getCommandQueue() {
-               return getInfoAny(CTYPE_COMMAND_QUEUE, CL_EVENT_COMMAND_QUEUE);
+               return getInfoAny(CL_EVENT_COMMAND_QUEUE, clGetEventInfo, CLCommandQueue::new);
        }
 
        /**
@@ -79,7 +139,7 @@ public class CLEvent extends CLObject {
         * @return
         */
        public CLContext getContext() {
-               return getInfoAny(CTYPE_CONTEXT, CL_EVENT_CONTEXT);
+               return getInfoAny(CL_EVENT_CONTEXT, clGetEventInfo, CLContext::new);
        }
 
        /**
@@ -103,22 +163,22 @@ public class CLEvent extends CLObject {
        }
 
        public long getProfilingCommandQueued() {
-               return getInfoLong(TYPE_EVENT_PROFILING, CL_PROFILING_COMMAND_QUEUED);
+               return getInfoLong(CL_PROFILING_COMMAND_QUEUED, clGetEventProfilingInfo);
        }
 
        public long getProfilingCommandSubmit() {
-               return getInfoLong(TYPE_EVENT_PROFILING, CL_PROFILING_COMMAND_SUBMIT);
+               return getInfoLong(CL_PROFILING_COMMAND_SUBMIT, clGetEventProfilingInfo);
        }
 
        public long getProfilingCommandStart() {
-               return getInfoLong(TYPE_EVENT_PROFILING, CL_PROFILING_COMMAND_START);
+               return getInfoLong(CL_PROFILING_COMMAND_START, clGetEventProfilingInfo);
        }
 
        public long getProfilingCommandEnd() {
-               return getInfoLong(TYPE_EVENT_PROFILING, CL_PROFILING_COMMAND_END);
+               return getInfoLong(CL_PROFILING_COMMAND_END, clGetEventProfilingInfo);
        }
 
        public long getProfilingCommandComplete() {
-               return getInfoLong(TYPE_EVENT_PROFILING, CL_PROFILING_COMMAND_COMPLETE);
+               return getInfoLong(CL_PROFILING_COMMAND_COMPLETE, clGetEventProfilingInfo);
        }
 }
index 96f6c5d..bb38677 100644 (file)
  */
 package au.notzed.zcl;
 
+import jdk.incubator.foreign.MemoryAddress;
+import api.Memory;
+import api.Allocator;
+import api.Native;
+import static au.notzed.zcl.CLLib.*;
+
 /**
  * Manages a list of cl_events.
  * <p>
@@ -38,10 +44,6 @@ package au.notzed.zcl;
 public class CLEventList {
 
        /**
-        * C-accessible copy of jevets[i].p
-        */
-       final long[] events;
-       /**
         * Event references.
         */
        final CLEvent[] jevents;
@@ -53,7 +55,6 @@ public class CLEventList {
         * @param capacity Sets the event list capacity. This
         */
        public CLEventList(int capacity) {
-               this.events = new long[capacity];
                this.jevents = new CLEvent[capacity];
        }
 
@@ -65,7 +66,6 @@ public class CLEventList {
        public void reset() {
                for (int i = 0; i < index; i++) {
                        jevents[i] = null;
-                       events[i] = 0;
                }
                index = 0;
        }
@@ -86,20 +86,10 @@ public class CLEventList {
         * @param event
         */
        public void add(CLEvent event) {
-               events[index] = event.getP();
                jevents[index++] = event;
        }
 
        /**
-        * Interface for JNI to add an event after a successful enqueue operation.
-        *
-        * @param eid
-        */
-       private void add(long eid) {
-               add(CLObject.resolve(CLEvent.class, eid));
-       }
-
-       /**
         * Get the number of active events.
         *
         * @return
@@ -108,10 +98,38 @@ public class CLEventList {
                return index;
        }
 
+       static MemoryAddress toWaitList(Allocator frame, CLEventList list) {
+               if (list != null) {
+                       MemoryAddress addr = frame.alloca(list.index * 8);
+
+                       for (int i=0;i<list.index;i++)
+                               Native.setAddr(addr, i, list.jevents[i].addr());
+
+                       return addr;
+               } else {
+                       return MemoryAddress.NULL;
+               }
+       }
+
        /**
         * Calls clWaitForEvents with all active events in this list.
         *
         * @throws CLException
         */
-       public native void waitForEvents() throws CLException;
+       public void waitForEvents() throws CLException {
+               if (size() > 0) {
+                       try (Allocator frame = Memory.stack()) {
+                               MemoryAddress events = toWaitList(frame, this);
+                               int res;
+
+                               res = clWaitForEvents(size(), events);
+                               if (res != 0)
+                                       throw new CLException(res);
+                       } catch (CLException | RuntimeException | Error t) {
+                               throw t;
+                       } catch (Throwable t) {
+                               throw new RuntimeException(t);
+                       }
+               }
+       }
 }
index 1c793a1..95331d1 100644 (file)
  */
 package au.notzed.zcl;
 
+import api.Callback;
+import api.Native;
+
 /**
  * Callback for CLEvent.
  */
+@FunctionalInterface
 public interface CLEventNotify {
 
        /**
@@ -28,4 +32,15 @@ public interface CLEventNotify {
         * @param status status.
         */
        public void notify(CLEvent event, int status);
+
+       @SuppressWarnings("unchecked")
+       static Callback<CLEventNotify> call(CLEventNotify notify) {
+               if (notify != null) {
+                       return Native.resolve(
+                               Call_pLIpv_v.stub((cevent, status, dummy) -> notify.notify(Native.resolve(cevent, CLEvent::new), status)),
+                               (p) -> new Callback<>(p, notify));
+               } else {
+                       return Callback.NULL;
+               }
+       }
 }
index fccafe9..fd6a922 100644 (file)
-/**
- * *****************************************************************************
- * Copyright (c) 2008 - 2012 The Khronos Group Inc.
- *
- * Permission is hereby granted, free of charge, to any person obtaining a copy
- * of this software and/or associated documentation files (the "Materials"), to
- * deal in the Materials without restriction, including without limitation the
- * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
- * sell copies of the Materials, and to permit persons to whom the Materials are
- * furnished to do so, subject to the following conditions:
- *
- * The above copyright notice and this permission notice shall be included in
- * all copies or substantial portions of the Materials.
- *
- * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
- * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
- * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
- * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
- * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
- * OUT OF OR IN CONNECTION WITH THE MATERIALS OR THE USE OR OTHER DEALINGS IN
- * THE MATERIALS.
- * ****************************************************************************
- */
+
 package au.notzed.zcl;
 
-/**
- * Exception for OpenCL error values.
- */
 public class CLException extends Exception {
+       int res;
 
-       int id;
-
-       public CLException(int id) {
-               this.id = id;
-       }
-
-       /**
-        * Get the error code.
-        *
-        * @return OpenCL error code.
-        * @see #codeToError
-        */
-       public int getID() {
-               return id;
+       public CLException() {
        }
 
-       @Override
-       public String getMessage() {
-               return String.format("Error (%d): '%s'", id, codeToError(id));
-       }
-
-       /**
-        * Convert the error code to a description.
-        *
-        * @param code
-        * @return
-        */
-       public static String codeToError(int code) {
-               switch (code) {
-               case CL.CL_SUCCESS:
-                       return "CL_SUCCESS";
-               case CL.CL_DEVICE_NOT_FOUND:
-                       return "CL_DEVICE_NOT_FOUND";
-               case CL.CL_DEVICE_NOT_AVAILABLE:
-                       return "CL_DEVICE_NOT_AVAILABLE";
-               case CL.CL_COMPILER_NOT_AVAILABLE:
-                       return "CL_COMPILER_NOT_AVAILABLE";
-               case CL.CL_MEM_OBJECT_ALLOCATION_FAILURE:
-                       return "CL_MEM_OBJECT_ALLOCATION_FAILURE";
-               case CL.CL_OUT_OF_RESOURCES:
-                       return "CL_OUT_OF_RESOURCES";
-               case CL.CL_OUT_OF_HOST_MEMORY:
-                       return "CL_OUT_OF_HOST_MEMORY";
-               case CL.CL_PROFILING_INFO_NOT_AVAILABLE:
-                       return "CL_PROFILING_INFO_NOT_AVAILABLE";
-               case CL.CL_MEM_COPY_OVERLAP:
-                       return "CL_MEM_COPY_OVERLAP";
-               case CL.CL_IMAGE_FORMAT_MISMATCH:
-                       return "CL_IMAGE_FORMAT_MISMATCH";
-               case CL.CL_IMAGE_FORMAT_NOT_SUPPORTED:
-                       return "CL_IMAGE_FORMAT_NOT_SUPPORTED";
-               case CL.CL_BUILD_PROGRAM_FAILURE:
-                       return "CL_BUILD_PROGRAM_FAILURE";
-               case CL.CL_MAP_FAILURE:
-                       return "CL_MAP_FAILURE";
-               case CL.CL_MISALIGNED_SUB_BUFFER_OFFSET:
-                       return "CL_MISALIGNED_SUB_BUFFER_OFFSET";
-               case CL.CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST:
-                       return "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST";
-               case CL.CL_COMPILE_PROGRAM_FAILURE:
-                       return "CL_COMPILE_PROGRAM_FAILURE";
-               case CL.CL_LINKER_NOT_AVAILABLE:
-                       return "CL_LINKER_NOT_AVAILABLE";
-               case CL.CL_LINK_PROGRAM_FAILURE:
-                       return "CL_LINK_PROGRAM_FAILURE";
-               case CL.CL_DEVICE_PARTITION_FAILED:
-                       return "CL_DEVICE_PARTITION_FAILED";
-               case CL.CL_KERNEL_ARG_INFO_NOT_AVAILABLE:
-                       return "CL_KERNEL_ARG_INFO_NOT_AVAILABLE";
-
-               case CL.CL_INVALID_VALUE:
-                       return "CL_INVALID_VALUE";
-               case CL.CL_INVALID_DEVICE_TYPE:
-                       return "CL_INVALID_DEVICE_TYPE";
-               case CL.CL_INVALID_PLATFORM:
-                       return "CL_INVALID_PLATFORM";
-               case CL.CL_INVALID_DEVICE:
-                       return "CL_INVALID_DEVICE";
-               case CL.CL_INVALID_CONTEXT:
-                       return "CL_INVALID_CONTEXT";
-               case CL.CL_INVALID_QUEUE_PROPERTIES:
-                       return "CL_INVALID_QUEUE_PROPERTIES";
-               case CL.CL_INVALID_COMMAND_QUEUE:
-                       return "CL_INVALID_COMMAND_QUEUE";
-               case CL.CL_INVALID_HOST_PTR:
-                       return "CL_INVALID_HOST_PTR";
-               case CL.CL_INVALID_MEM_OBJECT:
-                       return "CL_INVALID_MEM_OBJECT";
-               case CL.CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:
-                       return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR";
-               case CL.CL_INVALID_IMAGE_SIZE:
-                       return "CL_INVALID_IMAGE_SIZE";
-               case CL.CL_INVALID_SAMPLER:
-                       return "CL_INVALID_SAMPLER";
-               case CL.CL_INVALID_BINARY:
-                       return "CL_INVALID_BINARY";
-               case CL.CL_INVALID_BUILD_OPTIONS:
-                       return "CL_INVALID_BUILD_OPTIONS";
-               case CL.CL_INVALID_PROGRAM:
-                       return "CL_INVALID_PROGRAM";
-               case CL.CL_INVALID_PROGRAM_EXECUTABLE:
-                       return "CL_INVALID_PROGRAM_EXECUTABLE";
-               case CL.CL_INVALID_KERNEL_NAME:
-                       return "CL_INVALID_KERNEL_NAME";
-               case CL.CL_INVALID_KERNEL_DEFINITION:
-                       return "CL_INVALID_KERNEL_DEFINITION";
-               case CL.CL_INVALID_KERNEL:
-                       return "CL_INVALID_KERNEL";
-               case CL.CL_INVALID_ARG_INDEX:
-                       return "CL_INVALID_ARG_INDEX";
-               case CL.CL_INVALID_ARG_VALUE:
-                       return "CL_INVALID_ARG_VALUE";
-               case CL.CL_INVALID_ARG_SIZE:
-                       return "CL_INVALID_ARG_SIZE";
-               case CL.CL_INVALID_KERNEL_ARGS:
-                       return "CL_INVALID_KERNEL_ARGS";
-               case CL.CL_INVALID_WORK_DIMENSION:
-                       return "CL_INVALID_WORK_DIMENSION";
-               case CL.CL_INVALID_WORK_GROUP_SIZE:
-                       return "CL_INVALID_WORK_GROUP_SIZE";
-               case CL.CL_INVALID_WORK_ITEM_SIZE:
-                       return "CL_INVALID_WORK_ITEM_SIZE";
-               case CL.CL_INVALID_GLOBAL_OFFSET:
-                       return "CL_INVALID_GLOBAL_OFFSET";
-               case CL.CL_INVALID_EVENT_WAIT_LIST:
-                       return "CL_INVALID_EVENT_WAIT_LIST";
-               case CL.CL_INVALID_EVENT:
-                       return "CL_INVALID_EVENT";
-               case CL.CL_INVALID_OPERATION:
-                       return "CL_INVALID_OPERATION";
-               case CL.CL_INVALID_GL_OBJECT:
-                       return "CL_INVALID_GL_OBJECT";
-               case CL.CL_INVALID_BUFFER_SIZE:
-                       return "CL_INVALID_BUFFER_SIZE";
-               case CL.CL_INVALID_MIP_LEVEL:
-                       return "CL_INVALID_MIP_LEVEL";
-               case CL.CL_INVALID_GLOBAL_WORK_SIZE:
-                       return "CL_INVALID_GLOBAL_WORK_SIZE";
-               case CL.CL_INVALID_PROPERTY:
-                       return "CL_INVALID_PROPERTY";
-               case CL.CL_INVALID_IMAGE_DESCRIPTOR:
-                       return "CL_INVALID_IMAGE_DESCRIPTOR";
-               case CL.CL_INVALID_COMPILER_OPTIONS:
-                       return "CL_INVALID_COMPILER_OPTIONS";
-               case CL.CL_INVALID_LINKER_OPTIONS:
-                       return "CL_INVALID_LINKER_OPTIONS";
-               case CL.CL_INVALID_DEVICE_PARTITION_COUNT:
-                       return "CL_INVALID_DEVICE_PARTITION_COUNT";
-               default:
-                       return "Unknown error";
-               }
+       public CLException(int res) {
+               this.res = res;
        }
 }
index 85e073e..6686295 100644 (file)
@@ -16,6 +16,8 @@
  */
 package au.notzed.zcl;
 
+import jdk.incubator.foreign.MemoryAddress;
+
 /**
  * Extendable object. These keep track of the platform and api revision to be
  * able to lookup extension pointers efficiently.
@@ -28,7 +30,7 @@ public abstract class CLExtendable extends CLObject {
         */
        final int apiVersion;
 
-       public CLExtendable(long p) {
+       public CLExtendable(MemoryAddress p) {
                super(p);
 
                platform = initPlatform();
@@ -51,6 +53,15 @@ public abstract class CLExtendable extends CLObject {
                return apiVersion;
        }
 
+       public void requireAPIVersion(int version) throws UnsupportedOperationException {
+               if (apiVersion < version)
+                       throw new UnsupportedOperationException("Requires version " + ((apiVersion >> 8) & 0xff) + "." + (apiVersion & 0xff));
+       }
+
+       public boolean haveAPIVersion(int version) {
+               return apiVersion >= version;
+       }
+
        /**
         * Retrieve an extension interface for this object.Used by implementors of
         * CLExtenable.
@@ -61,6 +72,7 @@ public abstract class CLExtendable extends CLObject {
         * @return
         */
        protected <T extends CLExtension> T getExtension(Class<T> klass, int id) {
-               return platform.getExtension(klass, id);
+               //return platform.getExtension(klass, id);
+               return null;
        }
 }
index 16e4520..e4e165e 100644 (file)
@@ -16,6 +16,9 @@
  */
 package au.notzed.zcl;
 
+import jdk.incubator.foreign.MemoryAddress;
+import java.lang.invoke.MethodHandle;
+
 /**
  * Experimental code for extension support.
  * <p>
@@ -30,15 +33,15 @@ package au.notzed.zcl;
  */
 public abstract class CLExtension extends CLObject {
 
-       protected CLExtension(long p) {
+       protected CLExtension(MemoryAddress p) {
                super(p);
        }
 
        public abstract String getName();
 
        @Override
-       int getInfoType() {
+       MethodHandle getInfoFunc() {
                throw new UnsupportedOperationException();
        }
-       
+
 }
index b4117fb..60d4509 100644 (file)
 package au.notzed.zcl;
 
 import static au.notzed.zcl.CL.*;
+import static au.notzed.zcl.CLLib.*;
+import jdk.incubator.foreign.*;
+import api.Native;
+
+//import static au.au.notzed.zcl.CL.*;
 import java.nio.ByteBuffer;
+import java.lang.invoke.MethodHandle;
 
 /**
  * Interface for cl_image.
@@ -31,21 +37,28 @@ public class CLImage<T> extends CLMemory {
         * image object.
         *
         * @param p Native pointer.
+        * @param seg A segment backing the image on the host.
         */
-       public CLImage(long p) {
-               super(p);
+       public CLImage(MemoryAddress p) {
+               this(p, null);
+       }
+
+       public CLImage(MemoryAddress p, MemorySegment seg) {
+               super(p, seg);
        }
 
-       static void release(long p) {
+       static void release(MemoryAddress p) {
                CLMemory.release(p);
        }
 
        @Override
-       int getInfoType() {
-               return TYPE_IMAGE;
+       MethodHandle getInfoFunc() {
+               return clGetImageInfo;
        }
 
-       public native CLImageFormat getFormat();
+       public CLImageFormat getFormat() {
+               return getInfoJava(CL_IMAGE_BUFFER, clGetImageInfo, CLImageFormat::fromNative);
+       }
 
        public long getElementSize() {
                return getInfoSizeT(CL_IMAGE_ELEMENT_SIZE);
@@ -76,7 +89,7 @@ public class CLImage<T> extends CLMemory {
        }
 
        public CLBuffer getBuffer() {
-               return getInfoAny(CTYPE_BUFFER, CL_IMAGE_BUFFER);
+               return getInfoAny(CL_IMAGE_BUFFER, clGetImageInfo, CLBuffer::new);
        }
 
        public int getNumMipLevels() {
@@ -89,7 +102,7 @@ public class CLImage<T> extends CLMemory {
 
        @Override
        public String toString() {
-               return String.format("[%s: %dx%dx%d  0x%x]", getClass().getSimpleName(), getWidth(), getHeight(), getDepth(), getP());
+               return String.format("[%s: %dx%dx%d  0x%x]", getClass().getSimpleName(), getWidth(), getHeight(), getDepth(), addr().offset());
        }
 
        /**
index aacc985..f12e299 100644 (file)
  */
 package au.notzed.zcl;
 
+import jdk.incubator.foreign.*;
+import api.Native;
+import api.Allocator;
+
 /**
  * Holder for cl_image_desc equivalent.
+ * <h2>panama notes</h2>
+   To maintain compatability with the
+ *  previous api this remains as a simple pojo and marshalling is
+ *  done as required.
+<p>
+This also means it doesn't have to deal with allocation and deallocation and so on.
  */
 public class CLImageDesc {
 
        public int imageType;
-       public int imageWidth;
-       public int imageHeight;
-       public int imageDepth;
-       public int imageArraySize;
-       public int imageRowPitch;
-       public int imageSlicePitch;
-       public int numMipLevels;
-       public int numSamples;
+       public long imageWidth;
+       public long imageHeight;
+       public long imageDepth;
+       public long imageArraySize;
+       public long imageRowPitch;
+       public long imageSlicePitch;
+       public long numMipLevels;
+       public long numSamples;
        public CLMemory memObject;
 
        public CLImageDesc() {
        }
 
-       public CLImageDesc(int imageType, int imageWidth, int imageHeight, int imageDepth, int imageArraySize, int imageRowPitch, int imageSlicePitch, int numMipLevels, int numSamples, CLMemory memObject) {
+       public CLImageDesc(int imageType, long imageWidth, long imageHeight, long imageDepth, long imageArraySize, long imageRowPitch, long imageSlicePitch, long numMipLevels, long numSamples, CLMemory memObject) {
                this.imageType = imageType;
                this.imageWidth = imageWidth;
                this.imageHeight = imageHeight;
@@ -70,4 +80,35 @@ public class CLImageDesc {
        public static CLImageDesc create2D(int w, int h, int count) {
                return new CLImageDesc(CL.CL_MEM_OBJECT_IMAGE2D_ARRAY, w, h, 1, count, 0, 0, 0, 0, null);
        }
+
+
+       /*
+        * 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 <T extends CLProperty> MemoryAddress toNative(Allocator frame, CLImageDesc d) {
+               MemoryAddress addr = frame.alloca(sizeof);
+
+               Native.setInt(addr, d.imageType);
+               Native.setLong(addr, 1, d.imageWidth);
+               Native.setLong(addr, 2, d.imageHeight);
+               Native.setLong(addr, 3, d.imageDepth);
+               Native.setLong(addr, 4, d.imageArraySize);
+               Native.setLong(addr, 5, d.imageRowPitch);
+               Native.setLong(addr, 6, d.imageSlicePitch);
+               Native.setLong(addr, 7, d.numMipLevels);
+               Native.setLong(addr, 8, d.numSamples);
+               Native.setAddr(addr, 9, Native.addr(d.memObject));
+
+               return addr;
+       }
+
+       public static CLImageFormat fromNative(MemoryAddress addr) {
+               return new CLImageFormat(Native.getInt(addr), Native.getInt(addr, 1));
+       }
+
+       public static final long sizeof = 72;
+
+       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}]"); }
+
 }
index d870024..f419bc6 100644 (file)
  */
 package au.notzed.zcl;
 
+import jdk.incubator.foreign.*;
+import api.Native;
+import api.Allocator;
+
 /**
  * Holder for cl_image_format equivalent.
+ *
+ * <h2>panama notes</h2>
+   To maintain compatability with the
+ *  previous api this remains as a simple pojo and marshalling is
+ *  done as required.
+<p>
+This also means it doesn't have to deal with allocation and deallocation and so on.
+
  */
 public class CLImageFormat {
 
@@ -80,4 +92,26 @@ public class CLImageFormat {
                                getChannelOrder(), getChannelDataType());
        }
 
+       /*
+        * 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 <T extends CLProperty> MemoryAddress toNative(Allocator frame, CLImageFormat fmt) {
+               MemoryAddress addr = frame.alloca(2*4);
+
+               Native.setInt(addr, fmt.channelOrder);
+               Native.setInt(addr, 1, fmt.channelDataType);
+
+               return addr;
+       }
+
+       static CLImageFormat fromNative(MemoryAddress addr) {
+               return new CLImageFormat(Native.getInt(addr), Native.getInt(addr, 1));
+       }
+
+       static CLImageFormat fromNative(MemorySegment addr) {
+               return fromNative(addr.baseAddress());
+       }
+
+       static MemoryLayout layout() { return Native.parseStruct("[u32(image_channel_order)u32(image_channel_data_type)]"); }
 }
index 899e206..aaade1e 100644 (file)
  */
 package au.notzed.zcl;
 
-import java.nio.Buffer;
 import static au.notzed.zcl.CL.*;
+import static au.notzed.zcl.CLLib.*;
+import jdk.incubator.foreign.*;
+import api.Native;
+import api.Allocator;
+import api.Memory;
+
+import java.nio.Buffer;
 import java.nio.ByteBuffer;
+import java.lang.invoke.MethodHandle;
 
 /**
  * Interface for cl_kernel.
@@ -33,15 +40,24 @@ public class CLKernel extends CLObject {
         *
         * @param p Native pointer.
         */
-       public CLKernel(long p) {
+       CLKernel(MemoryAddress p) {
                super(p);
        }
 
-       private native static void release(long p);
+       static CLKernel create(MemoryAddress p) {
+               return Native.resolve(p, CLKernel::new);
+       }
+
+       private static void release(MemoryAddress p) {
+               try {
+                       clReleaseKernel(p);
+               } catch (Throwable t) {
+               }
+       }
 
        @Override
-       int getInfoType() {
-               return TYPE_KERNEL;
+       MethodHandle getInfoFunc() {
+               return clGetKernelInfo;
        }
 
        /**
@@ -51,7 +67,32 @@ public class CLKernel extends CLObject {
         * @throws CLRuntimeException
         * @since OpenCL 2.1
         */
-       public native CLKernel cloneKernel() throws CLRuntimeException;
+       public CLKernel cloneKernel() throws CLRuntimeException {
+               if (clCloneKernel != null) {
+                       // ??
+               }
+               throw new UnsupportedOperationException();
+       }
+
+       /**
+        * Calls clGetKernelArg.
+        * pval is a  pointer to the value.
+        */
+       private void setKernelArg(int index, long size, MemoryAddress pval) {
+               try {
+                       int res;
+
+                       res = clSetKernelArg(addr(), index, size, pval);
+
+                       if (res != 0)
+                               throw new CLRuntimeException(res);
+               } catch (CLRuntimeException t) {
+                       // oh joy
+                       throw t;
+               } catch (Throwable t) {
+                       throw new RuntimeException(t);
+               }
+       }
 
        /**
         * Calls clSetKernelArg with the arguments.
@@ -62,7 +103,7 @@ public class CLKernel extends CLObject {
         * ignored. May be null to setArg the size of a local scope parameter.
         * @param offset Offset in buffer.
         */
-       public native void setArg(int index, long size, Buffer buffer, long offset);
+       //public native void setArg(int index, long size, Buffer buffer, long offset);
 
        /**
         * Set the size of a parameter that is of a local scope.
@@ -71,7 +112,7 @@ public class CLKernel extends CLObject {
         * @param size The size to reserve in bytes.
         */
        public void setArgLDS(int index, long size) {
-               setArg(index, size, null, 0);
+               setKernelArg(index, size, MemoryAddress.NULL);
        }
 
        /**
@@ -80,7 +121,14 @@ public class CLKernel extends CLObject {
         * @param index
         * @param o
         */
-       public native void setArg(int index, CLObject o);
+       public void setArg(int index, CLObject o) {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress pval = frame.alloca(8);
+
+                       setAddr(pval, o.addr());
+                       setKernelArg(index, 8, pval);
+               }
+       }
 
        /**
         * Set a byte-valued argument. Equates to OpenCL types char, unsigned char,
@@ -89,7 +137,14 @@ public class CLKernel extends CLObject {
         * @param index
         * @param val
         */
-       public native void setArg(int index, byte val);
+       public void setArg(int index, byte val) {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress pval = frame.alloca(8);
+
+                       setByte(pval, val);
+                       setKernelArg(index, 1, pval);
+               }
+       }
 
        /**
         * Set a short-valued argument. Equates to OpenCL types short, unsigned
@@ -98,7 +153,14 @@ public class CLKernel extends CLObject {
         * @param index
         * @param val
         */
-       public native void setArg(int index, short val);
+       public void setArg(int index, short val) {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress pval = frame.alloca(8);
+
+                       setShort(pval, val);
+                       setKernelArg(index, 2, pval);
+               }
+       }
 
        /**
         * Set an integer-valued argument. Equates to OpenCL types int, unsigned
@@ -107,7 +169,14 @@ public class CLKernel extends CLObject {
         * @param index
         * @param val
         */
-       public native void setArg(int index, int val);
+       public void setArg(int index, int val) {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress pval = frame.alloca(8);
+
+                       setInt(pval, val);
+                       setKernelArg(index, 4, pval);
+               }
+       }
 
        /**
         * Set a long-valued argument. Equates (somewhat conveniently) to the OpenCL
@@ -116,7 +185,14 @@ public class CLKernel extends CLObject {
         * @param index
         * @param val
         */
-       public native void setArg(int index, long val);
+       public void setArg(int index, long val) {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress pval = frame.alloca(8);
+
+                       setLong(pval, val);
+                       setKernelArg(index, 8, pval);
+               }
+       }
 
        /**
         * Set a float-valued argument.
@@ -124,7 +200,14 @@ public class CLKernel extends CLObject {
         * @param index
         * @param val
         */
-       public native void setArg(int index, float val);
+       public void setArg(int index, float val) {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress pval = frame.alloca(8);
+
+                       setFloat(pval, val);
+                       setKernelArg(index, 4, pval);
+               }
+       }
 
        /**
         * Set a double-valued argument.
@@ -132,7 +215,15 @@ public class CLKernel extends CLObject {
         * @param index
         * @param val
         */
-       public native void setArg(int index, double val);
+       public void setArg(int index, double val) {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress pval = frame.alloca(8);
+
+                       setDouble(pval, val);
+                       setKernelArg(index, 8, pval);
+               }
+       }
+
 
        /**
         * Set SVM argument.
@@ -140,7 +231,9 @@ public class CLKernel extends CLObject {
         * @param index
         * @param svm MUST have been allocated using CLContext.SVMAlloc().
         */
-       public native void setArg(int index, ByteBuffer svm);
+       public void setArg(int index, ByteBuffer svm) {
+               throw new UnsupportedOperationException("not yet");
+       }
 
        /**
         * Set a multi-element byte argument. This may be used to setArg vector
@@ -149,7 +242,15 @@ public class CLKernel extends CLObject {
         * @param index
         * @param val
         */
-       public native void setArg(int index, byte... val);
+       public void setArg(int index, byte... val) {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress pval = frame.alloca(val.length);
+
+                       for (int i=0;i<val.length;i++)
+                               setByte(pval, i, val[i]);
+                       setKernelArg(index, val.length, pval);
+               }
+       }
 
        /**
         * Set a multi-element short argument. This may be used to setArg vector
@@ -158,7 +259,15 @@ public class CLKernel extends CLObject {
         * @param index
         * @param val
         */
-       public native void setArg(int index, short... val);
+       public void setArg(int index, short... val)  {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress pval = frame.alloca(val.length * 2);
+
+                       for (int i=0;i<val.length;i++)
+                               setShort(pval, i, val[i]);
+                       setKernelArg(index, val.length * 2, pval);
+               }
+       }
 
        /**
         * Set a multi-element integer argument. This may be used to setArg vector
@@ -167,7 +276,15 @@ public class CLKernel extends CLObject {
         * @param index
         * @param val
         */
-       public native void setArg(int index, int... val);
+       public void setArg(int index, int... val) {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress pval = frame.alloca(val.length * 4);
+
+                       for (int i=0;i<val.length;i++)
+                               setInt(pval, i, val[i]);
+                       setKernelArg(index, val.length * 4, pval);
+               }
+       }
 
        /**
         * Set a multi-element long argument. This may be used to setArg vector
@@ -176,7 +293,15 @@ public class CLKernel extends CLObject {
         * @param index
         * @param val
         */
-       public native void setArg(int index, long... val);
+       public void setArg(int index, long... val) {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress pval = frame.alloca(val.length * 8);
+
+                       for (int i=0;i<val.length;i++)
+                               setLong(pval, i, val[i]);
+                       setKernelArg(index, val.length * 8, pval);
+               }
+       }
 
        /**
         * Set a multi-element float argument. This may be used to setArg vector
@@ -185,7 +310,15 @@ public class CLKernel extends CLObject {
         * @param index
         * @param val
         */
-       public native void setArg(int index, float... val);
+       public void setArg(int index, float... val) {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress pval = frame.alloca(val.length * 4);
+
+                       for (int i=0;i<val.length;i++)
+                               setFloat(pval, i, val[i]);
+                       setKernelArg(index, val.length * 4, pval);
+               }
+       }
 
        /**
         * Set a multi-element double argument. This may be used to setArg vector
@@ -194,7 +327,15 @@ public class CLKernel extends CLObject {
         * @param index
         * @param val
         */
-       public native void setArg(int index, double... val);
+       public void setArg(int index, double... val) {
+               try (Allocator frame = Memory.stack()) {
+                       MemoryAddress pval = frame.alloca(val.length * 8);
+
+                       for (int i=0;i<val.length;i++)
+                               setDouble(pval, i, val[i]);
+                       setKernelArg(index, val.length * 8, pval);
+               }
+       }
 
        /**
         * Sets a number of arguments.
@@ -243,22 +384,17 @@ public class CLKernel extends CLObject {
        }
 
        public CLContext getContext() {
-               return getInfoAny(CTYPE_CONTEXT, CL_KERNEL_CONTEXT);
+               return getInfoAny(CL_KERNEL_CONTEXT, clGetKernelInfo, CLContext::new);
        }
 
        public CLProgram getProgram() {
-               return getInfoAny(CTYPE_PROGRAM, CL_KERNEL_PROGRAM);
+               return getInfoAny(CL_KERNEL_PROGRAM, clGetKernelInfo, CLProgram::new);
        }
 
        public String getAttributes() {
                return getInfoString(CL_KERNEL_ATTRIBUTES);
        }
 
-       /* cl_kernel_work_group_info */
-       native <T> T getWorkGroupInfoAny(CLDevice dev, int otype, int param_name) throws CLRuntimeException;
-
-       native <T> T getWorkGroupInfoAnyV(CLDevice dev, int otype, int param_name) throws CLRuntimeException;
-
        /**
         * gets CL_KERNEL_GLOBAL_WORK_SIZE
         *
@@ -266,58 +402,56 @@ public class CLKernel extends CLObject {
         * @return
         */
        public long[] getGlobalWorkSize(CLDevice device) {
-               return getWorkGroupInfoAnyV(device, CTYPE_SIZE_T, CL_KERNEL_GLOBAL_WORK_SIZE);
+               return getInfoLongV(device, CL_KERNEL_GLOBAL_WORK_SIZE, clGetKernelWorkGroupInfo);
        }
 
-       public long getWorkGroupSize(CLDevice device) {
-               return getWorkGroupInfoAny(device, CTYPE_SIZE_T, CL_KERNEL_WORK_GROUP_SIZE);
+        public long getWorkGroupSize(CLDevice device) {
+                return getInfoSizeT(device, CL_KERNEL_WORK_GROUP_SIZE, clGetKernelWorkGroupInfo);
        }
 
        public long[] getCompileWorkGroupSize(CLDevice device) {
-               return getWorkGroupInfoAnyV(device, CTYPE_SIZE_T, CL_KERNEL_COMPILE_WORK_GROUP_SIZE);
+               return getInfoLongV(device, CL_KERNEL_COMPILE_WORK_GROUP_SIZE, clGetKernelWorkGroupInfo);
        }
 
        public long getLocalMemSize(CLDevice device) {
-               return getWorkGroupInfoAny(device, CTYPE_LONG, CL_KERNEL_LOCAL_MEM_SIZE);
+               return getInfoSizeT(device, CL_KERNEL_LOCAL_MEM_SIZE, clGetKernelWorkGroupInfo);
        }
 
        public long getPreferredWorkGroupSizeMultiple(CLDevice device) {
-               return getWorkGroupInfoAny(device, CTYPE_SIZE_T, CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE);
+               return getInfoSizeT(device, CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE, clGetKernelWorkGroupInfo);
        }
 
        public long getPrivateMemSize(CLDevice device) {
-               return getWorkGroupInfoAny(device, CTYPE_SIZE_T, CL_KERNEL_PRIVATE_MEM_SIZE);
+               return getInfoSizeT(device, CL_KERNEL_PRIVATE_MEM_SIZE, clGetKernelWorkGroupInfo);
        }
 
-       /* cl_kernel_sub_group_info */
-       // this assumes all size_t's atm.
-       native <T> T getSubGroupInfoAny(CLDevice dev, int otype, int param_name, long[] input) throws CLRuntimeException;
-
-       native <T> T getSubGroupInfoAnyV(CLDevice dev, int otype, int param_name, long[] input) throws CLRuntimeException;
-
        /**
         * @since OpenCL 2.1
         */
        public long getMaxSubGroupSizeForNDRange(CLDevice device, long[] range) {
-               return getSubGroupInfoAny(device, CTYPE_SIZE_T, CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE, range);
+               // wrong, these all need to pass a range too
+               //return getInfoSizeT(CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE, device, clGetKernelSubGroupInfo);
+               throw new UnsupportedOperationException();
        }
 
        /**
         * @since OpenCL 2.1
         */
        public long getSubGroupCountForNDRange(CLDevice device, long[] range) {
-               return getSubGroupInfoAny(device, CTYPE_SIZE_T, CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE, range);
+               //return getInfoSizeT(CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE, device, clGetKernelSubGroupInfo);
+               throw new UnsupportedOperationException();
        }
 
-       /**
-        * @since OpenCL 2.1
-        */
-       public long[] getLocalSizeForSubGroupCount(CLDevice device, long count) {
-               return getSubGroupInfoAnyV(device, CTYPE_SIZE_T, CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT, new long[]{count});
-       }
+       // /**
+       //  * @since OpenCL 2.1
+       //  */
+       // public long[] getLocalSizeForSubGroupCount(CLDevice device, long count) {
+       //      return getSubGroupInfoAnyV(device, CTYPE_SIZE_T, CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT, new long[]{count});
+       // }
 
        public long getMaxNumSubGroups(CLDevice device) {
-               return getSubGroupInfoAny(device, CTYPE_SIZE_T, CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE, null);
+               throw new UnsupportedOperationException();
+               //return getInfoSizeT(CL_KERNEL_MAX_NUM_SUB_GROUPS, device, clGetKernelSubGroupInfo);
        }
 
        /* cl_kernel_exec_info */
@@ -330,11 +464,6 @@ public class CLKernel extends CLObject {
         */
        public final static int CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM = 0x11B7;
 
-       /* cl_kernel_arg_info */
-       native <T> T getArgInfoAny(int index, int otype, int param_name) throws CLRuntimeException;
-
-       native <T> T getArgInfoAnyV(int index, int otype, int param_name) throws CLRuntimeException;
-
        /**
         * gets CL_KERNEL_ARG_ADDRESS_QUALIFIER.
         *
@@ -344,7 +473,7 @@ public class CLKernel extends CLObject {
         * @since OpenCL 1.2
         */
        public int getArgAddressQualifier(int index) throws UnsupportedOperationException {
-               return getArgInfoAny(index, CTYPE_INT, CL_KERNEL_ARG_ADDRESS_QUALIFIER);
+               return getInfoInt(index, CL_KERNEL_ARG_ADDRESS_QUALIFIER, clGetKernelArgInfo);
        }
 
        /**
@@ -356,7 +485,7 @@ public class CLKernel extends CLObject {
         * @since OpenCL 1.2
         */
        public int getArgAccessQualifier(int index) throws UnsupportedOperationException {
-               return getArgInfoAny(index, CTYPE_INT, CL_KERNEL_ARG_ACCESS_QUALIFIER);
+               return getInfoInt(index, CL_KERNEL_ARG_ACCESS_QUALIFIER, clGetKernelArgInfo);
        }
 
        /**
@@ -368,7 +497,7 @@ public class CLKernel extends CLObject {
         * @since OpenCL 1.2
         */
        public String getArgTypeName(int index) throws UnsupportedOperationException {
-               return fromInfoString(getArgInfoAnyV(index, CTYPE_BYTE, CL_KERNEL_ARG_TYPE_NAME));
+               return getInfoString(index, CL_KERNEL_ARG_TYPE_NAME, clGetKernelArgInfo);
        }
 
        /**
@@ -380,7 +509,7 @@ public class CLKernel extends CLObject {
         * @since OpenCL 1.2
         */
        public long getArgTypeQualifier(int index) throws UnsupportedOperationException {
-               return getArgInfoAny(index, CTYPE_LONG, CL_KERNEL_ARG_TYPE_QUALIFIER);
+               return getInfoLong(index, CL_KERNEL_ARG_TYPE_QUALIFIER, clGetKernelArgInfo);
        }
 
        /**
@@ -392,6 +521,6 @@ public class CLKernel extends CLObject {
         * @since OpenCL 1.2
         */
        public String getArgName(int index) throws UnsupportedOperationException {
-               return fromInfoString(getArgInfoAnyV(index, CTYPE_BYTE, CL_KERNEL_ARG_NAME));
+               return getInfoString(index, CL_KERNEL_ARG_NAME, clGetKernelArgInfo);
        }
 }
index c4bd264..0a2923b 100644 (file)
  */
 package au.notzed.zcl;
 
-import java.nio.ByteBuffer;
 import static au.notzed.zcl.CL.*;
+import static au.notzed.zcl.CLLib.*;
+import jdk.incubator.foreign.*;
+import api.Native;
+import api.Callback;
+import api.Memory;
+import api.Allocator;
+import java.lang.invoke.MethodHandle;
+
+import java.nio.ByteBuffer;
+//import static au.au.notzed.zcl.CL.*;
 import java.nio.ByteOrder;
+import jdk.incubator.foreign.*;
 
 /**
  * Interface for cl_mem.
@@ -58,11 +68,81 @@ import java.nio.ByteOrder;
  */
 public abstract class CLMemory extends CLObject {
 
-       CLMemory(long p) {
+       /**
+        * If use USE_HOST_PTR was used then this keeps track of the
+        * host ptr reference both to avoid java freeing it and so it
+        * can be retrieved later.
+        */
+       MemorySegment seg;
+
+       /**
+        * If set, retain reference for lifetime of memory.
+        */
+       Callback<CLNotify<CLMemory>> destroyCallback;
+
+       CLMemory(MemoryAddress p, MemorySegment seg) {
                super(p);
+
+               this.seg = seg;
        }
 
-       native static void release(long p);
+       @Override
+       MethodHandle getInfoFunc() {
+               return clGetMemObjectInfo;
+       }
+
+       public static CLMemory create(MemoryAddress p) {
+               if (p.offset() == 0)
+                       return null;
+
+               // This is basically a workaround so that setMemObjectDestructorCallback passes
+               // the right type without changing the api or using reflection.
+               // Probably better solutions.
+
+               try (Allocator a = Memory.stack()) {
+                       MemoryAddress addr = getInfo(p, CL_MEM_TYPE, clGetMemObjectInfo, a, 4);
+                       int type = getInt(addr);
+
+                       switch (type) {
+                       case CL_MEM_OBJECT_BUFFER:
+                               return Native.resolve(p, CLBuffer::new);
+                       case CL_MEM_OBJECT_IMAGE2D:
+                       case CL_MEM_OBJECT_IMAGE3D:
+                       case CL_MEM_OBJECT_IMAGE2D_ARRAY:
+                       case CL_MEM_OBJECT_IMAGE1D:
+                       case CL_MEM_OBJECT_IMAGE1D_ARRAY:
+                       case CL_MEM_OBJECT_IMAGE1D_BUFFER:
+                               return Native.resolve(p, CLImage::new);
+                       case CL_MEM_OBJECT_PIPE:
+                               return Native.resolve(p, CLPipe::new);
+                       default:
+                               throw new UnsupportedOperationException();
+                       }
+               }
+       }
+
+       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) {
+               }
+       }
+
+       /**
+        * 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
+        * on the original thread of creation.
+        */
+       @Override
+       public void release() {
+               if (seg != null) {
+                       seg.close();
+                       seg = null;
+               }
+               super.release();
+       }
 
        /**
         * Call clSetMemObjectDestructorCallback.
@@ -72,14 +152,32 @@ public abstract class CLMemory extends CLObject {
         * @throws UnsupportedOperationException
         * @since OpenCL 1.1
         */
-       public native void setMemObjectDestructorCallback(CLNotify<CLMemory> notify) throws CLException, UnsupportedOperationException;
+       public void setMemObjectDestructorCallback(CLNotify<CLMemory> notify) throws CLException, UnsupportedOperationException {
+               //a bit costly perhaps
+               getContext().requireAPIVersion(CLPlatform.VERSION_1_1);
 
-       @Override
-       int getInfoType() {
-               return TYPE_MEM_OBJECT;
+               Native.release(destroyCallback);
+               if (notify != null) {
+                       destroyCallback = CLNotify.call(notify, CLMemory::create);
+
+                       try {
+                               int res = clSetMemObjectDestructorCallback(addr(), destroyCallback.addr(), MemoryAddress.NULL);
+                               if (res != 0)
+                                       throw new CLException(res);
+                       } catch (RuntimeException | Error t) {
+                               throw t;
+                       } catch (Throwable t) {
+                               throw new RuntimeException(t);
+                       }
+               } else {
+                       destroyCallback = null;
+               }
        }
 
-       public long getType() {
+       /**
+        * Get CL_MEM_TYPE
+        */
+       public int getType() {
                return getInfoInt(CL_MEM_TYPE);
        }
 
@@ -104,9 +202,15 @@ public abstract class CLMemory extends CLObject {
        /**
         * Get CL_MEM_HOST_PTR.