From 26fd23c53f35e374618b660a82d9de1755490a62 Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Mon, 22 Nov 2021 14:01:25 +0100 Subject: Replace texref by texobj in cone_fp --- cuda/3d/cone_fp.cu | 78 +++++++++++++++++++++++++++++------------------------- 1 file changed, 42 insertions(+), 36 deletions(-) (limited to 'cuda') diff --git a/cuda/3d/cone_fp.cu b/cuda/3d/cone_fp.cu index 2c3d1f6..e49ea24 100644 --- a/cuda/3d/cone_fp.cu +++ b/cuda/3d/cone_fp.cu @@ -35,10 +35,6 @@ along with the ASTRA Toolbox. If not, see . #include -typedef texture texture3D; - -static texture3D gT_coneVolumeTexture; - namespace astraCUDA3d { static const unsigned int g_anglesPerBlock = 4; @@ -63,21 +59,26 @@ __constant__ float gC_DetVY[g_MaxAngles]; __constant__ float gC_DetVZ[g_MaxAngles]; -bool bindVolumeDataTexture(const cudaArray* array) +bool bindVolumeDataTexture(cudaArray* array, cudaTextureObject_t& texObj) { - cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(); - - gT_coneVolumeTexture.addressMode[0] = cudaAddressModeBorder; - gT_coneVolumeTexture.addressMode[1] = cudaAddressModeBorder; - gT_coneVolumeTexture.addressMode[2] = cudaAddressModeBorder; - gT_coneVolumeTexture.filterMode = cudaFilterModeLinear; - gT_coneVolumeTexture.normalized = false; - - cudaBindTextureToArray(gT_coneVolumeTexture, array, channelDesc); - - // TODO: error value? - - return true; + 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), "cone_fp texture"); } @@ -89,7 +90,7 @@ struct DIR_X { __device__ float c0(float x, float y, float z) const { return x; } __device__ float c1(float x, float y, float z) const { return y; } __device__ float c2(float x, float y, float z) const { return z; } - __device__ float tex(float f0, float f1, float f2) const { return tex3D(gT_coneVolumeTexture, f0, f1, f2); } + __device__ float tex(cudaTextureObject_t tex, float f0, float f1, float f2) const { return tex3D(tex, f0, f1, f2); } __device__ float x(float f0, float f1, float f2) const { return f0; } __device__ float y(float f0, float f1, float f2) const { return f1; } __device__ float z(float f0, float f1, float f2) const { return f2; } @@ -103,7 +104,7 @@ struct DIR_Y { __device__ float c0(float x, float y, float z) const { return y; } __device__ float c1(float x, float y, float z) const { return x; } __device__ float c2(float x, float y, float z) const { return z; } - __device__ float tex(float f0, float f1, float f2) const { return tex3D(gT_coneVolumeTexture, f1, f0, f2); } + __device__ float tex(cudaTextureObject_t tex, float f0, float f1, float f2) const { return tex3D(tex, f1, f0, f2); } __device__ float x(float f0, float f1, float f2) const { return f1; } __device__ float y(float f0, float f1, float f2) const { return f0; } __device__ float z(float f0, float f1, float f2) const { return f2; } @@ -117,7 +118,7 @@ struct DIR_Z { __device__ float c0(float x, float y, float z) const { return z; } __device__ float c1(float x, float y, float z) const { return x; } __device__ float c2(float x, float y, float z) const { return y; } - __device__ float tex(float f0, float f1, float f2) const { return tex3D(gT_coneVolumeTexture, f1, f2, f0); } + __device__ float tex(cudaTextureObject_t tex, float f0, float f1, float f2) const { return tex3D(tex, f1, f2, f0); } __device__ float x(float f0, float f1, float f2) const { return f1; } __device__ float y(float f0, float f1, float f2) const { return f2; } __device__ float z(float f0, float f1, float f2) const { return f0; } @@ -144,6 +145,7 @@ struct SCALE_NONCUBE { template __global__ void cone_FP_t(float* D_projData, unsigned int projPitch, + cudaTextureObject_t tex, unsigned int startSlice, unsigned int startAngle, unsigned int endAngle, const SDimensions3D dims, @@ -208,7 +210,7 @@ __global__ void cone_FP_t(float* D_projData, unsigned int projPitch, for (int s = startSlice; s < endSlice; ++s) { - fVal += c.tex(f0, f1, f2); + fVal += c.tex(tex, f0, f1, f2); f0 += 1.0f; f1 += a1; f2 += a2; @@ -222,6 +224,7 @@ __global__ void cone_FP_t(float* D_projData, unsigned int projPitch, template __global__ void cone_FP_SS_t(float* D_projData, unsigned int projPitch, + cudaTextureObject_t tex, unsigned int startSlice, unsigned int startAngle, unsigned int endAngle, const SDimensions3D dims, int iRaysPerDetDim, @@ -295,7 +298,7 @@ __global__ void cone_FP_SS_t(float* D_projData, unsigned int projPitch, for (int s = startSlice; s < endSlice; ++s) { - fVal += c.tex(f0, f1, f2); + fVal += c.tex(tex, f0, f1, f2); f0 += 1.0f; f1 += a1; f2 += a2; @@ -313,7 +316,9 @@ __global__ void cone_FP_SS_t(float* D_projData, unsigned int projPitch, bool ConeFP_Array_internal(cudaPitchedPtr D_projData, - const SDimensions3D& dims, unsigned int angleCount, const SConeProjection* angles, + cudaTextureObject_t D_texObj, + const SDimensions3D& dims, + unsigned int angleCount, const SConeProjection* angles, const SProjectorParams3D& params) { // transfer angles to constant memory @@ -419,29 +424,29 @@ bool ConeFP_Array_internal(cudaPitchedPtr D_projData, for (unsigned int i = 0; i < dims.iVolX; i += g_blockSlices) if (params.iRaysPerDetDim == 1) if (cube) - cone_FP_t<<>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, scube); + cone_FP_t<<>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, scube); else - cone_FP_t<<>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, snoncubeX); + cone_FP_t<<>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, snoncubeX); else - cone_FP_SS_t<<>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeX); + cone_FP_SS_t<<>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeX); } else if (blockDirection == 1) { for (unsigned int i = 0; i < dims.iVolY; i += g_blockSlices) if (params.iRaysPerDetDim == 1) if (cube) - cone_FP_t<<>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, scube); + cone_FP_t<<>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, scube); else - cone_FP_t<<>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, snoncubeY); + cone_FP_t<<>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, snoncubeY); else - cone_FP_SS_t<<>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeY); + cone_FP_SS_t<<>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeY); } else if (blockDirection == 2) { for (unsigned int i = 0; i < dims.iVolZ; i += g_blockSlices) if (params.iRaysPerDetDim == 1) if (cube) - cone_FP_t<<>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, scube); + cone_FP_t<<>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, scube); else - cone_FP_t<<>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, snoncubeZ); + cone_FP_t<<>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, snoncubeZ); else - cone_FP_SS_t<<>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeZ); + cone_FP_SS_t<<>>((float*)D_projData.ptr, D_projData.pitch/sizeof(float), D_texObj, i, blockStart, blockEnd, dims, params.iRaysPerDetDim, snoncubeZ); } } @@ -469,11 +474,12 @@ bool ConeFP(cudaPitchedPtr D_volumeData, const SDimensions3D& dims, const SConeProjection* angles, const SProjectorParams3D& params) { - // transfer volume to array + cudaTextureObject_t D_texObj; + // transfer volume to array cudaArray* cuArray = allocateVolumeArray(dims); transferVolumeToArray(D_volumeData, cuArray, dims); - bindVolumeDataTexture(cuArray); + bindVolumeDataTexture(cuArray, D_texObj); bool ret; @@ -485,7 +491,7 @@ bool ConeFP(cudaPitchedPtr D_volumeData, cudaPitchedPtr D_subprojData = D_projData; D_subprojData.ptr = (char*)D_projData.ptr + iAngle * D_projData.pitch; - ret = ConeFP_Array_internal(D_subprojData, + ret = ConeFP_Array_internal(D_subprojData, D_texObj, dims, iEndAngle - iAngle, angles + iAngle, params); if (!ret) -- cgit v1.2.3