diff options
| -rw-r--r-- | cuda/2d/fan_fp.cu | 2 | ||||
| -rw-r--r-- | cuda/2d/par_fp.cu | 2 | ||||
| -rw-r--r-- | cuda/2d/util.cu | 17 | ||||
| -rw-r--r-- | include/astra/cuda/2d/util.h | 2 |
4 files changed, 16 insertions, 7 deletions
diff --git a/cuda/2d/fan_fp.cu b/cuda/2d/fan_fp.cu index ed99e08..5aa9674 100644 --- a/cuda/2d/fan_fp.cu +++ b/cuda/2d/fan_fp.cu @@ -201,7 +201,7 @@ bool FanFP_internal(float* D_volumeData, unsigned int volumePitch, cudaArray* D_dataArray; cudaTextureObject_t D_texObj; - if (!createTextureObject2D(D_volumeData, D_dataArray, D_texObj, volumePitch, dims.iVolWidth, dims.iVolHeight)) + if (!createArrayAndTextureObject2D(D_volumeData, D_dataArray, D_texObj, volumePitch, dims.iVolWidth, dims.iVolHeight)) return false; // transfer angles to constant memory diff --git a/cuda/2d/par_fp.cu b/cuda/2d/par_fp.cu index 6035e0c..3d9e842 100644 --- a/cuda/2d/par_fp.cu +++ b/cuda/2d/par_fp.cu @@ -238,7 +238,7 @@ bool FP_simple_internal(float* D_volumeData, unsigned int volumePitch, cudaArray* D_dataArray; cudaTextureObject_t D_texObj; - if (!createTextureObject2D(D_volumeData, D_dataArray, D_texObj, volumePitch, dims.iVolWidth, dims.iVolHeight)) + if (!createArrayAndTextureObject2D(D_volumeData, D_dataArray, D_texObj, volumePitch, dims.iVolWidth, dims.iVolHeight)) return false; 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) diff --git a/include/astra/cuda/2d/util.h b/include/astra/cuda/2d/util.h index cacf0a9..2cf0639 100644 --- a/include/astra/cuda/2d/util.h +++ b/include/astra/cuda/2d/util.h @@ -66,7 +66,7 @@ bool zeroProjectionData(float* D_ptr, unsigned int pitch, const SDimensions& dim void duplicateVolumeData(float* D_dst, float* D_src, unsigned int pitch, const SDimensions& dims); void duplicateProjectionData(float* D_dst, float* D_src, unsigned int pitch, const SDimensions& dims); -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); bool createTextureObjectPitch2D(float* data, cudaTextureObject_t& texObj, unsigned int pitch, unsigned int width, unsigned int height, cudaTextureAddressMode mode = cudaAddressModeBorder); |