diff options
Diffstat (limited to 'src/com/mbien/opencl')
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 ([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/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 { |