summaryrefslogtreecommitdiff
path: root/src
diff options
context:
space:
mode:
Diffstat (limited to 'src')
-rw-r--r--src/alpha.h4
-rw-r--r--src/intersect.h6
-rw-r--r--src/kernel.cu34
-rw-r--r--src/mesh.h29
4 files changed, 62 insertions, 11 deletions
diff --git a/src/alpha.h b/src/alpha.h
index a9963eb..e4a3a40 100644
--- a/src/alpha.h
+++ b/src/alpha.h
@@ -57,7 +57,7 @@ __device__ int get_color_alpha(const float3 &origin, const float3& direction)
float distance;
- if (!intersect_node(origin, direction, g_start_node))
+ if (!intersect_node(origin, direction, g_start_node, -1.0f))
return 0;
unsigned int stack[STACK_SIZE];
@@ -85,7 +85,7 @@ __device__ int get_color_alpha(const float3 &origin, const float3& direction)
{
for (i=first_child; i < stop; i++)
{
- if (intersect_node(origin, direction, i))
+ if (intersect_node(origin, direction, i, -1.0f))
{
*node = i;
node++;
diff --git a/src/intersect.h b/src/intersect.h
index c400f1c..d2f54ce 100644
--- a/src/intersect.h
+++ b/src/intersect.h
@@ -60,7 +60,7 @@ __device__ bool intersect_triangle(const float3 &origin, const float3 &direction
angle between the ray and the plane normal to determine the brightness.
`direction` must be normalized. */
-__device__ int get_color(const float3 &direction, const float3 &v0, const float3& v1, const float3 &v2, const int base_color=0xFFFFFFFF)
+__device__ unsigned int get_color(const float3 &direction, const float3 &v0, const float3& v1, const float3 &v2, const unsigned int base_color=0xFFFFFFFF)
{
float scale = dot(normalize(cross(v1-v0,v2-v1)),-direction);
@@ -82,7 +82,7 @@ __device__ int get_color(const float3 &direction, const float3 &v0, const float3
direction `direction` and the axis-aligned box defined by the opposite
vertices `lower_bound` and `upper_bound`. If the ray intersects the box
return True, else return False. */
-__device__ bool intersect_box(const float3 &origin, const float3 &direction, const float3 &lower_bound, const float3 &upper_bound)
+__device__ bool intersect_box(const float3 &origin, const float3 &direction, const float3 &lower_bound, const float3 &upper_bound, float& distance_to_box)
{
float kmin, kmax, kymin, kymax, kzmin, kzmax;
@@ -149,6 +149,8 @@ __device__ bool intersect_box(const float3 &origin, const float3 &direction, con
if (kmin > kmax)
return false;
+ distance_to_box = kmin;
+
return true;
}
diff --git a/src/kernel.cu b/src/kernel.cu
index f60ecb1..cc567b8 100644
--- a/src/kernel.cu
+++ b/src/kernel.cu
@@ -239,13 +239,45 @@ __global__ void process_image(int nthreads, float3 *image, int *pixels, int nima
} // process_image
+__global__ void ray_trace(int nthreads, float3 *positions, float3 *directions, int *pixels)
+{
+ int id = blockIdx.x*blockDim.x + threadIdx.x;
+
+ if (id >= nthreads)
+ return;
+
+ float3 position = positions[id];
+ float3 direction = directions[id];
+ direction /= norm(direction);
+
+ float distance;
+
+ int triangle_index = intersect_mesh(position, direction, distance);
+
+ if (triangle_index == -1)
+ {
+ pixels[id] = 0;
+ }
+ else
+ {
+ uint4 triangle_data = g_triangles[triangle_index];
+
+ float3 v0 = g_vertices[triangle_data.x];
+ float3 v1 = g_vertices[triangle_data.y];
+ float3 v2 = g_vertices[triangle_data.z];
+
+ pixels[id] = get_color(direction, v0, v1, v2, g_colors[triangle_index]);
+ }
+
+} // ray_trace
+
/* Trace the rays starting at `positions` traveling in the direction
`directions` to their intersection with the global mesh. If the ray
intersects the mesh set the pixel associated with the ray to a 32 bit
color whose brightness is determined by the cosine of the angle between
the ray and the normal of the triangle it intersected, else set the pixel
to 0. */
-__global__ void ray_trace(int nthreads, float3 *positions, float3 *directions, int *pixels)
+__global__ void ray_trace_alpha(int nthreads, float3 *positions, float3 *directions, int *pixels)
{
int id = blockIdx.x*blockDim.x + threadIdx.x;
diff --git a/src/mesh.h b/src/mesh.h
index bb30bef..a2bd207 100644
--- a/src/mesh.h
+++ b/src/mesh.h
@@ -36,13 +36,29 @@ __device__ int convert(int c)
/* Test the intersection between a ray starting at `origin` traveling in the
direction `direction` and the bounding box around node `i`. If the ray
intersects the bounding box return true, else return false. */
-__device__ bool intersect_node(const float3 &origin, const float3 &direction, const int &i)
+__device__ bool intersect_node(const float3 &origin, const float3 &direction, const int &i, const float &min_distance)
{
float3 lower_bound = g_lower_bounds[i];
float3 upper_bound = g_upper_bounds[i];
- return intersect_box(origin, direction, lower_bound, upper_bound);
-}
+ float distance_to_box;
+
+ if (intersect_box(origin, direction, lower_bound, upper_bound, distance_to_box))
+ {
+ if (min_distance < 0.0f)
+ return true;
+
+ if (distance_to_box > min_distance)
+ return false;
+
+ return true;
+ }
+ else
+ {
+ return false;
+ }
+
+} // intersect_node
/* Find the intersection between a ray starting at `origin` traveling in the
direction `direction` and the global mesh texture. If the ray intersects
@@ -53,8 +69,9 @@ __device__ int intersect_mesh(const float3 &origin, const float3& direction, flo
int triangle_index = -1;
float distance;
+ min_distance = -1.0f;
- if (!intersect_node(origin, direction, g_start_node))
+ if (!intersect_node(origin, direction, g_start_node, min_distance))
return -1;
unsigned int stack[STACK_SIZE];
@@ -82,7 +99,7 @@ __device__ int intersect_mesh(const float3 &origin, const float3& direction, flo
{
for (i=first_child; i < stop; i++)
{
- if (intersect_node(origin, direction, i))
+ if (intersect_node(origin, direction, i, min_distance))
{
*node = i;
node++;
@@ -134,7 +151,7 @@ __device__ int intersect_mesh(const float3 &origin, const float3& direction, flo
extern "C"
{
- __global__ void set_global_mesh_variables(uint4 *triangles, float3 *vertices, unsigned int *colors, unsigned int start_node, unsigned int first_node, float3 *lower_bounds, float3 *upper_bounds)
+__global__ void set_global_mesh_variables(uint4 *triangles, float3 *vertices, unsigned int *colors, unsigned int start_node, unsigned int first_node, float3 *lower_bounds, float3 *upper_bounds)
{
g_triangles = triangles;
g_vertices = vertices;