diff options
author | Anthony LaTorre <telatorre@gmail.com> | 2011-05-09 11:49:25 -0400 |
---|---|---|
committer | Anthony LaTorre <telatorre@gmail.com> | 2011-05-09 11:49:25 -0400 |
commit | 4d3bf194aa395e818fc67230b219f41e454c435f (patch) | |
tree | ed412e97788a470660f148e974af38b10840b96a /src | |
parent | 152240722edb92f84de18818c137804b7db63e6a (diff) | |
download | chroma-4d3bf194aa395e818fc67230b219f41e454c435f.tar.gz chroma-4d3bf194aa395e818fc67230b219f41e454c435f.tar.bz2 chroma-4d3bf194aa395e818fc67230b219f41e454c435f.zip |
slight speed increase
Diffstat (limited to 'src')
-rw-r--r-- | src/intersect.cu | 33 |
1 files changed, 23 insertions, 10 deletions
diff --git a/src/intersect.cu b/src/intersect.cu index 844f6fd..aa4f732 100644 --- a/src/intersect.cu +++ b/src/intersect.cu @@ -1,6 +1,11 @@ //-*-c-*- -__device__ __host__ Matrix inv(const Matrix&m, float& determinant) +__device__ float3 make_float3(const float4 &a) +{ + return make_float3(a.x, a.y, a.z); +} + +__device__ Matrix inv(const Matrix&m, float& determinant) { determinant = det(m); @@ -15,10 +20,14 @@ __device__ __host__ Matrix inv(const Matrix&m, float& determinant) m.a00*m.a11 - m.a01*m.a10)/determinant; } -__device__ bool intersect_triangle(const float3 &x, const float3 &p, float3 *vertex, float3 &intersection) +__device__ bool intersect_triangle(const float3 &x, const float3 &p, float3 *triangle, float3 &intersection) { + float3 v0 = triangle[0]; + float3 v1 = triangle[1]; + float3 v2 = triangle[2]; + float determinant; - float3 u = inv(make_matrix(vertex[1]-vertex[0],vertex[2]-vertex[0],-p), determinant)*(x-vertex[0]); + float3 u = inv(make_matrix(v1-v0,v2-v0,-p), determinant)*(x-v0); if (determinant == 0.0) return false; @@ -71,13 +80,17 @@ __global__ void rotate(int max_idx, float3 *x, float phi, float3 axis) 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) +__global__ void intersect_triangle_mesh(int max_idx, float3 *xarr, float3 *parr, int n, float3* mesh, int *pixelarr) { int idx = blockIdx.x*blockDim.x + threadIdx.x; if (idx > max_idx) return; + float3 x = xarr[idx]; + float3 p = parr[idx]; + int *pixel = pixelarr+idx; + bool hit = false; float distance, min_distance; @@ -86,23 +99,23 @@ __global__ void intersect_triangle_mesh(int max_idx, float3 *x, float3 *p, int n int i; for (i=0; i < n; i++) { - if (intersect_triangle(x[idx], p[idx], mesh+3*i, intersection)) + if (intersect_triangle(x, p, mesh+3*i, intersection)) { if (!hit) { - pixel[idx] = get_color(p[idx], mesh+3*i); + *pixel = get_color(p, mesh+3*i); - min_distance = norm(intersection-x[idx]); + min_distance = norm(intersection-x); min_intersection = intersection; hit = true; continue; } - distance = norm(intersection-x[idx]); + distance = norm(intersection-x); if (distance < min_distance) { - pixel[idx] = get_color(p[idx], mesh+3*i); + *pixel = get_color(p, mesh+3*i); min_distance = distance; min_intersection = intersection; @@ -111,7 +124,7 @@ __global__ void intersect_triangle_mesh(int max_idx, float3 *x, float3 *p, int n } if (!hit) - pixel[idx] = 0; + pixelarr[idx] = 0; } |