diff options
author | Anthony LaTorre <tlatorre9@gmail.com> | 2011-06-20 16:33:59 -0400 |
---|---|---|
committer | Anthony LaTorre <tlatorre9@gmail.com> | 2011-06-20 16:33:59 -0400 |
commit | 5b6ddaadfbcea436dfdc1e736f7da7763438dc45 (patch) | |
tree | 27d3f4d3903d2cd00449026acd23ad6de1419cd5 /src | |
parent | a149f96a766c4d8d63919535cc468c539036165e (diff) | |
download | chroma-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.cu | 29 |
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); |