summaryrefslogtreecommitdiff
path: root/src
diff options
context:
space:
mode:
Diffstat (limited to 'src')
-rw-r--r--src/daq.cu8
1 files changed, 6 insertions, 2 deletions
diff --git a/src/daq.cu b/src/daq.cu
index 7c5e6a5..2b95560 100644
--- a/src/daq.cu
+++ b/src/daq.cu
@@ -34,7 +34,9 @@ __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_q,
+ unsigned int *channel_histories)
{
int id = threadIdx.x + blockDim.x * blockIdx.x;
@@ -48,7 +50,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 +58,8 @@ __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);
}
}