From 0716c8af5be04ea1709ca178d5341e5ee3df607b Mon Sep 17 00:00:00 2001
From: Anthony LaTorre <tlatorre9@gmail.com>
Date: Mon, 25 Jul 2011 16:41:34 -0400
Subject: 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.

---
 src/alpha.h | 157 ++++++++++++++++++++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 157 insertions(+)
 create mode 100644 src/alpha.h

(limited to 'src/alpha.h')

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
-- 
cgit