projects / occt.git / blame
? search:
summary | shortlog | log | commit | commitdiff | tree
blob | blame (incremental) | history | HEAD
0024288: Provide flipping text for AIS_Dimensions
[occt.git] / src / OpenGl / OpenGl_RaytraceSource.cxx
CommitLineData
e276548b 1// Created on: 2013-10-16
2// Created by: Denis BOGOLEPOV
3// Copyright (c) 2013 OPEN CASCADE SAS
4//
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.
9//
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.
12//
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.
19
20#ifdef HAVE_CONFIG_H
21 #include <config.h>
22#endif
23
24#ifdef HAVE_OPENCL
25
26#define EOL "\n"
27
28extern const char THE_RAY_TRACE_OPENCL_SOURCE[] =
29
30 /////////////////////////////////////////////////////////////////////////////////////////
31 // Specific data types
32 EOL
33 //! Stores ray parameters.
34 EOL" typedef struct __SRay"
35 EOL" {"
36 EOL" float4 Origin;"
37 EOL" float4 Direct;"
38 EOL" }"
39 EOL" SRay;"
40 EOL
41 //! Stores parameters of intersection point.
42 EOL" typedef struct __SIntersect"
43 EOL" {"
44 EOL" float4 Normal;"
45 EOL" float Time;"
46 EOL" float U;"
47 EOL" float V;"
48 EOL" }"
49 EOL" SIntersect;"
50 EOL
51 EOL
52 /////////////////////////////////////////////////////////////////////////////////////////
53 // Some useful vector constants
54 EOL
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 )"
57 EOL
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 )"
61 EOL
62 EOL
63 /////////////////////////////////////////////////////////////////////////////////////////
64 // Support functions
65 EOL
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)"
77 EOL" {"
78 EOL" float2 aPixel = (float2) (theX / (float)theSizeX,"
79 EOL" theY / (float)theSizeY);"
80 EOL
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);"
83 EOL
84 EOL" theRay->Origin = mix (aP0, aP1, aPixel.y);"
85 EOL
86 EOL" aP0 = mix (theDirects.lo.lo, theDirects.lo.hi, aPixel.x);"
87 EOL" aP1 = mix (theDirects.hi.lo, theDirects.hi.hi, aPixel.x);"
88 EOL
89 EOL" theRay->Direct = mix (aP0, aP1, aPixel.y);"
90 EOL" }"
91 EOL
92 EOL
93 /////////////////////////////////////////////////////////////////////////////////////////
94 // Functions for compute ray-object intersection
95 EOL
96 EOL" #define _OOEPS_ exp2( -80.0f )"
97 EOL
98 // =======================================================================
99 // function : IntersectSphere
100 // purpose : Computes ray-sphere intersection
101 // =======================================================================
102 EOL" bool IntersectSphere (const SRay* theRay, float theRadius, float* theTime)"
103 EOL" {"
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);"
107 EOL
108 EOL" float aD = aDdotO * aDdotO - aDdotD * (aOdotO - theRadius * theRadius);"
109 EOL
110 EOL" if (aD > 0.f)"
111 EOL" {"
112 EOL" *theTime = (-aDdotO + native_sqrt (aD)) * (1.f / aDdotD);"
113 EOL
114 EOL" return *theTime > 0.f;"
115 EOL" }"
116 EOL
117 EOL" return false;"
118 EOL" }"
119 EOL
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)"
129 EOL" {"
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)),"
137 EOL" 0.f);"
138 EOL
139 EOL" const float4 aTime0 = (theMinPoint - theRay->Origin) * aInvDirect;"
140 EOL" const float4 aTime1 = (theMaxPoint - theRay->Origin) * aInvDirect;"
141 EOL
142 EOL" const float4 aTimeMax = max (aTime0, aTime1);"
143 EOL" const float4 aTimeMin = min (aTime0, aTime1);"
144 EOL
145 EOL" *theTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
146 EOL" *theTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
147 EOL
148 EOL" return (*theTimeStart <= *theTimeFinal) & (*theTimeFinal >= 0.f);"
149 EOL" }"
150 EOL
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)"
163 EOL" {"
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)),"
171 EOL" 0.f);"
172 EOL
173 EOL" float4 aTime0 = (theMinPoint0 - theRay->Origin) * aInvDirect;"
174 EOL" float4 aTime1 = (theMaxPoint0 - theRay->Origin) * aInvDirect;"
175 EOL
176 EOL" float4 aTimeMax = max (aTime0, aTime1);"
177 EOL" float4 aTimeMin = min (aTime0, aTime1);"
178 EOL
179 EOL" aTime0 = (theMinPoint1 - theRay->Origin) * aInvDirect;"
180 EOL" aTime1 = (theMaxPoint1 - theRay->Origin) * aInvDirect;"
181 EOL
182 EOL" float aTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
183 EOL" float aTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
184 EOL
185 EOL" aTimeMax = max (aTime0, aTime1);"
186 EOL" aTimeMin = min (aTime0, aTime1);"
187 EOL
188 EOL" *theTimeStart0 = (aTimeStart <= aTimeFinal) & (aTimeFinal >= 0.f) & (aTimeStart <= theMaxTime)"
189 EOL" ? aTimeStart : -MAXFLOAT;"
190 EOL
191 EOL" aTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
192 EOL" aTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
193 EOL
194 EOL" *theTimeStart1 = (aTimeStart <= aTimeFinal) & (aTimeFinal >= 0.f) & (aTimeStart <= theMaxTime)"
195 EOL" ? aTimeStart : -MAXFLOAT;"
196 EOL" }"
197 EOL
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,"
208 EOL" float* theU,"
209 EOL" float* theV)"
210 EOL" {"
211 EOL" const float4 aEdge0 = thePoint1 - thePoint0;"
212 EOL" const float4 aEdge1 = thePoint0 - thePoint2;"
213 EOL
214 EOL" *theNormal = cross (aEdge1, aEdge0);"
215 EOL
216 EOL" const float4 aEdge2 = (1.f / dot (*theNormal, theRay->Direct)) * (thePoint0 - theRay->Origin);"
217 EOL
218 EOL" *theTime = dot (*theNormal, aEdge2);"
219 EOL
220 EOL" const float4 theInc = cross (theRay->Direct, aEdge2);"
221 EOL
222 EOL" *theU = dot (theInc, aEdge1);"
223 EOL" *theV = dot (theInc, aEdge0);"
224 EOL
225 EOL" return (*theTime > 0) & (*theU >= 0.f) & (*theV >= 0.f) & (*theU + *theV <= 1.f);"
226 EOL" }"
227 EOL
228 /////////////////////////////////////////////////////////////////////////////////////////
229 // Support shading functions
230 EOL
231 EOL" const sampler_t EnvironmentSampler ="
232 EOL" CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_REPEAT | CLK_FILTER_LINEAR;"
233 EOL
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)"
241 EOL" {"
242 EOL" float4 aNormal0 = theNormals[theIndices.x],"
243 EOL" aNormal1 = theNormals[theIndices.y],"
244 EOL" aNormal2 = theNormals[theIndices.z];"
245 EOL
246 EOL" return fast_normalize (aNormal1 * theHit->U +"
247 EOL" aNormal2 * theHit->V +"
248 EOL" aNormal0 * (1.f - theHit->U - theHit->V));"
249 EOL" }"
250 EOL
251 // =======================================================================
252 // function : Shade
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)"
262 EOL" {"
263 EOL" float aLambert = dot (theNormal, theLight);"
264 EOL
265 EOL" aLambert = theTranspr > 0.f ? fabs (aLambert) : aLambert;"
266 EOL
267 EOL" if (aLambert > 0.f)"
268 EOL" {"
269 EOL" const float4 aMatDiff = theMaterials[7 * theMatIndex + 1];"
270 EOL" const float4 aMatSpec = theMaterials[7 * theMatIndex + 2];"
271 EOL
272 EOL" const float4 aReflect = 2.f * dot (theLight, theNormal) * theNormal - theLight;"
273 EOL
274 EOL" const float aSpecular = pow (max (dot (aReflect.xyz, theView.xyz), 0.f), aMatSpec.w);"
275 EOL
276 EOL" return theIntens * (aMatDiff * aLambert + aMatSpec * aSpecular);"
277 EOL" }"
278 EOL
279 EOL" return ZERO;"
280 EOL" }"
281 EOL
282 // =======================================================================
283 // function : Lat-long
284 // purpose : Converts world direction to environment texture coordinates
285 // =======================================================================
286 EOL" float2 Latlong (const float4 theDirect)"
287 EOL" {"
288 EOL" float aPsi = acos( -theDirect.y );"
289 EOL" float aPhi = atan2( theDirect.z, theDirect.x );"
290 EOL
291 EOL" aPhi = (aPhi < 0) ? (aPhi + 2.f * M_PI_F) : aPhi;"
292 EOL
293 EOL" return (float2) (aPhi / (2.f * M_PI_F), aPsi / M_PI_F);"
294 EOL" }"
295 EOL
296 /////////////////////////////////////////////////////////////////////////////////////////
297 // Core ray tracing function
298 EOL
299 // =======================================================================
300 // function : push
301 // purpose : Pushes BVH node index to local stack
302 // =======================================================================
303 EOL" void push (uint* theStack, char* thePos, const uint theValue)"
304 EOL" {"
305 EOL" (*thePos)++;"
306 EOL" theStack[*thePos] = theValue;"
307 EOL" }"
308 EOL
309 // =======================================================================
310 // function : pop
311 // purpose : Pops BVH node index from local stack
312 // =======================================================================
313 EOL" void pop (uint* theStack, char* thePos, uint* theValue)"
314 EOL" {"
315 EOL" *theValue = theStack[*thePos];"
316 EOL" (*thePos)--;"
317 EOL" }"
318 EOL
319 // #define BVH_MINIMIZE_MEM_LOADS
320 EOL
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)"
332 EOL" {"
333 EOL" uint aStack [32];"
334 EOL" char aHead = -1;"
335 EOL
336 EOL" uint aNode = 0;" // root node
337 EOL
338 EOL" float aTimeMin1;"
339 EOL" float aTimeMin2;"
340 EOL
341 EOL" float4 aNodeMinLft;"
342 EOL" float4 aNodeMaxLft;"
343 EOL" float4 aNodeMinRgh;"
344 EOL" float4 aNodeMaxRgh;"
345 EOL
346 EOL" #ifdef BVH_MINIMIZE_MEM_LOADS"
347 EOL" aNodeMinLft = theNodeMinPoints[aNode];"
348 EOL" aNodeMaxLft = theNodeMaxPoints[aNode];"
349 EOL" #endif"
350 EOL
351 EOL" int4 aTriangleIndex = (int4) (-1);"
352 EOL
353 EOL" theHit->Time = MAXFLOAT;"
354 EOL
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));"
359 EOL
360 EOL" aData = aData.y < 0 ? -aData : aData;"
361 EOL" #endif"
362 EOL
363 EOL" while (true)"
364 EOL" {"
365 EOL" #ifndef BVH_MINIMIZE_MEM_LOADS"
366 EOL" int3 aData = theNodeDataRecords[aNode].xyz;"
367 EOL" #endif"
368 EOL
369 EOL" if (aData.x != 1)" // if inner node
370 EOL" {"
371 EOL" aNodeMinLft = theNodeMinPoints[aData.y];"
372 EOL" aNodeMinRgh = theNodeMinPoints[aData.z];"
373 EOL" aNodeMaxLft = theNodeMaxPoints[aData.y];"
374 EOL" aNodeMaxRgh = theNodeMaxPoints[aData.z];"
375 EOL
376 EOL" IntersectNodes (theRay,"
377 EOL" aNodeMinLft,"
378 EOL" aNodeMaxLft,"
379 EOL" aNodeMinRgh,"
380 EOL" aNodeMaxRgh,"
381 EOL" &aTimeMin1,"
382 EOL" &aTimeMin2,"
383 EOL" theHit->Time);"
384 EOL
385 EOL" const bool aHitLft = (aTimeMin1 != -MAXFLOAT);"
386 EOL" const bool aHitRgh = (aTimeMin2 != -MAXFLOAT);"
387 EOL
388 EOL" if (aHitLft & aHitRgh)"
389 EOL" {"
390 EOL" aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;"
391 EOL
392 EOL" push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);"
393 EOL
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));"
398 EOL
399 EOL" aData = aData.y < 0 ? -aData : aData;"
400 EOL" #endif"
401 EOL" }"
402 EOL" else"
403 EOL" {"
404 EOL" if (aHitLft | aHitRgh)"
405 EOL" {"
406 EOL" aNode = aHitLft ? aData.y : aData.z;"
407 EOL
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));"
412 EOL
413 EOL" aData = aData.y < 0 ? -aData : aData;"
414 EOL" #endif"
415 EOL" }"
416 EOL" else"
417 EOL" {"
418 EOL" if (aHead < 0)"
419 EOL" return aTriangleIndex;"
420 EOL
421 EOL" pop (aStack, &aHead, &aNode);"
422 EOL
423 EOL" #ifdef BVH_MINIMIZE_MEM_LOADS"
424 EOL" aData = theNodeDataRecords[aNode].xyz;"
425 EOL" #endif"
426 EOL" }"
427 EOL" }"
428 EOL" }"
429 EOL" else " // if leaf node
430 EOL" {"
431 EOL" for (int nTri = aData.y; nTri <= aData.z; ++nTri)"
432 EOL" {"
433 EOL" int4 anIndex = theIndices[nTri];"
434 EOL
435 EOL" const float4 aP0 = theVertices[anIndex.x];"
436 EOL" const float4 aP1 = theVertices[anIndex.y];"
437 EOL" const float4 aP2 = theVertices[anIndex.z];"
438 EOL
439 EOL" float4 aNormal;"
440 EOL
441 EOL" float aTime, aU, aV;"
442 EOL
443 EOL" if (IntersectTriangle (theRay, aP0, aP1, aP2, &aNormal, &aTime, &aU, &aV) & (aTime < theHit->Time))"
444 EOL" {"
445 EOL" aTriangleIndex = anIndex;"
446 EOL" theHit->Normal = aNormal;"
447 EOL" theHit->Time = aTime;"
448 EOL" theHit->U = aU;"
449 EOL" theHit->V = aV;"
450 EOL" }"
451 EOL" }"
452 EOL
453 EOL" if (aHead < 0)"
454 EOL" return aTriangleIndex;"
455 EOL
456 EOL" pop (aStack, &aHead, &aNode);"
457 EOL
458 EOL" #ifdef BVH_MINIMIZE_MEM_LOADS"
459 EOL" aData = theNodeDataRecords[aNode].xyz;"
460 EOL" #endif"
461 EOL" }"
462 EOL" }"
463 EOL
464 EOL" return aTriangleIndex;"
465 EOL" }"
466 EOL
467 EOL" #define TRANSPARENT_SHADOW_"
468 EOL
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)"
481 EOL" {"
482 EOL" uint aStack [32];"
483 EOL" char aHead = -1;"
484 EOL
485 EOL" uint aNode = 0;" // root node
486 EOL
487 EOL" float aFactor = 1.f;" // light attenuation factor
488 EOL
489 EOL" float aTimeMin1;"
490 EOL" float aTimeMin2;"
491 EOL
492 EOL" while (true)"
493 EOL" {"
494 EOL" int3 aData = theNodeDataRecords[aNode].xyz;"
495 EOL
496 EOL" if (aData.x != 1)" // if inner node
497 EOL" {"
498 EOL" IntersectNodes (theRay,"
499 EOL" theNodeMinPoints[aData.y],"
500 EOL" theNodeMaxPoints[aData.y],"
501 EOL" theNodeMinPoints[aData.z],"
502 EOL" theNodeMaxPoints[aData.z],"
503 EOL" &aTimeMin1,"
504 EOL" &aTimeMin2,"
505 EOL" theDistance);"
506 EOL
507 EOL" const bool aHitLft = (aTimeMin1 != -MAXFLOAT);"
508 EOL" const bool aHitRgh = (aTimeMin2 != -MAXFLOAT);"
509 EOL
510 EOL" if (aHitLft & aHitRgh)"
511 EOL" {"
512 EOL" aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;"
513 EOL
514 EOL" push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);"
515 EOL" }"
516 EOL" else"
517 EOL" {"
518 EOL" if (aHitLft | aHitRgh)"
519 EOL" {"
520 EOL" aNode = aHitLft ? aData.y : aData.z;"
521 EOL" }"
522 EOL" else"
523 EOL" {"
524 EOL" if (aHead < 0)"
525 EOL" return aFactor;"
526 EOL
527 EOL" pop (aStack, &aHead, &aNode);"
528 EOL" }"
529 EOL" }"
530 EOL" }"
531 EOL" else " // if leaf node
532 EOL" {"
533 EOL" for (int nTri = aData.y; nTri <= aData.z; ++nTri)"
534 EOL" {"
535 EOL" int4 anIndex = theIndices[nTri];"
536 EOL
537 EOL" const float4 aP0 = theVertices[anIndex.x];"
538 EOL" const float4 aP1 = theVertices[anIndex.y];"
539 EOL" const float4 aP2 = theVertices[anIndex.z];"
540 EOL
541 EOL" float4 aNormal;"
542 EOL
543 EOL" float aTime, aU, aV;"
544 EOL
545 EOL" if (IntersectTriangle (theRay, aP0, aP1, aP2, &aNormal, &aTime, &aU, &aV) & (aTime < theDistance))"
546 EOL" {"
547 EOL" #ifdef TRANSPARENT_SHADOW"
548 EOL" aFactor *= materials[7 * index.w + 6].x;"
549 EOL
550 EOL" if (aFactor < 0.1f)"
551 EOL" return aFactor;"
552 EOL" #else"
553 EOL" return 0.f;"
554 EOL" #endif"
555 EOL" }"
556 EOL" }"
557 EOL
558 EOL" if (aHead < 0)"
559 EOL" return aFactor;"
560 EOL
561 EOL" pop (aStack, &aHead, &aNode);"
562 EOL" }"
563 EOL" }"
564 EOL
565 EOL" return aFactor;"
566 EOL" }"
567 EOL
568 EOL" #define _MAX_DEPTH_ 5"
569 EOL
570 EOL" #define _MAT_SIZE_ 7"
571 EOL
572 EOL" #define _LGH_SIZE_ 3"
573 EOL
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)"
593 EOL" {"
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);"
596 EOL
597 EOL" SIntersect aHit;"
598 EOL
599 EOL" for (int aDepth = 0; aDepth < _MAX_DEPTH_; ++aDepth)"
600 EOL" {"
601 EOL" int4 aTriangle = Traverse (theRay,"
602 EOL" theIndices,"
603 EOL" theVertices,"
604 EOL" theNodeMinPoints,"
605 EOL" theNodeMaxPoints,"
606 EOL" theNodeDataRecords,"
607 EOL" &aHit);"
608 EOL
609 EOL" if (aTriangle.x < 0.f)"
610 EOL" {"
611 EOL" float aTime;"
612 EOL
613 EOL" if (aWeight.w != 0.f || !IntersectSphere (theRay, theRadius, &aTime))"
614 EOL" break;"
615 EOL
616 EOL" float2 aTexCoord = Latlong (fma (theRay->Direct, (float4) (aTime), theRay->Origin) * (1.f / theRadius));"
617 EOL
618 EOL" aResult += aWeight * read_imagef (theEnvMap, EnvironmentSampler, aTexCoord);"
619 EOL
620 EOL" return (float4) (aResult.x,"
621 EOL" aResult.y,"
622 EOL" aResult.z,"
623 EOL" aWeight.w);"
624 EOL" }"
625 EOL
626 EOL" " // Compute geometric normal
627 EOL" float4 aGeomNormal = aHit.Normal; aGeomNormal = fast_normalize (aGeomNormal);"
628 EOL
629 EOL" " // Compute interpolated normal
630 EOL" float4 aNormal = SmoothNormal (theNormals, &aHit, aTriangle);"
631 EOL
632 EOL" " // Compute intersection point
633 EOL" float4 aPoint = theRay->Direct * aHit.Time + theRay->Origin;"
634 EOL
635 EOL" float4 aMaterAmb = theMaterials [_MAT_SIZE_ * aTriangle.w + 0];"
636 EOL" float4 aMaterTrn = theMaterials [_MAT_SIZE_ * aTriangle.w + 6];"
637 EOL
638 EOL" for (int nLight = 0; nLight < theLightCount; ++nLight)"
639 EOL" {"
640 EOL" float4 aLightAmbient = theLightSources [_LGH_SIZE_ * nLight];"
641 EOL
642 EOL" aResult += aWeight * aLightAmbient * aMaterAmb *"
643 EOL" (aMaterTrn.x * max (fabs (dot (theRay->Direct, aNormal)), 0.5f));"
644 EOL
645 EOL" if (aLightAmbient.w < 0.f)" // 'ambient' light
646 EOL" {"
647 EOL" continue;" // 'ambient' light has no another luminances
648 EOL" }"
649 EOL
650 EOL" float4 aLightPosition = theLightSources [_LGH_SIZE_ * nLight + 2];"
651 EOL
652 EOL" SRay aShadow;"
653 EOL" aShadow.Direct = aLightPosition;"
654 EOL
655 EOL" float aLightDistance = MAXFLOAT;"
656 EOL" if (aLightPosition.w != 0.f)"
657 EOL" {"
658 EOL" aLightDistance = length (aLightPosition - aPoint);"
659 EOL" aShadow.Direct = (aLightPosition - aPoint) * (1.f / aLightDistance);"
660 EOL" }"
661 EOL
662 EOL" aShadow.Origin = aPoint + aShadow.Direct * theEpsilon +"
663 EOL" aGeomNormal * copysign (theEpsilon, dot (aGeomNormal, aShadow.Direct));"
664 EOL
665 EOL" float aFactor = 1.f;"
666 EOL
667 EOL" if (isShadows)"
668 EOL" {"
669 EOL" aFactor = TraverseShadow (&aShadow,"
670 EOL" theIndices,"
671 EOL" theVertices,"
672 EOL" theMaterials,"
673 EOL" theNodeMinPoints,"
674 EOL" theNodeMaxPoints,"
675 EOL" theNodeDataRecords,"
676 EOL" aLightDistance);"
677 EOL" }"
678 EOL
679 EOL" aResult += (aMaterTrn.x * aFactor) * aWeight * Shade (theMaterials,"
680 EOL" aShadow.Direct,"
681 EOL" -theRay->Direct,"
682 EOL" aNormal,"
683 EOL" UNIT,"
684 EOL" aMaterTrn.y,"
685 EOL" aTriangle.w);"
686 EOL" }"
687 EOL
688 EOL" if (aMaterTrn.y > 0.f)"
689 EOL" {"
690 EOL" aWeight *= aMaterTrn.y;"
691 EOL" }"
692 EOL" else"
693 EOL" {"
694 EOL" float4 aMaterRef = theMaterials [_MAT_SIZE_ * aTriangle.w + 4];"
695 EOL" aWeight *= isReflect ? aMaterRef : ZERO;"
696 EOL
697 EOL" theRay->Direct -= 2.f * dot (theRay->Direct, aNormal) * aNormal;"
698 EOL
699 EOL" float aDdotN = dot (theRay->Direct, aGeomNormal);"
700 EOL" if (aDdotN < 0.f)"
701 EOL" theRay->Direct -= aDdotN * aGeomNormal;"
702 EOL" }"
703 EOL
704 EOL" if (aWeight.x < 0.1f && aWeight.y < 0.1f && aWeight.z < 0.1f)"
705 EOL" {"
706 EOL" return (float4) (aResult.x,"
707 EOL" aResult.y,"
708 EOL" aResult.z,"
709 EOL" aWeight.w);"
710 EOL" }"
711 EOL
712 EOL" theRay->Origin = theRay->Direct * theEpsilon + aPoint;"
713 EOL" }"
714 EOL
715 EOL" return (float4) (aResult.x,"
716 EOL" aResult.y,"
717 EOL" aResult.z,"
718 EOL" aWeight.w);"
719 EOL" }"
720 EOL
721 EOL
722 ///////////////////////////////////////////////////////////////////////////////
723 // Ray tracing kernel functions
724 EOL
725 // =======================================================================
726 // function : Main
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)"
748 EOL" {"
749 EOL" const int aX = get_global_id (0);"
750 EOL" const int aY = get_global_id (1);"
751 EOL
752 EOL" if (aX >= theSizeX || aY >= theSizeY)"
753 EOL" return;"
754 EOL
755 EOL" private SRay aRay;"
756 EOL
757 EOL" GenerateRay (&aRay,"
758 EOL" aX,"
759 EOL" aY,"
760 EOL" theSizeX,"
761 EOL" theSizeY,"
762 EOL" theOrigins,"
763 EOL" theDirects);"
764 EOL
765 EOL" float4 aBoxMin = theNodeMinPoints[0] - (float4) (theEpsilon, theEpsilon, theEpsilon, 0.f);"
766 EOL" float4 aBoxMax = theNodeMaxPoints[0] + (float4) (theEpsilon, theEpsilon, theEpsilon, 0.f);"
767 EOL
768 EOL" float aTimeStart;"
769 EOL" float aTimeFinal;"
770 EOL
771 EOL" float4 aColor = (float4) (0.f, 0.f, 0.f, 1.f);"
772 EOL
773 EOL" if (IntersectBox (&aRay, aBoxMin, aBoxMax, &aTimeStart, &aTimeFinal))"
774 EOL" {"
775 EOL" aRay.Origin = fma (aRay.Direct, (float4) (aTimeStart), aRay.Origin);"
776 EOL
777 EOL" aColor = Raytrace (&aRay,"
778 EOL" theEnvMap,"
779 EOL" theNodeMinPoints,"
780 EOL" theNodeMaxPoints,"
781 EOL" theNodeDataRecords,"
782 EOL" theLightSources,"
783 EOL" theMaterials,"
784 EOL" theVertices,"
785 EOL" theNormals,"
786 EOL" theIndices,"
787 EOL" theLightCount,"
788 EOL" theEpsilon,"
789 EOL" theRadius,"
790 EOL" isShadows,"
791 EOL" isReflect);"
792 EOL" }"
793 EOL
794 EOL" write_imagef (theOutput, (int2) (aX, aY), aColor);"
795 EOL" }"
796 EOL
797 EOL" const sampler_t OutputSampler ="
798 EOL" CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;"
799 EOL
800 EOL" #define _LUM_DELTA_ 0.075f"
801 EOL
802 EOL" #define AA_MAX 0.559017f"
803 EOL" #define AA_MIN 0.186339f"
804 EOL
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 )"
829 EOL" {"
830 EOL" const int aX = get_global_id (0);"
831 EOL" const int aY = get_global_id (1);"
832 EOL
833 EOL" if (aX >= theSizeX || aY >= theSizeY)"
834 EOL" return;"
835 EOL
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));"
839 EOL
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));"
843 EOL
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));"
847 EOL
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;"
856 EOL
857 EOL" float aLum = (aClr0.w == 1.f) ? -1.f : (0.2126f * aClr0.x + 0.7152f * aClr0.y + 0.0722f * aClr0.z);"
858 EOL
859 EOL
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_;"
868 EOL
869 EOL" float4 aColor = aClr0;"
870 EOL
871 EOL" private SRay aRay;"
872 EOL
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);"
875 EOL
876 EOL" if (render)"
877 EOL" {"
878 EOL" for (int aSample = 0; aSample <= 3; ++aSample)"
879 EOL" {"
880 EOL" float fX = aX, fY = aY;"
881 EOL
882 EOL" if (aSample == 0)"
883 EOL" {"
884 EOL" fX -= AA_MIN; fY -= AA_MAX;"
885 EOL" }"
886 EOL" else if (aSample == 1)"
887 EOL" {"
888 EOL" fX -= AA_MAX; fY += AA_MIN;"
889 EOL" }"
890 EOL" else if (aSample == 2)"
891 EOL" {"
892 EOL" fX += AA_MIN; fY += AA_MAX;"
893 EOL" }"
894 EOL" else"
895 EOL" {"
896 EOL" fX += AA_MAX; fY -= AA_MIN;"
897 EOL" }"
898 EOL
899 EOL" GenerateRay (&aRay,"
900 EOL" fX,"
901 EOL" fY,"
902 EOL" theSizeX,"
903 EOL" theSizeY,"
904 EOL" theOrigins,"
905 EOL" theDirects);"
906 EOL
907 EOL" float aTimeStart;"
908 EOL" float aTimeFinal;"
909 EOL
910 EOL" if (IntersectBox (&aRay, aBoxMin, aBoxMax, &aTimeStart, &aTimeFinal))"
911 EOL" {"
912 EOL" aRay.Origin = fma (aRay.Direct, (float4) (aTimeStart), aRay.Origin);"
913 EOL
914 EOL" aColor += Raytrace (&aRay,"
915 EOL" theEnvMap,"
916 EOL" theNodeMinPoints,"
917 EOL" theNodeMaxPoints,"
918 EOL" theNodeDataRecords,"
919 EOL" theLightSources,"
920 EOL" theMaterials,"
921 EOL" theVertices,"
922 EOL" theNormals,"
923 EOL" theIndices,"
924 EOL" theLightCount,"
925 EOL" theEpsilon,"
926 EOL" theRadius,"
927 EOL" isShadows,"
928 EOL" isReflect);"
929 EOL" }"
930 EOL" else"
931 EOL" aColor += (float4) (0.f, 0.f, 0.f, 1.f);"
932 EOL" }"
933 EOL
934 EOL" aColor *= 1.f / 5.f;"
935 EOL" }"
936 EOL
937 EOL" write_imagef (theOutput, (int2) (aX, aY), aColor);"
938 EOL" }";
939
940#endif
Open CASCADE Technology repository