diff options
| author | Suren A. Chilingaryan <csa@ipecompute4.ands.kit.edu> | 2022-07-26 23:29:12 +0200 |
|---|---|---|
| committer | Suren A. Chilingaryan <csa@ipecompute4.ands.kit.edu> | 2022-07-26 23:29:12 +0200 |
| commit | 9888385e475a747250ab802ee137b76f2b87925d (patch) | |
| tree | 9fd968221e25d71dec56cb285e0d5942b5751c38 /cuda/3d/par3d_bp.cu | |
| parent | 0bad868e66827f94ae93a909dc8ea330a8540a11 (diff) | |
| download | astra-9888385e475a747250ab802ee137b76f2b87925d.tar.gz astra-9888385e475a747250ab802ee137b76f2b87925d.tar.bz2 astra-9888385e475a747250ab802ee137b76f2b87925d.tar.xz astra-9888385e475a747250ab802ee137b76f2b87925d.zip | |
Half-precision back-/forward-projection for parallel geometryhalf
Diffstat (limited to 'cuda/3d/par3d_bp.cu')
| -rw-r--r-- | cuda/3d/par3d_bp.cu | 20 |
1 files changed, 19 insertions, 1 deletions
diff --git a/cuda/3d/par3d_bp.cu b/cuda/3d/par3d_bp.cu index 27d95fe..7958ac9 100644 --- a/cuda/3d/par3d_bp.cu +++ b/cuda/3d/par3d_bp.cu @@ -54,6 +54,8 @@ __constant__ DevPar3DParams gC_C[g_MaxAngles]; __constant__ float gC_scale[g_MaxAngles]; +#include "rounding.h" + template<unsigned int ZSIZE> __global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, cudaTextureObject_t tex, int startAngle, int angleOffset, const SDimensions3D dims, float fOutputScale) { @@ -100,10 +102,26 @@ __global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, cudaTexture float fU = fCu.w + fX * fCu.x + fY * fCu.y + fZ * fCu.z; float fV = fCv.w + fX * fCv.x + fY * fCv.y + fZ * fCv.z; +// printf("%f %f\n", fU, fV); for (int idx = 0; idx < ZSIZE; ++idx) { + float fVal; + textype h5 = texto(0.5f); + textype fU_ = texto(fU); + textype fUf_ = texto(floor(fU)); + float fUf = floor(fU); + + if ((fU - fUf) < 0.5f) { + textype fVal1 = texto(tex3D<float>(tex, fUf - 0.5f, fAngle, fV)); + textype fVal2 = texto(tex3D<float>(tex, fUf + 0.5f, fAngle, fV)); + fVal = texfrom(fVal1 + (fU_ + h5 - fUf_) * (fVal2 - fVal1)); + } else { + textype fVal1 = texto(tex3D<float>(tex, fUf + 0.5f, fAngle, fV)); + textype fVal2 = texto(tex3D<float>(tex, fUf + 1.5f, fAngle, fV)); + fVal = texfrom(fVal1 + (fU_ - h5 - fUf_) * (fVal2 - fVal1)); + } - float fVal = tex3D<float>(tex, fU, fAngle, fV); +// float fVal = tex3D<float>(tex, fU, fAngle, fV); Z[idx] += fVal * fS; fU += fCu.z; |