An Efficient Parametric Algorithm for Octree Traversal

__device__ void ProcSubtree(Ray & ray,
                            float tx0, float ty0, float tz0,
                            float tx1, float ty1, float tz1,
                            int a, float4 & dstColor)
{

    ActualState arrayStack[25];
    int stackTop = 0;

    int currNode = 0;
    float txm, tym, tzm;

    ActualState state;
    ActualState newState;

    state.tx0 = tx0; state.ty0 = ty0; state.tz0 = tz0;
    state.tx1 = tx1; state.ty1 = ty1; state.tz1 = tz1;
    state.index = 0;
    state.depth = 0;
    arrayStack[stackTop] = state;

    float invRGBA = 1.0f / 255.0f;
    uint8 stateLookUp = 0;


    while(stackTop >= 0)
    {


        state = arrayStack[stackTop];
        stackTop--;


        if (state.tx1 < 0.0f || state.ty1 < 0.0f || state.tz1 < 0.0f)
        {
            continue;
        }

        if (state.index < 0)
        {
            if (state.index != EMPTY_NODE)
            {


                //dstColor = make_float4(0);
                dstColor.x = 0.0f;
                dstColor.y = 0.5f;
                dstColor.z = 0.0f;
                break;
            }

            continue;
        }



        txm = 0.5f * (state.tx0 + state.tx1);
        tym = 0.5f * (state.ty0 + state.ty1);
        tzm = 0.5f * (state.tz0 + state.tz1);

        currNode = FirstNode(state.tx0, state.ty0, state.tz0, txm, tym, tzm);

        state.index++;
        do
        {
            switch (currNode)
            {
                case FLD_NODE:
                    newState.tx0 = state.tx0; newState.ty0 = state.ty0; newState.tz0 = state.tz0;
                    newState.tx1 = txm;       newState.ty1 = tym;       newState.tz1 = tzm;

                    break;
                case BLD_NODE:
                    newState.tx0 = state.tx0; newState.ty0 = state.ty0; newState.tz0 = tzm;
                    newState.tx1 = txm;       newState.ty1 = tym;       newState.tz1 = state.tz1;

                    break;
                case FLT_NODE:
                    newState.tx0 = state.tx0; newState.ty0 = tym;       newState.tz0 = state.tz0;
                    newState.tx1 = txm;       newState.ty1 = state.ty1; newState.tz1 = tzm;

                    break;
                case BLT_NODE:
                    newState.tx0 = state.tx0; newState.ty0 = tym;       newState.tz0 = tzm;
                    newState.tx1 = txm;       newState.ty1 = state.ty1; newState.tz1 = state.tz1;

                    break;
                case FRD_NODE:
                    newState.tx0 = txm;       newState.ty0 = state.ty0; newState.tz0 = state.tz0;
                    newState.tx1 = state.tx1; newState.ty1 = tym;       newState.tz1 = tzm;

                    break;
                case BRD_NODE:
                    newState.tx0 = txm;       newState.ty0 = state.ty0; newState.tz0 = tzm;
                    newState.tx1 = state.tx1; newState.ty1 = tym;       newState.tz1 = state.tz1;

                    break;
                case FRT_NODE:
                    newState.tx0 = txm;       newState.ty0 = tym;       newState.tz0 = state.tz0;
                    newState.tx1 = state.tx1; newState.ty1 = state.ty1; newState.tz1 = tzm;

                    break;
                case BRT_NODE:
                    newState.tx0 = txm;       newState.ty0 = tym;       newState.tz0 = tzm;
                    newState.tx1 = state.tx1; newState.ty1 = state.ty1; newState.tz1 = state.tz1;

                    break;
            }



            switch(currNode)
            {
                case FLD_NODE:
                    stateLookUp = FLD_NODE ^ a;
                    break;
                case BLD_NODE:
                    stateLookUp = BLD_NODE ^ a;
                    break;
                case FLT_NODE:
                    stateLookUp = FLT_NODE ^ a;
                    break;
                case BLT_NODE:
                    stateLookUp = BLT_NODE ^ a;
                    break;
                case FRD_NODE:
                    stateLookUp = FRD_NODE ^ a;
                    break;
                case BRD_NODE:
                    stateLookUp = BRD_NODE ^ a;
                    break;
                case FRT_NODE:
                    stateLookUp = FRT_NODE ^ a;
                    break;
                case BRT_NODE:
                    stateLookUp = BRT_NODE ^ a;
                    break;
                default:
                    break;
            }



            newState.index  = tex1Dfetch(dataTex, state.index + stateLookUp);

            if (newState.index == 0)
            {
                continue;
            }



            uint8 case1 = 8;
            uint8 case2 = 8;
            uint8 case3 = 8;

            switch(currNode)
            {
                case FLD_NODE:
                    case1 = FRD_NODE; case2 = FLT_NODE; case3 = BLD_NODE;
                    break;
                case BLD_NODE:
                    case1 = BRD_NODE; case2 = BLT_NODE;
                    break;
                case FLT_NODE:
                    case1 = FRT_NODE; case3 = BLT_NODE;
                    break;
                case BLT_NODE:
                    case1 = BRT_NODE;
                    break;
                case FRD_NODE:
                    case2 = FRT_NODE; case3 = BRD_NODE;
                    break;
                case BRD_NODE:
                    case2 = BRT_NODE;
                    break;
                case FRT_NODE:
                    case3 = BRT_NODE;
                    break;
                default:
                    break;
            }

            currNode = NewNode(newState.tx1, case1,
                               newState.ty1, case2,
                               newState.tz1, case3);


            newState.depth = state.depth + 1;

            stackTop++;
            arrayStack[stackTop] = newState;


        } while (currNode < 8);
    }

}

