qjulia_kernel.cl
// |
// File: qjulia.c |
// |
// Abstract: This example shows how to use OpenCL to raytrace a 4d Quaternion Julia-Set |
// Fractal and intermix the results of a compute kernel with OpenGL for rendering. |
// |
// Version: <1.0> |
// |
// Disclaimer: IMPORTANT: This Apple software is supplied to you by Apple Inc. ("Apple") |
// in consideration of your agreement to the following terms, and your use, |
// installation, modification or redistribution of this Apple software |
// constitutes acceptance of these terms. If you do not agree with these |
// terms, please do not use, install, modify or redistribute this Apple |
// software. |
// |
// In consideration of your agreement to abide by the following terms, and |
// subject to these terms, Apple grants you a personal, non - exclusive |
// license, under Apple's copyrights in this original Apple software ( the |
// "Apple Software" ), to use, reproduce, modify and redistribute the Apple |
// Software, with or without modifications, in source and / or binary forms; |
// provided that if you redistribute the Apple Software in its entirety and |
// without modifications, you must retain this notice and the following text |
// and disclaimers in all such redistributions of the Apple Software. Neither |
// the name, trademarks, service marks or logos of Apple Inc. may be used to |
// endorse or promote products derived from the Apple Software without specific |
// prior written permission from Apple. Except as expressly stated in this |
// notice, no other rights or licenses, express or implied, are granted by |
// Apple herein, including but not limited to any patent rights that may be |
// infringed by your derivative works or by other works in which the Apple |
// Software may be incorporated. |
// |
// The Apple Software is provided by Apple on an "AS IS" basis. APPLE MAKES NO |
// WARRANTIES, EXPRESS OR IMPLIED, INCLUDING WITHOUT LIMITATION THE IMPLIED |
// WARRANTIES OF NON - INFRINGEMENT, MERCHANTABILITY AND FITNESS FOR A |
// PARTICULAR PURPOSE, REGARDING THE APPLE SOFTWARE OR ITS USE AND OPERATION |
// ALONE OR IN COMBINATION WITH YOUR PRODUCTS. |
// |
// IN NO EVENT SHALL APPLE BE LIABLE FOR ANY SPECIAL, INDIRECT, INCIDENTAL OR |
// CONSEQUENTIAL DAMAGES ( INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF |
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS |
// INTERRUPTION ) ARISING IN ANY WAY OUT OF THE USE, REPRODUCTION, MODIFICATION |
// AND / OR DISTRIBUTION OF THE APPLE SOFTWARE, HOWEVER CAUSED AND WHETHER |
// UNDER THEORY OF CONTRACT, TORT ( INCLUDING NEGLIGENCE ), STRICT LIABILITY OR |
// OTHERWISE, EVEN IF APPLE HAS BEEN ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. |
// |
// Copyright ( C ) 2008 Apple Inc. All Rights Reserved. |
// |
//////////////////////////////////////////////////////////////////////////////////////////////////// |
#ifndef WIDTH |
#define WIDTH (512) |
#endif |
#ifndef HEIGHT |
#define HEIGHT (512) |
#endif |
#define ASPECT ((float)WIDTH / (float)HEIGHT) |
#define SQR(x) ((x)*(x)) |
#define BOUNDING_RADIUS (2.0f) |
#define BOUNDING_RADIUS_SQR (SQR(BOUNDING_RADIUS)) |
#define ESCAPE_THRESHOLD (BOUNDING_RADIUS * 1.5f) |
#define DELTA (1e-5f) |
#define ITERATIONS (10) |
#define EPSILON (0.003f) |
#define SHADOWS (0) |
//////////////////////////////////////////////////////////////////////////////////////////////////// |
float4 qmult( float4 q1, float4 q2 ) |
{ |
float4 r; |
float3 t; |
float3 q1yzw = (float3)(q1.y, q1.z, q1.w); |
float3 q2yzw = (float3)(q2.y, q2.z, q2.w); |
float3 c = cross( q1yzw, q2yzw ); |
t = q2yzw * q1.x + q1yzw * q2.x + c; |
r.x = q1.x * q2.x - dot( q1yzw, q2yzw ); |
r.yzw = t.xyz; |
return r; |
} |
float4 qsqr( float4 q ) |
{ |
float4 r; |
float3 t; |
float3 qyzw = (float3)(q.y, q.z, q.w); |
t = 2.0f * q.x * qyzw; |
r.x = q.x * q.x - dot( qyzw, qyzw ); |
r.yzw = t.xyz; |
return r; |
} |
//////////////////////////////////////////////////////////////////////////////////////////////////// |
float3 |
EstimateNormalQJulia( |
float3 p, |
float4 c, |
int iterations ) |
{ |
float4 qp = (float4)( p.x, p.y, p.z, 0.0f ); |
float4 gx1 = qp - (float4)( DELTA, 0.0f, 0.0f, 0.0f ); |
float4 gx2 = qp + (float4)( DELTA, 0.0f, 0.0f, 0.0f ); |
float4 gy1 = qp - (float4)( 0.0f, DELTA, 0.0f, 0.0f ); |
float4 gy2 = qp + (float4)( 0.0f, DELTA, 0.0f, 0.0f ); |
float4 gz1 = qp - (float4)( 0.0f, 0.0f, DELTA, 0.0f ); |
float4 gz2 = qp + (float4)( 0.0f, 0.0f, DELTA, 0.0f ); |
for ( int i = 0; i < iterations; i++ ) |
{ |
gx1 = qsqr( gx1 ) + c; |
gx2 = qsqr( gx2 ) + c; |
gy1 = qsqr( gy1 ) + c; |
gy2 = qsqr( gy2 ) + c; |
gz1 = qsqr( gz1 ) + c; |
gz2 = qsqr( gz2 ) + c; |
} |
float nx = fast_length(gx2) - fast_length(gx1); |
float ny = fast_length(gy2) - fast_length(gy1); |
float nz = fast_length(gz2) - fast_length(gz1); |
float3 normal = fast_normalize((float3)( nx, ny, nz )); |
return normal; |
} |
//////////////////////////////////////////////////////////////////////////////////////////////////// |
float4 |
IntersectQJulia( |
float3 rO, |
float3 rD, |
float4 c, |
float epsilon, |
float escape) |
{ |
float rd = 0.0f; |
float dist = epsilon; |
while ( dist >= epsilon && rd < escape) |
{ |
float4 z = (float4)( rO.x, rO.y, rO.z, 0.0f ); |
float4 zp = (float4)( 1.0f, 0.0f, 0.0f, 0.0f ); |
float zd = 0.0f; |
uint count = 0; |
while(zd < escape && count < ITERATIONS) |
{ |
zp = 2.0f * qmult(z, zp); |
z = qsqr(z) + c; |
zd = dot(z, z); |
count++; |
} |
float normZ = fast_length( z ); |
dist = 0.5f * normZ * half_log( normZ ) / fast_length( zp ); |
rO += rD * dist; |
rd = dot(rO, rO); |
} |
return (float4)(rO.x, rO.y, rO.z, dist); |
} |
//////////////////////////////////////////////////////////////////////////////////////////////////// |
float3 |
Phong( |
float3 light, |
float3 eye, |
float3 pt, |
float3 normal, |
float3 base) |
{ |
const float SpecularExponent = 10.0f; |
const float Specularity = 0.45f; |
float3 light_dir = fast_normalize( light - pt ); |
float3 eye_dir = fast_normalize( eye - pt ); |
float NdotL = dot( normal, light_dir ); |
float3 reflect_dir = light_dir - 2.0f * NdotL * normal; |
base += fabs(normal) * 0.5f; |
float3 diffuse = base * fmax(NdotL, 0.0f); |
float3 specular = Specularity * half_powr( fmax( dot(eye_dir, reflect_dir), 0.0f), SpecularExponent ); |
return diffuse + specular; |
} |
//////////////////////////////////////////////////////////////////////////////////////////////////// |
float |
IntersectSphere( |
float3 rO, |
float3 rD, |
float radius ) |
{ |
float fB = 2.0f * dot( rO, rD ); |
float fB2 = fB * fB; |
float fC = dot( rO, rO ) - radius; |
float fT = (fB2 - 4.0f * fC); |
if (fT <= 0.0f) |
return 0.0f; |
float fD = half_sqrt( fT ); |
float fT0 = ( -fB + fD ) * 0.5f; |
float fT1 = ( -fB - fD ) * 0.5f; |
fT = fmin(fT0, fT1); |
return fT; |
} |
//////////////////////////////////////////////////////////////////////////////////////////////////// |
float4 |
RaytraceQJulia( |
float3 rO, |
float3 rD, |
float4 mu, |
float epsilon, |
float3 eye, |
float3 light, |
float3 diffuse, |
float radius, |
bool shadows, |
int iterations ) |
{ |
const float4 background = (float4)( 0.15f, 0.15f, 0.15f, 0.0f ); |
float4 color = background; |
rD = fast_normalize( rD ); |
float t = IntersectSphere( rO, rD, radius ); |
if ( t <= 0.0f ) |
return color; |
rO += rD * t; |
float4 hit = IntersectQJulia( rO, rD, mu, epsilon, ESCAPE_THRESHOLD ); |
float dist = hit.w; |
if (dist >= epsilon) |
return color; |
rO.xyz = hit.xyz; |
float3 normal = EstimateNormalQJulia( rO, mu, iterations ); |
float3 rgb = Phong( light, rD, rO, normal, diffuse ); |
color.xyz = rgb.xyz; |
color.w = 1.0f; |
if (SHADOWS) |
{ |
float3 light_dir = fast_normalize( light - rO ); |
rO += normal * epsilon * 2.0f; |
hit = IntersectQJulia( rO, light_dir, mu, epsilon, ESCAPE_THRESHOLD ); |
dist = hit.w; |
color.xyz *= (dist < epsilon) ? (0.4f) : (1.0f); |
} |
return color; |
} |
/////////////////////////////////////////////////////////////////////////////////////////// |
float4 |
QJulia( |
float4 coord, |
float4 mu, |
float4 diffuse, |
float epsilon, |
float iterations, |
int shadows, |
uint width, |
uint height) |
{ |
float zoom = BOUNDING_RADIUS_SQR; |
float radius = BOUNDING_RADIUS_SQR; |
float2 size = (float2)((float)width, (float)height); |
float scale = max(size.x, size.y); |
float2 h = (float2)(0.5f, 0.5f); // half |
float2 position = (coord.xy - h * size) / scale; |
float2 frame = (position) * zoom; |
float3 light = (float3)(1.5f, 0.5f, 4.0f); |
float3 eye = (float3)(0.0f, 0.0f, 4.0f); |
float3 ray = (float3)(frame.x, frame.y, 0.0f); |
float3 base = (float3)(diffuse.x, diffuse.y, diffuse.z); |
float3 rO = eye; |
float3 rD = (ray - rO); |
float4 color = RaytraceQJulia( rO, rD, mu, epsilon, eye, light, base, radius, shadows, iterations); |
return color; |
} |
//////////////////////////////////////////////////////////////////////////////////////////////////// |
__kernel void |
QJuliaKernel( |
__global uchar4 *result, |
const float4 mu, |
const float4 diffuse, |
const float epsilon) |
{ |
int tx = get_global_id(0); |
int ty = get_global_id(1); |
int sx = get_global_size(0); |
int sy = get_global_size(1); |
int index = ty * WIDTH + tx; |
bool valid = (tx < WIDTH) && (ty < HEIGHT); |
float4 coord = (float4)((float)tx, (float)ty, 0.0f, 0.0f); |
if(valid) |
{ |
float4 color = QJulia(coord, mu, diffuse, epsilon, ITERATIONS, SHADOWS, WIDTH, HEIGHT); |
uchar4 output = convert_uchar4_sat_rte(color * 255.0f); |
result[index] = output; |
} |
} |
//////////////////////////////////////////////////////////////////////////////////////////////////// |
Copyright © 2011 Apple Inc. All Rights Reserved. Terms of Use | Privacy Policy | Updated: 2011-10-03