diff options
| -rw-r--r-- | Core/regularisers_GPU/Diffus_4thO_GPU_core.cu | 21 | ||||
| -rw-r--r-- | Core/regularisers_GPU/Diffus_4thO_GPU_core.h | 2 | ||||
| -rw-r--r-- | Wrappers/Python/src/gpu_regularisers.pyx | 8 |
3 files changed, 18 insertions, 13 deletions
diff --git a/Core/regularisers_GPU/Diffus_4thO_GPU_core.cu b/Core/regularisers_GPU/Diffus_4thO_GPU_core.cu index fd586ef..f453377 100644 --- a/Core/regularisers_GPU/Diffus_4thO_GPU_core.cu +++ b/Core/regularisers_GPU/Diffus_4thO_GPU_core.cu @@ -39,13 +39,15 @@ limitations under the License. #define CHECK(call) \ { \ const cudaError_t error = call; \ + const int ERR=1; \ + const int OK=0; \ if (error != cudaSuccess) \ { \ fprintf(stderr, "Error: %s:%d, ", __FILE__, __LINE__); \ fprintf(stderr, "code: %d, reason: %s\n", error, \ cudaGetErrorString(error)); \ - exit(1); \ - } \ + ERR; \ + } else {OK;} \ } #define BLKXSIZE 8 @@ -228,21 +230,21 @@ __global__ void Diffusion_update_step3D_kernel(float *Output, float *Input, floa /*%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%*/ /********************* MAIN HOST FUNCTION ******************/ /*%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%%*/ -extern "C" void Diffus4th_GPU_main(float *Input, float *Output, float lambdaPar, float sigmaPar, int iterationsNumb, float tau, int N, int M, int Z) +extern "C" int Diffus4th_GPU_main(float *Input, float *Output, float lambdaPar, float sigmaPar, int iterationsNumb, float tau, int N, int M, int Z) { int dimTotal, dev = 0; - CHECK(cudaSetDevice(dev)); + if (CHECK(cudaSetDevice(dev))>0) return 1; float *d_input, *d_output, *d_W_Lapl; float sigmaPar2; sigmaPar2 = sigmaPar*sigmaPar; dimTotal = N*M*Z; - CHECK(cudaMalloc((void**)&d_input,dimTotal*sizeof(float))); - CHECK(cudaMalloc((void**)&d_output,dimTotal*sizeof(float))); - CHECK(cudaMalloc((void**)&d_W_Lapl,dimTotal*sizeof(float))); + if (CHECK(cudaMalloc((void**)&d_input,dimTotal*sizeof(float)))>0) return 1; + if (CHECK(cudaMalloc((void**)&d_output,dimTotal*sizeof(float)))>0) return 1; + if (CHECK(cudaMalloc((void**)&d_W_Lapl,dimTotal*sizeof(float)))>0) return 1; - CHECK(cudaMemcpy(d_input,Input,dimTotal*sizeof(float),cudaMemcpyHostToDevice)); - CHECK(cudaMemcpy(d_output,Input,dimTotal*sizeof(float),cudaMemcpyHostToDevice)); + if (CHECK(cudaMemcpy(d_input,Input,dimTotal*sizeof(float),cudaMemcpyHostToDevice))>0) return 1; + if (CHECK(cudaMemcpy(d_output,Input,dimTotal*sizeof(float),cudaMemcpyHostToDevice))>0) return 1; if (Z == 1) { /*2D case */ @@ -275,4 +277,5 @@ extern "C" void Diffus4th_GPU_main(float *Input, float *Output, float lambdaPar, CHECK(cudaFree(d_input)); CHECK(cudaFree(d_output)); CHECK(cudaFree(d_W_Lapl)); + return 0; } diff --git a/Core/regularisers_GPU/Diffus_4thO_GPU_core.h b/Core/regularisers_GPU/Diffus_4thO_GPU_core.h index 6314c1f..77d5d79 100644 --- a/Core/regularisers_GPU/Diffus_4thO_GPU_core.h +++ b/Core/regularisers_GPU/Diffus_4thO_GPU_core.h @@ -3,6 +3,6 @@ #include "CCPiDefines.h" #include <stdio.h> -extern "C" CCPI_EXPORT void Diffus4th_GPU_main(float *Input, float *Output, float lambdaPar, float sigmaPar, int iterationsNumb, float tau, int N, int M, int Z); +extern "C" CCPI_EXPORT int Diffus4th_GPU_main(float *Input, float *Output, float lambdaPar, float sigmaPar, int iterationsNumb, float tau, int N, int M, int Z); #endif diff --git a/Wrappers/Python/src/gpu_regularisers.pyx b/Wrappers/Python/src/gpu_regularisers.pyx index 302727e..3457796 100644 --- a/Wrappers/Python/src/gpu_regularisers.pyx +++ b/Wrappers/Python/src/gpu_regularisers.pyx @@ -25,7 +25,7 @@ cdef extern void TGV_GPU_main(float *Input, float *Output, float lambdaPar, floa cdef extern void LLT_ROF_GPU_main(float *Input, float *Output, float lambdaROF, float lambdaLLT, int iterationsNumb, float tau, int N, int M, int Z); cdef extern void NonlDiff_GPU_main(float *Input, float *Output, float lambdaPar, float sigmaPar, int iterationsNumb, float tau, int penaltytype, int N, int M, int Z); cdef extern void dTV_FGP_GPU_main(float *Input, float *InputRef, float *Output, float lambdaPar, int iterationsNumb, float epsil, float eta, int methodTV, int nonneg, int printM, int N, int M, int Z); -cdef extern void Diffus4th_GPU_main(float *Input, float *Output, float lambdaPar, float sigmaPar, int iterationsNumb, float tau, int N, int M, int Z); +cdef extern int Diffus4th_GPU_main(float *Input, float *Output, float lambdaPar, float sigmaPar, int iterationsNumb, float tau, int N, int M, int Z); cdef extern void PatchSelect_GPU_main(float *Input, unsigned short *H_i, unsigned short *H_j, float *Weights, int N, int M, int SearchWindow, int SimilarWin, int NumNeighb, float h); # Total-variation Rudin-Osher-Fatemi (ROF) @@ -522,7 +522,8 @@ def Diff4th_2D(np.ndarray[np.float32_t, ndim=2, mode="c"] inputData, # Run Anisotropic Fourth-Order diffusion for 2D data # Running CUDA code here - Diffus4th_GPU_main(&inputData[0,0], &outputData[0,0], regularisation_parameter, edge_parameter, iterationsNumb, time_marching_parameter, dims[1], dims[0], 1) + if (Diffus4th_GPU_main(&inputData[0,0], &outputData[0,0], regularisation_parameter, edge_parameter, iterationsNumb, time_marching_parameter, dims[1], dims[0], 1)>0): + raise RuntimeError('Runtime error!') return outputData def Diff4th_3D(np.ndarray[np.float32_t, ndim=3, mode="c"] inputData, @@ -540,7 +541,8 @@ def Diff4th_3D(np.ndarray[np.float32_t, ndim=3, mode="c"] inputData, # Run Anisotropic Fourth-Order diffusion for 3D data # Running CUDA code here - Diffus4th_GPU_main(&inputData[0,0,0], &outputData[0,0,0], regularisation_parameter, edge_parameter, iterationsNumb, time_marching_parameter, dims[2], dims[1], dims[0]) + if (Diffus4th_GPU_main(&inputData[0,0,0], &outputData[0,0,0], regularisation_parameter, edge_parameter, iterationsNumb, time_marching_parameter, dims[2], dims[1], dims[0])>0): + RuntimeError('Runtime error!') return outputData #****************************************************************# |