Remove angle limits in fan_fp, fan_bp

2 files changed, 69 insertions, 7 deletions

diff --git a/cuda/2d/fan_bp.cu b/cuda/2d/fan_bp.cu
index dd2a7b6..8983a9c 100644
--- a/cuda/2d/fan_bp.cu
+++ b/cuda/2d/fan_bp.cu
@@ -277,11 +277,10 @@ __global__ void devFanBP_FBPWeighted(float* D_volData, unsigned int volPitch, un
 }
 
 
-bool FanBP(float* D_volumeData, unsigned int volumePitch,
+bool FanBP_internal(float* D_volumeData, unsigned int volumePitch,
            float* D_projData, unsigned int projPitch,
            const SDimensions& dims, const SFanProjection* angles)
 {
-	// TODO: process angles block by block
 	assert(dims.iProjAngles <= g_MaxAngles);
 
 	bindProjDataTexture(D_projData, projPitch, dims.iProjDets, dims.iProjAngles);
@@ -324,11 +323,10 @@ bool FanBP(float* D_volumeData, unsigned int volumePitch,
 	return true;
 }
 
-bool FanBP_FBPWeighted(float* D_volumeData, unsigned int volumePitch,
+bool FanBP_FBPWeighted_internal(float* D_volumeData, unsigned int volumePitch,
            float* D_projData, unsigned int projPitch,
            const SDimensions& dims, const SFanProjection* angles)
 {
-	// TODO: process angles block by block
 	assert(dims.iProjAngles <= g_MaxAngles);
 
 	bindProjDataTexture(D_projData, projPitch, dims.iProjDets, dims.iProjAngles);
@@ -368,7 +366,6 @@ bool FanBP_FBPWeighted(float* D_volumeData, unsigned int volumePitch,
 	return true;
 }
 
-
 // D_projData is a pointer to one padded sinogram line
 bool FanBP_SART(float* D_volumeData, unsigned int volumePitch,
                 float* D_projData, unsigned int projPitch,
@@ -402,6 +399,49 @@ bool FanBP_SART(float* D_volumeData, unsigned int volumePitch,
 	return true;
 }
 
+bool FanBP(float* D_volumeData, unsigned int volumePitch,
+           float* D_projData, unsigned int projPitch,
+           const SDimensions& dims, const SFanProjection* angles)
+{
+	for (unsigned int iAngle = 0; iAngle < dims.iProjAngles; iAngle += g_MaxAngles) {
+		SDimensions subdims = dims;
+		unsigned int iEndAngle = iAngle + g_MaxAngles;
+		if (iEndAngle >= dims.iProjAngles)
+			iEndAngle = dims.iProjAngles;
+		subdims.iProjAngles = iEndAngle - iAngle;
+
+		bool ret;
+		ret = FanBP_internal(D_volumeData, volumePitch,
+		                  D_projData + iAngle * projPitch, projPitch,
+		                  subdims, angles + iAngle);
+		if (!ret)
+			return false;
+	}
+	return true;
+}
+
+bool FanBP_FBPWeighted(float* D_volumeData, unsigned int volumePitch,
+           float* D_projData, unsigned int projPitch,
+           const SDimensions& dims, const SFanProjection* angles)
+{
+	for (unsigned int iAngle = 0; iAngle < dims.iProjAngles; iAngle += g_MaxAngles) {
+		SDimensions subdims = dims;
+		unsigned int iEndAngle = iAngle + g_MaxAngles;
+		if (iEndAngle >= dims.iProjAngles)
+			iEndAngle = dims.iProjAngles;
+		subdims.iProjAngles = iEndAngle - iAngle;
+
+		bool ret;
+		ret = FanBP_FBPWeighted_internal(D_volumeData, volumePitch,
+		                  D_projData + iAngle * projPitch, projPitch,
+		                  subdims, angles + iAngle);
+
+		if (!ret)
+			return false;
+	}
+	return true;
+}
+
 
 }
 
diff --git a/cuda/2d/fan_fp.cu b/cuda/2d/fan_fp.cu
index b24029c..5f1ccdf 100644
--- a/cuda/2d/fan_fp.cu
+++ b/cuda/2d/fan_fp.cu
@@ -224,12 +224,11 @@ __global__ void FanFPvertical(float* D_projData, unsigned int projPitch, unsigne
 	projData[angle*projPitch+detector] += fVal;
 }
 
-bool FanFP(float* D_volumeData, unsigned int volumePitch,
+bool FanFP_internal(float* D_volumeData, unsigned int volumePitch,
            float* D_projData, unsigned int projPitch,
            const SDimensions& dims, const SFanProjection* angles,
            float outputScale)
 {
-	// TODO: load angles into constant memory in smaller blocks
 	assert(dims.iProjAngles <= g_MaxAngles);
 
 	cudaArray* D_dataArray;
@@ -286,6 +285,29 @@ bool FanFP(float* D_volumeData, unsigned int volumePitch,
 	return true;
 }
 
+bool FanFP(float* D_volumeData, unsigned int volumePitch,
+           float* D_projData, unsigned int projPitch,
+           const SDimensions& dims, const SFanProjection* angles,
+           float outputScale)
+{
+	for (unsigned int iAngle = 0; iAngle < dims.iProjAngles; iAngle += g_MaxAngles) {
+		SDimensions subdims = dims;
+		unsigned int iEndAngle = iAngle + g_MaxAngles;
+		if (iEndAngle >= dims.iProjAngles)
+			iEndAngle = dims.iProjAngles;
+		subdims.iProjAngles = iEndAngle - iAngle;
+
+		bool ret;
+		ret = FanFP_internal(D_volumeData, volumePitch,
+		                         D_projData + iAngle * projPitch, projPitch,
+		                         subdims, angles + iAngle,
+		                         outputScale);
+		if (!ret)
+			return false;
+	}
+	return true;
+}
+
 }
 
 #ifdef STANDALONE