diff options
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); |