initial import of David Bucciarelli's Julia3d demo ported to java using OpenCL and JOGL2.

15 files changed, 1662 insertions, 14 deletions

diff --git a/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java b/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java
index 9dfaf5a..8febd4c 100644
--- a/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java
+++ b/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java
@@ -6,8 +6,6 @@ import com.mbien.opencl.CLGLContext;
 import com.mbien.opencl.CLKernel;
 import com.mbien.opencl.CLProgram;
 import com.sun.opengl.util.Animator;
-import java.awt.event.WindowAdapter;
-import java.awt.event.WindowEvent;
 import java.io.IOException;
 import javax.media.opengl.DebugGL2;
 import javax.media.opengl.GL2;
@@ -102,13 +100,7 @@ public class GLCLInteroperabilityDemo implements GLEventListener {
         usi.init(canvas);
 
         JFrame frame = new JFrame("JOGL-JOCL Interoperability Example");
-        frame.addWindowListener(new WindowAdapter() {
-            @Override
-            public void windowClosed(WindowEvent e) {
-                deinit();
-            }
-
-        });
+        frame.setDefaultCloseOperation(JFrame.EXIT_ON_CLOSE);
         frame.add(canvas);
         frame.setSize(width, height);
 
@@ -262,11 +254,6 @@ public class GLCLInteroperabilityDemo implements GLEventListener {
 
     public void dispose(GLAutoDrawable drawable) {  }
 
-    private void deinit() {
-        clContext.release(); // only for demonstration purposes, JVM will cleanup on exit anyway
-        System.exit(0);
-    }
-
     public static void main(String[] args) {
         new GLCLInteroperabilityDemo();
     }
diff --git a/src/com/mbien/opencl/demos/julia3d/Julia3d.java b/src/com/mbien/opencl/demos/julia3d/Julia3d.java
new file mode 100644
index 0000000..dbe2824
--- /dev/null
+++ b/src/com/mbien/opencl/demos/julia3d/Julia3d.java
@@ -0,0 +1,213 @@
+package com.mbien.opencl.demos.julia3d;
+
+import com.mbien.opencl.CLBuffer;
+import com.mbien.opencl.CLCommandQueue;
+import com.mbien.opencl.CLContext;
+import com.mbien.opencl.CLDevice;
+import com.mbien.opencl.CLKernel;
+import com.mbien.opencl.CLPlatform;
+import com.mbien.opencl.CLProgram;
+import com.mbien.opencl.demos.julia3d.structs.Camera;
+import com.mbien.opencl.demos.julia3d.structs.RenderingConfig;
+import com.mbien.opencl.demos.julia3d.structs.Vec;
+import java.io.IOException;
+import java.nio.Buffer;
+import java.nio.ByteBuffer;
+import java.nio.FloatBuffer;
+import javax.swing.SwingUtilities;
+
+import static com.mbien.opencl.CLMemory.Mem.*;
+import static com.mbien.opencl.CLDevice.Type.*;
+import static com.mbien.opencl.CLProgram.CompilerOptions.*;
+import static com.mbien.opencl.demos.julia3d.UserSceneController.*;
+
+/**
+ * This sample has been ported from David Buciarelli's juliaGPU v1.2 written in C.
+ * @author Michael Bien
+ */
+public class Julia3d {
+
+    private final CLContext context;
+    private CLBuffer<FloatBuffer> pixelBuffer;
+    private final CLBuffer<ByteBuffer> configBuffer;
+    private final CLCommandQueue commandQueue;
+    private final CLProgram program;
+    private final CLKernel julia;
+    private final CLKernel multiply;
+
+    private final int workGroupSize;
+    private final String kernelFileName = "rendering_kernel.cl";
+
+    final RenderingConfig config;
+
+    private Julia3d(RenderingConfig renderConfig) {
+        this.config = renderConfig;
+        updateCamera();
+
+        //setup
+        CLDevice gpu = CLPlatform.getDefault().getMaxFlopsDevice(GPU);
+        context = CLContext.create(gpu);
+
+        workGroupSize = 256;
+
+        //allocate buffers
+        configBuffer = context.createBuffer(config.getBuffer(), READ_ONLY, USE_BUFFER);
+        commandQueue = gpu.createCommandQueue();
+//        update(true);
+
+        try {
+            program = context.createProgram(Julia3d.class.getResourceAsStream(kernelFileName))
+                             .build(FAST_RELAXED_MATH);
+        } catch (IOException ex) {
+            throw new RuntimeException("unable to load program from source", ex);
+        }
+
+        julia = program.createCLKernel("JuliaGPU");
+        multiply = program.createCLKernel("multiply");
+        System.out.println(program.getBuildStatus(gpu));
+        System.out.println(program.getBuildLog());
+
+    }
+
+    void update(boolean reallocate) {
+
+        updateCamera();
+
+        int bufferSize = config.getWidth() * config.getHeight() * 3;
+        if(reallocate) {
+            if(pixelBuffer != null) {
+                pixelBuffer.release();
+            }
+
+            pixelBuffer = context.createFloatBuffer(bufferSize, READ_WRITE, USE_BUFFER);
+        }
+
+        commandQueue.putWriteBuffer(configBuffer, true);
+        
+        julia.putArg(pixelBuffer)
+             .putArg(configBuffer)
+             .rewind();
+
+        multiply.putArg(pixelBuffer)
+                .putArg(bufferSize)
+                .rewind();
+    }
+
+
+    void compute(boolean fastRendering) {
+
+        // calculate workgroup size
+        int globalThreads = config.getWidth() * config.getHeight();
+        if(globalThreads % workGroupSize != 0)
+            globalThreads = (globalThreads / workGroupSize + 1) * workGroupSize;
+
+        int localThreads = workGroupSize;
+        int superSamplingSize = config.getSuperSamplingSize();
+
+        if (!fastRendering && superSamplingSize > 1) {
+            
+            for (int y = 0; y < superSamplingSize; ++y) {
+                for (int x = 0; x < superSamplingSize; ++x) {
+                    
+                    float sampleX = (x + 0.5f) / superSamplingSize;
+                    float sampleY = (y + 0.5f) / superSamplingSize;
+                    
+                    if (x == 0 && y == 0) {
+                        // First pass
+                        julia.setArg(2, 0)
+                             .setArg(3, sampleX)
+                             .setArg(4, sampleY);
+
+                        commandQueue.put1DRangeKernel(julia, 0, globalThreads, localThreads);
+                            
+                    } else if (x == (superSamplingSize - 1) && y == (superSamplingSize - 1)) {
+                        // Last pass
+                        julia.setArg(2, 1)
+                             .setArg(3, sampleX)
+                             .setArg(4, sampleY);
+
+                        // normalize the values we accumulated
+                        multiply.setArg(2, 1.0f/(superSamplingSize*superSamplingSize));
+                        
+                        commandQueue.put1DRangeKernel(julia,    0, globalThreads,   localThreads)
+                                    .put1DRangeKernel(multiply, 0, globalThreads*3, localThreads);
+                    } else {
+                        julia.setArg(2, 1)
+                             .setArg(3, sampleX)
+                             .setArg(4, sampleY);
+                        
+                        commandQueue.put1DRangeKernel(julia, 0, globalThreads, localThreads);
+                        
+                    }
+                }
+            }
+            
+        }else{
+            
+            //fast rendering
+            julia.setArg(2, 0)
+                 .setArg(3, 0.0f)
+                 .setArg(4, 0.0f);
+
+            commandQueue.put1DRangeKernel(julia, 0, globalThreads, localThreads);
+        }
+        
+        commandQueue.putBarrier()
+                    .putReadBuffer(pixelBuffer, true);
+
+    }
+
+    private void updateCamera() {
+
+        Camera camera = config.getCamera();
+
+        Vec dir    = camera.getDir();
+        Vec target = camera.getTarget();
+        Vec camX   = camera.getX();
+        Vec camY   = camera.getY();
+        Vec orig   = camera.getOrig();
+
+        vsub(dir, target, orig);
+        vnorm(dir);
+
+        Vec up = Vec.create().setX(0).setY(1).setZ(0);
+        vxcross(camX, dir, up);
+        vnorm(camX);
+        vmul(camX, config.getWidth() * .5135f / config.getHeight(), camX);
+
+        vxcross(camY, camX, dir);
+        vnorm(camY);
+        vmul(camY, .5135f, camY);
+    }
+
+
+    public static void main(String[] args) {
+
+        RenderingConfig config = RenderingConfig.create()
+            .setWidth(640).setHeight(480)
+            .setEnableShadow(1)
+            .setSuperSamplingSize(2)
+            .setActvateFastRendering(1)
+            .setMaxIterations(9)
+            .setEpsilon(0.003f * 0.75f)
+            .setLight(new float[] {5, 10, 15})
+            .setMu(new float[] {-0.2f, 0.4f, -0.4f, -0.4f});
+
+        config.getCamera().getOrig()  .setX(1).setY(2).setZ(8);
+        config.getCamera().getTarget().setX(0).setY(0).setZ(0);
+
+        final Julia3d julia3d = new Julia3d(config);
+
+        SwingUtilities.invokeLater(new Runnable() {
+            public void run() {
+                new Renderer(julia3d);
+            }
+        });
+    }
+
+    Buffer getPixelBuffer() {
+        return pixelBuffer.getBuffer();
+    }
+
+
+}
diff --git a/src/com/mbien/opencl/demos/julia3d/Renderer.java b/src/com/mbien/opencl/demos/julia3d/Renderer.java
new file mode 100644
index 0000000..92ca2d7
--- /dev/null
+++ b/src/com/mbien/opencl/demos/julia3d/Renderer.java
@@ -0,0 +1,203 @@
+package com.mbien.opencl.demos.julia3d;
+
+import com.mbien.opencl.demos.julia3d.structs.RenderingConfig;
+import com.sun.opengl.util.BufferUtil;
+import com.sun.opengl.util.awt.TextRenderer;
+import java.awt.Dimension;
+import java.awt.Font;
+import java.nio.FloatBuffer;
+import java.util.Timer;
+import java.util.TimerTask;
+import javax.media.opengl.GL2;
+import javax.media.opengl.GLAutoDrawable;
+import javax.media.opengl.GLCapabilities;
+import javax.media.opengl.GLEventListener;
+import javax.media.opengl.GLProfile;
+import javax.media.opengl.awt.GLCanvas;
+import javax.swing.JFrame;
+
+import static javax.media.opengl.GL2.*;
+import static java.lang.String.*;
+
+/**
+ * JOGL renderer for displaying the julia set.
+ * @author Michael Bien
+ */
+public class Renderer implements GLEventListener {
+
+    public final static int MU_RECT_SIZE = 80;
+
+    private final Julia3d julia3d;
+    private final GLCanvas canvas;
+    private final RenderingConfig config;
+    private final FloatBuffer juliaSlice;
+    private final UserSceneController usi;
+    private final TextRenderer textRenderer;
+
+    private TimerTask task;
+    private final Timer timer;
+
+    public Renderer(Julia3d julia3d) {
+        this.julia3d = julia3d;
+        this.config = julia3d.config;
+
+        timer = new Timer();
+
+        juliaSlice = BufferUtil.newFloatBuffer(MU_RECT_SIZE * MU_RECT_SIZE * 4);
+
+        canvas = new GLCanvas(new GLCapabilities(GLProfile.get(GLProfile.GL2)));
+        canvas.addGLEventListener(this);
+
+        usi = new UserSceneController();
+        usi.init(this, canvas, config);
+
+        JFrame frame = new JFrame("Java OpenCL - Julia3D GPU");
+        frame.setDefaultCloseOperation(JFrame.EXIT_ON_CLOSE);
+        canvas.setPreferredSize(new Dimension(config.getWidth(), config.getHeight()));
+        frame.add(canvas);
+        frame.pack();
+
+        textRenderer = new TextRenderer(frame.getFont().deriveFont(Font.BOLD, 14), true, true, null, false);
+
+        frame.setVisible(true);
+    }
+
+    public void init(GLAutoDrawable drawable) {
+        drawable.getGL().getGL2().glMatrixMode(GL_PROJECTION);
+    }
+
+    void update() {
+        julia3d.update(false);
+        canvas.display();
+    }
+
+    public void display(GLAutoDrawable drawable) {
+
+        //compute
+        julia3d.compute(config.getActvateFastRendering() == 1);
+
+        GL2 gl = drawable.getGL().getGL2();
+        gl.glClear(GL_COLOR_BUFFER_BIT);
+
+        // draw julia set
+	gl.glRasterPos2i(0, 0);
+	gl.glDrawPixels(config.getWidth(), config.getHeight(), GL_RGB, GL_FLOAT, julia3d.getPixelBuffer());
+
+
+        // Draw Mu constant
+        int width = config.getWidth();
+        int height = config.getHeight();
+        float[] mu = config.getMu();
+
+	gl.glEnable(GL_BLEND);
+            gl.glBlendFunc(GL_SRC_ALPHA, GL_ONE_MINUS_SRC_ALPHA);
+            int baseMu1 = width - MU_RECT_SIZE - 2;
+            int baseMu2 = 1;
+            drawJuliaSlice(gl, baseMu1, baseMu2, mu[0], mu[1]);
+            int baseMu3 = width - MU_RECT_SIZE - 2;
+            int baseMu4 = MU_RECT_SIZE + 2;
+            drawJuliaSlice(gl, baseMu3, baseMu4, mu[2], mu[3]);
+	gl.glDisable(GL_BLEND);
+
+	gl.glColor3f(1, 1, 1);
+	int mu1 = (int) (baseMu1 + MU_RECT_SIZE * (mu[0] + 1.5f) / 3.f);
+	int mu2 = (int) (baseMu2 + MU_RECT_SIZE * (mu[1] + 1.5f) / 3.f);
+	gl.glBegin(GL_LINES);
+            gl.glVertex2i(mu1 - 4, mu2);
+            gl.glVertex2i(mu1 + 4, mu2);
+            gl.glVertex2i(mu1, mu2 - 4);
+            gl.glVertex2i(mu1, mu2 + 4);
+	gl.glEnd();
+
+	int mu3 = (int) (baseMu3 + MU_RECT_SIZE * (mu[2] + 1.5f) / 3.f);
+	int mu4 = (int) (baseMu4 + MU_RECT_SIZE * (mu[3] + 1.5f) / 3.f);
+	gl.glBegin(GL_LINES);
+            gl.glVertex2i(mu3 - 4, mu4);
+            gl.glVertex2i(mu3 + 4, mu4);
+            gl.glVertex2i(mu3, mu4 - 4);
+            gl.glVertex2i(mu3, mu4 + 4);
+	gl.glEnd();
+
+        // info text
+        textRenderer.beginRendering(width, height);
+        textRenderer.draw(format("Epsilon %.5f - Max. Iter. %d", config.getEpsilon(), config.getMaxIterations()), 8, 10);
+        textRenderer.draw(format("Mu = (%.3f, %.3f, %.3f, %.3f)", mu[0], mu[1], mu[2], mu[3]), 8, 25);
+        textRenderer.draw(format("Shadow %s - SuperSampling %dx%d - Fast rendering %s",
+			config.getEnableShadow() == 1 ? "on" : "off",
+                        config.getSuperSamplingSize(), config.getSuperSamplingSize(),
+			config.getActvateFastRendering() == 1 ? "on" : "off"), 8, 40);
+        textRenderer.endRendering();
+
+        // timer task scheduling, delay gpu intensive high quality rendering
+        if(task != null) {
+            task.cancel();
+        }
+        if(config.getActvateFastRendering() == 1) {
+            task = new TimerTask() {
+                @Override
+                public void run() {
+                    config.setActvateFastRendering(0);
+                    update();
+                    config.setActvateFastRendering(1);
+                }
+            };
+            timer.schedule(task, 2000);
+        }
+    }
+
+    private void drawJuliaSlice(GL2 gl, int origX, int origY, float cR, float cI) {
+
+        int index = 0;
+        float invSize = 3.0f / MU_RECT_SIZE;
+        for (int i = 0; i < MU_RECT_SIZE; ++i) {
+            for (int j = 0; j < MU_RECT_SIZE; ++j) {
+
+                float x = i * invSize - 1.5f;
+                float y = j * invSize - 1.5f;
+
+                int iter;
+                for (iter = 0; iter < 64; ++iter) {
+                    float x2 = x * x;
+                    float y2 = y * y;
+                    if (x2 + y2 > 4.0f) {
+                        break;
+                    }
+
+                    float newx = x2 - y2 + cR;
+                    float newy = 2.f * x * y + cI;
+                    x = newx;
+                    y = newy;
+                }
+                
+                juliaSlice.put(index++, iter / 64.0f);
+                juliaSlice.put(index++, 0.0f);
+                juliaSlice.put(index++, 0.0f);
+                juliaSlice.put(index++, 0.5f);
+            }
+        }
+
+	gl.glRasterPos2i(origX, origY);
+	gl.glDrawPixels(MU_RECT_SIZE, MU_RECT_SIZE, GL_RGBA, GL_FLOAT, juliaSlice);
+    }
+
+
+    public void reshape(GLAutoDrawable drawable, int x, int y, int newWidth, int newHeight) {
+
+        config.setWidth(newWidth);
+	config.setHeight(newHeight);
+
+        GL2 gl = drawable.getGL().getGL2();
+
+	gl.glViewport(0, 0, newWidth, newHeight);
+	gl.glLoadIdentity();
+	gl.glOrtho(-0.5f, newWidth - 0.5f, -0.5f, newHeight - 0.5f, -1.0f, 1.0f);
+
+        julia3d.update(true);
+
+    }
+
+    public void dispose(GLAutoDrawable drawable) {
+    }
+
+
+}
diff --git a/src/com/mbien/opencl/demos/julia3d/UserSceneController.java b/src/com/mbien/opencl/demos/julia3d/UserSceneController.java
new file mode 100644
index 0000000..849de3a
--- /dev/null
+++ b/src/com/mbien/opencl/demos/julia3d/UserSceneController.java
@@ -0,0 +1,249 @@
+package com.mbien.opencl.demos.julia3d;
+
+import com.mbien.opencl.demos.julia3d.structs.RenderingConfig;
+import com.mbien.opencl.demos.julia3d.structs.Vec;
+import java.awt.Component;
+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;
+
+import static java.lang.Math.*;
+import static com.mbien.opencl.demos.julia3d.Renderer.*;
+
+/**
+ * Utility class for interacting with a scene. Supports rotation and zoom around origin.
+ * @author Michael Bien
+ */
+public class UserSceneController {
+
+    private Point dragstart;
+    private RenderingConfig model;
+    private Renderer view;
+
+    private enum MOUSE_MODE { DRAG_ROTATE, DRAG_ZOOM }
+    private MOUSE_MODE dragmode = MOUSE_MODE.DRAG_ROTATE;
+
+
+    public void init(Renderer view, Component component, RenderingConfig model) {
+        initMouseListeners(component);
+        this.view = view;
+        this.model = model;
+    }
+
+    private void initMouseListeners(Component component) {
+
+        MouseAdapter mouseAdapter = new MouseAdapter() {
+            @Override
+            public void mouseDragged(MouseEvent e) {
+
+                int x = e.getX();
+                int y = e.getY();
+
+                switch (dragmode) {
+                    case DRAG_ROTATE:
+                        if (dragstart != null) {
+                            int height = model.getHeight();
+                            int width = model.getWidth();
+
+                            int ry = height - y - 1;
+                            int baseMu1 = width - MU_RECT_SIZE - 2;
+                            int baseMu2 = 1;
+                            int baseMu3 = width - MU_RECT_SIZE - 2;
+                            int baseMu4 = MU_RECT_SIZE + 2;
+
+                            if ((x >= baseMu1 && x <= baseMu1 + MU_RECT_SIZE) && (ry >= baseMu2 && ry <= baseMu2 + MU_RECT_SIZE)) {
+                                float[] mu = model.getMu();
+                                mu[0] = 3.f * ( x - baseMu1) / (float)MU_RECT_SIZE - 1.5f;
+                                mu[1] = 3.f * (ry - baseMu2) / (float)MU_RECT_SIZE - 1.5f;
+                                model.setMu(mu);
+                            } else if ((x >= baseMu3 && x <= baseMu3 + MU_RECT_SIZE) && (ry >= baseMu4 && ry <= baseMu4 + MU_RECT_SIZE)) {
+                                float[] mu = model.getMu();
+                                mu[2] = 3.f * ( x - baseMu3) / (float)MU_RECT_SIZE - 1.5f;
+                                mu[3] = 3.f * (ry - baseMu4) / (float)MU_RECT_SIZE - 1.5f;
+                                model.setMu(mu);
+                            } else {
+                                rotateCameraYbyOrig(0.01f * (x - dragstart.getX()));
+                                rotateCameraXbyOrig(0.01f * (y - dragstart.getY()));
+                            }
+                        }
+                        dragstart = e.getPoint();
+                        view.update();
+                        break;
+                    case DRAG_ZOOM:
+                        if (dragstart != null) {
+                            float zoom = (float) ((y - dragstart.getY()) / 10.0f);
+                            zoom(zoom);
+                        }
+                        dragstart = e.getPoint();
+                        view.update();
+                        break;
+                }
+
+            }
+
+            @Override
+            public void mousePressed(MouseEvent e) {
+                switch (e.getButton()) {
+                    case (MouseEvent.BUTTON1):
+                        dragmode = MOUSE_MODE.DRAG_ROTATE;
+                        break;
+                    case (MouseEvent.BUTTON2):
+                        dragmode = MOUSE_MODE.DRAG_ZOOM;
+                        break;
+                    case (MouseEvent.BUTTON3):
+                        dragmode = MOUSE_MODE.DRAG_ZOOM;
+                        break;
+                }
+            }
+
+            @Override
+            public void mouseReleased(MouseEvent e) {
+                switch (e.getButton()) {
+                    case (MouseEvent.BUTTON1):
+                        dragmode = MOUSE_MODE.DRAG_ZOOM;
+                        break;
+                    case (MouseEvent.BUTTON2):
+                        dragmode = MOUSE_MODE.DRAG_ROTATE;
+                        break;
+                    case (MouseEvent.BUTTON3):
+                        dragmode = MOUSE_MODE.DRAG_ROTATE;
+                        break;
+                }
+
+                dragstart = null;
+            }
+
+            @Override
+            public void mouseWheelMoved(MouseWheelEvent e) {
+                float zoom = e.getWheelRotation() * 0.1f;
+                zoom(zoom);
+                view.update();
+            }
+
+        };
+
+        KeyAdapter keyAdapter = new KeyAdapter() {
+
+            @Override
+            public void keyPressed(KeyEvent e) {
+
+                switch (e.getKeyChar()) {
+                    case 'l':
+                        model.setEnableShadow(model.getEnableShadow()==0 ? 1 : 0);
+                        break;
+                    case '1':
+                        model.setEpsilon(model.getEpsilon() * 0.75f);
+                        break;
+                    case '2':
+                        model.setEpsilon(model.getEpsilon() * 1.f / 0.75f);
+                        break;
+                    case '3':
+                        model.setMaxIterations(max(1, model.getMaxIterations() -1));
+                        break;
+                    case '4':
+                        model.setMaxIterations(min(12, model.getMaxIterations()+1));
+                        break;
+                    case '5':
+                        model.setSuperSamplingSize(max(1, model.getSuperSamplingSize() -1));
+                        break;
+                    case '6':
+                        model.setSuperSamplingSize(min(5, model.getSuperSamplingSize() +1));
+                        break;
+                    default:
+                        break;
+                }
+                view.update();
+
+            }
+
+        };
+
+        component.addKeyListener(keyAdapter);
+
+        component.addMouseListener(mouseAdapter);
+        component.addMouseMotionListener(mouseAdapter);
+        component.addMouseWheelListener(mouseAdapter);
+
+    }
+    private void zoom(float zoom) {
+        Vec orig = model.getCamera().getOrig();
+        orig.setX(orig.getX()+zoom)
+            .setY(orig.getY()+zoom)
+            .setZ(orig.getZ()+zoom);
+    }
+
+    private void rotateLightX(float k) {
+        float[] light = model.getLight();
+        float y = light[1];
+        float z = light[2];
+        light[1] = (float) ( y * cos(k) + z * sin(k));
+        light[2] = (float) (-y * sin(k) + z * cos(k));
+        model.setLight(light);
+    }
+
+    private void rotateLightY(float k) {
+        float[] light = model.getLight();
+        float x = light[0];
+        float z = light[2];
+        light[0] = (float) (x * cos(k) - z * sin(k));
+        light[2] = (float) (x * sin(k) + z * cos(k));
+        model.setLight(light);
+    }
+
+    private void rotateCameraXbyOrig(double k) {
+        Vec orig = model.getCamera().getOrig();
+        float y = orig.getY();
+        float z = orig.getZ();
+        orig.setY((float) ( y * cos(k) + z * sin(k)));
+        orig.setZ((float) (-y * sin(k) + z * cos(k)));
+    }
+
+    private void rotateCameraYbyOrig(double k) {
+        Vec orig = model.getCamera().getOrig();
+        float x = orig.getX();
+        float z = orig.getZ();
+        orig.setX((float) (x * cos(k) - z * sin(k)));
+        orig.setZ((float) (x * sin(k) + z * cos(k)));
+    }
+
+
+    public final static void vadd(Vec v, Vec a, Vec b) {
+        v.setX(a.getX() + b.getX());
+        v.setY(a.getY() + b.getY());
+        v.setZ(a.getZ() + b.getZ());
+    }
+
+    public final static void vsub(Vec v, Vec a, Vec b) {
+        v.setX(a.getX() - b.getX());
+        v.setY(a.getY() - b.getY());
+        v.setZ(a.getZ() - b.getZ());
+    }
+
+    public final static void vmul(Vec v, float s, Vec b) {
+        v.setX(s * b.getX());
+        v.setY(s * b.getY());
+        v.setZ(s * b.getZ());
+    }
+
+    public final static float vdot(Vec a, Vec b) {
+        return a.getX() * b.getX()
+             + a.getY() * b.getY()
+             + a.getZ() * b.getZ();
+    }
+
+    public final static void vnorm(Vec v) {
+        float s = (float) (1.0f / sqrt(vdot(v, v)));
+        vmul(v, s, v);
+    }
+
+    public final static void vxcross(Vec v, Vec a, Vec b) {
+        v.setX(a.getY() * b.getZ() - a.getZ() * b.getY());
+        v.setY(a.getZ() * b.getX() - a.getX() * b.getZ());
+        v.setZ(a.getX() * b.getY() - a.getY() * b.getX());
+    }
+
+
+}
\ No newline at end of file
diff --git a/src/com/mbien/opencl/demos/julia3d/config.h b/src/com/mbien/opencl/demos/julia3d/config.h
new file mode 100644
index 0000000..72df3ff
--- /dev/null
+++ b/src/com/mbien/opencl/demos/julia3d/config.h
@@ -0,0 +1,24 @@
+
+typedef struct {
+	float x, y, z; // position, also color (r,g,b)
+} Vec;
+
+typedef struct {
+	/* User defined values */
+	Vec orig, target;
+	/* Calculated values */
+	Vec dir, x, y;
+} Camera;
+
+typedef struct {
+	unsigned int width, height;
+	int superSamplingSize;
+	int actvateFastRendering;
+	int enableShadow;
+
+	unsigned int maxIterations;
+	float epsilon;
+	float mu[4];
+	float light[3];
+	Camera camera;
+} RenderingConfig;
diff --git a/src/com/mbien/opencl/demos/julia3d/rendering_kernel.cl b/src/com/mbien/opencl/demos/julia3d/rendering_kernel.cl
new file mode 100644
index 0000000..9c25c1b
--- /dev/null
+++ b/src/com/mbien/opencl/demos/julia3d/rendering_kernel.cl
@@ -0,0 +1,382 @@
+/*
+Copyright (c) 2009 David Bucciarelli ([email protected])
+
+Permission is hereby granted, free of charge, to any person obtaining
+a copy of this software and associated documentation files (the
+"Software"), to deal in the Software without restriction, including
+without limitation the rights to use, copy, modify, merge, publish,
+distribute, sublicense, and/or sell copies of the Software, and to
+permit persons to whom the Software is furnished to do so, subject to
+the following conditions:
+
+The above copyright notice and this permission notice shall be included
+in all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE.
+*/
+
+#define GPU_KERNEL
+
+
+#define BOUNDING_RADIUS_2 4.f
+#define ESCAPE_THRESHOLD 1e1f
+#define DELTA 1e-4f
+
+typedef struct {
+    float x, y, z; // position, also color (r,g,b)
+} Vec;
+
+typedef struct {
+    Vec orig, target;
+    Vec dir, x, y;
+} Camera;
+
+typedef struct {
+    unsigned int width, height;
+    int superSamplingSize;
+    int actvateFastRendering;
+    int enableShadow;
+
+    unsigned int maxIterations;
+    float epsilon;
+    float mu[4];
+    float light[3];
+    Camera camera;
+} RenderingConfig;
+
+
+static float4 QuatMult(const float4 q1, const float4 q2) {
+    float4 r;
+
+    // a1a2 - b1b2 - c1c2 - d1d2
+    r.x = q1.x * q2.x - q1.y * q2.y - q1.z * q2.z - q1.w * q2.w;
+    // a1b2 + b1a2 + c1d2 - d1c2
+    r.y = q1.x * q2.y + q1.y * q2.x + q1.z * q2.w - q1.w * q2.z;
+    // a1c2 - b1d2 + c1a2 + d1b2
+    r.z = q1.x * q2.z - q1.y * q2.w + q1.z * q2.x + q1.w * q2.y;
+    // a1d2 + b1c2 - c1b2 + d1a2
+    r.w = q1.x * q2.w + q1.y * q2.z - q1.z * q2.y + q1.w * q2.x;
+
+    return r;
+}
+
+static float4 QuatSqr(const float4 q) {
+    float4 r;
+
+    r.x = q.x * q.x - q.y * q.y - q.z * q.z - q.w * q.w;
+    r.y = 2.f * q.x * q.y;
+    r.z = 2.f * q.x * q.z;
+    r.w = 2.f * q.x * q.w;
+
+    return r;
+}
+
+static void IterateIntersect(float4 *q, float4 *qp, const float4 c, const uint maxIterations) {
+    float4 q0 = *q;
+    float4 qp0 = *qp;
+
+    for (uint i = 0; i < maxIterations; ++i) {
+        qp0 = 2.f * QuatMult(q0, qp0);
+        q0 = QuatSqr(q0) + c;
+
+        if (dot(q0, q0) > ESCAPE_THRESHOLD)
+            break;
+    }
+
+    *q = q0;
+    *qp = qp0;
+}
+
+static float IntersectJulia(const float4 eyeRayOrig, const float4 eyeRayDir,
+        const float4 c, const uint maxIterations, const float epsilon,
+        float4 *hitPoint, uint *steps) {
+    float dist;
+    float4 r0 = eyeRayOrig;
+
+    uint s = 0;
+    do {
+        float4 z = r0;
+        float4 zp = (float4) (1.f, 0.f, 0.f, 0.f);
+
+        IterateIntersect(&z, &zp, c, maxIterations);
+
+        const float normZP = length(zp);
+
+        // We are inside
+        if (normZP == 0.f)
+            break;
+
+        const float normZ = length(z);
+        dist = 0.5f * normZ * log(normZ) / normZP;
+
+        r0 += eyeRayDir * dist;
+        s++;
+    } while ((dist > epsilon) && (dot(r0, r0) < BOUNDING_RADIUS_2));
+
+    *hitPoint = r0;
+    *steps = s;
+    return dist;
+}
+
+#define WORLD_RADIUS 1000.f
+#define WORLD_CENTER ((float4)(0.f, -WORLD_RADIUS - 2.f, 0.f, 0.f))
+
+float IntersectFloorSphere(const float4 eyeRayOrig, const float4 eyeRayDir) {
+    const float4 op = WORLD_CENTER - eyeRayOrig;
+    const float b = dot(op, eyeRayDir);
+    float det = b * b - dot(op, op) + WORLD_RADIUS * WORLD_RADIUS;
+
+    if (det < 0.f)
+        return -1.f;
+    else
+        det = sqrt(det);
+
+    float t = b - det;
+    if (t > 0.f)
+        return t;
+    else {
+        // We are inside, avoid the hit
+        return -1.f;
+    }
+}
+
+float IntersectBoundingSphere(const float4 eyeRayOrig, const float4 eyeRayDir) {
+    const float4 op = -eyeRayOrig;
+    const float b = dot(op, eyeRayDir);
+    float det = b * b - dot(op, op) + BOUNDING_RADIUS_2;
+
+    if (det < 0.f)
+        return -1.f;
+    else
+        det = sqrt(det);
+
+    float t = b - det;
+    if (t > 0.f)
+        return t;
+    else {
+        t = b + det;
+
+        if (t > 0.f) {
+            // We are inside, start from the ray origin
+            return 0.0f;
+        } else
+            return -1.f;
+    }
+}
+
+static float4 NormEstimate(const float4 p, const float4 c,
+        const float delta, const uint maxIterations) {
+    float4 N;
+    float4 qP = p;
+    float gradX, gradY, gradZ;
+
+    float4 gx1 = qP - (float4) (DELTA, 0.f, 0.f, 0.f);
+    float4 gx2 = qP + (float4) (DELTA, 0.f, 0.f, 0.f);
+    float4 gy1 = qP - (float4) (0.f, DELTA, 0.f, 0.f);
+    float4 gy2 = qP + (float4) (0.f, DELTA, 0.f, 0.f);
+    float4 gz1 = qP - (float4) (0.f, 0.f, DELTA, 0.f);
+    float4 gz2 = qP + (float4) (0.f, 0.f, DELTA, 0.f);
+
+    for (uint i = 0; i < maxIterations; ++i) {
+        gx1 = QuatSqr(gx1) + c;
+        gx2 = QuatSqr(gx2) + c;
+        gy1 = QuatSqr(gy1) + c;
+        gy2 = QuatSqr(gy2) + c;
+        gz1 = QuatSqr(gz1) + c;
+        gz2 = QuatSqr(gz2) + c;
+    }
+
+    gradX = length(gx2) - length(gx1);
+    gradY = length(gy2) - length(gy1);
+    gradZ = length(gz2) - length(gz1);
+
+    N = normalize((float4) (gradX, gradY, gradZ, 0.f));
+
+    return N;
+}
+
+static float4 Phong(const float4 light, const float4 eye, const float4 pt, const float4 N, const float4 diffuse) {
+
+    const float4 ambient = (float4) (0.05f, 0.05f, 0.05f, 0.f);
+    float4 L = normalize(light - pt);
+    float NdotL = dot(N, L);
+    if (NdotL < 0.f)
+        return diffuse * ambient;
+
+    const float specularExponent = 30.f;
+    const float specularity = 0.65f;
+
+    float4 E = normalize(eye - pt);
+    float4 H = (L + E) * (float) 0.5f;
+
+    return diffuse * NdotL +
+           specularity * pow(dot(N, H), specularExponent) +
+           diffuse * ambient;
+}
+
+kernel void JuliaGPU(   global float *pixels,
+                        const global RenderingConfig *config,
+                        int enableAccumulation,
+                        float sampleX,
+                        float sampleY ) {
+
+    const int gid = get_global_id(0);
+    unsigned width = config->width;
+    unsigned height = config->height;
+
+    const unsigned int x = gid % width;
+    const int y = gid / width;
+
+    // Check if we have to do something
+    if (y >= height)
+        return;
+
+    const float epsilon = config->actvateFastRendering ? (config->epsilon * (1.f / 0.75f)) : config->epsilon;
+    const uint maxIterations = max(1u, config->actvateFastRendering ? (config->maxIterations - 1) : config->maxIterations);
+
+    const float4 mu = (float4)(config->mu[0], config->mu[1], config->mu[2], config->mu[3]);
+    const float4 light = (float4) (config->light[0], config->light[1], config->light[2], 0.f);
+    const global Camera *camera = &config->camera;
+
+    //--------------------------------------------------------------------------
+    // Calculate eye ray
+    //--------------------------------------------------------------------------
+
+    const float invWidth = 1.f / width;
+    const float invHeight = 1.f / height;
+    const float kcx = (x + sampleX) * invWidth - .5f;
+    const float4 kcx4 = (float4) kcx;
+    const float kcy = (y + sampleY) * invHeight - .5f;
+    const float4 kcy4 = (float4) kcy;
+
+    const float4 cameraX = (float4) (camera->x.x, camera->x.y, camera->x.z, 0.f);
+    const float4 cameraY = (float4) (camera->y.x, camera->y.y, camera->y.z, 0.f);
+    const float4 cameraDir = (float4) (camera->dir.x, camera->dir.y, camera->dir.z, 0.f);
+    const float4 cameraOrig = (float4) (camera->orig.x, camera->orig.y, camera->orig.z, 0.f);
+
+    const float4 eyeRayDir = normalize(cameraX * kcx4 + cameraY * kcy4 + cameraDir);
+    const float4 eyeRayOrig = eyeRayDir * (float4) 0.1f + cameraOrig;
+
+    //--------------------------------------------------------------------------
+    // Check if we hit the bounding sphere
+    //--------------------------------------------------------------------------
+
+    float distSet = IntersectBoundingSphere(eyeRayOrig, eyeRayDir);
+    float4 hitPoint;
+    if (distSet >= 0.f) {
+        //--------------------------------------------------------------------------
+        // Find the intersection with the set
+        //--------------------------------------------------------------------------
+
+        uint steps;
+        float4 rayOrig = eyeRayOrig + eyeRayDir * (float4) distSet;
+        distSet = IntersectJulia(rayOrig, eyeRayDir, mu, maxIterations,
+                epsilon, &hitPoint, &steps);
+        if (distSet > epsilon)
+            distSet = -1.f;
+    }
+
+    //--------------------------------------------------------------------------
+    // Check if we hit the floor
+    //--------------------------------------------------------------------------
+
+    float distFloor = IntersectFloorSphere(eyeRayOrig, eyeRayDir);
+
+    //--------------------------------------------------------------------------
+    // Select the hit point
+    //--------------------------------------------------------------------------
+
+    int doShade = 0;
+    int useAO = 1;
+    float4 diffuse, n, color;
+    if ((distSet < 0.f) && (distFloor < 0.f)) {
+        // Sky hit
+        color = (float4) (0.f, 0.1f, 0.3f, 0.f);
+    } else if ((distSet >= 0.f) && ((distFloor < 0.f) || (distSet <= distFloor))) {
+        // Set hit
+        diffuse = (float4) (1.f, 0.35f, 0.15f, 0.f);
+        n = NormEstimate(hitPoint, mu, distSet, maxIterations);
+        doShade = 1;
+    } else if ((distFloor >= 0.f) && ((distSet < 0.f) || (distFloor <= distSet))) {
+        // Floor hit
+        hitPoint = eyeRayOrig + eyeRayDir * (float4) distFloor;
+        n = hitPoint - WORLD_CENTER;
+        n = normalize(n);
+        // The most important feature in a ray tracer: a checker texture !
+        const int ix = (hitPoint.x > 0.f) ? hitPoint.x : (1.f - hitPoint.x);
+        const int iz = (hitPoint.z > 0.f) ? hitPoint.z : (1.f - hitPoint.z);
+        if ((ix + iz) % 2)
+            diffuse = (float4) (0.75f, 0.75f, 0.75f, 0.f);
+        else
+            diffuse = (float4) (0.75f, 0.f, 0.f, 0.f);
+        doShade = 1;
+        useAO = 0;
+    }
+
+    //--------------------------------------------------------------------------
+    // Select the shadow pass
+    //--------------------------------------------------------------------------
+
+    if (doShade) {
+        float shadowFactor = 1.f;
+        if (config->enableShadow) {
+            float4 L = normalize(light - hitPoint);
+            float4 rO = hitPoint + n * 1e-2f;
+            float4 shadowHitPoint;
+
+            // Check bounding sphere
+            float shadowDistSet = IntersectBoundingSphere(rO, L);
+            if (shadowDistSet >= 0.f) {
+                uint steps;
+
+                rO = rO + L * (float4) shadowDistSet;
+                shadowDistSet = IntersectJulia(rO, L, mu, maxIterations, epsilon,
+                        &shadowHitPoint, &steps);
+                if (shadowDistSet < epsilon) {
+                    if (useAO) {
+                        // Use steps count to simulate ambient occlusion
+                        shadowFactor = 0.6f - min(steps / 255.f, 0.5f);
+                    } else
+                        shadowFactor = 0.6f;
+                }
+            } else
+                shadowDistSet = -1.f;
+        }
+
+        //--------------------------------------------------------------------------
+        // Direct lighting of hit point
+        //--------------------------------------------------------------------------
+
+        color = Phong(light, eyeRayOrig, hitPoint, n, diffuse) * shadowFactor;
+    }
+
+    //--------------------------------------------------------------------------
+    // Write pixel
+    //--------------------------------------------------------------------------
+
+    int offset = 3 * (x + y * width);
+    color = clamp(color, (float4) (0.f, 0.f, 0.f, 0.f), (float4) (1.f, 1.f, 1.f, 0.f));
+    if (enableAccumulation) {
+        pixels[offset++] += color.s0;
+        pixels[offset++] += color.s1;
+        pixels[offset  ] += color.s2;
+    } else {
+        pixels[offset++] = color.s0;
+        pixels[offset++] = color.s1;
+        pixels[offset  ] = color.s2;
+    }
+}
+
+kernel void multiply(global float *array, const int numElements, const float s) {
+    const int gid = get_global_id(0);
+    if (gid >= numElements)  {
+        return;
+    }
+    array[gid] *= s;
+}
diff --git a/src/com/mbien/opencl/demos/julia3d/structs/Camera.java b/src/com/mbien/opencl/demos/julia3d/structs/Camera.java
new file mode 100644
index 0000000..8b820e1
--- /dev/null
+++ b/src/com/mbien/opencl/demos/julia3d/structs/Camera.java
@@ -0,0 +1,52 @@
+/* !---- DO NOT EDIT: This file autogenerated by com/sun/gluegen/JavaEmitter.java on Tue Feb 09 18:20:26 CET 2010 ----! */
+
+
+package com.mbien.opencl.demos.julia3d.structs;
+
+import java.nio.*;
+
+import com.sun.gluegen.runtime.*;
+
+
+public abstract class Camera {
+
+  StructAccessor accessor;
+
+  public static int size() {
+    if (CPU.is32Bit()) {
+      return Camera32.size();
+    } else {
+      return Camera64.size();
+    }
+  }
+
+  public static Camera create() {
+    return create(BufferFactory.newDirectByteBuffer(size()));
+  }
+
+  public static Camera create(java.nio.ByteBuffer buf) {
+    if (CPU.is32Bit()) {
+      return new Camera32(buf);
+    } else {
+      return new Camera64(buf);
+    }
+  }
+
+  Camera(java.nio.ByteBuffer buf) {
+    accessor = new StructAccessor(buf);
+  }
+
+  public java.nio.ByteBuffer getBuffer() {
+    return accessor.getBuffer();
+  }
+
+  public abstract Vec getOrig();
+
+  public abstract Vec getTarget();
+
+  public abstract Vec getDir();
+
+  public abstract Vec getX();
+
+  public abstract Vec getY();
+}
diff --git a/src/com/mbien/opencl/demos/julia3d/structs/Camera32.java b/src/com/mbien/opencl/demos/julia3d/structs/Camera32.java
new file mode 100644
index 0000000..2d09540
--- /dev/null
+++ b/src/com/mbien/opencl/demos/julia3d/structs/Camera32.java
@@ -0,0 +1,41 @@
+/* !---- DO NOT EDIT: This file autogenerated by com/sun/gluegen/JavaEmitter.java on Tue Feb 09 18:20:26 CET 2010 ----! */
+
+
+package com.mbien.opencl.demos.julia3d.structs;
+
+import java.nio.*;
+
+import com.sun.gluegen.runtime.*;
+
+
+class Camera32 extends Camera {
+
+  public static int size() {
+    return 76;
+  }
+
+  Camera32(java.nio.ByteBuffer buf) {
+    super(buf);
+  }
+
+
+  public Vec getOrig() {
+    return Vec.create(accessor.slice(0, 12));
+  }
+
+  public Vec getTarget() {
+    return Vec.create(accessor.slice(16, 12));
+  }
+
+  public Vec getDir() {
+    return Vec.create(accessor.slice(32, 12));
+  }
+
+  public Vec getX() {
+    return Vec.create(accessor.slice(48, 12));
+  }
+
+  public Vec getY() {
+    return Vec.create(accessor.slice(64, 12));
+  }
+}
diff --git a/src/com/mbien/opencl/demos/julia3d/structs/Camera64.java b/src/com/mbien/opencl/demos/julia3d/structs/Camera64.java
new file mode 100644
index 0000000..fe5a729
--- /dev/null
+++ b/src/com/mbien/opencl/demos/julia3d/structs/Camera64.java
@@ -0,0 +1,52 @@
+/* !---- DO NOT EDIT: This file autogenerated by com/sun/gluegen/JavaEmitter.java on Tue Feb 09 18:20:26 CET 2010 ----! */
+
+
+package com.mbien.opencl.demos.julia3d.structs;
+
+import java.nio.*;
+
+import com.sun.gluegen.runtime.*;
+
+
+class Camera64 extends Camera {
+
+  private final Vec orig;
+  private final Vec target;
+  private final Vec dir;
+  private final Vec x;
+  private final Vec y;
+
+  public static int size() {
+    return 60;
+  }
+
+  Camera64(java.nio.ByteBuffer buf) {
+    super(buf);
+    orig = Vec.create(accessor.slice(0, 12));
+    target = Vec.create(accessor.slice(12, 12));
+    dir = Vec.create(accessor.slice(24, 12));
+    x = Vec.create(accessor.slice(36, 12));
+    y = Vec.create(accessor.slice(48, 12));
+  }
+
+
+  public Vec getOrig() {
+    return orig;
+  }
+
+  public Vec getTarget() {
+    return target;
+  }
+
+  public Vec getDir() {
+    return dir;
+  }
+
+  public Vec getX() {
+    return x;
+  }
+
+  public Vec getY() {
+    return y;
+  }
+}
diff --git a/src/com/mbien/opencl/demos/julia3d/structs/RenderingConfig.java b/src/com/mbien/opencl/demos/julia3d/structs/RenderingConfig.java
new file mode 100644
index 0000000..f96210f
--- /dev/null
+++ b/src/com/mbien/opencl/demos/julia3d/structs/RenderingConfig.java
@@ -0,0 +1,80 @@
+/* !---- DO NOT EDIT: This file autogenerated by com/sun/gluegen/JavaEmitter.java on Tue Feb 09 18:20:26 CET 2010 ----! */
+
+
+package com.mbien.opencl.demos.julia3d.structs;
+
+import java.nio.*;
+
+import com.sun.gluegen.runtime.*;
+
+
+public abstract class RenderingConfig {
+
+  StructAccessor accessor;
+
+  public static int size() {
+    if (CPU.is32Bit()) {
+      return RenderingConfig32.size();
+    } else {
+      return RenderingConfig64.size();
+    }
+  }
+
+  public static RenderingConfig create() {
+    return create(BufferFactory.newDirectByteBuffer(size()));
+  }
+
+  public static RenderingConfig create(java.nio.ByteBuffer buf) {
+    if (CPU.is32Bit()) {
+      return new RenderingConfig32(buf);
+    } else {
+      return new RenderingConfig64(buf);
+    }
+  }
+
+  RenderingConfig(java.nio.ByteBuffer buf) {
+    accessor = new StructAccessor(buf);
+  }
+
+  public java.nio.ByteBuffer getBuffer() {
+    return accessor.getBuffer();
+  }
+
+  public abstract RenderingConfig setWidth(int val);
+
+  public abstract int getWidth();
+
+  public abstract RenderingConfig setHeight(int val);
+
+  public abstract int getHeight();
+
+  public abstract RenderingConfig setSuperSamplingSize(int val);
+
+  public abstract int getSuperSamplingSize();
+
+  public abstract RenderingConfig setActvateFastRendering(int val);
+
+  public abstract int getActvateFastRendering();
+
+  public abstract RenderingConfig setEnableShadow(int val);
+
+  public abstract int getEnableShadow();
+
+  public abstract RenderingConfig setMaxIterations(int val);
+
+  public abstract int getMaxIterations();
+
+  public abstract RenderingConfig setEpsilon(float val);
+
+  public abstract float getEpsilon();
+
+  public abstract RenderingConfig setMu(float[] val);
+
+  public abstract float[] getMu();
+
+  public abstract RenderingConfig setLight(float[] val);
+
+  public abstract float[] getLight();
+
+  public abstract Camera getCamera();
+}
diff --git a/src/com/mbien/opencl/demos/julia3d/structs/RenderingConfig32.java b/src/com/mbien/opencl/demos/julia3d/structs/RenderingConfig32.java
new file mode 100644
index 0000000..ec32613
--- /dev/null
+++ b/src/com/mbien/opencl/demos/julia3d/structs/RenderingConfig32.java
@@ -0,0 +1,106 @@
+/* !---- DO NOT EDIT: This file autogenerated by com/sun/gluegen/JavaEmitter.java on Tue Feb 09 18:20:26 CET 2010 ----! */
+
+
+package com.mbien.opencl.demos.julia3d.structs;
+
+import java.nio.*;
+
+import com.sun.gluegen.runtime.*;
+
+
+class RenderingConfig32 extends RenderingConfig {
+
+  public static int size() {
+    return 140;
+  }
+
+  RenderingConfig32(java.nio.ByteBuffer buf) {
+    super(buf);
+  }
+
+
+  public RenderingConfig setWidth(int val) {
+    accessor.setIntAt(0, val);
+    return this;
+  }
+
+  public int getWidth() {
+    return accessor.getIntAt(0);
+  }
+
+  public RenderingConfig setHeight(int val) {
+    accessor.setIntAt(1, val);
+    return this;
+  }
+
+  public int getHeight() {
+    return accessor.getIntAt(1);
+  }
+
+  public RenderingConfig setSuperSamplingSize(int val) {
+    accessor.setIntAt(2, val);
+    return this;
+  }
+
+  public int getSuperSamplingSize() {
+    return accessor.getIntAt(2);
+  }
+
+  public RenderingConfig setActvateFastRendering(int val) {
+    accessor.setIntAt(3, val);
+    return this;
+  }
+
+  public int getActvateFastRendering() {
+    return accessor.getIntAt(3);
+  }
+
+  public RenderingConfig setEnableShadow(int val) {
+    accessor.setIntAt(4, val);
+    return this;
+  }
+
+  public int getEnableShadow() {
+    return accessor.getIntAt(4);
+  }
+
+  public RenderingConfig setMaxIterations(int val) {
+    accessor.setIntAt(5, val);
+    return this;
+  }
+
+  public int getMaxIterations() {
+    return accessor.getIntAt(5);
+  }
+
+  public RenderingConfig setEpsilon(float val) {
+    accessor.setFloatAt(6, val);
+    return this;
+  }
+
+  public float getEpsilon() {
+    return accessor.getFloatAt(6);
+  }
+
+  public RenderingConfig setMu(float[] val) {
+    accessor.setFloatsAt(8, val);
+    return this;
+  }
+
+  public float[] getMu() {
+    return accessor.getFloatsAt(8, new float[4]);
+  }
+
+  public RenderingConfig setLight(float[] val) {
+    accessor.setFloatsAt(12, val);
+    return this;
+  }
+
+  public float[] getLight() {
+    return accessor.getFloatsAt(12, new float[3]);
+  }
+
+  public Camera getCamera() {
+    return Camera.create(accessor.slice(64, 76));
+  }
+}
diff --git a/src/com/mbien/opencl/demos/julia3d/structs/RenderingConfig64.java b/src/com/mbien/opencl/demos/julia3d/structs/RenderingConfig64.java
new file mode 100644
index 0000000..452fbd3
--- /dev/null
+++ b/src/com/mbien/opencl/demos/julia3d/structs/RenderingConfig64.java
@@ -0,0 +1,109 @@
+/* !---- DO NOT EDIT: This file autogenerated by com/sun/gluegen/JavaEmitter.java on Tue Feb 09 18:20:26 CET 2010 ----! */
+
+
+package com.mbien.opencl.demos.julia3d.structs;
+
+import java.nio.*;
+
+import com.sun.gluegen.runtime.*;
+
+
+class RenderingConfig64 extends RenderingConfig {
+
+  private final Camera camera;
+
+  public static int size() {
+    return 116;
+  }
+
+  RenderingConfig64(java.nio.ByteBuffer buf) {
+    super(buf);
+    camera = Camera.create(accessor.slice(56, 60));
+  }
+
+
+  public RenderingConfig setWidth(int val) {
+    accessor.setIntAt(0, val);
+    return this;
+  }
+
+  public int getWidth() {
+    return accessor.getIntAt(0);
+  }
+
+  public RenderingConfig setHeight(int val) {
+    accessor.setIntAt(1, val);
+    return this;
+  }
+
+  public int getHeight() {
+    return accessor.getIntAt(1);
+  }
+
+  public RenderingConfig setSuperSamplingSize(int val) {
+    accessor.setIntAt(2, val);
+    return this;
+  }
+
+  public int getSuperSamplingSize() {
+    return accessor.getIntAt(2);
+  }
+
+  public RenderingConfig setActvateFastRendering(int val) {
+    accessor.setIntAt(3, val);
+    return this;
+  }
+
+  public int getActvateFastRendering() {
+    return accessor.getIntAt(3);
+  }
+
+  public RenderingConfig setEnableShadow(int val) {
+    accessor.setIntAt(4, val);
+    return this;
+  }
+
+  public int getEnableShadow() {
+    return accessor.getIntAt(4);
+  }
+
+  public RenderingConfig setMaxIterations(int val) {
+    accessor.setIntAt(5, val);
+    return this;
+  }
+
+  public int getMaxIterations() {
+    return accessor.getIntAt(5);
+  }
+
+  public RenderingConfig setEpsilon(float val) {
+    accessor.setFloatAt(6, val);
+    return this;
+  }
+
+  public float getEpsilon() {
+    return accessor.getFloatAt(6);
+  }
+
+  public RenderingConfig setMu(float[] val) {
+    accessor.setFloatsAt(7, val);
+    return this;
+  }
+
+  public float[] getMu() {
+    return accessor.getFloatsAt(7, new float[4]);
+  }
+
+  public RenderingConfig setLight(float[] val) {
+    accessor.setFloatsAt(11, val);
+    return this;
+  }
+
+  public float[] getLight() {
+    return accessor.getFloatsAt(11, new float[3]);
+  }
+
+  public Camera getCamera() {
+    return camera;
+  }
+}
diff --git a/src/com/mbien/opencl/demos/julia3d/structs/Vec.java b/src/com/mbien/opencl/demos/julia3d/structs/Vec.java
new file mode 100644
index 0000000..1344db4
--- /dev/null
+++ b/src/com/mbien/opencl/demos/julia3d/structs/Vec.java
@@ -0,0 +1,54 @@
+/* !---- DO NOT EDIT: This file autogenerated by com/sun/gluegen/JavaEmitter.java on Tue Feb 09 18:20:26 CET 2010 ----! */
+
+
+package com.mbien.opencl.demos.julia3d.structs;
+
+import java.nio.*;
+
+import com.sun.gluegen.runtime.*;
+
+
+public abstract class Vec {
+
+  StructAccessor accessor;
+
+  public static int size() {
+    if (CPU.is32Bit()) {
+      return Vec32.size();
+    } else {
+      return Vec64.size();
+    }
+  }
+
+  public static Vec create() {
+    return create(BufferFactory.newDirectByteBuffer(size()));
+  }
+
+  public static Vec create(java.nio.ByteBuffer buf) {
+    if (CPU.is32Bit()) {
+      return new Vec32(buf);
+    } else {
+      return new Vec64(buf);
+    }
+  }
+
+  Vec(java.nio.ByteBuffer buf) {
+    accessor = new StructAccessor(buf);
+  }
+
+  public java.nio.ByteBuffer getBuffer() {
+    return accessor.getBuffer();
+  }
+
+  public abstract Vec setX(float val);
+
+  public abstract float getX();
+
+  public abstract Vec setY(float val);
+
+  public abstract float getY();
+
+  public abstract Vec setZ(float val);
+
+  public abstract float getZ();
+}
diff --git a/src/com/mbien/opencl/demos/julia3d/structs/Vec32.java b/src/com/mbien/opencl/demos/julia3d/structs/Vec32.java
new file mode 100644
index 0000000..970dd35
--- /dev/null
+++ b/src/com/mbien/opencl/demos/julia3d/structs/Vec32.java
@@ -0,0 +1,48 @@
+/* !---- DO NOT EDIT: This file autogenerated by com/sun/gluegen/JavaEmitter.java on Tue Feb 09 18:20:26 CET 2010 ----! */
+
+
+package com.mbien.opencl.demos.julia3d.structs;
+
+import java.nio.*;
+
+import com.sun.gluegen.runtime.*;
+
+
+class Vec32 extends Vec {
+
+  public static int size() {
+    return 12;
+  }
+
+  Vec32(java.nio.ByteBuffer buf) {
+    super(buf);
+  }
+
+
+  public Vec setX(float val) {
+    accessor.setFloatAt(0, val);
+    return this;
+  }
+
+  public float getX() {
+    return accessor.getFloatAt(0);
+  }
+
+  public Vec setY(float val) {
+    accessor.setFloatAt(1, val);
+    return this;
+  }
+
+  public float getY() {
+    return accessor.getFloatAt(1);
+  }
+
+  public Vec setZ(float val) {
+    accessor.setFloatAt(2, val);
+    return this;
+  }
+
+  public float getZ() {
+    return accessor.getFloatAt(2);
+  }
+}
diff --git a/src/com/mbien/opencl/demos/julia3d/structs/Vec64.java b/src/com/mbien/opencl/demos/julia3d/structs/Vec64.java
new file mode 100644
index 0000000..a52148e
--- /dev/null
+++ b/src/com/mbien/opencl/demos/julia3d/structs/Vec64.java
@@ -0,0 +1,48 @@
+/* !---- DO NOT EDIT: This file autogenerated by com/sun/gluegen/JavaEmitter.java on Tue Feb 09 18:20:26 CET 2010 ----! */
+
+
+package com.mbien.opencl.demos.julia3d.structs;
+
+import java.nio.*;
+
+import com.sun.gluegen.runtime.*;
+
+
+class Vec64 extends Vec {
+
+  public static int size() {
+    return 12;
+  }
+
+  Vec64(java.nio.ByteBuffer buf) {
+    super(buf);
+  }
+
+
+  public Vec setX(float val) {
+    accessor.setFloatAt(0, val);
+    return this;
+  }
+
+  public float getX() {
+    return accessor.getFloatAt(0);
+  }
+
+  public Vec setY(float val) {
+    accessor.setFloatAt(1, val);
+    return this;
+  }
+
+  public float getY() {
+    return accessor.getFloatAt(1);
+  }
+
+  public Vec setZ(float val) {
+    accessor.setFloatAt(2, val);
+    return this;
+  }
+
+  public float getZ() {
+    return accessor.getFloatAt(2);
+  }
+}