From 7ce0b7cca179e903e8011cd96c9910cbdf62ae00 Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Wed, 16 Apr 2014 11:13:01 +0000 Subject: Remove padding in 3D cuda in favour of Border mode --- cuda/3d/cone_bp.cu | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) (limited to 'cuda/3d/cone_bp.cu') diff --git a/cuda/3d/cone_bp.cu b/cuda/3d/cone_bp.cu index 7f8e320..1b0154d 100644 --- a/cuda/3d/cone_bp.cu +++ b/cuda/3d/cone_bp.cu @@ -73,9 +73,9 @@ bool bindProjDataTexture(const cudaArray* array) { cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(); - gT_coneProjTexture.addressMode[0] = cudaAddressModeClamp; - gT_coneProjTexture.addressMode[1] = cudaAddressModeClamp; - gT_coneProjTexture.addressMode[2] = cudaAddressModeClamp; + gT_coneProjTexture.addressMode[0] = cudaAddressModeBorder; + gT_coneProjTexture.addressMode[1] = cudaAddressModeBorder; + gT_coneProjTexture.addressMode[2] = cudaAddressModeBorder; gT_coneProjTexture.filterMode = cudaFilterModeLinear; gT_coneProjTexture.normalized = false; @@ -148,8 +148,8 @@ __global__ void dev_cone_BP(void* D_volData, unsigned int volPitch, int startAng const float fVNum = fCvc + fX * fCvx + fY * fCvy + fZ * fCvz; const float fDen = fCdc + fX * fCdx + fY * fCdy + fZ * fCdz; - const float fU = fUNum / fDen + 1.0f; - const float fV = fVNum / fDen + 1.0f; + const float fU = fUNum / fDen; + const float fV = fVNum / fDen; fVal += tex3D(gT_coneProjTexture, fU, fAngle, fV); @@ -230,8 +230,8 @@ __global__ void dev_cone_BP_SS(void* D_volData, unsigned int volPitch, int start const float fVNum = fCvc + fXs * fCvx + fYs * fCvy + fZs * fCvz; const float fDen = fCdc + fXs * fCdx + fYs * fCdy + fZs * fCdz; - const float fU = fUNum / fDen + 1.0f; - const float fV = fVNum / fDen + 1.0f; + const float fU = fUNum / fDen; + const float fV = fVNum / fDen; fVal += tex3D(gT_coneProjTexture, fU, fAngle, fV); -- cgit v1.2.3