summaryrefslogtreecommitdiff
path: root/src/daq.cu
diff options
context:
space:
mode:
Diffstat (limited to 'src/daq.cu')
-rw-r--r--src/daq.cu26
1 files changed, 9 insertions, 17 deletions
diff --git a/src/daq.cu b/src/daq.cu
index a7f5120..07e2d5e 100644
--- a/src/daq.cu
+++ b/src/daq.cu
@@ -16,7 +16,8 @@ __device__ float sortable_int_to_float(unsigned int i)
//return __int_as_float(i ^ mask);
}
-extern "C" {
+extern "C"
+{
__global__ void reset_earliest_time_int(float maxtime, int ntime_ints, unsigned int *time_ints)
{
@@ -27,8 +28,10 @@ __global__ void reset_earliest_time_int(float maxtime, int ntime_ints, unsigned
}
}
-__global__ void run_daq(curandState *s, int detection_state, float time_rms,
- int nphotons, float *photon_times, int *photon_states,
+__global__ void run_daq(curandState *s, unsigned int detection_state,
+ float time_rms,
+ int nphotons, float *photon_times,
+ unsigned int *photon_histories,
int *last_hit_triangles, int *solid_map,
int nsolids, unsigned int *earliest_time_int)
{
@@ -39,37 +42,26 @@ __global__ void run_daq(curandState *s, int detection_state, float time_rms,
{
curandState rng = s[id];
-
-
int triangle_id = last_hit_triangles[id];
if (triangle_id > -1)
{
int solid_id = solid_map[triangle_id];
- int state = photon_states[id];
+ int history = photon_histories[id];
float time = photon_times[id] + curand_normal(&rng) * time_rms;
unsigned int time_int = float_to_sortable_int(time);
-
-
- if (solid_id < nsolids && state == detection_state) {
+ if (solid_id < nsolids && (history & detection_state))
+ {
atomicMin(earliest_time_int + solid_id, time_int);
}
-
}
-
-
s[id] = rng;
-
-
}
-
-
-
}
__global__ void convert_sortable_int_to_float(int n,