summaryrefslogtreecommitdiff
path: root/src/intersect.cu
diff options
context:
space:
mode:
authorAnthony LaTorre <telatorre@gmail.com>2011-05-07 20:32:35 -0400
committerAnthony LaTorre <telatorre@gmail.com>2011-05-07 20:32:35 -0400
commit16736eeae08d9627ad751c65919900ae9191a08f (patch)
tree602af6b7fdfd32c9cdc7fc3b85e6c8e247209197 /src/intersect.cu
parent76a6dd33cdaf4b583b7e8353f198925ddb4a4685 (diff)
downloadchroma-16736eeae08d9627ad751c65919900ae9191a08f.tar.gz
chroma-16736eeae08d9627ad751c65919900ae9191a08f.tar.bz2
chroma-16736eeae08d9627ad751c65919900ae9191a08f.zip
tie fighter
Diffstat (limited to 'src/intersect.cu')
-rw-r--r--src/intersect.cu118
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"