diff options
Diffstat (limited to 'src/mesh.h')
| -rw-r--r-- | src/mesh.h | 39 |
1 files changed, 20 insertions, 19 deletions
@@ -18,7 +18,8 @@ texture<float4, 1, cudaReadModeElementType> lower_bounds; /* map to child nodes/triangles and the number of child nodes/triangles */ texture<unsigned int, 1, cudaReadModeElementType> node_map; -texture<unsigned int, 1, cudaReadModeElementType> node_length; +texture<unsigned int, 1, cudaReadModeElementType> node_map_end; +//texture<unsigned int, 1, cudaReadModeElementType> node_length; __device__ float3 make_float3(const float4 &a) { @@ -48,7 +49,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__ __noinline__ int intersect_mesh(const float3 &origin, const float3& direction, float &min_distance, int last_hit_triangle = -1) +__device__ int intersect_mesh(const float3 &origin, const float3& direction, float &min_distance, int last_hit_triangle = -1) { int triangle_index = -1; @@ -57,34 +58,34 @@ __device__ __noinline__ int intersect_mesh(const float3 &origin, const float3& d if (!intersect_node(origin, direction, g_start_node)) return -1; - int stack[STACK_SIZE]; + unsigned int stack[STACK_SIZE]; - int *head = &stack[0]; - int *node = &stack[1]; - int *tail = &stack[STACK_SIZE-1]; + unsigned int *head = &stack[0]; + unsigned int *node = &stack[1]; + unsigned int *tail = &stack[STACK_SIZE-1]; *node = g_start_node; - int i; + unsigned int i; do { - int first_child = tex1Dfetch(node_map, *node); - int child_count = tex1Dfetch(node_length, *node); + unsigned int first_child = tex1Dfetch(node_map, *node); + unsigned int stop = tex1Dfetch(node_map_end, *node); - while (*node >= g_first_node && child_count == 1) + while (*node >= g_first_node && stop == first_child+1) { *node = first_child; first_child = tex1Dfetch(node_map, *node); - child_count = tex1Dfetch(node_length, *node); + stop = tex1Dfetch(node_map_end, *node); } if (*node >= g_first_node) { - for (i=0; i < child_count; i++) + for (i=first_child; i < stop; i++) { - if (intersect_node(origin, direction, first_child+i)) + if (intersect_node(origin, direction, i)) { - *node = first_child+i; + *node = i; node++; } } @@ -93,12 +94,12 @@ __device__ __noinline__ int intersect_mesh(const float3 &origin, const float3& d } else // node is a leaf { - for (i=0; i < child_count; i++) + for (i=first_child; i < stop; i++) { - if (last_hit_triangle == first_child+i) + if (last_hit_triangle == i) continue; - uint4 triangle_data = g_triangles[first_child+i]; + uint4 triangle_data = g_triangles[i]; float3 v0 = g_vertices[triangle_data.x]; float3 v1 = g_vertices[triangle_data.y]; @@ -108,14 +109,14 @@ __device__ __noinline__ int intersect_mesh(const float3 &origin, const float3& d { if (triangle_index == -1) { - triangle_index = first_child + i; + triangle_index = i; min_distance = distance; continue; } if (distance < min_distance) { - triangle_index = first_child + i; + triangle_index = i; min_distance = distance; } } |