diff options
Diffstat (limited to 'src/mesh.h')
| -rw-r--r-- | src/mesh.h | 273 |
1 files changed, 125 insertions, 148 deletions
@@ -2,187 +2,164 @@ #define __MESH_H__ #include "intersect.h" +#include "geometry.h" #define STACK_SIZE 500 -/* flattened triangle mesh */ -__device__ float3 *g_vertices; -__device__ uint4 *g_triangles; -__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 */ -__device__ float3 *g_lower_bounds; -__device__ float3 *g_upper_bounds; - -/* map to child node/triangle indices */ -texture<unsigned int, 1, cudaReadModeElementType> node_map; -texture<unsigned int, 1, cudaReadModeElementType> node_map_end; - -__device__ float3 make_float3(const float4 &a) +/* Tests the intersection between a ray and a node in the bounding volume + hierarchy. If the ray intersects the bounding volume and `min_distance` + is less than zero or the distance from `origin` to the intersection is + less than `min_distance`, return true, else return false. */ +__device__ bool +intersect_node(const float3 &origin, const float3 &direction, + Geometry *geometry, const int &i, const float &min_distance) { - return make_float3(a.x, a.y, a.z); -} + /* assigning these to local variables is faster for some reason */ + float3 lower_bound = geometry->lower_bounds[i]; + float3 upper_bound = geometry->upper_bounds[i]; -__device__ int convert(int c) -{ - if (c & 0x80) - return (0xFFFFFF00 | c); - else - return c; -} + float distance_to_box; -/* Test the intersection between a ray starting at `origin` traveling in the - direction `direction` and the bounding box around node `i`. If the ray - intersects the bounding box return true, else return false. */ -__device__ bool intersect_node(const float3 &origin, const float3 &direction, const int &i, const float &min_distance) -{ - float3 lower_bound = g_lower_bounds[i]; - float3 upper_bound = g_upper_bounds[i]; + if (intersect_box(origin, direction, lower_bound, upper_bound, + distance_to_box)) { + if (min_distance < 0.0f) + return true; - float distance_to_box; + if (distance_to_box > min_distance) + return false; - if (intersect_box(origin, direction, lower_bound, upper_bound, distance_to_box)) - { - if (min_distance < 0.0f) - return true; + return true; + } + else { + return false; + } - if (distance_to_box > min_distance) - return false; - - return true; - } - else - { - return false; - } - -} // intersect_node +} -/* Find the intersection between a ray starting at `origin` traveling in the - 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) +/* Finds the intersection between a ray and `geometry`. If the ray does + intersect the mesh and the index of the intersected triangle is not equal + to `last_hit_triangle`, set `min_distance` to the distance from `origin` to + the intersection and return the index of the triangle which the ray + intersected, else return -1. */ +__device__ int +intersect_mesh(const float3 &origin, const float3& direction, + Geometry *geometry, float &min_distance, + int last_hit_triangle = -1) { - int triangle_index = -1; + int triangle_index = -1; - float distance; - min_distance = -1.0f; + float distance; + min_distance = -1.0f; - if (!intersect_node(origin, direction, g_start_node, min_distance)) - return -1; + if (!intersect_node(origin, direction, geometry, geometry->start_node, + min_distance)) + return -1; - unsigned int stack[STACK_SIZE]; + unsigned int stack[STACK_SIZE]; - unsigned int *head = &stack[0]; - unsigned int *node = &stack[1]; - unsigned int *tail = &stack[STACK_SIZE-1]; - *node = g_start_node; + unsigned int *head = &stack[0]; + unsigned int *node = &stack[1]; + unsigned int *tail = &stack[STACK_SIZE-1]; + *node = geometry->start_node; - unsigned int i; + unsigned int i; - do - { - unsigned int first_child = tex1Dfetch(node_map, *node); - unsigned int stop = tex1Dfetch(node_map_end, *node); + do + { + unsigned int first_child = geometry->node_map[*node]; + unsigned int stop = geometry->node_map_end[*node]; - while (*node >= g_first_node && stop == first_child+1) - { - *node = first_child; - first_child = tex1Dfetch(node_map, *node); - stop = tex1Dfetch(node_map_end, *node); - } + while (*node >= geometry->first_node && stop == first_child+1) { + *node = first_child; + first_child = geometry->node_map[*node]; + stop = geometry->node_map_end[*node]; + } - if (*node >= g_first_node) - { - for (i=first_child; i < stop; i++) - { - if (intersect_node(origin, direction, i, min_distance)) - { - *node = i; - node++; - } - } - - node--; + if (*node >= geometry->first_node) { + for (i=first_child; i < stop; i++) { + if (intersect_node(origin, direction, geometry, i, + min_distance)) { + *node = i; + node++; + } + } + + node--; + } + else { + // node is a leaf + for (i=first_child; i < stop; i++) { + if (last_hit_triangle == i) + continue; + + Triangle triangle = get_triangle(geometry, i); + + if (intersect_triangle(origin, direction, triangle, + distance)) { + if (triangle_index == -1) { + triangle_index = i; + min_distance = distance; + continue; + } + + if (distance < min_distance) { + triangle_index = i; + min_distance = distance; + } } - else // node is a leaf - { - for (i=first_child; i < stop; i++) - { - if (last_hit_triangle == i) - continue; - - uint4 triangle_data = g_triangles[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)) - { - if (triangle_index == -1) - { - triangle_index = i; - min_distance = distance; - continue; - } - - if (distance < min_distance) - { - triangle_index = i; - min_distance = distance; - } - } - } // triangle loop - - node--; - - } // node is a leaf - - } // while loop - while (node != head); - - return triangle_index; + } // triangle loop + + node--; + + } // node is a leaf + + } // while loop + while (node != head); + + return triangle_index; } extern "C" { -__global__ void set_global_mesh_variables(uint4 *triangles, float3 *vertices, unsigned int *colors, unsigned int start_node, unsigned int first_node, float3 *lower_bounds, float3 *upper_bounds) +__global__ void +distance_to_mesh(int nthreads, float3 *positions, float3 *directions, + Geometry *g, float *distances) { - g_triangles = triangles; - g_vertices = vertices; - g_colors = colors; - g_start_node = start_node; - g_first_node = first_node; - g_lower_bounds = lower_bounds; - g_upper_bounds = upper_bounds; -} + int id = blockIdx.x*blockDim.x + threadIdx.x; -__global__ void set_colors(unsigned int *colors) -{ - g_colors = colors; + if (id >= nthreads) + return; + + float3 position = positions[id]; + float3 direction = directions[id]; + direction /= norm(direction); + + float distance; + + int triangle_index = intersect_mesh(position, direction, g, distance); + + if (triangle_index == -1) + distances[id] = 1e9; + else + distances[id] = distance; } -__global__ void color_solids(int first_triangle, int nthreads, - int *solid_id_map, - bool *solid_hit, - unsigned int *solid_colors) +__global__ void +color_solids(int first_triangle, int nthreads, int *solid_id_map, + bool *solid_hit, unsigned int *solid_colors, Geometry *geometry) { - int id = blockIdx.x*blockDim.x + threadIdx.x; + int id = blockIdx.x*blockDim.x + threadIdx.x; - if (id >= nthreads) - return; + if (id >= nthreads) + return; - int triangle_id = first_triangle + id; - int solid_id = solid_id_map[triangle_id]; - if (solid_hit[solid_id]) - g_colors[triangle_id] = solid_colors[solid_id]; + int triangle_id = first_triangle + id; + int solid_id = solid_id_map[triangle_id]; + if (solid_hit[solid_id]) + geometry->colors[triangle_id] = solid_colors[solid_id]; } -} // extern "c" +} // extern "C" #endif |