diff options
| author | Anthony LaTorre <tlatorre9@gmail.com> | 2011-08-09 16:39:59 -0400 |
|---|---|---|
| committer | Anthony LaTorre <tlatorre9@gmail.com> | 2011-08-09 16:39:59 -0400 |
| commit | 6870f1f5695c7cca3ed8e0a1cf559245991c80b5 (patch) | |
| tree | 184fbbbb7cf2d44607dd192b34c0c5455d8cdcf5 /src | |
| parent | 4f3e0b7709bb64ffc24f2a003509d5f480848239 (diff) | |
| parent | 5b478fe72e600e06cd7b2e8a05a600f30c44d5c0 (diff) | |
| download | chroma-6870f1f5695c7cca3ed8e0a1cf559245991c80b5.tar.gz chroma-6870f1f5695c7cca3ed8e0a1cf559245991c80b5.tar.bz2 chroma-6870f1f5695c7cca3ed8e0a1cf559245991c80b5.zip | |
merge heads
Diffstat (limited to 'src')
| -rw-r--r-- | src/daq.cu | 8 |
1 files changed, 6 insertions, 2 deletions
@@ -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); } } |