From be2da43a560a7241c56e727fb481f1389e9f7fdf Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Mon, 22 Nov 2021 14:43:07 +0100 Subject: De-duplicate 2D texture object creation --- cuda/2d/util.cu | 60 +++++++++++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 60 insertions(+) (limited to 'cuda/2d/util.cu') diff --git a/cuda/2d/util.cu b/cuda/2d/util.cu index ac360f0..2ad3c0f 100644 --- a/cuda/2d/util.cu +++ b/cuda/2d/util.cu @@ -126,6 +126,66 @@ void duplicateProjectionData(float* D_dst, float* D_src, unsigned int pitch, con cudaMemcpy2D(D_dst, sizeof(float)*pitch, D_src, sizeof(float)*pitch, sizeof(float)*dims.iProjDets, dims.iProjAngles, cudaMemcpyDeviceToDevice); } +bool createTextureObject2D(float* data, cudaArray*& dataArray, cudaTextureObject_t& texObj, unsigned int pitch, unsigned int width, unsigned int height) +{ + // TODO: For very small sizes (roughly <=512x128) with few angles (<=180) + // not using an array is more efficient. + + cudaChannelFormatDesc channelDesc = + cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat); + + dataArray = 0; + cudaMallocArray(&dataArray, &channelDesc, width, height); + cudaMemcpy2DToArray(dataArray, 0, 0, data, pitch*sizeof(float), width*sizeof(float), height, cudaMemcpyDeviceToDevice); + + cudaResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = cudaResourceTypeArray; + resDesc.res.array.array = dataArray; + + cudaTextureDesc texDesc; + memset(&texDesc, 0, sizeof(texDesc)); + texDesc.addressMode[0] = cudaAddressModeBorder; + texDesc.addressMode[1] = cudaAddressModeBorder; + texDesc.filterMode = cudaFilterModeLinear; + texDesc.readMode = cudaReadModeElementType; + texDesc.normalizedCoords = 0; + + texObj = 0; + + return checkCuda(cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL), "createTextureObject2D"); +} + +bool createTextureObjectPitch2D(float* data, cudaTextureObject_t& texObj, unsigned int pitch, unsigned int width, unsigned int height, cudaTextureAddressMode mode) +{ + cudaChannelFormatDesc channelDesc = + cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat); + + cudaResourceDesc resDesc; + memset(&resDesc, 0, sizeof(resDesc)); + resDesc.resType = cudaResourceTypePitch2D; + resDesc.res.pitch2D.devPtr = (void*)data; + resDesc.res.pitch2D.desc = channelDesc; + resDesc.res.pitch2D.width = width; + resDesc.res.pitch2D.height = height; + resDesc.res.pitch2D.pitchInBytes = sizeof(float)*pitch; + + cudaTextureDesc texDesc; + memset(&texDesc, 0, sizeof(texDesc)); + texDesc.addressMode[0] = mode; + texDesc.addressMode[1] = mode; + texDesc.filterMode = cudaFilterModeLinear; + texDesc.readMode = cudaReadModeElementType; + texDesc.normalizedCoords = 0; + + texObj = 0; + + return checkCuda(cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL), "createTextureObjectPitch2D"); +} + + + + template __global__ void reduce1D(float *g_idata, float *g_odata, unsigned int n) { -- cgit v1.2.3 From 7cad7b813838ed2ddb65a4c9ea1c08c625c50043 Mon Sep 17 00:00:00 2001 From: Willem Jan Palenstijn Date: Mon, 22 Nov 2021 14:44:50 +0100 Subject: Fix memleak in error handling --- cuda/2d/util.cu | 17 +++++++++++++---- 1 file changed, 13 insertions(+), 4 deletions(-) (limited to 'cuda/2d/util.cu') diff --git a/cuda/2d/util.cu b/cuda/2d/util.cu index 2ad3c0f..4a58880 100644 --- a/cuda/2d/util.cu +++ b/cuda/2d/util.cu @@ -126,7 +126,7 @@ void duplicateProjectionData(float* D_dst, float* D_src, unsigned int pitch, con cudaMemcpy2D(D_dst, sizeof(float)*pitch, D_src, sizeof(float)*pitch, sizeof(float)*dims.iProjDets, dims.iProjAngles, cudaMemcpyDeviceToDevice); } -bool createTextureObject2D(float* data, cudaArray*& dataArray, cudaTextureObject_t& texObj, unsigned int pitch, unsigned int width, unsigned int height) +bool createArrayAndTextureObject2D(float* data, cudaArray*& dataArray, cudaTextureObject_t& texObj, unsigned int pitch, unsigned int width, unsigned int height) { // TODO: For very small sizes (roughly <=512x128) with few angles (<=180) // not using an array is more efficient. @@ -135,8 +135,12 @@ bool createTextureObject2D(float* data, cudaArray*& dataArray, cudaTextureObject cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat); dataArray = 0; - cudaMallocArray(&dataArray, &channelDesc, width, height); - cudaMemcpy2DToArray(dataArray, 0, 0, data, pitch*sizeof(float), width*sizeof(float), height, cudaMemcpyDeviceToDevice); + if (!checkCuda(cudaMallocArray(&dataArray, &channelDesc, width, height), "createTextureObject2D malloc")) + return false; + if (!checkCuda(cudaMemcpy2DToArray(dataArray, 0, 0, data, pitch*sizeof(float), width*sizeof(float), height, cudaMemcpyDeviceToDevice), "createTextureObject2D memcpy")) { + cudaFreeArray(dataArray); + return false; + } cudaResourceDesc resDesc; memset(&resDesc, 0, sizeof(resDesc)); @@ -153,7 +157,12 @@ bool createTextureObject2D(float* data, cudaArray*& dataArray, cudaTextureObject texObj = 0; - return checkCuda(cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL), "createTextureObject2D"); + if (!checkCuda(cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL), "createTextureObject2D")) { + cudaFreeArray(dataArray); + return false; + } + + return true; } bool createTextureObjectPitch2D(float* data, cudaTextureObject_t& texObj, unsigned int pitch, unsigned int width, unsigned int height, cudaTextureAddressMode mode) -- cgit v1.2.3