diff options
Diffstat (limited to 'src')
| -rw-r--r-- | src/alpha.h | 99 | ||||
| -rw-r--r-- | src/kernel.cu | 16 | ||||
| -rw-r--r-- | src/sorting.h | 103 |
3 files changed, 158 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 diff --git a/src/sorting.h b/src/sorting.h new file mode 100644 index 0000000..ec16bbe --- /dev/null +++ b/src/sorting.h @@ -0,0 +1,103 @@ +template <class T> +__device__ void swap(T &a, T &b) +{ + T tmp = a; + a = b; + b = tmp; +} + +template <class T> +__device__ void reverse(int n, T *a) +{ + for (int i=0; i < n/2; i++) + swap(a[i],a[n-1-i]); +} + +template <class T> +__device__ void piksrt(int n, T *arr) +{ + int i,j; + T a; + + for (j=1; j < n; j++) + { + a = arr[j]; + i = j-1; + while (i >= 0 && arr[i] > a) + { + arr[i+1] = arr[i]; + i--; + } + arr[i+1] = a; + } +} + +template <class T, class U> +__device__ void piksrt2(int n, T *arr, U *brr) +{ + int i,j; + T a; + U b; + + for (j=1; j < n; j++) + { + a = arr[j]; + b = brr[j]; + i = j-1; + while (i >= 0 && arr[i] > a) + { + arr[i+1] = arr[i]; + brr[i+1] = brr[i]; + i--; + } + arr[i+1] = a; + brr[i+1] = b; + } +} + +template <class T> +__device__ unsigned long searchsorted(unsigned long n, T *arr, const T &x) +{ + unsigned long ju,jm,jl; + int ascnd; + + jl = 0; + ju = n; + + ascnd = (arr[n-1] > arr[0]); + + while (ju-jl > 1) + { + jm = (ju+jl) >> 1; + + if (x >= arr[jm] == ascnd) + jl = jm; + else + ju = jm; + } + + if (x <= arr[0]) + return 0; + else if (x == arr[n-1]) + return n-1; + else + return ju; +} + +template <class T> +__device__ void insert(unsigned long n, T *arr, unsigned long i, const T &x) +{ + unsigned long j; + for (j=n-1; j > i; j--) + arr[j] = arr[j-1]; + arr[i] = x; +} + +template <class T> +__device__ void add_sorted(unsigned long n, T *arr, const T &x) +{ + unsigned long i = searchsorted(n, arr, x); + + if (i < n) + insert(n, arr, i, x); +} |