Moved cudaGetLastError() call to NVBENCH_CUDA_CALL macro

NVIDIA · May 30, 2024 · 9f8818b · 9f8818b
1 parent b333e5c
commit 9f8818b
Show file tree

Hide file tree

Showing 5 changed files with 9 additions and 21 deletions.
diff --git a/nvbench/cuda_call.cuh b/nvbench/cuda_call.cuh
@@ -24,12 +24,14 @@
 #include <string>
 
 /// Throws a std::runtime_error if `call` doesn't return `cudaSuccess`.
+/// Resets the error with cudaGetLastError().
 #define NVBENCH_CUDA_CALL(call)                                                                    \
   do                                                                                               \
   {                                                                                                \
     const cudaError_t nvbench_cuda_call_error = call;                                              \
     if (nvbench_cuda_call_error != cudaSuccess)                                                    \
     {                                                                                              \
+      cudaGetLastError();                                                                          \
       nvbench::cuda_call::throw_error(__FILE__, __LINE__, #call, nvbench_cuda_call_error);         \
     }                                                                                              \
   } while (false)
@@ -57,21 +59,6 @@
     }                                                                                              \
   } while (false)
 
-/// Throws a std::runtime_error if `call` doesn't return `cudaSuccess`.
-/// Resets the error with cudaGetLastError().
-#define NVBENCH_CUDA_CALL_RESET_ERROR(call)                                                        \
-  do                                                                                               \
-  {                                                                                                \
-    const cudaError_t nvbench_cuda_call_error = call;                                              \
-    if (nvbench_cuda_call_error != cudaSuccess)                                                    \
-    {                                                                                              \
-      cudaGetLastError();                                                                          \
-      nvbench::cuda_call::throw_error(__FILE__, __LINE__, #call, nvbench_cuda_call_error);         \
-    }                                                                                              \
-  }                                                                                                \
-  while (false)
-
-
 namespace nvbench::cuda_call
 {
 

diff --git a/nvbench/detail/measure_cold.cuh b/nvbench/detail/measure_cold.cuh
@@ -71,7 +71,7 @@ protected:
 
   __forceinline__ void sync_stream() const
   {
-    NVBENCH_CUDA_CALL_RESET_ERROR(cudaStreamSynchronize(m_launch.get_stream()));
+    NVBENCH_CUDA_CALL(cudaStreamSynchronize(m_launch.get_stream()));
   }
 
   void block_stream();

diff --git a/nvbench/detail/measure_cupti.cuh b/nvbench/detail/measure_cupti.cuh
@@ -65,7 +65,7 @@ protected:
 
   __forceinline__ void sync_stream() const
   {
-    NVBENCH_CUDA_CALL_RESET_ERROR(cudaStreamSynchronize(m_launch.get_stream()));
+    NVBENCH_CUDA_CALL(cudaStreamSynchronize(m_launch.get_stream()));
   }
 
   nvbench::state &m_state;

diff --git a/nvbench/detail/measure_hot.cuh b/nvbench/detail/measure_hot.cuh
@@ -204,7 +204,7 @@ private:
 
   __forceinline__ void sync_stream() const
   {
-    NVBENCH_CUDA_CALL_RESET_ERROR(cudaStreamSynchronize(m_launch.get_stream()));
+    NVBENCH_CUDA_CALL(cudaStreamSynchronize(m_launch.get_stream()));
   }
 
   KernelLauncher &m_kernel_launcher;

diff --git a/testing/reset_error.cu b/testing/reset_error.cu
@@ -2,6 +2,7 @@
 
 #include "test_asserts.cuh"
 
+
 namespace
 {
     __global__ void multiply5(const int32_t* __restrict__ a, int32_t* __restrict__ b)
@@ -17,13 +18,13 @@ int main()
 
   try
   {
-    NVBENCH_CUDA_CALL_RESET_ERROR(cudaStreamSynchronize(0));
-    return 1;
+    NVBENCH_CUDA_CALL(cudaStreamSynchronize(0));
+    ASSERT(false);
   }
   catch (const std::runtime_error &)
   {
     ASSERT(cudaGetLastError() == cudaError_t::cudaSuccess);
   }
 
   return 0;
-}
+}