summaryrefslogtreecommitdiff
path: root/src
diff options
context:
space:
mode:
authorAnthony LaTorre <tlatorre9@gmail.com>2011-07-27 15:09:12 -0400
committerAnthony LaTorre <tlatorre9@gmail.com>2011-07-27 15:09:12 -0400
commit368524c5007a1f6ce2ef7be8d5dc1156c41a9dc2 (patch)
tree62992d86956f0f7282bd69c0f160710ff0489b2a /src
parentb94392f8ed2ce7bf364d0549bcda8e632d56de05 (diff)
downloadchroma-368524c5007a1f6ce2ef7be8d5dc1156c41a9dc2.tar.gz
chroma-368524c5007a1f6ce2ef7be8d5dc1156c41a9dc2.tar.bz2
chroma-368524c5007a1f6ce2ef7be8d5dc1156c41a9dc2.zip
camera automatically increases render quality when it is not being moved
Diffstat (limited to 'src')
-rw-r--r--src/kernel.cu160
1 files changed, 78 insertions, 82 deletions
diff --git a/src/kernel.cu b/src/kernel.cu
index ffd2d88..a36ff91 100644
--- a/src/kernel.cu
+++ b/src/kernel.cu
@@ -87,26 +87,27 @@ __global__ void rotate(int nthreads, float3 *points, float phi, float3 axis)
points[id] = rotate(points[id], phi, axis);
}
-__global__ void build_rgb_lookup(int nthreads, curandState *rng_states, float3 *positions, float3 *directions, float3 *rgb_lookup1, float3 *rgb_lookup2, int runs, int max_steps)
+ __global__ void build_rgb_lookup(int nthreads, float3 *positions, float3 *directions, float3 *rgb_lookup1, float3 *rgb_lookup2, unsigned long long seed, int runs, int max_steps)
{
int id = blockIdx.x*blockDim.x + threadIdx.x;
if (id >= nthreads)
return;
- curandState rng = rng_states[id];
+ curandState rng;
+ curand_init(seed, id, 0, &rng);
- Photon seed;
- seed.position = positions[id];
- seed.direction = directions[id];
- seed.direction /= norm(seed.direction);
- seed.polarization = uniform_sphere(&rng);
- seed.time = 0.0f;
- seed.history = 0x0;
+ Photon p0;
+ p0.position = positions[id];
+ p0.direction = directions[id];
+ p0.direction /= norm(p0.direction);
+ p0.polarization = uniform_sphere(&rng);
+ p0.time = 0.0f;
+ p0.history = 0x0;
float distance;
- int hit_triangle = intersect_mesh(seed.position, seed.direction, distance);
+ int hit_triangle = intersect_mesh(p0.position, p0.direction, distance);
if (hit_triangle != id)
return;
@@ -119,17 +120,17 @@ __global__ void build_rgb_lookup(int nthreads, curandState *rng_states, float3 *
float3 v1 = g_vertices[triangle_data.y];
float3 v2 = g_vertices[triangle_data.z];
- float cos_theta = dot(normalize(cross(v1-v0, v2-v1)), -seed.direction);
+ float cos_theta = dot(normalize(cross(v1-v0, v2-v1)), -p0.direction);
if (cos_theta < 0.0f)
- cos_theta = dot(-normalize(cross(v1-v0, v2-v1)), -seed.direction);
+ cos_theta = dot(-normalize(cross(v1-v0, v2-v1)), -p0.direction);
Photon p;
State s;
for (int i=0; i < runs; i++)
{
- p = seed;
+ p = p0;
p.wavelength = RED_WAVELENGTH;
to_diffuse(p, s, rng, max_steps);
@@ -146,7 +147,7 @@ __global__ void build_rgb_lookup(int nthreads, curandState *rng_states, float3 *
}
}
- p = seed;
+ p = p0;
p.wavelength = GREEN_WAVELENGTH;
to_diffuse(p, s, rng, max_steps);
@@ -163,7 +164,7 @@ __global__ void build_rgb_lookup(int nthreads, curandState *rng_states, float3 *
}
}
- p = seed;
+ p = p0;
p.wavelength = BLUE_WAVELENGTH;
to_diffuse(p, s, rng, max_steps);
@@ -181,11 +182,9 @@ __global__ void build_rgb_lookup(int nthreads, curandState *rng_states, float3 *
}
}
- rng_states[id] = rng;
-
} // build_rgb_lookup
-__global__ void render(int nthreads, curandState *rng_states, float3 *positions, float3 *directions, float3 *rgb_lookup1, float3 *rgb_lookup2, int runs, int *pixels, int max_steps)
+__global__ void render(int nthreads, curandState *rng_states, float3 *positions, float3 *directions, float *red, float *green, float *blue, float3 *rgb_lookup1, float3 *rgb_lookup2, int max_steps)
{
int id = blockIdx.x*blockDim.x + threadIdx.x;
@@ -194,81 +193,78 @@ __global__ void render(int nthreads, curandState *rng_states, float3 *positions,
curandState rng = rng_states[id];
- Photon seed;
- seed.position = positions[id];
- seed.direction = directions[id];
- seed.direction /= norm(seed.direction);
- seed.polarization = uniform_sphere(&rng);
- seed.time = 0.0f;
- seed.history = 0x0;
-
- float3 rgb = make_float3(0.0, 0.0, 0.0);
+ Photon p0;
+ p0.position = positions[id];
+ p0.direction = directions[id];
+ p0.direction /= norm(p0.direction);
+ p0.polarization = uniform_sphere(&rng);
+ p0.time = 0.0f;
+ p0.history = 0x0;
- Photon p;
State s;
+ Photon p = p0;
+ p.wavelength = RED_WAVELENGTH;
+
+ to_diffuse(p, s, rng, max_steps);
- for (int i=0; i < runs; i++)
+ if (p.history & REFLECT_DIFFUSE)
{
- p = seed;
- p.wavelength = RED_WAVELENGTH;
-
- to_diffuse(p, s, rng, max_steps);
-
- if (p.history & REFLECT_DIFFUSE)
+ if (s.inside_to_outside)
{
- if (s.inside_to_outside)
- {
- rgb.x += rgb_lookup1[p.last_hit_triangle].x;
- }
- else
- {
- rgb.x += rgb_lookup2[p.last_hit_triangle].x;
- }
+ red[id] = rgb_lookup1[p.last_hit_triangle].x;
}
-
- p = seed;
- p.wavelength = GREEN_WAVELENGTH;
-
- to_diffuse(p, s, rng, max_steps);
-
- if (p.history & REFLECT_DIFFUSE)
+ else
{
- if (s.inside_to_outside)
- {
- rgb.y += rgb_lookup1[p.last_hit_triangle].y;
- }
- else
- {
- rgb.y += rgb_lookup2[p.last_hit_triangle].y;
- }
+ red[id] = rgb_lookup2[p.last_hit_triangle].x;
}
-
- p = seed;
- p.wavelength = BLUE_WAVELENGTH;
-
- to_diffuse(p, s, rng, max_steps);
-
- if (p.history & REFLECT_DIFFUSE)
+ }
+ else
+ {
+ red[id] = 0.0;
+ }
+
+ p = p0;
+ p.wavelength = GREEN_WAVELENGTH;
+
+ to_diffuse(p, s, rng, max_steps);
+
+ if (p.history & REFLECT_DIFFUSE)
+ {
+ if (s.inside_to_outside)
{
- if (s.inside_to_outside)
- {
- rgb.z += rgb_lookup1[p.last_hit_triangle].z;
- }
- else
- {
- rgb.z += rgb_lookup2[p.last_hit_triangle].z;
- }
+ green[id] = rgb_lookup1[p.last_hit_triangle].y;
+ }
+ else
+ {
+ green[id] = rgb_lookup2[p.last_hit_triangle].y;
}
}
-
- rgb /= runs;
-
- unsigned int r = floorf(rgb.x*255);
- unsigned int g = floorf(rgb.y*255);
- unsigned int b = floorf(rgb.z*255);
-
- pixels[id] = r << 16 | g << 8 | b;
-
+ else
+ {
+ green[id] = 0.0;
+ }
+
+ p = p0;
+ p.wavelength = BLUE_WAVELENGTH;
+
+ to_diffuse(p, s, rng, max_steps);
+
+ if (p.history & REFLECT_DIFFUSE)
+ {
+ if (s.inside_to_outside)
+ {
+ blue[id] = rgb_lookup1[p.last_hit_triangle].z;
+ }
+ else
+ {
+ blue[id] = rgb_lookup2[p.last_hit_triangle].z;
+ }
+ }
+ else
+ {
+ blue[id] = 0.0;
+ }
+
rng_states[id] = rng;
} // render