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, 13 insertions, 6 deletions

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