summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
authorMichael Bien <mbien@fh-landshut.de>2010-04-01 00:45:19 +0200
committerMichael Bien <mbien@fh-landshut.de>2010-04-01 00:45:19 +0200
commita09a4fcc59fd0a4fb716f4fcd9c779575045ab15 (patch)
treeec7173b3361406b1697698a0b09a18b07f87f2f5
parentfe567f8d89d6b4aa45f34b4f2557db5db0b158da (diff)
modifications due to class movement in gluegen.
-rw-r--r--src/com/mbien/opencl/demos/fractal/MultiDeviceFractal.java2
-rw-r--r--src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java2
-rw-r--r--src/com/mbien/opencl/demos/julia3d/Renderer.java2
-rw-r--r--src/com/mbien/opencl/demos/julia3d/mandelbrot_kernel.cl357
-rw-r--r--src/com/mbien/opencl/demos/julia3d/structs/Camera.java2
-rw-r--r--src/com/mbien/opencl/demos/julia3d/structs/RenderingConfig.java2
-rw-r--r--src/com/mbien/opencl/demos/julia3d/structs/Vec.java3
7 files changed, 363 insertions, 7 deletions
diff --git a/src/com/mbien/opencl/demos/fractal/MultiDeviceFractal.java b/src/com/mbien/opencl/demos/fractal/MultiDeviceFractal.java
index c7fd23d..265add8 100644
--- a/src/com/mbien/opencl/demos/fractal/MultiDeviceFractal.java
+++ b/src/com/mbien/opencl/demos/fractal/MultiDeviceFractal.java
@@ -38,7 +38,7 @@ import javax.media.opengl.awt.GLCanvas;
import javax.swing.JFrame;
import javax.swing.SwingUtilities;
-import static com.jogamp.gluegen.runtime.Buffers.*;
+import static com.jogamp.common.nio.Buffers.*;
import static javax.media.opengl.GL2.*;
import static com.mbien.opencl.CLMemory.Mem.*;
import static com.mbien.opencl.CLEvent.ProfilingCommand.*;
diff --git a/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java b/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java
index 36e67bb..7913650 100644
--- a/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java
+++ b/src/com/mbien/opencl/demos/joglinterop/GLCLInteroperabilityDemo.java
@@ -18,7 +18,7 @@ import javax.media.opengl.glu.gl2.GLUgl2;
import javax.swing.JFrame;
import javax.swing.SwingUtilities;
-import static com.jogamp.gluegen.runtime.Buffers.*;
+import static com.jogamp.common.nio.Buffers.*;
/**
* JOCL - JOGL interoperability example.
diff --git a/src/com/mbien/opencl/demos/julia3d/Renderer.java b/src/com/mbien/opencl/demos/julia3d/Renderer.java
index 5eaf219..9266aba 100644
--- a/src/com/mbien/opencl/demos/julia3d/Renderer.java
+++ b/src/com/mbien/opencl/demos/julia3d/Renderer.java
@@ -15,7 +15,7 @@ import javax.media.opengl.GLProfile;
import javax.media.opengl.awt.GLCanvas;
import javax.swing.JFrame;
-import static com.jogamp.gluegen.runtime.Buffers.*;
+import static com.jogamp.common.nio.Buffers.*;
import static javax.media.opengl.GL2.*;
import static java.lang.String.*;
diff --git a/src/com/mbien/opencl/demos/julia3d/mandelbrot_kernel.cl b/src/com/mbien/opencl/demos/julia3d/mandelbrot_kernel.cl
new file mode 100644
index 0000000..d5acd02
--- /dev/null
+++ b/src/com/mbien/opencl/demos/julia3d/mandelbrot_kernel.cl
@@ -0,0 +1,357 @@
+/*
+Copyright (c) 2009 David Bucciarelli (davibu@interfree.it)
+
+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/structs/Camera.java b/src/com/mbien/opencl/demos/julia3d/structs/Camera.java
index a160986..50d5868 100644
--- a/src/com/mbien/opencl/demos/julia3d/structs/Camera.java
+++ b/src/com/mbien/opencl/demos/julia3d/structs/Camera.java
@@ -5,7 +5,7 @@ package com.mbien.opencl.demos.julia3d.structs;
import java.nio.*;
-import com.jogamp.gluegen.runtime.*;
+import com.jogamp.common.nio.*;
public abstract class Camera {
diff --git a/src/com/mbien/opencl/demos/julia3d/structs/RenderingConfig.java b/src/com/mbien/opencl/demos/julia3d/structs/RenderingConfig.java
index 0f29bc7..5bc65b5 100644
--- a/src/com/mbien/opencl/demos/julia3d/structs/RenderingConfig.java
+++ b/src/com/mbien/opencl/demos/julia3d/structs/RenderingConfig.java
@@ -5,7 +5,7 @@ package com.mbien.opencl.demos.julia3d.structs;
import java.nio.*;
-import com.jogamp.gluegen.runtime.*;
+import com.jogamp.common.nio.*;
public abstract class RenderingConfig {
diff --git a/src/com/mbien/opencl/demos/julia3d/structs/Vec.java b/src/com/mbien/opencl/demos/julia3d/structs/Vec.java
index 89a79dd..3bd6b18 100644
--- a/src/com/mbien/opencl/demos/julia3d/structs/Vec.java
+++ b/src/com/mbien/opencl/demos/julia3d/structs/Vec.java
@@ -3,9 +3,8 @@
package com.mbien.opencl.demos.julia3d.structs;
-import java.nio.*;
-import com.jogamp.gluegen.runtime.*;
+import com.jogamp.common.nio.*;
public abstract class Vec {