From 12f86554bafb66e6afc6193f181527ba0749de92 Mon Sep 17 00:00:00 2001
From: Willem Jan Palenstijn <wjp@usecode.org>
Date: Sat, 30 Mar 2019 21:21:16 +0100
Subject: Adjust SART to line integral scaling

---
 cuda/2d/fan_bp.cu | 9 ++++-----
 cuda/2d/par_bp.cu | 5 +++--
 cuda/2d/sart.cu   | 5 +++--
 3 files changed, 10 insertions(+), 9 deletions(-)

(limited to 'cuda/2d')

diff --git a/cuda/2d/fan_bp.cu b/cuda/2d/fan_bp.cu
index 2072cd4..0d21897 100644
--- a/cuda/2d/fan_bp.cu
+++ b/cuda/2d/fan_bp.cu
@@ -217,17 +217,16 @@ __global__ void devFanBP_SART(float* D_volData, unsigned int volPitch, const SDi
 	const float fDetSY = gC_DetSY[0];
 	const float fDetUX = gC_DetUX[0];
 	const float fDetUY = gC_DetUY[0];
-	const float fScale = gC_Scale[0];
+
+	// NB: The 'scale' constant in devBP is cancelled out by the SART weighting
 
 	const float fXD = fSrcX - fX;
 	const float fYD = fSrcY - fY;
 
 	const float fNum = fDetSY * fXD - fDetSX * fYD + fX*fSrcY - fY*fSrcX;
 	const float fDen = fDetUX * fYD - fDetUY * fXD;
-	const float fr = __fdividef(1.0f, fDen);
-
-	const float fT = fNum * fr;
-	const float fVal = tex2D(gT_FanProjTexture, fT, 0.5f) * fScale * fr;
+	const float fT = fNum / fDen;
+	const float fVal = tex2D(gT_FanProjTexture, fT, 0.5f);
 
 	volData[Y*volPitch+X] += fVal * fOutputScale;
 }
diff --git a/cuda/2d/par_bp.cu b/cuda/2d/par_bp.cu
index 21484b1..3bf78ee 100644
--- a/cuda/2d/par_bp.cu
+++ b/cuda/2d/par_bp.cu
@@ -176,7 +176,7 @@ __global__ void devBP_SART(float* D_volData, unsigned int volPitch, float offset
 	const float fT = fX * angle_cos - fY * angle_sin + offset;
 	const float fVal = tex2D(gT_projTexture, fT, 0.5f);
 
-	// NB: The 'scale' constant in devBP has been folded into fOutputScale by the caller here
+	// NB: The 'scale' constant in devBP is cancelled out by the SART weighting
 
 	D_volData[Y*volPitch+X] += fVal * fOutputScale;
 }
@@ -280,7 +280,8 @@ bool BP_SART(float* D_volumeData, unsigned int volumePitch,
 	float angle_scaled_cos = angles[angle].fRayY / d;
 	float angle_scaled_sin = -angles[angle].fRayX / d; // TODO: Check signs
 	float angle_offset = (angles[angle].fDetSY * angles[angle].fRayX - angles[angle].fDetSX * angles[angle].fRayY) / d;
-	fOutputScale *= sqrt(angles[angle].fRayX * angles[angle].fRayX + angles[angle].fRayY * angles[angle].fRayY) / abs(d);
+	// NB: The adjoint scaling factor from regular BP is cancelled out by the SART weighting
+	//fOutputScale *= sqrt(angles[angle].fRayX * angles[angle].fRayX + angles[angle].fRayY * angles[angle].fRayY) / abs(d);
 
 	dim3 dimBlock(g_blockSlices, g_blockSliceSize);
 	dim3 dimGrid((dims.iVolWidth+g_blockSlices-1)/g_blockSlices,
diff --git a/cuda/2d/sart.cu b/cuda/2d/sart.cu
index aff77a3..12ad6df 100644
--- a/cuda/2d/sart.cu
+++ b/cuda/2d/sart.cu
@@ -266,14 +266,15 @@ bool SART::callBP_SART(float* D_volumeData, unsigned int volumePitch,
                        float* D_projData, unsigned int projPitch,
                        unsigned int angle, float outputScale)
 {
+	// NB: No fProjectorScale here, as that it is cancelled out in the SART weighting
 	if (parProjs) {
 		assert(!fanProjs);
 		return BP_SART(D_volumeData, volumePitch, D_projData, projPitch,
-		               angle, dims, parProjs, outputScale * fProjectorScale);
+		               angle, dims, parProjs, outputScale);
 	} else {
 		assert(fanProjs);
 		return FanBP_SART(D_volumeData, volumePitch, D_projData, projPitch,
-		                  angle, dims, fanProjs, outputScale * fProjectorScale);
+		                  angle, dims, fanProjs, outputScale);
 	}
 
 }
-- 
cgit v1.2.3