0024778: Convertation of the generic classes to the non-generic. Part 9
[occt.git] / src / OpenGl / OpenGl_RaytraceSource.cxx
CommitLineData
e276548b 1// Created on: 2013-10-16
2// Created by: Denis BOGOLEPOV
d5f74e42 3// Copyright (c) 2013-2014 OPEN CASCADE SAS
e276548b 4//
973c2be1 5// This file is part of Open CASCADE Technology software library.
e276548b 6//
d5f74e42 7// This library is free software; you can redistribute it and/or modify it under
8// the terms of the GNU Lesser General Public License version 2.1 as published
973c2be1 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.
e276548b 12//
973c2be1 13// Alternatively, this file may be used under the terms of Open CASCADE
14// commercial license or contractual agreement.
e276548b 15
16#ifdef HAVE_CONFIG_H
17 #include <config.h>
18#endif
19
20#ifdef HAVE_OPENCL
21
22#define EOL "\n"
23
24extern const char THE_RAY_TRACE_OPENCL_SOURCE[] =
25
26 /////////////////////////////////////////////////////////////////////////////////////////
27 // Specific data types
28 EOL
29 //! Stores ray parameters.
30 EOL" typedef struct __SRay"
31 EOL" {"
32 EOL" float4 Origin;"
33 EOL" float4 Direct;"
34 EOL" }"
35 EOL" SRay;"
36 EOL
37 //! Stores parameters of intersection point.
38 EOL" typedef struct __SIntersect"
39 EOL" {"
40 EOL" float4 Normal;"
41 EOL" float Time;"
42 EOL" float U;"
43 EOL" float V;"
44 EOL" }"
45 EOL" SIntersect;"
46 EOL
47 EOL
48 /////////////////////////////////////////////////////////////////////////////////////////
49 // Some useful vector constants
50 EOL
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 )"
53 EOL
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 )"
57 EOL
e276548b 58 /////////////////////////////////////////////////////////////////////////////////////////
59 // Support functions
60 EOL
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)"
72 EOL" {"
73 EOL" float2 aPixel = (float2) (theX / (float)theSizeX,"
74 EOL" theY / (float)theSizeY);"
75 EOL
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);"
78 EOL
79 EOL" theRay->Origin = mix (aP0, aP1, aPixel.y);"
80 EOL
81 EOL" aP0 = mix (theDirects.lo.lo, theDirects.lo.hi, aPixel.x);"
82 EOL" aP1 = mix (theDirects.hi.lo, theDirects.hi.hi, aPixel.x);"
83 EOL
84 EOL" theRay->Direct = mix (aP0, aP1, aPixel.y);"
85 EOL" }"
86 EOL
87 EOL
88 /////////////////////////////////////////////////////////////////////////////////////////
89 // Functions for compute ray-object intersection
90 EOL
265d4508 91 EOL" #define _OOEPS_ exp2 (-80.0f)"
e276548b 92 EOL
93 // =======================================================================
94 // function : IntersectSphere
95 // purpose : Computes ray-sphere intersection
96 // =======================================================================
265d4508 97 EOL" float IntersectSphere (const SRay* theRay, float theRadius)"
e276548b 98 EOL" {"
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);"
102 EOL
103 EOL" float aD = aDdotO * aDdotO - aDdotD * (aOdotO - theRadius * theRadius);"
104 EOL
105 EOL" if (aD > 0.f)"
106 EOL" {"
265d4508 107 EOL" float aTime = (-aDdotO + native_sqrt (aD)) * (1.f / aDdotD);"
e276548b 108 EOL
265d4508 109 EOL" return aTime > 0.f ? aTime : MAXFLOAT;"
e276548b 110 EOL" }"
111 EOL
265d4508 112 EOL" return MAXFLOAT;"
e276548b 113 EOL" }"
114 EOL
115 // =======================================================================
116 // function : IntersectBox
117 // purpose : Computes ray-box intersection (slab test)
118 // =======================================================================
265d4508 119 EOL" float IntersectBox (const SRay* theRay,"
120 EOL" float4 theMinPoint,"
121 EOL" float4 theMaxPoint)"
e276548b 122 EOL" {"
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)),"
130 EOL" 0.f);"
131 EOL
132 EOL" const float4 aTime0 = (theMinPoint - theRay->Origin) * aInvDirect;"
133 EOL" const float4 aTime1 = (theMaxPoint - theRay->Origin) * aInvDirect;"
134 EOL
135 EOL" const float4 aTimeMax = max (aTime0, aTime1);"
136 EOL" const float4 aTimeMin = min (aTime0, aTime1);"
137 EOL
265d4508 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));"
e276548b 140 EOL
265d4508 141 EOL" return (theTimeStart <= theTimeFinal) && (theTimeFinal >= 0.f) ? theTimeStart : MAXFLOAT;"
e276548b 142 EOL" }"
143 EOL
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)"
156 EOL" {"
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)),"
164 EOL" 0.f);"
165 EOL
166 EOL" float4 aTime0 = (theMinPoint0 - theRay->Origin) * aInvDirect;"
167 EOL" float4 aTime1 = (theMaxPoint0 - theRay->Origin) * aInvDirect;"
168 EOL
169 EOL" float4 aTimeMax = max (aTime0, aTime1);"
170 EOL" float4 aTimeMin = min (aTime0, aTime1);"
171 EOL
172 EOL" aTime0 = (theMinPoint1 - theRay->Origin) * aInvDirect;"
173 EOL" aTime1 = (theMaxPoint1 - theRay->Origin) * aInvDirect;"
174 EOL
175 EOL" float aTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
176 EOL" float aTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
177 EOL
178 EOL" aTimeMax = max (aTime0, aTime1);"
179 EOL" aTimeMin = min (aTime0, aTime1);"
180 EOL
181 EOL" *theTimeStart0 = (aTimeStart <= aTimeFinal) & (aTimeFinal >= 0.f) & (aTimeStart <= theMaxTime)"
182 EOL" ? aTimeStart : -MAXFLOAT;"
183 EOL
184 EOL" aTimeFinal = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
185 EOL" aTimeStart = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
186 EOL
187 EOL" *theTimeStart1 = (aTimeStart <= aTimeFinal) & (aTimeFinal >= 0.f) & (aTimeStart <= theMaxTime)"
188 EOL" ? aTimeStart : -MAXFLOAT;"
189 EOL" }"
190 EOL
191 // =======================================================================
192 // function : IntersectTriangle
193 // purpose : Computes ray-triangle intersection (branchless version)
194 // =======================================================================
265d4508 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,"
200 EOL" float* theU,"
201 EOL" float* theV)"
e276548b 202 EOL" {"
203 EOL" const float4 aEdge0 = thePoint1 - thePoint0;"
204 EOL" const float4 aEdge1 = thePoint0 - thePoint2;"
205 EOL
206 EOL" *theNormal = cross (aEdge1, aEdge0);"
207 EOL
208 EOL" const float4 aEdge2 = (1.f / dot (*theNormal, theRay->Direct)) * (thePoint0 - theRay->Origin);"
209 EOL
265d4508 210 EOL" const float aTime = dot (*theNormal, aEdge2);"
e276548b 211 EOL
265d4508 212 EOL" const float4 theVec = cross (theRay->Direct, aEdge2);"
e276548b 213 EOL
265d4508 214 EOL" *theU = dot (theVec, aEdge1);"
215 EOL" *theV = dot (theVec, aEdge0);"
e276548b 216 EOL
265d4508 217 EOL" return (aTime >= 0.f) & (*theU >= 0.f) & (*theV >= 0.f) & (*theU + *theV <= 1.f) ? aTime : MAXFLOAT;"
e276548b 218 EOL" }"
219 EOL
220 /////////////////////////////////////////////////////////////////////////////////////////
221 // Support shading functions
222 EOL
223 EOL" const sampler_t EnvironmentSampler ="
224 EOL" CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_REPEAT | CLK_FILTER_LINEAR;"
225 EOL
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)"
233 EOL" {"
234 EOL" float4 aNormal0 = theNormals[theIndices.x],"
235 EOL" aNormal1 = theNormals[theIndices.y],"
236 EOL" aNormal2 = theNormals[theIndices.z];"
237 EOL
238 EOL" return fast_normalize (aNormal1 * theHit->U +"
239 EOL" aNormal2 * theHit->V +"
240 EOL" aNormal0 * (1.f - theHit->U - theHit->V));"
241 EOL" }"
242 EOL
243 // =======================================================================
244 // function : Shade
245 // purpose : Computes Phong-based illumination
246 // =======================================================================
265d4508 247 EOL" float4 Shade (const float4 theMatDiff,"
248 EOL" const float4 theMatSpec,"
e276548b 249 EOL" const float4 theLight,"
250 EOL" const float4 theView,"
251 EOL" const float4 theNormal,"
252 EOL" const float4 theIntens,"
265d4508 253 EOL" const float theTranspr)"
e276548b 254 EOL" {"
255 EOL" float aLambert = dot (theNormal, theLight);"
256 EOL
257 EOL" aLambert = theTranspr > 0.f ? fabs (aLambert) : aLambert;"
258 EOL
259 EOL" if (aLambert > 0.f)"
260 EOL" {"
e276548b 261 EOL" const float4 aReflect = 2.f * dot (theLight, theNormal) * theNormal - theLight;"
262 EOL
265d4508 263 EOL" const float aSpecular = pow (max (dot (aReflect.xyz, theView.xyz), 0.f), theMatSpec.w);"
e276548b 264 EOL
265d4508 265 EOL" return theIntens * (theMatDiff * aLambert + theMatSpec * aSpecular);"
e276548b 266 EOL" }"
267 EOL
268 EOL" return ZERO;"
269 EOL" }"
270 EOL
271 // =======================================================================
265d4508 272 // function : Latlong
e276548b 273 // purpose : Converts world direction to environment texture coordinates
274 // =======================================================================
265d4508 275 EOL" float2 Latlong (const float4 thePoint, const float theRadius)"
e276548b 276 EOL" {"
265d4508 277 EOL" float aPsi = acospi (-thePoint.y / theRadius);"
278 EOL" float aPhi = atan2pi (thePoint.z, thePoint.x);"
e276548b 279 EOL
265d4508 280 EOL" aPhi = (aPhi < 0.f) ? aPhi + 2.f : aPhi;"
e276548b 281 EOL
265d4508 282 EOL" return (float2) (aPhi * 0.5f, aPsi);"
e276548b 283 EOL" }"
284 EOL
285 /////////////////////////////////////////////////////////////////////////////////////////
286 // Core ray tracing function
287 EOL
288 // =======================================================================
289 // function : push
290 // purpose : Pushes BVH node index to local stack
291 // =======================================================================
292 EOL" void push (uint* theStack, char* thePos, const uint theValue)"
293 EOL" {"
294 EOL" (*thePos)++;"
295 EOL" theStack[*thePos] = theValue;"
296 EOL" }"
297 EOL
298 // =======================================================================
299 // function : pop
300 // purpose : Pops BVH node index from local stack
301 // =======================================================================
302 EOL" void pop (uint* theStack, char* thePos, uint* theValue)"
303 EOL" {"
304 EOL" *theValue = theStack[*thePos];"
305 EOL" (*thePos)--;"
306 EOL" }"
307 EOL
e276548b 308 // =======================================================================
265d4508 309 // function : ObjectNearestHit
310 // purpose : Finds intersection with nearest object triangle
e276548b 311 // =======================================================================
265d4508 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)"
e276548b 319 EOL" {"
320 EOL" uint aStack [32];"
e276548b 321 EOL
265d4508 322 EOL" char aHead = -1;" // stack pointer
323 EOL" uint aNode = 0;" // node to visit
324 EOL
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)),"
332 EOL" 0.f);"
333 EOL
334 EOL" int4 aTriangleIndex = (int4) (-1);"
e276548b 335 EOL
265d4508 336 EOL" float aTimeExit;"
e276548b 337 EOL" float aTimeMin1;"
338 EOL" float aTimeMin2;"
339 EOL
265d4508 340 EOL" while (true)"
341 EOL" {"
342 EOL" const int3 aData = theObjectNodeInfoBuffer[aNode].xyz;"
e276548b 343 EOL
265d4508 344 EOL" if (aData.x == 0)" // if inner node
345 EOL" {"
346 EOL" float4 aNodeMin = theObjectMinPointBuffer[aData.y];"
347 EOL" float4 aNodeMax = theObjectMaxPointBuffer[aData.y];"
e276548b 348 EOL
265d4508 349 EOL" float4 aTime0 = (aNodeMin - theRay->Origin) * aInvDirect;"
350 EOL" float4 aTime1 = (aNodeMax - theRay->Origin) * aInvDirect;"
e276548b 351 EOL
265d4508 352 EOL" float4 aTimeMax = max (aTime0, aTime1);"
353 EOL" float4 aTimeMin = min (aTime0, aTime1);"
e276548b 354 EOL
265d4508 355 EOL" aTimeExit = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
356 EOL" aTimeMin1 = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
e276548b 357 EOL
265d4508 358 EOL" const bool aHitLft = (aTimeMin1 <= aTimeExit) & (aTimeExit >= 0.f) & (aTimeMin1 <= theIntersect->Time);"
e276548b 359 EOL
265d4508 360 EOL" aNodeMin = theObjectMinPointBuffer[aData.z];"
361 EOL" aNodeMax = theObjectMaxPointBuffer[aData.z];"
e276548b 362 EOL
265d4508 363 EOL" aTime0 = (aNodeMin - theRay->Origin) * aInvDirect;"
364 EOL" aTime1 = (aNodeMax - theRay->Origin) * aInvDirect;"
e276548b 365 EOL
265d4508 366 EOL" aTimeMax = max (aTime0, aTime1);"
367 EOL" aTimeMin = min (aTime0, aTime1);"
e276548b 368 EOL
265d4508 369 EOL" aTimeExit = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
370 EOL" aTimeMin2 = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
371 EOL
372 EOL" const bool aHitRgh = (aTimeMin2 <= aTimeExit) & (aTimeExit >= 0.f) & (aTimeMin2 <= theIntersect->Time);"
e276548b 373 EOL
374 EOL" if (aHitLft & aHitRgh)"
375 EOL" {"
376 EOL" aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;"
377 EOL
378 EOL" push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);"
e276548b 379 EOL" }"
380 EOL" else"
381 EOL" {"
382 EOL" if (aHitLft | aHitRgh)"
383 EOL" {"
384 EOL" aNode = aHitLft ? aData.y : aData.z;"
e276548b 385 EOL" }"
386 EOL" else"
387 EOL" {"
388 EOL" if (aHead < 0)"
389 EOL" return aTriangleIndex;"
390 EOL
391 EOL" pop (aStack, &aHead, &aNode);"
e276548b 392 EOL" }"
393 EOL" }"
394 EOL" }"
395 EOL" else " // if leaf node
396 EOL" {"
265d4508 397 EOL" for (int aTriIdx = aData.y; aTriIdx <= aData.z; ++aTriIdx)"
e276548b 398 EOL" {"
265d4508 399 EOL" const int4 aTestTriangle = theGeometryTriangBuffer[aTriIdx];"
e276548b 400 EOL
265d4508 401 EOL" const float4 aPoint0 = theGeometryVertexBuffer[aTestTriangle.x];"
402 EOL" const float4 aPoint1 = theGeometryVertexBuffer[aTestTriangle.y];"
403 EOL" const float4 aPoint2 = theGeometryVertexBuffer[aTestTriangle.z];"
e276548b 404 EOL
265d4508 405 EOL" float4 aNormal; float aU, aV;"
e276548b 406 EOL
265d4508 407 EOL" float aTime = IntersectTriangle (theRay,"
408 EOL" aPoint0,"
409 EOL" aPoint1,"
410 EOL" aPoint2,"
411 EOL" &aNormal,"
412 EOL" &aU,"
413 EOL" &aV);"
e276548b 414 EOL
265d4508 415 EOL" if (aTime < theIntersect->Time)"
e276548b 416 EOL" {"
265d4508 417 EOL" aTriangleIndex = aTestTriangle;"
418 EOL" theIntersect->Normal = aNormal;"
419 EOL" theIntersect->Time = aTime;"
420 EOL" theIntersect->U = aU;"
421 EOL" theIntersect->V = aV;"
e276548b 422 EOL" }"
423 EOL" }"
424 EOL
425 EOL" if (aHead < 0)"
426 EOL" return aTriangleIndex;"
427 EOL
428 EOL" pop (aStack, &aHead, &aNode);"
e276548b 429 EOL" }"
430 EOL" }"
431 EOL
432 EOL" return aTriangleIndex;"
265d4508 433 EOL" }"
e276548b 434 EOL
435 // =======================================================================
265d4508 436 // function : ObjectAnyHit
437 // purpose : Finds intersection with any object triangle
e276548b 438 // =======================================================================
265d4508 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)"
e276548b 446 EOL" {"
447 EOL" uint aStack [32];"
e276548b 448 EOL
265d4508 449 EOL" char aHead = -1;" // stack pointer
450 EOL" uint aNode = 0;" // node to visit
e276548b 451 EOL
265d4508 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)),"
459 EOL" 0.f);"
e276548b 460 EOL
265d4508 461 EOL" float aTimeExit;"
e276548b 462 EOL" float aTimeMin1;"
463 EOL" float aTimeMin2;"
464 EOL
465 EOL" while (true)"
466 EOL" {"
265d4508 467 EOL" const int3 aData = theObjectNodeInfoBuffer[aNode].xyz;"
468 EOL
469 EOL" if (aData.x == 0)" // if inner node
470 EOL" {"
471 EOL" float4 aNodeMin = theObjectMinPointBuffer[aData.y];"
472 EOL" float4 aNodeMax = theObjectMaxPointBuffer[aData.y];"
473 EOL
474 EOL" float4 aTime0 = (aNodeMin - theRay->Origin) * aInvDirect;"
475 EOL" float4 aTime1 = (aNodeMax - theRay->Origin) * aInvDirect;"
476 EOL
477 EOL" float4 aTimeMax = max (aTime0, aTime1);"
478 EOL" float4 aTimeMin = min (aTime0, aTime1);"
479 EOL
480 EOL" aTimeExit = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
481 EOL" aTimeMin1 = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
482 EOL
483 EOL" const bool aHitLft = (aTimeMin1 <= aTimeExit) & (aTimeExit >= 0.f) & (aTimeMin1 <= theDistance);"
484 EOL
485 EOL" aNodeMin = theObjectMinPointBuffer[aData.z];"
486 EOL" aNodeMax = theObjectMaxPointBuffer[aData.z];"
487 EOL
488 EOL" aTime0 = (aNodeMin - theRay->Origin) * aInvDirect;"
489 EOL" aTime1 = (aNodeMax - theRay->Origin) * aInvDirect;"
490 EOL
491 EOL" aTimeMax = max (aTime0, aTime1);"
492 EOL" aTimeMin = min (aTime0, aTime1);"
493 EOL
494 EOL" aTimeExit = min (aTimeMax.x, min (aTimeMax.y, aTimeMax.z));"
495 EOL" aTimeMin2 = max (aTimeMin.x, max (aTimeMin.y, aTimeMin.z));"
496 EOL
497 EOL" const bool aHitRgh = (aTimeMin2 <= aTimeExit) & (aTimeExit >= 0.f) & (aTimeMin2 <= theDistance);"
498 EOL
499 EOL" if (aHitLft & aHitRgh)"
500 EOL" {"
501 EOL" aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;"
502 EOL
503 EOL" push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);"
504 EOL" }"
505 EOL" else"
506 EOL" {"
507 EOL" if (aHitLft | aHitRgh)"
508 EOL" {"
509 EOL" aNode = aHitLft ? aData.y : aData.z;"
510 EOL" }"
511 EOL" else"
512 EOL" {"
513 EOL" if (aHead < 0)"
514 EOL" return 1.f;"
515 EOL
516 EOL" pop (aStack, &aHead, &aNode);"
517 EOL" }"
518 EOL" }"
519 EOL" }"
520 EOL" else " // if leaf node
521 EOL" {"
522 EOL" for (int aTriIdx = aData.y; aTriIdx <= aData.z; ++aTriIdx)"
523 EOL" {"
524 EOL" const int4 aTestTriangle = theGeometryTriangBuffer[aTriIdx];"
525 EOL
526 EOL" const float4 aPoint0 = theGeometryVertexBuffer[aTestTriangle.x];"
527 EOL" const float4 aPoint1 = theGeometryVertexBuffer[aTestTriangle.y];"
528 EOL" const float4 aPoint2 = theGeometryVertexBuffer[aTestTriangle.z];"
529 EOL
530 EOL" float4 aNormal; float aU, aV;"
531 EOL
532 EOL" float aTime = IntersectTriangle (theRay,"
533 EOL" aPoint0,"
534 EOL" aPoint1,"
535 EOL" aPoint2,"
536 EOL" &aNormal,"
537 EOL" &aU,"
538 EOL" &aV);"
539 EOL
540 EOL" if (aTime < theDistance)"
541 EOL" {"
542 EOL" return 0.f;"
543 EOL" }"
544 EOL" }"
545 EOL
546 EOL" if (aHead < 0)"
547 EOL" return 1.f;"
548 EOL
549 EOL" pop (aStack, &aHead, &aNode);"
550 EOL" }"
551 EOL" }"
552 EOL
553 EOL" return 1.f;"
554 EOL" }"
555 EOL
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)"
570 EOL" {"
571 EOL" theIntersect->Time = MAXFLOAT;"
572 EOL
573 EOL" uint aStack [16];"
574 EOL
575 EOL" char aHead = -1;" // stack pointer
576 EOL" uint aNode = 0;" // node to visit
577 EOL
578 EOL" int4 aNearestTriangle = (int4) (-1);"
579 EOL
580 EOL" while (true)"
581 EOL" {"
582 EOL" const int4 aData = theSceneNodeInfoBuffer[aNode];"
e276548b 583 EOL
265d4508 584 EOL" if (aData.x != 0)" // if leaf node
e276548b 585 EOL" {"
265d4508 586 EOL" const float4 aNodeMin = theSceneMinPointBuffer[aNode];"
587 EOL" const float4 aNodeMax = theSceneMaxPointBuffer[aNode];"
588 EOL
589 EOL" if (IntersectBox (theRay, aNodeMin, aNodeMax) <= theIntersect->Time)"
590 EOL" {"
591 EOL" int4 anIndex = ObjectNearestHit (theRay,"
592 EOL" theIntersect,"
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);"
598 EOL
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,"
603 EOL" anIndex.w);"
604 EOL" }"
605 EOL
606 EOL" if (aHead < 0)"
607 EOL" return aNearestTriangle;"
608 EOL
609 EOL" pop (aStack, &aHead, &aNode);"
610 EOL" }"
611 EOL" else " // if inner node
612 EOL" {"
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];"
617 EOL
618 EOL" float aTimeMin1;"
619 EOL" float aTimeMin2;"
620 EOL
e276548b 621 EOL" IntersectNodes (theRay,"
265d4508 622 EOL" aNodeMinLft,"
623 EOL" aNodeMaxLft,"
624 EOL" aNodeMinRgh,"
625 EOL" aNodeMaxRgh,"
e276548b 626 EOL" &aTimeMin1,"
627 EOL" &aTimeMin2,"
265d4508 628 EOL" theIntersect->Time);"
e276548b 629 EOL
630 EOL" const bool aHitLft = (aTimeMin1 != -MAXFLOAT);"
631 EOL" const bool aHitRgh = (aTimeMin2 != -MAXFLOAT);"
632 EOL
633 EOL" if (aHitLft & aHitRgh)"
634 EOL" {"
635 EOL" aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;"
636 EOL
637 EOL" push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);"
638 EOL" }"
639 EOL" else"
640 EOL" {"
641 EOL" if (aHitLft | aHitRgh)"
642 EOL" {"
643 EOL" aNode = aHitLft ? aData.y : aData.z;"
644 EOL" }"
645 EOL" else"
646 EOL" {"
647 EOL" if (aHead < 0)"
265d4508 648 EOL" return aNearestTriangle;"
e276548b 649 EOL
650 EOL" pop (aStack, &aHead, &aNode);"
651 EOL" }"
652 EOL" }"
653 EOL" }"
265d4508 654 EOL" }"
655 EOL
656 EOL" return aNearestTriangle;"
657 EOL" }"
658 EOL
659 // =======================================================================
660 // function : AnyHit
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)"
673 EOL" {"
674 EOL" uint aStack [16];"
e276548b 675 EOL
265d4508 676 EOL" char aHead = -1;" // stack pointer
677 EOL" uint aNode = 0;" // node to visit
e276548b 678 EOL
265d4508 679 EOL" while (true)"
680 EOL" {"
681 EOL" const int4 aData = theSceneNodeInfoBuffer[aNode];"
e276548b 682 EOL
265d4508 683 EOL" if (aData.x != 0)" // if leaf node
684 EOL" {"
685 EOL" const float4 aNodeMin = theSceneMinPointBuffer[aNode];"
686 EOL" const float4 aNodeMax = theSceneMaxPointBuffer[aNode];"
e276548b 687 EOL
265d4508 688 EOL" if (IntersectBox (theRay, aNodeMin, aNodeMax) <= theDistance)"
689 EOL" {"
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,"
696 EOL" theDistance))"
e276548b 697 EOL" {"
e276548b 698 EOL" return 0.f;"
e276548b 699 EOL" }"
700 EOL" }"
701 EOL
702 EOL" if (aHead < 0)"
265d4508 703 EOL" return 1.f;"
e276548b 704 EOL
705 EOL" pop (aStack, &aHead, &aNode);"
706 EOL" }"
265d4508 707 EOL" else" // if inner node
708 EOL" {"
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];"
713 EOL
714 EOL" float aTimeMin1;"
715 EOL" float aTimeMin2;"
716 EOL
717 EOL" IntersectNodes (theRay,"
718 EOL" aNodeMinLft,"
719 EOL" aNodeMaxLft,"
720 EOL" aNodeMinRgh,"
721 EOL" aNodeMaxRgh,"
722 EOL" &aTimeMin1,"
723 EOL" &aTimeMin2,"
724 EOL" theDistance);"
725 EOL
726 EOL" const bool aHitLft = (aTimeMin1 != -MAXFLOAT);"
727 EOL" const bool aHitRgh = (aTimeMin2 != -MAXFLOAT);"
e276548b 728 EOL
265d4508 729 EOL" if (aHitLft & aHitRgh)"
730 EOL" {"
731 EOL" aNode = (aTimeMin1 < aTimeMin2) ? aData.y : aData.z;"
732 EOL
733 EOL" push (aStack, &aHead, (aTimeMin1 < aTimeMin2) ? aData.z : aData.y);"
734 EOL" }"
735 EOL" else"
736 EOL" {"
737 EOL" if (aHitLft | aHitRgh)"
738 EOL" {"
739 EOL" aNode = aHitLft ? aData.y : aData.z;"
740 EOL" }"
741 EOL" else"
742 EOL" {"
743 EOL" if (aHead < 0)"
744 EOL" return 1.f;"
745 EOL
746 EOL" pop (aStack, &aHead, &aNode);"
747 EOL" }"
748 EOL" }"
749 EOL" }"
750 EOL" }"
e276548b 751 EOL" }"
752 EOL
753 EOL" #define _MAX_DEPTH_ 5"
754 EOL
265d4508 755 EOL" #define THRESHOLD (float4) (0.1f, 0.1f, 0.1f, 1.f)"
e276548b 756 EOL
265d4508 757 EOL" #define LIGHT_POS(Buffer, LightID) Buffer[2 * LightID + 1]"
758 EOL" #define LIGHT_RAD(Buffer, LightID) Buffer[2 * LightID + 0]"
759 EOL
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]"
e276548b 767 EOL
768 // =======================================================================
265d4508 769 // function : Radiance
e276548b 770 // purpose : Computes color of specified ray
771 // =======================================================================
265d4508 772 EOL" float4 Radiance (SRay* theRay,"
e276548b 773 EOL" __read_only image2d_t theEnvMap,"
265d4508 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)"
e276548b 791 EOL" {"
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);"
794 EOL
795 EOL" SIntersect aHit;"
796 EOL
797 EOL" for (int aDepth = 0; aDepth < _MAX_DEPTH_; ++aDepth)"
798 EOL" {"
265d4508 799 EOL" int4 aTriangle = NearestHit (theRay,"
800 EOL" &aHit,"
801 EOL" theSceneNodeInfoBuffer,"
802 EOL" theSceneMinPointBuffer,"
803 EOL" theSceneMaxPointBuffer,"
804 EOL" theObjectNodeInfoBuffer,"
805 EOL" theObjectMinPointBuffer,"
806 EOL" theObjectMaxPointBuffer,"
807 EOL" theGeometryTriangBuffer,"
808 EOL" theGeometryVertexBuffer);"
e276548b 809 EOL
810 EOL" if (aTriangle.x < 0.f)"
811 EOL" {"
265d4508 812 EOL" if (aWeight.w != 0.f)"
e276548b 813 EOL" break;"
814 EOL
265d4508 815 EOL" float aTime = IntersectSphere (theRay, theSceneRadius);"
e276548b 816 EOL
265d4508 817 EOL" if (aTime != MAXFLOAT)"
818 EOL" {"
819 EOL" aResult += aWeight * read_imagef (theEnvMap, EnvironmentSampler,"
820 EOL" Latlong (theRay->Origin + theRay->Direct * aTime, theSceneRadius));"
821 EOL" }"
e276548b 822 EOL
823 EOL" return (float4) (aResult.x,"
824 EOL" aResult.y,"
825 EOL" aResult.z,"
826 EOL" aWeight.w);"
827 EOL" }"
828 EOL
265d4508 829 EOL // Compute geometric normal
e276548b 830 EOL" float4 aGeomNormal = aHit.Normal; aGeomNormal = fast_normalize (aGeomNormal);"
831 EOL
265d4508 832 EOL // Compute interpolated normal
833 EOL" float4 aNormal = SmoothNormal (theGeometryNormalBuffer, &aHit, aTriangle);"
e276548b 834 EOL
265d4508 835 EOL // Compute intersection point
e276548b 836 EOL" float4 aPoint = theRay->Direct * aHit.Time + theRay->Origin;"
837 EOL
265d4508 838 EOL" float4 aMaterAmb = MATERIAL_AMBN (theMaterialBuffer, aTriangle);"
839 EOL" float4 aMaterTrn = MATERIAL_TRAN (theMaterialBuffer, aTriangle);"
e276548b 840 EOL
265d4508 841 EOL" aResult += aWeight * theGlobalAmbient * aMaterAmb *"
e276548b 842 EOL" (aMaterTrn.x * max (fabs (dot (theRay->Direct, aNormal)), 0.5f));"
843 EOL
265d4508 844 EOL" for (int nLight = 0; nLight < theLightBufferSize; ++nLight)"
845 EOL" {"
846 EOL" float4 aLightPosition = LIGHT_POS (theLightSourceBuffer, nLight);"
e276548b 847 EOL
848 EOL" SRay aShadow;"
849 EOL" aShadow.Direct = aLightPosition;"
850 EOL
851 EOL" float aLightDistance = MAXFLOAT;"
852 EOL" if (aLightPosition.w != 0.f)"
853 EOL" {"
854 EOL" aLightDistance = length (aLightPosition - aPoint);"
855 EOL" aShadow.Direct = (aLightPosition - aPoint) * (1.f / aLightDistance);"
856 EOL" }"
857 EOL
265d4508 858 EOL" aShadow.Origin = aPoint + aShadow.Direct * theSceneEpsilon +"
859 EOL" aGeomNormal * copysign (theSceneEpsilon, dot (aGeomNormal, aShadow.Direct));"
e276548b 860 EOL
265d4508 861 EOL" float aVisibility = 1.f;"
e276548b 862 EOL
265d4508 863 EOL" if (theShadowsEnabled)"
e276548b 864 EOL" {"
265d4508 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);"
e276548b 875 EOL" }"
876 EOL
265d4508 877 EOL" if (aVisibility > 0.f)"
878 EOL" {"
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,"
883 EOL" aNormal,"
884 EOL" LIGHT_RAD (theLightSourceBuffer, nLight),"
885 EOL" aMaterTrn.y);"
886 EOL" }"
e276548b 887 EOL" }"
888 EOL
889 EOL" if (aMaterTrn.y > 0.f)"
890 EOL" {"
891 EOL" aWeight *= aMaterTrn.y;"
892 EOL" }"
893 EOL" else"
894 EOL" {"
265d4508 895 EOL" aWeight *= theReflectEnabled ? MATERIAL_REFL (theMaterialBuffer, aTriangle) : ZERO;"
e276548b 896 EOL
265d4508 897 EOL" float4 aDirect = theRay->Direct - 2.f * dot (theRay->Direct, aNormal) * aNormal;"
e276548b 898 EOL
265d4508 899 EOL" float aDdotN = dot (aDirect, aGeomNormal);"
e276548b 900 EOL" if (aDdotN < 0.f)"
265d4508 901 EOL" theRay->Direct -= 2.f * dot (theRay->Direct, aGeomNormal) * aGeomNormal;"
902 EOL" else"
903 EOL" theRay->Direct = aDirect;"
e276548b 904 EOL" }"
905 EOL
265d4508 906 EOL" if (all (islessequal (aWeight, THRESHOLD)))"
e276548b 907 EOL" {"
908 EOL" return (float4) (aResult.x,"
909 EOL" aResult.y,"
910 EOL" aResult.z,"
911 EOL" aWeight.w);"
912 EOL" }"
913 EOL
265d4508 914 EOL" theRay->Origin = theRay->Direct * theSceneEpsilon + aPoint;"
e276548b 915 EOL" }"
916 EOL
917 EOL" return (float4) (aResult.x,"
918 EOL" aResult.y,"
919 EOL" aResult.z,"
920 EOL" aWeight.w);"
921 EOL" }"
922 EOL
e276548b 923 ///////////////////////////////////////////////////////////////////////////////
924 // Ray tracing kernel functions
925 EOL
926 // =======================================================================
265d4508 927 // function : RaytraceRender
e276548b 928 // purpose : Computes pixel color using ray-tracing
929 // =======================================================================
265d4508 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)"
e276548b 953 EOL" {"
265d4508 954 EOL" const int aPixelX = get_global_id (0);"
955 EOL" const int aPixelY = get_global_id (1);"
e276548b 956 EOL
265d4508 957 EOL" if (aPixelX >= theSizeX || aPixelY >= theSizeY)"
e276548b 958 EOL" return;"
959 EOL
960 EOL" private SRay aRay;"
961 EOL
962 EOL" GenerateRay (&aRay,"
265d4508 963 EOL" aPixelX,"
964 EOL" aPixelY,"
e276548b 965 EOL" theSizeX,"
966 EOL" theSizeY,"
967 EOL" theOrigins,"
968 EOL" theDirects);"
969 EOL
e276548b 970 EOL" float4 aColor = (float4) (0.f, 0.f, 0.f, 1.f);"
971 EOL
265d4508 972 EOL" float aTimeStart = IntersectBox (&aRay, theSceneMinPointBuffer[0], theSceneMaxPointBuffer[0]);"
973 EOL
974 EOL" if (aTimeStart != MAXFLOAT)"
e276548b 975 EOL" {"
265d4508 976 EOL" aRay.Origin += aRay.Direct * max (aTimeStart - theSceneEpsilon, 0.f);"
977 EOL
978 EOL" aColor = clamp (Radiance (&aRay,"
979 EOL" theEnvMap,"
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);"
e276548b 997 EOL" }"
998 EOL
265d4508 999 EOL" write_imagef (theOutput, (int2) (aPixelX, aPixelY), aColor);"
e276548b 1000 EOL" }"
1001 EOL
1002 EOL" const sampler_t OutputSampler ="
1003 EOL" CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;"
1004 EOL
265d4508 1005 EOL" #define _LUM_DELTA_ 0.085f"
e276548b 1006 EOL
1007 EOL" #define AA_MAX 0.559017f"
1008 EOL" #define AA_MIN 0.186339f"
1009 EOL
1010 // =======================================================================
265d4508 1011 // function : RaytraceSmooth
e276548b 1012 // purpose : Performs adaptive sub-pixel rendering
1013 // =======================================================================
265d4508 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)"
e276548b 1038 EOL" {"
265d4508 1039 EOL" const int aPixelX = get_global_id (0);"
1040 EOL" const int aPixelY = get_global_id (1);"
e276548b 1041 EOL
265d4508 1042 EOL" if (aPixelX >= theSizeX || aPixelY >= theSizeY)"
e276548b 1043 EOL" return;"
1044 EOL
265d4508 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));"
e276548b 1048 EOL
265d4508 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));"
e276548b 1052 EOL
265d4508 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));"
e276548b 1056 EOL
265d4508 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_;"
e276548b 1065 EOL
265d4508 1066 EOL" if (!render)"
1067 EOL" {"
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;"
1076 EOL
1077 EOL" float aLum = (aClr0.w == 1.f) ? -1.f : (0.2126f * aClr0.x + 0.7152f * aClr0.y + 0.0722f * aClr0.z);"
1078 EOL
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_;"
1087 EOL" }"
e276548b 1088 EOL
265d4508 1089 EOL" float4 aColor = clamp (aClr0, 0.f, 1.f);"
e276548b 1090 EOL
1091 EOL" private SRay aRay;"
1092 EOL
265d4508 1093 EOL" const float4 aBoxMin = theSceneMinPointBuffer[0];"
1094 EOL" const float4 aBoxMax = theSceneMaxPointBuffer[0];"
e276548b 1095 EOL
1096 EOL" if (render)"
1097 EOL" {"
1098 EOL" for (int aSample = 0; aSample <= 3; ++aSample)"
1099 EOL" {"
265d4508 1100 EOL" float fX = aPixelX, fY = aPixelY;"
e276548b 1101 EOL
1102 EOL" if (aSample == 0)"
1103 EOL" {"
1104 EOL" fX -= AA_MIN; fY -= AA_MAX;"
1105 EOL" }"
1106 EOL" else if (aSample == 1)"
1107 EOL" {"
1108 EOL" fX -= AA_MAX; fY += AA_MIN;"
1109 EOL" }"
1110 EOL" else if (aSample == 2)"
1111 EOL" {"
1112 EOL" fX += AA_MIN; fY += AA_MAX;"
1113 EOL" }"
1114 EOL" else"
1115 EOL" {"
1116 EOL" fX += AA_MAX; fY -= AA_MIN;"
1117 EOL" }"
1118 EOL
1119 EOL" GenerateRay (&aRay,"
1120 EOL" fX,"
1121 EOL" fY,"
1122 EOL" theSizeX,"
1123 EOL" theSizeY,"
1124 EOL" theOrigins,"
1125 EOL" theDirects);"
1126 EOL
265d4508 1127 EOL" float aTimeStart = IntersectBox (&aRay, aBoxMin, aBoxMax);"
e276548b 1128 EOL
265d4508 1129 EOL" if (aTimeStart != MAXFLOAT)"
e276548b 1130 EOL" {"
265d4508 1131 EOL" aRay.Origin += aRay.Direct * max (aTimeStart - theSceneEpsilon, 0.f);"
1132 EOL
1133 EOL" aColor += clamp (Radiance (&aRay,"
1134 EOL" theEnvMap,"
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);"
e276548b 1152 EOL" }"
1153 EOL" else"
1154 EOL" aColor += (float4) (0.f, 0.f, 0.f, 1.f);"
1155 EOL" }"
1156 EOL
1157 EOL" aColor *= 1.f / 5.f;"
1158 EOL" }"
1159 EOL
265d4508 1160 EOL" write_imagef (theOutput, (int2) (aPixelX, aPixelY), aColor);"
e276548b 1161 EOL" }";
1162
1163#endif