moved triangle colors to a separate global device array so that the ray tracer and photon simulation can be run in the same context. added the ability to add alpha channel to triangle color so that triangles can be made transparent. added __noinline__ modifier to certain device functions to speed up kernel compilation.

4 files changed, 175 insertions, 27 deletions

diff --git a/src/alpha.h b/src/alpha.h
new file mode 100644
index 0000000..afd7ea5
--- /dev/null
+++ b/src/alpha.h
@@ -0,0 +1,157 @@
+#ifndef __ALPHA_H__
+#define __ALPHA_H__
+
+#include "linalg.h"
+#include "intersect.h"
+#include "mesh.h"
+
+template <class T>
+__device__ void swap(T &a, T &b)
+{
+	T temp = a;
+	a = b;
+	b = temp;
+}
+
+#define ALPHA_DEPTH 5
+
+struct HitList
+{
+	int size;
+	int indices[ALPHA_DEPTH];
+	float distances[ALPHA_DEPTH];
+};
+
+__device__ void add_to_hit_list(const float &distance, const int &index, HitList &h)
+{
+	int i;
+	if (h.size >= ALPHA_DEPTH)
+	{
+		if (distance > h.distances[ALPHA_DEPTH-1])
+			return;
+
+		i = ALPHA_DEPTH-1;
+	}
+	else
+	{
+		i = h.size;
+		h.size += 1;
+	}
+
+	h.indices[i] = index;
+	h.distances[i] = distance;
+
+	while (i > 0 && h.distances[i-1] > h.distances[i])
+	{
+		swap(h.distances[i-1], h.distances[i]);
+		swap(h.indices[i-1], h.indices[i]);
+
+		i -= 1;
+	}
+}
+
+__device__ __noinline__ int get_color_alpha(const float3 &origin, const float3& direction)
+{
+	HitList h;
+	h.size = 0;
+
+	float distance;
+
+	if (!intersect_node(origin, direction, g_start_node))
+		return 0;
+
+	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++)
+			{
+				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))
+				{
+					add_to_hit_list(distance, first_child+i, h);
+				}
+
+			} // triangle loop
+
+			node--;
+
+		} // node is a leaf
+
+	} // while loop
+	while (node != head);
+
+	if (h.size < 1)
+		return 0;
+
+	float scale = 1.0;
+	unsigned int r = 0;
+	unsigned int g = 0;
+	unsigned int b = 0;
+	for (i=0; i < h.size; i++)
+	{
+		uint4 triangle_data = g_triangles[h.indices[i]];
+
+		float3 v0 = g_vertices[triangle_data.x];
+		float3 v1 = g_vertices[triangle_data.y];
+		float3 v2 = g_vertices[triangle_data.z];
+
+		float cos_theta = dot(normalize(cross(v1-v0,v2-v1)),-direction);
+
+		if (cos_theta < 0.0f)
+			cos_theta = dot(-normalize(cross(v1-v0,v2-v1)),-direction);
+
+		unsigned int r0 = 0xff & (g_colors[h.indices[i]] >> 16);
+		unsigned int g0 = 0xff & (g_colors[h.indices[i]] >> 8);
+		unsigned int b0 = 0xff & g_colors[h.indices[i]];
+
+		float alpha = (255 - (0xff & (g_colors[h.indices[i]] >> 24)))/255.0f;
+
+		r += floorf(r0*scale*cos_theta*alpha);
+		g += floorf(g0*scale*cos_theta*alpha);
+		b += floorf(b0*scale*cos_theta*alpha);
+
+		scale *= (1.0f-alpha);
+	}
+
+	return r << 16 | g << 8 | b;
+}
+
+#endif
diff --git a/src/intersect.h b/src/intersect.h
index b1713d7..c400f1c 100644
--- a/src/intersect.h
+++ b/src/intersect.h
@@ -9,7 +9,7 @@
 #include "matrix.h"
 #include "rotate.h"
 
