From 06ece999c3866f2d19acfd0f23b2f62d02b50577 Mon Sep 17 00:00:00 2001 From: Stan Seibert Date: Tue, 9 Aug 2011 15:43:36 -0400 Subject: Store a photon history for each hit channel. If multiple photons hit the same channel, their history bits are OR'ed together. --- src/daq.cu | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) (limited to 'src') 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); } } -- cgit From 5b478fe72e600e06cd7b2e8a05a600f30c44d5c0 Mon Sep 17 00:00:00 2001 From: Stan Seibert Date: Tue, 9 Aug 2011 15:53:35 -0400 Subject: Put number of detected photons into charge value for channel. --- src/daq.cu | 2 ++ 1 file changed, 2 insertions(+) (limited to 'src') diff --git a/src/daq.cu b/src/daq.cu index 9824d7d..2b95560 100644 --- a/src/daq.cu +++ b/src/daq.cu @@ -35,6 +35,7 @@ __global__ void run_daq(curandState *s, unsigned int detection_state, unsigned int *photon_histories, int *last_hit_triangles, int *solid_map, int nsolids, unsigned int *earliest_time_int, + unsigned int *channel_q, unsigned int *channel_histories) { @@ -57,6 +58,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); + atomicAdd(channel_q + solid_id, 1); atomicOr(channel_histories + solid_id, history); } -- cgit