removed __ in opencl code, modified sample to use real GL-CL interoperability.

3 files changed, 55 insertions, 52 deletions

diff --git a/src/com/mbien/opencl/demos/hellojocl/VectorAdd.cl b/src/com/mbien/opencl/demos/hellojocl/VectorAdd.cl
index e8023b1..ac9dde2 100644
--- a/src/com/mbien/opencl/demos/hellojocl/VectorAdd.cl
+++ b/src/com/mbien/opencl/demos/hellojocl/VectorAdd.cl
@@ -1,6 +1,6 @@
 
     // OpenCL Kernel Function for element by element vector addition
-    __kernel void VectorAdd(__global const float* a, __global const float* b, __global float* c, int numElements) {
+    kernel void VectorAdd(global const float* a, global const float* b, global float* c, int numElements) {
 
         // get index into global data array
         int iGID = get_global_id(0);
diff --git a/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java b/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java
index 4581643..06fba82 100644
--- a/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java
+++ b/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java
@@ -2,8 +2,7 @@ package com.mbien.opencl.demos.joglinterop;
 
 import com.mbien.opencl.CLBuffer;
 import com.mbien.opencl.CLCommandQueue;
-import com.mbien.opencl.CLContext;
-import com.mbien.opencl.CLException;
+import com.mbien.opencl.CLGLContext;
 import com.mbien.opencl.CLKernel;
 import com.mbien.opencl.CLProgram;
 import com.sun.opengl.util.Animator;
@@ -26,7 +25,7 @@ import javax.swing.SwingUtilities;
 import static com.sun.opengl.util.BufferUtil.*;
 
 /**
- * Sample for interoperability between JOCL and JOGL.
+ * JOCL - JOGL interoperability example.
  * @author Michael Bien
  */
 public class GLCLInteroperabilityDemo implements GLEventListener {
@@ -47,17 +46,20 @@ public class GLCLInteroperabilityDemo implements GLEventListener {
 
     private final UserSceneInteraction usi;
 
-    private CLContext clContext;
+    private CLGLContext clContext;
     private CLKernel kernel;
     private CLCommandQueue commandQueue;
     private CLBuffer<FloatBuffer> clBuffer;
 
     private float step = 0;
 
-    public GLCLInteroperabilityDemo() throws IOException {
+    private boolean initialized = false;
+
+    public GLCLInteroperabilityDemo() {
 
         this.usi = new UserSceneInteraction();
 
+        // create direct memory buffers
         vb = newFloatBuffer(MESH_SIZE * MESH_SIZE * 4);
         ib = newIntBuffer((MESH_SIZE - 1) * (MESH_SIZE - 1) * 2 * 3);
 
@@ -69,12 +71,12 @@ public class GLCLInteroperabilityDemo implements GLEventListener {
             for (int w = 0; w < MESH_SIZE - 1; w++) {
 
                 // 0 - 3 - 2
-                ib.put(w * 6 + h * (MESH_SIZE - 1) * 6,      w + (h) * (MESH_SIZE)        );
-                ib.put(w * 6 + h * (MESH_SIZE - 1) * 6 + 1,  w + (h) * (MESH_SIZE) + 1    );
+                ib.put(w * 6 + h * (MESH_SIZE - 1) * 6,      w + (h    ) * (MESH_SIZE)    );
+                ib.put(w * 6 + h * (MESH_SIZE - 1) * 6 + 1,  w + (h    ) * (MESH_SIZE) + 1);
                 ib.put(w * 6 + h * (MESH_SIZE - 1) * 6 + 2,  w + (h + 1) * (MESH_SIZE) + 1);
 
                 // 0 - 2 - 1
-                ib.put(w * 6 + h * (MESH_SIZE - 1) * 6 + 3,  w + (h) * (MESH_SIZE)        );
+                ib.put(w * 6 + h * (MESH_SIZE - 1) * 6 + 3,  w + (h    ) * (MESH_SIZE)    );
                 ib.put(w * 6 + h * (MESH_SIZE - 1) * 6 + 4,  w + (h + 1) * (MESH_SIZE) + 1);
                 ib.put(w * 6 + h * (MESH_SIZE - 1) * 6 + 5,  w + (h + 1) * (MESH_SIZE)    );
 
@@ -132,20 +134,31 @@ public class GLCLInteroperabilityDemo implements GLEventListener {
 
         gl.glGenBuffers(glObjects.length, glObjects, 0);
 
-        gl.glBindBuffer(GL2.GL_ELEMENT_ARRAY_BUFFER, glObjects[INDICES]);
-        gl.glBufferData(GL2.GL_ELEMENT_ARRAY_BUFFER, ib.capacity() * SIZEOF_INT, ib, GL2.GL_STATIC_DRAW);
-        gl.glBindBuffer(GL2.GL_ELEMENT_ARRAY_BUFFER, 0);
+//        gl.glBindBuffer(GL2.GL_ELEMENT_ARRAY_BUFFER, glObjects[INDICES]);
+//        gl.glBufferData(GL2.GL_ELEMENT_ARRAY_BUFFER, ib.capacity() * SIZEOF_INT, ib, GL2.GL_STATIC_DRAW);
+//        gl.glBindBuffer(GL2.GL_ELEMENT_ARRAY_BUFFER, 0);
 
         gl.glEnableClientState(GL2.GL_VERTEX_ARRAY);
             gl.glBindBuffer(GL2.GL_ARRAY_BUFFER, glObjects[VERTICES]);
             gl.glBufferData(GL2.GL_ARRAY_BUFFER, vb.capacity() * SIZEOF_FLOAT, vb, GL2.GL_DYNAMIC_DRAW);
-            gl.glBindBuffer(GL2.GL_ARRAY_BUFFER, 0);
+//            gl.glBindBuffer(GL2.GL_ARRAY_BUFFER, 0);
         gl.glDisableClientState(GL2.GL_VERTEX_ARRAY);
 
-        // OpenCL
+        gl.glFinish();
+        initCL(drawable);
+
+        pushPerspectiveView(gl);
+
+        // start rendering thread
+        Animator animator = new Animator(drawable);
+        animator.start();
+    }
+
+    private void initCL(GLAutoDrawable drawable) {
+
         CLProgram program;
         try {
-            clContext = CLContext.create();
+            clContext = CLGLContext.create(drawable.getContext());
             program = clContext.createProgram(getClass().getResourceAsStream("JoglInterop.cl"));
             program.build();
             System.out.println(program.getBuildLog());
@@ -156,25 +169,22 @@ public class GLCLInteroperabilityDemo implements GLEventListener {
 
         commandQueue = clContext.getMaxFlopsDevice().createCommandQueue();
 
-        kernel = program.getCLKernel("sineWave");
-//        clBuffer = clContext.createFromGLBuffer(vb, glObjects[VERTICES], CLBuffer.Mem.WRITE_ONLY);
-        clBuffer = clContext.createBuffer(vb, CLBuffer.Mem.WRITE_ONLY);
-        kernel.setArg(0, clBuffer);
-        kernel.setArg(1, MESH_SIZE);
+        clBuffer = clContext.createFromGLBuffer(vb, glObjects[VERTICES], CLBuffer.Mem.WRITE_ONLY);
 
-        pushPerspectiveView(gl);
+        kernel = program.getCLKernel("sineWave")
+                        .setArg(0, clBuffer)
+                        .setArg(1, MESH_SIZE);
 
-        // start rendering thread
-        Animator animator = new Animator(drawable);
-        animator.start();
+//        computeHeightfield(drawable.getGL().getGL2());
+
+        System.out.println("cl initialised");
     }
 
+
     public void display(GLAutoDrawable drawable) {
 
         GL2 gl = drawable.getGL().getGL2();
 
-        compute(gl);
-
         gl.glClear(GL2.GL_COLOR_BUFFER_BIT | GL2.GL_DEPTH_BUFFER_BIT);
         gl.glLoadIdentity();
 
@@ -185,38 +195,31 @@ public class GLCLInteroperabilityDemo implements GLEventListener {
             gl.glBindBuffer(GL2.GL_ARRAY_BUFFER, glObjects[VERTICES]);
             gl.glVertexPointer(4, GL2.GL_FLOAT, 0, 0);
 
-            gl.glBindBuffer(GL2.GL_ELEMENT_ARRAY_BUFFER, glObjects[INDICES]);
-            gl.glDrawElements(GL2.GL_TRIANGLES, ib.capacity(), GL2.GL_UNSIGNED_INT, 0);
+//            gl.glBindBuffer(GL2.GL_ELEMENT_ARRAY_BUFFER, glObjects[INDICES]);
+//            gl.glDrawElements(GL2.GL_TRIANGLES, ib.capacity(), GL2.GL_UNSIGNED_INT, 0);
 
-//            gl.glDrawArrays(GL2.GL_POINTS, 0, vb.capacity()/4);
+            gl.glDrawArrays(GL2.GL_POINTS, 0, vb.capacity()/4);
 
-            gl.glBindBuffer(GL2.GL_ARRAY_BUFFER, 0);
+//            gl.glBindBuffer(GL2.GL_ARRAY_BUFFER, 0);
 
         gl.glDisableClientState(GL2.GL_VERTEX_ARRAY);
 
-
+        gl.glFinish();
+        computeHeightfield();
+        
     }
 
-    private void compute(GL2 gl) {
-
-        // not yet supported by OpenCL implementation
-//        commandQueue.putAcquireGLObject(clBuffer.ID);
-
-        //start computation
-        kernel.setArg(2, step += 0.1f);
-        commandQueue.putNDRangeKernel(kernel, 2, null, new long[] {MESH_SIZE, MESH_SIZE}, null);
-
-//        commandQueue.putReleaseGLObject(clBuffer.ID);
-
-        // workaround until full OpenCL-OpenGL interop. is supported
-        // copy from cl buffer to gl vbo
-        gl.glBindBuffer(GL2.GL_ARRAY_BUFFER, glObjects[VERTICES]);
-        gl.glVertexPointer(4, GL2.GL_FLOAT, 0, 0);
+    /*
+     * Computes a heightfield using a OpenCL kernel.
+     */
+    private void computeHeightfield() {
 
-        //blocking read of resultbuffer
-        commandQueue.putReadBuffer(clBuffer, vb, true);
+        kernel.setArg(2, step += 0.05f);
 
-        gl.glBufferData(GL2.GL_ARRAY_BUFFER, vb.capacity() * SIZEOF_FLOAT, vb, GL2.GL_DYNAMIC_DRAW);
+        commandQueue.putAcquireGLObject(clBuffer.ID)
+                    .put2DRangeKernel(kernel, 0, 0, MESH_SIZE, MESH_SIZE, 0, 0)
+                    .putReleaseGLObject(clBuffer.ID)
+                    .finish();
 
     }
 
@@ -257,11 +260,11 @@ public class GLCLInteroperabilityDemo implements GLEventListener {
     public void dispose(GLAutoDrawable drawable) {  }
 
     private void deinit() {
-        clContext.release();
+        clContext.release(); // only for demonstration purposes, JVM will cleanup on exit anyway
         System.exit(0);
     }
 
-    public static void main(String[] args) throws IOException {
+    public static void main(String[] args) {
         new GLCLInteroperabilityDemo();
     }
 
diff --git a/src/com/mbien/opencl/demos/joglinterop/JoglInterop.cl b/src/com/mbien/opencl/demos/joglinterop/JoglInterop.cl
index 5b42eca..0f0bcfc 100644
--- a/src/com/mbien/opencl/demos/joglinterop/JoglInterop.cl
+++ b/src/com/mbien/opencl/demos/joglinterop/JoglInterop.cl
@@ -2,7 +2,7 @@
 /**
 * animated 2D sine pattern.
 */
-__kernel void sineWave(__global float4 * vertex, int size, float time) {
+kernel void sineWave(global float4 * vertex, int size, float time) {
 
     unsigned int x = get_global_id(0);
     unsigned int y = get_global_id(1);