1 // Created on: 2013-10-16
2 // Created by: Denis BOGOLEPOV
3 // Copyright (c) 2013 OPEN CASCADE SAS
5 // This file is part of Open CASCADE Technology software library.
7 // This library is free software; you can redistribute it and / or modify it
8 // under the terms of the GNU Lesser General Public version 2.1 as published
9 // by the Free Software Foundation, with special exception defined in the file
10 // OCCT_LGPL_EXCEPTION.txt. Consult the file LICENSE_LGPL_21.txt included in OCCT
11 // distribution for complete text of the license and disclaimer of any warranty.
13 // Alternatively, this file may be used under the terms of Open CASCADE
14 // commercial license or contractual agreement.
24 extern const char THE_RAY_TRACE_OPENCL_SOURCE[] =
26 /////////////////////////////////////////////////////////////////////////////////////////
27 // Specific data types
29 //! Stores ray parameters.
30 EOL" typedef struct __SRay"
37 //! Stores parameters of intersection point.
38 EOL" typedef struct __SIntersect"
48 /////////////////////////////////////////////////////////////////////////////////////////
49 // Some useful vector constants
51 EOL" #define ZERO ( float4 )( 0.f, 0.f, 0.f, 0.f )"
52 EOL" #define UNIT ( float4 )( 1.f, 1.f, 1.f, 0.f )"
54 EOL" #define AXIS_X ( float4 )( 1.f, 0.f, 0.f, 0.f )"
55 EOL" #define AXIS_Y ( float4 )( 0.f, 1.f, 0.f, 0.f )"
56 EOL" #define AXIS_Z ( float4 )( 0.f, 0.f, 1.f, 0.f )"
58 /////////////////////////////////////////////////////////////////////////////////////////
61 // =======================================================================
62 // function : GenerateRay
63 // purpose : Generates primary ray for current work item
64 // =======================================================================
65 EOL" void GenerateRay (SRay* theRay,"
66 EOL" const float theX,"
67 EOL" const float theY,"
68 EOL" const int theSizeX,"
69 EOL" const int theSizeY,"
70 EOL" const float16 theOrigins,"
71 EOL" const float16 theDirects)"
73 EOL" float2 aPixel = (float2) (theX / (float)theSizeX,"
74 EOL" theY / (float)theSizeY);"
76 EOL" float4 aP0 = mix (theOrigins.lo.lo, theOrigins.lo.hi, aPixel.x);"
77 EOL" float4 aP1 = mix (theOrigins.hi.lo, theOrigins.hi.hi, aPixel.x);"
79 EOL" theRay->Origin = mix (aP0, aP1, aPixel.y);"
81 EOL" aP0 = mix (theDirects.lo.lo, theDirects.lo.hi, aPixel.x);"
82 EOL" aP1 = mix (theDirects.hi.lo, theDirects.hi.hi, aPixel.x);"
84 EOL" theRay->Direct = mix (aP0, aP1, aPixel.y);"
88 /////////////////////////////////////////////////////////////////////////////////////////
89 // Functions for compute ray-object intersection
91 EOL" #define _OOEPS_ exp2 (-80.0f)"
93 // =======================================================================
94 // function : IntersectSphere
95 // purpose : Computes ray-sphere intersection
96 // =======================================================================
97 EOL" float IntersectSphere (const SRay* theRay, float theRadius)"
99 EOL" float aDdotD = dot (theRay->Direct.xyz, theRay->Direct.xyz);"
100 EOL" float aDdotO = dot (theRay->Direct.xyz, theRay->Origin.xyz);"
101 EOL" float aOdotO = dot (theRay->Origin.xyz, theRay->Origin.xyz);"
103 EOL" float aD = aDdotO * aDdotO - aDdotD * (aOdotO - theRadius * theRadius);"
107 EOL" float aTime = (-aDdotO + native_sqrt (aD)) * (1.f / aDdotD);"
109 EOL" return aTime > 0.f ? aTime : MAXFLOAT;"
112 EOL" return MAXFLOAT;"
115 // =======================================================================
116 // function : IntersectBox
117 // purpose : Computes ray-box intersection (slab test)
118 // =======================================================================
119 EOL" float IntersectBox (const SRay* theRay,"
120 EOL" float4 theMinPoint,"
121 EOL" float4 theMaxPoint)"
123 EOL" const float4 aInvDirect = (float4)("
124 EOL" 1.f / (fabs (theRay->Direct.x) > _OOEPS_ ?"
125 EOL" theRay->Direct.x : copysign (_OOEPS_, theRay->Direct.x)),"
126 EOL" 1.f / (fabs (theRay->Direct.y) > _OOEPS_ ?"
127 EOL" theRay->Direct.y : copysign (_OOEPS_, theRay->Direct.y)),"
128 EOL" 1.f / (fabs (theRay->Direct.z) > _OOEPS_ ?"
129 EOL" theRay->Direct.z : copysign (_OOEPS_, theRay->Direct.z)),"
132 EOL" const float4 aTime0 = (theMinPoint - theRay->Origin) * aInvDirect;"
133 EOL" const float4 aTime1 = (theMaxPoint - theRay->Origin) * aInvDirect;"
135 EOL" const float4 aTimeMax = max (aTime0, aTime1);"
136 EOL" const float4 aTimeMin = min (aTime0, aTime1);"
138 EOL" const float theTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
139 EOL" const float theTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
141 EOL" return (theTimeStart <= theTimeFinal) && (theTimeFinal >= 0.f) ? theTimeStart : MAXFLOAT;"
144 // =======================================================================
145 // function : IntersectNodes
146 // purpose : Computes intersection of ray with two child nodes (boxes)
147 // =======================================================================
148 EOL" void IntersectNodes (const SRay* theRay,"
149 EOL" float4 theMinPoint0,"
150 EOL" float4 theMaxPoint0,"
151 EOL" float4 theMinPoint1,"
152 EOL" float4 theMaxPoint1,"
153 EOL" float* theTimeStart0,"
154 EOL" float* theTimeStart1,"
155 EOL" float theMaxTime)"
157 EOL" const float4 aInvDirect = (float4)("
158 EOL" 1.f / (fabs (theRay->Direct.x) > _OOEPS_ ?"
159 EOL" theRay->Direct.x : copysign (_OOEPS_, theRay->Direct.x)),"
160 EOL" 1.f / (fabs (theRay->Direct.y) > _OOEPS_ ?"
161 EOL" theRay->Direct.y : copysign (_OOEPS_, theRay->Direct.y)),"
162 EOL" 1.f / (fabs (theRay->Direct.z) > _OOEPS_ ?"
163 EOL" theRay->Direct.z : copysign (_OOEPS_, theRay->Direct.z)),"
166 EOL" float4 aTime0 = (theMinPoint0 - theRay->Origin) * aInvDirect;"
167 EOL" float4 aTime1 = (theMaxPoint0 - theRay->Origin) * aInvDirect;"
169 EOL" float4 aTimeMax = max (aTime0, aTime1);"
170 EOL" float4 aTimeMin = min (aTime0, aTime1);"
172 EOL" aTime0 = (theMinPoint1 - theRay->Origin) * aInvDirect;"
173 EOL" aTime1 = (theMaxPoint1 - theRay->Origin) * aInvDirect;"
175 EOL" float aTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
176 EOL" float aTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
178 EOL" aTimeMax = max (aTime0, aTime1);"
179 EOL" aTimeMin = min (aTime0, aTime1);"
181 EOL" *theTimeStart0 = (aTimeStart <= aTimeFinal) & (aTimeFinal >= 0.f) & (aTimeStart <= theMaxTime)"
182 EOL" ? aTimeStart : -MAXFLOAT;"
184 EOL" aTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
185 EOL" aTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
187 EOL" *theTimeStart1 = (aTimeStart <= aTimeFinal) & (aTimeFinal >= 0.f) & (aTimeStart <= theMaxTime)"
188 EOL" ? aTimeStart : -MAXFLOAT;"
191 // =======================================================================
192 // function : IntersectTriangle
193 // purpose : Computes ray-triangle intersection (branchless version)
194 // =======================================================================
195 EOL" float IntersectTriangle (const SRay* theRay,"
196 EOL" const float4 thePoint0,"
197 EOL" const float4 thePoint1,"
198 EOL" const float4 thePoint2,"
199 EOL" float4* theNormal,"
203 EOL" const float4 aEdge0 = thePoint1 - thePoint0;"
204 EOL" const float4 aEdge1 = thePoint0 - thePoint2;"
206 EOL" *theNormal = cross (aEdge1, aEdge0);"
208 EOL" const float4 aEdge2 = (1.f / dot (*theNormal, theRay->Direct)) * (thePoint0 - theRay->Origin);"
210 EOL" const float aTime = dot (*theNormal, aEdge2);"
212 EOL" const float4 theVec = cross (theRay->Direct, aEdge2);"
214 EOL" *theU = dot (theVec, aEdge1);"
215 EOL" *theV = dot (theVec, aEdge0);"
217 EOL" return (aTime >= 0.f) & (*theU >= 0.f) & (*theV >= 0.f) & (*theU + *theV <= 1.f) ? aTime : MAXFLOAT;"
220 /////////////////////////////////////////////////////////////////////////////////////////
221 // Support shading functions
223 EOL" const sampler_t EnvironmentSampler ="
224 EOL" CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_REPEAT | CLK_FILTER_LINEAR;"
226 // =======================================================================
227 // function : SmoothNormal
228 // purpose : Interpolates normal across the triangle
229 // =======================================================================
230 EOL" float4 SmoothNormal (__global float4* theNormals,"
231 EOL" const SIntersect* theHit,"
232 EOL" const int4 theIndices)"
234 EOL" float4 aNormal0 = theNormals[theIndices.x],"
235 EOL" aNormal1 = theNormals[theIndices.y],"
236 EOL" aNormal2 = theNormals[theIndices.z];"
238 EOL" return fast_normalize (aNormal1 * theHit->U +"
239 EOL" aNormal2 * theHit->V +"
240 EOL" aNormal0 * (1.f - theHit->U - theHit->V));"
243 // =======================================================================
245 // purpose : Computes Phong-based illumination
246 // =======================================================================
247 EOL" float4 Shade (const float4 theMatDiff,"
248 EOL" const float4 theMatSpec,"
249 EOL" const float4 theLight,"
250 EOL" const float4 theView,"
251 EOL" const float4 theNormal,"
252 EOL" const float4 theIntens,"
253 EOL" const float theTranspr)"
255 EOL" float aLambert = dot (theNormal, theLight);"
257 EOL" aLambert = theTranspr > 0.f ? fabs (aLambert) : aLambert;"
259 EOL" if (aLambert > 0.f)"
261 EOL" const float4 aReflect = 2.f * dot (theLight, theNormal) * theNormal - theLight;"
263 EOL" const float aSpecular = pow (max (dot (aReflect.xyz, theView.xyz), 0.f), theMatSpec.w);"
265 EOL" return theIntens * (theMatDiff * aLambert + theMatSpec * aSpecular);"
271 // =======================================================================
272 // function : Latlong
273 // purpose : Converts world direction to environment texture coordinates
274 // =======================================================================
275 EOL" float2 Latlong (const float4 thePoint, const float theRadius)"
277 EOL" float aPsi = acospi (-thePoint.y / theRadius);"
278 EOL" float aPhi = atan2pi (thePoint.z, thePoint.x);"
280 EOL" aPhi = (aPhi < 0.f) ? aPhi + 2.f : aPhi;"
282 EOL" return (float2) (aPhi * 0.5f, aPsi);"
285 /////////////////////////////////////////////////////////////////////////////////////////
286 // Core ray tracing function
288 // =======================================================================
290 // purpose : Pushes BVH node index to local stack
291 // =======================================================================
292 EOL" void push (uint* theStack, char* thePos, const uint theValue)"
295 EOL" theStack[*thePos] = theValue;"
298 // =======================================================================
300 // purpose : Pops BVH node index from local stack
301 // =======================================================================
302 EOL" void pop (uint* theStack, char* thePos, uint* theValue)"
304 EOL" *theValue = theStack[*thePos];"
308 // =======================================================================
309 // function : ObjectNearestHit
310 // purpose : Finds intersection with nearest object triangle
311 // =======================================================================
312 EOL" int4 ObjectNearestHit (const SRay* theRay,"
313 EOL" SIntersect* theIntersect,"
314 EOL" __global int4* theObjectNodeInfoBuffer,"
315 EOL" __global float4* theObjectMinPointBuffer,"
316 EOL" __global float4* theObjectMaxPointBuffer,"
317 EOL" __global int4* theGeometryTriangBuffer,"
318 EOL" __global float4* theGeometryVertexBuffer)"
320 EOL" uint aStack [32];"
322 EOL" char aHead = -1;" // stack pointer
323 EOL" uint aNode = 0;" // node to visit
325 EOL" const float4 aInvDirect = (float4) ("
326 EOL" 1.f / (fabs (theRay->Direct.x) > _OOEPS_ ?"
327 EOL" theRay->Direct.x : copysign (_OOEPS_, theRay->Direct.x)),"
328 EOL" 1.f / (fabs (theRay->Direct.y) > _OOEPS_ ?"
329 EOL" theRay->Direct.y : copysign (_OOEPS_, theRay->Direct.y)),"
330 EOL" 1.f / (fabs (theRay->Direct.z) > _OOEPS_ ?"
331 EOL" theRay->Direct.z : copysign (_OOEPS_, theRay->Direct.z)),"
334 EOL" int4 aTriangleIndex = (int4) (-1);"
336 EOL" float aTimeExit;"
337 EOL" float aTimeMin1;"
338 EOL" float aTimeMin2;"
342 EOL" const int3 aData = theObjectNodeInfoBuffer[aNode].xyz;"
344 EOL" if (aData.x == 0)" // if inner node
346 EOL" float4 aNodeMin = theObjectMinPointBuffer[aData.y];"
347 EOL" float4 aNodeMax = theObjectMaxPointBuffer[aData.y];"
349 EOL" float4 aTime0 = (aNodeMin - theRay->Origin) * aInvDirect;"
350 EOL" float4 aTime1 = (aNodeMax - theRay->Origin) * aInvDirect;"
352 EOL" float4 aTimeMax = max (aTime0, aTime1);"
353 EOL" float4 aTimeMin = min (aTime0, aTime1);"
355 EOL" aTimeExit = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
356 EOL" aTimeMin1 = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
358 EOL" const bool aHitLft = (aTimeMin1 <= aTimeExit) & (aTimeExit >= 0.f) & (aTimeMin1 <= theIntersect->Time);"
360 EOL" aNodeMin = theObjectMinPointBuffer[aData.z];"
361 EOL" aNodeMax = theObjectMaxPointBuffer[aData.z];"
363 EOL" aTime0 = (aNodeMin - theRay->Origin) * aInvDirect;"
364 EOL" aTime1 = (aNodeMax - theRay->Origin) * aInvDirect;"
366 EOL" aTimeMax = max (aTime0, aTime1);"
367 EOL" aTimeMin = min (aTime0, aTime1);"
369 EOL" aTimeExit = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
370 EOL" aTimeMin2 = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
372 EOL" const bool aHitRgh = (aTimeMin2 <= aTimeExit) & (aTimeExit >= 0.f) & (aTimeMin2 <= theIntersect->Time);"
374 EOL" if (aHitLft & aHitRgh)"
376 EOL" aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;"
378 EOL" push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);"
382 EOL" if (aHitLft | aHitRgh)"
384 EOL" aNode = aHitLft ? aData.y : aData.z;"
389 EOL" return aTriangleIndex;"
391 EOL" pop (aStack, &aHead, &aNode);"
395 EOL" else " // if leaf node
397 EOL" for (int aTriIdx = aData.y; aTriIdx <= aData.z; ++aTriIdx)"
399 EOL" const int4 aTestTriangle = theGeometryTriangBuffer[aTriIdx];"
401 EOL" const float4 aPoint0 = theGeometryVertexBuffer[aTestTriangle.x];"
402 EOL" const float4 aPoint1 = theGeometryVertexBuffer[aTestTriangle.y];"
403 EOL" const float4 aPoint2 = theGeometryVertexBuffer[aTestTriangle.z];"
405 EOL" float4 aNormal; float aU, aV;"
407 EOL" float aTime = IntersectTriangle (theRay,"
415 EOL" if (aTime < theIntersect->Time)"
417 EOL" aTriangleIndex = aTestTriangle;"
418 EOL" theIntersect->Normal = aNormal;"
419 EOL" theIntersect->Time = aTime;"
420 EOL" theIntersect->U = aU;"
421 EOL" theIntersect->V = aV;"
426 EOL" return aTriangleIndex;"
428 EOL" pop (aStack, &aHead, &aNode);"
432 EOL" return aTriangleIndex;"
435 // =======================================================================
436 // function : ObjectAnyHit
437 // purpose : Finds intersection with any object triangle
438 // =======================================================================
439 EOL" float ObjectAnyHit (const SRay* theRay,"
440 EOL" __global int4* theObjectNodeInfoBuffer,"
441 EOL" __global float4* theObjectMinPointBuffer,"
442 EOL" __global float4* theObjectMaxPointBuffer,"
443 EOL" __global int4* theGeometryTriangBuffer,"
444 EOL" __global float4* theGeometryVertexBuffer,"
445 EOL" const float theDistance)"
447 EOL" uint aStack [32];"
449 EOL" char aHead = -1;" // stack pointer
450 EOL" uint aNode = 0;" // node to visit
452 EOL" const float4 aInvDirect = (float4) ("
453 EOL" 1.f / (fabs (theRay->Direct.x) > _OOEPS_ ?"
454 EOL" theRay->Direct.x : copysign (_OOEPS_, theRay->Direct.x)),"
455 EOL" 1.f / (fabs (theRay->Direct.y) > _OOEPS_ ?"
456 EOL" theRay->Direct.y : copysign (_OOEPS_, theRay->Direct.y)),"
457 EOL" 1.f / (fabs (theRay->Direct.z) > _OOEPS_ ?"
458 EOL" theRay->Direct.z : copysign (_OOEPS_, theRay->Direct.z)),"
461 EOL" float aTimeExit;"
462 EOL" float aTimeMin1;"
463 EOL" float aTimeMin2;"
467 EOL" const int3 aData = theObjectNodeInfoBuffer[aNode].xyz;"
469 EOL" if (aData.x == 0)" // if inner node
471 EOL" float4 aNodeMin = theObjectMinPointBuffer[aData.y];"
472 EOL" float4 aNodeMax = theObjectMaxPointBuffer[aData.y];"
474 EOL" float4 aTime0 = (aNodeMin - theRay->Origin) * aInvDirect;"
475 EOL" float4 aTime1 = (aNodeMax - theRay->Origin) * aInvDirect;"
477 EOL" float4 aTimeMax = max (aTime0, aTime1);"
478 EOL" float4 aTimeMin = min (aTime0, aTime1);"
480 EOL" aTimeExit = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
481 EOL" aTimeMin1 = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
483 EOL" const bool aHitLft = (aTimeMin1 <= aTimeExit) & (aTimeExit >= 0.f) & (aTimeMin1 <= theDistance);"
485 EOL" aNodeMin = theObjectMinPointBuffer[aData.z];"
486 EOL" aNodeMax = theObjectMaxPointBuffer[aData.z];"
488 EOL" aTime0 = (aNodeMin - theRay->Origin) * aInvDirect;"
489 EOL" aTime1 = (aNodeMax - theRay->Origin) * aInvDirect;"
491 EOL" aTimeMax = max (aTime0, aTime1);"
492 EOL" aTimeMin = min (aTime0, aTime1);"
494 EOL" aTimeExit = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
495 EOL" aTimeMin2 = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
497 EOL" const bool aHitRgh = (aTimeMin2 <= aTimeExit) & (aTimeExit >= 0.f) & (aTimeMin2 <= theDistance);"
499 EOL" if (aHitLft & aHitRgh)"
501 EOL" aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;"
503 EOL" push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);"
507 EOL" if (aHitLft | aHitRgh)"
509 EOL" aNode = aHitLft ? aData.y : aData.z;"
516 EOL" pop (aStack, &aHead, &aNode);"
520 EOL" else " // if leaf node
522 EOL" for (int aTriIdx = aData.y; aTriIdx <= aData.z; ++aTriIdx)"
524 EOL" const int4 aTestTriangle = theGeometryTriangBuffer[aTriIdx];"
526 EOL" const float4 aPoint0 = theGeometryVertexBuffer[aTestTriangle.x];"
527 EOL" const float4 aPoint1 = theGeometryVertexBuffer[aTestTriangle.y];"
528 EOL" const float4 aPoint2 = theGeometryVertexBuffer[aTestTriangle.z];"
530 EOL" float4 aNormal; float aU, aV;"
532 EOL" float aTime = IntersectTriangle (theRay,"
540 EOL" if (aTime < theDistance)"
549 EOL" pop (aStack, &aHead, &aNode);"
556 // =======================================================================
557 // function : NearestHit
558 // purpose : Finds intersection with nearest scene triangle
559 // =======================================================================
560 EOL" int4 NearestHit (const SRay* theRay,"
561 EOL" SIntersect* theIntersect,"
562 EOL" __global int4* theSceneNodeInfoBuffer,"
563 EOL" __global float4* theSceneMinPointBuffer,"
564 EOL" __global float4* theSceneMaxPointBuffer,"
565 EOL" __global int4* theObjectNodeInfoBuffer,"
566 EOL" __global float4* theObjectMinPointBuffer,"
567 EOL" __global float4* theObjectMaxPointBuffer,"
568 EOL" __global int4* theGeometryTriangBuffer,"
569 EOL" __global float4* theGeometryVertexBuffer)"
571 EOL" theIntersect->Time = MAXFLOAT;"
573 EOL" uint aStack [16];"
575 EOL" char aHead = -1;" // stack pointer
576 EOL" uint aNode = 0;" // node to visit
578 EOL" int4 aNearestTriangle = (int4) (-1);"
582 EOL" const int4 aData = theSceneNodeInfoBuffer[aNode];"
584 EOL" if (aData.x != 0)" // if leaf node
586 EOL" const float4 aNodeMin = theSceneMinPointBuffer[aNode];"
587 EOL" const float4 aNodeMax = theSceneMaxPointBuffer[aNode];"
589 EOL" if (IntersectBox (theRay, aNodeMin, aNodeMax) <= theIntersect->Time)"
591 EOL" int4 anIndex = ObjectNearestHit (theRay,"
593 EOL" theObjectNodeInfoBuffer + aData.y,"
594 EOL" theObjectMinPointBuffer + aData.y,"
595 EOL" theObjectMaxPointBuffer + aData.y,"
596 EOL" theGeometryTriangBuffer + aData.w,"
597 EOL" theGeometryVertexBuffer + aData.z);"
599 EOL" if (anIndex.x != -1)"
600 EOL" aNearestTriangle = (int4) (anIndex.x + aData.z,"
601 EOL" anIndex.y + aData.z,"
602 EOL" anIndex.z + aData.z,"
607 EOL" return aNearestTriangle;"
609 EOL" pop (aStack, &aHead, &aNode);"
611 EOL" else " // if inner node
613 EOL" float4 aNodeMinLft = theSceneMinPointBuffer[aData.y];"
614 EOL" float4 aNodeMinRgh = theSceneMinPointBuffer[aData.z];"
615 EOL" float4 aNodeMaxLft = theSceneMaxPointBuffer[aData.y];"
616 EOL" float4 aNodeMaxRgh = theSceneMaxPointBuffer[aData.z];"
618 EOL" float aTimeMin1;"
619 EOL" float aTimeMin2;"
621 EOL" IntersectNodes (theRay,"
628 EOL" theIntersect->Time);"
630 EOL" const bool aHitLft = (aTimeMin1 != -MAXFLOAT);"
631 EOL" const bool aHitRgh = (aTimeMin2 != -MAXFLOAT);"
633 EOL" if (aHitLft & aHitRgh)"
635 EOL" aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;"
637 EOL" push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);"
641 EOL" if (aHitLft | aHitRgh)"
643 EOL" aNode = aHitLft ? aData.y : aData.z;"
648 EOL" return aNearestTriangle;"
650 EOL" pop (aStack, &aHead, &aNode);"
656 EOL" return aNearestTriangle;"
659 // =======================================================================
661 // purpose : Finds intersection with any scene triangle
662 // =======================================================================
663 EOL" float AnyHit (const SRay* theRay,"
664 EOL" __global int4* theSceneNodeInfoBuffer,"
665 EOL" __global float4* theSceneMinPointBuffer,"
666 EOL" __global float4* theSceneMaxPointBuffer,"
667 EOL" __global int4* theObjectNodeInfoBuffer,"
668 EOL" __global float4* theObjectMinPointBuffer,"
669 EOL" __global float4* theObjectMaxPointBuffer,"
670 EOL" __global int4* theGeometryTriangBuffer,"
671 EOL" __global float4* theGeometryVertexBuffer,"
672 EOL" const float theDistance)"
674 EOL" uint aStack [16];"
676 EOL" char aHead = -1;" // stack pointer
677 EOL" uint aNode = 0;" // node to visit
681 EOL" const int4 aData = theSceneNodeInfoBuffer[aNode];"
683 EOL" if (aData.x != 0)" // if leaf node
685 EOL" const float4 aNodeMin = theSceneMinPointBuffer[aNode];"
686 EOL" const float4 aNodeMax = theSceneMaxPointBuffer[aNode];"
688 EOL" if (IntersectBox (theRay, aNodeMin, aNodeMax) <= theDistance)"
690 EOL" if (0.f == ObjectAnyHit (theRay,"
691 EOL" theObjectNodeInfoBuffer + aData.y,"
692 EOL" theObjectMinPointBuffer + aData.y,"
693 EOL" theObjectMaxPointBuffer + aData.y,"
694 EOL" theGeometryTriangBuffer + aData.w,"
695 EOL" theGeometryVertexBuffer + aData.z,"
705 EOL" pop (aStack, &aHead, &aNode);"
707 EOL" else" // if inner node
709 EOL" float4 aNodeMinLft = theSceneMinPointBuffer[aData.y];"
710 EOL" float4 aNodeMinRgh = theSceneMinPointBuffer[aData.z];"
711 EOL" float4 aNodeMaxLft = theSceneMaxPointBuffer[aData.y];"
712 EOL" float4 aNodeMaxRgh = theSceneMaxPointBuffer[aData.z];"
714 EOL" float aTimeMin1;"
715 EOL" float aTimeMin2;"
717 EOL" IntersectNodes (theRay,"
726 EOL" const bool aHitLft = (aTimeMin1 != -MAXFLOAT);"
727 EOL" const bool aHitRgh = (aTimeMin2 != -MAXFLOAT);"
729 EOL" if (aHitLft & aHitRgh)"
731 EOL" aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;"
733 EOL" push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);"
737 EOL" if (aHitLft | aHitRgh)"
739 EOL" aNode = aHitLft ? aData.y : aData.z;"
746 EOL" pop (aStack, &aHead, &aNode);"
753 EOL" #define _MAX_DEPTH_ 5"
755 EOL" #define THRESHOLD (float4) (0.1f, 0.1f, 0.1f, 1.f)"
757 EOL" #define LIGHT_POS(Buffer, LightID) Buffer[2 * LightID + 1]"
758 EOL" #define LIGHT_RAD(Buffer, LightID) Buffer[2 * LightID + 0]"
760 EOL" #define MATERIAL_AMBN(Buffer, TriangleID) Buffer[7 * TriangleID.w + 0]"
761 EOL" #define MATERIAL_DIFF(Buffer, TriangleID) Buffer[7 * TriangleID.w + 1]"
762 EOL" #define MATERIAL_SPEC(Buffer, TriangleID) Buffer[7 * TriangleID.w + 2]"
763 EOL" #define MATERIAL_EMIS(Buffer, TriangleID) Buffer[7 * TriangleID.w + 3]"
764 EOL" #define MATERIAL_REFL(Buffer, TriangleID) Buffer[7 * TriangleID.w + 4]"
765 EOL" #define MATERIAL_REFR(Buffer, TriangleID) Buffer[7 * TriangleID.w + 5]"
766 EOL" #define MATERIAL_TRAN(Buffer, TriangleID) Buffer[7 * TriangleID.w + 6]"
768 // =======================================================================
769 // function : Radiance
770 // purpose : Computes color of specified ray
771 // =======================================================================
772 EOL" float4 Radiance (SRay* theRay,"
773 EOL" __read_only image2d_t theEnvMap,"
774 EOL" __global int4* theSceneNodeInfoBuffer,"
775 EOL" __global float4* theSceneMinPointBuffer,"
776 EOL" __global float4* theSceneMaxPointBuffer,"
777 EOL" __global int4* theObjectNodeInfoBuffer,"
778 EOL" __global float4* theObjectMinPointBuffer,"
779 EOL" __global float4* theObjectMaxPointBuffer,"
780 EOL" __global int4* theGeometryTriangBuffer,"
781 EOL" __global float4* theGeometryVertexBuffer,"
782 EOL" __global float4* theGeometryNormalBuffer,"
783 EOL" __global float4* theLightSourceBuffer,"
784 EOL" __global float4* theMaterialBuffer,"
785 EOL" const float4 theGlobalAmbient,"
786 EOL" const int theLightBufferSize,"
787 EOL" const int theShadowsEnabled,"
788 EOL" const int theReflectEnabled,"
789 EOL" const float theSceneEpsilon,"
790 EOL" const float theSceneRadius)"
792 EOL" float4 aResult = (float4) (0.f, 0.f, 0.f, 0.f);"
793 EOL" float4 aWeight = (float4) (1.f, 1.f, 1.f, 1.f);"
795 EOL" SIntersect aHit;"
797 EOL" for (int aDepth = 0; aDepth < _MAX_DEPTH_; ++aDepth)"
799 EOL" int4 aTriangle = NearestHit (theRay,"
801 EOL" theSceneNodeInfoBuffer,"
802 EOL" theSceneMinPointBuffer,"
803 EOL" theSceneMaxPointBuffer,"
804 EOL" theObjectNodeInfoBuffer,"
805 EOL" theObjectMinPointBuffer,"
806 EOL" theObjectMaxPointBuffer,"
807 EOL" theGeometryTriangBuffer,"
808 EOL" theGeometryVertexBuffer);"
810 EOL" if (aTriangle.x < 0.f)"
812 EOL" if (aWeight.w != 0.f)"
815 EOL" float aTime = IntersectSphere (theRay, theSceneRadius);"
817 EOL" if (aTime != MAXFLOAT)"
819 EOL" aResult += aWeight * read_imagef (theEnvMap, EnvironmentSampler,"
820 EOL" Latlong (theRay->Origin + theRay->Direct * aTime, theSceneRadius));"
823 EOL" return (float4) (aResult.x,"
829 EOL // Compute geometric normal
830 EOL" float4 aGeomNormal = aHit.Normal; aGeomNormal = fast_normalize (aGeomNormal);"
832 EOL // Compute interpolated normal
833 EOL" float4 aNormal = SmoothNormal (theGeometryNormalBuffer, &aHit, aTriangle);"
835 EOL // Compute intersection point
836 EOL" float4 aPoint = theRay->Direct * aHit.Time + theRay->Origin;"
838 EOL" float4 aMaterAmb = MATERIAL_AMBN (theMaterialBuffer, aTriangle);"
839 EOL" float4 aMaterTrn = MATERIAL_TRAN (theMaterialBuffer, aTriangle);"
841 EOL" aResult += aWeight * theGlobalAmbient * aMaterAmb *"
842 EOL" (aMaterTrn.x * max (fabs (dot (theRay->Direct, aNormal)), 0.5f));"
844 EOL" for (int nLight = 0; nLight < theLightBufferSize; ++nLight)"
846 EOL" float4 aLightPosition = LIGHT_POS (theLightSourceBuffer, nLight);"
849 EOL" aShadow.Direct = aLightPosition;"
851 EOL" float aLightDistance = MAXFLOAT;"
852 EOL" if (aLightPosition.w != 0.f)"
854 EOL" aLightDistance = length (aLightPosition - aPoint);"
855 EOL" aShadow.Direct = (aLightPosition - aPoint) * (1.f / aLightDistance);"
858 EOL" aShadow.Origin = aPoint + aShadow.Direct * theSceneEpsilon +"
859 EOL" aGeomNormal * copysign (theSceneEpsilon, dot (aGeomNormal, aShadow.Direct));"
861 EOL" float aVisibility = 1.f;"
863 EOL" if (theShadowsEnabled)"
865 EOL" aVisibility = AnyHit (&aShadow,"
866 EOL" theSceneNodeInfoBuffer,"
867 EOL" theSceneMinPointBuffer,"
868 EOL" theSceneMaxPointBuffer,"
869 EOL" theObjectNodeInfoBuffer,"
870 EOL" theObjectMinPointBuffer,"
871 EOL" theObjectMaxPointBuffer,"
872 EOL" theGeometryTriangBuffer,"
873 EOL" theGeometryVertexBuffer,"
874 EOL" aLightDistance);"
877 EOL" if (aVisibility > 0.f)"
879 EOL" aResult += aMaterTrn.x * aWeight * Shade (MATERIAL_DIFF (theMaterialBuffer, aTriangle),"
880 EOL" MATERIAL_SPEC (theMaterialBuffer, aTriangle),"
881 EOL" aShadow.Direct,"
882 EOL" -theRay->Direct,"
884 EOL" LIGHT_RAD (theLightSourceBuffer, nLight),"
889 EOL" if (aMaterTrn.y > 0.f)"
891 EOL" aWeight *= aMaterTrn.y;"
895 EOL" aWeight *= theReflectEnabled ? MATERIAL_REFL (theMaterialBuffer, aTriangle) : ZERO;"
897 EOL" float4 aDirect = theRay->Direct - 2.f * dot (theRay->Direct, aNormal) * aNormal;"
899 EOL" float aDdotN = dot (aDirect, aGeomNormal);"
900 EOL" if (aDdotN < 0.f)"
901 EOL" theRay->Direct -= 2.f * dot (theRay->Direct, aGeomNormal) * aGeomNormal;"
903 EOL" theRay->Direct = aDirect;"
906 EOL" if (all (islessequal (aWeight, THRESHOLD)))"
908 EOL" return (float4) (aResult.x,"
914 EOL" theRay->Origin = theRay->Direct * theSceneEpsilon + aPoint;"
917 EOL" return (float4) (aResult.x,"
923 ///////////////////////////////////////////////////////////////////////////////
924 // Ray tracing kernel functions
926 // =======================================================================
927 // function : RaytraceRender
928 // purpose : Computes pixel color using ray-tracing
929 // =======================================================================
930 EOL" __kernel void RaytraceRender (const int theSizeX,"
931 EOL" const int theSizeY,"
932 EOL" const float16 theOrigins,"
933 EOL" const float16 theDirects,"
934 EOL" __read_only image2d_t theEnvMap,"
935 EOL" __write_only image2d_t theOutput,"
936 EOL" __global int4* theSceneNodeInfoBuffer,"
937 EOL" __global float4* theSceneMinPointBuffer,"
938 EOL" __global float4* theSceneMaxPointBuffer,"
939 EOL" __global int4* theObjectNodeInfoBuffer,"
940 EOL" __global float4* theObjectMinPointBuffer,"
941 EOL" __global float4* theObjectMaxPointBuffer,"
942 EOL" __global int4* theGeometryTriangBuffer,"
943 EOL" __global float4* theGeometryVertexBuffer,"
944 EOL" __global float4* theGeometryNormalBuffer,"
945 EOL" __global float4* theLightSourceBuffer,"
946 EOL" __global float4* theMaterialBuffer,"
947 EOL" const float4 theGlobalAmbient,"
948 EOL" const int theLightBufferSize,"
949 EOL" const int theShadowsEnabled,"
950 EOL" const int theReflectEnabled,"
951 EOL" const float theSceneEpsilon,"
952 EOL" const float theSceneRadius)"
954 EOL" const int aPixelX = get_global_id (0);"
955 EOL" const int aPixelY = get_global_id (1);"
957 EOL" if (aPixelX >= theSizeX || aPixelY >= theSizeY)"
960 EOL" private SRay aRay;"
962 EOL" GenerateRay (&aRay,"
970 EOL" float4 aColor = (float4) (0.f, 0.f, 0.f, 1.f);"
972 EOL" float aTimeStart = IntersectBox (&aRay, theSceneMinPointBuffer[0], theSceneMaxPointBuffer[0]);"
974 EOL" if (aTimeStart != MAXFLOAT)"
976 EOL" aRay.Origin += aRay.Direct * max (aTimeStart - theSceneEpsilon, 0.f);"
978 EOL" aColor = clamp (Radiance (&aRay,"
980 EOL" theSceneNodeInfoBuffer,"
981 EOL" theSceneMinPointBuffer,"
982 EOL" theSceneMaxPointBuffer,"
983 EOL" theObjectNodeInfoBuffer,"
984 EOL" theObjectMinPointBuffer,"
985 EOL" theObjectMaxPointBuffer,"
986 EOL" theGeometryTriangBuffer,"
987 EOL" theGeometryVertexBuffer,"
988 EOL" theGeometryNormalBuffer,"
989 EOL" theLightSourceBuffer,"
990 EOL" theMaterialBuffer,"
991 EOL" theGlobalAmbient,"
992 EOL" theLightBufferSize,"
993 EOL" theShadowsEnabled,"
994 EOL" theReflectEnabled,"
995 EOL" theSceneEpsilon,"
996 EOL" theSceneRadius), 0.f, 1.f);"
999 EOL" write_imagef (theOutput, (int2) (aPixelX, aPixelY), aColor);"
1002 EOL" const sampler_t OutputSampler ="
1003 EOL" CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;"
1005 EOL" #define _LUM_DELTA_ 0.085f"
1007 EOL" #define AA_MAX 0.559017f"
1008 EOL" #define AA_MIN 0.186339f"
1010 // =======================================================================
1011 // function : RaytraceSmooth
1012 // purpose : Performs adaptive sub-pixel rendering
1013 // =======================================================================
1014 EOL" __kernel void RaytraceSmooth (const int theSizeX,"
1015 EOL" const int theSizeY,"
1016 EOL" const float16 theOrigins,"
1017 EOL" const float16 theDirects,"
1018 EOL" __read_only image2d_t theInput,"
1019 EOL" __read_only image2d_t theEnvMap,"
1020 EOL" __write_only image2d_t theOutput,"
1021 EOL" __global int4* theSceneNodeInfoBuffer,"
1022 EOL" __global float4* theSceneMinPointBuffer,"
1023 EOL" __global float4* theSceneMaxPointBuffer,"
1024 EOL" __global int4* theObjectNodeInfoBuffer,"
1025 EOL" __global float4* theObjectMinPointBuffer,"
1026 EOL" __global float4* theObjectMaxPointBuffer,"
1027 EOL" __global int4* theGeometryTriangBuffer,"
1028 EOL" __global float4* theGeometryVertexBuffer,"
1029 EOL" __global float4* theGeometryNormalBuffer,"
1030 EOL" __global float4* theLightSourceBuffer,"
1031 EOL" __global float4* theMaterialBuffer,"
1032 EOL" const float4 theGlobalAmbient,"
1033 EOL" const int theLightBufferSize,"
1034 EOL" const int theShadowsEnabled,"
1035 EOL" const int theReflectEnabled,"
1036 EOL" const float theSceneEpsilon,"
1037 EOL" const float theSceneRadius)"
1039 EOL" const int aPixelX = get_global_id (0);"
1040 EOL" const int aPixelY = get_global_id (1);"
1042 EOL" if (aPixelX >= theSizeX || aPixelY >= theSizeY)"
1045 EOL" float4 aClr0 = read_imagef (theInput, OutputSampler, (float2) (aPixelX + 0, aPixelY + 0));"
1046 EOL" float4 aClr1 = read_imagef (theInput, OutputSampler, (float2) (aPixelX + 0, aPixelY - 1));"
1047 EOL" float4 aClr2 = read_imagef (theInput, OutputSampler, (float2) (aPixelX + 0, aPixelY + 1));"
1049 EOL" float4 aClr3 = read_imagef (theInput, OutputSampler, (float2) (aPixelX + 1, aPixelY + 0));"
1050 EOL" float4 aClr4 = read_imagef (theInput, OutputSampler, (float2) (aPixelX + 1, aPixelY - 1));"
1051 EOL" float4 aClr5 = read_imagef (theInput, OutputSampler, (float2) (aPixelX + 1, aPixelY + 1));"
1053 EOL" float4 aClr6 = read_imagef (theInput, OutputSampler, (float2) (aPixelX - 1, aPixelY + 0));"
1054 EOL" float4 aClr7 = read_imagef (theInput, OutputSampler, (float2) (aPixelX - 1, aPixelY - 1));"
1055 EOL" float4 aClr8 = read_imagef (theInput, OutputSampler, (float2) (aPixelX - 1, aPixelY + 1));"
1057 EOL" bool render = fabs (aClr1.w - aClr0.w) > _LUM_DELTA_ ||"
1058 EOL" fabs (aClr2.w - aClr0.w) > _LUM_DELTA_ ||"
1059 EOL" fabs (aClr3.w - aClr0.w) > _LUM_DELTA_ ||"
1060 EOL" fabs (aClr4.w - aClr0.w) > _LUM_DELTA_ ||"
1061 EOL" fabs (aClr5.w - aClr0.w) > _LUM_DELTA_ ||"
1062 EOL" fabs (aClr6.w - aClr0.w) > _LUM_DELTA_ ||"
1063 EOL" fabs (aClr7.w - aClr0.w) > _LUM_DELTA_ ||"
1064 EOL" fabs (aClr8.w - aClr0.w) > _LUM_DELTA_;"
1068 EOL" aClr1 = (aClr1.w == 1.f) ? -UNIT : aClr1;"
1069 EOL" aClr2 = (aClr2.w == 1.f) ? -UNIT : aClr2;"
1070 EOL" aClr3 = (aClr3.w == 1.f) ? -UNIT : aClr3;"
1071 EOL" aClr4 = (aClr4.w == 1.f) ? -UNIT : aClr4;"
1072 EOL" aClr5 = (aClr5.w == 1.f) ? -UNIT : aClr5;"
1073 EOL" aClr6 = (aClr6.w == 1.f) ? -UNIT : aClr6;"
1074 EOL" aClr7 = (aClr7.w == 1.f) ? -UNIT : aClr7;"
1075 EOL" aClr8 = (aClr8.w == 1.f) ? -UNIT : aClr8;"
1077 EOL" float aLum = (aClr0.w == 1.f) ? -1.f : (0.2126f * aClr0.x + 0.7152f * aClr0.y + 0.0722f * aClr0.z);"
1079 EOL" render = fabs (0.2126f * aClr1.x + 0.7152f * aClr1.y + 0.0722f * aClr1.z - aLum) > _LUM_DELTA_ ||"
1080 EOL" fabs (0.2126f * aClr2.x + 0.7152f * aClr2.y + 0.0722f * aClr2.z - aLum) > _LUM_DELTA_ ||"
1081 EOL" fabs (0.2126f * aClr3.x + 0.7152f * aClr3.y + 0.0722f * aClr3.z - aLum) > _LUM_DELTA_ ||"
1082 EOL" fabs (0.2126f * aClr4.x + 0.7152f * aClr4.y + 0.0722f * aClr4.z - aLum) > _LUM_DELTA_ ||"
1083 EOL" fabs (0.2126f * aClr5.x + 0.7152f * aClr5.y + 0.0722f * aClr5.z - aLum) > _LUM_DELTA_ ||"
1084 EOL" fabs (0.2126f * aClr6.x + 0.7152f * aClr6.y + 0.0722f * aClr6.z - aLum) > _LUM_DELTA_ ||"
1085 EOL" fabs (0.2126f * aClr7.x + 0.7152f * aClr7.y + 0.0722f * aClr7.z - aLum) > _LUM_DELTA_ ||"
1086 EOL" fabs (0.2126f * aClr8.x + 0.7152f * aClr8.y + 0.0722f * aClr8.z - aLum) > _LUM_DELTA_;"
1089 EOL" float4 aColor = clamp (aClr0, 0.f, 1.f);"
1091 EOL" private SRay aRay;"
1093 EOL" const float4 aBoxMin = theSceneMinPointBuffer[0];"
1094 EOL" const float4 aBoxMax = theSceneMaxPointBuffer[0];"
1098 EOL" for (int aSample = 0; aSample <= 3; ++aSample)"
1100 EOL" float fX = aPixelX, fY = aPixelY;"
1102 EOL" if (aSample == 0)"
1104 EOL" fX -= AA_MIN; fY -= AA_MAX;"
1106 EOL" else if (aSample == 1)"
1108 EOL" fX -= AA_MAX; fY += AA_MIN;"
1110 EOL" else if (aSample == 2)"
1112 EOL" fX += AA_MIN; fY += AA_MAX;"
1116 EOL" fX += AA_MAX; fY -= AA_MIN;"
1119 EOL" GenerateRay (&aRay,"
1127 EOL" float aTimeStart = IntersectBox (&aRay, aBoxMin, aBoxMax);"
1129 EOL" if (aTimeStart != MAXFLOAT)"
1131 EOL" aRay.Origin += aRay.Direct * max (aTimeStart - theSceneEpsilon, 0.f);"
1133 EOL" aColor += clamp (Radiance (&aRay,"
1135 EOL" theSceneNodeInfoBuffer,"
1136 EOL" theSceneMinPointBuffer,"
1137 EOL" theSceneMaxPointBuffer,"
1138 EOL" theObjectNodeInfoBuffer,"
1139 EOL" theObjectMinPointBuffer,"
1140 EOL" theObjectMaxPointBuffer,"
1141 EOL" theGeometryTriangBuffer,"
1142 EOL" theGeometryVertexBuffer,"
1143 EOL" theGeometryNormalBuffer,"
1144 EOL" theLightSourceBuffer,"
1145 EOL" theMaterialBuffer,"
1146 EOL" theGlobalAmbient,"
1147 EOL" theLightBufferSize,"
1148 EOL" theShadowsEnabled,"
1149 EOL" theReflectEnabled,"
1150 EOL" theSceneEpsilon,"
1151 EOL" theSceneRadius), 0.f, 1.f);"
1154 EOL" aColor += (float4) (0.f, 0.f, 0.f, 1.f);"
1157 EOL" aColor *= 1.f / 5.f;"
1160 EOL" write_imagef (theOutput, (int2) (aPixelX, aPixelY), aColor);"