Skip to content

Commit

Permalink
Mandelbrot example uses now all available GPUs in parallel, paralleli…
Browse files Browse the repository at this point in the history 
…sm level is modifiable at runtime.

Small modifications due to api changes.
  • Loading branch information
@mbien
mbien committed Jan 29, 2010
1 parent da609ba commit 58e06a1
Show file tree
Hide file tree
Showing 4 changed files with 137 additions and 53 deletions.
6 changes: 3 additions & 3 deletions src/com/mbien/opencl/demos/fractal/Mandelbrot.cl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
180 changes: 132 additions & 48 deletions ...m/mbien/opencl/demos/fractal/Fractal.java → ...ncl/demos/fractal/MultiDeviceFractal.java
View file
Original file line number Diff line number Diff line change
Expand Up @@ -2,15 +2,20 @@

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;
Expand All @@ -34,22 +39,29 @@
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;
Expand All @@ -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;
Expand All @@ -68,13 +87,15 @@ public Fractal(int width, int height) {
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) {
Expand All @@ -98,29 +119,41 @@ public void init(GLAutoDrawable drawable) {

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);
}

}
Expand Down Expand Up @@ -164,51 +197,84 @@ private void initView(GL2 gl, int width, int height) {

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) {
Expand All @@ -219,15 +285,17 @@ public void reshape(GLAutoDrawable drawable, int x, int y, int width, int height
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());
}

private void initSceneInteraction() {

MouseAdapter adapter = new MouseAdapter() {
MouseAdapter mouseAdapter = new MouseAdapter() {

Point lastpos = new Point();

Expand Down Expand Up @@ -268,25 +336,41 @@ public void mouseWheelMoved(MouseWheelEvent e) {
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) {
}

public static void main(String args[]) {
SwingUtilities.invokeLater(new Runnable() {

public void run() {
new Fractal(500, 500);
new MultiDeviceFractal(512, 512);
}
});
}
Expand Down
2 changes: 1 addition & 1 deletion src/com/mbien/opencl/demos/hellojocl/HelloJOCL.java
View file
Original file line number Diff line number Diff line change
Expand Up @@ -45,7 +45,7 @@ public static void main(String[] args) throws IOException {

// 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.
Expand Down
2 changes: 1 addition & 1 deletion src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java
View file
Original file line number Diff line number Diff line change
Expand Up @@ -176,7 +176,7 @@ private void initCL() {
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();
Expand Down

0 comments on commit 58e06a1

Please sign in to comment.