Enable creation of charge and time PDF on GPU

1 files changed, 27 insertions, 0 deletions

diff --git a/src/daq.cu b/src/daq.cu
index 2b95560..a34fbbe 100644
--- a/src/daq.cu
+++ b/src/daq.cu
@@ -80,4 +80,31 @@ __global__ void convert_sortable_int_to_float(int n,
 		float_output[id] = sortable_int_to_float(sortable_ints[id]);
 }
 
+__global__ void bin_hits(int nchannels,
+			 unsigned int *channel_q, float *channel_time,
+			 unsigned int *hitcount, 
+			 int tbins, float tmin, float tmax,
+			 int qbins, float qmin, float qmax,
+			 unsigned int *pdf)
+{
+	int id = threadIdx.x + blockDim.x * blockIdx.x;
+
+	if (id >= nchannels)
+	  return;
+
+	unsigned int q = channel_q[id];
+	float t = channel_time[id];
+	
+	if (t < 1e8 && t >= tmin && t < tmax && q >= qmin && q < qmax) {
+		hitcount[id] += 1;
+		
+		int tbin = (t - tmin) / (tmax - tmin) * tbins;
+		int qbin = (q - qmin) / (qmax - qmin) * qbins;
+	    
+		// row major order (channel, t, q)
+		int bin = id * (tbins * qbins) + tbin * qbins + qbin;
+		pdf[bin] += 1;
+	}
+}
+
 } // extern "C"