From: Bruno Turcksin <bruno.turcksin@gmail.com>
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=a5350c8c361362ed65a9bafe0cfcbff351455800;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 <tt>char*</tt>, 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<n_elements; ++i)
diff --git a/source/lac/cuda_vector.cu b/source/lac/cuda_vector.cu
index a5b74c84ea..1e3b295edf 100644
--- a/source/lac/cuda_vector.cu
+++ b/source/lac/cuda_vector.cu
@@ -463,11 +463,11 @@ namespace LinearAlgebra
     {
       // Allocate the memory
       cudaError_t error_code = cudaMalloc(&val, n_elements*sizeof(Number));
-      CudaAssert(error_code);
+      AssertCuda(error_code);
       // Copy the values.
       error_code = cudaMemcpy(val, V.val,n_elements*sizeof(Number),
                               cudaMemcpyDeviceToDevice);
-      CudaAssert(error_code);
+      AssertCuda(error_code);
     }
 
 
@@ -479,7 +479,7 @@ namespace LinearAlgebra
     {
       // Allocate the memory
       cudaError_t error_code = cudaMalloc(&val, n_elements*sizeof(Number));
-      CudaAssert(error_code);
+      AssertCuda(error_code);
     }
 
 
@@ -490,7 +490,7 @@ namespace LinearAlgebra
       if (val != nullptr)
         {
           cudaError_t error_code = cudaFree(val);
-          CudaAssert(error_code);
+          AssertCuda(error_code);
           val = nullptr;
           n_elements = 0;
         }
@@ -508,7 +508,7 @@ namespace LinearAlgebra
           if (val != nullptr)
             {
               cudaError_t error_code = cudaFree(val);
-              CudaAssert(error_code);
+              AssertCuda(error_code);
               val = nullptr;
             }
         }
@@ -517,18 +517,18 @@ namespace LinearAlgebra
           if (n_elements != n)
             {
               cudaError_t error_code = cudaFree(val);
-              CudaAssert(error_code);
+              AssertCuda(error_code);
             }
 
           cudaError_t error_code = cudaMalloc(&val, n*sizeof(Number));
-          CudaAssert(error_code);
+          AssertCuda(error_code);
 
           // If necessary set the elements to zero
           if (omit_zeroing_entries == false)
             {
               cudaError_t error_code = cudaMemset(val, 0,
                                                   n_elements*sizeof(Number));
-              CudaAssert(error_code);
+              AssertCuda(error_code);
             }
         }
       n_elements = n;
@@ -546,19 +546,19 @@ namespace LinearAlgebra
           cudaError_t error_code = cudaMemcpy(val, V.begin(),
                                               n_elements*sizeof(Number),
                                               cudaMemcpyHostToDevice);
-          CudaAssert(error_code);
+          AssertCuda(error_code);
         }
       else
         {
           // Create a temporary vector on the device
           Number *tmp;
           cudaError_t error_code = cudaMalloc(&tmp, n_elements*sizeof(Number));
-          CudaAssert(error_code);
+          AssertCuda(error_code);
 
           // Copy the vector from the host to the temporary vector on the device
           error_code = cudaMemcpy(&tmp[0], V.begin(), n_elements*sizeof(Number),
                                   cudaMemcpyHostToDevice);
-          CudaAssert(error_code);
+          AssertCuda(error_code);
 
           // Add the two vectors
           const int n_blocks = 1 + (n_elements-1)/(CHUNK_SIZE*BLOCK_SIZE);
@@ -566,13 +566,13 @@ namespace LinearAlgebra
           internal::vector_bin_op<Number,internal::Binop_Addition>
           <<<n_blocks,BLOCK_SIZE>>>(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
       <<<n_blocks,BLOCK_SIZE>>>(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
       <<<n_blocks,BLOCK_SIZE>>>(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<Number> <<<dim3(n_blocks,1),dim3(BLOCK_SIZE)>>>(
@@ -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<n_elements; ++i)
         out << cpu_val[i] << std::endl;
       out << std::flush;