-#define EPSILON 1.0e-3f
+#define EPSILON 0.0f
 
 /* Test the intersection between a ray starting from `origin` traveling in the
    direction `direction` and a triangle defined by the vertices `v0`, `v1`, and
@@ -69,7 +69,7 @@ __device__ int get_color(const float3 &direction, const float3 &v0, const float3
 	unsigned int b = 0xFF & base_color;
 
 	if (scale < 0.0f)
-		scale = dot(-normalize(cross(v1-v0,v2-v0)),-direction);
+		scale = dot(-normalize(cross(v1-v0,v2-v1)),-direction);
 
 	r = floorf(r*scale);
 	g = floorf(g*scale);
diff --git a/src/kernel.cu b/src/kernel.cu
index 711ae80..0dc5639 100644
--- a/src/kernel.cu
+++ b/src/kernel.cu
@@ -7,6 +7,7 @@
 #include "rotate.h"
 #include "mesh.h"
 #include "photon.h"
+#include "alpha.h"
 
 #define RED_WAVELENGTH 685
 #define BLUE_WAVELENGTH 465
@@ -18,7 +19,7 @@ __device__ void fAtomicAdd(float *addr, float data)
 		data = atomicExch(addr, data+atomicExch(addr, 0.0f));
 }
 
-__device__ void to_diffuse(Photon &p, State &s, curandState &rng, const int &max_steps)
+__device__ __noinline__ void to_diffuse(Photon &p, State &s, curandState &rng, const int &max_steps)
 {
 	int steps = 0;
 	while (steps < max_steps)
@@ -289,24 +290,7 @@ __global__ void ray_trace(int nthreads, float3 *positions, float3 *directions, i
 	float3 direction = directions[id];
 	direction /= norm(direction);
 
-	float distance;
-
-	int triangle_index = intersect_mesh(position, direction, distance);
-
-	if (triangle_index == -1)
-	{
-		pixels[id] = 0x000000;
-	}
-	else
-	{
-		uint4 triangle_data = g_triangles[triangle_index];
-
-		float3 v0 = g_vertices[triangle_data.x];
-		float3 v1 = g_vertices[triangle_data.y];
-		float3 v2 = g_vertices[triangle_data.z];
-
-		pixels[id] = get_color(direction, v0, v1, v2, triangle_data.w);
-	}
+	pixels[id] = get_color_alpha(position, direction);
 
 } // ray_trace
 
diff --git a/src/mesh.h b/src/mesh.h
index 6f678c0..ec5a508 100644
--- a/src/mesh.h
+++ b/src/mesh.h
@@ -8,8 +8,9 @@
 /* flattened triangle mesh */
 __device__ float3 *g_vertices;
 __device__ uint4 *g_triangles;
-__device__ int g_start_node;
-__device__ int g_first_node;
+__device__ unsigned int *g_colors;
+__device__ unsigned int g_start_node;
+__device__ unsigned int g_first_node;
 
 /* lower/upper bounds for the bounding box associated with each node/leaf */
 texture<float4, 1, cudaReadModeElementType> upper_bounds;
@@ -47,7 +48,7 @@ __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, float &min_distance, int last_hit_triangle = -1)
+__device__ __noinline__ int intersect_mesh(const float3 &origin, const float3& direction, float &min_distance, int last_hit_triangle = -1)
 {
 	int triangle_index = -1;
 
@@ -133,14 +134,20 @@ __device__ int intersect_mesh(const float3 &origin, const float3& direction, flo
 extern "C"
 {
 
-__global__ void set_global_mesh_variables(uint4 *triangle_ptr, float3 *vertex_ptr, int start_node, int first_node)
+__global__ void set_global_mesh_variables(uint4 *triangles, float3 *vertices, unsigned int *colors, unsigned int start_node, unsigned int first_node)
 {
-	g_triangles = triangle_ptr;
-	g_vertices = vertex_ptr;
+	g_triangles = triangles;
+	g_vertices = vertices;
+	g_colors = colors;
 	g_start_node = start_node;
 	g_first_node = first_node;
 }
 
+__global__ void set_colors(unsigned int *colors)
+{
+	g_colors = colors;
+}
+
 } // extern "c"
 
 #endif