diff options
author | Stan Seibert <stan@mtrr.org> | 2011-08-09 15:43:36 -0400 |
---|---|---|
committer | Stan Seibert <stan@mtrr.org> | 2011-08-09 15:43:36 -0400 |
commit | 06ece999c3866f2d19acfd0f23b2f62d02b50577 (patch) | |
tree | f9bb7d591c69b7848d764e5055b2d167382c6337 /src/daq.cu | |
parent | 135ad2ca5b5ade4c52ae98f8fc7545fcc88fb449 (diff) | |
download | chroma-06ece999c3866f2d19acfd0f23b2f62d02b50577.tar.gz chroma-06ece999c3866f2d19acfd0f23b2f62d02b50577.tar.bz2 chroma-06ece999c3866f2d19acfd0f23b2f62d02b50577.zip |
Store a photon history for each hit channel. If multiple photons hit the
same channel, their history bits are OR'ed together.
Diffstat (limited to 'src/daq.cu')
-rw-r--r-- | src/daq.cu | 6 |
1 files changed, 4 insertions, 2 deletions
@@ -34,7 +34,8 @@ __global__ void run_daq(curandState *s, unsigned int detection_state, int nphotons, float *photon_times, unsigned int *photon_histories, int *last_hit_triangles, int *solid_map, - int nsolids, unsigned int *earliest_time_int) + int nsolids, unsigned int *earliest_time_int, + unsigned int *channel_histories) { int id = threadIdx.x + blockDim.x * blockIdx.x; @@ -48,7 +49,7 @@ __global__ void run_daq(curandState *s, unsigned int detection_state, if (triangle_id > -1) { int solid_id = solid_map[triangle_id]; - int history = photon_histories[photon_id]; + unsigned int history = photon_histories[photon_id]; if (solid_id < nsolids && (history & detection_state)) { @@ -56,6 +57,7 @@ __global__ void run_daq(curandState *s, unsigned int detection_state, unsigned int time_int = float_to_sortable_int(time); atomicMin(earliest_time_int + solid_id, time_int); + atomicOr(channel_histories + solid_id, history); } } |