From 7ef60d0d9ef25bf93e3cf62c54debd7cc22f3d42 Mon Sep 17 00:00:00 2001
From: Anthony LaTorre <telatorre@gmail.com>
Date: Fri, 20 May 2011 11:16:00 -0400
Subject: faster bounding volume hierarchy construction

---
 src/kernel.cu | 40 ++++++++++++++++++++--------------------
 1 file changed, 20 insertions(+), 20 deletions(-)

(limited to 'src')

diff --git a/src/kernel.cu b/src/kernel.cu
index 0d5e54b..f14ed69 100644
--- a/src/kernel.cu
+++ b/src/kernel.cu
@@ -7,7 +7,7 @@
 #include "rotate.h"
 #include "intersect.h"
 
-#define STACK_SIZE 100
+#define STACK_SIZE 500
 
 /* flattened triangle mesh */
 texture<float4, 1, cudaReadModeElementType> mesh;
@@ -17,8 +17,8 @@ texture<float4, 1, cudaReadModeElementType> upper_bounds;
 texture<float4, 1, cudaReadModeElementType> lower_bounds;
 
 /* map to child nodes/triangles and the number of child nodes/triangles */
-texture<uint, 1, cudaReadModeElementType> child_map_arr;
-texture<uint, 1, cudaReadModeElementType> child_len_arr;
+texture<unsigned int, 1, cudaReadModeElementType> node_map;
+texture<unsigned int, 1, cudaReadModeElementType> node_length;
 
 __device__ float3 make_float3(const float4 &a)
 {
@@ -40,14 +40,14 @@ __device__ bool intersect_node(const float3 &origin, const float3 &direction, co
    direction `direction` and the global mesh texture. If the ray intersects
    the texture return the index of the triangle which the ray intersected,
    else return -1. */
-__device__ int intersect_mesh(const float3 &origin, const float3& direction, const int first_leaf)
+__device__ int intersect_mesh(const float3 &origin, const float3& direction, const int start_node, const int first_node)
 {
 	int triangle_idx = -1;
 
 	float distance;
 	float min_distance;
 
-	if (!intersect_node(origin, direction, 0))
+	if (!intersect_node(origin, direction, start_node))
 		return -1;
 
 	int stack[STACK_SIZE];
@@ -55,20 +55,20 @@ __device__ int intersect_mesh(const float3 &origin, const float3& direction, con
 	int *head = &stack[0];
 	int *node = &stack[1];
 	int *tail = &stack[STACK_SIZE-1];
-	*node = 0;
+	*node = start_node;
 
 	int i;
 
 	do
 	{
-		int child_map = tex1Dfetch(child_map_arr, *node);
-		int child_len = tex1Dfetch(child_len_arr, *node);
+		int index = tex1Dfetch(node_map, *node);
+		int length = tex1Dfetch(node_length, *node);
 
-		if (*node < first_leaf)
+		if (*node >= first_node)
 		{
-			for (i=0; i < child_len; i++)
-				if (intersect_node(origin, direction, child_map+i))
-					*node++ = child_map+i;
+			for (i=0; i < length; i++)
+				if (intersect_node(origin, direction, index+i))
+					*node++ = index+i;
 
 			if (node == head)
 				break;
@@ -77,9 +77,9 @@ __device__ int intersect_mesh(const float3 &origin, const float3& direction, con
 		}
 		else // node is a leaf
 		{
-			for (i=0; i < child_len; i++)
+			for (i=0; i < length; i++)
 			{
-				int mesh_idx = 3*(child_map + i);
+				int mesh_idx = 3*(index + i);
 
 				float3 v0 = make_float3(tex1Dfetch(mesh, mesh_idx));
 				float3 v1 = make_float3(tex1Dfetch(mesh, mesh_idx+1));
@@ -89,14 +89,14 @@ __device__ int intersect_mesh(const float3 &origin, const float3& direction, con
 				{
 					if (triangle_idx == -1)
 					{
-						triangle_idx = child_map + i;
+						triangle_idx = index + i;
 						min_distance = distance;
 						continue;
 					}
 
 					if (distance < min_distance)
 					{
-						triangle_idx = child_map + i;
+						triangle_idx = index + i;
 						min_distance = distance;
 					}
 				}
@@ -164,7 +164,7 @@ __global__ void rotate(int max_idx, float3 *points, float phi, float3 axis)
    set the pixel associated with the ray to a 32 bit color whose brightness is
    determined by the cosine of the angle between the ray and the normal of the
    triangle it intersected, else set the pixel to 0. */
-__global__ void ray_trace(int max_idx, float3 *origins, float3 *directions, int first_leaf, int *pixels)
+__global__ void ray_trace(int max_idx, float3 *origins, float3 *directions, int start_node, int first_node, int *pixels)
 {
 	int idx = blockIdx.x*blockDim.x + threadIdx.x;
 
@@ -175,7 +175,7 @@ __global__ void ray_trace(int max_idx, float3 *origins, float3 *directions, int
 	float3 direction = *(directions+idx);
 	direction /= norm(direction);
 
-	int intersection_idx = intersect_mesh(origin, direction, first_leaf);
+	int intersection_idx = intersect_mesh(origin, direction, start_node, first_node);
 
 	if (intersection_idx == -1)
 	{
@@ -198,7 +198,7 @@ __global__ void ray_trace(int max_idx, float3 *origins, float3 *directions, int
    intersects the mesh set the hit_solid array value associated with the
    photon to the triangle index of the triangle the photon intersected, else
    set the hit_solid array value to -1. */
-__global__ void propagate(int max_idx, float3 *origins, float3 *directions, int first_leaf, int *hit_triangles)
+__global__ void propagate(int max_idx, float3 *origins, float3 *directions, int start_node, int first_node, int *hit_triangles)
 {
 	int idx = blockIdx.x*blockDim.x + threadIdx.x;
 
@@ -209,7 +209,7 @@ __global__ void propagate(int max_idx, float3 *origins, float3 *directions, int
 	float3 direction = *(directions+idx);
 	direction /= norm(direction);
 
-	*(hit_triangles+idx) = intersect_mesh(origin, direction, first_leaf);
+	*(hit_triangles+idx) = intersect_mesh(origin, direction, start_node, first_node);
 
 } // propagate
 
-- 
cgit