finished JOGL-JOCL interoperability demo.

3 files changed, 88 insertions, 61 deletions

diff --git a/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java b/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java
index 79e1d26..769774c 100644
--- a/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java
+++ b/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java
@@ -1,14 +1,12 @@
 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.CLDevice;
 import com.mbien.opencl.CLException;
 import com.mbien.opencl.CLKernel;
-import com.mbien.opencl.CLPlatform;
 import com.mbien.opencl.CLProgram;
 import com.sun.opengl.util.Animator;
-import com.sun.opengl.util.BufferUtil;
 import java.awt.event.WindowAdapter;
 import java.awt.event.WindowEvent;
 import java.io.IOException;
@@ -35,7 +33,7 @@ public class GLCLInteroperabilityDemo implements GLEventListener {
 
     private final GLUgl2 glu = new GLUgl2();
 
-    private final int GRID_SIZE = 100;
+    private final int MESH_SIZE = 256;
     
     private int width;
     private int height;
@@ -43,62 +41,47 @@ public class GLCLInteroperabilityDemo implements GLEventListener {
     private final FloatBuffer vb;
     private final IntBuffer ib;
 
-    private final int[] buffer = new int[2];
-    private final int INDICES  = 0;
-    private final int VERTICES = 1;
+    private final int[] glObjects = new int[2];
+    private final int VERTICES = 0;
+    private final int INDICES  = 1;
 
     private final UserSceneInteraction usi;
 
     private CLContext clContext;
     private CLKernel kernel;
     private CLCommandQueue commandQueue;
-    private final CLProgram program;
+    private CLBuffer<FloatBuffer> clBuffer;
+
+    private float step = 0;
 
     public GLCLInteroperabilityDemo() throws IOException {
 
         this.usi = new UserSceneInteraction();
 
-        vb = newFloatBuffer(GRID_SIZE * GRID_SIZE * 4);
-        ib = newIntBuffer((GRID_SIZE - 1) * (GRID_SIZE - 1) * 2 * 3);
+        vb = newFloatBuffer(MESH_SIZE * MESH_SIZE * 4);
+        ib = newIntBuffer((MESH_SIZE - 1) * (MESH_SIZE - 1) * 2 * 3);
 
         // build indices
         //    0---3
         //    | \ |
         //    1---2
-        for (int h = 0; h < GRID_SIZE - 1; h++) {
-            for (int w = 0; w < GRID_SIZE - 1; w++) {
+        for (int h = 0; h < MESH_SIZE - 1; h++) {
+            for (int w = 0; w < MESH_SIZE - 1; w++) {
 
                 // 0 - 3 - 2
-                ib.put(w * 6 + h * (GRID_SIZE - 1) * 6,      w + (h) * (GRID_SIZE)        );
-                ib.put(w * 6 + h * (GRID_SIZE - 1) * 6 + 1,  w + (h) * (GRID_SIZE) + 1    );
-                ib.put(w * 6 + h * (GRID_SIZE - 1) * 6 + 2,  w + (h + 1) * (GRID_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 * (GRID_SIZE - 1) * 6 + 3,  w + (h) * (GRID_SIZE)        );
-                ib.put(w * 6 + h * (GRID_SIZE - 1) * 6 + 4,  w + (h + 1) * (GRID_SIZE) + 1);
-                ib.put(w * 6 + h * (GRID_SIZE - 1) * 6 + 5,  w + (h + 1) * (GRID_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)    );
 
             }
         }
         ib.rewind();
 
-        // build grid
-        for (int w = 0; w < GRID_SIZE; w++) {
-            for (int h = GRID_SIZE; h > 0; h--) {
-                vb.put(w - GRID_SIZE / 2).put(h - GRID_SIZE / 2).put(0).put(1);
-            }
-        }
-        vb.rewind();
-
-        try {
-            clContext = CLContext.create();
-            program = clContext.createProgram(getClass().getResourceAsStream("JoglInterop.cl"));
-//            System.out.println(program.getSource());
-//            program.build();
-        } catch (IOException ex) {
-            throw new CLException("can not handle exception", ex);
-        }
-
         SwingUtilities.invokeLater(new Runnable() {
             public void run() {
                 initUI();
@@ -143,27 +126,41 @@ public class GLCLInteroperabilityDemo implements GLEventListener {
 
         GL2 gl = drawable.getGL().getGL2();
 
+        gl.setSwapInterval(1);
+
         gl.glPolygonMode(GL2.GL_FRONT_AND_BACK, GL2.GL_LINE);
         
-        gl.glGenBuffers(buffer.length, buffer, 0);
+        gl.glGenBuffers(glObjects.length, glObjects, 0);
         
-        gl.glBindBuffer(GL2.GL_ARRAY_BUFFER, buffer[INDICES]);
-        gl.glBufferData(GL2.GL_ARRAY_BUFFER, ib.capacity() * SIZEOF_FLOAT, ib, GL2.GL_STATIC_DRAW);
-        gl.glBindBuffer(GL2.GL_ARRAY_BUFFER, 0);
-
-        gl.glBindBuffer(GL2.GL_ELEMENT_ARRAY_BUFFER, buffer[VERTICES]);
-        gl.glBufferData(GL2.GL_ELEMENT_ARRAY_BUFFER, vb.capacity() * SIZEOF_FLOAT, vb, GL2.GL_DYNAMIC_DRAW);
+        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.glDisableClientState(GL2.GL_VERTEX_ARRAY);
+        
         // OpenCL
-//        commandQueue = clContext.getMaxFlopsDevice().createCommandQueue();
+        CLProgram program;
+        try {
+            clContext = CLContext.create();
+            program = clContext.createProgram(getClass().getResourceAsStream("JoglInterop.cl"));
+            program.build();
+            System.out.println(program.getBuildLog());
+            System.out.println(program.getBuildStatus());
+        } catch (IOException ex) {
+            throw new CLException("can not handle exception", ex);
+        }
 
-//            kernel = program.getCLKernel("sineWave");
-//        CLBuffer<FloatBuffer> clBuffer = clContext.createFromGLBuffer(vb, buffer[VERTICES], CLBuffer.Mem.WRITE_ONLY);
-//        kernel.setArg(0, clBuffer);
-//        kernel.setArg(1, GRID_SIZE);
-//        kernel.setArg(2, GRID_SIZE);
-        
+        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);
 
         pushPerspectiveView(gl);
 
@@ -175,6 +172,9 @@ public class GLCLInteroperabilityDemo implements GLEventListener {
     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();
 
@@ -182,15 +182,41 @@ public class GLCLInteroperabilityDemo implements GLEventListener {
 
         gl.glEnableClientState(GL2.GL_VERTEX_ARRAY);
 
-            gl.glBindBuffer(GL2.GL_ARRAY_BUFFER, buffer[VERTICES]);
+            gl.glBindBuffer(GL2.GL_ARRAY_BUFFER, glObjects[VERTICES]);
             gl.glVertexPointer(4, GL2.GL_FLOAT, 0, 0);
 
-            gl.glBindBuffer(GL2.GL_ELEMENT_ARRAY_BUFFER, buffer[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_POINTS, ib.capacity(), GL2.GL_UNSIGNED_INT, 0);
+
+            gl.glDrawArrays(GL2.GL_POINTS, 0, vb.capacity()/4);
+
+            gl.glBindBuffer(GL2.GL_ARRAY_BUFFER, 0);
 
         gl.glDisableClientState(GL2.GL_VERTEX_ARRAY);
 
-        gl.glBindBuffer(GL2.GL_ARRAY_BUFFER, 0);
+
+    }
+
+    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);
+
+        //blocking read of resultbuffer
+        commandQueue.putReadBuffer(clBuffer, vb, true);
+
+        gl.glBufferData(GL2.GL_ARRAY_BUFFER, vb.capacity() * SIZEOF_FLOAT, vb, GL2.GL_DYNAMIC_DRAW);
 
     }
 
diff --git a/src/com/mbien/opencl/demos/joglinterop/JoglInterop.cl b/src/com/mbien/opencl/demos/joglinterop/JoglInterop.cl
index 02356b2..cd92f14 100644
--- a/src/com/mbien/opencl/demos/joglinterop/JoglInterop.cl
+++ b/src/com/mbien/opencl/demos/joglinterop/JoglInterop.cl
@@ -1,14 +1,15 @@
-/*
- * Simple kernel to modify vertex positions in sine wave pattern
- */
-__kernel void sineWave(__global float4 * pos, unsigned int width, unsigned int height, float time) {
+
+/**
+* animated 2D sine pattern.
+*/
+__kernel void sineWave(__global float4 * vertex, int size, float time) {
 
     unsigned int x = get_global_id(0);
     unsigned int y = get_global_id(1);
 
     // calculate uv coordinates
-    float u = x / (float) width;
-    float v = y / (float) height;
+    float u = x / (float) size;
+    float v = y / (float) size;
 
     u = u*2.0f - 1.0f;
     v = v*2.0f - 1.0f;
@@ -18,6 +19,6 @@ __kernel void sineWave(__global float4 * pos, unsigned int width, unsigned int h
     float w = sin(u*freq + time) * cos(v*freq + time) * 0.5f;
 
     // write output vertex
-    pos[y*width+x] = (float4)(u, w, v, 1.0f);
+    vertex[y*size + x] = (float4)(u*10.0f, w*10.0f, v*10.0f, 1.0f);
 }
 
diff --git a/src/com/mbien/opencl/demos/joglinterop/UserSceneInteraction.java b/src/com/mbien/opencl/demos/joglinterop/UserSceneInteraction.java
index 50ffd6b..ab36f5e 100644
--- a/src/com/mbien/opencl/demos/joglinterop/UserSceneInteraction.java
+++ b/src/com/mbien/opencl/demos/joglinterop/UserSceneInteraction.java
@@ -15,7 +15,7 @@ import javax.media.opengl.GL2;
  */
 public class UserSceneInteraction {
 
-    private float z = -150;
+    private float z = -20;
     private float rotx;
     private float roty;