introduced CLGLContext, refactored dependencies, cleanup in opencl code.

10 files changed, 154 insertions, 48 deletions

diff --git a/nbproject/project.properties b/nbproject/project.properties
index 204a0320..cda7bf22 100644
--- a/nbproject/project.properties
+++ b/nbproject/project.properties
@@ -21,14 +21,23 @@ dist.dir=dist
 dist.jar=${dist.dir}/jocl.jar
 dist.javadoc.dir=${dist.dir}/javadoc
 excludes=
+
 #default value, overwrite this or pass it as -D ant property if required
 gluegen.root=../gluegen
+jogl.root=../jogl
 file.reference.gluegen-rt.jar=${gluegen.root}/build/gluegen-rt.jar
 file.reference.gluegen.jar=${gluegen.root}/build/gluegen.jar
+file.reference.jogl.all.jar=${jogl.root}/build/jogl/jogl.all.jar
+file.reference.nativewindow.all.jar=${jogl.root}/build/nativewindow/nativewindow.all.jar
+file.reference.newt.all.jar=${jogl.root}/build/newt/newt.all.jar
+
 includes=**
 jar.compress=false
 javac.classpath=\
-    ${file.reference.gluegen-rt.jar}
+    ${file.reference.gluegen-rt.jar}:\
+    ${file.reference.jogl.all.jar}:\
+    ${file.reference.newt.all.jar}:\
+    ${file.reference.nativewindow.all.jar}
 # Space-separated list of extra javac options
 javac.compilerargs=
 javac.deprecation=true
diff --git a/resources/cl-impl.cfg b/resources/cl-impl.cfg
index 9dff59bc..c2af12cf 100644
--- a/resources/cl-impl.cfg
+++ b/resources/cl-impl.cfg
@@ -30,3 +30,6 @@ Ignore clEnqueueNativeKernel
 #include custom code
 IncludeAs CustomJavaCode CLImpl clImplCustomCode.java
 IncludeAs CustomCCode clImplCustomCode.c
+
+#JavaEpilogue clCreateKernelsInProgram if(kernels!=null && CPU.is32Bit() && kernels.lenght > 1) { convert32To64(kernels); }
+
diff --git a/resources/clImplCustomCode.java b/resources/clImplCustomCode.java
index 91b1fb48..cffe9a72 100644
--- a/resources/clImplCustomCode.java
+++ b/resources/clImplCustomCode.java
@@ -51,3 +51,14 @@
         /** Entry point to C language function: <code> int32_t clBuildProgram(cl_program, uint32_t, cl_device_id * , const char * , void * ); </code>    */
         private native int clBuildProgram1(long program, int devices, Object deviceList, String options, BuildProgramCallback cb, Object userData);
 
+        
+        private final static void convert32To64(long[] values) {
+            if(values.length%2 == 1) {
+                values[values.length-1] = values[values.length/2]>>>32;
+            }
+            for (int i = values.length - 1 - values.length%2; i >= 0; i-=2) {
+                long temp = values[i/2];
+                values[i-1] = temp>>>32;
+                values[i  ] = temp & 0x00000000FFFFFFFFL;
+            }
+		}
diff --git a/resources/opencl.h b/resources/opencl.h
index 8e46611c..d1b0c795 100644
--- a/resources/opencl.h
+++ b/resources/opencl.h
@@ -1,7 +1,7 @@
 
 #include    <CL/cl_platform.h>
 #include    <CL/cl.h>
-//#include    <CL/cl_ext.h>
+#include    <CL/cl_ext.h>
 
 #ifdef _WIN32
 #include    <windows.h>    
diff --git a/src/com/mbien/opencl/CLCommandQueue.java b/src/com/mbien/opencl/CLCommandQueue.java
index d488d1c4..737e8c81 100644
--- a/src/com/mbien/opencl/CLCommandQueue.java
+++ b/src/com/mbien/opencl/CLCommandQueue.java
@@ -154,6 +154,16 @@ public class CLCommandQueue implements CLResource {
                 localWorkSize   ==0 ? null : new long[] {localWorkSize   }  );
     }
 
