summaryrefslogtreecommitdiffstats
path: root/cuda/3d/par3d_bp.cu
diff options
context:
space:
mode:
Diffstat (limited to 'cuda/3d/par3d_bp.cu')
-rw-r--r--cuda/3d/par3d_bp.cu21
1 files changed, 18 insertions, 3 deletions
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;