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

}