diff options
Diffstat (limited to 'cuda')
| -rw-r--r-- | cuda/2d/fan_bp.cu | 9 | ||||
| -rw-r--r-- | cuda/2d/par_bp.cu | 5 | ||||
| -rw-r--r-- | cuda/2d/sart.cu | 5 |
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); } } |