diff options
Diffstat (limited to 'src/kernel.cu')
-rw-r--r-- | src/kernel.cu | 31 |
1 files changed, 14 insertions, 17 deletions
diff --git a/src/kernel.cu b/src/kernel.cu index 02c52cf..d8f2300 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -7,8 +7,6 @@ texture<float4, 1, cudaReadModeElementType> mesh; - - texture<float4, 1, cudaReadModeElementType> upper_bound_arr; texture<float4, 1, cudaReadModeElementType> lower_bound_arr; @@ -20,7 +18,7 @@ __device__ float3 make_float3(const float4 &a) return make_float3(a.x, a.y, a.z); } -__device__ bool intersect_triangle(const float3 &origin, const float3 &direction, const float3 &v0, const float3 &v1, const float3 &v2, float3 &intersection) +__device__ bool intersect_triangle(const float3 &origin, const float3 &direction, const float3 &v0, const float3 &v1, const float3 &v2, float &distance) { Matrix m = make_matrix(v1-v0, v2-v0, -direction); @@ -46,7 +44,7 @@ __device__ bool intersect_triangle(const float3 &origin, const float3 &direction if (u3 < 0.0f || (1-u1-u2) < 0.0f) return false; - intersection = origin + direction*u3; + distance = norm(direction*u3); return true; } @@ -54,13 +52,12 @@ __device__ bool intersect_triangle(const float3 &origin, const float3 &direction __device__ int get_color(const float3 &direction, const float3 &v0, const float3& v1, const float3 &v2) { float scale = 255*dot(normalize(cross(v1-v0,v2-v0)),-direction); + if (scale < 0.0f) scale = 255*dot(-normalize(cross(v1-v0,v2-v0)),-direction); - int rgb = floorf(scale); - return rgb << 16 | rgb << 8 | rgb; } @@ -166,7 +163,7 @@ __global__ void rotate(int max_idx, float3 *pt, float phi, float3 axis) pt[idx] = rotate(pt[idx], phi, axis); } -__global__ void intersect_mesh(int max_idx, float3 *origin_arr, float3 *direction_arr, int *pixel_arr, int first_leaf) +__global__ void intersect_mesh(int max_idx, float3 *origin_arr, float3 *direction_arr, int *pixel_arr, int first_leaf, int *state_arr) { int idx = blockIdx.x*blockDim.x + threadIdx.x; @@ -177,11 +174,12 @@ __global__ void intersect_mesh(int max_idx, float3 *origin_arr, float3 *directio float3 direction = direction_arr[idx]; int *pixel = pixel_arr+idx; + int *state = state_arr+idx; bool hit = false; - float distance, min_distance; - float3 intersection, min_intersection; + float distance; + float min_distance; if (!intersect_node(origin, direction, 0)) { @@ -226,32 +224,31 @@ __global__ void intersect_mesh(int max_idx, float3 *origin_arr, float3 *directio { for (i=0; i < child_len; i++) { - int mesh_idx = child_map + 3*i; + int mesh_idx = 3*(child_map + i); float3 v0 = make_float3(tex1Dfetch(mesh, mesh_idx)); float3 v1 = make_float3(tex1Dfetch(mesh, mesh_idx+1)); float3 v2 = make_float3(tex1Dfetch(mesh, mesh_idx+2)); - if (intersect_triangle(origin, direction, v0, v1, v2, intersection)) + if (intersect_triangle(origin, direction, v0, v1, v2, distance)) { if (!hit) { *pixel = get_color(direction, v0, v1, v2); - - min_distance = norm(intersection-origin); - min_intersection = intersection; + *state = mesh_idx; + + min_distance = distance; + hit = true; continue; } - distance = norm(intersection-origin); - if (distance < min_distance) { *pixel = get_color(direction, v0, v1, v2); + *state = mesh_idx; min_distance = distance; - min_intersection = intersection; } } |