diff options
Diffstat (limited to 'src')
| -rw-r--r-- | src/alpha.h | 99 | ||||
| -rw-r--r-- | src/kernel.cu | 16 |
2 files changed, 55 insertions, 60 deletions
diff --git a/src/alpha.h b/src/alpha.h index e4a3a40..263fa1e 100644 --- a/src/alpha.h +++ b/src/alpha.h @@ -4,59 +4,14 @@ #include "linalg.h" #include "intersect.h" #include "mesh.h" +#include "sorting.h" -template <class T> -__device__ void swap(T &a, T &b) -{ - T temp = a; - a = b; - b = temp; -} +#include "stdio.h" #define ALPHA_DEPTH 10 -struct HitList -{ - unsigned int size; - unsigned int indices[ALPHA_DEPTH]; - float distances[ALPHA_DEPTH]; -}; - -__device__ void add_to_hit_list(const float &distance, const int &index, HitList &h) -{ - unsigned int i; - if (h.size >= ALPHA_DEPTH) - { - if (distance > h.distances[ALPHA_DEPTH-1]) - return; - - i = ALPHA_DEPTH-1; - } - else - { - i = h.size; - h.size += 1; - } - - h.indices[i] = index; - h.distances[i] = distance; - - while (i > 0 && h.distances[i-1] > h.distances[i]) - { - swap(h.distances[i-1], h.distances[i]); - swap(h.indices[i-1], h.indices[i]); - - i -= 1; - } -} - -__device__ int get_color_alpha(const float3 &origin, const float3& direction) +__device__ int get_color_alpha(const float3 &origin, const float3& direction, bool &hit, float &distance) { - HitList h; - h.size = 0; - - float distance; - if (!intersect_node(origin, direction, g_start_node, -1.0f)) return 0; @@ -69,6 +24,10 @@ __device__ int get_color_alpha(const float3 &origin, const float3& direction) unsigned int i; + float distances[ALPHA_DEPTH]; + unsigned int indices[ALPHA_DEPTH]; + unsigned int n=0; + do { unsigned int first_child = tex1Dfetch(node_map, *node); @@ -106,9 +65,26 @@ __device__ int get_color_alpha(const float3 &origin, const float3& direction) if (intersect_triangle(origin, direction, v0, v1, v2, distance)) { - add_to_hit_list(distance, i, h); + if (n < 1) + { + distances[0] = distance; + indices[0] = i; + } + else + { + unsigned long j = searchsorted(n, distances, distance); + + if (j <= ALPHA_DEPTH-1) + { + insert(ALPHA_DEPTH, distances, j, distance); + insert(ALPHA_DEPTH, indices, j, i); + } + } + + if (n < ALPHA_DEPTH) + n++; } - + } // triangle loop node--; @@ -118,16 +94,25 @@ __device__ int get_color_alpha(const float3 &origin, const float3& direction) } // while loop while (node != head); - if (h.size < 1) + if (n < 1) + { + hit = false; return 0; + } + + hit = true; + distance = distances[0]; float scale = 1.0f; float fr = 0.0f; float fg = 0.0f; float fb = 0.0f; - for (i=0; i < h.size; i++) + unsigned int index; + for (i=0; i < n; i++) { - uint4 triangle_data = g_triangles[h.indices[i]]; + index = indices[i]; + + uint4 triangle_data = g_triangles[index]; float3 v0 = g_vertices[triangle_data.x]; float3 v1 = g_vertices[triangle_data.y]; @@ -138,11 +123,11 @@ __device__ int get_color_alpha(const float3 &origin, const float3& direction) if (cos_theta < 0.0f) cos_theta = dot(-normalize(cross(v1-v0,v2-v1)),-direction); - unsigned int r0 = 0xff & (g_colors[h.indices[i]] >> 16); - unsigned int g0 = 0xff & (g_colors[h.indices[i]] >> 8); - unsigned int b0 = 0xff & g_colors[h.indices[i]]; + unsigned int r0 = 0xff & (g_colors[index] >> 16); + unsigned int g0 = 0xff & (g_colors[index] >> 8); + unsigned int b0 = 0xff & g_colors[index]; - float alpha = (255 - (0xff & (g_colors[h.indices[i]] >> 24)))/255.0f; + float alpha = (255 - (0xff & (g_colors[index] >> 24)))/255.0f; fr += r0*scale*cos_theta*alpha; fg += g0*scale*cos_theta*alpha; 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 |