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$