__device__ void TraverseRay(Ray & ray, const float3 & dataSize, float4 & dstColor)
{
    int a = 0;
    if (ray.dir.x < 0.0f)
    {
        ray.origin.x = dataSize.x - ray.origin.x;
        ray.dir.x *= -1;
        a |= 4;
    }
    if (ray.dir.y < 0.0f)
    {
        ray.origin.y = dataSize.y - ray.origin.y;
        ray.dir.y *= -1;
        a |= 2;
    }
    if (ray.dir.z < 0.0f)
    {
        ray.origin.z = dataSize.z - ray.origin.z;
        ray.dir.z *= -1;
        a |= 1;
    }


    float invDir = 1.0f / ray.dir.x;

    float tx0 = -ray.origin.x * invDir;
    float tx1 = (dataSize.x - ray.origin.x) * invDir;

    invDir = 1.0f / ray.dir.y;
    float ty0 = -ray.origin.y * invDir;
    float ty1 = (dataSize.y - ray.origin.y) * invDir;

    invDir = 1.0f / ray.dir.z;
    float tz0 = -ray.origin.z * invDir;
    float tz1 = (dataSize.z - ray.origin.z) * invDir;

    /*
    float invDir = __fdividef(1.0f, ray.dir.x);

    float tx0 = (0 - ray.origin.x) * invDir;
    float tx1 = ( 64 - ray.origin.x) * invDir;

    invDir = __fdividef(1.0f, ray.dir.y);
    float ty0 = (0 - ray.origin.y) * invDir;
    float ty1 = ( 64 - ray.origin.y) * invDir;

    invDir = __fdividef(1.0f, ray.dir.z);
    float tz0 = (0 - ray.origin.z) * invDir;
    float tz1 = ( 64 - ray.origin.z) * invDir;
    */

    if (ray.dir.x == 0.0f)
    {
        if (tx0 < 0) tx0 = -INT_MAX;
        else tx0 = INT_MAX;
        if (tx1 < 0) tx1 = -INT_MAX;
        else tx1 = INT_MAX;
    }
    if (ray.dir.y == 0.0f)
    {
        if (ty0 < 0) ty0 = -INT_MAX;
        else ty0 = INT_MAX;
        if (ty1 < 0) ty1 = -INT_MAX;
        else ty1 = INT_MAX;
    }
    if (ray.dir.z == 0.0f)
    {
        if (tz0 < 0) tz0 = -INT_MAX;
        else tz0 = INT_MAX;
        if (tz1 < 0) tz1 = -INT_MAX;
        else tz1 = INT_MAX;
    }

    if (Max(tx0, ty0, tz0) < Min(tx1, ty1, tz1))
    {
        //dstColor.x = 0.5f;
        ProcSubtree(ray, tx0, ty0, tz0, tx1, ty1, tz1, a, dstColor);
    }


}