summaryrefslogtreecommitdiffstats
path: root/cuda/2d
diff options
context:
space:
mode:
authorWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2019-03-29 15:03:57 +0100
committerWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2019-09-25 14:10:08 +0200
commit35957b6ef72749cdc520ded67a0eb8cdfd7ea655 (patch)
tree8f9f35c04c2731fe4e20139171ede3dac9d4d894 /cuda/2d
parent0f4ceb4c7f3f63fddf8fbf44c59fcd8f415e3847 (diff)
downloadastra-35957b6ef72749cdc520ded67a0eb8cdfd7ea655.tar.gz
astra-35957b6ef72749cdc520ded67a0eb8cdfd7ea655.tar.bz2
astra-35957b6ef72749cdc520ded67a0eb8cdfd7ea655.tar.xz
astra-35957b6ef72749cdc520ded67a0eb8cdfd7ea655.zip
Adjust linear/cuda kernels to line integral scaling
Diffstat (limited to 'cuda/2d')
-rw-r--r--cuda/2d/astra.cu3
-rw-r--r--cuda/2d/par_fp.cu10
2 files changed, 6 insertions, 7 deletions
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;