summaryrefslogtreecommitdiff
path: root/src
diff options
context:
space:
mode:
authorAnthony LaTorre <tlatorre9@gmail.com>2011-06-24 10:30:32 -0400
committerAnthony LaTorre <tlatorre9@gmail.com>2011-06-24 10:30:32 -0400
commitd63d0d0cd84519b7c0bd7f6b576bd0f24ff22dc2 (patch)
treeea28a98ad484e7dce963b2f330107bb4fc593372 /src
parentac9f568017e9a7cf8da04acdacc342b4db6576fa (diff)
downloadchroma-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.cu46
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