diff options
| author | Willem Jan Palenstijn <WillemJan.Palenstijn@uantwerpen.be> | 2014-04-16 11:13:40 +0000 |
|---|---|---|
| committer | wpalenst <WillemJan.Palenstijn@uantwerpen.be> | 2014-04-16 11:13:40 +0000 |
| commit | c72bc7cd47ecb5665a287fb88e101f88118f5232 (patch) | |
| tree | 367c19f29647f4256783acfce9db4e8431bd0039 /cuda/2d/sart.cu | |
| parent | bcff7884bb89dbecd394c75a8f57b0a54a743b52 (diff) | |
| download | astra-c72bc7cd47ecb5665a287fb88e101f88118f5232.tar.gz astra-c72bc7cd47ecb5665a287fb88e101f88118f5232.tar.bz2 astra-c72bc7cd47ecb5665a287fb88e101f88118f5232.tar.xz astra-c72bc7cd47ecb5665a287fb88e101f88118f5232.zip | |
Split up processVol in Vol/Sino cases
Diffstat (limited to 'cuda/2d/sart.cu')
| -rw-r--r-- | cuda/2d/sart.cu | 14 |
1 files changed, 7 insertions, 7 deletions
diff --git a/cuda/2d/sart.cu b/cuda/2d/sart.cu index 79c00ef..048661f 100644 --- a/cuda/2d/sart.cu +++ b/cuda/2d/sart.cu @@ -150,14 +150,14 @@ bool SART::precomputeWeights() zeroVolumeData(D_tmpData, tmpPitch, dims); - processVol<opSet>(D_tmpData, 1.0f, tmpPitch, dims.iVolWidth, dims.iVolHeight); + processVol<opSet>(D_tmpData, 1.0f, tmpPitch, dims); callFP(D_tmpData, tmpPitch, D_lineWeight, linePitch, 1.0f); cudaFree(D_tmpData); D_tmpData = 0; } - processVol<opInvert>(D_lineWeight, linePitch, dims.iProjDets, dims.iProjAngles); + processSino<opInvert>(D_lineWeight, linePitch, dims); return true; } @@ -185,7 +185,7 @@ bool SART::iterate(unsigned int iterations) // do FP, subtracting projection from sinogram if (useVolumeMask) { cudaMemcpy2D(D_tmpData, sizeof(float)*tmpPitch, D_volumeData, sizeof(float)*volumePitch, sizeof(float)*(dims.iVolWidth), dims.iVolHeight, cudaMemcpyDeviceToDevice); - processVol<opMul>(D_tmpData, D_maskData, tmpPitch, dims.iVolWidth, dims.iVolHeight); + processVol<opMul>(D_tmpData, D_maskData, tmpPitch, dims); callFP_SART(D_tmpData, tmpPitch, D_projData, projPitch, angle, -1.0f); } else { callFP_SART(D_volumeData, volumePitch, D_projData, projPitch, angle, -1.0f); @@ -198,15 +198,15 @@ bool SART::iterate(unsigned int iterations) // TODO: Try putting the masking directly in the BP zeroVolumeData(D_tmpData, tmpPitch, dims); callBP_SART(D_tmpData, tmpPitch, D_projData, projPitch, angle); - processVol<opAddMul>(D_volumeData, D_maskData, D_tmpData, volumePitch, dims.iVolWidth, dims.iVolHeight); + processVol<opAddMul>(D_volumeData, D_maskData, D_tmpData, volumePitch, dims); } else { callBP_SART(D_volumeData, volumePitch, D_projData, projPitch, angle); } if (useMinConstraint) - processVol<opClampMin>(D_volumeData, fMinConstraint, volumePitch, dims.iVolWidth, dims.iVolHeight); + processVol<opClampMin>(D_volumeData, fMinConstraint, volumePitch, dims); if (useMaxConstraint) - processVol<opClampMax>(D_volumeData, fMaxConstraint, volumePitch, dims.iVolWidth, dims.iVolHeight); + processVol<opClampMax>(D_volumeData, fMaxConstraint, volumePitch, dims); iteration++; @@ -228,7 +228,7 @@ float SART::computeDiffNorm() // do FP, subtracting projection from sinogram if (useVolumeMask) { cudaMemcpy2D(D_tmpData, sizeof(float)*tmpPitch, D_volumeData, sizeof(float)*volumePitch, sizeof(float)*(dims.iVolWidth), dims.iVolHeight, cudaMemcpyDeviceToDevice); - processVol<opMul>(D_tmpData, D_maskData, tmpPitch, dims.iVolWidth, dims.iVolHeight); + processVol<opMul>(D_tmpData, D_maskData, tmpPitch, dims); callFP(D_tmpData, tmpPitch, D_projData, projPitch, -1.0f); } else { callFP(D_volumeData, volumePitch, D_projData, projPitch, -1.0f); |