__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);
}
}