From c72bc7cd47ecb5665a287fb88e101f88118f5232 Mon Sep 17 00:00:00 2001
From: Willem Jan Palenstijn <WillemJan.Palenstijn@uantwerpen.be>
Date: Wed, 16 Apr 2014 11:13:40 +0000
Subject: Split up processVol in Vol/Sino cases

---
 cuda/2d/cgls.cu | 14 +++++++-------
 1 file changed, 7 insertions(+), 7 deletions(-)

(limited to 'cuda/2d/cgls.cu')

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);
-- 
cgit v1.2.3