From fa8d1082f9d989f2a3819540a9bf30dc67618709 Mon Sep 17 00:00:00 2001 From: Anthony LaTorre Date: Thu, 25 Aug 2011 21:42:04 -0400 Subject: add 3d support to camera views by displaying images as an anaglyph. alpha coloring is now calculated using the new searching/sorting algorithms. --- src/kernel.cu | 16 +++++++++++++--- 1 file changed, 13 insertions(+), 3 deletions(-) (limited to 'src/kernel.cu') diff --git a/src/kernel.cu b/src/kernel.cu index 307412e..71b4153 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -239,7 +239,7 @@ __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) +__global__ void ray_trace(int nthreads, float3 *positions, float3 *directions, int *pixels, float *distances) { int id = blockIdx.x*blockDim.x + threadIdx.x; @@ -257,6 +257,7 @@ __global__ void ray_trace(int nthreads, float3 *positions, float3 *directions, i if (triangle_index == -1) { pixels[id] = 0; + distances[id] = 1e9f; } else { @@ -267,6 +268,7 @@ __global__ void ray_trace(int nthreads, float3 *positions, float3 *directions, i float3 v2 = g_vertices[triangle_data.z]; pixels[id] = get_color(direction, v0, v1, v2, g_colors[triangle_index]); + distances[id] = distance; } } // ray_trace @@ -277,7 +279,7 @@ __global__ void ray_trace(int nthreads, float3 *positions, float3 *directions, i 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_alpha(int nthreads, float3 *positions, float3 *directions, int *pixels) +__global__ void ray_trace_alpha(int nthreads, float3 *positions, float3 *directions, int *pixels, float *distances) { int id = blockIdx.x*blockDim.x + threadIdx.x; @@ -288,7 +290,15 @@ __global__ void ray_trace_alpha(int nthreads, float3 *positions, float3 *directi float3 direction = directions[id]; direction /= norm(direction); - pixels[id] = get_color_alpha(position, direction); + bool hit; + float distance; + + pixels[id] = get_color_alpha(position, direction, hit, distance); + + if (hit) + distances[id] = distance; + else + distances[id] = 1e9; } // ray_trace -- cgit From 46f7f58dd6cf3c008e3ef0496f0ee60b52db6941 Mon Sep 17 00:00:00 2001 From: Anthony LaTorre Date: Fri, 26 Aug 2011 01:27:43 -0400 Subject: no more 3d headache! 3d viewing angle changes depending on the distance to the object in the center of the screen. --- src/kernel.cu | 35 +++++++++++++++++++++++++---------- 1 file changed, 25 insertions(+), 10 deletions(-) (limited to 'src/kernel.cu') diff --git a/src/kernel.cu b/src/kernel.cu index 71b4153..d36d260 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -239,7 +239,29 @@ __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, float *distances) +__global__ void distance_to_mesh(int nthreads, float3 *positions, float3 *directions, float *distances) +{ + 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) + distances[id] = 1e9; + else + distances[id] = distance; + +} // distance_to_mesh + +__global__ void ray_trace(int nthreads, float3 *positions, float3 *directions, int *pixels) { int id = blockIdx.x*blockDim.x + threadIdx.x; @@ -257,7 +279,6 @@ __global__ void ray_trace(int nthreads, float3 *positions, float3 *directions, i if (triangle_index == -1) { pixels[id] = 0; - distances[id] = 1e9f; } else { @@ -268,7 +289,6 @@ __global__ void ray_trace(int nthreads, float3 *positions, float3 *directions, i float3 v2 = g_vertices[triangle_data.z]; pixels[id] = get_color(direction, v0, v1, v2, g_colors[triangle_index]); - distances[id] = distance; } } // ray_trace @@ -279,7 +299,7 @@ __global__ void ray_trace(int nthreads, float3 *positions, float3 *directions, i 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_alpha(int nthreads, float3 *positions, float3 *directions, int *pixels, float *distances) +__global__ void ray_trace_alpha(int nthreads, float3 *positions, float3 *directions, int *pixels) { int id = blockIdx.x*blockDim.x + threadIdx.x; @@ -293,12 +313,7 @@ __global__ void ray_trace_alpha(int nthreads, float3 *positions, float3 *directi bool hit; float distance; - pixels[id] = get_color_alpha(position, direction, hit, distance); - - if (hit) - distances[id] = distance; - else - distances[id] = 1e9; + pixels[id] = get_color_alpha(position, direction); } // ray_trace -- cgit From 707ca1b366f11032682cc864ca2848905e6b485c Mon Sep 17 00:00:00 2001 From: Anthony LaTorre Date: Fri, 2 Sep 2011 12:12:38 -0400 Subject: update event structure. break gpu.GPU class into separate smaller independent classes. --- src/kernel.cu | 35 ----------------------------------- 1 file changed, 35 deletions(-) (limited to 'src/kernel.cu') diff --git a/src/kernel.cu b/src/kernel.cu index d36d260..4418305 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -62,41 +62,6 @@ __device__ void to_diffuse(Photon &p, State &s, curandState &rng, const int &max extern "C" { -/* Translate the points `a` by the vector `v` */ -__global__ void translate(int nthreads, float3 *a, float3 v) -{ - int id = blockIdx.x*blockDim.x + threadIdx.x; - - if (id >= nthreads) - return; - - a[id] += v; -} - -/* Rotate the points `a` through an angle `phi` counter-clockwise about the - axis `axis` (when looking towards +infinity). */ -__global__ void rotate(int nthreads, float3 *a, float phi, float3 axis) -{ - int id = blockIdx.x*blockDim.x + threadIdx.x; - - if (id >= nthreads) - return; - - a[id] = rotate(a[id], phi, axis); -} - -__global__ void rotate_around_point(int nthreads, float3 *a, float phi, float3 axis, float3 point) -{ - int id = blockIdx.x*blockDim.x + threadIdx.x; - - if (id >= nthreads) - return; - - a[id] -= point; - a[id] = rotate(a[id], phi, axis); - a[id] += point; -} - __global__ void update_xyz_lookup(int nthreads, int total_threads, int offset, float3 position, curandState *rng_states, float wavelength, float3 xyz, float3 *xyz_lookup1, float3 *xyz_lookup2, int max_steps) { int kernel_id = blockIdx.x*blockDim.x + threadIdx.x; -- cgit