summaryrefslogtreecommitdiff
path: root/src/kernel.cu
diff options
context:
space:
mode:
authorStan Seibert <stan@mtrr.org>2011-09-03 09:21:36 -0400
committerStan Seibert <stan@mtrr.org>2011-09-03 09:21:36 -0400
commit38f05bf761490def1886016524f328528b08f549 (patch)
treee0ee6555ee8bdf02a9e0b832a33707bcee06a3fa /src/kernel.cu
parent48550062440c5b7f1479ecbe17fd4b024a90fca2 (diff)
parent707ca1b366f11032682cc864ca2848905e6b485c (diff)
downloadchroma-38f05bf761490def1886016524f328528b08f549.tar.gz
chroma-38f05bf761490def1886016524f328528b08f549.tar.bz2
chroma-38f05bf761490def1886016524f328528b08f549.zip
merge
Diffstat (limited to 'src/kernel.cu')
-rw-r--r--src/kernel.cu60
1 files changed, 25 insertions, 35 deletions
diff --git a/src/kernel.cu b/src/kernel.cu
index 307412e..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;
@@ -239,6 +204,28 @@ __global__ void process_image(int nthreads, float3 *image, int *pixels, int nima
} // process_image
+__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;
@@ -288,6 +275,9 @@ __global__ void ray_trace_alpha(int nthreads, float3 *positions, float3 *directi
float3 direction = directions[id];
direction /= norm(direction);
+ bool hit;
+ float distance;
+
pixels[id] = get_color_alpha(position, direction);
} // ray_trace