diff options
| author | Suren A. Chilingaryan <csa@ipecompute4.ands.kit.edu> | 2022-07-26 23:32:36 +0200 |
|---|---|---|
| committer | Suren A. Chilingaryan <csa@ipecompute4.ands.kit.edu> | 2022-07-26 23:32:36 +0200 |
| commit | accc4439d9dd035765ed77d94a0ceece3270cc0b (patch) | |
| tree | 886507357f95680e8a6a7dc80aefe04ca42c5b2d /cuda/3d/par3d_fp.cu | |
| parent | bfceef4da377c32cd59d1e51efb9aa1a21c7f4bd (diff) | |
| download | astra-tnv.tar.gz astra-tnv.tar.bz2 astra-tnv.tar.xz astra-tnv.zip | |
Half-precision back-/forward-projection for parallel geometrytnv
Diffstat (limited to 'cuda/3d/par3d_fp.cu')
| -rw-r--r-- | cuda/3d/par3d_fp.cu | 19 |
1 files changed, 18 insertions, 1 deletions
diff --git a/cuda/3d/par3d_fp.cu b/cuda/3d/par3d_fp.cu index a99308f..3ad9f0d 100644 --- a/cuda/3d/par3d_fp.cu +++ b/cuda/3d/par3d_fp.cu @@ -146,6 +146,7 @@ struct SCALE_NONCUBE { // blockIdx: x = u/v detector // y = angle block +#include "rounding.h" template<class COORD, class SCALE> __global__ void par3D_FP_t(float* D_projData, unsigned int projPitch, @@ -212,7 +213,23 @@ __global__ void par3D_FP_t(float* D_projData, unsigned int projPitch, for (int s = startSlice; s < endSlice; ++s) { - fVal += c.tex(f0, f1, f2); + textype h5 = texto(0.5f); + textype f1_ = texto(f1); + textype f1f_ = texto(floor(f1)); + float f1f = floor(f1); + + if ((f1 - f1f) < 0.5f) { + textype fVal1 = texto(c.tex(f0, f1f - 0.5f, f2)); + textype fVal2 = texto(c.tex(f0, f1f + 0.5f, f2)); + fVal += texfrom(fVal1 + (f1_ + h5 - f1f_) * (fVal2 - fVal1)); +// fVal += texfrom(__hfma(__hadd(h5,__hsub(f1_, f1f_)), __hsub(fVal2, fVal1), fVal1)); + } else { + textype fVal1 = texto(c.tex(f0, f1f + 0.5f, f2)); + textype fVal2 = texto(c.tex(f0, f1f + 1.5f, f2)); + fVal += texfrom(fVal1 + (f1_ - h5 - f1f_) * (fVal2 - fVal1)); + } + +// fVal += c.tex(f0, f1, f2); f0 += 1.0f; f1 += a1; f2 += a2; |