diff options
Diffstat (limited to 'src/com/mbien/opencl/demos')
| -rw-r--r-- | src/com/mbien/opencl/demos/fractal/Mandelbrot.cl | 6 | ||||
| -rw-r--r-- | src/com/mbien/opencl/demos/fractal/MultiDeviceFractal.java (renamed from src/com/mbien/opencl/demos/fractal/Fractal.java) | 180 | ||||
| -rw-r--r-- | src/com/mbien/opencl/demos/hellojocl/HelloJOCL.java | 2 | ||||
| -rw-r--r-- | src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java | 2 |
4 files changed, 137 insertions, 53 deletions
diff --git a/src/com/mbien/opencl/demos/fractal/Mandelbrot.cl b/src/com/mbien/opencl/demos/fractal/Mandelbrot.cl index e6583f4..f4f39a1 100644 --- a/src/com/mbien/opencl/demos/fractal/Mandelbrot.cl +++ b/src/com/mbien/opencl/demos/fractal/Mandelbrot.cl @@ -7,14 +7,14 @@ kernel void mandelbrot( global uint *output, const int width, const int height, const float x0, const float y0, - const float x1, const float y1, + const float rangeX, const float rangeY, global uint *colorMap, const int colorMapSize, const int maxIterations) { unsigned int ix = get_global_id(0); unsigned int iy = get_global_id(1); - float r = x0 + ix * (x1-x0) / width; - float i = y0 + iy * (y1-y0) / height; + float r = x0 + ix * rangeX / width; + float i = y0 + iy * rangeY / height; float x = 0; float y = 0; diff --git a/src/com/mbien/opencl/demos/fractal/Fractal.java b/src/com/mbien/opencl/demos/fractal/MultiDeviceFractal.java index 06cde84..ee7e11d 100644 --- a/src/com/mbien/opencl/demos/fractal/Fractal.java +++ b/src/com/mbien/opencl/demos/fractal/MultiDeviceFractal.java @@ -2,15 +2,20 @@ package com.mbien.opencl.demos.fractal; import com.mbien.opencl.CLBuffer; import com.mbien.opencl.CLCommandQueue; +import com.mbien.opencl.CLDevice; import com.mbien.opencl.CLException; import com.mbien.opencl.CLGLBuffer; import com.mbien.opencl.CLGLContext; import com.mbien.opencl.CLKernel; import com.mbien.opencl.CLProgram; import com.mbien.opencl.CLProgram.CompilerOptions; +import com.sun.opengl.util.awt.TextRenderer; import java.awt.Color; import java.awt.Dimension; +import java.awt.Font; import java.awt.Point; +import java.awt.event.KeyAdapter; +import java.awt.event.KeyEvent; import java.awt.event.MouseAdapter; import java.awt.event.MouseEvent; import java.awt.event.MouseWheelEvent; @@ -34,22 +39,29 @@ import static com.sun.gluegen.runtime.BufferFactory.*; import static javax.media.opengl.GL2.*; import static com.mbien.opencl.CLMemory.Mem.*; import static com.mbien.opencl.CLDevice.Type.*; +import static java.lang.Math.*; /** - * Computes the Mandelbrot set with OpenCL and renders the result with OpenGL. + * Computes the Mandelbrot set with OpenCL using multiple GPUs and renders the result with OpenGL. * A shared PBO is used as storage for the fractal image. * http://en.wikipedia.org/wiki/Mandelbrot_set * @author Michael Bien */ -public class Fractal implements GLEventListener { +public class MultiDeviceFractal implements GLEventListener { + + // max number of used GPUs + private static final int MAX_PARRALLELISM_LEVEL = 8; + + // max per pixel iterations to compute the fractal + private static final int MAX_ITERATIONS = 500; private GLCanvas canvas; private CLGLContext clContext; - private CLCommandQueue commandQueue; - private CLKernel kernel; + private CLCommandQueue[] queues; + private CLKernel[] kernels; - private CLGLBuffer<IntBuffer> pboBuffer; + private CLGLBuffer<IntBuffer>[] pboBuffers; private int width = 0; private int height = 0; @@ -59,7 +71,14 @@ public class Fractal implements GLEventListener { private float maxX = 0.6f; private float maxY = 1.3f; - public Fractal(int width, int height) { + private int slices; + + private boolean drawSeperator = false; + private boolean initialized = false; + + private final TextRenderer textRenderer; + + public MultiDeviceFractal(int width, int height) { this.width = width; this.height = height; @@ -68,13 +87,15 @@ public class Fractal implements GLEventListener { canvas.addGLEventListener(this); initSceneInteraction(); - JFrame frame = new JFrame("JOCL Mandelbrot Set"); + JFrame frame = new JFrame("JOCL Multi GPU Mandelbrot Set"); frame.setDefaultCloseOperation(JFrame.EXIT_ON_CLOSE); canvas.setPreferredSize(new Dimension(width, height)); frame.add(canvas); frame.pack(); frame.setVisible(true); + + textRenderer = new TextRenderer(frame.getFont().deriveFont(Font.BOLD, 14), true, true, null, false); } public void init(GLAutoDrawable drawable) { @@ -98,29 +119,41 @@ public class Fractal implements GLEventListener { private void initCL(GLContext glCtx){ try { - clContext = CLGLContext.create(glCtx); + // create context managing all available GPUs + clContext = CLGLContext.create(glCtx, GPU); // load and build program CLProgram program = clContext.createProgram(getClass().getResourceAsStream("Mandelbrot.cl")) .build(CompilerOptions.FAST_RELAXED_MATH); - // setup colormap and create command queue on fastest GPU - CLBuffer<IntBuffer> colorMap = clContext.createIntBuffer(32*2, READ_ONLY); - initColorMap(colorMap.getBuffer(), 32, Color.BLUE, Color.GREEN, Color.RED); + CLDevice[] devices = clContext.getCLDevices(); + + slices = min(devices.length, MAX_PARRALLELISM_LEVEL); + + // create command queues for every GPU, setup colormap and init kernels + queues = new CLCommandQueue[slices]; + kernels = new CLKernel[slices]; + + for (int i = 0; i < slices; i++) { - commandQueue = clContext.getMaxFlopsDevice(GPU).createCommandQueue() - .putWriteBuffer(colorMap, true); // blocking upload + CLBuffer<IntBuffer> colorMap = clContext.createIntBuffer(32*2, READ_ONLY); + initColorMap(colorMap.getBuffer(), 32, Color.BLUE, Color.GREEN, Color.RED); - // init kernel with constants - kernel = program.getCLKernel("mandelbrot") - .setArg(7, colorMap) - .setArg(8, colorMap.getBuffer().capacity()) - .setArg(9, 200); // maxIterations + // create command queue and upload color map buffer on each used device + queues[i] = devices[i].createCommandQueue().putWriteBuffer(colorMap, true); // blocking upload + + // init kernel with constants + kernels[i] = program.createCLKernel("mandelbrot") + .setArg(7, colorMap) + .setArg(8, colorMap.getBuffer().capacity()) + .setArg(9, MAX_ITERATIONS); + + } } catch (IOException ex) { - Logger.getLogger(Fractal.class.getName()).log(Level.SEVERE, "can not find OpenCL program.", ex); + Logger.getLogger(getClass().getName()).log(Level.SEVERE, "can not find 'Mandelbrot.cl' in classpath.", ex); } catch (CLException ex) { - Logger.getLogger(Fractal.class.getName()).log(Level.SEVERE, "something went wrong, hopefully no one got hurt", ex); + Logger.getLogger(getClass().getName()).log(Level.SEVERE, "something went wrong, hopefully no one got hurt", ex); } } @@ -164,51 +197,84 @@ public class Fractal implements GLEventListener { gl.glMatrixMode(GL_PROJECTION); gl.glLoadIdentity(); - gl.glOrtho(0.0, 1.0, 0.0, 1.0, 0.0, 1.0); + gl.glOrtho(0.0, width, 0.0, height, 0.0, 1.0); } + @SuppressWarnings("unchecked") private void initPBO(GL gl) { - int[] pbo = new int[1]; - gl.glGenBuffers(1, pbo, 0); + pboBuffers = new CLGLBuffer[kernels.length]; - // setup empty PBO - gl.glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pbo[0]); - gl.glBufferData(GL_PIXEL_UNPACK_BUFFER, width * height * SIZEOF_INT, null, GL_STREAM_DRAW); - gl.glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0); + int[] pbo = new int[pboBuffers.length]; + gl.glGenBuffers(pboBuffers.length, pbo, 0); + + // setup one empty PBO per slice + for (int i = 0; i < slices; i++) { - pboBuffer = clContext.createFromGLBuffer(null, pbo[0], WRITE_ONLY); + gl.glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pbo[i]); + gl.glBufferData(GL_PIXEL_UNPACK_BUFFER, width*height * SIZEOF_INT / slices, null, GL_STREAM_DRAW); + gl.glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0); + pboBuffers[i] = clContext.createFromGLBuffer(null, pbo[i], WRITE_ONLY); + } + + initialized = true; } public void display(GLAutoDrawable drawable) { + if(!initialized) { + initPBO(drawable.getGL()); + } compute(); render(drawable.getGL().getGL2()); } + // OpenCL private void compute() { - kernel.putArg(pboBuffer) - .putArg(width).putArg(height) - .putArg(minX).putArg(minY) - .putArg(maxX).putArg(maxY) - .rewind(); + int sliceWidth = width / slices; + float rangeX = (maxX - minX) / slices; + float rangeY = (maxY - minY); + + for (int i = 0; i < slices; i++) { + + kernels[i].putArg(pboBuffers[i]) + .putArg(sliceWidth).putArg(height) + .putArg(minX + rangeX*i).putArg(minY) + .putArg( rangeX ).putArg(rangeY) + .rewind(); + + queues[i].putAcquireGLObject(pboBuffers[i].ID) + .put2DRangeKernel(kernels[i], 0, 0, sliceWidth, height, 0, 0) + .putReleaseGLObject(pboBuffers[i].ID); + } - commandQueue.putAcquireGLObject(pboBuffer.ID) - .put2DRangeKernel(kernel, 0, 0, width, height, 0, 0) - .putReleaseGLObject(pboBuffer.ID) - .putBarrier(); } + // OpenGL private void render(GL2 gl) { gl.glClear(GL_COLOR_BUFFER_BIT); - gl.glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pboBuffer.getGLObjectID()); - gl.glRasterPos2i(0, 0); - gl.glDrawPixels(width, height, GL_BGRA, GL_UNSIGNED_BYTE, 0); + //draw slices + int sliceWidth = width / slices; + + for (int i = 0; i < slices; i++) { + + int seperatorOffset = drawSeperator?i:0; + + gl.glBindBuffer(GL_PIXEL_UNPACK_BUFFER, pboBuffers[i].GLID); + gl.glRasterPos2i(sliceWidth*i + seperatorOffset, 0); + + gl.glDrawPixels(sliceWidth, height, GL_BGRA, GL_UNSIGNED_BYTE, 0); + + } gl.glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0); + //draw info text + textRenderer.beginRendering(width, height, false); + textRenderer.draw("#GPUs: "+slices, 10, 10); + textRenderer.endRendering(); } public void reshape(GLAutoDrawable drawable, int x, int y, int width, int height) { @@ -219,7 +285,9 @@ public class Fractal implements GLEventListener { this.width = width; this.height = height; - pboBuffer.release(); + for (CLGLBuffer<IntBuffer> buffer : pboBuffers) { + buffer.release(); + } initPBO(drawable.getGL()); initView(drawable.getGL().getGL2(), drawable.getWidth(), drawable.getHeight()); @@ -227,7 +295,7 @@ public class Fractal implements GLEventListener { private void initSceneInteraction() { - MouseAdapter adapter = new MouseAdapter() { + MouseAdapter mouseAdapter = new MouseAdapter() { Point lastpos = new Point(); @@ -268,15 +336,32 @@ public class Fractal implements GLEventListener { minX += deltaX+offsetX; minY += deltaY-offsetY; - maxX +=-deltaX+offsetX; - maxY +=-deltaY-offsetY; + maxX +=-deltaX+offsetX; + maxY +=-deltaY-offsetY; canvas.display(); } }; - canvas.addMouseMotionListener(adapter); - canvas.addMouseWheelListener(adapter); + KeyAdapter keyAdapter = new KeyAdapter() { + + @Override + public void keyPressed(KeyEvent e) { + if(e.getKeyCode() == KeyEvent.VK_SPACE) { + drawSeperator = !drawSeperator; + }else if(e.getKeyChar() > '0' && e.getKeyChar() < '9') { + int number = e.getKeyChar()-'0'; + slices = min(number, min(queues.length, MAX_PARRALLELISM_LEVEL)); + initialized = false; + } + canvas.display(); + } + + }; + + canvas.addMouseMotionListener(mouseAdapter); + canvas.addMouseWheelListener(mouseAdapter); + canvas.addKeyListener(keyAdapter); } public void dispose(GLAutoDrawable drawable) { @@ -284,9 +369,8 @@ public class Fractal implements GLEventListener { public static void main(String args[]) { SwingUtilities.invokeLater(new Runnable() { - public void run() { - new Fractal(500, 500); + new MultiDeviceFractal(512, 512); } }); } diff --git a/src/com/mbien/opencl/demos/hellojocl/HelloJOCL.java b/src/com/mbien/opencl/demos/hellojocl/HelloJOCL.java index 115278a..acbd84c 100644 --- a/src/com/mbien/opencl/demos/hellojocl/HelloJOCL.java +++ b/src/com/mbien/opencl/demos/hellojocl/HelloJOCL.java @@ -45,7 +45,7 @@ public class HelloJOCL { // get a reference to the kernel functon with the name 'VectorAdd' // and map the buffers to its input parameters. - CLKernel kernel = program.getCLKernel("VectorAdd"); + CLKernel kernel = program.createCLKernel("VectorAdd"); kernel.putArgs(clBufferA, clBufferB, clBufferC).putArg(elementCount); // create command queue on fastest device. diff --git a/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java b/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java index 7fceabb..9dfaf5a 100644 --- a/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java +++ b/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java @@ -176,7 +176,7 @@ public class GLCLInteroperabilityDemo implements GLEventListener { System.out.println("cl buffer type: " + clBuffer.getGLObjectType()); System.out.println("shared with gl buffer: " + clBuffer.getGLObjectID()); - kernel = program.getCLKernel("sineWave") + kernel = program.createCLKernel("sineWave") .putArg(clBuffer) .putArg(MESH_SIZE) .rewind(); |