From: Bruno Turcksin Date: Mon, 29 Aug 2016 15:21:50 +0000 (-0400) Subject: Rename CudaAssert to AssertCuda and improves its implementation. X-Git-Tag: v8.5.0-rc1~695^2 X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=refs%2Fpull%2F3005%2Fhead;p=dealii.git Rename CudaAssert to AssertCuda and improves its implementation. --- diff --git a/include/deal.II/base/exceptions.h b/include/deal.II/base/exceptions.h index 5eca2b25d8..92e4017549 100644 --- a/include/deal.II/base/exceptions.h +++ b/include/deal.II/base/exceptions.h @@ -1065,7 +1065,17 @@ namespace StandardExceptions "find a valid muparser library on your system and also did " "not choose the one that comes bundled with deal.II."); - +#ifdef DEAL_II_WITH_CUDA + /** + * This exception is raised if an error happened in a CUDA kernel. + * + * The constructor takes a single char*, the output of + * cudaGetErrorString. + */ + DeclException1 (ExcCudaError, + char *, + << arg1); +#endif //@} } /*namespace StandardExceptions*/ @@ -1142,15 +1152,10 @@ namespace StandardExceptions * @ingroup Exceptions * @author Bruno Turcksin, 2016 */ -#define CudaAssert(error_code) \ - { \ - if (error_code != cudaSuccess) \ - { \ - fprintf(stderr,"Error in %s (%d): %s\n",__FILE__, \ - __LINE__,cudaGetErrorString(error_code)); \ - exit(1); \ - } \ - } +// For now use AssertThrow instead of Assert because macros are not passed +// correctly to nvcc. +#define AssertCuda(error_code) AssertThrow(error_code == cudaSuccess, \ + dealii::ExcCudaError(cudaGetErrorString(error_code))) #endif using namespace StandardExceptions; diff --git a/include/deal.II/lac/read_write_vector.templates.h b/include/deal.II/lac/read_write_vector.templates.h index 8496f35c68..4220502e1b 100644 --- a/include/deal.II/lac/read_write_vector.templates.h +++ b/include/deal.II/lac/read_write_vector.templates.h @@ -393,7 +393,7 @@ namespace LinearAlgebra cudaError_t error_code = cudaMemcpy(&val[0], cuda_vec.get_values(), n_elements*sizeof(Number), cudaMemcpyDeviceToHost); - CudaAssert(error_code); + AssertCuda(error_code); } else { @@ -402,7 +402,7 @@ namespace LinearAlgebra cudaError_t error_code = cudaMemcpy(&tmp[0], cuda_vec.get_values(), n_elements*sizeof(Number), cudaMemcpyDeviceToHost); - CudaAssert(error_code); + AssertCuda(error_code); // Add the two vectors for (unsigned int i=0; i <<>>(val, tmp, n_elements); // Check that the kernel was launched correctly - CudaAssert(cudaGetLastError()); + AssertCuda(cudaGetLastError()); // Check that there was no problem during the execution of the kernel - CudaAssert(cudaDeviceSynchronize()); + AssertCuda(cudaDeviceSynchronize()); // Delete the temporary vector error_code = cudaFree(tmp); - CudaAssert(error_code); + AssertCuda(error_code); } } @@ -587,9 +587,9 @@ namespace LinearAlgebra factor, n_elements); // Check that the kernel was launched correctly - CudaAssert(cudaGetLastError()); + AssertCuda(cudaGetLastError()); // Check that there was no problem during the execution of the kernel - CudaAssert(cudaDeviceSynchronize()); + AssertCuda(cudaDeviceSynchronize()); return *this; } @@ -606,9 +606,9 @@ namespace LinearAlgebra 1./factor, n_elements); // Check that the kernel was launched correctly - CudaAssert(cudaGetLastError()); + AssertCuda(cudaGetLastError()); // Check that there was no problem during the execution of the kernel - CudaAssert(cudaDeviceSynchronize()); + AssertCuda(cudaDeviceSynchronize()); return *this; } @@ -633,9 +633,9 @@ namespace LinearAlgebra <<>>(val, down_V.val, n_elements); // Check that the kernel was launched correctly - CudaAssert(cudaGetLastError()); + AssertCuda(cudaGetLastError()); // Check that there was no problem during the execution of the kernel - CudaAssert(cudaDeviceSynchronize()); + AssertCuda(cudaDeviceSynchronize()); return *this; } @@ -660,9 +660,9 @@ namespace LinearAlgebra <<>>(val, down_V.val, n_elements); // Check that the kernel was launched correctly - CudaAssert(cudaGetLastError()); + AssertCuda(cudaGetLastError()); // Check that there was no problem during the execution of the kernel - CudaAssert(cudaDeviceSynchronize()); + AssertCuda(cudaDeviceSynchronize()); return *this; } @@ -683,7 +683,7 @@ namespace LinearAlgebra Number *result_device; cudaError_t error_code = cudaMalloc(&result_device, n_elements*sizeof(Number)); - CudaAssert(error_code); + AssertCuda(error_code); error_code = cudaMemset(result_device, Number(), sizeof(Number)); const int n_blocks = 1 + (n_elements-1)/(CHUNK_SIZE*BLOCK_SIZE); @@ -696,10 +696,10 @@ namespace LinearAlgebra Number result; error_code = cudaMemcpy(&result, result_device, sizeof(Number), cudaMemcpyDeviceToHost); - CudaAssert(error_code); + AssertCuda(error_code); // Free the memory on the device error_code = cudaFree(result_device); - CudaAssert(error_code); + AssertCuda(error_code); return result; } @@ -715,9 +715,9 @@ namespace LinearAlgebra n_elements); // Check that the kernel was launched correctly - CudaAssert(cudaGetLastError()); + AssertCuda(cudaGetLastError()); // Check that there was no problem during the execution of the kernel - CudaAssert(cudaDeviceSynchronize()); + AssertCuda(cudaDeviceSynchronize()); } @@ -741,9 +741,9 @@ namespace LinearAlgebra a, down_V.val, n_elements); // Check that the kernel was launched correctly - CudaAssert(cudaGetLastError()); + AssertCuda(cudaGetLastError()); // Check that there was no problem during the execution of the kernel - CudaAssert(cudaDeviceSynchronize()); + AssertCuda(cudaDeviceSynchronize()); } @@ -778,9 +778,9 @@ namespace LinearAlgebra a, down_V.val, b, down_W.val, n_elements); // Check that the kernel was launched correctly - CudaAssert(cudaGetLastError()); + AssertCuda(cudaGetLastError()); // Check that there was no problem during the execution of the kernel - CudaAssert(cudaDeviceSynchronize()); + AssertCuda(cudaDeviceSynchronize()); } @@ -806,9 +806,9 @@ namespace LinearAlgebra a, down_V.val, n_elements); // Check that the kernel was launched correctly - CudaAssert(cudaGetLastError()); + AssertCuda(cudaGetLastError()); // Check that there was no problem during the execution of the kernel - CudaAssert(cudaDeviceSynchronize()); + AssertCuda(cudaDeviceSynchronize()); } @@ -831,9 +831,9 @@ namespace LinearAlgebra down_scaling_factors.val, n_elements); // Check that the kernel was launched correctly - CudaAssert(cudaGetLastError()); + AssertCuda(cudaGetLastError()); // Check that there was no problem during the execution of the kernel - CudaAssert(cudaDeviceSynchronize()); + AssertCuda(cudaDeviceSynchronize()); } @@ -857,9 +857,9 @@ namespace LinearAlgebra down_V.val, n_elements); // Check that the kernel was launched correctly - CudaAssert(cudaGetLastError()); + AssertCuda(cudaGetLastError()); // Check that there was no problem during the execution of the kernel - CudaAssert(cudaDeviceSynchronize()); + AssertCuda(cudaDeviceSynchronize()); } @@ -869,7 +869,7 @@ namespace LinearAlgebra { Number *result_device; cudaError_t error_code = cudaMalloc(&result_device, sizeof(Number)); - CudaAssert(error_code); + AssertCuda(error_code); error_code = cudaMemset(result_device, Number(), sizeof(Number)); const int n_blocks = 1 + (n_elements-1)/(CHUNK_SIZE*BLOCK_SIZE); @@ -882,10 +882,10 @@ namespace LinearAlgebra Number result; error_code = cudaMemcpy(&result, result_device, sizeof(Number), cudaMemcpyDeviceToHost); - CudaAssert(error_code); + AssertCuda(error_code); // Free the memory on the device error_code = cudaFree(result_device); - CudaAssert(error_code); + AssertCuda(error_code); return result; } @@ -905,7 +905,7 @@ namespace LinearAlgebra { Number *result_device; cudaError_t error_code = cudaMalloc(&result_device, sizeof(Number)); - CudaAssert(error_code); + AssertCuda(error_code); error_code = cudaMemset(result_device, Number(), sizeof(Number)); const int n_blocks = 1 + (n_elements-1)/(CHUNK_SIZE*BLOCK_SIZE); @@ -918,10 +918,10 @@ namespace LinearAlgebra Number result; error_code = cudaMemcpy(&result, result_device, sizeof(Number), cudaMemcpyDeviceToHost); - CudaAssert(error_code); + AssertCuda(error_code); // Free the memory on the device error_code = cudaFree(result_device); - CudaAssert(error_code); + AssertCuda(error_code); return result; } @@ -951,9 +951,9 @@ namespace LinearAlgebra Number *res_d; cudaError_t error_code = cudaMalloc(&res_d, sizeof(Number)); - CudaAssert(error_code); + AssertCuda(error_code); error_code = cudaMemset(res_d, 0., sizeof(Number)); - CudaAssert(error_code); + AssertCuda(error_code); const int n_blocks = 1 + (n_elements-1)/(CHUNK_SIZE*BLOCK_SIZE); internal::add_and_dot <<>>( @@ -961,7 +961,7 @@ namespace LinearAlgebra Number res; error_code = cudaMemcpy(&res, res_d, sizeof(Number), cudaMemcpyDeviceToHost); - CudaAssert(error_code); + AssertCuda(error_code); error_code = cudaFree(res_d); return res; @@ -994,7 +994,7 @@ namespace LinearAlgebra cudaError_t error_code = cudaMemcpy(cpu_val, val, n_elements*sizeof(Number), cudaMemcpyHostToDevice); - CudaAssert(error_code); + AssertCuda(error_code); for (unsigned int i=0; i