|
RISA
|
|
RISA
|
Namespaces | |
| detail | |
Classes | |
| class | Attenuation |
| This stage computes the attenuation coefficients in the fan beam sinogram. More... | |
| class | Backprojection |
| This stage back projects a parallel beam sinogram and returns the reconstructed image. More... | |
| class | D2H |
| This stage transfer a data element from device to host. More... | |
| class | DetectorInterpolation |
| This stage interpolates the defect detectors in the raw data sinogram. More... | |
| class | Fan2Para |
| This stage performs the fan to parallel beam rebinning. More... | |
| class | Filter |
| This stage filters the projections in the parallel beam sinogram with a precomputed filter function. More... | |
| class | H2D |
| This stage transfers the input data element from device to host. More... | |
| struct | hashTable |
| collects all precomputed hash table values More... | |
| class | Masking |
| This stage multiplies a precomputed mask with the reconstructed image. More... | |
| struct | parameters |
| collects all parameters that are needed in the fan to parallel beam interpolation kernel More... | |
| class | Reordering |
| This stage restructures the unordered input data received from the detector modules. More... | |
| class | Template |
| This stage transfer a data element from device to host. More... | |
Functions | |
| __global__ void | computeAttenuation (const unsigned short *__restrict__ sinogram_in, const float *__restrict__ mask, float *__restrict__ sinogram_out, const float *__restrict__ avgReference, const float *__restrict__ avgDark, const float temp, const int numberOfDetectors, const int numberOfProjections, const int planeId) |
| CUDA kernel to compute the attenuation coefficients. More... | |
| __global__ void | backProjectLinear (const float *const __restrict__ sinogram, float *__restrict__ image, const int numberOfPixels, const int numberOfProjections, const int numberOfDetectors) |
| This function performs the back projection operation with linear interpolation. More... | |
| __global__ void | backProjectTex (cudaTextureObject_t tex, float *__restrict__ image, const int numberOfPixels, const int numberOfProjections, const int numberOfDetectors) |
| This function performs the back projection with nearest neighbor interpolation using texture memory. More... | |
| __global__ void | backProjectNearest (const float *const __restrict__ sinogram, float *__restrict__ image, const int numberOfPixels, const int numberOfProjections, const int numberOfDetectors) |
| This function performs the back projection with nearest neighbor interpolation. More... | |
| template<typename T > | |
| auto | ellipse_kreis_uwe (T alpha, T DX, T DZ, T SourceRingDiam) -> T |
| __global__ void | interpolation (int k, const float *__restrict__ SinFan_data, float *__restrict__ SinPar_data, const float *__restrict__ Gamma, const float *__restrict__ Teta, const float *__restrict__ alpha_kreis, const float *__restrict__ s, const int *__restrict__ teta_nach_ray_1, const int *__restrict__ teta_nach_ray_2, const int *__restrict__ teta_vor_ray_1, const int *__restrict__ teta_vor_ray_2, const int *__restrict__ gamma_vor_ray_1, const int *__restrict__ gamma_vor_ray_2, const int *__restrict__ gamma_nach_ray_1, const int *__restrict__ gamma_nach_ray_2, const float *__restrict__ teta_ziel_ray_1, const float *__restrict__ teta_ziel_ray_2, const float *__restrict__ gamma_ziel_ray_1, const float *__restrict__ gamma_ziel_ray_2, const int *__restrict__ ray_1, const int *__restrict__ ray_2, const parameters *__restrict__ params) |
| template<typename T > | |
| __global__ void | setValue (T *data, T value, int size) |
| __global__ void | filterSL (int x, int y, cufftComplex *data) |
| __global__ void | filterRamp (int x, int y, cufftComplex *data) |
| __global__ void | filterHamming (int x, int y, cufftComplex *data) |
| __global__ void | filterHanning (int x, int y, cufftComplex *data) |
| template<typename T > | |
| auto | sheppLogan (const T w, const T d) -> T |
| computes the value of the Shepp-Logan-filter function More... | |
| template<typename T > | |
| auto | cosine (const T w, const T d) -> T |
| computes the value of the Cosine-filter function More... | |
| template<typename T > | |
| auto | hamming (const T w, const T d) -> T |
| computes the value of the Hamming-filter function More... | |
| template<typename T > | |
| auto | hanning (const T w, const T d) -> T |
| computes the value of the Hanning-filter function More... | |
| __global__ void | applyFilter (const int x, const int y, cufftComplex *data, const float *const __restrict__ filter) |
| __global__ void | mask (float *__restrict__ img, const float value, const int numberOfPixels) |
| This CUDA kernel multiplies the mask and the reconstructed image. More... | |
| __global__ void | reorder (const unsigned short *__restrict__ unorderedSino, unsigned short *__restrict__ orderedSino, const int *__restrict__ hashTable, const int numberOfProjections, const int numberOfDetectors) |
Variables | |
| __constant__ float | sinLookup [2048] |
| __constant__ float | cosLookup [2048] |
| __constant__ float | normalizationFactor [1] |
| __constant__ float | scale [1] |
| __constant__ float | imageCenter [1] |
| __global__ void risa::cuda::applyFilter | ( | const int | x, |
| const int | y, | ||
| cufftComplex * | data, | ||
| const float *const __restrict__ | filter | ||
| ) |
The variable i represents the detector index in the parallel beam sinogram, the variable j represents the projection in the parallel beam sinogram.
| [in] | x | the number of detectors in the parallel beam sinogram |
| [in] | y | the number of projections in the parallel beam sinogram |
| [in,out] | data | the inverse transformed parallel ray sinogram |
| [in] | filter | pointer to the precomputed filter function |
| __global__ void risa::cuda::backProjectLinear | ( | const float *const __restrict__ | sinogram, |
| float *__restrict__ | image, | ||
| const int | numberOfPixels, | ||
| const int | numberOfProjections, | ||
| const int | numberOfDetectors | ||
| ) |
This function performs the back projection operation with linear interpolation.
With a pixel-driven back projection approach, this CUDA kernel spans number of pixels times number of pixels, in the reconstruction grid, threads. Starting at the pixel center, each thread follows the ray path and computes the intersection with the detector pixels. The intersection does not line up, and thus, a linear interpolation is performed.
| [in] | sinogram | linearized sinogram data. Each projection is stored linearly after each other |
| [out] | image | the reconstruction grid, in which the reconstructed image is stored |
| [in] | numberOfPixels | the number of pixels in the reconstruction grid in one dimension |
| [in] | numberOfProjections | the number of projections in the parallel beam sinogram over 180 degrees |
| [in] | numberOfDetectors | the number of detectors in the parallel beam sinogram |
Definition at line 238 of file Backprojection.cu.
| __global__ void risa::cuda::backProjectNearest | ( | const float *const __restrict__ | sinogram, |
| float *__restrict__ | image, | ||
| const int | numberOfPixels, | ||
| const int | numberOfProjections, | ||
| const int | numberOfDetectors | ||
| ) |
This function performs the back projection with nearest neighbor interpolation.
With a pixel-driven back projection approach, this CUDA kernel spans number of pixels times number of pixels, in the reconstruction grid, threads. Starting at the pixel center, each thread follows the ray path and computes the intersection with the detector pixels. The intersection does not line up, and thus, a nearest neigbor interpolation is performed.
| [in] | sinogram | linearized sinogram data. Each projection is stored linearly after each other |
| [out] | image | the reconstruction grid, in which the reconstructed image is stored |
| [in] | numberOfPixels | the number of pixels in the reconstruction grid in one dimension |
| [in] | numberOfProjections | the number of projections in the parallel beam sinogram over 180 degrees |
| [in] | numberOfDetectors | the number of detectors in the parallel beam sinogram |
Definition at line 305 of file Backprojection.cu.
| __global__ void risa::cuda::backProjectTex | ( | cudaTextureObject_t | tex, |
| float *__restrict__ | image, | ||
| const int | numberOfPixels, | ||
| const int | numberOfProjections, | ||
| const int | numberOfDetectors | ||
| ) |
This function performs the back projection with nearest neighbor interpolation using texture memory.
With a pixel-driven back projection approach, this CUDA kernel spans number of pixels times number of pixels, in the reconstruction grid, threads. Starting at the pixel center, each thread follows the ray path and computes the intersection with the detector pixels. The intersection does not line up, and thus, a nearest neigbor interpolation using te texture memory is performed. This back projection kernel is the fastest one.
| [in] | tex | linearized sinogram data stored in texture memory. Each projection is stored linearly after each other |
| [out] | image | the reconstruction grid, in which the reconstructed image is stored |
| [in] | numberOfPixels | the number of pixels in the reconstruction grid in one dimension |
| [in] | numberOfProjections | the number of projections in the parallel beam sinogram over 180 degrees |
| [in] | numberOfDetectors | the number of detectors in the parallel beam sinogram |
Definition at line 273 of file Backprojection.cu.
| __global__ void risa::cuda::computeAttenuation | ( | const unsigned short *__restrict__ | sinogram_in, |
| const float *__restrict__ | mask, | ||
| float *__restrict__ | sinogram_out, | ||
| const float *__restrict__ | avgReference, | ||
| const float *__restrict__ | avgDark, | ||
| const float | temp, | ||
| const int | numberOfDetectors, | ||
| const int | numberOfProjections, | ||
| const int | planeId | ||
| ) |
CUDA kernel to compute the attenuation coefficients.
This CUDA kernel computes the attenuation coefficient for the fan to parallel beam sinogram. Furthermore, it multiplies the resulting values with a precomputed mask to hide unrelevant areas, that are previously known by the geometry of the measurement system.
| [in] | sinogram_in | the pointer to the raw data sinogram of size numberOfDetectors*numberOfProjections |
| [in] | mask | the pointer to the mask values, that is multiplied with the attenuation coefficient |
| [out] | sinogram_out | pointer to the fan to parallel beam sinogram |
| [in] | avgReference | pointer to the averaged reference measurement on device |
| [in] | avgDark | pointer to the averaged dark measurement on device |
| [in] | temp | |
| [in] | numberOfDetectors | the number of detectors in the fan beam sinogram |
| [in] | numberOfProjections | the number of projections in the fan beam sinogram |
| [in] | planeId | the id of the sinogram's plane |
Definition at line 416 of file Attenuation.cu.
|
inline |
computes the value of the Cosine-filter function
| [in] | w | the coordinate at the frequency axis |
| [in] | d | the cutoff-fraction |
Definition at line 113 of file cuda_kernels_filter.h.
| auto risa::cuda::ellipse_kreis_uwe | ( | T | alpha, |
| T | DX, | ||
| T | DZ, | ||
| T | SourceRingDiam | ||
| ) | -> T |
Definition at line 39 of file cuda_kernels_fan2para.h.
| __global__ void risa::cuda::filterHamming | ( | int | x, |
| int | y, | ||
| cufftComplex * | data | ||
| ) |
Definition at line 60 of file cuda_kernels_filter.h.
| __global__ void risa::cuda::filterHanning | ( | int | x, |
| int | y, | ||
| cufftComplex * | data | ||
| ) |
Definition at line 76 of file cuda_kernels_filter.h.
| __global__ void risa::cuda::filterRamp | ( | int | x, |
| int | y, | ||
| cufftComplex * | data | ||
| ) |
Definition at line 49 of file cuda_kernels_filter.h.
| __global__ void risa::cuda::filterSL | ( | int | x, |
| int | y, | ||
| cufftComplex * | data | ||
| ) |
Definition at line 33 of file cuda_kernels_filter.h.
|
inline |
computes the value of the Hamming-filter function
| [in] | w | the coordinate at the frequency axis |
| [in] | d | the cutoff-fraction |
Definition at line 126 of file cuda_kernels_filter.h.
|
inline |
computes the value of the Hanning-filter function
| [in] | w | the coordinate at the frequency axis |
| [in] | d | the cutoff-fraction |
Definition at line 139 of file cuda_kernels_filter.h.
| __global__ void risa::cuda::interpolation | ( | int | k, |
| const float *__restrict__ | SinFan_data, | ||
| float *__restrict__ | SinPar_data, | ||
| const float *__restrict__ | Gamma, | ||
| const float *__restrict__ | Teta, | ||
| const float *__restrict__ | alpha_kreis, | ||
| const float *__restrict__ | s, | ||
| const int *__restrict__ | teta_nach_ray_1, | ||
| const int *__restrict__ | teta_nach_ray_2, | ||
| const int *__restrict__ | teta_vor_ray_1, | ||
| const int *__restrict__ | teta_vor_ray_2, | ||
| const int *__restrict__ | gamma_vor_ray_1, | ||
| const int *__restrict__ | gamma_vor_ray_2, | ||
| const int *__restrict__ | gamma_nach_ray_1, | ||
| const int *__restrict__ | gamma_nach_ray_2, | ||
| const float *__restrict__ | teta_ziel_ray_1, | ||
| const float *__restrict__ | teta_ziel_ray_2, | ||
| const float *__restrict__ | gamma_ziel_ray_1, | ||
| const float *__restrict__ | gamma_ziel_ray_2, | ||
| const int *__restrict__ | ray_1, | ||
| const int *__restrict__ | ray_2, | ||
| const parameters *__restrict__ | params | ||
| ) |
Definition at line 64 of file cuda_kernels_fan2para.h.
| __global__ void risa::cuda::mask | ( | float *__restrict__ | img, |
| const float | value, | ||
| const int | numberOfPixels | ||
| ) |
This CUDA kernel multiplies the mask and the reconstructed image.
| [in,out] | img | the reconstructed image, that is multiplied with the mask in-place |
| [in] | value | the value, the pixels shall be replaced with |
| [in] | numberOfPixels | the number of pixels in the reconstruction grid in one dimension |
Definition at line 157 of file Masking.cu.
| __global__ void risa::cuda::reorder | ( | const unsigned short *__restrict__ | unorderedSino, |
| unsigned short *__restrict__ | orderedSino, | ||
| const int *__restrict__ | hashTable, | ||
| const int | numberOfProjections, | ||
| const int | numberOfDetectors | ||
| ) |
Definition at line 195 of file Reordering.cu.
| __global__ void risa::cuda::setValue | ( | T * | data, |
| T | value, | ||
| int | size | ||
| ) |
Definition at line 180 of file cuda_kernels_fan2para.h.
|
inline |
computes the value of the Shepp-Logan-filter function
| [in] | w | the coordinate at the frequency axis |
| [in] | d | the cutoff-fraction |
Definition at line 100 of file cuda_kernels_filter.h.
| __constant__ float risa::cuda::cosLookup[2048] |
Definition at line 41 of file Backprojection.cu.
| __constant__ float risa::cuda::imageCenter[1] |
Definition at line 44 of file Backprojection.cu.
| __constant__ float risa::cuda::normalizationFactor[1] |
Definition at line 42 of file Backprojection.cu.
| __constant__ float risa::cuda::scale[1] |
Definition at line 43 of file Backprojection.cu.
| __constant__ float risa::cuda::sinLookup[2048] |
Definition at line 40 of file Backprojection.cu.
1.8.11 
