diff options
author | Anthony LaTorre <telatorre@gmail.com> | 2011-05-07 20:32:35 -0400 |
---|---|---|
committer | Anthony LaTorre <telatorre@gmail.com> | 2011-05-07 20:32:35 -0400 |
commit | 16736eeae08d9627ad751c65919900ae9191a08f (patch) | |
tree | 602af6b7fdfd32c9cdc7fc3b85e6c8e247209197 /src/intersect.cu | |
parent | 76a6dd33cdaf4b583b7e8353f198925ddb4a4685 (diff) | |
download | chroma-16736eeae08d9627ad751c65919900ae9191a08f.tar.gz chroma-16736eeae08d9627ad751c65919900ae9191a08f.tar.bz2 chroma-16736eeae08d9627ad751c65919900ae9191a08f.zip |
tie fighter
Diffstat (limited to 'src/intersect.cu')
-rw-r--r-- | src/intersect.cu | 118 |
1 files changed, 118 insertions, 0 deletions
diff --git a/src/intersect.cu b/src/intersect.cu new file mode 100644 index 0000000..844f6fd --- /dev/null +++ b/src/intersect.cu @@ -0,0 +1,118 @@ +//-*-c-*- + +__device__ __host__ Matrix inv(const Matrix&m, float& determinant) +{ + determinant = det(m); + + return make_matrix(m.a11*m.a22 - m.a12*m.a21, + m.a02*m.a21 - m.a01*m.a22, + m.a01*m.a12 - m.a02*m.a11, + m.a12*m.a20 - m.a10*m.a22, + m.a00*m.a22 - m.a02*m.a20, + m.a02*m.a10 - m.a00*m.a12, + m.a10*m.a21 - m.a11*m.a20, + m.a01*m.a20 - m.a00*m.a21, + m.a00*m.a11 - m.a01*m.a10)/determinant; +} + +__device__ bool intersect_triangle(const float3 &x, const float3 &p, float3 *vertex, float3 &intersection) +{ + float determinant; + float3 u = inv(make_matrix(vertex[1]-vertex[0],vertex[2]-vertex[0],-p), determinant)*(x-vertex[0]); + + if (determinant == 0.0) + return false; + + if (u.x < 0.0 || u.y < 0.0 || u.z < 0.0 || (1-u.x-u.y) < 0.0) + return false; + + intersection = x + p*u.z; + + return true; +} + +__device__ int get_color(const float3 &p, float3 *vertex) +{ + float3 v1 = vertex[1] - vertex[0]; + float3 v2 = vertex[2] - vertex[0]; + + float3 normal = cross(v1,v2); + + float scale; + scale = dot(normal,-p)/(norm(normal)*norm(p)); + if (scale < 0.0) + scale = dot(-normal,-p)/(norm(normal)*norm(p)); + + int rgb = floorf(255*scale); + + return rgb*65536 + rgb*256 + rgb; +} + +extern "C" +{ + +__global__ void translate(int max_idx, float3 *x, float3 v) +{ + int idx = blockIdx.x*blockDim.x + threadIdx.x; + + if (idx > max_idx) + return; + + x[idx] += v; +} + +__global__ void rotate(int max_idx, float3 *x, float phi, float3 axis) +{ + int idx = blockIdx.x*blockDim.x + threadIdx.x; + + if (idx > max_idx) + return; + + x[idx] = rotate(x[idx], phi, axis); +} + +__global__ void intersect_triangle_mesh(int max_idx, float3 *x, float3 *p, int n, float3 *mesh, int *pixel) +{ + int idx = blockIdx.x*blockDim.x + threadIdx.x; + + if (idx > max_idx) + return; + + bool hit = false; + + float distance, min_distance; + float3 intersection, min_intersection; + + int i; + for (i=0; i < n; i++) + { + if (intersect_triangle(x[idx], p[idx], mesh+3*i, intersection)) + { + if (!hit) + { + pixel[idx] = get_color(p[idx], mesh+3*i); + + min_distance = norm(intersection-x[idx]); + min_intersection = intersection; + hit = true; + continue; + } + + distance = norm(intersection-x[idx]); + + if (distance < min_distance) + { + pixel[idx] = get_color(p[idx], mesh+3*i); + + min_distance = distance; + min_intersection = intersection; + } + } + } + + if (!hit) + pixel[idx] = 0; + +} + +} // extern "c" |