Chroma CUDA Photon handling

geometry from CUDA photon propagation, in photon.h:

584 __device__ int
585 propagate_at_surface(Photon &p, State &s, curandState &rng, Geometry *geometry,
586                      bool use_weights=false)
587 {
588     Surface *surface = geometry->surfaces[s.surface_index];
589
590     if (surface->model == SURFACE_COMPLEX)
591         return propagate_complex(p, s, rng, surface, use_weights);
592     else if (surface->model == SURFACE_WLS)
593         return propagate_at_wls(p, s, rng, surface, use_weights);
594     else {
595         // use default surface model: do a combination of specular and
596         // diffuse reflection, detection, and absorption based on relative
597         // probabilties
598
599         // since the surface properties are interpolated linearly, we are
600         // guaranteed that they still sum to 1.0.
601         float detect = interp_property(surface, p.wavelength, surface->detect);
602         float absorb = interp_property(surface, p.wavelength, surface->absorb);
603         float reflect_diffuse = interp_property(surface, p.wavelength, surface->reflect_diffuse);
604         float reflect_specular = interp_property(surface, p.wavelength, surface->reflect_specular);
605
606         float uniform_sample = curand_uniform(&rng);

simon:cuda blyth$ grep __shared__ *.*
bvh.cu:    __shared__ unsigned long long min_area[128];
bvh.cu:    __shared__ unsigned long long adjacent_area;
daq.cu:    __shared__ int photon_id;
daq.cu:    __shared__ int triangle_id;
daq.cu:    __shared__ int solid_id;
daq.cu:    __shared__ int channel_index;
daq.cu:    __shared__ unsigned int history;
daq.cu:    __shared__ float photon_time;
daq.cu:    __shared__ float weight;
mesh.h:    __shared__ Geometry sg;
pdf.cu:    __shared__ float distance_table[1000];
pdf.cu:    __shared__ unsigned int *work_queue;
pdf.cu:    __shared__ int queue_items;
pdf.cu:    __shared__ int channel_id;
pdf.cu:    __shared__ float channel_event_time;
pdf.cu:    __shared__ int distance_table_len;
pdf.cu:    __shared__ int offset;
propagate.cu:    __shared__ unsigned int counter;
propagate.cu:    __shared__ Geometry sg;
render.cu:    __shared__ Geometry sg;
simon:cuda blyth$

simon:cuda blyth$ grep sync *.*
bvh.cu:    __syncthreads();
bvh.cu:    __syncthreads();
bvh.cu:    __syncthreads();
daq.cu:    __syncthreads();
mesh.h:    __syncthreads();
pdf.cu:    __syncthreads();
pdf.cu:    __syncthreads();
pdf.cu:    __syncthreads();
pdf.cu:    __syncthreads();
propagate.cu:    __syncthreads();
propagate.cu:    __syncthreads();
propagate.cu:    __syncthreads();
render.cu:    __syncthreads();
simon:cuda blyth$