summaryrefslogtreecommitdiffstats
path: root/cuda/2d
diff options
context:
space:
mode:
authorWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2017-09-27 13:57:04 +0200
committerWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2017-09-27 13:57:04 +0200
commit1a8243ed0311c3074a79b97e1730bf3409774b8d (patch)
treebc1c499f211c832ab8f28f332f616deb2403b5e9 /cuda/2d
parente1a3f0ba1fe7455d0c9183ad07f106aebc1c821f (diff)
downloadastra-1a8243ed0311c3074a79b97e1730bf3409774b8d.tar.gz
astra-1a8243ed0311c3074a79b97e1730bf3409774b8d.tar.bz2
astra-1a8243ed0311c3074a79b97e1730bf3409774b8d.tar.xz
astra-1a8243ed0311c3074a79b97e1730bf3409774b8d.zip
Unify some parallel_vec parameter computations
Diffstat (limited to 'cuda/2d')
-rw-r--r--cuda/2d/par_fp.cu30
-rw-r--r--cuda/2d/sirt.cu5
2 files changed, 6 insertions, 29 deletions
diff --git a/cuda/2d/par_fp.cu b/cuda/2d/par_fp.cu
index 10ce683..a0b04ee 100644
--- a/cuda/2d/par_fp.cu
+++ b/cuda/2d/par_fp.cu
@@ -251,36 +251,14 @@ __global__ void FPvertical_simple(float* D_projData, unsigned int projPitch, uns
static void convertAndUploadAngles(const SParProjection *projs, unsigned int nth, unsigned int ndets)
{
float *angles = new float[nth];
- float *offset = new float[nth];
+ float *offsets = new float[nth];
float *detsizes = new float[nth];
- for (int i = 0; i < nth; ++i) {
-
-#warning TODO: Replace by getParParamaters
-
- // Take part of DetU orthogonal to Ray
- double ux = projs[i].fDetUX;
- double uy = projs[i].fDetUY;
-
- double t = (ux * projs[i].fRayX + uy * projs[i].fRayY) / (projs[i].fRayX * projs[i].fRayX + projs[i].fRayY * projs[i].fRayY);
-
- ux -= t * projs[i].fRayX;
- uy -= t * projs[i].fRayY;
-
- double angle = atan2(uy, ux);
-
- angles[i] = angle;
-
- double norm2 = uy * uy + ux * ux;
-
- detsizes[i] = sqrt(norm2);
-
- // CHECKME: SIGNS?
- offset[i] = -0.5*ndets - (projs[i].fDetSY*uy + projs[i].fDetSX*ux) / norm2;
- }
+ for (int i = 0; i < nth; ++i)
+ getParParameters(projs[i], ndets, angles[i], detsizes[i], offsets[i]);
cudaMemcpyToSymbol(gC_angle, angles, nth*sizeof(float), 0, cudaMemcpyHostToDevice);
- cudaMemcpyToSymbol(gC_angle_offset, offset, nth*sizeof(float), 0, cudaMemcpyHostToDevice);
+ cudaMemcpyToSymbol(gC_angle_offset, offsets, nth*sizeof(float), 0, cudaMemcpyHostToDevice);
cudaMemcpyToSymbol(gC_angle_detsize, detsizes, nth*sizeof(float), 0, cudaMemcpyHostToDevice);
}
diff --git a/cuda/2d/sirt.cu b/cuda/2d/sirt.cu
index 2516c6c..b393d7f 100644
--- a/cuda/2d/sirt.cu
+++ b/cuda/2d/sirt.cu
@@ -184,9 +184,8 @@ bool SIRT::doSlabCorrections()
float bound = cosf(1.3963f);
float* t = (float*)D_sinoData;
for (int i = 0; i < dims.iProjAngles; ++i) {
- // TODO: Checkme
- // TODO: Replace by getParParameters
- double angle = atan2(parProjs[i].fRayX, -parProjs[i].fRayY);
+ float angle, detsize, offset;
+ getParParameters(parProjs[i], dims.iProjDets, angle, detsize, offset);
float f = fabs(cosf(angle));
if (f < bound)