summaryrefslogtreecommitdiff
path: root/src/daq.cu
diff options
context:
space:
mode:
authorStan Seibert <stan@mtrr.org>2011-08-09 15:43:36 -0400
committerStan Seibert <stan@mtrr.org>2011-08-09 15:43:36 -0400
commit06ece999c3866f2d19acfd0f23b2f62d02b50577 (patch)
treef9bb7d591c69b7848d764e5055b2d167382c6337 /src/daq.cu
parent135ad2ca5b5ade4c52ae98f8fc7545fcc88fb449 (diff)
downloadchroma-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.cu6
1 files changed, 4 insertions, 2 deletions
diff --git a/src/daq.cu b/src/daq.cu
index 7c5e6a5..9824d7d 100644
--- a/src/daq.cu
+++ b/src/daq.cu
@@ -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);
}
}