1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
|
// -*-c++-*-
#include <curand_kernel.h>
__device__ unsigned int float_to_sortable_int(float f)
{
return __float_as_int(f);
//int i = __float_as_int(f);
//unsigned int mask = -(int)(i >> 31) | 0x80000000;
//return i ^ mask;
}
__device__ float sortable_int_to_float(unsigned int i)
{
return __int_as_float(i);
//unsigned int mask = ((i >> 31) - 1) | 0x80000000;
//return __int_as_float(i ^ mask);
}
__device__ curandState daq_rng_states[100000];
extern "C" {
__global__ void init_daq_rng(int nthreads,
unsigned long long seed, unsigned long long offset)
{
int id = blockIdx.x*blockDim.x + threadIdx.x;
if (id >= nthreads)
return;
curand_init(seed, id, offset, daq_rng_states+id);
}
__global__ void reset_earliest_time_int(float maxtime,
int ntime_ints, unsigned int *time_ints)
{
int id = threadIdx.x + blockDim.x * blockIdx.x;
if (id < ntime_ints) {
unsigned int maxtime_int = float_to_sortable_int(maxtime);
time_ints[id] = maxtime_int;
}
}
__global__ void run_daq(int detection_state, float time_rms,
int nphotons, float *photon_times, int *photon_states,
int *last_hit_triangles, int *solid_map,
int nsolids, unsigned int *earliest_time_int)
{
int id = threadIdx.x + blockDim.x * blockIdx.x;
curandState_t rng = daq_rng_states[id];
if (id < nphotons) {
int triangle_id = last_hit_triangles[id];
if (triangle_id > -1) {
int solid_id = solid_map[triangle_id];
int state = photon_states[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)
atomicMin(earliest_time_int + solid_id, time_int);
}
}
daq_rng_states[id] = rng;
}
__global__ void convert_sortable_int_to_float(int n,
unsigned int *sortable_ints,
float *float_output)
{
int id = threadIdx.x + blockDim.x * blockIdx.x;
if (id < n)
float_output[id] = sortable_int_to_float(sortable_ints[id]);
}
} // extern "C"
|
