summaryrefslogtreecommitdiffstats
path: root/cuda/2d/par_fp.cu
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/par_fp.cu
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/par_fp.cu')
-rw-r--r--cuda/2d/par_fp.cu30
1 files changed, 4 insertions, 26 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);
}