summaryrefslogtreecommitdiffstats
path: root/cuda/2d
diff options
context:
space:
mode:
authorWillem Jan Palenstijn <wjp@usecode.org>2019-03-30 21:21:16 +0100
committerWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2019-09-25 14:10:09 +0200
commit12f86554bafb66e6afc6193f181527ba0749de92 (patch)
tree7d89d128c3011251d19be98333eefca450e99c5b /cuda/2d
parent11114e33fb504b0b74f3d239e77fe248a027cc23 (diff)
downloadastra-12f86554bafb66e6afc6193f181527ba0749de92.tar.gz
astra-12f86554bafb66e6afc6193f181527ba0749de92.tar.bz2
astra-12f86554bafb66e6afc6193f181527ba0749de92.tar.xz
astra-12f86554bafb66e6afc6193f181527ba0749de92.zip
Adjust SART to line integral scaling
Diffstat (limited to 'cuda/2d')
-rw-r--r--cuda/2d/fan_bp.cu9
-rw-r--r--cuda/2d/par_bp.cu5
-rw-r--r--cuda/2d/sart.cu5
3 files changed, 10 insertions, 9 deletions
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);
}
}