Remove angle limits in fan_fp, fan_bp

1 files changed, 24 insertions, 2 deletions

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