summaryrefslogtreecommitdiff
path: root/src/kernel.cu
diff options
context:
space:
mode:
Diffstat (limited to 'src/kernel.cu')
-rw-r--r--src/kernel.cu49
1 files changed, 28 insertions, 21 deletions
diff --git a/src/kernel.cu b/src/kernel.cu
index 2cd9347..711ae80 100644
--- a/src/kernel.cu
+++ b/src/kernel.cu
@@ -18,7 +18,7 @@ __device__ void fAtomicAdd(float *addr, float data)
data = atomicExch(addr, data+atomicExch(addr, 0.0f));
}
-__device__ void to_diffuse(Photon &p, State &s, const int &max_steps)
+__device__ void to_diffuse(Photon &p, State &s, curandState &rng, const int &max_steps)
{
int steps = 0;
while (steps < max_steps)
@@ -27,12 +27,12 @@ __device__ void to_diffuse(Photon &p, State &s, const int &max_steps)
int command;
- command = fill_state(s, p);
+ command = fill_state(s, p, rng);
if (command == BREAK)
break;
- command = propagate_to_boundary(p, s);
+ command = propagate_to_boundary(p, s, rng);
if (command == BREAK)
break;
@@ -42,7 +42,7 @@ __device__ void to_diffuse(Photon &p, State &s, const int &max_steps)
if (s.surface_index != -1)
{
- command = propagate_at_surface(p, s);
+ command = propagate_at_surface(p, s, rng);
if (p.history & REFLECT_DIFFUSE)
break;
@@ -54,7 +54,7 @@ __device__ void to_diffuse(Photon &p, State &s, const int &max_steps)
continue;
}
- propagate_at_boundary(p, s);
+ propagate_at_boundary(p, s, rng);
} // while (steps < max_steps)
@@ -93,12 +93,13 @@ __global__ void build_rgb_lookup(int nthreads, curandState *rng_states, float3 *
if (id >= nthreads)
return;
+ curandState rng = rng_states[id];
+
Photon seed;
- seed.rng = rng_states[id];
seed.position = positions[id];
seed.direction = directions[id];
seed.direction /= norm(seed.direction);
- seed.polarization = uniform_sphere(&seed.rng);
+ seed.polarization = uniform_sphere(&rng);
seed.time = 0.0f;
seed.history = 0x0;
@@ -130,7 +131,7 @@ __global__ void build_rgb_lookup(int nthreads, curandState *rng_states, float3 *
p = seed;
p.wavelength = RED_WAVELENGTH;
- to_diffuse(p, s, max_steps);
+ to_diffuse(p, s, rng, max_steps);
if (p.history & REFLECT_DIFFUSE)
{
@@ -147,7 +148,7 @@ __global__ void build_rgb_lookup(int nthreads, curandState *rng_states, float3 *
p = seed;
p.wavelength = BLUE_WAVELENGTH;
- to_diffuse(p, s, max_steps);
+ to_diffuse(p, s, rng, max_steps);
if (p.history & REFLECT_DIFFUSE)
{
@@ -164,7 +165,7 @@ __global__ void build_rgb_lookup(int nthreads, curandState *rng_states, float3 *
p = seed;
p.wavelength = GREEN_WAVELENGTH;
- to_diffuse(p, s, max_steps);
+ to_diffuse(p, s, rng, max_steps);
if (p.history & REFLECT_DIFFUSE)
{
@@ -179,6 +180,8 @@ __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)
@@ -188,12 +191,13 @@ __global__ void render(int nthreads, curandState *rng_states, float3 *positions,
if (id >= nthreads)
return;
+ curandState rng = rng_states[id];
+
Photon seed;
- seed.rng = rng_states[id];
seed.position = positions[id];
seed.direction = directions[id];
seed.direction /= norm(seed.direction);
- seed.polarization = uniform_sphere(&seed.rng);
+ seed.polarization = uniform_sphere(&rng);
seed.time = 0.0f;
seed.history = 0x0;
@@ -207,7 +211,7 @@ __global__ void render(int nthreads, curandState *rng_states, float3 *positions,
p = seed;
p.wavelength = RED_WAVELENGTH;
- to_diffuse(p, s, max_steps);
+ to_diffuse(p, s, rng, max_steps);
if (p.history & REFLECT_DIFFUSE)
{
@@ -224,7 +228,7 @@ __global__ void render(int nthreads, curandState *rng_states, float3 *positions,
p = seed;
p.wavelength = BLUE_WAVELENGTH;
- to_diffuse(p, s, max_steps);
+ to_diffuse(p, s, rng, max_steps);
if (p.history & REFLECT_DIFFUSE)
{
@@ -241,7 +245,7 @@ __global__ void render(int nthreads, curandState *rng_states, float3 *positions,
p = seed;
p.wavelength = GREEN_WAVELENGTH;
- to_diffuse(p, s, max_steps);
+ to_diffuse(p, s, rng, max_steps);
if (p.history & REFLECT_DIFFUSE)
{
@@ -264,6 +268,8 @@ __global__ void render(int nthreads, curandState *rng_states, float3 *positions,
pixels[id] = r << 16 | g << 8 | b;
+ rng_states[id] = rng;
+
} // render
/* Trace the rays starting at `positions` traveling in the direction
@@ -311,8 +317,9 @@ __global__ void propagate(int nthreads, curandState *rng_states, float3 *positio
if (id >= nthreads)
return;
+ curandState rng = rng_states[id];
+
Photon p;
- p.rng = rng_states[id];
p.position = positions[id];
p.direction = directions[id];
p.direction /= norm(p.direction);
@@ -335,12 +342,12 @@ __global__ void propagate(int nthreads, curandState *rng_states, float3 *positio
int command;
- command = fill_state(s, p);
+ command = fill_state(s, p, rng);
if (command == BREAK)
break;
- command = propagate_to_boundary(p, s);
+ command = propagate_to_boundary(p, s, rng);
if (command == BREAK)
break;
@@ -350,7 +357,7 @@ __global__ void propagate(int nthreads, curandState *rng_states, float3 *positio
if (s.surface_index != -1)
{
- command = propagate_at_surface(p, s);
+ command = propagate_at_surface(p, s, rng);
if (command == BREAK)
break;
@@ -359,11 +366,11 @@ __global__ void propagate(int nthreads, curandState *rng_states, float3 *positio
continue;
}
- propagate_at_boundary(p, s);
+ propagate_at_boundary(p, s, rng);
} // while (steps < max_steps)
- rng_states[id] = p.rng;
+ rng_states[id] = rng;
positions[id] = p.position;
directions[id] = p.direction;
polarizations[id] = p.polarization;
generated by cgit (git 2.34.1) at 2025-11-04 05:49:06 +0000