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.

1 files changed, 157 insertions, 0 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