summaryrefslogtreecommitdiffstats
path: root/cuda
diff options
context:
space:
mode:
Diffstat (limited to 'cuda')
-rw-r--r--cuda/2d/fan_bp.cu11
-rw-r--r--cuda/2d/fbp.cu6
2 files changed, 9 insertions, 8 deletions
diff --git a/cuda/2d/fan_bp.cu b/cuda/2d/fan_bp.cu
index 0d21897..428485c 100644
--- a/cuda/2d/fan_bp.cu
+++ b/cuda/2d/fan_bp.cu
@@ -271,11 +271,14 @@ __global__ void devFanBP_FBPWeighted(float* D_volData, unsigned int volPitch, un
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);
+
+ // fDen = || u (x-s) ||
+ // Required scale factor is ( || u s || / || u (x-s) || ) ^ 2.
+ // The factor || u s || ^ 2 is handled by the preweighting
- const float fWeight = fXD*fXD + fYD*fYD;
-
- const float fT = fNum / fDen;
- fVal += tex2D(gT_FanProjTexture, fT, fA) / fWeight;
+ const float fT = fNum * fr;
+ fVal += tex2D(gT_FanProjTexture, fT, fA) * fr * fr;
fA += 1.0f;
}
diff --git a/cuda/2d/fbp.cu b/cuda/2d/fbp.cu
index f0edc19..28cdd92 100644
--- a/cuda/2d/fbp.cu
+++ b/cuda/2d/fbp.cu
@@ -301,7 +301,7 @@ bool FBP::iterate(unsigned int iterations)
astraCUDA3d::FDK_PreWeight(tmp, fOriginSource,
fOriginDetector, 0.0f,
- fFanDetSize, 1.0f, /* fPixelSize */ 1.0f,
+ fFanDetSize, 1.0f, /* fPixelSize, but is normalized */ 1.0f,
m_bShortScan, dims3d, pfAngles);
} else {
// TODO: How should different detector pixel size in different
@@ -328,9 +328,7 @@ bool FBP::iterate(unsigned int iterations)
}
if (fanProjs) {
- float fOutputScale = 1.0 / (/*fPixelSize * fPixelSize * fPixelSize * */ fFanDetSize * fFanDetSize);
-
- ok = FanBP_FBPWeighted(D_volumeData, volumePitch, D_sinoData, sinoPitch, dims, fanProjs, fOutputScale * fProjectorScale * fReconstructionScale);
+ ok = FanBP_FBPWeighted(D_volumeData, volumePitch, D_sinoData, sinoPitch, dims, fanProjs, fProjectorScale * fReconstructionScale);
} else {
// scale by number of angles. For the fan-beam case, this is already