diff options
Diffstat (limited to 'src')
| -rw-r--r-- | src/intersect.cu (renamed from src/kernel.cu) | 44 |
1 files changed, 37 insertions, 7 deletions
diff --git a/src/kernel.cu b/src/intersect.cu index c2b3fb2..1402aa1 100644 --- a/src/kernel.cu +++ b/src/intersect.cu @@ -5,11 +5,14 @@ #include "matrix.h" #include "rotate.h" +/* flattened triangle mesh */ texture<float4, 1, cudaReadModeElementType> mesh; +/* lower/upper bounds for the bounding box associated with each node/leaf */ texture<float4, 1, cudaReadModeElementType> upper_bound_arr; texture<float4, 1, cudaReadModeElementType> lower_bound_arr; +/* map to child nodes/triangles and the number of child nodes/triangles */ texture<uint, 1, cudaReadModeElementType> child_map_arr; texture<uint, 1, cudaReadModeElementType> child_len_arr; @@ -18,6 +21,12 @@ __device__ float3 make_float3(const float4 &a) return make_float3(a.x, a.y, a.z); } +/* Test the intersection between a ray starting from `origin` traveling in the + direction `direction` and a triangle defined by the vertices `v0`, `v1`, and + `v2`. If the ray intersects the triangle, set `distance` to the distance + between `origin` and the intersection and return true, else return false. + + `direction` must be normalized. */ __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); @@ -29,26 +38,38 @@ __device__ bool intersect_triangle(const float3 &origin, const float3 &direction float3 b = origin-v0; - float u1 = ((m.a11*m.a22 - m.a12*m.a21)*b.x + (m.a02*m.a21 - m.a01*m.a22)*b.y + (m.a01*m.a12 - m.a02*m.a11)*b.z)/determinant; + float u1 = ((m.a11*m.a22 - m.a12*m.a21)*b.x + + (m.a02*m.a21 - m.a01*m.a22)*b.y + + (m.a01*m.a12 - m.a02*m.a11)*b.z)/determinant; if (u1 < 0.0f) return false; - float u2 = ((m.a12*m.a20 - m.a10*m.a22)*b.x + (m.a00*m.a22 - m.a02*m.a20)*b.y + (m.a02*m.a10 - m.a00*m.a12)*b.z)/determinant; + float u2 = ((m.a12*m.a20 - m.a10*m.a22)*b.x + + (m.a00*m.a22 - m.a02*m.a20)*b.y + + (m.a02*m.a10 - m.a00*m.a12)*b.z)/determinant; if (u2 < 0.0f) return false; - float u3 = ((m.a10*m.a21 - m.a11*m.a20)*b.x + (m.a01*m.a20 - m.a00*m.a21)*b.y + (m.a00*m.a11 - m.a01*m.a10)*b.z)/determinant; + float u3 = ((m.a10*m.a21 - m.a11*m.a20)*b.x + + (m.a01*m.a20 - m.a00*m.a21)*b.y + + (m.a00*m.a11 - m.a01*m.a10)*b.z)/determinant; if (u3 < 0.0f || (1-u1-u2) < 0.0f) return false; - distance = norm(direction*u3); + distance = u3; return true; } +/* Return the 32 bit color associated with the intersection between a ray + starting from `origin` traveling in the direction `direction` and the + plane defined by the points `v0`, `v1`, and `v2` using the cosine of the + angle between the ray and the plane normal to determine the brightness. + + `direction` must be normalized. */ __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); @@ -61,7 +82,10 @@ __device__ int get_color(const float3 &direction, const float3 &v0, const float3 return rgb << 16 | rgb << 8 | rgb; } - +/* Test the intersection between a ray starting from `origin` traveling in the + direction `direction` and the axis-aligned box defined by the opposite + vertices `lower_bound` and `upper_bound`. If the ray intersects the box + return True, else return False. */ __device__ bool intersect_box(const float3 &origin, const float3 &direction, const float3 &lower_bound, const float3 &upper_bound) { float kmin, kmax, kymin, kymax, kzmin, kzmax; @@ -132,6 +156,9 @@ __device__ bool intersect_box(const float3 &origin, const float3 &direction, con return true; } +/* 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) { float3 lower_bound = make_float3(tex1Dfetch(lower_bound_arr, i)); @@ -236,7 +263,7 @@ __global__ void intersect_mesh(int max_idx, float3 *origin_arr, float3 *directio if (!hit) { *pixel = get_color(direction, v0, v1, v2); - *state = mesh_idx; + *state = child_map + i; min_distance = distance; @@ -247,7 +274,7 @@ __global__ void intersect_mesh(int max_idx, float3 *origin_arr, float3 *directio if (distance < min_distance) { *pixel = get_color(direction, v0, v1, v2); - *state = mesh_idx; + *state = child_map + i; min_distance = distance; } @@ -263,7 +290,10 @@ __global__ void intersect_mesh(int max_idx, float3 *origin_arr, float3 *directio while (node != head); if (!hit) + { + *state = -1; *pixel = 0; + } } // intersect mesh |