diff options
| author | Willem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl> | 2017-09-27 13:57:04 +0200 |
|---|---|---|
| committer | Willem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl> | 2017-09-27 13:57:04 +0200 |
| commit | 1a8243ed0311c3074a79b97e1730bf3409774b8d (patch) | |
| tree | bc1c499f211c832ab8f28f332f616deb2403b5e9 /cuda | |
| parent | e1a3f0ba1fe7455d0c9183ad07f106aebc1c821f (diff) | |
| download | astra-1a8243ed0311c3074a79b97e1730bf3409774b8d.tar.gz astra-1a8243ed0311c3074a79b97e1730bf3409774b8d.tar.bz2 astra-1a8243ed0311c3074a79b97e1730bf3409774b8d.tar.xz astra-1a8243ed0311c3074a79b97e1730bf3409774b8d.zip | |
Unify some parallel_vec parameter computations
Diffstat (limited to 'cuda')
| -rw-r--r-- | cuda/2d/par_fp.cu | 30 | ||||
| -rw-r--r-- | cuda/2d/sirt.cu | 5 |
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) |