From accc4439d9dd035765ed77d94a0ceece3270cc0b Mon Sep 17 00:00:00 2001 From: "Suren A. Chilingaryan" Date: Tue, 26 Jul 2022 23:32:36 +0200 Subject: Half-precision back-/forward-projection for parallel geometry --- cuda/3d/par3d_bp.cu | 21 ++++++++++++++++++--- 1 file changed, 18 insertions(+), 3 deletions(-) (limited to 'cuda/3d/par3d_bp.cu') diff --git a/cuda/3d/par3d_bp.cu b/cuda/3d/par3d_bp.cu index e43479a..75cbf2b 100644 --- a/cuda/3d/par3d_bp.cu +++ b/cuda/3d/par3d_bp.cu @@ -75,7 +75,7 @@ static bool bindProjDataTexture(const cudaArray* array) return true; } - +#include "rounding.h" __global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, int startAngle, int angleOffset, const SDimensions3D dims, float fOutputScale) { float* volData = (float*)D_volData; @@ -122,8 +122,23 @@ __global__ void dev_par3D_BP(void* D_volData, unsigned int volPitch, int startAn float fV = fCv.w + fX * fCv.x + fY * fCv.y + fZ * fCv.z; for (int idx = 0; idx < ZSIZE; ++idx) { - - float fVal = tex3D(gT_par3DProjTexture, fU, fAngle, fV); + 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(gT_par3DProjTexture, fUf - 0.5f, fAngle, fV)); + textype fVal2 = texto(tex3D(gT_par3DProjTexture, fUf + 0.5f, fAngle, fV)); + fVal = texfrom(fVal1 + (fU_ + h5 - fUf_) * (fVal2 - fVal1)); + } else { + textype fVal1 = texto(tex3D(gT_par3DProjTexture, fUf + 0.5f, fAngle, fV)); + textype fVal2 = texto(tex3D(gT_par3DProjTexture, fUf + 1.5f, fAngle, fV)); + fVal = texfrom(fVal1 + (fU_ - h5 - fUf_) * (fVal2 - fVal1)); + } + +// float fVal = tex3D(gT_par3DProjTexture, fU, fAngle, fV); Z[idx] += fVal; fU += fCu.z; -- cgit v1.2.3