1 files changed, 146 insertions, 0 deletions
diff --git a/src/mesh.h b/src/mesh.h
new file mode 100644
index 0000000..6f678c0
--- /dev/null
+++ b/src/mesh.h
@@ -0,0 +1,146 @@
+#ifndef __MESH_H__
+#define __MESH_H__
+
+#include "intersect.h"
+
+#define STACK_SIZE 500
+
+/* flattened triangle mesh */
+__device__ float3 *g_vertices;
+__device__ uint4 *g_triangles;
+__device__ int g_start_node;
+__device__ int g_first_node;
+
+/* lower/upper bounds for the bounding box associated with each node/leaf */
+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<unsigned int, 1, cudaReadModeElementType> node_map;
+texture<unsigned int, 1, cudaReadModeElementType> node_length;
+
+__device__ float3 make_float3(const float4 &a)
+{
+	return make_float3(a.x, a.y, a.z);
+}
+
+__device__ int convert(int c)
+{
+	if (c & 0x80)
+		return (0xFFFFFF00 | c);
+	else
+		return c;
+}
+
+/* Test the intersection between a ray starting at `origin` traveling in the
+   direction `direction` and the bounding box around node `i`. If the ray
+   intersects the bounding box return true, else return false. */
+__device__ bool intersect_node(const float3 &origin, const float3 &direction, const int &i)
+{
+	float3 lower_bound = make_float3(tex1Dfetch(lower_bounds, i));
+	float3 upper_bound = make_float3(tex1Dfetch(upper_bounds, i));
+
+	return intersect_box(origin, direction, lower_bound, upper_bound);
+}
+
+/* Find the intersection between a ray starting at `origin` traveling in the
+   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, float &min_distance, int last_hit_triangle = -1)
+{
+	int triangle_index = -1;
+
+	float distance;
+
+	if (!intersect_node(origin, direction, g_start_node))
+		return -1;
+
+	int stack[STACK_SIZE];
+
+	int *head = &stack[0];
+	int *node = &stack[1];
+	int *tail = &stack[STACK_SIZE-1];
+	*node = g_start_node;
+
+	int i;
+
+	do
+	{
+		int first_child = tex1Dfetch(node_map, *node);
+		int child_count = tex1Dfetch(node_length, *node);
+
+		while (*node >= g_first_node && child_count == 1)
+		{
+			*node = first_child;
+			first_child = tex1Dfetch(node_map, *node);
+			child_count = tex1Dfetch(node_length, *node);
+		}
+		
+		if (*node >= g_first_node)
+		{
+			for (i=0; i < child_count; i++)
+			{
+				if (intersect_node(origin, direction, first_child+i))
+				{
+					*node = first_child+i;
+					node++;
+				}
+			}
+
+			node--;
+		}
+		else // node is a leaf
+		{
+			for (i=0; i < child_count; i++)
+			{
+				if (last_hit_triangle == first_child+i)
+					continue;
+
+				uint4 triangle_data = g_triangles[first_child+i];
+
+				float3 v0 = g_vertices[triangle_data.x];
+				float3 v1 = g_vertices[triangle_data.y];
+				float3 v2 = g_vertices[triangle_data.z];
+
+				if (intersect_triangle(origin, direction, v0, v1, v2, distance))
+				{
+					if (triangle_index == -1)
+					{
+						triangle_index = first_child + i;
+						min_distance = distance;
+						continue;
+					}
+
+					if (distance < min_distance)
+					{
+						triangle_index = first_child + i;
+						min_distance = distance;
+					}
+				}
+			} // triangle loop
+
+			node--;
+
+		} // node is a leaf
+
+	} // while loop
+	while (node != head);
+
+	return triangle_index;
+}
+
+extern "C"
+{
+
+__global__ void set_global_mesh_variables(uint4 *triangle_ptr, float3 *vertex_ptr, int start_node, int first_node)
+{
+	g_triangles = triangle_ptr;
+	g_vertices = vertex_ptr;
+	g_start_node = start_node;
+	g_first_node = first_node;
+}
+
+} // extern "c"
+
+#endif