diff options
author | Willem Jan Palenstijn <WillemJan.Palenstijn@uantwerpen.be> | 2014-04-16 11:13:22 +0000 |
---|---|---|
committer | wpalenst <WillemJan.Palenstijn@uantwerpen.be> | 2014-04-16 11:13:22 +0000 |
commit | 0c77eee16e2f4161c1ebc110b15ce6563d4a96c2 (patch) | |
tree | d824a763da5578391839c57057d5ecee26554873 /cuda/3d/fdk.cu | |
parent | 8b046691e7cf5ba603b690e5a820869f7aba0bb6 (diff) | |
download | astra-0c77eee16e2f4161c1ebc110b15ce6563d4a96c2.tar.gz astra-0c77eee16e2f4161c1ebc110b15ce6563d4a96c2.tar.bz2 astra-0c77eee16e2f4161c1ebc110b15ce6563d4a96c2.tar.xz astra-0c77eee16e2f4161c1ebc110b15ce6563d4a96c2.zip |
Add fan beam support to FBP_CUDA
Diffstat (limited to 'cuda/3d/fdk.cu')
-rw-r--r-- | cuda/3d/fdk.cu | 23 |
1 files changed, 20 insertions, 3 deletions
diff --git a/cuda/3d/fdk.cu b/cuda/3d/fdk.cu index 883187d..fd5a8a2 100644 --- a/cuda/3d/fdk.cu +++ b/cuda/3d/fdk.cu @@ -298,8 +298,7 @@ __global__ void devFDK_ParkerWeight(void* D_projData, unsigned int projPitch, un // Perform the FDK pre-weighting and filtering -bool FDK_Filter(cudaPitchedPtr D_projData, - cufftComplex * D_filter, +bool FDK_PreWeight(cudaPitchedPtr D_projData, float fSrcOrigin, float fDetOrigin, float fSrcZ, float fDetZ, float fDetUSize, float fDetVSize, bool bShortScan, @@ -335,13 +334,22 @@ bool FDK_Filter(cudaPitchedPtr D_projData, } cudaTextForceKernelsCompletion(); + return true; +} +bool FDK_Filter(cudaPitchedPtr D_projData, + cufftComplex * D_filter, + float fSrcOrigin, float fDetOrigin, + float fSrcZ, float fDetZ, + float fDetUSize, float fDetVSize, bool bShortScan, + const SDimensions3D& dims, const float* angles) +{ // The filtering is a regular ramp filter per detector line. int iPaddedDetCount = calcNextPowerOfTwo(2 * dims.iProjU); int iHalfFFTSize = calcFFTFourSize(iPaddedDetCount); - + int projPitch = D_projData.pitch/sizeof(float); // We process one sinogram at a time. @@ -390,11 +398,18 @@ bool FDK(cudaPitchedPtr D_volumeData, int iPaddedDetCount = calcNextPowerOfTwo(2 * dims.iProjU); int iHalfFFTSize = calcFFTFourSize(iPaddedDetCount); + ok = FDK_PreWeight(D_projData, fSrcOrigin, fDetOrigin, + fSrcZ, fDetZ, fDetUSize, fDetVSize, + bShortScan, dims, angles); + if (!ok) + return false; + cufftComplex *pHostFilter = new cufftComplex[dims.iProjAngles * iHalfFFTSize]; memset(pHostFilter, 0, sizeof(cufftComplex) * dims.iProjAngles * iHalfFFTSize); genFilter(FILTER_RAMLAK, 1.0f, dims.iProjAngles, pHostFilter, iPaddedDetCount, iHalfFFTSize); + allocateComplexOnDevice(dims.iProjAngles, iHalfFFTSize, &D_filter); uploadComplexArrayToDevice(dims.iProjAngles, iHalfFFTSize, pHostFilter, D_filter); @@ -403,6 +418,8 @@ bool FDK(cudaPitchedPtr D_volumeData, // Perform filtering + + ok = FDK_Filter(D_projData, D_filter, fSrcOrigin, fDetOrigin, fSrcZ, fDetZ, fDetUSize, fDetVSize, bShortScan, dims, angles); |