summaryrefslogtreecommitdiff
path: root/src/kernel.cu
diff options
context:
space:
mode:
Diffstat (limited to 'src/kernel.cu')
-rw-r--r--src/kernel.cu34
1 files changed, 33 insertions, 1 deletions
diff --git a/src/kernel.cu b/src/kernel.cu
index fe518f6..a26460a 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;