+    public CLCommandQueue put2DRangeKernel(CLKernel kernel, long globalWorkOffsetX, long globalWorkOffsetY,
+                                                            long globalWorkSizeX, long globalWorkSizeY,
+                                                            long localWorkSizeX, long localWorkSizeY) {
+        return this.putNDRangeKernel(
+                kernel, 2,
+                globalWorkOffsetX==0 && globalWorkOffsetY==0 ? null : new long[] {globalWorkOffsetX, globalWorkOffsetY},
+                globalWorkSizeX  ==0 && globalWorkSizeY  ==0 ? null : new long[] {globalWorkSizeX,   globalWorkSizeY  },
+                localWorkSizeX   ==0 && localWorkSizeY   ==0 ? null : new long[] {localWorkSizeX,    localWorkSizeY   }  );
+    }
+
     public CLCommandQueue putNDRangeKernel(CLKernel kernel, int workDimension, long[] globalWorkOffset, long[] globalWorkSize, long[] localWorkSize) {
 
        int ret = cl.clEnqueueNDRangeKernel(
diff --git a/src/com/mbien/opencl/CLContext.java b/src/com/mbien/opencl/CLContext.java
index 04df962b..75e4d35a 100644
--- a/src/com/mbien/opencl/CLContext.java
+++ b/src/com/mbien/opencl/CLContext.java
@@ -29,19 +29,19 @@ import static com.sun.gluegen.runtime.BufferFactory.*;
  * specified in the context.
  * @author Michael Bien
  */
-public final class CLContext implements CLResource {
+public class CLContext implements CLResource {
 
     final CL cl;
     public final long ID;
 
-    private CLDevice[] devices;
+    protected CLDevice[] devices;
 
-    private final List<CLProgram> programs;
-    private final List<CLBuffer<? extends Buffer>> buffers;
-    private final Map<CLDevice, List<CLCommandQueue>> queuesMap;
+    protected final List<CLProgram> programs;
+    protected final List<CLBuffer<? extends Buffer>> buffers;
+    protected final Map<CLDevice, List<CLCommandQueue>> queuesMap;
 
 
-    private CLContext(long contextID) {
+    protected CLContext(long contextID) {
         this.cl = CLPlatform.getLowLevelBinding();
         this.ID = contextID;
         this.programs = new ArrayList<CLProgram>();
@@ -75,7 +75,7 @@ public final class CLContext implements CLResource {
      * The platform to be used is implementation dependent.
      */
     public static final CLContext create() {
-        return createContextFromType(null, CL.CL_DEVICE_TYPE_ALL);
+        return new CLContext(createContextFromType(null, CL.CL_DEVICE_TYPE_ALL));
     }
 
     /**
@@ -121,7 +121,7 @@ public final class CLContext implements CLResource {
             properties.rewind();
         }
 
-        return createContextFromType(properties, type);
+        return new CLContext(createContextFromType(properties, type));
     }
 
     /**
@@ -143,27 +143,27 @@ public final class CLContext implements CLResource {
             properties.rewind();
         }
 
-        return createContext(properties, deviceIDs);
+        return new CLContext(createContext(properties, deviceIDs));
     }
 
-    private static final CLContext createContextFromType(LongBuffer properties, long deviceType) {
+    protected static final long createContextFromType(LongBuffer properties, long deviceType) {
 
         IntBuffer status = IntBuffer.allocate(1);
         long context = CLPlatform.getLowLevelBinding().clCreateContextFromType(properties, deviceType, null, null, status);
 
         checkForError(status.get(), "can not create CL context");
 
-        return new CLContext(context);
+        return context;
     }
 
-    private static final CLContext createContext(LongBuffer properties, long[] devices) {
+    protected static final long createContext(LongBuffer properties, long[] devices) {
 
         IntBuffer status = IntBuffer.allocate(1);
         long context = CLPlatform.getLowLevelBinding().clCreateContext(properties, devices, null, null, status);
 
         checkForError(status.get(), "can not create CL context");
 
-        return new CLContext(context);
+        return context;
     }
 
     /**
@@ -263,16 +263,6 @@ public final class CLContext implements CLResource {
         return buffer;
     }
 
-    public final <B extends Buffer> CLBuffer<B> createFromGLBuffer(B directBuffer, int glBuffer, Mem... flags) {
-        return createFromGLBuffer(directBuffer, glBuffer, Mem.flagsToInt(flags));
-    }
-    
-    public final <B extends Buffer> CLBuffer<B> createFromGLBuffer(B directBuffer, int glBuffer, int flags) {
-        CLBuffer<B> buffer = new CLBuffer<B>(this, directBuffer, glBuffer, flags);
-        buffers.add(buffer);
-        return buffer;
-    }
-
     CLCommandQueue createCommandQueue(CLDevice device, long properties) {
 
         CLCommandQueue queue = new CLCommandQueue(this, device, properties);
diff --git a/src/com/mbien/opencl/CLGLContext.java b/src/com/mbien/opencl/CLGLContext.java
new file mode 100644
index 00000000..9699c1ba
--- /dev/null
+++ b/src/com/mbien/opencl/CLGLContext.java
@@ -0,0 +1,83 @@
+package com.mbien.opencl;
+
+import java.nio.Buffer;
+import com.mbien.opencl.CLBuffer.Mem;
+import com.sun.opengl.impl.GLContextImpl;
+import com.sun.opengl.impl.macosx.cgl.MacOSXCGLContext;
+import com.sun.opengl.impl.windows.wgl.WindowsWGLContext;
+import com.sun.opengl.impl.x11.glx.X11GLXContext;
+import com.sun.opengl.impl.x11.glx.X11GLXGraphicsConfiguration;
+import java.nio.LongBuffer;
+import javax.media.nativewindow.DefaultGraphicsConfiguration;
+import javax.media.opengl.GLContext;
+
+/**
+ *
+ * @author Michael Bien
+ */
+public final class CLGLContext extends CLContext {
+
+    final long glContextID;
+
+    private CLGLContext(long clContextID, long glContextID) {
+        super(clContextID);
+        this.glContextID = glContextID;
+    }
+
+    public static CLGLContext create(GLContext glContext) {
+
+        long glID = glContext.getContext();
+
+//UNIX
+//cl_context_properties props[] = {
+//         CL_GL_CONTEXT_KHR, (cl_context_properties)glXGetCurrentContext(),
+//         CL_GLX_DISPLAY_KHR,  (cl_context_properties) glXGetCurrentDisplay(),
+//         CL_CONTEXT_PLATFORM, (cl_context_properties)cpPlatform, 0};
+
+//WIN32
+//cl_context_properties props[] = {
+//         CL_GL_CONTEXT_KHR, (cl_context_properties)TODO0,
+//         CL_WGL_HDC_KHR, (cl_context_properties)TODO 0,
+//         CL_CONTEXT_PLATFORM, (cl_context_properties)cpPlatform, 0};
+
+//MACOSX
+//cl_context_properties props[] = {
+//         CL_CGL_SHAREGROUP_KHR, (cl_context_properties)TODO 0,
+//         CL_CONTEXT_PLATFORM, (cl_context_properties)cpPlatform, 0};
+
+        GLContextImpl ctxImpl = (GLContextImpl)glContext;
+
+        DefaultGraphicsConfiguration config = (DefaultGraphicsConfiguration)ctxImpl.getDrawableImpl()
+             .getNativeWindow().getGraphicsConfiguration().getNativeGraphicsConfiguration();
+
+        LongBuffer properties = LongBuffer.allocate(5);
+        if(glContext instanceof X11GLXContext) {
+            long handle = config.getScreen().getDevice().getHandle();
+            properties.put(CLGLI.CL_GL_CONTEXT_KHR).put(glID)
+                      .put(CLGLI.CL_GLX_DISPLAY_KHR).put(handle);
+        }else if(glContext instanceof WindowsWGLContext) {
+            // TODO test on windows
+            throw new RuntimeException("cl-gl interoperability on windows not yet implemented");
+        }else if(glContext instanceof MacOSXCGLContext) {
+            // TODO test on mac
+            throw new RuntimeException("cl-gl interoperability on mac not yet implemented");
+        }
+        
+        properties.put(0).rewind(); // 0 terminated array
+
+        long clID = createContextFromType(properties, CL.CL_DEVICE_TYPE_ALL);
+
+        return new CLGLContext(clID, glID);
+    }
+
+
+    public final <B extends Buffer> CLBuffer<B> createFromGLBuffer(B directBuffer, int glBuffer, Mem... flags) {
+        return createFromGLBuffer(directBuffer, glBuffer, Mem.flagsToInt(flags));
+    }
+
+    public final <B extends Buffer> CLBuffer<B> createFromGLBuffer(B directBuffer, int glBuffer, int flags) {
+        CLBuffer<B> buffer = new CLBuffer<B>(this, directBuffer, glBuffer, flags);
+        buffers.add(buffer);
+        return buffer;
+    }
+}
diff --git a/test/com/mbien/opencl/HighLevelBindingTest.java b/test/com/mbien/opencl/HighLevelBindingTest.java
index bd6d7179..42e74667 100644
--- a/test/com/mbien/opencl/HighLevelBindingTest.java
+++ b/test/com/mbien/opencl/HighLevelBindingTest.java
@@ -241,7 +241,7 @@ public class HighLevelBindingTest {
             fail("expected exception but got none :(");
         }catch(CLException ex) {
             out.println("got expected exception:\n"+ex.getMessage());
-            assertTrue(ex.errorcode == CL.CL_INVALID_PROGRAM_EXECUTABLE);
+            assertEquals(ex.errorcode, CL.CL_INVALID_PROGRAM_EXECUTABLE);
         }
 
         program.build();
diff --git a/test/com/mbien/opencl/LowLevelBindingTest.java b/test/com/mbien/opencl/LowLevelBindingTest.java
index 83e0ca45..0a600102 100644
--- a/test/com/mbien/opencl/LowLevelBindingTest.java
+++ b/test/com/mbien/opencl/LowLevelBindingTest.java
@@ -20,26 +20,26 @@ import static com.sun.gluegen.runtime.BufferFactory.*;
 public class LowLevelBindingTest {
 
     private final static String programSource =
-              " // OpenCL Kernel Function for element by element vector addition                                                \n"
-            + "__kernel void VectorAdd(__global const int* a, __global const int* b, __global int* c, int iNumElements) {       \n"
-            + "    // get index into global data array                                                                          \n"
-            + "    int iGID = get_global_id(0);                                                                                 \n"
-            + "    // bound check (equivalent to the limit on a 'for' loop for standard/serial C code                           \n"
-            + "    if (iGID >= iNumElements)  {                                                                                 \n"
-            + "        return;                                                                                                  \n"
-            + "    }                                                                                                            \n"
-            + "    // add the vector elements                                                                                   \n"
-            + "    c[iGID] = a[iGID] + b[iGID];                                                                                 \n"
-            + "}                                                                                                                \n"
-            + "__kernel void Test(__global const int* a, __global const int* b, __global int* c, int iNumElements) {            \n"
-            + "    // get index into global data array                                                                          \n"
-            + "    int iGID = get_global_id(0);                                                                                 \n"
-            + "    // bound check (equivalent to the limit on a 'for' loop for standard/serial C code                           \n"
-            + "    if (iGID >= iNumElements)  {                                                                                 \n"
-            + "        return;                                                                                                  \n"
-            + "    }                                                                                                            \n"
-            + "    c[iGID] = iGID;                                                                                              \n"
-            + "}                                                                                                                \n";
+              " // OpenCL Kernel Function for element by element vector addition                                  \n"
+            + "kernel void VectorAdd(global const int* a, global const int* b, global int* c, int iNumElements) { \n"
+            + "    // get index into global data array                                                            \n"
+            + "    int iGID = get_global_id(0);                                                                   \n"
+            + "    // bound check (equivalent to the limit on a 'for' loop for standard/serial C code             \n"
+            + "    if (iGID >= iNumElements)  {                                                                   \n"
+            + "        return;                                                                                    \n"
+            + "    }                                                                                              \n"
+            + "    // add the vector elements                                                                     \n"
+            + "    c[iGID] = a[iGID] + b[iGID];                                                                   \n"
+            + "}                                                                                                  \n"
+            + "kernel void Test(global const int* a, global const int* b, global int* c, int iNumElements) {      \n"
+            + "    // get index into global data array                                                            \n"
+            + "    int iGID = get_global_id(0);                                                                   \n"
+            + "    // bound check (equivalent to the limit on a 'for' loop for standard/serial C code             \n"
+            + "    if (iGID >= iNumElements)  {                                                                   \n"
+            + "        return;                                                                                    \n"
+            + "    }                                                                                              \n"
+            + "    c[iGID] = iGID;                                                                                \n"
+            + "}                                                                                                  \n";
 
 
     @BeforeClass
diff --git a/test/com/mbien/opencl/testkernels.cl b/test/com/mbien/opencl/testkernels.cl
index 0790cb32..ec7e8bf6 100644
--- a/test/com/mbien/opencl/testkernels.cl
+++ b/test/com/mbien/opencl/testkernels.cl
@@ -1,6 +1,6 @@
 
     // OpenCL Kernel Function for element by element vector addition
-    __kernel void VectorAddGM(__global const int* a, __global const int* b, __global int* c, int iNumElements) {
+    kernel void VectorAddGM(global const int* a, global const int* b, global int* c, int iNumElements) {
         // get index into global data array
         int iGID = get_global_id(0);
         // bound check (equivalent to the limit on a 'for' loop for standard/serial C code
@@ -11,7 +11,7 @@
         c[iGID] = a[iGID] + b[iGID];
     }
 
-    __kernel void Test(__global const int* a, __global const int* b, __global int* c, int iNumElements) {
+    kernel void Test(global const int* a, global const int* b, global int* c, int iNumElements) {
         // get index into global data array
         int iGID = get_global_id(0);
         // bound check (equivalent to the limit on a 'for' loop for standard/serial C code