diff options
Diffstat (limited to 'src/com')
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); + } +} |