diff options
Diffstat (limited to 'cuda/3d')
| -rw-r--r-- | cuda/3d/fdk.cu | 23 | ||||
| -rw-r--r-- | cuda/3d/fdk.h | 8 |
2 files changed, 28 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); diff --git a/cuda/3d/fdk.h b/cuda/3d/fdk.h index 5443b19..c9123b9 100644 --- a/cuda/3d/fdk.h +++ b/cuda/3d/fdk.h @@ -29,8 +29,16 @@ $Id$ #ifndef _CUDA_FDK_H #define _CUDA_FDK_H +#include "dims3d.h" + namespace astraCUDA3d { +bool FDK_PreWeight(cudaPitchedPtr D_projData, + float fSrcOrigin, float fDetOrigin, + float fSrcZ, float fDetZ, + float fDetUSize, float fDetVSize, bool bShortScan, + const SDimensions3D& dims, const float* angles); + bool FDK(cudaPitchedPtr D_volumeData, cudaPitchedPtr D_projData, float fSrcOrigin, float fDetOrigin, |