summaryrefslogtreecommitdiff
path: root/src
diff options
context:
space:
mode:
authorAnthony LaTorre <telatorre@gmail.com>2011-05-15 18:35:41 -0400
committerAnthony LaTorre <telatorre@gmail.com>2011-05-15 18:35:41 -0400
commit2aa342458d1278487f9ca47ff0111e74b20d63ef (patch)
tree63f2bfbbc1f971525cf1f79c06e74d969484eaa2 /src
parent8ea783d053e817568b3e7d04f28a6fd2583f18cf (diff)
downloadchroma-2aa342458d1278487f9ca47ff0111e74b20d63ef.tar.gz
chroma-2aa342458d1278487f9ca47ff0111e74b20d63ef.tar.bz2
chroma-2aa342458d1278487f9ca47ff0111e74b20d63ef.zip
cleanup. fixed tests
Diffstat (limited to 'src')
-rw-r--r--src/kernel.cu31
-rw-r--r--src/matrix.h4
-rw-r--r--src/rotate.h3
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)