Add fan beam support to FBP_CUDA

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,