diff options
author | Anthony LaTorre <telatorre@gmail.com> | 2011-05-15 18:35:41 -0400 |
---|---|---|
committer | Anthony LaTorre <telatorre@gmail.com> | 2011-05-15 18:35:41 -0400 |
commit | 2aa342458d1278487f9ca47ff0111e74b20d63ef (patch) | |
tree | 63f2bfbbc1f971525cf1f79c06e74d969484eaa2 /src | |
parent | 8ea783d053e817568b3e7d04f28a6fd2583f18cf (diff) | |
download | chroma-2aa342458d1278487f9ca47ff0111e74b20d63ef.tar.gz chroma-2aa342458d1278487f9ca47ff0111e74b20d63ef.tar.bz2 chroma-2aa342458d1278487f9ca47ff0111e74b20d63ef.zip |
cleanup. fixed tests
Diffstat (limited to 'src')
-rw-r--r-- | src/kernel.cu | 31 | ||||
-rw-r--r-- | src/matrix.h | 4 | ||||
-rw-r--r-- | src/rotate.h | 3 |
3 files changed, 18 insertions, 20 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; } } diff --git a/src/matrix.h b/src/matrix.h index 5c72344..a3e480b 100644 --- a/src/matrix.h +++ b/src/matrix.h @@ -197,9 +197,7 @@ __device__ __host__ Matrix operator/ (const float &c, const Matrix &m) __device__ __host__ float det(const Matrix &m) { - return m.a00*(m.a11*m.a22 - m.a12*m.a21) - \ - m.a10*(m.a01*m.a22 - m.a02*m.a21) + \ - m.a20*(m.a01*m.a12 - m.a02*m.a11); + return m.a00*(m.a11*m.a22 - m.a12*m.a21) - m.a10*(m.a01*m.a22 - m.a02*m.a21) + m.a20*(m.a01*m.a12 - m.a02*m.a11); } __device__ __host__ Matrix inv(const Matrix &m) diff --git a/src/rotate.h b/src/rotate.h index fec76a8..7fc6f7e 100644 --- a/src/rotate.h +++ b/src/rotate.h @@ -1,6 +1,9 @@ #ifndef __ROTATE_H__ #define __ROTATE_H__ +#include "linalg.h" +#include "matrix.h" + __device__ const Matrix IDENTITY_MATRIX = {1,0,0,0,1,0,0,0,1}; __device__ __host__ Matrix make_rotation_matrix(float phi, const float3 &n) |