changes due to package renaming in jocl.

15 files changed, 0 insertions, 2016 deletions

diff --git a/src/com/mbien/opencl/demos/julia3d/Julia3d.java b/src/com/mbien/opencl/demos/julia3d/Julia3d.java
deleted file mode 100644
index 5d5378f..0000000
--- a/src/com/mbien/opencl/demos/julia3d/Julia3d.java
+++ /dev/null
@@ -1,212 +0,0 @@
-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.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();
-        context = CLContext.create(gpu);
-
-        workGroupSize = 256;
-
-        //allocate buffers
-        configBuffer = context.createBuffer(config.getBuffer(), READ_ONLY);
-        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
deleted file mode 100644
index 9266aba..0000000
--- a/src/com/mbien/opencl/demos/julia3d/Renderer.java
+++ /dev/null
@@ -1,203 +0,0 @@
-package com.mbien.opencl.demos.julia3d;
-
-import com.mbien.opencl.demos.julia3d.structs.RenderingConfig;
-import com.jogamp.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 com.jogamp.common.nio.Buffers.*;
-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 = newDirectFloatBuffer(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
deleted file mode 100644
index 849de3a..0000000
--- a/src/com/mbien/opencl/demos/julia3d/UserSceneController.java
+++ /dev/null
@@ -1,249 +0,0 @@
-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
deleted file mode 100644
index 72df3ff..0000000
--- a/src/com/mbien/opencl/demos/julia3d/config.h
+++ /dev/null
@@ -1,24 +0,0 @@
-
-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/mandelbrot_kernel.cl b/src/com/mbien/opencl/demos/julia3d/mandelbrot_kernel.cl
deleted file mode 100644
index d5acd02..0000000
--- a/src/com/mbien/opencl/demos/julia3d/mandelbrot_kernel.cl
+++ /dev/null
@@ -1,357 +0,0 @@
-/*
-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
-
-
-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;
-
-#define BOUNDING_RADIUS_2 4.f
-
-// Scalar derivative approach by Enforcer:
-// http://www.fractalforums.com/mandelbulb-implementation/realtime-renderingoptimisations/
-static float IterateIntersect(const float4 z0, const float4 c0, const uint maxIterations) {
-	float4 z = z0;
-	float4 c = c0;
-
-	float dr = 1.0f;
-	float r2 = dot(z, z);
-	float r = sqrt(r2);
-	for (int n = 0; (n < maxIterations) && (r < 2.f); ++n) {
-		const float zo0 = asin(z.z / r);
-		const float zi0 = atan2(z.y, z.x);
-		float zr = r2 * r2 * r2 * r;
-		const float zo = zo0 * 7.f;
-		const float zi = zi0 * 7.f;
-		const float czo = cos(zo);
-
-		dr = zr * dr * 7.f + 1.f;
-		zr *= r;
-
-		z = zr * (float4)(czo * cos(zi), czo * sin(zi), sin(zo), 0.f);
-		z += c;
-
-		r2 = dot(z, z);
-		r = sqrt(r2);
-	}
-
-	return 0.5f * log(r) * r / dr;
-}
-
-static float IntersectBulb(const float4 eyeRayOrig, const float4 eyeRayDir,
-		const float4 c, const uint maxIterations, const float epsilon,
-		const float maxDist, float4 *hitPoint, uint *steps) {
-	float dist;
-	float4 r0 = eyeRayOrig;
-	float distDone = 0.f;
-
-	uint s = 0;
-	do {
-		dist = IterateIntersect(r0, c, maxIterations);
-		distDone += dist;
-		// We are inside
-		if (dist <= 0.f)
-			break;
-
-		r0 += eyeRayDir * dist;
-		s++;
-	} while ((dist > epsilon) && (distDone < maxDist));
-
-	*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;
-	}
-}
-
-int IntersectBoundingSphere(const float4 eyeRayOrig, const float4 eyeRayDir,
-		float *tmin, float*tmax) {
-	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 0;
-	else
-		det = sqrt(det);
-
-	float t1 = b - det;
-	float t2 = b + det;
-	if (t1 > 0.f) {
-		*tmin = t1;
-		*tmax = t2;
-		return 1;
-	} else {
-		if (t2 > 0.f) {
-			// We are inside, start from the ray origin
-			*tmin = 0.f;
-			*tmax = t2;
-
-			return 1;
-		} else
-			return 0;
-	}
-}
-
-static float4 NormEstimate(const float4 p, const float4 c,
-		const float delta, const uint maxIterations) {
-	const float4 qP = p;
-	const float4 gx1 = qP - (float4)(delta, 0.f, 0.f, 0.f);
-	const float4 gx2 = qP + (float4)(delta, 0.f, 0.f, 0.f);
-	const float4 gy1 = qP - (float4)(0.f, delta, 0.f, 0.f);
-	const float4 gy2 = qP + (float4)(0.f, delta, 0.f, 0.f);
-	const float4 gz1 = qP - (float4)(0.f, 0.f, delta, 0.f);
-	const float4 gz2 = qP + (float4)(0.f, 0.f, delta, 0.f);
-
-	const float gradX = length(IterateIntersect(gx2, c, maxIterations)) -
-		length(IterateIntersect(gx1, c, maxIterations));
-	const float gradY = length(IterateIntersect(gy2, c, maxIterations)) -
-		length(IterateIntersect(gy1, c, maxIterations));
-	const float gradZ = length(IterateIntersect(gz2, c, maxIterations)) -
-		length(IterateIntersect(gz1, c, maxIterations));
-
-	const float4 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 MandelbulbGPU(
-	__global float *pixels,
-	const __global RenderingConfig *config,
-	const int enableAccumulation,
-	const float sampleX,
-	const float sampleY) {
-    const int gid = get_global_id(0);
-	const unsigned width = config->width;
-	const 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.5f / 0.75f)) : config->epsilon;
-	const uint maxIterations = config->actvateFastRendering ? (max(3u, config->maxIterations) - 2u) : 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
-	//--------------------------------------------------------------------------
-
-	int useAO = 1;
-	float4 diffuse, n, color;
-
-	float4 hitPoint;
-	float dist, tmin, tmax;
-	if (IntersectBoundingSphere(eyeRayOrig, eyeRayDir, &tmin, &tmax)) {
-		//--------------------------------------------------------------------------
-		// Find the intersection with the set
-		//--------------------------------------------------------------------------
-
-		uint steps;
-		float4 rayOrig = eyeRayOrig + eyeRayDir * (float4)tmin;
-		dist = IntersectBulb(rayOrig, eyeRayDir, mu, maxIterations,
-				epsilon, tmax - tmin, &hitPoint, &steps);
-
-		if (dist <= epsilon) {
-			// Set hit
-			diffuse = (float4) (1.f, 0.35f, 0.15f, 0.f);
-			n = NormEstimate(hitPoint, mu, dist, maxIterations);
-		} else
-			dist = -1.f;
-	} else
-		dist = -1.f;
-
-	//--------------------------------------------------------------------------
-	// Check if we hit the floor
-	//--------------------------------------------------------------------------
-
-	if (dist < 0.f) {
-		dist = IntersectFloorSphere(eyeRayOrig, eyeRayDir);
-
-		if (dist >= 0.f) {
-			// Floor hit
-			hitPoint = eyeRayOrig + eyeRayDir * (float4)dist;
-			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);
-			useAO = 0;
-		} else {
-			// Sky hit
-			color = (float4)(0.f, 0.1f, 0.3f, 0.f);
-		}
-	} else {
-		// Sky hit
-		color = (float4)(0.f, 0.1f, 0.3f, 0.f);
-	}
-
-	//--------------------------------------------------------------------------
-	// Select the shadow pass
-	//--------------------------------------------------------------------------
-
-	if (dist >= 0.f) {
-		float shadowFactor = 1.f;
-		if (config->enableShadow) {
-			float4 L = normalize(light -  hitPoint);
-			float4 rO = hitPoint + n * 1e-2f;
-			float4 shadowHitPoint;
-
-			// Check bounding sphere
-			if (IntersectBoundingSphere(rO, L, &tmin, &tmax)) {
-				float shadowDistSet = tmin;
-				uint steps;
-
-				rO = rO + L * (float4)shadowDistSet;
-				shadowDistSet = IntersectBulb(rO, L, mu, maxIterations, epsilon,
-						tmax - tmin, &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;
-				}
-			}
-		}
-
-		//--------------------------------------------------------------------------
-		// 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/rendering_kernel.cl b/src/com/mbien/opencl/demos/julia3d/rendering_kernel.cl
deleted file mode 100644
index 9c25c1b..0000000
--- a/src/com/mbien/opencl/demos/julia3d/rendering_kernel.cl
+++ /dev/null
@@ -1,382 +0,0 @@
-/*
-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
deleted file mode 100644
index 50d5868..0000000
--- a/src/com/mbien/opencl/demos/julia3d/structs/Camera.java
+++ /dev/null
@@ -1,52 +0,0 @@
-/* !---- 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.jogamp.common.nio.*;
-
-
-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(Buffers.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
deleted file mode 100644
index 85dea51..0000000
--- a/src/com/mbien/opencl/demos/julia3d/structs/Camera32.java
+++ /dev/null
@@ -1,41 +0,0 @@
-/* !---- 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.jogamp.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
deleted file mode 100644
index afe4acd..0000000
--- a/src/com/mbien/opencl/demos/julia3d/structs/Camera64.java
+++ /dev/null
@@ -1,52 +0,0 @@
-/* !---- 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.jogamp.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
deleted file mode 100644
index 5bc65b5..0000000
--- a/src/com/mbien/opencl/demos/julia3d/structs/RenderingConfig.java
+++ /dev/null
@@ -1,80 +0,0 @@
-/* !---- 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.jogamp.common.nio.*;
-
-
-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(Buffers.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
deleted file mode 100644
index 3898d60..0000000
--- a/src/com/mbien/opencl/demos/julia3d/structs/RenderingConfig32.java
+++ /dev/null
@@ -1,106 +0,0 @@
-/* !---- 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.jogamp.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
deleted file mode 100644
index e3a0267..0000000
--- a/src/com/mbien/opencl/demos/julia3d/structs/RenderingConfig64.java
+++ /dev/null
@@ -1,109 +0,0 @@
-/* !---- 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.jogamp.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
deleted file mode 100644
index 3bd6b18..0000000
--- a/src/com/mbien/opencl/demos/julia3d/structs/Vec.java
+++ /dev/null
@@ -1,53 +0,0 @@
-/* !---- 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 com.jogamp.common.nio.*;
-
-
-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(Buffers.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
deleted file mode 100644
index 3c0ee5d..0000000
--- a/src/com/mbien/opencl/demos/julia3d/structs/Vec32.java
+++ /dev/null
@@ -1,48 +0,0 @@
-/* !---- 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.jogamp.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
deleted file mode 100644
index 5952021..0000000
--- a/src/com/mbien/opencl/demos/julia3d/structs/Vec64.java
+++ /dev/null
@@ -1,48 +0,0 @@
-/* !---- 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.jogamp.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);
-  }
-}