From 35957b6ef72749cdc520ded67a0eb8cdfd7ea655 Mon Sep 17 00:00:00 2001
From: Willem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>
Date: Fri, 29 Mar 2019 15:03:57 +0100
Subject: Adjust linear/cuda kernels to line integral scaling

---
 cuda/2d/astra.cu  |  3 ++-
 cuda/2d/par_fp.cu | 10 ++++------
 2 files changed, 6 insertions(+), 7 deletions(-)

(limited to 'cuda/2d')

diff --git a/cuda/2d/astra.cu b/cuda/2d/astra.cu
index ec03517..7ff1c95 100644
--- a/cuda/2d/astra.cu
+++ b/cuda/2d/astra.cu
@@ -302,7 +302,8 @@ static bool convertAstraGeometry_internal(const CVolumeGeometry2D* pVolGeom,
 		pProjs[i].scale(factor);
 	}
 	// CHECKME: Check factor
-	fOutputScale *= pVolGeom->getPixelLengthX() * pVolGeom->getPixelLengthY();
+	// NB: Only valid for square pixels
+	fOutputScale *= pVolGeom->getPixelLengthX();
 
 	return true;
 }
diff --git a/cuda/2d/par_fp.cu b/cuda/2d/par_fp.cu
index 0835301..e03381c 100644
--- a/cuda/2d/par_fp.cu
+++ b/cuda/2d/par_fp.cu
@@ -115,10 +115,9 @@ __global__ void FPhorizontal_simple(float* D_projData, unsigned int projPitch, u
 	float fSliceStep = cos_theta / sin_theta;
 	float fDistCorr;
 	if (sin_theta > 0.0f)
-		fDistCorr = -fDetStep;
+		fDistCorr = outputScale / sin_theta;
 	else
-		fDistCorr = fDetStep;
-	fDistCorr *= outputScale;
+		fDistCorr = -outputScale / sin_theta;
 
 	float fVal = 0.0f;
 	// project detector on slice
@@ -193,10 +192,9 @@ __global__ void FPvertical_simple(float* D_projData, unsigned int projPitch, uns
 	float fSliceStep = sin_theta / cos_theta;
 	float fDistCorr;
 	if (cos_theta < 0.0f)
-		fDistCorr = -fDetStep;
+		fDistCorr = -outputScale / cos_theta; 
 	else
-		fDistCorr = fDetStep;
-	fDistCorr *= outputScale;
+		fDistCorr = outputScale / cos_theta;
 
 	float fVal = 0.0f;
 	float fP = (detector - 0.5f*dims.iProjDets + 0.5f - gC_angle_offset[angle]) * fDetStep + (startSlice - 0.5f*dims.iVolHeight + 0.5f) * fSliceStep + 0.5f*dims.iVolWidth - 0.5f + 0.5f;
-- 
cgit v1.2.3