summaryrefslogtreecommitdiffstats
path: root/cuda/3d/par3d_fp.cu
diff options
context:
space:
mode:
authorWillem Jan Palenstijn <WillemJan.Palenstijn@uantwerpen.be>2014-04-16 11:13:01 +0000
committerwpalenst <WillemJan.Palenstijn@uantwerpen.be>2014-04-16 11:13:01 +0000
commit7ce0b7cca179e903e8011cd96c9910cbdf62ae00 (patch)
tree2ebfa5687c8126d1d2b345a1bfc8c374a62a227b /cuda/3d/par3d_fp.cu
parent3a6769465bee7d56d0ddff36613b886446421e07 (diff)
downloadastra-7ce0b7cca179e903e8011cd96c9910cbdf62ae00.tar.gz
astra-7ce0b7cca179e903e8011cd96c9910cbdf62ae00.tar.bz2
astra-7ce0b7cca179e903e8011cd96c9910cbdf62ae00.tar.xz
astra-7ce0b7cca179e903e8011cd96c9910cbdf62ae00.zip
Remove padding in 3D cuda in favour of Border mode
Diffstat (limited to 'cuda/3d/par3d_fp.cu')
-rw-r--r--cuda/3d/par3d_fp.cu38
1 files changed, 19 insertions, 19 deletions
diff --git a/cuda/3d/par3d_fp.cu b/cuda/3d/par3d_fp.cu
index 6bf9037..e0f743c 100644
--- a/cuda/3d/par3d_fp.cu
+++ b/cuda/3d/par3d_fp.cu
@@ -72,9 +72,9 @@ static bool bindVolumeDataTexture(const cudaArray* array)
{
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
- gT_par3DVolumeTexture.addressMode[0] = cudaAddressModeClamp;
- gT_par3DVolumeTexture.addressMode[1] = cudaAddressModeClamp;
- gT_par3DVolumeTexture.addressMode[2] = cudaAddressModeClamp;
+ gT_par3DVolumeTexture.addressMode[0] = cudaAddressModeBorder;
+ gT_par3DVolumeTexture.addressMode[1] = cudaAddressModeBorder;
+ gT_par3DVolumeTexture.addressMode[2] = cudaAddressModeBorder;
gT_par3DVolumeTexture.filterMode = cudaFilterModeLinear;
gT_par3DVolumeTexture.normalized = false;
@@ -144,9 +144,9 @@ static bool bindVolumeDataTexture(const cudaArray* array)
\
float fVal = 0.0f; \
\
- float f##c0 = startSlice + 1.5f; \
- float f##c1 = a##c1 * (startSlice - 0.5f*dims.iVol##c0 + 0.5f) + b##c1 + 0.5f*dims.iVol##c1 - 0.5f + 1.5f;\
- float f##c2 = a##c2 * (startSlice - 0.5f*dims.iVol##c0 + 0.5f) + b##c2 + 0.5f*dims.iVol##c2 - 0.5f + 1.5f;\
+ float f##c0 = startSlice + 0.5f; \
+ float f##c1 = a##c1 * (startSlice - 0.5f*dims.iVol##c0 + 0.5f) + b##c1 + 0.5f*dims.iVol##c1 - 0.5f + 0.5f;\
+ float f##c2 = a##c2 * (startSlice - 0.5f*dims.iVol##c0 + 0.5f) + b##c2 + 0.5f*dims.iVol##c2 - 0.5f + 0.5f;\
\
for (int s = startSlice; s < endSlice; ++s) \
{ \
@@ -226,9 +226,9 @@ static bool bindVolumeDataTexture(const cudaArray* array)
\
float fVal = 0.0f; \
\
- float f##c0 = startSlice + 1.5f; \
- float f##c1 = a##c1 * (startSlice - 0.5f*dims.iVol##c0 + 0.5f) + b##c1 + 0.5f*dims.iVol##c1 - 0.5f + 1.5f;\
- float f##c2 = a##c2 * (startSlice - 0.5f*dims.iVol##c0 + 0.5f) + b##c2 + 0.5f*dims.iVol##c2 - 0.5f + 1.5f;\
+ float f##c0 = startSlice + 0.5f; \
+ float f##c1 = a##c1 * (startSlice - 0.5f*dims.iVol##c0 + 0.5f) + b##c1 + 0.5f*dims.iVol##c1 - 0.5f + 0.5f;\
+ float f##c2 = a##c2 * (startSlice - 0.5f*dims.iVol##c0 + 0.5f) + b##c2 + 0.5f*dims.iVol##c2 - 0.5f + 0.5f;\
\
for (int s = startSlice; s < endSlice; ++s) \
{ \
@@ -281,16 +281,16 @@ PAR3D_FP_SS_BODY(Z,X,Y)
__device__ float dirWeights(float fX, float fN) {
- if (fX <= 0.5f) // outside image on left
+ if (fX <= -0.5f) // outside image on left
return 0.0f;
- if (fX <= 1.5f) // half outside image on left
- return (fX - 0.5f) * (fX - 0.5f);
- if (fX <= fN + 0.5f) { // inside image
- float t = fX - 0.5f - floorf(fX - 0.5f);
+ if (fX <= 0.5f) // half outside image on left
+ return (fX + 0.5f) * (fX + 0.5f);
+ if (fX <= fN - 0.5f) { // inside image
+ float t = fX + 0.5f - floorf(fX + 0.5f);
return t*t + (1-t)*(1-t);
}
- if (fX <= fN + 1.5f) // half outside image on right
- return (fN + 1.5f - fX) * (fN + 1.5f - fX);
+ if (fX <= fN + 0.5f) // half outside image on right
+ return (fN + 0.5f - fX) * (fN + 0.5f - fX);
return 0.0f; // outside image on right
}
@@ -346,9 +346,9 @@ __device__ float dirWeights(float fX, float fN) {
\
float fVal = 0.0f; \
\
- float f##c0 = startSlice + 1.5f; \
- float f##c1 = a##c1 * (startSlice - 0.5f*dims.iVol##c0 + 0.5f) + b##c1 + 0.5f*dims.iVol##c1 - 0.5f + 1.5f;\
- float f##c2 = a##c2 * (startSlice - 0.5f*dims.iVol##c0 + 0.5f) + b##c2 + 0.5f*dims.iVol##c2 - 0.5f + 1.5f;\
+ float f##c0 = startSlice + 0.5f; \
+ float f##c1 = a##c1 * (startSlice - 0.5f*dims.iVol##c0 + 0.5f) + b##c1 + 0.5f*dims.iVol##c1 - 0.5f + 0.5f;\
+ float f##c2 = a##c2 * (startSlice - 0.5f*dims.iVol##c0 + 0.5f) + b##c2 + 0.5f*dims.iVol##c2 - 0.5f + 0.5f;\
\
for (int s = startSlice; s < endSlice; ++s) \
{ \