Add fan beam support to FBP_CUDA

4 files changed, 131 insertions, 4 deletions

diff --git a/cuda/2d/astra.cu b/cuda/2d/astra.cu
index 1c2e623..e317f6c 100644
--- a/cuda/2d/astra.cu
+++ b/cuda/2d/astra.cu
@@ -33,6 +33,7 @@ $Id$
 #include "par_fp.h"
 #include "fan_fp.h"
 #include "par_bp.h"
+#include "fan_bp.h"
 #include "arith.h"
 #include "astra.h"
 
@@ -43,6 +44,9 @@ $Id$
 
 #include "../../include/astra/Logger.h"
 
+// For fan beam FBP weighting
+#include "../3d/fdk.h"
+
 using namespace astraCUDA;
 using namespace std;
 
@@ -61,8 +65,14 @@ public:
 	float* angles;
 	float* TOffsets;
 
+	float fOriginSourceDistance;
+	float fOriginDetectorDistance;
+
 	float fPixelSize;
 
+	bool bFanBeam;
+	bool bShortScan;
+
 	bool initialized;
 	bool setStartReconstruction;
 
@@ -152,9 +162,33 @@ bool AstraFBP::setProjectionGeometry(unsigned int iProjAngles,
 	pData->angles = new float[iProjAngles];
 	memcpy(pData->angles, pfAngles, iProjAngles * sizeof(pfAngles[0]));
 
+	pData->bFanBeam = false;
+
+	return true;
+}
+
+bool AstraFBP::setFanGeometry(unsigned int iProjAngles,
+                              unsigned int iProjDets,
+                              const float* pfAngles,
+                              float fOriginSourceDistance,
+                              float fOriginDetectorDistance,
+                              float fDetSize,
+                              bool bShortScan)
+{
+	// Slightly abusing setProjectionGeometry for this...
+	if (!setProjectionGeometry(iProjAngles, iProjDets, pfAngles, fDetSize))
+		return false;
+
+	pData->fOriginSourceDistance = fOriginSourceDistance;
+	pData->fOriginDetectorDistance = fOriginDetectorDistance;
+
+	pData->bFanBeam = true;
+	pData->bShortScan = bShortScan;
+
 	return true;
 }
 
+
 bool AstraFBP::setPixelSuperSampling(unsigned int iPixelSuperSampling)
 {
 	if (pData->initialized)
@@ -274,6 +308,34 @@ bool AstraFBP::run()
 
 	bool ok = false;
 
+	if (pData->bFanBeam) {
+		// Call FDK_PreWeight to handle fan beam geometry. We treat
+		// this as a cone beam setup of a single slice:
+
+		// TODO: TOffsets...
+
+		// We create a fake cudaPitchedPtr
+		cudaPitchedPtr tmp;
+		tmp.ptr = pData->D_sinoData;
+		tmp.pitch = pData->sinoPitch * sizeof(float);
+		tmp.xsize = pData->dims.iProjDets;
+		tmp.ysize = pData->dims.iProjAngles;
+		// and a fake Dimensions3D
+		astraCUDA3d::SDimensions3D dims3d;
+		dims3d.iVolX = pData->dims.iVolWidth;
+		dims3d.iVolY = pData->dims.iVolHeight;
+		dims3d.iVolZ = 1;
+		dims3d.iProjAngles = pData->dims.iProjAngles;
+		dims3d.iProjU = pData->dims.iProjDets;
+		dims3d.iProjV = 1;
+		dims3d.iRaysPerDetDim = dims3d.iRaysPerVoxelDim = 1;
+
+		astraCUDA3d::FDK_PreWeight(tmp, pData->fOriginSourceDistance,
+		              pData->fOriginDetectorDistance, 0.0f, 0.0f,
+		              pData->dims.fDetScale, 1.0f, // TODO: Are these correct?
+		              pData->bShortScan, dims3d, pData->angles);
+	}
+
 	if (pData->m_pDevFilter) {
 
 		int iFFTRealDetCount = calcNextPowerOfTwo(2 * pData->dims.iProjDets);
@@ -293,7 +355,34 @@ bool AstraFBP::run()
 
 	}
 
-	ok = BP(pData->D_volumeData, pData->volumePitch, pData->D_sinoData, pData->sinoPitch, pData->dims, pData->angles, pData->TOffsets);
+	if (pData->bFanBeam) {
+		// TODO: TOffsets?
+		// TODO: Remove this code duplication with CudaReconstructionAlgorithm
+		SFanProjection* projs;
+		projs = new SFanProjection[pData->dims.iProjAngles];
+
+		float fSrcX0 = 0.0f;
+		float fSrcY0 = -pData->fOriginSourceDistance / pData->fPixelSize;
+		float fDetUX0 = pData->dims.fDetScale;
+		float fDetUY0 = 0.0f;
+		float fDetSX0 = pData->dims.iProjDets * fDetUX0 / -2.0f;
+		float fDetSY0 = pData->fOriginDetectorDistance / pData->fPixelSize;
+
+#define ROTATE0(name,i,alpha) do { projs[i].f##name##X = f##name##X0 * cos(alpha) - f##name##Y0 * sin(alpha); projs[i].f##name##Y = f##name##X0 * sin(alpha) + f##name##Y0 * cos(alpha); } while(0)
+		for (unsigned int i = 0; i < pData->dims.iProjAngles; ++i) {
+			ROTATE0(Src, i, pData->angles[i]);
+			ROTATE0(DetS, i, pData->angles[i]);
+			ROTATE0(DetU, i, pData->angles[i]);
+		}
+
+#undef ROTATE0
+		ok = FanBP_FBPWeighted(pData->D_volumeData, pData->volumePitch, pData->D_sinoData, pData->sinoPitch, pData->dims, projs);
+
+		delete[] projs;
+
+	} else {
+		ok = BP(pData->D_volumeData, pData->volumePitch, pData->D_sinoData, pData->sinoPitch, pData->dims, pData->angles, pData->TOffsets);
+	}
 	if(!ok)
 	{
 		return false;
diff --git a/cuda/2d/astra.h b/cuda/2d/astra.h
index 9e58301..42b15e8 100644
--- a/cuda/2d/astra.h
+++ b/cuda/2d/astra.h
@@ -68,6 +68,19 @@ public:
 	                           unsigned int iProjDets,
 	                           const float *pfAngles,
 	                           float fDetSize = 1.0f);
+	// Set the projection angles and number of detector pixels per angle.
+	// pfAngles must be a float array of length iProjAngles.
+	// fDetSize indicates the size of a detector pixel compared to a
+	// volume pixel edge.
+	//
+	// pfAngles will only be read from during this call.
+	bool setFanGeometry(unsigned int iProjAngles,
+	                    unsigned int iProjDets,
+	                    const float *pfAngles,
+	                    float fOriginSourceDistance,
+	                    float fOriginDetectorDistance,
+	                    float fDetSize = 1.0f,
+	                    bool bShortScan = false);
 
 	// Set linear supersampling factor for the BP.
 	// (The number of rays is the square of this)
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,