Voxel
 All Classes Namespaces Files Functions Typedefs Enumerations Enumerator Macros Pages
trace.cl
Go to the documentation of this file.
1 /**
2  \file
3  \brief Contains OpenCL code for raytracing a scene.
4 */
5 
6 
7 
8 
9 /*==============================================================================
10  Node structure
11 ==============================================================================*/
12 
13 
14 struct Node {
15 
16  uint texture;
17  uint children;
18 };
19 
20 
21 
22 
23 /*==============================================================================
24  Morton Z-order functions
25 ==============================================================================*/
26 
27 
28 /**
29  \brief Maps 2D coordinates to Morton Z-order indices.
30 
31  \param[in] coordinate The 2D coordinates.
32  \param[in] dimension Dimensions of the region to which the coordinates belong.
33  \return The Morton Z-order index.
34 
35  Suppose that the coordinates represent positions in a buffer of the given
36  dimensions. This function will return the (row-major) Morton Z-order index
37  corresponding to these coordinates. The \ref InverseMortonZOrder function will
38  perform the reverse operation (i.e. it is the inverse of this function). The
39  \p dimension argument does not need to be a power of two.
40 
41  For example, if \p dimension.{x,y}={9,6}, then coordinates will map to
42  indices as follows:
43 
44  \f[
45  \begin{array}{ccccccccc}
46  0 & 1 & 4 & 5 & 16 & 17 & 20 & 21 & 48 \\
47  2 & 3 & 6 & 7 & 18 & 19 & 22 & 23 & 49 \\
48  8 & 9 & 12 & 13 & 24 & 25 & 28 & 29 & 50 \\
49  10 & 11 & 14 & 15 & 26 & 27 & 30 & 31 & 51 \\
50  32 & 33 & 36 & 37 & 40 & 41 & 44 & 45 & 52 \\
51  34 & 35 & 38 & 39 & 42 & 43 & 46 & 47 & 53
52  \end{array}
53  \f]
54 
55  \see \ref InverseMortonZOrder
56 */
57 uint MortonZOrder( uint2 coordinate, uint2 dimension ) {
58 
59  uint result = 0;
60 
61  int shift = 0;
62  for ( ; any( ( 1u << shift ) < dimension ); ++shift );
63  for ( --shift; shift >= 0; --shift ) {
64 
65  uint const bit = ( 1u << shift );
66 
67  if ( coordinate.y & bit ) {
68 
69  dimension.y -= bit;
70  result += ( dimension.x << shift );
71  }
72  else if ( dimension.y > bit )
73  dimension.y = bit;
74 
75  if ( coordinate.x & bit ) {
76 
77  dimension.x -= bit;
78  result += ( dimension.y << shift );
79  }
80  else if ( dimension.x > bit )
81  dimension.x = bit;
82  }
83 
84  return result;
85 }
86 
87 
88 /**
89  \brief Maps Morton Z-order indices to 2D coordinates.
90 
91  \param[in] index The Morton Z-order index.
92  \param[in] dimension Dimensions of the region to which the coordinates belong.
93  \return The 2D coordinates.
94 
95  This function will return the coordinates which should be considered by the
96  "index"th thread, if it wants to work in Morton Z-order, over a buffer of the
97  given size. In other words, this is an inverse Z-order mapping (and this
98  function is the inverse of \ref MortonZOrder). The \p dimension argument does
99  not need to be a power of two.
100 
101  \see \ref MortonZOrder
102 */
103 uint2 InverseMortonZOrder( uint index, uint2 dimension ) {
104 
105  uint2 result = ( uint2 )( 0, 0 );
106 
107  int shift = 0;
108  for ( ; any( ( 1u << shift ) < dimension ); ++shift );
109  for ( --shift; shift >= 0; --shift ) {
110 
111  uint const bit = ( 1u << shift );
112 
113  { uint const shifted = ( dimension.x << shift );
114  if ( index >= shifted ) {
115 
116  result.y += bit;
117  dimension.y -= bit;
118  index -= shifted;
119  }
120  else if ( dimension.y > bit )
121  dimension.y = bit;
122  }
123 
124  { uint const shifted = ( dimension.y << shift );
125  if ( index >= shifted ) {
126 
127  result.x += bit;
128  dimension.x -= bit;
129  index -= shifted;
130  }
131  else if ( dimension.x > bit )
132  dimension.x = bit;
133  }
134  }
135 
136  return result;
137 }
138 
139 
140 
141 
142 /*==============================================================================
143  Project function
144 ==============================================================================*/
145 
146 
147 /**
148  \brief Projects a ray onto the unit cube.
149 
150  \param[in,out] pProjected Origin of the ray on input, intersection with unit cube on output.
151  \param[in] ray Direction in which the ray propagates.
152  \return True if the ray intersects the unit cube, false otherwise.
153 
154  If the ray never intersects with the unit cube \f$[0,1]^3\f$, then this
155  function will return false. Otherwise it will store the first intersection
156  point in \p *pProjected, and return true. If the origin of the ray is
157  inside the cube, then \p *pProjected will not be changed, but the function
158  will return true.
159 */
160 float Project( __private float3* pProjected, float3 const ray ) {
161 
162  if (
163  ( pProjected->x >= 0 ) && ( pProjected->x <= 1 ) &&
164  ( pProjected->y >= 0 ) && ( pProjected->y <= 1 ) &&
165  ( pProjected->z >= 0 ) && ( pProjected->z <= 1 )
166  )
167  {
168  return 0; // **NOTE: return
169  }
170  else {
171 
172  if ( pProjected->x < 0 ) {
173 
174  if ( ray.x > 0 ) {
175 
176  float const lambda = -pProjected->x / ray.x;
177  float3 projected = *pProjected + lambda * ray;
178  if (
179  ( projected.y >= 0 ) && ( projected.y <= 1 ) &&
180  ( projected.z >= 0 ) && ( projected.z <= 1 )
181  )
182  {
183  projected.x = 0; // just to make sure
184  *pProjected = projected;
185  return lambda; // **NOTE: return
186  }
187  }
188  }
189  else if ( pProjected->x > 1 ) {
190 
191  if ( ray.x < 0 ) {
192 
193  float const lambda = ( 1 - pProjected->x ) / ray.x;
194  float3 projected = *pProjected + lambda * ray;
195  if (
196  ( projected.y >= 0 ) && ( projected.y <= 1 ) &&
197  ( projected.z >= 0 ) && ( projected.z <= 1 )
198  )
199  {
200  projected.x = 1; // just to make sure
201  *pProjected = projected;
202  return lambda; // **NOTE: return
203  }
204  }
205  }
206 
207  if ( pProjected->y < 0 ) {
208 
209  if ( ray.y > 0 ) {
210 
211  float const lambda = -pProjected->y / ray.y;
212  float3 projected = *pProjected + lambda * ray;
213  if (
214  ( projected.x >= 0 ) && ( projected.x <= 1 ) &&
215  ( projected.z >= 0 ) && ( projected.z <= 1 )
216  )
217  {
218  projected.y = 0; // just to make sure
219  *pProjected = projected;
220  return lambda; // **NOTE: return
221  }
222  }
223  }
224  else if ( pProjected->y > 1 ) {
225 
226  if ( ray.y < 0 ) {
227 
228  float const lambda = ( 1 - pProjected->y ) / ray.y;
229  float3 projected = *pProjected + lambda * ray;
230  if (
231  ( projected.x >= 0 ) && ( projected.x <= 1 ) &&
232  ( projected.z >= 0 ) && ( projected.z <= 1 )
233  )
234  {
235  projected.y = 1; // just to make sure
236  *pProjected = projected;
237  return lambda; // **NOTE: return
238  }
239  }
240  }
241 
242  if ( pProjected->z < 0 ) {
243 
244  if ( ray.z > 0 ) {
245 
246  float const lambda = -pProjected->z / ray.z;
247  float3 projected = *pProjected + lambda * ray;
248  if (
249  ( projected.x >= 0 ) && ( projected.x <= 1 ) &&
250  ( projected.y >= 0 ) && ( projected.y <= 1 )
251  )
252  {
253  projected.z = 0; // just to make sure
254  *pProjected = projected;
255  return lambda; // **NOTE: return
256  }
257  }
258  }
259  else if ( pProjected->z > 1 ) {
260 
261  if ( ray.z < 0 ) {
262 
263  float const lambda = ( 1 - pProjected->z ) / ray.z;
264  float3 projected = *pProjected + lambda * ray;
265  if (
266  ( projected.x >= 0 ) && ( projected.x <= 1 ) &&
267  ( projected.y >= 0 ) && ( projected.y <= 1 )
268  )
269  {
270  projected.z = 1; // just to make sure
271  *pProjected = projected;
272  return lambda; // **NOTE: return
273  }
274  }
275  }
276  }
277 
278  return INFINITY;
279 }
280 
281 
282 
283 
284 /*==============================================================================
285  TraceTexture function
286 ==============================================================================*/
287 
288 
289 int TraceTexture(
290  __private uint* const pColor,
291  __private float* const pLambda,
292  __read_only image3d_t image,
293  uint index,
294  __private float3* const pRemaining,
295  __private float3* const pRemainingMaximum,
296  int const origin
297 )
298 {
299  sampler_t const sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;
300 
301  uint const mask = ( ( origin & 4 ) ? ( 3 << 22 ) : 0 ) | ( ( origin & 2 ) ? ( 3 << 11 ) : 0 ) | ( ( origin & 1 ) ? 3 : 0 );
302 
303  *pRemainingMaximum /= 2;
304  if ( pRemaining->x >= pRemainingMaximum->x ) {
305 
306  index += 2;
307  pRemaining->x -= pRemainingMaximum->x;
308  }
309  if ( pRemaining->y >= pRemainingMaximum->y ) {
310 
311  index += ( 2 << 11 );
312  pRemaining->y -= pRemainingMaximum->y;
313  }
314  if ( pRemaining->z >= pRemainingMaximum->z ) {
315 
316  index += ( 2 << 22 );
317  pRemaining->z -= pRemainingMaximum->z;
318  }
319 
320  *pRemainingMaximum /= 2;
321  if ( pRemaining->x >= pRemainingMaximum->x ) {
322 
323  index += 1;
324  pRemaining->x -= pRemainingMaximum->x;
325  }
326  if ( pRemaining->y >= pRemainingMaximum->y ) {
327 
328  index += ( 1u << 11 );
329  pRemaining->y -= pRemainingMaximum->y;
330  }
331  if ( pRemaining->z >= pRemainingMaximum->z ) {
332 
333  index += ( 1u << 22 );
334  pRemaining->z -= pRemainingMaximum->z;
335  }
336 
337  int comparisons = 0;
338  for ( ; ; ) {
339 
340  uint const source = ( index ^ mask );
341  int4 coordinate = ( int4 )( source & 0x7ff, ( source >> 11 ) & 0x7ff, source >> 22, 0 );
342  float4 const color = read_imagef(
343  image,
344  sampler,
345  coordinate
346  );
347  if ( ( color.x + color.y + color.z ) > 0 ) { // **NOTE: black is transparent
348 
349  comparisons = -1;
350  *pColor = 0xff000000 + ( ( uint )( color.x * 255.99 ) << 16 ) + ( ( uint )( color.y * 255.99 ) << 8 ) + ( uint )( color.z * 255.99 ); // **NOTE: BGRA
351  goto texture_loop_end_label; // **NOTE: goto
352  }
353 
354  comparisons = (
355  ( ( pRemaining->x < pRemaining->y ) ? 1 : 0 ) |
356  ( ( pRemaining->y < pRemaining->z ) ? 2 : 0 ) |
357  ( ( pRemaining->z < pRemaining->x ) ? 4 : 0 )
358  );
359  switch( comparisons ) {
360 
361  case 0: case 7: // any are fine--fall through
362  case 1: case 3: {
363 
364  *pLambda += pRemaining->x;
365  pRemaining->y -= pRemaining->x;
366  pRemaining->z -= pRemaining->x;
367  pRemaining->x = pRemainingMaximum->x;
368 
369  if ( ( index & 3 ) == 0 ) {
370 
371  index |= 3;
372  goto texture_loop_end_label; // **NOTE: goto
373  }
374 
375  index -= 1;
376 
377  break;
378  }
379 
380  case 2: case 6: {
381 
382  *pLambda += pRemaining->y;
383  pRemaining->x -= pRemaining->y;
384  pRemaining->z -= pRemaining->y;
385  pRemaining->y = pRemainingMaximum->y;
386 
387  if ( ( ( index >> 11 ) & 3 ) == 0 ) {
388 
389  index |= ( 3u << 11 );
390  goto texture_loop_end_label; // **NOTE: goto
391  }
392 
393  index -= ( 1u << 11 );
394 
395  break;
396  }
397 
398  case 4: case 5: {
399 
400  *pLambda += pRemaining->z;
401  pRemaining->x -= pRemaining->z;
402  pRemaining->y -= pRemaining->z;
403  pRemaining->z = pRemainingMaximum->z;
404 
405  if ( ( ( index >> 22 ) & 3 ) == 0 ) {
406 
407  index |= ( 3u << 22 );
408  goto texture_loop_end_label; // **NOTE: goto
409  }
410 
411  index -= ( 1u << 22 );
412 
413  break;
414  }
415  }
416  }
417 texture_loop_end_label:;
418 
419  pRemaining->x += pRemainingMaximum->x * ( index & 3 );
420  pRemaining->y += pRemainingMaximum->y * ( ( index >> 11 ) & 3 );
421  pRemaining->z += pRemainingMaximum->z * ( ( index >> 22 ) & 3 );
422 
423  *pRemainingMaximum *= 4;
424 
425  return comparisons;
426 }
427 
428 
429 
430 
431 /*==============================================================================
432  Trace function
433 ==============================================================================*/
434 
435 
436 void Trace(
437  __global uint* const colorDestination,
438  __global float* const depthDestination,
439  float const spread,
440  __global void const* const heap,
441  __read_only image3d_t image,
442  uint const root,
443  float3 const pp,
444  float3 const ray
445 )
446 {
447  float const epsilon = 1e-6;
448 
449  uint color = 0;
450  float3 projected = pp;
451  float lambda = Project( &projected, ray );
452 
453  if ( isfinite( lambda ) ) {
454 
455  int origin = 7;
456  float3 remaining = 1 - projected;
457  float3 delta = ray;
458  if ( delta.x < 0 ) {
459 
460  origin ^= 1;
461  remaining.x = projected.x;
462  delta.x = -delta.x;
463  }
464  if ( delta.y < 0 ) {
465 
466  origin ^= 2;
467  remaining.y = projected.y;
468  delta.y = -delta.y;
469  }
470  if ( delta.z < 0 ) {
471 
472  origin ^= 4;
473  remaining.z = projected.z;
474  delta.z = -delta.z;
475  }
476  delta = max( delta, epsilon );
477 
478 #if ( SETTING_MIPMAP != 0 )
479  float3 const mipmapCoordinate = fabs( pp - projected ) + remaining;
480  float const spreadInverse = 0.125 / spread; // 0.5 because we want to compare to the size at the *next* depth, and 0.25 because we draw 4x4x4 textures
481 #endif // SETTING_MIPMAP
482 
483  float3 remainingMaximum = 1 / delta;
484  remaining *= remainingMaximum;
485 
486  uint3 coordinate = ( uint3 )( 0, 0, 0 );
487  int coordinateDepth = 0;
488 
489  __global struct Node const* pNode = ( __global struct Node const* )( ( __global uchar const* )heap + root );
490 
491  __global struct Node const* childrenStack[ SETTING_TRACE_STACK_SIZE ];
492  int depth = 0;
493 
494  for ( ; ; ) {
495 
496  // ==== align the tree depth with the coordinate depth ====
497 
498  // descend until we reach the coordinate depth or a leaf
499  while ( ( depth < coordinateDepth ) && ( pNode->children != OPENCL_NULL ) ) {
500 
501 #if ( SETTING_MIPMAP != 0 )
502  uint3 const shiftedCoordinate = ( coordinate >> ( coordinateDepth - depth ) );
503 #if 0
504  float3 const mipmapDistance = ldexp( mipmapCoordinate, depth ) - ( float3 )( shiftedCoordinate.x, shiftedCoordinate.y, shiftedCoordinate.z );
505 #else // 0/1
506  float const sideLengthInverse = ( 1 << depth );
507  float3 const mipmapDistance = mipmapCoordinate * sideLengthInverse - ( float3 )( shiftedCoordinate.x, shiftedCoordinate.y, shiftedCoordinate.z );
508 #endif // 0/1
509  if ( any( mipmapDistance >= spreadInverse ) )
510  break; // **NOTE: break
511 #endif // SETTING_MIPMAP
512 
513  pNode = ( __global struct Node const* )( ( __global uchar const* )heap + pNode->children );
514  childrenStack[ depth ] = pNode;
515  ++depth;
516 
517  uint const mask = ( 1 << ( coordinateDepth - depth ) );
518  pNode += origin ^ (
519  ( ( coordinate.x & mask ) ? 1 : 0 ) |
520  ( ( coordinate.y & mask ) ? 2 : 0 ) |
521  ( ( coordinate.z & mask ) ? 4 : 0 )
522  );
523  }
524 
525  // shrink coordinate depth to match tree depth
526  while ( depth < coordinateDepth ) {
527 
528 #if 0
529  int3 const comparison = ( ( coordinate & 1 ) != 0 );
530  remaining += as_float3( comparison & as_int3( remainingMaximum ) );
531 #else // 0/1
532  if ( coordinate.x & 1 )
533  remaining.x += remainingMaximum.x;
534  if ( coordinate.y & 1 )
535  remaining.y += remainingMaximum.y;
536  if ( coordinate.z & 1 )
537  remaining.z += remainingMaximum.z;
538 #endif // 0/1
539 
540  remainingMaximum *= 2;
541 
542  coordinate /= 2;
543  --coordinateDepth;
544  }
545 
546  // ==== descend until we reach a leaf node ====
547 
548  while ( pNode->children != OPENCL_NULL ) {
549 
550 #if ( SETTING_MIPMAP != 0 )
551 #if 0
552  float3 const mipmapDistance = ldexp( mipmapCoordinate, depth ) - ( float3 )( coordinate.x, coordinate.y, coordinate.z );
553 #else // 0/1
554  float const sideLengthInverse = ( 1 << depth );
555  float3 const mipmapDistance = mipmapCoordinate * sideLengthInverse - ( float3 )( coordinate.x, coordinate.y, coordinate.z );
556 #endif // 0/1
557  if ( any( mipmapDistance >= spreadInverse ) )
558  break; // **NOTE: break
559 #endif // SETTING_MIPMAP
560 
561  pNode = ( __global struct Node const* )( ( __global uchar const* )heap + pNode->children );
562  childrenStack[ depth ] = pNode;
563  ++depth;
564 
565  coordinate *= 2;
566  ++coordinateDepth;
567 
568  remainingMaximum /= 2;
569 
570 #if 0
571  int3 const comparison = ( remaining >= remainingMaximum );
572  remaining -= as_float3( comparison & as_int3( remainingMaximum ) );
573  coordinate ^= ( comparison & 1 );
574 #else // 0/1
575  if ( remaining.x >= remainingMaximum.x ) {
576 
577  remaining.x -= remainingMaximum.x;
578  coordinate.x ^= 1;
579  }
580  if ( remaining.y >= remainingMaximum.y ) {
581 
582  remaining.y -= remainingMaximum.y;
583  coordinate.y ^= 1;
584  }
585  if ( remaining.z >= remainingMaximum.z ) {
586 
587  remaining.z -= remainingMaximum.z;
588  coordinate.z ^= 1;
589  }
590 #endif // 0/1
591 
592  pNode += origin ^ (
593  ( ( coordinate.x & 1 ) ? 1 : 0 ) |
594  ( ( coordinate.y & 1 ) ? 2 : 0 ) |
595  ( ( coordinate.z & 1 ) ? 4 : 0 )
596  );
597  }
598 
599  // ==== step through the current voxel (texture) ====
600 
601  int comparisons = -1;
602  if ( pNode->texture != OPENCL_NULL ) {
603 
604  comparisons = TraceTexture( // returns -1 if drawn
605  &color,
606  &lambda,
607  image,
608  pNode->texture,
609  &remaining,
610  &remainingMaximum,
611  origin
612  );
613 
614  if ( comparisons < 0 )
615  goto main_loop_end_label; // **NOTE: goto
616  }
617  else {
618 
619  comparisons = (
620  ( ( remaining.x < remaining.y ) ? 1 : 0 ) |
621  ( ( remaining.y < remaining.z ) ? 2 : 0 ) |
622  ( ( remaining.z < remaining.x ) ? 4 : 0 )
623  );
624  switch( comparisons ) {
625 
626  case 0: case 7: // any are fine--fall through
627  case 1: case 3: {
628 
629  lambda += remaining.x;
630  remaining.y -= remaining.x;
631  remaining.z -= remaining.x;
632  remaining.x = remainingMaximum.x;
633 
634  break;
635  }
636 
637  case 2: case 6: {
638 
639  lambda += remaining.y;
640  remaining.x -= remaining.y;
641  remaining.z -= remaining.y;
642  remaining.y = remainingMaximum.y;
643 
644  break;
645  }
646 
647  case 4: case 5: {
648 
649  lambda += remaining.z;
650  remaining.x -= remaining.z;
651  remaining.y -= remaining.z;
652  remaining.z = remainingMaximum.z;
653 
654  break;
655  }
656  }
657  }
658 
659  // ==== step to the next adjacent voxel, ascending as necessary ====
660 
661  switch( comparisons ) {
662 
663  case 0: case 7: // any are fine--fall through
664  case 1: case 3: {
665 
666  if ( coordinate.x == 0 ) { // underflow indicates that we've left the octtree
667 
668  lambda = INFINITY;
669  goto main_loop_end_label; // **NOTE: goto
670  }
671 
672  --coordinate.x;
673  for ( uint mask = 1; ( coordinate.x & mask ) != 0; mask *= 2, --depth );
674  // **NOTE: pNode will be updated below
675 
676  break;
677  }
678 
679  case 2: case 6: {
680 
681  if ( coordinate.y == 0 ) { // underflow indicates that we've left the octtree
682 
683  lambda = INFINITY;
684  goto main_loop_end_label; // **NOTE: goto
685  }
686 
687  --coordinate.y;
688  for ( uint mask = 1; ( coordinate.y & mask ) != 0; mask *= 2, --depth );
689  // **NOTE: pNode will be updated below
690 
691  break;
692  }
693 
694  case 4: case 5: {
695 
696  if ( coordinate.z == 0 ) { // underflow indicates that we've left the octtree
697 
698  lambda = INFINITY;
699  goto main_loop_end_label; // **NOTE: goto
700  }
701 
702  --coordinate.z;
703  for ( uint mask = 1; ( coordinate.z & mask ) != 0; mask *= 2, --depth );
704  // **NOTE: pNode will be updated below
705 
706  break;
707  }
708  }
709 
710  // move pNode to the correct position at the current depth
711  pNode = childrenStack[ depth - 1 ];
712  uint const mask = ( 1 << ( coordinateDepth - depth ) );
713  pNode += origin ^ (
714  ( ( coordinate.x & mask ) ? 1 : 0 ) |
715  ( ( coordinate.y & mask ) ? 2 : 0 ) |
716  ( ( coordinate.z & mask ) ? 4 : 0 )
717  );
718  }
719 main_loop_end_label:;
720  }
721 
722  *colorDestination = color;
723  *depthDestination = lambda;
724 }
725 
726 
727 
728 
729 /*==============================================================================
730  TraceScreen function
731 ==============================================================================*/
732 
733 
734 __kernel void TraceScreen(
735  __global uint* const colors,
736  __global float* const depths,
737  uint2 const dimension,
738  uint const colorsPitch,
739  uint const depthsPitch,
740  float2 const fieldOfView,
741  float const spread,
742  __global void const* const heap,
743  __read_only image3d_t image,
744  uint const root,
745  float3 const xx,
746  float3 const yy,
747  float3 const zz,
748  float3 const pp
749 )
750 {
751  uint const index = get_global_id( 0 );
752  if ( index < dimension.x * dimension.y ) {
753 
754  /*
755  To maximize the chances of threads in the same warp accessing the
756  same data (by working on nearby rays), we work in Morton Z-order.
757  */
758  uint2 const coordinate = InverseMortonZOrder( index, dimension );
759 
760  float3 const spreadX = ( 2 * fieldOfView.x / dimension.x ) * xx;
761  float3 const spreadY = ( 2 * fieldOfView.y / dimension.y ) * yy;
762 
763  float3 const ray = (
764  ( coordinate.x + 0.5 - 0.5 * dimension.x ) * spreadX -
765  ( coordinate.y + 0.5 - 0.5 * dimension.y ) * spreadY -
766  zz
767  );
768 
769  __global uint* const colorDestination = ( __global uint* )( ( __global uchar* )colors + coordinate.y * colorsPitch + coordinate.x * sizeof( uint ) );
770  __global float* const depthDestination = ( __global float* )( ( __global uchar* )depths + coordinate.y * depthsPitch + coordinate.x * sizeof( float ) );
771  Trace( colorDestination, depthDestination, spread, heap, image, root, pp, ray );
772  }
773 }
Generated on Thu Jun 20 2013 11:05:27 for Voxel by   doxygen 1.8.1.2