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