diff options
Diffstat (limited to 'src/propagate.cu')
| -rw-r--r-- | src/propagate.cu | 18 |
1 files changed, 14 insertions, 4 deletions
diff --git a/src/propagate.cu b/src/propagate.cu index ecc32c0..5e98c9f 100644 --- a/src/propagate.cu +++ b/src/propagate.cu @@ -5,6 +5,8 @@ #include "geometry.h" #include "photon.h" +#include "stdio.h" + extern "C" { @@ -45,7 +47,6 @@ photon_duplicate(int first_photon, int nthreads, } } - __global__ void propagate(int first_photon, int nthreads, unsigned int *input_queue, unsigned int *output_queue, curandState *rng_states, @@ -53,13 +54,22 @@ propagate(int first_photon, int nthreads, unsigned int *input_queue, float *wavelengths, float3 *polarizations, float *times, unsigned int *histories, int *last_hit_triangles, int max_steps, - Geometry *geometry) + Geometry *g) { + __shared__ Geometry sg; + + if (threadIdx.x == 0) + sg = *g; + + __syncthreads(); + int id = blockIdx.x*blockDim.x + threadIdx.x; if (id >= nthreads) return; + g = &sg; + curandState rng = rng_states[id]; int photon_id = input_queue[first_photon + id]; @@ -92,7 +102,7 @@ propagate(int first_photon, int nthreads, unsigned int *input_queue, break; } - fill_state(s, p, geometry); + fill_state(s, p, g); if (p.last_hit_triangle == -1) break; @@ -106,7 +116,7 @@ propagate(int first_photon, int nthreads, unsigned int *input_queue, continue; if (s.surface_index != -1) { - command = propagate_at_surface(p, s, rng, geometry); + command = propagate_at_surface(p, s, rng, g); if (command == BREAK) break; |