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/cgls.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/cgls.cu')
| -rw-r--r-- | cuda/2d/cgls.cu | 14 |
1 files changed, 7 insertions, 7 deletions
diff --git a/cuda/2d/cgls.cu b/cuda/2d/cgls.cu index f4175e1..fce8beb 100644 --- a/cuda/2d/cgls.cu +++ b/cuda/2d/cgls.cu @@ -126,7 +126,7 @@ bool CGLS::iterate(unsigned int iterations) if (useVolumeMask) { // Use z as temporary storage here since it is unused cudaMemcpy2D(D_z, sizeof(float)*zPitch, D_volumeData, sizeof(float)*volumePitch, sizeof(float)*(dims.iVolWidth), dims.iVolHeight, cudaMemcpyDeviceToDevice); - processVol<opMul>(D_z, D_maskData, zPitch, dims.iVolWidth, dims.iVolHeight); + processVol<opMul>(D_z, D_maskData, zPitch, dims); callFP(D_z, zPitch, D_r, rPitch, -1.0f); } else { callFP(D_volumeData, volumePitch, D_r, rPitch, -1.0f); @@ -137,7 +137,7 @@ bool CGLS::iterate(unsigned int iterations) zeroVolumeData(D_p, pPitch, dims); callBP(D_p, pPitch, D_r, rPitch); if (useVolumeMask) - processVol<opMul>(D_p, D_maskData, pPitch, dims.iVolWidth, dims.iVolHeight); + processVol<opMul>(D_p, D_maskData, pPitch, dims); gamma = dotProduct2D(D_p, pPitch, dims.iVolWidth, dims.iVolHeight); @@ -158,24 +158,24 @@ bool CGLS::iterate(unsigned int iterations) float alpha = gamma / ww; // x += alpha*p - processVol<opAddScaled>(D_volumeData, D_p, alpha, volumePitch, dims.iVolWidth, dims.iVolHeight); + processVol<opAddScaled>(D_volumeData, D_p, alpha, volumePitch, dims); // r -= alpha*w - processVol<opAddScaled>(D_r, D_w, -alpha, rPitch, dims.iProjDets, dims.iProjAngles); + processSino<opAddScaled>(D_r, D_w, -alpha, rPitch, dims); // z = A'*r zeroVolumeData(D_z, zPitch, dims); callBP(D_z, zPitch, D_r, rPitch); if (useVolumeMask) - processVol<opMul>(D_z, D_maskData, zPitch, dims.iVolWidth, dims.iVolHeight); + processVol<opMul>(D_z, D_maskData, zPitch, dims); float beta = 1.0f / gamma; gamma = dotProduct2D(D_z, zPitch, dims.iVolWidth, dims.iVolHeight); beta *= gamma; // p = z + beta*p - processVol<opScaleAndAdd>(D_p, D_z, beta, pPitch, dims.iVolWidth, dims.iVolHeight); + processVol<opScaleAndAdd>(D_p, D_z, beta, pPitch, dims); } @@ -194,7 +194,7 @@ float CGLS::computeDiffNorm() // do FP, subtracting projection from sinogram if (useVolumeMask) { cudaMemcpy2D(D_z, sizeof(float)*zPitch, D_volumeData, sizeof(float)*volumePitch, sizeof(float)*(dims.iVolWidth), dims.iVolHeight, cudaMemcpyDeviceToDevice); - processVol<opMul>(D_z, D_maskData, zPitch, dims.iVolWidth, dims.iVolHeight); + processVol<opMul>(D_z, D_maskData, zPitch, dims); callFP(D_z, zPitch, D_w, wPitch, -1.0f); } else { callFP(D_volumeData, volumePitch, D_w, wPitch, -1.0f); |