summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
Diffstat
-rw-r--r--src/com/mbien/opencl/demos/fractal/Mandelbrot.cl6
-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.java2
-rw-r--r--src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java2
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();
generated by cgit v1.2.3 (git 2.39.1) at 2024-12-01 19:17:54 +0000