1#ifndef __PARALLELPROJ_CUDA_PROJECTOR_KERNELS_H__
2#define __PARALLELPROJ_CUDA_PROJECTOR_KERNELS_H__
3
4/** @brief Calculate whether a ray and a cube intersect and the possible intersection points
5 *
6 * The ray is given by origin + t*rdir (vector notation)
7 * if the ray intersects the cube, the two t values t1 and t2
8 * for the intersection points are computed.
9 *
10 * This algorithm assumes that the IEEE floating point arithmetic standard 754 is followed
11 * which handles divions by 0 and -0 correctly.
12 *
13 * @param orig0 ... 0 cordinate of the ray origin
14 * @param orig1 ... 1 cordinate of the ray origin
15 * @param orig2 ... 2 cordinate of the ray origin
16 * @param bounds0_min ... 0 cordinate of the start of the cube bounding box
17 * @param bounds1_min ... 1 cordinate of the start of the cube bounding box
18 * @param bounds2_min ... 2 cordinate of the start of the cube bounding box
19 * @param bounds0_max ... 0 cordinate of the end of the cube bounding box
20 * @param bounds1_max ... 1 cordinate of the end of the cube bounding box
21 * @param bounds2_max ... 2 cordinate of the end of the cube bounding box
22 * @param rdir0 ... 0 cordinate of the ray directional vector
23 * @param rdir1 ... 1 cordinate of the ray directional vector
24 * @param rdir2 ... 2 cordinate of the ray directional vector
25 * @param t1 ... (output) ray parameter of 1st intersection point
26 * @param t2 ... (output) ray parameter of 2nd intersection point
27 *
28 * @return ... unsigned char (0 or 1) whether ray intersects cube
29 */
30extern "C" __device__ unsigned char ray_cube_intersection_cuda(float orig0,
31 float orig1,
32 float orig2,
33 float bounds0_min,
34 float bounds1_min,
35 float bounds2_min,
36 float bounds0_max,
37 float bounds1_max,
38 float bounds2_max,
39 float rdir0,
40 float rdir1,
41 float rdir2,
42 float* t1,
43 float* t2);
44
45
46
47/**
48 * @brief Calculate the TOF bins along an LOR to which a voxel contributes
49 *
50 * @param x_m0 0-component of center of LOR
51 * @param x_m1 1-component of center of LOR
52 * @param x_m2 2-component of center of LOR
53 * @param x_v0 0-component of voxel
54 * @param x_v1 1-component of voxel
55 * @param x_v2 2-component of voxel
56 * @param u0 0-component of unit vector that points from start to end of LOR
57 * @param u1 1-component of unit vector that points from start to end of LOR
58 * @param u2 2-component of unit vector that points from start to end of LOR
59 * @param tofbin_width width of the TOF bins in spatial units
60 * @param tofcenter_offset offset of the central tofbin from the midpoint of the LOR in spatial units
61 * @param sigma_tof TOF resolution in spatial coordinates
62 * @param n_sigmas number of sigmas considered to be relevant
63 * @param n_half n_tofbins // 2
64 * @param it1 (output) lower relevant tof bin
65 * @param it2 (output) upper relevant tof bin
66 */
67extern "C" __device__ void relevant_tof_bins_cuda(float x_m0,
68 float x_m1,
69 float x_m2,
70 float x_v0,
71 float x_v1,
72 float x_v2,
73 float u0,
74 float u1,
75 float u2,
76 float tofbin_width,
77 float tofcenter_offset,
78 float sigma_tof,
79 float n_sigmas,
80 int n_half,
81 int *it1,
82 int *it2);
83
84
85
86/** @brief 3D non-tof joseph back projector CUDA kernel
87 *
88 * @param xstart array of shape [3*nlors] with the coordinates of the start points of the LORs.
89 * The start coordinates of the n-th LOR are at xstart[n*3 + i] with i = 0,1,2.
90 * Units are the ones of voxsize.
91 * @param xend array of shape [3*nlors] with the coordinates of the end points of the LORs.
92 * The start coordinates of the n-th LOR are at xstart[n*3 + i] with i = 0,1,2.
93 * Units are the ones of voxsize.
94 * @param img array of shape [n0*n1*n2] for the back projection image (output).
95 * The pixel [i,j,k] ist stored at [n1*n2*i + n2*j + k].
96 * @param img_origin array [x0_0,x0_1,x0_2] of coordinates of the center of the [0,0,0] voxel
97 * @param voxsize array [vs0, vs1, vs2] of the voxel sizes
98 * @param p array of length nlors containg the values to be back projected
99 * @param nlors number of projections (length of p array)
100 * @param img_dim array with dimensions of image [n0,n1,n2]
101 */
102extern "C" __global__ void joseph3d_back_cuda_kernel(float *xstart,
103 float *xend,
104 float *img,
105 float *img_origin,
106 float *voxsize,
107 float *p,
108 long long nlors,
109 int *img_dim);
110
111
112
113/** @brief 3D non-tof joseph forward projector CUDA kernel
114 *
115 * @param xstart array of shape [3*nlors] with the coordinates of the start points of the LORs.
116 * The start coordinates of the n-th LOR are at xstart[n*3 + i] with i = 0,1,2.
117 * Units are the ones of voxsize.
118 * @param xend array of shape [3*nlors] with the coordinates of the end points of the LORs.
119 * The start coordinates of the n-th LOR are at xstart[n*3 + i] with i = 0,1,2.
120 * Units are the ones of voxsize.
121 * @param img array of shape [n0*n1*n2] containing the 3D image to be projected.
122 * The pixel [i,j,k] ist stored at [n1*n2*i + n2*j + k].
123 * @param img_origin array [x0_0,x0_1,x0_2] of coordinates of the center of the [0,0,0] voxel
124 * @param voxsize array [vs0, vs1, vs2] of the voxel sizes
125 * @param p array of length nlors (output) used to store the projections
126 * @param nlors number of projections (length of p array)
127 * @param img_dim array with dimensions of image [n0,n1,n2]
128 */
129extern "C" __global__ void joseph3d_fwd_cuda_kernel(float *xstart,
130 float *xend,
131 float *img,
132 float *img_origin,
133 float *voxsize,
134 float *p,
135 long long nlors,
136 int *img_dim);
137
138
139
140
141/** @brief 3D sinogram tof cuda joseph back projector kernel
142 *
143 * @param xstart array of shape [3*nlors] with the coordinates of the start points of the LORs.
144 * The start coordinates of the n-th LOR are at xstart[n*3 + i] with i = 0,1,2.
145 * Units are the ones of voxsize.
146 * @param xend array of shape [3*nlors] with the coordinates of the end points of the LORs.
147 * The start coordinates of the n-th LOR are at xstart[n*3 + i] with i = 0,1,2.
148 * Units are the ones of voxsize.
149 * @param img array of shape [n0*n1*n2] containing the 3D image used for back projection (output).
150 * The pixel [i,j,k] ist stored at [n1*n2*i + n2*j + k].
151 * @param img_origin array [x0_0,x0_1,x0_2] of coordinates of the center of the [0,0,0] voxel
152 * @param voxsize array [vs0, vs1, vs2] of the voxel sizes
153 * @param p array of length nlors with the values to be back projected
154 * @param nlors number of geometrical LORs
155 * @param img_dim array with dimensions of image [n0,n1,n2]
156 * @param n_tofbins number of TOF bins
157 * @param tofbin_width width of the TOF bins in spatial units (units of xstart and xend)
158 * @param sigma_tof array of length 1 or nlors (depending on lor_dependent_sigma_tof)
159 * with the TOF resolution (sigma) for each LOR in
160 * spatial units (units of xstart and xend)
161 * @param tofcenter_offset array of length 1 or nlors (depending on lor_dependent_tofcenter_offset)
162 * with the offset of the central TOF bin from the
163 * midpoint of each LOR in spatial units (units of xstart and xend).
164 * A positive value means a shift towards the end point of the LOR.
165 * @param n_sigmas number of sigmas to consider for calculation of TOF kernel
166 * @param lor_dependent_sigma_tof unsigned char 0 or 1
167 * 1 means that the TOF sigmas are LOR dependent
168 * any other value means that the first value in the sigma_tof
169 * array is used for all LORs
170 * @param lor_dependent_tofcenter_offset unsigned char 0 or 1
171 * 1 means that the TOF center offsets are LOR dependent
172 * any other value means that the first value in the
173 * tofcenter_offset array is used for all LORs
174 */
175extern "C" __global__ void joseph3d_back_tof_sino_cuda_kernel(float *xstart,
176 float *xend,
177 float *img,
178 float *img_origin,
179 float *voxsize,
180 float *p,
181 long long nlors,
182 int *img_dim,
183 short n_tofbins,
184 float tofbin_width,
185 float *sigma_tof,
186 float *tofcenter_offset,
187 float n_sigmas,
188 unsigned char lor_dependent_sigma_tof,
189 unsigned char lor_dependent_tofcenter_offset);
190
191
192
193
194/** @brief 3D sinogram tof joseph forward projector kernel
195 *
196 * @param xstart array of shape [3*nlors] with the coordinates of the start points of the LORs.
197 * The start coordinates of the n-th LOR are at xstart[n*3 + i] with i = 0,1,2.
198 * Units are the ones of voxsize.
199 * @param xend array of shape [3*nlors] with the coordinates of the end points of the LORs.
200 * The start coordinates of the n-th LOR are at xstart[n*3 + i] with i = 0,1,2.
201 * Units are the ones of voxsize.
202 * @param img array of shape [n0*n1*n2] containing the 3D image to be projected.
203 * The pixel [i,j,k] ist stored at [n1*n2*i + n2*j + k].
204 * @param img_origin array [x0_0,x0_1,x0_2] of coordinates of the center of the [0,0,0] voxel
205 * @param voxsize array [vs0, vs1, vs2] of the voxel sizes
206 * @param p array of length nlors*n_tofbins (output) used to store the projections
207 * @param nlors number of geomtrical LORs
208 * @param img_dim array with dimensions of image [n0,n1,n2]
209 * @param n_tofbins number of TOF bins
210 * @param tofbin_width width of the TOF bins in spatial units (units of xstart and xend)
211 * @param sigma_tof array of length 1 or nlors (depending on lor_dependent_sigma_tof)
212 * with the TOF resolution (sigma) for each LOR in
213 * spatial units (units of xstart and xend)
214 * @param tofcenter_offset array of length 1 or nlors (depending on lor_dependent_tofcenter_offset)
215 * with the offset of the central TOF bin from the
216 * midpoint of each LOR in spatial units (units of xstart and xend).
217 * A positive value means a shift towards the end point of the LOR.
218 * @param n_sigmas number of sigmas to consider for calculation of TOF kernel
219 * @param lor_dependent_sigma_tof unsigned char 0 or 1
220 * 1 means that the TOF sigmas are LOR dependent
221 * any other value means that the first value in the sigma_tof
222 * array is used for all LORs
223 * @param lor_dependent_tofcenter_offset unsigned char 0 or 1
224 * 1 means that the TOF center offsets are LOR dependent
225 * any other value means that the first value in the
226 * tofcenter_offset array is used for all LORs
227 */
228extern "C" __global__ void joseph3d_fwd_tof_sino_cuda_kernel(float *xstart,
229 float *xend,
230 float *img,
231 float *img_origin,
232 float *voxsize,
233 float *p,
234 long long nlors,
235 int *img_dim,
236 short n_tofbins,
237 float tofbin_width,
238 float *sigma_tof,
239 float *tofcenter_offset,
240 float n_sigmas,
241 unsigned char lor_dependent_sigma_tof,
242 unsigned char lor_dependent_tofcenter_offset);
243
244
245
246
247/** @brief 3D listmode tof cuda joseph back projector kernel
248 *
249 * @param xstart array of shape [3*nlors] with the coordinates of the start points of the LORs.
250 * The start coordinates of the n-th LOR are at xstart[n*3 + i] with i = 0,1,2.
251 * Units are the ones of voxsize.
252 * @param xend array of shape [3*nlors] with the coordinates of the end points of the LORs.
253 * The start coordinates of the n-th LOR are at xstart[n*3 + i] with i = 0,1,2.
254 * Units are the ones of voxsize.
255 * @param img array of shape [n0*n1*n2] containing the 3D image used for back projection (output).
256 * The pixel [i,j,k] ist stored at [n1*n2*i + n2*j + k].
257 * @param img_origin array [x0_0,x0_1,x0_2] of coordinates of the center of the [0,0,0] voxel
258 * @param voxsize array [vs0, vs1, vs2] of the voxel sizes
259 * @param p array of length nlors with the values to be back projected
260 * @param nlors number of geometrical LORs
261 * @param img_dim array with dimensions of image [n0,n1,n2]
262 * @param tofbin_width width of the TOF bins in spatial units (units of xstart and xend)
263 * @param sigma_tof array of length 1 or nlors (depending on lor_dependent_sigma_tof)
264 * with the TOF resolution (sigma) for each LOR in
265 * spatial units (units of xstart and xend)
266 * @param tofcenter_offset array of length 1 or nlors (depending on lor_dependent_tofcenter_offset)
267 * with the offset of the central TOF bin from the
268 * midpoint of each LOR in spatial units (units of xstart and xend).
269 * A positive value means a shift towards the end point of the LOR.
270 * @param n_sigmas number of sigmas to consider for calculation of TOF kernel
271 * @param tof_bin array containing the TOF bin of each event
272 * @param lor_dependent_sigma_tof unsigned char 0 or 1
273 * 1 means that the TOF sigmas are LOR dependent
274 * any other value means that the first value in the sigma_tof
275 * array is used for all LORs
276 * @param lor_dependent_tofcenter_offset unsigned char 0 or 1
277 * 1 means that the TOF center offsets are LOR dependent
278 * any other value means that the first value in the
279 * tofcenter_offset array is used for all LORs
280 */
281extern "C" __global__ void joseph3d_back_tof_lm_cuda_kernel(float *xstart,
282 float *xend,
283 float *img,
284 float *img_origin,
285 float *voxsize,
286 float *p,
287 long long nlors,
288 int *img_dim,
289 float tofbin_width,
290 float *sigma_tof,
291 float *tofcenter_offset,
292 float n_sigmas,
293 short *tof_bin,
294 unsigned char lor_dependent_sigma_tof,
295 unsigned char lor_dependent_tofcenter_offset);
296
297
298
299/** @brief 3D listmode tof joseph forward projector kernel
300 *
301 * @param xstart array of shape [3*nlors] with the coordinates of the start points of the LORs.
302 * The start coordinates of the n-th LOR are at xstart[n*3 + i] with i = 0,1,2.
303 * Units are the ones of voxsize.
304 * @param xend array of shape [3*nlors] with the coordinates of the end points of the LORs.
305 * The start coordinates of the n-th LOR are at xstart[n*3 + i] with i = 0,1,2.
306 * Units are the ones of voxsize.
307 * @param img array of shape [n0*n1*n2] containing the 3D image to be projected.
308 * The pixel [i,j,k] ist stored at [n1*n2*i + n2*j + k].
309 * @param img_origin array [x0_0,x0_1,x0_2] of coordinates of the center of the [0,0,0] voxel
310 * @param voxsize array [vs0, vs1, vs2] of the voxel sizes
311 * @param p array of length nlors (output) used to store the projections
312 * @param nlors number of geomtrical LORs
313 * @param img_dim array with dimensions of image [n0,n1,n2]
314 * @param tofbin_width width of the TOF bins in spatial units (units of xstart and xend)
315 * @param sigma_tof array of length 1 or nlors (depending on lor_dependent_sigma_tof)
316 * with the TOF resolution (sigma) for each LOR in
317 * spatial units (units of xstart and xend)
318 * @param tofcenter_offset array of length 1 or nlors (depending on lor_dependent_tofcenter_offset)
319 * with the offset of the central TOF bin from the
320 * midpoint of each LOR in spatial units (units of xstart and xend).
321 * A positive value means a shift towards the end point of the LOR.
322 * @param n_sigmas number of sigmas to consider for calculation of TOF kernel
323 * @param tof_bin array containing the TOF bin of each event
324 * @param lor_dependent_sigma_tof unsigned char 0 or 1
325 * 1 means that the TOF sigmas are LOR dependent
326 * any other value means that the first value in the sigma_tof
327 * array is used for all LORs
328 * @param lor_dependent_tofcenter_offset unsigned char 0 or 1
329 * 1 means that the TOF center offsets are LOR dependent
330 * any other value means that the first value in the
331 * tofcenter_offset array is used for all LORs
332 */
333extern "C" __global__ void joseph3d_fwd_tof_lm_cuda_kernel(float *xstart,
334 float *xend,
335 float *img,
336 float *img_origin,
337 float *voxsize,
338 float *p,
339 long long nlors,
340 int *img_dim,
341 float tofbin_width,
342 float *sigma_tof,
343 float *tofcenter_offset,
344 float n_sigmas,
345 short *tof_bin,
346 unsigned char lor_dependent_sigma_tof,
347 unsigned char lor_dependent_tofcenter_offset);
348#endif