diff options
Diffstat (limited to 'src/com/mbien/opencl')
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; |