summaryrefslogtreecommitdiffstats
path: root/cuda/3d/par3d_bp.cu
diff options
context:
space:
mode:
authorWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2021-11-22 14:14:55 +0100
committerWillem Jan Palenstijn <wjp@usecode.org>2021-11-26 12:09:09 +0100
commitd4645875801d29402b1e3f2a42a5d9902a37a718 (patch)
treeb81497857784e28ca0743b6ecbebb3299b8ef557 /cuda/3d/par3d_bp.cu
parent64d42baf1bd15df8e0ecd00d89bb5c7be787cac2 (diff)
downloadastra-d4645875801d29402b1e3f2a42a5d9902a37a718.tar.gz
astra-d4645875801d29402b1e3f2a42a5d9902a37a718.tar.bz2
astra-d4645875801d29402b1e3f2a42a5d9902a37a718.tar.xz
astra-d4645875801d29402b1e3f2a42a5d9902a37a718.zip
De-duplicate 3D texture object creation
Diffstat (limited to 'cuda/3d/par3d_bp.cu')
-rw-r--r--cuda/3d/par3d_bp.cu26
1 files changed, 2 insertions, 24 deletions
diff --git a/cuda/3d/par3d_bp.cu b/cuda/3d/par3d_bp.cu
index 748086e..27d95fe 100644
--- a/cuda/3d/par3d_bp.cu
+++ b/cuda/3d/par3d_bp.cu
@@ -54,29 +54,6 @@ __constant__ DevPar3DParams gC_C[g_MaxAngles];
__constant__ float gC_scale[g_MaxAngles];
-static bool bindProjDataTexture(cudaArray* array, cudaTextureObject_t& texObj)
-{
- cudaChannelFormatDesc channelDesc =
- cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat);
-
- cudaResourceDesc resDesc;
- memset(&resDesc, 0, sizeof(resDesc));
- resDesc.resType = cudaResourceTypeArray;
- resDesc.res.array.array = array;
-
- cudaTextureDesc texDesc;
- memset(&texDesc, 0, sizeof(texDesc));
- texDesc.addressMode[0] = cudaAddressModeBorder;
- texDesc.addressMode[1] = cudaAddressModeBorder;
- texDesc.addressMode[2] = cudaAddressModeBorder;
- texDesc.filterMode = cudaFilterModeLinear;
- texDesc.readMode = cudaReadModeElementType;
- texDesc.normalizedCoords = 0;
-
- return checkCuda(cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL), "par3d_bp texture");
-}
-
-
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)
{
@@ -261,7 +238,8 @@ bool Par3DBP_Array(cudaPitchedPtr D_volumeData,
const SProjectorParams3D& params)
{
cudaTextureObject_t D_texObj;
- bindProjDataTexture(D_projArray, D_texObj);
+ if (!createTextureObject3D(D_projArray, D_texObj))
+ return false;
float fOutputScale = params.fOutputScale * params.fVolScaleX * params.fVolScaleY * params.fVolScaleZ;