summaryrefslogtreecommitdiff
path: root/src
diff options
context:
space:
mode:
authorAnthony LaTorre <telatorre@gmail.com>2011-05-09 11:49:25 -0400
committerAnthony LaTorre <telatorre@gmail.com>2011-05-09 11:49:25 -0400
commit4d3bf194aa395e818fc67230b219f41e454c435f (patch)
treeed412e97788a470660f148e974af38b10840b96a /src
parent152240722edb92f84de18818c137804b7db63e6a (diff)
downloadchroma-4d3bf194aa395e818fc67230b219f41e454c435f.tar.gz
chroma-4d3bf194aa395e818fc67230b219f41e454c435f.tar.bz2
chroma-4d3bf194aa395e818fc67230b219f41e454c435f.zip
slight speed increase
Diffstat (limited to 'src')
-rw-r--r--src/intersect.cu33
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;
}