API tools faq
paste
Advertisement
Guest User

Untitled

a guest
Jun 11th, 2010
155
0
Never
Add comment
Not a member of Pastebin yet? Sign Up, it unlocks many cool features!
C 40.02 KB | None | 0 0
raw download clone embed print report
  1. #pragma OPENCL EXTENSION cl_khr_fp64 : enable
  2. #pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable
  3. typedef unsigned char       U8;
  4. typedef unsigned short      U16;
  5. typedef unsigned int        U32;
  6. typedef signed char         S8;
  7. typedef signed short        S16;
  8. typedef signed int          S32;
  9. typedef float               F32;
  10. typedef double              F64;
  11. typedef void                (*FuncPtr)(void);
  12. typedef unsigned long       U64;
  13. typedef signed long         S64;
  14. typedef S64                 SPTR;
  15. typedef U64                 UPTR;
  16. inline U64    doubleToBits    (F64 a)         { return *(U64*)&a; }
  17. inline F64    bitsToDouble    (U64 a)         { return *(F64*)&a; }
  18. inline U32    floatToBits     (F32 a)         { return as_int(a); }
  19. inline F32    bitsToFloat     (U32 a)         { return as_float(a); }
  20. inline F32    scale           (F32 a, int b)  { return a * (exp2((float)b)); }
  21. inline int    popc8           (U32 mask);
  22. inline int    popc16          (U32 mask);
  23. inline int    popc32          (U32 mask);
  24. typedef double2 Vec2d;
  25. typedef int2 Vec2i;
  26. typedef float2 Vec2f;
  27. typedef struct Vec3d_
  28. {
  29.     double          x;
  30.     double          y;
  31.     double          z;
  32. } Vec3d;
  33. typedef struct Vec3i_
  34. {
  35.     int             x;
  36.     int             y;
  37.     int             z;
  38. } Vec3i;
  39. typedef struct Vec3f_
  40. {
  41.     float           x;
  42.     float           y;
  43.     float           z;
  44. } Vec3f;
  45. typedef double4 Vec4d;
  46. typedef int4 Vec4i;
  47. typedef float4 Vec4f;
  48. typedef struct Mat2f_
  49. {
  50.     float   m00, m10;
  51.     float   m01, m11;
  52. } Mat2f;
  53. typedef struct Mat3f_
  54. {
  55.     float   m00, m10, m20;
  56.     float   m01, m11, m21;
  57.     float   m02, m12, m22;
  58. } Mat3f;
  59. typedef struct Mat4f_
  60. {
  61.     float   m00, m10, m20, m30;
  62.     float   m01, m11, m21, m31;
  63.     float   m02, m12, m22, m32;
  64.     float   m03, m13, m23, m33;
  65. } Mat4f;
  66. __constant int c_popc8LUT[] =
  67. {
  68.     0, 1, 1, 2, 1, 2, 2, 3, 1, 2, 2, 3, 2, 3, 3, 4,
  69.     1, 2, 2, 3, 2, 3, 3, 4, 2, 3, 3, 4, 3, 4, 4, 5,
  70.     1, 2, 2, 3, 2, 3, 3, 4, 2, 3, 3, 4, 3, 4, 4, 5,
  71.     2, 3, 3, 4, 3, 4, 4, 5, 3, 4, 4, 5, 4, 5, 5, 6,
  72.     1, 2, 2, 3, 2, 3, 3, 4, 2, 3, 3, 4, 3, 4, 4, 5,
  73.     2, 3, 3, 4, 3, 4, 4, 5, 3, 4, 4, 5, 4, 5, 5, 6,
  74.     2, 3, 3, 4, 3, 4, 4, 5, 3, 4, 4, 5, 4, 5, 5, 6,
  75.     3, 4, 4, 5, 4, 5, 5, 6, 4, 5, 5, 6, 5, 6, 6, 7,
  76.     1, 2, 2, 3, 2, 3, 3, 4, 2, 3, 3, 4, 3, 4, 4, 5,
  77.     2, 3, 3, 4, 3, 4, 4, 5, 3, 4, 4, 5, 4, 5, 5, 6,
  78.     2, 3, 3, 4, 3, 4, 4, 5, 3, 4, 4, 5, 4, 5, 5, 6,
  79.     3, 4, 4, 5, 4, 5, 5, 6, 4, 5, 5, 6, 5, 6, 6, 7,
  80.     2, 3, 3, 4, 3, 4, 4, 5, 3, 4, 4, 5, 4, 5, 5, 6,
  81.     3, 4, 4, 5, 4, 5, 5, 6, 4, 5, 5, 6, 5, 6, 6, 7,
  82.     3, 4, 4, 5, 4, 5, 5, 6, 4, 5, 5, 6, 5, 6, 6, 7,
  83.     4, 5, 5, 6, 5, 6, 6, 7, 5, 6, 6, 7, 6, 7, 7, 8,
  84. };
  85. inline int popc8(U32 mask)
  86. {
  87.     ((void)0);
  88.     return c_popc8LUT[mask];
  89. }
  90. inline int popc16(U32 mask)
  91. {
  92.     ((void)0);
  93.     return c_popc8LUT[mask & 0xffu] + c_popc8LUT[mask >> 8];
  94. }
  95. inline int popc32(U32 mask)
  96. {
  97.     int result = c_popc8LUT[mask & 0xffu];
  98.     result += c_popc8LUT[(mask >> 8) & 0xffu];
  99.     result += c_popc8LUT[(mask >> 16) & 0xffu];
  100.     result += c_popc8LUT[mask >> 24];
  101.     return result;
  102. }
  103.  
  104. enum RenderFlags
  105. {
  106.     RenderFlags_CoarsePass              = 1 << 0,
  107.     RenderFlags_UseCoarseData           = 1 << 1,
  108.     RenderFlags_VisualizeIterations     = 1 << 2,
  109.     RenderFlags_VisualizeRaycastLevel   = 1 << 3,
  110. };
  111. enum AttachSlot
  112. {
  113.     AttachSlot_Contour = 0,
  114.     AttachSlot_Attribute,
  115.     AttachSlot_AO,
  116.     AttachSlot_Max
  117. };
  118. enum
  119. {
  120.     PageBytesLog2   = 13,
  121.     PageBytes       = 1 << PageBytesLog2,
  122.     PageSizeLog2    = PageBytesLog2 - 2,
  123.     PageSize        = 1 << PageSizeLog2,
  124.     TrunksPerBlock  = 512,
  125. };
  126. enum FindMode
  127. {
  128.     FindMode_Load = 0,
  129.     FindMode_Build,
  130.     FindMode_LoadOrBuild,
  131.     FindMode_Unload,
  132.     FindMode_UnloadDeepest,
  133.     FindMode_Max
  134. };
  135. enum BlockInfo
  136. {
  137.     BlockInfo_SliceID = 0,
  138.     BlockInfo_IndexInSlice,
  139.     BlockInfo_BlockPtr,
  140.     BlockInfo_NumAttach,
  141.     BlockInfo_End
  142. };
  143. enum AttachInfo
  144. {
  145.     AttachInfo_Type = 0,
  146.     AttachInfo_Ptr,
  147.     AttachInfo_End
  148. };
  149. enum AttachType
  150. {
  151.     VoidAttach                  = 0,    
  152.     ColorNormalDXTAttach        = 1,    
  153.     ColorNormalPaletteAttach    = 3,    
  154.     BuildDataAttach             = 4,    
  155.     ContourAttach               = 6,    
  156.     ColorNormalCornerAttach     = 7,    
  157.     AOAttach                    = 10,  
  158.     AttachType_Max
  159. };
  160. typedef struct OctreeMatrices_
  161. {
  162.     Mat4f           viewportToCamera;
  163.     Mat4f           cameraToOctree;
  164.     Mat4f           octreeToWorld;
  165.     F32             pixelInOctree;      
  166.     Mat4f           worldToOctree;
  167.     Mat3f           octreeToWorldN;    
  168.     Vec3f           cameraPosition;    
  169.     Mat4f           octreeToViewport;
  170.     Mat4f           viewportToOctreeN;  
  171. }OctreeMatrices;
  172. typedef struct StandardInput_
  173. {
  174.     Vec2i           frameSize;
  175.     U32             flags;
  176.     S32             batchSize;          
  177.     S32             aaRays;            
  178.     F32             maxVoxelSize;      
  179.     F32             brightness;
  180.     S32             coarseSize;        
  181.     Vec2i           coarseFrameSize;    
  182.     S32             numPrimaryRays;    
  183.     S32             totalWork;          
  184.     OctreeMatrices  octreeMatrices;
  185. }StandardInput;
  186.  
  187. inline F32          fmaxf3          (F32 a, F32 b, F32 c)           { return fmax(fmax(a, b), c); }
  188. inline F32          fminf3          (F32 a, F32 b, F32 c)           { return fmin(fmin(a, b), c); }
  189. inline F32          smoothstep01    (F32 f)                         { return smoothstep(0.f, 1.f, f); }
  190. inline float4   scaleVec4         (const float4* a, const float4* b)  { float4 res = {(*a).x * (*b).x, (*a).y * (*b).y, (*a).z* (*b).z, (*a).w * (*b).w}; return res; }
  191. inline float4    mulMatVec4     (const Mat4f* m, const float4* v);
  192. inline float4    mulMatVec3     (const Mat3f* m, const float4* v);
  193. inline float4    mulMatVec3const(__constant Mat3f* m, const float4* v);
  194. inline Mat3f     extractMat3f   (__constant Mat4f* m);
  195. inline float4    fromABGR        (U32 abgr);
  196. inline U32       toABGR          (float4 v);
  197. inline void      jenkinsMix      (U32* a, U32* b, U32* c);
  198. inline float4    perpendicular   (const float4* v);
  199. inline float4    mulMatVec4     (const Mat4f* m, const float4* v)
  200. {
  201.     float4 res = {
  202.         (*m).m00 * (*v).x + (*m).m01 * (*v).y + (*m).m02 * (*v).z + (*m).m03 * (*v).w,
  203.         (*m).m10 * (*v).x + (*m).m11 * (*v).y + (*m).m12 * (*v).z + (*m).m13 * (*v).w,
  204.         (*m).m20 * (*v).x + (*m).m21 * (*v).y + (*m).m22 * (*v).z + (*m).m23 * (*v).w,
  205.         (*m).m30 * (*v).x + (*m).m31 * (*v).y + (*m).m32 * (*v).z + (*m).m33 * (*v).w};
  206.     return res;
  207. }
  208. inline float4    mulMat4Vec3        (__constant Mat4f* m, const float4* v)
  209. {
  210.     float4 res = {
  211.         (*m).m00 * (*v).x + (*m).m01 * (*v).y + (*m).m02 * (*v).z + (*m).m03,
  212.         (*m).m10 * (*v).x + (*m).m11 * (*v).y + (*m).m12 * (*v).z + (*m).m13,
  213.         (*m).m20 * (*v).x + (*m).m21 * (*v).y + (*m).m22 * (*v).z + (*m).m23,
  214.         0.0f};
  215.     return res;
  216. }
  217. inline float4    mulMatVec3     (const Mat3f* m, const float4* v)
  218. {
  219.     float4 res = {
  220.         (*m).m00 * (*v).x + (*m).m01 * (*v).y + (*m).m02 * (*v).z,
  221.         (*m).m10 * (*v).x + (*m).m11 * (*v).y + (*m).m12 * (*v).z,
  222.         (*m).m20 * (*v).x + (*m).m21 * (*v).y + (*m).m22 * (*v).z,
  223.         0.0f};
  224.     return res;
  225. }
  226. inline float4    mulMatVec3local    (const Mat3f* m, __local const float4* v)
  227. {
  228.     float4 res ={
  229.         (*m).m00 * (*v).x + (*m).m01 * (*v).y + (*m).m02 * (*v).z,
  230.         (*m).m10 * (*v).x + (*m).m11 * (*v).y + (*m).m12 * (*v).z,
  231.         (*m).m20 * (*v).x + (*m).m21 * (*v).y + (*m).m22 * (*v).z,
  232.         0.0f};
  233.     return res;
  234. }
  235. inline float4    mulMatVec3const     (__constant Mat3f* m, const float4* v)
  236. {
  237.     float4 res = {
  238.         (*m).m00 * (*v).x + (*m).m01 * (*v).y + (*m).m02 * (*v).z,
  239.         (*m).m10 * (*v).x + (*m).m11 * (*v).y + (*m).m12 * (*v).z,
  240.         (*m).m20 * (*v).x + (*m).m21 * (*v).y + (*m).m22 * (*v).z,
  241.         0.0f};
  242.     return res;
  243. }
  244. inline Mat3f extractMat3f(__constant Mat4f* m)
  245. {
  246.     Mat3f n;
  247.     n.m00 = (*m).m00; n.m01 = (*m).m01; n.m02 = (*m).m02;
  248.     n.m10 = (*m).m10; n.m11 = (*m).m11; n.m12 = (*m).m12;
  249.     n.m20 = (*m).m20; n.m21 = (*m).m21; n.m22 = (*m).m22;
  250.     return n;
  251. }
  252. inline float4 fromABGR(U32 abgr)
  253. {
  254.     float4 res = {
  255.         (F32)(abgr & 0xFF),
  256.         (F32)((abgr >> 8) & 0xFF),
  257.         (F32)((abgr >> 16) & 0xFF),
  258.         (F32)(abgr >> 24)};
  259.     return res;
  260. }
  261. inline U32 toABGR(float4 v)
  262. {
  263.     v.x = fmin(fmax(v.x, 0.f), 255.f);
  264.     v.y = fmin(fmax(v.y, 0.f), 255.f);
  265.     v.z = fmin(fmax(v.z, 0.f), 255.f);
  266.     v.w = fmin(fmax(v.w, 0.f), 255.f);
  267.     U32 ir = (U32)(v.x);
  268.     U32 ig = (U32)(v.y);
  269.     U32 ib = (U32)(v.z);
  270.     U32 ia = (U32)(v.w);
  271.     return ir | (ig << 8) | (ib << 16) | (ia << 24);
  272. }
  273. inline void jenkinsMix(U32* a, U32* b, U32* c)
  274. {
  275.     *a -= *b; *a -= *c; *a ^= (*c>>13);
  276.     *b -= *c; *b -= *a; *b ^= (*a<<8);
  277.     *c -= *a; *c -= *b; *c ^= (*b>>13);
  278.     *a -= *b; *a -= *c; *a ^= (*c>>12);
  279.     *b -= *c; *b -= *a; *b ^= (*a<<16);
  280.     *c -= *a; *c -= *b; *c ^= (*b>>5);
  281.     *a -= *b; *a -= *c; *a ^= (*c>>3);
  282.     *b -= *c; *b -= *a; *b ^= (*a<<10);
  283.     *c -= *a; *c -= *b; *c ^= (*b>>15);
  284. }
  285. inline float4 perpendicular(const float4* v)
  286. {
  287.     float4 res;
  288.     float vmin = fmin(fabs((*v).x), fmin(fabs((*v).y), fabs((*v).z)));
  289.     if (vmin == fabs((*v).x))
  290.         res = (float4)(0.0f, (*v).z, -(*v).y, 0.0f);
  291.     else if (vmin == fabs((*v).y))
  292.         res = (float4)(-(*v).z, 0.0f, (*v).x, 0.0f);
  293.     else
  294.         res = (float4)((*v).y, -(*v).x, 0.0f, 0.0f);
  295.     return res;
  296. }
  297.  
  298. typedef struct Ray_
  299. {
  300.     float4      orig;  
  301.     float4      dir;   
  302. }Ray;
  303. typedef struct CastResult_
  304. {
  305.    
  306.     float4      pos;       
  307.     int         iter;
  308.     __global S32* node;    
  309.     int         childIdx;
  310.     int         stackPtr;
  311. }CastResult;
  312. typedef struct CastStack_
  313. {
  314.     F32         tmax[23 + 1];
  315.     intptr_t    node[23 + 1];
  316. }CastStack;
  317. __global S32* readStack(CastStack* castStack, int idx, F32* tmax)
  318. {
  319.     *tmax = (*castStack).tmax[idx];
  320.     intptr_t e = (*castStack).node[idx];
  321.     return (__global S32*)e;
  322. };
  323. void writeStack(CastStack* castStack, int idx, __global S32* node, F32 tmax)
  324. {
  325.     (*castStack).tmax[idx] = tmax;
  326.     (*castStack).node[idx] = (intptr_t) node;
  327. };
  328. void castRayLocal(__global int* g_rootNode, unsigned int rootNodeOfs, CastResult* res, CastStack* stack, volatile __local Ray* ray)
  329. {
  330.     const float epsilon = exp2((float)-23);
  331.     float ray_orig_sz = (*ray).orig.w;
  332.     int iter = 0;
  333.    
  334.     if (fabs((*ray).dir.x) < epsilon) (*ray).dir.x = copysign(epsilon, (*ray).dir.x);
  335.     if (fabs((*ray).dir.y) < epsilon) (*ray).dir.y = copysign(epsilon, (*ray).dir.y);
  336.     if (fabs((*ray).dir.z) < epsilon) (*ray).dir.z = copysign(epsilon, (*ray).dir.z);
  337.    
  338.    
  339.     float tx_coef = 1.0f / -fabs((*ray).dir.x);
  340.     float ty_coef = 1.0f / -fabs((*ray).dir.y);
  341.     float tz_coef = 1.0f / -fabs((*ray).dir.z);
  342.     float tx_bias = tx_coef * (*ray).orig.x;
  343.     float ty_bias = ty_coef * (*ray).orig.y;
  344.     float tz_bias = tz_coef * (*ray).orig.z;
  345.    
  346.    
  347.     int octant_mask = 7;
  348.     if ((*ray).dir.x > 0.0f) octant_mask ^= 1, tx_bias = 3.0f * tx_coef - tx_bias;
  349.     if ((*ray).dir.y > 0.0f) octant_mask ^= 2, ty_bias = 3.0f * ty_coef - ty_bias;
  350.     if ((*ray).dir.z > 0.0f) octant_mask ^= 4, tz_bias = 3.0f * tz_coef - tz_bias;
  351.    
  352.    
  353.     float t_min = fmax(fmax(2.0f * tx_coef - tx_bias, 2.0f * ty_coef - ty_bias), 2.0f * tz_coef - tz_bias);
  354.     float t_max = fmin(fmin(tx_coef - tx_bias, ty_coef - ty_bias), tz_coef - tz_bias);
  355.     float h = t_max;
  356.     t_min = fmax(t_min, 0.0f);
  357.     t_max = fmin(t_max, 1.0f);
  358.    
  359.     __global int*   parent  = g_rootNode + rootNodeOfs;
  360.     int2   child_descriptor = {0, 0};
  361.     int    idx              = 0;
  362.     float4 pos              = {1.0f, 1.0f, 1.0f, 0.0f};
  363.     int    scale            = 23 - 1;
  364.     float  scale_exp2       = 0.5f;
  365.     if (1.5f * tx_coef - tx_bias > t_min) idx ^= 1, pos.x = 1.5f;
  366.     if (1.5f * ty_coef - ty_bias > t_min) idx ^= 2, pos.y = 1.5f;
  367.     if (1.5f * tz_coef - tz_bias > t_min) idx ^= 4, pos.z = 1.5f;
  368.    
  369.    
  370.    bool bBreak;
  371.    
  372.     while (scale < 23)
  373.     {
  374.        
  375.        
  376.         iter++;
  377.         if (iter > 10000)
  378.             break;
  379.        
  380.         if (child_descriptor.x == 0)
  381.         {
  382.             child_descriptor = *((__global int2*)(parent));
  383.            
  384.         }
  385.        
  386.        
  387.         float tx_corner = pos.x * tx_coef - tx_bias;
  388.         float ty_corner = pos.y * ty_coef - ty_bias;
  389.         float tz_corner = pos.z * tz_coef - tz_bias;
  390.         float tc_max = fmin(fmin(tx_corner, ty_corner), tz_corner);
  391.        
  392.        
  393.         int child_shift = idx ^ octant_mask;
  394.         int child_masks = child_descriptor.x << child_shift;
  395.         if ((child_masks & 0x8000) != 0 && t_min <= t_max)
  396.         {
  397.            
  398.            
  399.             if (tc_max * (*ray).dir.w + ray_orig_sz >= scale_exp2)
  400.                 break;
  401.            
  402.            
  403.            
  404.            
  405.            
  406.             float tv_max = fmin(t_max, tc_max);
  407.             float half = scale_exp2 * 0.5f;
  408.             float tx_center = half * tx_coef + tx_corner;
  409.             float ty_center = half * ty_coef + ty_corner;
  410.             float tz_center = half * tz_coef + tz_corner;
  411.            
  412.             int contour_mask = child_descriptor.y << child_shift;
  413.             if ((contour_mask & 0x80) != 0)
  414.             {
  415.                
  416.                 int   ofs    = (unsigned int)child_descriptor.y >> 8;      
  417.                 int   value  = parent[ofs + popc8(contour_mask & 0x7F)];    
  418.                
  419.                 float cthick = (float)(unsigned int)value * scale_exp2 * 0.75f;
  420.                 float cpos   = (float)(value << 7) * scale_exp2 * 1.5f;    
  421.                 float cdirx  = (float)(value << 14) * (*ray).dir.x;            
  422.                 float cdiry  = (float)(value << 20) * (*ray).dir.y;            
  423.                 float cdirz  = (float)(value << 26) * (*ray).dir.z;            
  424.                 float tcoef  = 1.0f / (cdirx + cdiry + cdirz);
  425.                 float tavg   = tx_center * cdirx + ty_center * cdiry + tz_center * cdirz + cpos;
  426.                 float tdiff  = cthick * tcoef;
  427.                 t_min  = fmax(t_min,  tcoef * tavg - fabs(tdiff));
  428.                 tv_max = fmin(tv_max, tcoef * tavg + fabs(tdiff));
  429.             }
  430.            
  431.            
  432.             if (t_min <= tv_max)
  433.             {
  434.                
  435.                
  436.                 if ((child_masks & 0x0080) == 0)
  437.                  bBreak = true;        
  438.            else
  439.                 {
  440.                     if (tc_max < h)
  441.                     {
  442.                        
  443.                         writeStack(stack, scale, parent, t_max);
  444.                        
  445.                     }
  446.                     h = tc_max;
  447.                    
  448.                     int ofs = (unsigned int)child_descriptor.x >> 17;
  449.                     if ((child_descriptor.x & 0x10000) != 0)
  450.                     {
  451.                         ofs = parent[ofs * 2];
  452.                        
  453.                     }
  454.                     ofs += popc8(child_masks & 0x7F);
  455.                     parent += ofs * 2;
  456.                    
  457.                     idx = 0;
  458.                     scale--;
  459.                     scale_exp2 = half;
  460.                     if (tx_center > t_min) idx ^= 1, pos.x += scale_exp2;
  461.                     if (ty_center > t_min) idx ^= 2, pos.y += scale_exp2;
  462.                     if (tz_center > t_min) idx ^= 4, pos.z += scale_exp2;
  463.                    
  464.                     t_max = tv_max;
  465.                     child_descriptor.x = 0;
  466.                     //continue;
  467.                 }
  468.             }
  469.         }
  470.        
  471.         if(bBreak)
  472.          break;
  473.        
  474.        
  475.         int step_mask = 0;
  476.         if (tx_corner <= tc_max) step_mask ^= 1, pos.x -= scale_exp2;
  477.         if (ty_corner <= tc_max) step_mask ^= 2, pos.y -= scale_exp2;
  478.         if (tz_corner <= tc_max) step_mask ^= 4, pos.z -= scale_exp2;
  479.        
  480.         t_min = tc_max;
  481.         idx ^= step_mask;
  482.        
  483.         if ((idx & step_mask) != 0)
  484.         {
  485.            
  486.            
  487.            
  488.            
  489.             unsigned int differing_bits = 0;
  490.             if ((step_mask & 1) != 0) differing_bits |= as_int(pos.x) ^ as_int(pos.x + scale_exp2);
  491.             if ((step_mask & 2) != 0) differing_bits |= as_int(pos.y) ^ as_int(pos.y + scale_exp2);
  492.             if ((step_mask & 4) != 0) differing_bits |= as_int(pos.z) ^ as_int(pos.z + scale_exp2);
  493.             scale = (as_int((float)differing_bits) >> 23) - 127;
  494.             scale_exp2 = as_float((scale - 23 + 127) << 23);
  495.            
  496.             parent = readStack(stack, scale, &t_max);      
  497.            
  498.            
  499.             int shx = as_int(pos.x) >> scale;
  500.             int shy = as_int(pos.y) >> scale;
  501.             int shz = as_int(pos.z) >> scale;
  502.             pos.x = as_float(shx << scale);
  503.             pos.y = as_float(shy << scale);
  504.             pos.z = as_float(shz << scale);
  505.             idx  = (shx & 1) | ((shy & 1) << 1) | ((shz & 1) << 2);
  506.            
  507.             h = 0.0f;
  508.             child_descriptor.x = 0;
  509.         }
  510.        
  511.     }
  512.    
  513.     if (scale >= 23 || iter > 10000)
  514.     {
  515.         t_min = 2.0f;
  516.     }
  517.    
  518.     if ((octant_mask & 1) == 0) pos.x = 3.0f - scale_exp2 - pos.x;
  519.     if ((octant_mask & 2) == 0) pos.y = 3.0f - scale_exp2 - pos.y;
  520.     if ((octant_mask & 4) == 0) pos.z = 3.0f - scale_exp2 - pos.z;
  521.    
  522.     (*res).pos.w = t_min;
  523.     (*res).iter = iter;
  524.     (*res).pos.x = fmin(fmax((*ray).orig.x + t_min * (*ray).dir.x, pos.x + epsilon), pos.x + scale_exp2 - epsilon);
  525.     (*res).pos.y = fmin(fmax((*ray).orig.y + t_min * (*ray).dir.y, pos.y + epsilon), pos.y + scale_exp2 - epsilon);
  526.     (*res).pos.z = fmin(fmax((*ray).orig.z + t_min * (*ray).dir.z, pos.z + epsilon), pos.z + scale_exp2 - epsilon);
  527.     (*res).node = parent;
  528.     (*res).childIdx = idx ^ octant_mask ^ 7;
  529.     (*res).stackPtr = scale;
  530. }
  531. void castRay(__global int* g_rootNode, unsigned int rootNodeOfs, CastResult* res, CastStack* stack, volatile Ray* ray)
  532. {
  533.     const float epsilon = exp2((float)-23);
  534.     float ray_orig_sz = (*ray).orig.w;
  535.     int iter = 0;
  536.    
  537.     if (fabs((*ray).dir.x) < epsilon) (*ray).dir.x = copysign(epsilon, (*ray).dir.x);
  538.     if (fabs((*ray).dir.y) < epsilon) (*ray).dir.y = copysign(epsilon, (*ray).dir.y);
  539.     if (fabs((*ray).dir.z) < epsilon) (*ray).dir.z = copysign(epsilon, (*ray).dir.z);
  540.    
  541.    
  542.     float tx_coef = 1.0f / -fabs((*ray).dir.x);
  543.     float ty_coef = 1.0f / -fabs((*ray).dir.y);
  544.     float tz_coef = 1.0f / -fabs((*ray).dir.z);
  545.     float tx_bias = tx_coef * (*ray).orig.x;
  546.     float ty_bias = ty_coef * (*ray).orig.y;
  547.     float tz_bias = tz_coef * (*ray).orig.z;
  548.    
  549.    
  550.     int octant_mask = 7;
  551.     if ((*ray).dir.x > 0.0f) octant_mask ^= 1, tx_bias = 3.0f * tx_coef - tx_bias;
  552.     if ((*ray).dir.y > 0.0f) octant_mask ^= 2, ty_bias = 3.0f * ty_coef - ty_bias;
  553.     if ((*ray).dir.z > 0.0f) octant_mask ^= 4, tz_bias = 3.0f * tz_coef - tz_bias;
  554.    
  555.    
  556.     float t_min = fmax(fmax(2.0f * tx_coef - tx_bias, 2.0f * ty_coef - ty_bias), 2.0f * tz_coef - tz_bias);
  557.     float t_max = fmin(fmin(tx_coef - tx_bias, ty_coef - ty_bias), tz_coef - tz_bias);
  558.     float h = t_max;
  559.     t_min = fmax(t_min, 0.0f);
  560.     t_max = fmin(t_max, 1.0f);
  561.    
  562.     __global int*   parent          = g_rootNode + rootNodeOfs;        
  563.     int2   child_descriptor = {0, 0};
  564.     int    idx              = 0;
  565.     float4 pos              = {1.0f, 1.0f, 1.0f, 0.0f};
  566.     int    scale            = 23 - 1;
  567.     float  scale_exp2       = 0.5f;
  568.     if (1.5f * tx_coef - tx_bias > t_min) idx ^= 1, pos.x = 1.5f;
  569.     if (1.5f * ty_coef - ty_bias > t_min) idx ^= 2, pos.y = 1.5f;
  570.     if (1.5f * tz_coef - tz_bias > t_min) idx ^= 4, pos.z = 1.5f;
  571.    
  572.    
  573.    
  574.     while (scale < 23)
  575.     {
  576.        
  577.        
  578.         iter++;
  579.         if (iter > 10000)
  580.             break;
  581.        
  582.         if (child_descriptor.x == 0)
  583.         {
  584.             child_descriptor = *(__global int2*)(parent);
  585.            
  586.         }
  587.        
  588.        
  589.         float tx_corner = pos.x * tx_coef - tx_bias;
  590.         float ty_corner = pos.y * ty_coef - ty_bias;
  591.         float tz_corner = pos.z * tz_coef - tz_bias;
  592.         float tc_max = fmin(fmin(tx_corner, ty_corner), tz_corner);
  593.        
  594.        
  595.         int child_shift = idx ^ octant_mask;
  596.         int child_masks = child_descriptor.x << child_shift;
  597.         if ((child_masks & 0x8000) != 0 && t_min <= t_max)
  598.         {
  599.            
  600.            
  601.             if (tc_max * (*ray).dir.w + ray_orig_sz >= scale_exp2)
  602.                 break;
  603.            
  604.            
  605.            
  606.            
  607.            
  608.             float tv_max = fmin(t_max, tc_max);
  609.             float half = scale_exp2 * 0.5f;
  610.             float tx_center = half * tx_coef + tx_corner;
  611.             float ty_center = half * ty_coef + ty_corner;
  612.             float tz_center = half * tz_coef + tz_corner;
  613.            
  614.             int contour_mask = child_descriptor.y << child_shift;
  615.             if ((contour_mask & 0x80) != 0)
  616.             {
  617.                
  618.                 int   ofs    = (unsigned int)child_descriptor.y >> 8;      
  619.                 int   value  = parent[ofs + popc8(contour_mask & 0x7F)];    
  620.                
  621.                 float cthick = (float)(unsigned int)value * scale_exp2 * 0.75f;
  622.                 float cpos   = (float)(value << 7) * scale_exp2 * 1.5f;    
  623.                 float cdirx  = (float)(value << 14) * (*ray).dir.x;            
  624.                 float cdiry  = (float)(value << 20) * (*ray).dir.y;            
  625.                 float cdirz  = (float)(value << 26) * (*ray).dir.z;            
  626.                 float tcoef  = 1.0f / (cdirx + cdiry + cdirz);
  627.                 float tavg   = tx_center * cdirx + ty_center * cdiry + tz_center * cdirz + cpos;
  628.                 float tdiff  = cthick * tcoef;
  629.                 t_min  = fmax(t_min,  tcoef * tavg - fabs(tdiff));
  630.                 tv_max = fmin(tv_max, tcoef * tavg + fabs(tdiff));
  631.             }
  632.            
  633.            
  634.             if (t_min <= tv_max)
  635.             {
  636.                
  637.                
  638.                 if ((child_masks & 0x0080) == 0)
  639.                     break;
  640.                
  641.                
  642.                
  643.                
  644.                 if (tc_max < h)
  645.                 {
  646.                    
  647.                     writeStack(stack, scale, parent, t_max);
  648.                    
  649.                 }
  650.                 h = tc_max;
  651.                
  652.                 int ofs = (unsigned int)child_descriptor.x >> 17;
  653.                 if ((child_descriptor.x & 0x10000) != 0)
  654.                 {
  655.                     ofs = parent[ofs * 2];
  656.                    
  657.                 }
  658.                 ofs += popc8(child_masks & 0x7F);
  659.                 parent += ofs * 2;
  660.                
  661.                 idx = 0;
  662.                 scale--;
  663.                 scale_exp2 = half;
  664.                 if (tx_center > t_min) idx ^= 1, pos.x += scale_exp2;
  665.                 if (ty_center > t_min) idx ^= 2, pos.y += scale_exp2;
  666.                 if (tz_center > t_min) idx ^= 4, pos.z += scale_exp2;
  667.                
  668.                 t_max = tv_max;
  669.                 child_descriptor.x = 0;
  670.                 continue;
  671.             }
  672.         }
  673.        
  674.        
  675.        
  676.        
  677.         int step_mask = 0;
  678.         if (tx_corner <= tc_max) step_mask ^= 1, pos.x -= scale_exp2;
  679.         if (ty_corner <= tc_max) step_mask ^= 2, pos.y -= scale_exp2;
  680.         if (tz_corner <= tc_max) step_mask ^= 4, pos.z -= scale_exp2;
  681.        
  682.         t_min = tc_max;
  683.         idx ^= step_mask;
  684.        
  685.         if ((idx & step_mask) != 0)
  686.         {
  687.            
  688.            
  689.            
  690.            
  691.             unsigned int differing_bits = 0;
  692.             if ((step_mask & 1) != 0) differing_bits |= as_int(pos.x) ^ as_int(pos.x + scale_exp2);
  693.             if ((step_mask & 2) != 0) differing_bits |= as_int(pos.y) ^ as_int(pos.y + scale_exp2);
  694.             if ((step_mask & 4) != 0) differing_bits |= as_int(pos.z) ^ as_int(pos.z + scale_exp2);
  695.             scale = (as_int((float)differing_bits) >> 23) - 127;
  696.             scale_exp2 = as_float((scale - 23 + 127) << 23);
  697.            
  698.             parent = readStack(stack, scale, &t_max);      
  699.            
  700.            
  701.             int shx = as_int(pos.x) >> scale;
  702.             int shy = as_int(pos.y) >> scale;
  703.             int shz = as_int(pos.z) >> scale;
  704.             pos.x = as_float(shx << scale);
  705.             pos.y = as_float(shy << scale);
  706.             pos.z = as_float(shz << scale);
  707.             idx  = (shx & 1) | ((shy & 1) << 1) | ((shz & 1) << 2);
  708.            
  709.             h = 0.0f;
  710.             child_descriptor.x = 0;
  711.         }
  712.        
  713.     }
  714.    
  715.     if (scale >= 23 || iter > 10000)
  716.     {
  717.         t_min = 2.0f;
  718.     }
  719.    
  720.     if ((octant_mask & 1) == 0) pos.x = 3.0f - scale_exp2 - pos.x;
  721.     if ((octant_mask & 2) == 0) pos.y = 3.0f - scale_exp2 - pos.y;
  722.     if ((octant_mask & 4) == 0) pos.z = 3.0f - scale_exp2 - pos.z;
  723.    
  724.     (*res).pos.w = t_min;
  725.     (*res).iter = iter;
  726.     (*res).pos.x = fmin(fmax((*ray).orig.x + t_min * (*ray).dir.x, pos.x + epsilon), pos.x + scale_exp2 - epsilon);
  727.     (*res).pos.y = fmin(fmax((*ray).orig.y + t_min * (*ray).dir.y, pos.y + epsilon), pos.y + scale_exp2 - epsilon);
  728.     (*res).pos.z = fmin(fmax((*ray).orig.z + t_min * (*ray).dir.z, pos.z + epsilon), pos.z + scale_exp2 - epsilon);
  729.     (*res).node = parent;
  730.     (*res).childIdx = idx ^ octant_mask ^ 7;
  731.     (*res).stackPtr = scale;
  732. }
  733.  
  734. inline float4    decodeRawNormal         (U32 value);
  735. inline float4    decodeDXTColor          (U64 block, int texelIdx);
  736. inline float4    decodeDXTNormal         (U64 blockA, U64 blockB, int texelIdx);
  737. inline float     decodeAO                (U64 block, int texelIdx);
  738. void             lookupVoxelColorNormal  (float4* colorRes, float4* normalRes, const CastResult* castRes, const CastStack* stack);
  739. void             lookupVoxelAO           (float* res, const CastResult* castRes, const CastStack* stack);
  740. inline float4 decodeRawNormal(U32 value)
  741. {
  742.     S32 sign = (S32)value >> 31;
  743.     F32 t = (F32)(sign ^ 0x7fffffff);
  744.     F32 u = (F32)((S32)value << 3);
  745.     F32 v = (F32)((S32)value << 18);
  746.     float4 result = { t, u, v, 0.0f };
  747.     if ((value & 0x20000000) != 0)
  748.         result.x = v, result.y = t, result.z = u;
  749.     else if ((value & 0x40000000) != 0)
  750.         result.x = u, result.y = v, result.z = t;
  751.     return result;
  752. }
  753. __constant F32 c_dxtColorCoefs[4] =
  754. {
  755.     1.0f / (F32)(1 << 24),
  756.     0.0f,
  757.     2.0f / (F32)(3 << 24),
  758.     1.0f / (F32)(3 << 24),
  759. };
  760. inline float4 decodeDXTColor(U64 block, int texelIdx)
  761. {
  762.     U32 head = (U32)block;
  763.     U32 bits = (U32)(block >> 32);
  764.     F32 c0 = c_dxtColorCoefs[(bits >> (texelIdx * 2)) & 3];
  765.     F32 c1 = 1.0f / (F32)(1 << 24) - c0;
  766.     float4 res = {
  767.         c0 * (F32)(head << 16) + c1 * (F32)head,
  768.         c0 * (F32)(head << 21) + c1 * (F32)(head << 5),
  769.         c0 * (F32)(head << 27) + c1 * (F32)(head << 11),
  770.         0.0f};
  771.     return res;
  772. }
  773. __constant F32 c_dxtNormalCoefs[4] = {
  774.     -1.0f,
  775.     -1.0f / 3.0f,
  776.     +1.0f / 3.0f,
  777.     +1.0f,
  778. };
  779. inline float4 decodeDXTNormal(U64 blockA, U64 blockB, int texelIdx)
  780. {
  781.     U32 headBase = (U32)blockA;
  782.     U32 headUV   = (U32)blockB;
  783.     U32 bitsU    = (U32)(blockA >> 32);
  784.     U32 bitsV    = (U32)(blockB >> 32);
  785.     int shift = texelIdx * 2;
  786.     F32 cu = c_dxtNormalCoefs[(bitsU >> shift) & 3];
  787.     F32 cv = c_dxtNormalCoefs[(bitsV >> shift) & 3];
  788.     cu *= as_float(((headUV & 15) + (127 + 3 - 13)) << 23);
  789.     cv *= as_float((((headUV >> 16) & 15) + (127 + 3 - 13)) << 23);
  790.     float4 base = decodeRawNormal(headBase);
  791.     float4 res = {
  792.         base.x + cu * (F32)(S32)(headUV << 16) + cv * (F32)(S32)headUV,
  793.         base.y + cu * (F32)(S32)(headUV << 20) + cv * (F32)(S32)(headUV << 4),
  794.         base.z + cu * (F32)(S32)(headUV << 24) + cv * (F32)(S32)(headUV << 8),
  795.         0.0f};
  796.     return res;
  797. }
  798. __constant F32 c_dxtAOCoefs[8] = {
  799.     0.f,
  800.     1.f / ((F32)(1 << 24) * 255.f * 7.f),
  801.     2.f / ((F32)(1 << 24) * 255.f * 7.f),
  802.     3.f / ((F32)(1 << 24) * 255.f * 7.f),
  803.     4.f / ((F32)(1 << 24) * 255.f * 7.f),
  804.     5.f / ((F32)(1 << 24) * 255.f * 7.f),
  805.     6.f / ((F32)(1 << 24) * 255.f * 7.f),
  806.     1.f / ((F32)(1 << 24) * 255.f)
  807. };
  808. inline float decodeAO(U64 block, int texelIdx)
  809. {
  810.     F32 c0 = c_dxtAOCoefs[((U32)(block >> (texelIdx * 3 + 16))) & 7];
  811.     F32 c1 = c_dxtAOCoefs[7] - c0;
  812.     return c0 * (F32)((U32)block << 16) + c1 * (F32)((U32)block << 24);
  813. }
  814. void lookupVoxelColorNormal(float4* colorRes, float4* normalRes, const CastResult* castRes, const CastStack* stack)
  815. {
  816.    
  817.     __global S32* pageHeader  = (__global S32*)((intptr_t)(*castRes).node & -PageBytes);
  818.    
  819.     __global S32* blockInfo   = pageHeader + *pageHeader;
  820.     __global S32* blockStart  = blockInfo + blockInfo[BlockInfo_BlockPtr];
  821.     __global S32* attachInfos = blockInfo + BlockInfo_End;
  822.     __global S32* attachInfo  = attachInfos + AttachInfo_End * AttachSlot_Attribute;
  823.     __global S32* attachData  = blockInfo + attachInfo[AttachInfo_Ptr];
  824.     ptrdiff_t foo = (ptrdiff_t)((*castRes).node - blockStart); 
  825.     foo = foo >> 2;
  826.     foo *= 6;
  827.     __global S32* bar = attachData + foo;
  828.     __global U64* dxtBlock = (__global U64*) bar;
  829.    
  830.    
  831.     U64 colorBlock = dxtBlock[0];
  832.     U64 normalBlockA = dxtBlock[1];
  833.     U64 normalBlockB = dxtBlock[2];
  834.    
  835.    
  836.     int texelIdx = (*castRes).childIdx | (int)((((*castRes).node - pageHeader) & 2) << 2);
  837.     float4 tmp = decodeDXTColor(colorBlock, texelIdx);
  838.     *colorRes = (float4)(tmp.x, tmp.y, tmp.z, 255.0f);
  839.     *normalRes = decodeDXTNormal(normalBlockA, normalBlockB, texelIdx);
  840. }
  841. void lookupVoxelAO( float* res, const CastResult* castRes, const CastStack* stack)
  842. {
  843.    
  844.     __global S32* pageHeader  = (__global S32*)((S32)(*castRes).node & -PageBytes);
  845.     __global S32* blockInfo   = pageHeader + *pageHeader;
  846.     __global S32* blockStart  = blockInfo + blockInfo[BlockInfo_BlockPtr];
  847.     __global S32* attachInfos = blockInfo + BlockInfo_End;
  848.     __global S32* attachInfo  = attachInfos + AttachInfo_End * AttachSlot_AO;
  849.     __global S32* attachData  = blockInfo + attachInfo[AttachInfo_Ptr];
  850.    
  851.     __global U64* dxtBlock    = (__global U64*)(attachData + (((*castRes).node - blockStart) >> 2) * 2);
  852.    
  853.    
  854.     U64 block = dxtBlock[0];
  855.    
  856.     int texelIdx = (*castRes).childIdx | ((((*castRes).node - pageHeader) & 2) << 2);
  857.     *res = decodeAO(block, texelIdx);
  858. }
  859. typedef struct Aux_
  860. {
  861.     __global U32*        framePtr;
  862.     float       vSizeMultiplier;
  863.     union
  864.     {
  865.         S32     fetchWorkTemp;
  866.         Ray     ray;
  867.         struct
  868.         {
  869.             U32 color;
  870.             U32 alpha;
  871.         }       aa;
  872.     };
  873. }Aux;
  874. Ray constructPrimaryRay(__constant StandardInput* input, __global float* g_frameCoarse, int ppos, int ridx, volatile __local Aux* aux)
  875. {
  876.     float vsize = (*input).maxVoxelSize;
  877.     int xsize = (*input).frameSize.x;
  878.    
  879.     if ((*input).flags & RenderFlags_CoarsePass)
  880.     {
  881.         vsize = (float)(*input).coarseSize * 2.83f;
  882.         xsize = (*input).coarseFrameSize.x;
  883.     } else
  884.     {
  885.        
  886.         U32 a = ppos;
  887.         U32 b = ridx;
  888.         U32 c = 0x9e3779b9u;
  889.         jenkinsMix(&a, &b, &c);
  890.         float f = (float)c / ((float)(1u << 31) * 2.f);
  891.         f = .5f + .5f*f;
  892.         (*aux).vSizeMultiplier = f;
  893.         vsize *= f;
  894.     }
  895.    
  896.     int pixely = ppos / xsize;
  897.     int pixelx = ppos - (pixely * xsize);
  898.     F32 fx = pixelx;
  899.     F32 fy = pixely;
  900.     if ((*input).flags & RenderFlags_CoarsePass)
  901.     {
  902.         fx *= (float)(*input).coarseSize;
  903.         fy *= (float)(*input).coarseSize;
  904.     }
  905.     else
  906.     {
  907.         if ((*input).aaRays == 1)
  908.         {
  909.             fx += .5f;
  910.             fy += .5f;
  911.         }
  912.        
  913.        
  914.     }
  915.     F32 tmin = 0.f;
  916.     if ((*input).flags & RenderFlags_UseCoarseData)
  917.     {
  918.        
  919.         int bx = pixelx / (*input).coarseSize;
  920.         int by = pixely / (*input).coarseSize;
  921.         int bidx = bx + by * (*input).coarseFrameSize.x;
  922.         F32 tmin0 = g_frameCoarse[bidx];                                       
  923.         F32 tmin1 = g_frameCoarse[bidx+1];                                     
  924.         F32 tmin2 = g_frameCoarse[bidx+(*input).coarseFrameSize.x];            
  925.         F32 tmin3 = g_frameCoarse[bidx+(*input).coarseFrameSize.x+1];          
  926.         tmin = fmin(fmin(tmin0, tmin1), fmin(tmin2, tmin3));
  927.         tmin = fmin(tmin, 0.9999f);
  928.     }
  929.     __constant Mat4f* vtc = &(*input).octreeMatrices.viewportToCamera;
  930.     __constant Mat4f* cto = &(*input).octreeMatrices.cameraToOctree;
  931.     float4 pos = {
  932.         (*vtc).m00 * fx + (*vtc).m01 * fy + (*vtc).m03,
  933.         (*vtc).m10 * fx + (*vtc).m11 * fy + (*vtc).m13,
  934.         (*vtc).m20 * fx + (*vtc).m21 * fy + (*vtc).m23,
  935.         (*vtc).m30 * fx + (*vtc).m31 * fy + (*vtc).m33};
  936.     float4 near = {                            
  937.         pos.x - (*vtc).m02,
  938.         pos.y - (*vtc).m12,
  939.         pos.z - (*vtc).m22,
  940.         (*input).octreeMatrices.pixelInOctree * vsize};
  941.    
  942.     float4 diff = {                            
  943.         (*vtc).m32 * pos.x - (*vtc).m02 * pos.w,
  944.         (*vtc).m32 * pos.y - (*vtc).m12 * pos.w,
  945.         (*vtc).m32 * pos.z - (*vtc).m22 * pos.w,
  946.         near.w * (*vtc).m32};
  947.    
  948.     float a = 1.0f / (pos.w - (*vtc).m32);
  949.     float b = 2.0f * a / fmax(pos.w + (*vtc).m32, 1.0e-8f);
  950.     float c = tmin * b;
  951.     Ray ray;
  952.     ray.orig = near * a - diff * c;
  953.     ray.dir  = diff * (c - b);
  954.     ray.orig.xyz = (mulMat4Vec3(cto, &ray.orig)).xyz;
  955.     ray.dir = (float4)(
  956.         (*cto).m00 * ray.dir.x + (*cto).m01 * ray.dir.y + (*cto).m02 * ray.dir.z,
  957.         (*cto).m10 * ray.dir.x + (*cto).m11 * ray.dir.y + (*cto).m12 * ray.dir.z,
  958.         (*cto).m20 * ray.dir.x + (*cto).m21 * ray.dir.y + (*cto).m22 * ray.dir.z,
  959.         ray.dir.w);
  960.     return ray;
  961. }
  962. U32 processPrimaryRay(__constant StandardInput* input, __global int * g_rootNode, unsigned int rootNodeOfs,
  963.                     volatile __local Ray* ray, float vSizeMultiplier)
  964. {
  965.    
  966.     CastResult castRes;
  967.     CastStack stack;
  968.     castRayLocal(g_rootNode, rootNodeOfs, &castRes, &stack, ray);
  969.    
  970.     if ((*input).flags & RenderFlags_VisualizeIterations)  
  971.     {
  972.         F32 v = 255.0f * (F32)castRes.iter / 64.0f;
  973.         return toABGR((float4)(v, v, v, 0.0f));            
  974.     }
  975.     else if ((*input).flags & RenderFlags_VisualizeRaycastLevel)
  976.     {
  977.         F32 v = 0.0f;
  978.         if (castRes.pos.w <= 1.0f)
  979.             v = 255.0f - ((F32)23 - (F32)castRes.stackPtr) * (255.0f / 18.0f);
  980.         return toABGR((float4)(v * 0.5f, v, v * 0.5f, 0.0f));
  981.     }
  982.    
  983.     float4 L = {0.3643f, 0.3535f, 0.8616f, 0.0f};
  984.     Mat3f m3f = extractMat3f(&(*input).octreeMatrices.octreeToWorld);
  985.     float4 I = normalize(mulMatVec3local(&m3f, &(*ray).dir));
  986.    
  987.    
  988.     if (castRes.pos.w > 1.0f)
  989.     {
  990.         float4 c;
  991.         if (I.y >= 0.f)
  992.         {
  993.             float4 horz = { 179.0f, 205.0f, 253.0f, 0.0f };
  994.             float4 zen  = { 77.0f,  102.0f, 179.0f, 0.0f };
  995.             c = horz + (zen - horz) * I.y * I.y;
  996.             c *= 2.5f;
  997.         }
  998.         else
  999.         {
  1000.             float4 horz = { 192.0f, 154.0f, 102.0f, 0.0f };
  1001.             float4 zen  = { 128.0f, 102.0f, 77.0f, 0.0f };
  1002.             c = horz - (zen - horz) * I.y;
  1003.         }
  1004.         c *= fmax(L.y, 0.0f);
  1005.         float IL = dot(I, L);
  1006.         if (IL > 0.0f)
  1007.         {
  1008.             float4 f = {255.0f, 179.0f, 102.0f, 0.0f};
  1009.             c += f * pow(IL, 1000.0f);
  1010.         }
  1011.            
  1012.         return toABGR((float4)(c.x, c.y, c.z, 0.0f));
  1013.     }
  1014.    
  1015.     float4 voxelColor;
  1016.     float4 voxelNormal;
  1017.    
  1018.     lookupVoxelColorNormal(&voxelColor, &voxelNormal, &castRes, &stack);
  1019.    
  1020.     return toABGR(voxelColor);
  1021.     F32 voxelAmbient = 1.0f;
  1022.     lookupVoxelAO(&voxelAmbient, &castRes, &stack);
  1023.    
  1024.     float4 N  = normalize(mulMatVec3const(&(*input).octreeMatrices.octreeToWorldN, &voxelNormal));
  1025.     float4 R  = (I - N * (dot(N, I) * 2.0f));
  1026.     F32    LN = dot(L, N);
  1027.    
  1028.     bool shadow = (LN <= 0.0f);
  1029.    
  1030.     float4 shadedColor = voxelColor * (voxelAmbient * (0.25f + LN * ((LN < 0.0f) ? 0.15f : (shadow) ? 0.25f : 1.0f)));
  1031.     if (!shadow)
  1032.         shadedColor += (float4)(32.f, 32.f, 32.f, 0.0f) * pow(fmax(dot(L, R), 0.0f), 18.0f);
  1033.     shadedColor *= (*input).brightness;
  1034.     U32 color = toABGR(shadedColor);
  1035.    
  1036.     float vSize = (F32)(1 << castRes.stackPtr) / (F32)(1 << 23);
  1037.     float pSize = (*ray).orig.w + castRes.pos.w * (*ray).dir.w;
  1038.     vSize *= vSizeMultiplier;
  1039.     float blurRadius = max(vSize / pSize * (*input).maxVoxelSize, 1.0f);
  1040.    
  1041.     shadedColor.w = log2(blurRadius) * 32.0f + 0.5f;
  1042.     return toABGR(shadedColor);
  1043. }
  1044. void fetchWorkFirst(int* warp, int* batchCounter, __global int* g_warpCounter, int batchSize, volatile __local S32* sharedTemp)
  1045. {
  1046.     *warp = get_local_id(1) + get_group_id(0) * 2;
  1047.    
  1048.     *batchCounter = 0;
  1049. }
  1050. void fetchWorkNext(int* warp, int* batchCounter, __global int* g_warpCounter, int batchSize, volatile __local S32* sharedTemp)
  1051. {
  1052.     *warp = 0x03FFFFFF;
  1053. }
  1054. __kernel void renderKernel(__constant StandardInput* input, __global int* g_warpCounter, __global int* g_activeWarps,
  1055.                     __global int* g_coarseIndexToPixel, __global float* g_frameCoarse,
  1056.                     __global int* g_indexToPixel, __global unsigned int* g_frame,
  1057.                     __global unsigned int* g_aaSampleBuffer,
  1058.                     __global int* g_rootNode, unsigned int rootNodeOfs)
  1059. {
  1060.     int dummy = (*input).totalWork;
  1061.  
  1062.  
  1063.     __local Aux auxbuf[32 * 2];
  1064.     volatile __local Aux* aux0 = &auxbuf[32 * get_local_id(1)];
  1065.     volatile __local Aux* aux  = &auxbuf[get_local_id(0) + 32 * get_local_id(1)];
  1066.    
  1067.    
  1068.     int warp, batchCounter;
  1069.     fetchWorkFirst(&warp, &batchCounter, g_warpCounter, (*input).batchSize, (volatile __local S32*) &(*aux0).fetchWorkTemp);
  1070.     if (warp * 32 >= (*input).totalWork)
  1071.         return;
  1072.  
  1073.     for (;;)
  1074.     {
  1075.        
  1076.         int ridx = warp * 32 + get_local_id(0);
  1077.         if (ridx >= (*input).totalWork)
  1078.             return;
  1079.        
  1080.        
  1081.        
  1082.        
  1083.        
  1084.         int pidx = ridx;
  1085.         int ppos;
  1086.         if ((*input).flags & RenderFlags_CoarsePass)
  1087.         {
  1088.             ppos = g_coarseIndexToPixel[pidx];                 
  1089.             (*aux).framePtr = (__global U32*)g_frameCoarse + ppos;
  1090.         }
  1091.         else
  1092.         {
  1093.             ppos = g_indexToPixel[pidx];                       
  1094.             (*aux).framePtr = (__global U32*)g_frame + ppos;
  1095.         }
  1096.        
  1097.         Ray ray = constructPrimaryRay(input, g_frameCoarse, ppos, ridx, aux);
  1098.         (*aux).ray.orig.x   = ray.orig.x;
  1099.         (*aux).ray.orig.y   = ray.orig.y;
  1100.         (*aux).ray.orig.z   = ray.orig.z;
  1101.         (*aux).ray.dir.x    = ray.dir.x;
  1102.         (*aux).ray.dir.y    = ray.dir.y;
  1103.         (*aux).ray.dir.z    = ray.dir.z;
  1104.         (*aux).ray.orig.w   = ray.orig.w;
  1105.         (*aux).ray.dir.w    = ray.dir.w;
  1106.         if ((*input).flags & RenderFlags_CoarsePass)
  1107.         {
  1108.             CastResult castRes;
  1109.             CastStack stack;
  1110.             castRayLocal(g_rootNode, rootNodeOfs, &castRes, &stack, (volatile __local Ray*) &((*aux).ray));
  1111.             if (castRes.pos.w < 1.0f)
  1112.             {
  1113.                 F32 size = (F32)(1 << castRes.stackPtr) / (F32)(1 << 23);
  1114.                 castRes.pos.w -= size / length((float4)((*aux).ray.dir.xyz, 0.0f)) * 0.5f;
  1115.             }
  1116.             *((__global float*)(*aux).framePtr) = max(castRes.pos.w, 0.0f);
  1117.         }
  1118.         else
  1119.         {
  1120.            
  1121.             U32 color; // = processPrimaryRay(input, g_rootNode, rootNodeOfs, (volatile __local Ray*) &(*aux).ray, (*aux).vSizeMultiplier);
  1122.            
  1123.             if ((*input).aaRays == 1)
  1124.             {
  1125.                 *(*aux).framePtr = color;
  1126.             }
  1127.             else
  1128.             {
  1129.                
  1130.                 U32 resc = (color & 0xff) | ((color & 0xff00) << 2) | ((color & 0xff0000) << 4);
  1131.                 (*aux).aa.color = resc;  
  1132.                 (*aux).aa.alpha = color;
  1133.                
  1134.                 if ((get_local_id(0) & 3) == 0)
  1135.                 {
  1136.                    
  1137.                     U32 resc0 = aux[0].aa.color;
  1138.                     U32 resc1 = aux[1].aa.color;
  1139.                     U32 resc2 = aux[2].aa.color;
  1140.                     U32 resc3 = aux[3].aa.color;
  1141.                     resc = (resc0 + resc1 + resc2 + resc3);
  1142.                     resc = ((resc >> 2) & 0xff) | ((resc >> 4) & 0xff00) | ((resc >> 6) & 0xff0000);
  1143.                    
  1144.                     U32 resa0 = aux[0].aa.alpha;
  1145.                     U32 resa1 = aux[1].aa.alpha;
  1146.                     U32 resa2 = aux[2].aa.alpha;
  1147.                     U32 resa3 = aux[3].aa.alpha;
  1148.                     U32 resa = min(min(resa0, resa1), min(resa2, resa3));;
  1149.                    
  1150.                     *(*aux).framePtr = (resa & 0xff000000) | resc;
  1151.                 }
  1152.             }
  1153.         }
  1154.        
  1155.         fetchWorkNext(&warp, &batchCounter, g_warpCounter, (*input).batchSize, (volatile __local S32*) &(*aux0).fetchWorkTemp);
  1156.     }
  1157.  
  1158.  
  1159.     const int nWidth = get_global_size(0);
  1160.  
  1161.     const int xOut = get_global_id(0);
  1162.     const int yOut = get_global_id(1);
  1163.  
  1164.     const int idxOut = yOut * nWidth + xOut;
  1165.  
  1166.     g_frame[idxOut] = 1000000;
  1167. }
Advertisement
Add Comment
Please, Sign In to add comment
Advertisement
Public Pastes
Advertisement
We use cookies for various purposes including analytics. By continuing to use Pastebin, you agree to our use of cookies as described in the Cookies Policy.  OK, I Understand
Not a member of Pastebin yet?
Sign Up, it unlocks many cool features!