Octree Traversal

kernel void Draw(__global unsigned int *pPixels, __global unsigned int *pChildData, __global unsigned int *pColorData, __global float *pRayDirs, __global float *pCamPos, __global float *pRotData)
{
  const float largestFloat = 9999999999.f;
  const float smallestFloat = 0.000000001f;

  const int octreeLayers = 12;
  int id = get_global_id(0);

  float3 rayPos = (float3)(pCamPos[0], pCamPos[1], pCamPos[2]);
  float3 rayDir = (float3)(pRayDirs[id * 3 + 0], pRayDirs[id * 3 + 1], pRayDirs[id * 3 + 2]);
  rayDir = (float3)(rayDir.x, rayDir.y * pRotData[2] - rayDir.z * pRotData[3], rayDir.y * pRotData[3] + rayDir.z * pRotData[2]);
  rayDir = (float3)(rayDir.x * pRotData[0] - rayDir.y * pRotData[1], rayDir.x * pRotData[1] + rayDir.y * pRotData[0], rayDir.z);

  unsigned int stackPtr = 0;
  unsigned int nodeStack[64];
  unsigned int nodeLayers[64];
  float3 nodeBoxes[64];

  nodeStack[stackPtr] = 0;
  nodeLayers[stackPtr] = 1;
  nodeBoxes[stackPtr] = (float3)(0, 0, 0);
  stackPtr = 1;

  float dirfracx = 1.0f / rayDir.x;
  float dirfracy = 1.0f / rayDir.y;
  float dirfracz = 1.0f / rayDir.z;

  while (stackPtr > 0)
  {
    unsigned int nodeID = nodeStack[stackPtr - 1];
    unsigned int childID = pChildData[nodeID];
    unsigned int nodeLayer = nodeLayers[stackPtr - 1];
    float3 nodeBoxPos = nodeBoxes[stackPtr - 1];
    float nodeBoxSize = (float)(1 << (octreeLayers - nodeLayer));
    --stackPtr;

    bool hit = false;
    // Ray In Box
    {
      if (rayPos.x >= nodeBoxPos.x && rayPos.y >= nodeBoxPos.y && rayPos.z >= nodeBoxPos.z)
        if (rayPos.x < (nodeBoxPos.x + nodeBoxSize) && rayPos.y < (nodeBoxPos.y + nodeBoxSize) && rayPos.z < (nodeBoxPos.z + nodeBoxSize))
          hit = true;
    }
    if (!hit)
    {
      float t1 = ((nodeBoxPos.x) - rayPos.x) * dirfracx;
      float t2 = ((nodeBoxPos.x + nodeBoxSize) - rayPos.x) * dirfracx;
      float t3 = ((nodeBoxPos.y) - rayPos.y) * dirfracy;
      float t4 = ((nodeBoxPos.y + nodeBoxSize) - rayPos.y) * dirfracy;
      float t5 = ((nodeBoxPos.z) - rayPos.z) * dirfracz;
      float t6 = ((nodeBoxPos.z + nodeBoxSize) - rayPos.z) * dirfracz;
      float tmax = min(min(max(t1, t2), max(t3, t4)), max(t5, t6));
      if (tmax > 0)
      {
        float tmin = max(max(min(t1, t2), min(t3, t4)), min(t5, t6));
        if (tmin < tmax)
          hit = true;
      }
    }

    if (hit)
    {
      if (nodeLayer == octreeLayers)
      {
        pPixels[id] = pColorData[nodeID];
        return;
      }
      else
        if (childID > 0)
        {
          // Loop based
          for (unsigned int cz = 0; cz < 2; cz++) for (unsigned int cy = 0; cy < 2; cy++) for (unsigned int cx = 0; cx < 2; cx++)
          // Iterate children in order of intersection
          {
            unsigned int c = (cx ^ (rayDir.x > 0)) + (cy ^ (rayDir.y > 0)) * 2 + (cz ^ (rayDir.z > 0)) * 4;
            if (pColorData[childID + c] == 0) continue; // Skip empty children
            nodeStack[stackPtr] = childID + c;
            nodeLayers[stackPtr] = nodeLayer + 1;
            unsigned int subBoxSize = 1 << (octreeLayers - nodeLayer - 1);
            float3 boxStart = nodeBoxPos + (float3)((c & 1) > 0, (c & 2) > 0, (c & 4) > 0) * (float)subBoxSize;
            nodeBoxes[stackPtr] = boxStart;
            ++stackPtr;
          }
        }
    }
  }
  pPixels[id] = 0;
}