1 // Created on: 2013-10-16
2 // Created by: Denis BOGOLEPOV
3 // Copyright (c) 2013 OPEN CASCADE SAS
5 // The content of this file is subject to the Open CASCADE Technology Public
6 // License Version 6.5 (the "License"). You may not use the content of this file
7 // except in compliance with the License. Please obtain a copy of the License
8 // at http://www.opencascade.org and read it completely before using this file.
10 // The Initial Developer of the Original Code is Open CASCADE S.A.S., having its
11 // main offices at: 1, place des Freres Montgolfier, 78280 Guyancourt, France.
13 // The Original Code and all software distributed under the License is
14 // distributed on an "AS IS" basis, without warranty of any kind, and the
15 // Initial Developer hereby disclaims all such warranties, including without
16 // limitation, any warranties of merchantability, fitness for a particular
17 // purpose or non-infringement. Please see the License for the specific terms
18 // and conditions governing the rights and limitations under the License.
28 extern const char THE_RAY_TRACE_OPENCL_SOURCE[] =
30 /////////////////////////////////////////////////////////////////////////////////////////
31 // Specific data types
33 //! Stores ray parameters.
34 EOL" typedef struct __SRay"
41 //! Stores parameters of intersection point.
42 EOL" typedef struct __SIntersect"
52 /////////////////////////////////////////////////////////////////////////////////////////
53 // Some useful vector constants
55 EOL" #define ZERO ( float4 )( 0.f, 0.f, 0.f, 0.f )"
56 EOL" #define UNIT ( float4 )( 1.f, 1.f, 1.f, 0.f )"
58 EOL" #define AXIS_X ( float4 )( 1.f, 0.f, 0.f, 0.f )"
59 EOL" #define AXIS_Y ( float4 )( 0.f, 1.f, 0.f, 0.f )"
60 EOL" #define AXIS_Z ( float4 )( 0.f, 0.f, 1.f, 0.f )"
63 /////////////////////////////////////////////////////////////////////////////////////////
66 // =======================================================================
67 // function : GenerateRay
68 // purpose : Generates primary ray for current work item
69 // =======================================================================
70 EOL" void GenerateRay (SRay* theRay,"
71 EOL" const float theX,"
72 EOL" const float theY,"
73 EOL" const int theSizeX,"
74 EOL" const int theSizeY,"
75 EOL" const float16 theOrigins,"
76 EOL" const float16 theDirects)"
78 EOL" float2 aPixel = (float2) (theX / (float)theSizeX,"
79 EOL" theY / (float)theSizeY);"
81 EOL" float4 aP0 = mix (theOrigins.lo.lo, theOrigins.lo.hi, aPixel.x);"
82 EOL" float4 aP1 = mix (theOrigins.hi.lo, theOrigins.hi.hi, aPixel.x);"
84 EOL" theRay->Origin = mix (aP0, aP1, aPixel.y);"
86 EOL" aP0 = mix (theDirects.lo.lo, theDirects.lo.hi, aPixel.x);"
87 EOL" aP1 = mix (theDirects.hi.lo, theDirects.hi.hi, aPixel.x);"
89 EOL" theRay->Direct = mix (aP0, aP1, aPixel.y);"
93 /////////////////////////////////////////////////////////////////////////////////////////
94 // Functions for compute ray-object intersection
96 EOL" #define _OOEPS_ exp2( -80.0f )"
98 // =======================================================================
99 // function : IntersectSphere
100 // purpose : Computes ray-sphere intersection
101 // =======================================================================
102 EOL" bool IntersectSphere (const SRay* theRay, float theRadius, float* theTime)"
104 EOL" float aDdotD = dot (theRay->Direct.xyz, theRay->Direct.xyz);"
105 EOL" float aDdotO = dot (theRay->Direct.xyz, theRay->Origin.xyz);"
106 EOL" float aOdotO = dot (theRay->Origin.xyz, theRay->Origin.xyz);"
108 EOL" float aD = aDdotO * aDdotO - aDdotD * (aOdotO - theRadius * theRadius);"
112 EOL" *theTime = (-aDdotO + native_sqrt (aD)) * (1.f / aDdotD);"
114 EOL" return *theTime > 0.f;"
120 // =======================================================================
121 // function : IntersectBox
122 // purpose : Computes ray-box intersection (slab test)
123 // =======================================================================
124 EOL" bool IntersectBox (const SRay* theRay,"
125 EOL" float4 theMinPoint,"
126 EOL" float4 theMaxPoint,"
127 EOL" float* theTimeStart,"
128 EOL" float* theTimeFinal)"
130 EOL" const float4 aInvDirect = (float4)("
131 EOL" 1.f / (fabs (theRay->Direct.x) > _OOEPS_ ?"
132 EOL" theRay->Direct.x : copysign (_OOEPS_, theRay->Direct.x)),"
133 EOL" 1.f / (fabs (theRay->Direct.y) > _OOEPS_ ?"
134 EOL" theRay->Direct.y : copysign (_OOEPS_, theRay->Direct.y)),"
135 EOL" 1.f / (fabs (theRay->Direct.z) > _OOEPS_ ?"
136 EOL" theRay->Direct.z : copysign (_OOEPS_, theRay->Direct.z)),"
139 EOL" const float4 aTime0 = (theMinPoint - theRay->Origin) * aInvDirect;"
140 EOL" const float4 aTime1 = (theMaxPoint - theRay->Origin) * aInvDirect;"
142 EOL" const float4 aTimeMax = max (aTime0, aTime1);"
143 EOL" const float4 aTimeMin = min (aTime0, aTime1);"
145 EOL" *theTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
146 EOL" *theTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
148 EOL" return (*theTimeStart <= *theTimeFinal) & (*theTimeFinal >= 0.f);"
151 // =======================================================================
152 // function : IntersectNodes
153 // purpose : Computes intersection of ray with two child nodes (boxes)
154 // =======================================================================
155 EOL" void IntersectNodes (const SRay* theRay,"
156 EOL" float4 theMinPoint0,"
157 EOL" float4 theMaxPoint0,"
158 EOL" float4 theMinPoint1,"
159 EOL" float4 theMaxPoint1,"
160 EOL" float* theTimeStart0,"
161 EOL" float* theTimeStart1,"
162 EOL" float theMaxTime)"
164 EOL" const float4 aInvDirect = (float4)("
165 EOL" 1.f / (fabs (theRay->Direct.x) > _OOEPS_ ?"
166 EOL" theRay->Direct.x : copysign (_OOEPS_, theRay->Direct.x)),"
167 EOL" 1.f / (fabs (theRay->Direct.y) > _OOEPS_ ?"
168 EOL" theRay->Direct.y : copysign (_OOEPS_, theRay->Direct.y)),"
169 EOL" 1.f / (fabs (theRay->Direct.z) > _OOEPS_ ?"
170 EOL" theRay->Direct.z : copysign (_OOEPS_, theRay->Direct.z)),"
173 EOL" float4 aTime0 = (theMinPoint0 - theRay->Origin) * aInvDirect;"
174 EOL" float4 aTime1 = (theMaxPoint0 - theRay->Origin) * aInvDirect;"
176 EOL" float4 aTimeMax = max (aTime0, aTime1);"
177 EOL" float4 aTimeMin = min (aTime0, aTime1);"
179 EOL" aTime0 = (theMinPoint1 - theRay->Origin) * aInvDirect;"
180 EOL" aTime1 = (theMaxPoint1 - theRay->Origin) * aInvDirect;"
182 EOL" float aTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
183 EOL" float aTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
185 EOL" aTimeMax = max (aTime0, aTime1);"
186 EOL" aTimeMin = min (aTime0, aTime1);"
188 EOL" *theTimeStart0 = (aTimeStart <= aTimeFinal) & (aTimeFinal >= 0.f) & (aTimeStart <= theMaxTime)"
189 EOL" ? aTimeStart : -MAXFLOAT;"
191 EOL" aTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
192 EOL" aTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
194 EOL" *theTimeStart1 = (aTimeStart <= aTimeFinal) & (aTimeFinal >= 0.f) & (aTimeStart <= theMaxTime)"
195 EOL" ? aTimeStart : -MAXFLOAT;"
198 // =======================================================================
199 // function : IntersectTriangle
200 // purpose : Computes ray-triangle intersection (branchless version)
201 // =======================================================================
202 EOL" bool IntersectTriangle (const SRay* theRay,"
203 EOL" const float4 thePoint0,"
204 EOL" const float4 thePoint1,"
205 EOL" const float4 thePoint2,"
206 EOL" float4* theNormal,"
207 EOL" float* theTime,"
211 EOL" const float4 aEdge0 = thePoint1 - thePoint0;"
212 EOL" const float4 aEdge1 = thePoint0 - thePoint2;"
214 EOL" *theNormal = cross (aEdge1, aEdge0);"
216 EOL" const float4 aEdge2 = (1.f / dot (*theNormal, theRay->Direct)) * (thePoint0 - theRay->Origin);"
218 EOL" *theTime = dot (*theNormal, aEdge2);"
220 EOL" const float4 theInc = cross (theRay->Direct, aEdge2);"
222 EOL" *theU = dot (theInc, aEdge1);"
223 EOL" *theV = dot (theInc, aEdge0);"
225 EOL" return (*theTime > 0) & (*theU >= 0.f) & (*theV >= 0.f) & (*theU + *theV <= 1.f);"
228 /////////////////////////////////////////////////////////////////////////////////////////
229 // Support shading functions
231 EOL" const sampler_t EnvironmentSampler ="
232 EOL" CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_REPEAT | CLK_FILTER_LINEAR;"
234 // =======================================================================
235 // function : SmoothNormal
236 // purpose : Interpolates normal across the triangle
237 // =======================================================================
238 EOL" float4 SmoothNormal (__global float4* theNormals,"
239 EOL" const SIntersect* theHit,"
240 EOL" const int4 theIndices)"
242 EOL" float4 aNormal0 = theNormals[theIndices.x],"
243 EOL" aNormal1 = theNormals[theIndices.y],"
244 EOL" aNormal2 = theNormals[theIndices.z];"
246 EOL" return fast_normalize (aNormal1 * theHit->U +"
247 EOL" aNormal2 * theHit->V +"
248 EOL" aNormal0 * (1.f - theHit->U - theHit->V));"
251 // =======================================================================
253 // purpose : Computes Phong-based illumination
254 // =======================================================================
255 EOL" float4 Shade (__global float4* theMaterials,"
256 EOL" const float4 theLight,"
257 EOL" const float4 theView,"
258 EOL" const float4 theNormal,"
259 EOL" const float4 theIntens,"
260 EOL" const float theTranspr,"
261 EOL" const int theMatIndex)"
263 EOL" float aLambert = dot (theNormal, theLight);"
265 EOL" aLambert = theTranspr > 0.f ? fabs (aLambert) : aLambert;"
267 EOL" if (aLambert > 0.f)"
269 EOL" const float4 aMatDiff = theMaterials[7 * theMatIndex + 1];"
270 EOL" const float4 aMatSpec = theMaterials[7 * theMatIndex + 2];"
272 EOL" const float4 aReflect = 2.f * dot (theLight, theNormal) * theNormal - theLight;"
274 EOL" const float aSpecular = pow (max (dot (aReflect.xyz, theView.xyz), 0.f), aMatSpec.w);"
276 EOL" return theIntens * (aMatDiff * aLambert + aMatSpec * aSpecular);"
282 // =======================================================================
283 // function : Lat-long
284 // purpose : Converts world direction to environment texture coordinates
285 // =======================================================================
286 EOL" float2 Latlong (const float4 theDirect)"
288 EOL" float aPsi = acos( -theDirect.y );"
289 EOL" float aPhi = atan2( theDirect.z, theDirect.x );"
291 EOL" aPhi = (aPhi < 0) ? (aPhi + 2.f * M_PI_F) : aPhi;"
293 EOL" return (float2) (aPhi / (2.f * M_PI_F), aPsi / M_PI_F);"
296 /////////////////////////////////////////////////////////////////////////////////////////
297 // Core ray tracing function
299 // =======================================================================
301 // purpose : Pushes BVH node index to local stack
302 // =======================================================================
303 EOL" void push (uint* theStack, char* thePos, const uint theValue)"
306 EOL" theStack[*thePos] = theValue;"
309 // =======================================================================
311 // purpose : Pops BVH node index from local stack
312 // =======================================================================
313 EOL" void pop (uint* theStack, char* thePos, uint* theValue)"
315 EOL" *theValue = theStack[*thePos];"
319 // #define BVH_MINIMIZE_MEM_LOADS
321 // =======================================================================
322 // function : Traverse
323 // purpose : Finds intersection with nearest triangle
324 // =======================================================================
325 EOL" int4 Traverse (const SRay* theRay,"
326 EOL" __global int4* theIndices,"
327 EOL" __global float4* theVertices,"
328 EOL" __global float4* theNodeMinPoints,"
329 EOL" __global float4* theNodeMaxPoints,"
330 EOL" __global int4* theNodeDataRecords,"
331 EOL" SIntersect* theHit)"
333 EOL" uint aStack [32];"
334 EOL" char aHead = -1;"
336 EOL" uint aNode = 0;" // root node
338 EOL" float aTimeMin1;"
339 EOL" float aTimeMin2;"
341 EOL" float4 aNodeMinLft;"
342 EOL" float4 aNodeMaxLft;"
343 EOL" float4 aNodeMinRgh;"
344 EOL" float4 aNodeMaxRgh;"
346 EOL" #ifdef BVH_MINIMIZE_MEM_LOADS"
347 EOL" aNodeMinLft = theNodeMinPoints[aNode];"
348 EOL" aNodeMaxLft = theNodeMaxPoints[aNode];"
351 EOL" int4 aTriangleIndex = (int4) (-1);"
353 EOL" theHit->Time = MAXFLOAT;"
355 EOL" #ifdef BVH_MINIMIZE_MEM_LOADS"
356 EOL" int3 aData = (int3) (1,"
357 EOL" as_int (aNodeMinLft.w),"
358 EOL" as_int (aNodeMaxLft.w));"
360 EOL" aData = aData.y < 0 ? -aData : aData;"
365 EOL" #ifndef BVH_MINIMIZE_MEM_LOADS"
366 EOL" int3 aData = theNodeDataRecords[aNode].xyz;"
369 EOL" if (aData.x != 1)" // if inner node
371 EOL" aNodeMinLft = theNodeMinPoints[aData.y];"
372 EOL" aNodeMinRgh = theNodeMinPoints[aData.z];"
373 EOL" aNodeMaxLft = theNodeMaxPoints[aData.y];"
374 EOL" aNodeMaxRgh = theNodeMaxPoints[aData.z];"
376 EOL" IntersectNodes (theRay,"
385 EOL" const bool aHitLft = (aTimeMin1 != -MAXFLOAT);"
386 EOL" const bool aHitRgh = (aTimeMin2 != -MAXFLOAT);"
388 EOL" if (aHitLft & aHitRgh)"
390 EOL" aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;"
392 EOL" push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);"
394 EOL" #ifdef BVH_MINIMIZE_MEM_LOADS"
395 EOL" aData = (int3) (1,"
396 EOL" as_int (aTimeMin1 < aTimeMin2 ? aNodeMinLft.w : aNodeMinRgh.w),"
397 EOL" as_int (aTimeMin1 < aTimeMin2 ? aNodeMaxLft.w : aNodeMaxRgh.w));"
399 EOL" aData = aData.y < 0 ? -aData : aData;"
404 EOL" if (aHitLft | aHitRgh)"
406 EOL" aNode = aHitLft ? aData.y : aData.z;"
408 EOL" #ifdef BVH_MINIMIZE_MEM_LOADS"
409 EOL" aData = (int3) (1,"
410 EOL" as_int (aHitLft ? aNodeMinLft.w : aNodeMinRgh.w),"
411 EOL" as_int (aHitLft ? aNodeMaxLft.w : aNodeMaxRgh.w));"
413 EOL" aData = aData.y < 0 ? -aData : aData;"
419 EOL" return aTriangleIndex;"
421 EOL" pop (aStack, &aHead, &aNode);"
423 EOL" #ifdef BVH_MINIMIZE_MEM_LOADS"
424 EOL" aData = theNodeDataRecords[aNode].xyz;"
429 EOL" else " // if leaf node
431 EOL" for (int nTri = aData.y; nTri <= aData.z; ++nTri)"
433 EOL" int4 anIndex = theIndices[nTri];"
435 EOL" const float4 aP0 = theVertices[anIndex.x];"
436 EOL" const float4 aP1 = theVertices[anIndex.y];"
437 EOL" const float4 aP2 = theVertices[anIndex.z];"
439 EOL" float4 aNormal;"
441 EOL" float aTime, aU, aV;"
443 EOL" if (IntersectTriangle (theRay, aP0, aP1, aP2, &aNormal, &aTime, &aU, &aV) & (aTime < theHit->Time))"
445 EOL" aTriangleIndex = anIndex;"
446 EOL" theHit->Normal = aNormal;"
447 EOL" theHit->Time = aTime;"
448 EOL" theHit->U = aU;"
449 EOL" theHit->V = aV;"
454 EOL" return aTriangleIndex;"
456 EOL" pop (aStack, &aHead, &aNode);"
458 EOL" #ifdef BVH_MINIMIZE_MEM_LOADS"
459 EOL" aData = theNodeDataRecords[aNode].xyz;"
464 EOL" return aTriangleIndex;"
467 EOL" #define TRANSPARENT_SHADOW_"
469 // =======================================================================
470 // function : TraverseShadow
471 // purpose : Finds intersection with any triangle
472 // =======================================================================
473 EOL" float TraverseShadow (const SRay* theRay,"
474 EOL" __global int4* theIndices,"
475 EOL" __global float4* theVertices,"
476 EOL" __global float4* materials,"
477 EOL" __global float4* theNodeMinPoints,"
478 EOL" __global float4* theNodeMaxPoints,"
479 EOL" __global int4* theNodeDataRecords,"
480 EOL" float theDistance)"
482 EOL" uint aStack [32];"
483 EOL" char aHead = -1;"
485 EOL" uint aNode = 0;" // root node
487 EOL" float aFactor = 1.f;" // light attenuation factor
489 EOL" float aTimeMin1;"
490 EOL" float aTimeMin2;"
494 EOL" int3 aData = theNodeDataRecords[aNode].xyz;"
496 EOL" if (aData.x != 1)" // if inner node
498 EOL" IntersectNodes (theRay,"
499 EOL" theNodeMinPoints[aData.y],"
500 EOL" theNodeMaxPoints[aData.y],"
501 EOL" theNodeMinPoints[aData.z],"
502 EOL" theNodeMaxPoints[aData.z],"
507 EOL" const bool aHitLft = (aTimeMin1 != -MAXFLOAT);"
508 EOL" const bool aHitRgh = (aTimeMin2 != -MAXFLOAT);"
510 EOL" if (aHitLft & aHitRgh)"
512 EOL" aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;"
514 EOL" push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);"
518 EOL" if (aHitLft | aHitRgh)"
520 EOL" aNode = aHitLft ? aData.y : aData.z;"
525 EOL" return aFactor;"
527 EOL" pop (aStack, &aHead, &aNode);"
531 EOL" else " // if leaf node
533 EOL" for (int nTri = aData.y; nTri <= aData.z; ++nTri)"
535 EOL" int4 anIndex = theIndices[nTri];"
537 EOL" const float4 aP0 = theVertices[anIndex.x];"
538 EOL" const float4 aP1 = theVertices[anIndex.y];"
539 EOL" const float4 aP2 = theVertices[anIndex.z];"
541 EOL" float4 aNormal;"
543 EOL" float aTime, aU, aV;"
545 EOL" if (IntersectTriangle (theRay, aP0, aP1, aP2, &aNormal, &aTime, &aU, &aV) & (aTime < theDistance))"
547 EOL" #ifdef TRANSPARENT_SHADOW"
548 EOL" aFactor *= materials[7 * index.w + 6].x;"
550 EOL" if (aFactor < 0.1f)"
551 EOL" return aFactor;"
559 EOL" return aFactor;"
561 EOL" pop (aStack, &aHead, &aNode);"
565 EOL" return aFactor;"
568 EOL" #define _MAX_DEPTH_ 5"
570 EOL" #define _MAT_SIZE_ 7"
572 EOL" #define _LGH_SIZE_ 3"
574 // =======================================================================
575 // function : Raytrace
576 // purpose : Computes color of specified ray
577 // =======================================================================
578 EOL" float4 Raytrace (SRay* theRay,"
579 EOL" __read_only image2d_t theEnvMap,"
580 EOL" __global float4* theNodeMinPoints,"
581 EOL" __global float4* theNodeMaxPoints,"
582 EOL" __global int4* theNodeDataRecords,"
583 EOL" __global float4* theLightSources,"
584 EOL" __global float4* theMaterials,"
585 EOL" __global float4* theVertices,"
586 EOL" __global float4* theNormals,"
587 EOL" __global int4* theIndices,"
588 EOL" const int theLightCount,"
589 EOL" const float theEpsilon,"
590 EOL" const float theRadius,"
591 EOL" const int isShadows,"
592 EOL" const int isReflect)"
594 EOL" float4 aResult = (float4) (0.f, 0.f, 0.f, 0.f);"
595 EOL" float4 aWeight = (float4) (1.f, 1.f, 1.f, 1.f);"
597 EOL" SIntersect aHit;"
599 EOL" for (int aDepth = 0; aDepth < _MAX_DEPTH_; ++aDepth)"
601 EOL" int4 aTriangle = Traverse (theRay,"
604 EOL" theNodeMinPoints,"
605 EOL" theNodeMaxPoints,"
606 EOL" theNodeDataRecords,"
609 EOL" if (aTriangle.x < 0.f)"
613 EOL" if (aWeight.w != 0.f || !IntersectSphere (theRay, theRadius, &aTime))"
616 EOL" float2 aTexCoord = Latlong (fma (theRay->Direct, (float4) (aTime), theRay->Origin) * (1.f / theRadius));"
618 EOL" aResult += aWeight * read_imagef (theEnvMap, EnvironmentSampler, aTexCoord);"
620 EOL" return (float4) (aResult.x,"
626 EOL" " // Compute geometric normal
627 EOL" float4 aGeomNormal = aHit.Normal; aGeomNormal = fast_normalize (aGeomNormal);"
629 EOL" " // Compute interpolated normal
630 EOL" float4 aNormal = SmoothNormal (theNormals, &aHit, aTriangle);"
632 EOL" " // Compute intersection point
633 EOL" float4 aPoint = theRay->Direct * aHit.Time + theRay->Origin;"
635 EOL" float4 aMaterAmb = theMaterials [_MAT_SIZE_ * aTriangle.w + 0];"
636 EOL" float4 aMaterTrn = theMaterials [_MAT_SIZE_ * aTriangle.w + 6];"
638 EOL" for (int nLight = 0; nLight < theLightCount; ++nLight)"
640 EOL" float4 aLightAmbient = theLightSources [_LGH_SIZE_ * nLight];"
642 EOL" aResult += aWeight * aLightAmbient * aMaterAmb *"
643 EOL" (aMaterTrn.x * max (fabs (dot (theRay->Direct, aNormal)), 0.5f));"
645 EOL" if (aLightAmbient.w < 0.f)" // 'ambient' light
647 EOL" continue;" // 'ambient' light has no another luminances
650 EOL" float4 aLightPosition = theLightSources [_LGH_SIZE_ * nLight + 2];"
653 EOL" aShadow.Direct = aLightPosition;"
655 EOL" float aLightDistance = MAXFLOAT;"
656 EOL" if (aLightPosition.w != 0.f)"
658 EOL" aLightDistance = length (aLightPosition - aPoint);"
659 EOL" aShadow.Direct = (aLightPosition - aPoint) * (1.f / aLightDistance);"
662 EOL" aShadow.Origin = aPoint + aShadow.Direct * theEpsilon +"
663 EOL" aGeomNormal * copysign (theEpsilon, dot (aGeomNormal, aShadow.Direct));"
665 EOL" float aFactor = 1.f;"
669 EOL" aFactor = TraverseShadow (&aShadow,"
673 EOL" theNodeMinPoints,"
674 EOL" theNodeMaxPoints,"
675 EOL" theNodeDataRecords,"
676 EOL" aLightDistance);"
679 EOL" aResult += (aMaterTrn.x * aFactor) * aWeight * Shade (theMaterials,"
680 EOL" aShadow.Direct,"
681 EOL" -theRay->Direct,"
688 EOL" if (aMaterTrn.y > 0.f)"
690 EOL" aWeight *= aMaterTrn.y;"
694 EOL" float4 aMaterRef = theMaterials [_MAT_SIZE_ * aTriangle.w + 4];"
695 EOL" aWeight *= isReflect ? aMaterRef : ZERO;"
697 EOL" theRay->Direct -= 2.f * dot (theRay->Direct, aNormal) * aNormal;"
699 EOL" float aDdotN = dot (theRay->Direct, aGeomNormal);"
700 EOL" if (aDdotN < 0.f)"
701 EOL" theRay->Direct -= aDdotN * aGeomNormal;"
704 EOL" if (aWeight.x < 0.1f && aWeight.y < 0.1f && aWeight.z < 0.1f)"
706 EOL" return (float4) (aResult.x,"
712 EOL" theRay->Origin = theRay->Direct * theEpsilon + aPoint;"
715 EOL" return (float4) (aResult.x,"
722 ///////////////////////////////////////////////////////////////////////////////
723 // Ray tracing kernel functions
725 // =======================================================================
727 // purpose : Computes pixel color using ray-tracing
728 // =======================================================================
729 EOL" __kernel void Main (__write_only image2d_t theOutput,"
730 EOL" __read_only image2d_t theEnvMap,"
731 EOL" __global float4* theNodeMinPoints,"
732 EOL" __global float4* theNodeMaxPoints,"
733 EOL" __global int4* theNodeDataRecords,"
734 EOL" __global float4* theLightSources,"
735 EOL" __global float4* theMaterials,"
736 EOL" __global float4* theVertices,"
737 EOL" __global float4* theNormals,"
738 EOL" __global int4* theIndices,"
739 EOL" const float16 theOrigins,"
740 EOL" const float16 theDirects,"
741 EOL" const int theLightCount,"
742 EOL" const float theEpsilon,"
743 EOL" const float theRadius,"
744 EOL" const int isShadows,"
745 EOL" const int isReflect,"
746 EOL" const int theSizeX,"
747 EOL" const int theSizeY)"
749 EOL" const int aX = get_global_id (0);"
750 EOL" const int aY = get_global_id (1);"
752 EOL" if (aX >= theSizeX || aY >= theSizeY)"
755 EOL" private SRay aRay;"
757 EOL" GenerateRay (&aRay,"
765 EOL" float4 aBoxMin = theNodeMinPoints[0] - (float4) (theEpsilon, theEpsilon, theEpsilon, 0.f);"
766 EOL" float4 aBoxMax = theNodeMaxPoints[0] + (float4) (theEpsilon, theEpsilon, theEpsilon, 0.f);"
768 EOL" float aTimeStart;"
769 EOL" float aTimeFinal;"
771 EOL" float4 aColor = (float4) (0.f, 0.f, 0.f, 1.f);"
773 EOL" if (IntersectBox (&aRay, aBoxMin, aBoxMax, &aTimeStart, &aTimeFinal))"
775 EOL" aRay.Origin = fma (aRay.Direct, (float4) (aTimeStart), aRay.Origin);"
777 EOL" aColor = Raytrace (&aRay,"
779 EOL" theNodeMinPoints,"
780 EOL" theNodeMaxPoints,"
781 EOL" theNodeDataRecords,"
782 EOL" theLightSources,"
794 EOL" write_imagef (theOutput, (int2) (aX, aY), aColor);"
797 EOL" const sampler_t OutputSampler ="
798 EOL" CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;"
800 EOL" #define _LUM_DELTA_ 0.075f"
802 EOL" #define AA_MAX 0.559017f"
803 EOL" #define AA_MIN 0.186339f"
805 // =======================================================================
806 // function : MainAntialiased
807 // purpose : Performs adaptive sub-pixel rendering
808 // =======================================================================
809 EOL" __kernel void MainAntialiased ( __read_only image2d_t theInput,"
810 EOL" __write_only image2d_t theOutput,"
811 EOL" __read_only image2d_t theEnvMap,"
812 EOL" __global float4* theNodeMinPoints,"
813 EOL" __global float4* theNodeMaxPoints,"
814 EOL" __global int4* theNodeDataRecords,"
815 EOL" __global float4* theLightSources,"
816 EOL" __global float4* theMaterials,"
817 EOL" __global float4* theVertices,"
818 EOL" __global float4* theNormals,"
819 EOL" __global int4* theIndices,"
820 EOL" const float16 theOrigins,"
821 EOL" const float16 theDirects,"
822 EOL" const int theLightCount,"
823 EOL" const float theEpsilon,"
824 EOL" const float theRadius,"
825 EOL" const int isShadows,"
826 EOL" const int isReflect,"
827 EOL" const int theSizeX,"
828 EOL" const int theSizeY )"
830 EOL" const int aX = get_global_id (0);"
831 EOL" const int aY = get_global_id (1);"
833 EOL" if (aX >= theSizeX || aY >= theSizeY)"
836 EOL" float4 aClr0 = read_imagef (theInput, OutputSampler, (float2) (aX + 0, aY + 0));"
837 EOL" float4 aClr1 = read_imagef (theInput, OutputSampler, (float2) (aX + 0, aY - 1));"
838 EOL" float4 aClr2 = read_imagef (theInput, OutputSampler, (float2) (aX + 0, aY + 1));"
840 EOL" float4 aClr3 = read_imagef (theInput, OutputSampler, (float2) (aX + 1, aY + 0));"
841 EOL" float4 aClr4 = read_imagef (theInput, OutputSampler, (float2) (aX + 1, aY - 1));"
842 EOL" float4 aClr5 = read_imagef (theInput, OutputSampler, (float2) (aX + 1, aY + 1));"
844 EOL" float4 aClr6 = read_imagef (theInput, OutputSampler, (float2) (aX - 1, aY + 0));"
845 EOL" float4 aClr7 = read_imagef (theInput, OutputSampler, (float2) (aX - 1, aY - 1));"
846 EOL" float4 aClr8 = read_imagef (theInput, OutputSampler, (float2) (aX - 1, aY + 1));"
848 EOL" aClr1 = (aClr1.w == 1.f) ? -UNIT : aClr1;"
849 EOL" aClr2 = (aClr2.w == 1.f) ? -UNIT : aClr2;"
850 EOL" aClr3 = (aClr3.w == 1.f) ? -UNIT : aClr3;"
851 EOL" aClr4 = (aClr4.w == 1.f) ? -UNIT : aClr4;"
852 EOL" aClr5 = (aClr5.w == 1.f) ? -UNIT : aClr5;"
853 EOL" aClr6 = (aClr6.w == 1.f) ? -UNIT : aClr6;"
854 EOL" aClr7 = (aClr7.w == 1.f) ? -UNIT : aClr7;"
855 EOL" aClr8 = (aClr8.w == 1.f) ? -UNIT : aClr8;"
857 EOL" float aLum = (aClr0.w == 1.f) ? -1.f : (0.2126f * aClr0.x + 0.7152f * aClr0.y + 0.0722f * aClr0.z);"
860 EOL" bool render = fabs (0.2126f * aClr1.x + 0.7152f * aClr1.y + 0.0722f * aClr1.z - aLum) > _LUM_DELTA_ ||"
861 EOL" fabs (0.2126f * aClr2.x + 0.7152f * aClr2.y + 0.0722f * aClr2.z - aLum) > _LUM_DELTA_ ||"
862 EOL" fabs (0.2126f * aClr3.x + 0.7152f * aClr3.y + 0.0722f * aClr3.z - aLum) > _LUM_DELTA_ ||"
863 EOL" fabs (0.2126f * aClr4.x + 0.7152f * aClr4.y + 0.0722f * aClr4.z - aLum) > _LUM_DELTA_ ||"
864 EOL" fabs (0.2126f * aClr5.x + 0.7152f * aClr5.y + 0.0722f * aClr5.z - aLum) > _LUM_DELTA_ ||"
865 EOL" fabs (0.2126f * aClr6.x + 0.7152f * aClr6.y + 0.0722f * aClr6.z - aLum) > _LUM_DELTA_ ||"
866 EOL" fabs (0.2126f * aClr7.x + 0.7152f * aClr7.y + 0.0722f * aClr7.z - aLum) > _LUM_DELTA_ ||"
867 EOL" fabs (0.2126f * aClr8.x + 0.7152f * aClr8.y + 0.0722f * aClr8.z - aLum) > _LUM_DELTA_;"
869 EOL" float4 aColor = aClr0;"
871 EOL" private SRay aRay;"
873 EOL" const float4 aBoxMin = theNodeMinPoints[0] - (float4) (theEpsilon, theEpsilon, theEpsilon, 0.f);"
874 EOL" const float4 aBoxMax = theNodeMaxPoints[0] + (float4) (theEpsilon, theEpsilon, theEpsilon, 0.f);"
878 EOL" for (int aSample = 0; aSample <= 3; ++aSample)"
880 EOL" float fX = aX, fY = aY;"
882 EOL" if (aSample == 0)"
884 EOL" fX -= AA_MIN; fY -= AA_MAX;"
886 EOL" else if (aSample == 1)"
888 EOL" fX -= AA_MAX; fY += AA_MIN;"
890 EOL" else if (aSample == 2)"
892 EOL" fX += AA_MIN; fY += AA_MAX;"
896 EOL" fX += AA_MAX; fY -= AA_MIN;"
899 EOL" GenerateRay (&aRay,"
907 EOL" float aTimeStart;"
908 EOL" float aTimeFinal;"
910 EOL" if (IntersectBox (&aRay, aBoxMin, aBoxMax, &aTimeStart, &aTimeFinal))"
912 EOL" aRay.Origin = fma (aRay.Direct, (float4) (aTimeStart), aRay.Origin);"
914 EOL" aColor += Raytrace (&aRay,"
916 EOL" theNodeMinPoints,"
917 EOL" theNodeMaxPoints,"
918 EOL" theNodeDataRecords,"
919 EOL" theLightSources,"
931 EOL" aColor += (float4) (0.f, 0.f, 0.f, 1.f);"
934 EOL" aColor *= 1.f / 5.f;"
937 EOL" write_imagef (theOutput, (int2) (aX, aY), aColor);"