diff options
author | Anthony LaTorre <tlatorre9@gmail.com> | 2011-06-24 10:30:32 -0400 |
---|---|---|
committer | Anthony LaTorre <tlatorre9@gmail.com> | 2011-06-24 10:30:32 -0400 |
commit | d63d0d0cd84519b7c0bd7f6b576bd0f24ff22dc2 (patch) | |
tree | ea28a98ad484e7dce963b2f330107bb4fc593372 /src | |
parent | ac9f568017e9a7cf8da04acdacc342b4db6576fa (diff) | |
download | chroma-d63d0d0cd84519b7c0bd7f6b576bd0f24ff22dc2.tar.gz chroma-d63d0d0cd84519b7c0bd7f6b576bd0f24ff22dc2.tar.bz2 chroma-d63d0d0cd84519b7c0bd7f6b576bd0f24ff22dc2.zip |
speedup mesh intersection by skipping directly to the child node of nodes with only one child.
Diffstat (limited to 'src')
-rw-r--r-- | src/kernel.cu | 46 |
1 files changed, 24 insertions, 22 deletions
diff --git a/src/kernel.cu b/src/kernel.cu index 55fde70..f257c6c 100644 --- a/src/kernel.cu +++ b/src/kernel.cu @@ -81,34 +81,37 @@ __device__ int intersect_mesh(const float3 &origin, const float3& direction, con do { - int index = tex1Dfetch(node_map, *node); - int length = tex1Dfetch(node_length, *node); + int first_child = tex1Dfetch(node_map, *node); + int child_count = tex1Dfetch(node_length, *node); + while (*node >= first_node && child_count == 1) + { + *node = first_child; + first_child = tex1Dfetch(node_map, *node); + child_count = tex1Dfetch(node_length, *node); + } + if (*node >= first_node) { - for (i=0; i < length; i++) + for (i=0; i < child_count; i++) { - if (intersect_node(origin, direction, index+i)) + if (intersect_node(origin, direction, first_child+i)) { - //*node++ = index+i; - *node = index+i; + *node = first_child+i; node++; } } - if (node == head) - break; - node--; } else // node is a leaf { - for (i=0; i < length; i++) + for (i=0; i < child_count; i++) { - if (last_hit_triangle == index+i) + if (last_hit_triangle == first_child+i) continue; - uint4 triangle_data = triangles[index+i]; + uint4 triangle_data = triangles[first_child+i]; float3 v0 = vertices[triangle_data.x]; float3 v1 = vertices[triangle_data.y]; @@ -118,14 +121,14 @@ __device__ int intersect_mesh(const float3 &origin, const float3& direction, con { if (triangle_index == -1) { - triangle_index = index + i; + triangle_index = first_child + i; min_distance = distance; continue; } if (distance < min_distance) { - triangle_index = index + i; + triangle_index = first_child + i; min_distance = distance; } } @@ -245,7 +248,9 @@ __global__ void propagate(int nthreads, curandState *rng_states, float3 *positio curandState rng = rng_states[id]; float3 position = positions[id]; float3 direction = directions[id]; + direction /= norm(direction); float3 polarization = polarizations[id]; + polarization /= norm(polarization); float wavelength = wavelengths[id]; float time = times[id]; int last_hit_triangle = last_hit_triangles[id]; @@ -255,9 +260,6 @@ __global__ void propagate(int nthreads, curandState *rng_states, float3 *positio { steps++; - direction /= norm(direction); - polarization /= norm(polarization); - float distance; last_hit_triangle = intersect_mesh(position, direction, start_node, first_node, distance, last_hit_triangle); @@ -372,10 +374,8 @@ __global__ void propagate(int nthreads, curandState *rng_states, float3 *positio float reflection_diffuse = interp_property(wavelength, surface.reflection_diffuse); float reflection_specular = interp_property(wavelength, surface.reflection_specular); - float sum = absorption + reflection_diffuse + reflection_specular; - absorption /= sum; - reflection_diffuse /= sum; - reflection_specular /= sum; + // since the surface properties are interpolated linearly, + // we are guaranteed that they still sum to one. float uniform_sample = curand_uniform(&rng); @@ -402,9 +402,10 @@ __global__ void propagate(int nthreads, curandState *rng_states, float3 *positio else { // specularly reflect - direction = rotate(surface_normal, incident_angle, p); + // polarization ? + continue; } @@ -446,6 +447,7 @@ __global__ void propagate(int nthreads, curandState *rng_states, float3 *positio } polarization = cross(p, direction); + polarization /= norm(polarization); continue; } // photon polarization parallel to surface |