rt.cl

__kernel void raycastStream(__global Ray* rays, __global BVHNode* bvh, __global Triangle* triangles, __global RayHit* hits)
{
	int workID = get_global_id(0);

	Ray ray = rays[workID];

	int signs[3];
	signs[0] = ray.direction[0] < 0;
	signs[1] = ray.direction[1] < 0;
	signs[2] = ray.direction[2] < 0;
	float3 inverseDirection = (float3)(1.f/ray.direction[0],1.f/ray.direction[1],1.f/ray.direction[2]);

	int nodeStack[STACK_DEPTH];
	int todoNode = 0;
	nodeStack[0] = 0;
	float2 currentUV;
	float nearestDistance = FLT_MAX;
	RayHit hit;
	hit.t = -1;
	hit.index = -1;

	while(todoNode >= 0)
	{
		BVHNode node = bvh[nodeStack[todoNode]];
		todoNode--;
		if(!(node.nextNode[0] & (1<<(sizeof(unsigned int)*8-1))))
		{
			float aabbDistance[2];
			aabbDistance[0] = rayAABBIntersection(ray,signs,inverseDirection,bvh[node.nextNode[0]].aabb);
			aabbDistance[1] = rayAABBIntersection(ray,signs,inverseDirection,bvh[node.nextNode[1]].aabb);

			if(aabbDistance[0] >= 0)// && aabbDistance[0] < nearestDistance)
				nodeStack[++todoNode] = node.nextNode[0];

			if(aabbDistance[1] >= 0)// && aabbDistance[1] < nearestDistance)
				nodeStack[++todoNode] = node.nextNode[1];
		}
		else
		{
			int triangleIndex = node.nextNode[0] & (~(1<<(sizeof(unsigned int)*8-1)));
			const Triangle triangle = triangles[triangleIndex];
			float dist = findIntersection(ray, triangle, &currentUV);
			if(dist > 0 && dist < nearestDistance)
			{
				nearestDistance = dist;
				hit.t = nearestDistance;
				hit.index = triangleIndex;
				hit.u = currentUV.x;
				hit.v = currentUV.y;
			}
		}
	}

	hits[workID] = hit;
}