summaryrefslogtreecommitdiff
path: root/src
diff options
context:
space:
mode:
authorAnthony LaTorre <tlatorre9@gmail.com>2011-06-20 16:33:59 -0400
committerAnthony LaTorre <tlatorre9@gmail.com>2011-06-20 16:33:59 -0400
commit5b6ddaadfbcea436dfdc1e736f7da7763438dc45 (patch)
tree27d3f4d3903d2cd00449026acd23ad6de1419cd5 /src
parenta149f96a766c4d8d63919535cc468c539036165e (diff)
downloadchroma-5b6ddaadfbcea436dfdc1e736f7da7763438dc45.tar.gz
chroma-5b6ddaadfbcea436dfdc1e736f7da7763438dc45.tar.bz2
chroma-5b6ddaadfbcea436dfdc1e736f7da7763438dc45.zip
pack material and surface indices into the fourth byte of the triangle array on the GPU. you can now take a screenshot of an image rendered with view.py() by pressing the f12 key.
Diffstat (limited to 'src')
-rw-r--r--src/kernel.cu29
1 files changed, 20 insertions, 9 deletions
diff --git a/src/kernel.cu b/src/kernel.cu
index 6c3ef1b..405f06c 100644
--- a/src/kernel.cu
+++ b/src/kernel.cu
@@ -25,11 +25,6 @@ enum
texture<float4, 1, cudaReadModeElementType> vertices;
__device__ uint4 *triangles;
-/* material/surface index lookup for each triangle */
-texture<int, 1, cudaReadModeElementType> material1_lookup;
-texture<int, 1, cudaReadModeElementType> material2_lookup;
-texture<int, 1, cudaReadModeElementType> surface_lookup;
-
/* lower/upper bounds for the bounding box associated with each node/leaf */
texture<float4, 1, cudaReadModeElementType> upper_bounds;
texture<float4, 1, cudaReadModeElementType> lower_bounds;
@@ -43,6 +38,14 @@ __device__ float3 make_float3(const float4 &a)
return make_float3(a.x, a.y, a.z);
}
+__device__ int convert(int c)
+{
+ if (c & 0x80)
+ return (0xFFFFFF00 | c);
+ else
+ return c;
+}
+
/* Test the intersection between a ray starting at `origin` traveling in the
direction `direction` and the bounding box around node `i`. If the ray
intersects the bounding box return true, else return false. */
@@ -204,7 +207,7 @@ __global__ void ray_trace(int nthreads, float3 *positions, float3 *directions, i
if (triangle_index == -1)
{
- pixels[id] = 0;
+ pixels[id] = 0x000000;
}
else
{
@@ -263,9 +266,9 @@ __global__ void propagate(int nthreads, float3 *positions, float3 *directions, f
float3 v1 = make_float3(tex1Dfetch(vertices, triangle_data.y));
float3 v2 = make_float3(tex1Dfetch(vertices, triangle_data.z));
- int material_in_index = tex1Dfetch(material1_lookup, last_hit_triangle);
- int material_out_index = tex1Dfetch(material2_lookup, last_hit_triangle);
- int surface_index = tex1Dfetch(surface_lookup, last_hit_triangle);
+ 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));
float3 surface_normal = cross(v1-v0,v2-v1);
surface_normal /= norm(surface_normal);
@@ -408,9 +411,13 @@ __global__ void propagate(int nthreads, float3 *positions, float3 *directions, f
float reflection_probability = reflection_coefficient*reflection_coefficient;
if ((curand_uniform(&rng) < reflection_probability) || isnan(refracted_angle))
+ {
direction = rotate(surface_normal, incident_angle, p);
+ }
else
+ {
direction = rotate(surface_normal, PI-refracted_angle, p);
+ }
polarization = p;
@@ -422,9 +429,13 @@ __global__ void propagate(int nthreads, float3 *positions, float3 *directions, f
float reflection_probability = reflection_coefficient*reflection_coefficient;
if ((curand_uniform(&rng) < reflection_probability) || isnan(refracted_angle))
+ {
direction = rotate(surface_normal, incident_angle, p);
+ }
else
+ {
direction = rotate(surface_normal, PI-refracted_angle, p);
+ }
polarization = cross(p, direction);