Fix memleak in error handling

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);