summaryrefslogtreecommitdiff
path: root/src/alpha.h
diff options
context:
space:
mode:
Diffstat (limited to 'src/alpha.h')
-rw-r--r--src/alpha.h157
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