summaryrefslogtreecommitdiff
path: root/src/kernel.cu
diff options
context:
space:
mode:
Diffstat (limited to 'src/kernel.cu')
-rw-r--r--src/kernel.cu125
1 files changed, 4 insertions, 121 deletions
diff --git a/src/kernel.cu b/src/kernel.cu
index aae6e95..44a6c03 100644
--- a/src/kernel.cu
+++ b/src/kernel.cu
@@ -12,9 +12,12 @@
#define STACK_SIZE 500
+#define RED_WAVELENGTH 685
+#define BLUE_WAVELENGTH 465
+#define GREEN_WAVELENGTH 545
+
enum
{
- DIFFUSE_HIT = -3,
DEBUG = -2,
INIT = -1,
NO_HIT,
@@ -152,12 +155,6 @@ __device__ void myAtomicAdd(float *addr, float data)
data = atomicExch(addr, data+atomicExch(addr, 0.0f));
}
-
-
-
-
-
-
__device__ int to_diffuse(curandState &rng, float3 position, float3 direction, float wavelength, float3 polarization, int start_node, int first_node, int max_steps)
{
int last_hit_triangle = -1;
@@ -357,22 +354,6 @@ __device__ int to_diffuse(curandState &rng, float3 position, float3 direction, f
} // to_diffuse
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
-
extern "C"
{
@@ -405,11 +386,6 @@ __global__ void rotate(int nthreads, float3 *points, float phi, float3 axis)
points[id] = rotate(points[id], phi, axis);
}
-#define RED_WAVELENGTH 685
-#define BLUE_WAVELENGTH 465
-#define GREEN_WAVELENGTH 545
-
-
__global__ void build_rgb_lookup(int nthreads, curandState *rng_states, float3 *positions, float3 *directions, int start_node, int first_node, float3 *rgb_lookup, int runs, int max_steps)
{
int id = blockIdx.x*blockDim.x + threadIdx.x;
@@ -508,94 +484,6 @@ __global__ void render(int nthreads, curandState *rng_states, float3 *positions,
} // render
-#if 0
-
-__global__ void get_triangle(int nthreads, float3 *positions, float3 *directions, int start_node, int first_node, int *triangles)
-{
- 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;
-
- triangles[id] = intersect_mesh(position, direction, start_node, first_node, distance);
-}
-
-__global__ void get_cos_theta(int nthreads, float3 *positions, float3 *directions, int start_node, int first_node, int *triangle, float *cos_thetas)
-{
- int id = blockIdx.x*blockDim.x + threadIdx.x;
-
- if (id >= nthreads)
- return;
-
- float3 position = positions[id];
- float3 direction = directions[id];
- direction /= norm(direction);
-
- uint4 triangle_data = triangles[triangle[id]];
-
- float3 v0 = vertices[triangle_data.x];
- float3 v1 = vertices[triangle_data.y];
- float3 v2 = vertices[triangle_data.z];
-
- float cos_theta = dot(normalize(cross(v1-v0, v2-v1)), -direction);
-
- if (cos_theta < 0.0f)
- cos_theta = dot(-normalize(cross(v1-v0, v2-v1)), -direction);
-
- cos_thetas[id] = cos_theta;
-}
-
-#endif
-
-#if 0
-
-__global__ void to_diffuse(int nthreads, curandState *rng_states, float3 *positions, float3 *directions, float *wavelengths, float3 *polarizations, int *last_hit_triangles, int start_node, int first_node, int max_steps)
-{
- int id = blockIdx.x*blockDim.x + threadIdx.x;
-
- if (id >= nthreads)
- return;
-
- curandState rng = rng_states[id];
- float3 poisiton = positions[id];
- float3 direction = directions[id];
- direction /= norm(direction);
- float3 polarization = polarizations[id];
- polarization /= norm(polarization);
- float wavelength = wavelengths[id];
- float last_hit_triangle = last_hit_triangles[id];
-
- int steps = 0;
- while (steps < max_steps)
- {
- steps++;
-
- float distance;
-
- last_hit_triangle = intersect_mesh(position, direction, start_node, first_node, distance, last_hit_triangle);
-
- if (last_hit_triangle == -1)
- break;
-
- uint4 triangle_data = triangles[last_hit_triangle];
-
- float3 v0 = vertices[triangle_data.x];
- float3 v1 = vertices[triangle_data.y];
- float3 v2 = vertices[triangle_data.z];
-
- int material_in_index = convert(0xFF & (triangle_data.w >> 24));
- int material_out_index = convert(0xFF & (triangle_data.w >> 16));
- int surface_index = convert(0xFF & triangle_data.w >> 8);
-
-
-#endif
-
/* 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
@@ -872,9 +760,4 @@ __global__ void propagate(int nthreads, curandState *rng_states, float3 *positio
} // propagate
-#if 0
-
-
-#endif
-
} // extern "c"