Create and use NVBENCH_CUDA_CALL_RESET_ERROR.

NVIDIA · May 29, 2024 · b333e5c · b333e5c
1 parent 088c9ee
commit b333e5c
Show file tree

Hide file tree

Showing 6 changed files with 48 additions and 3 deletions.
diff --git a/nvbench/cuda_call.cuh b/nvbench/cuda_call.cuh
@@ -57,6 +57,21 @@
     }                                                                                              \
   } 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(cudaStreamSynchronize(m_launch.get_stream()));
+    NVBENCH_CUDA_CALL_RESET_ERROR(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(cudaStreamSynchronize(m_launch.get_stream()));
+    NVBENCH_CUDA_CALL_RESET_ERROR(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(cudaStreamSynchronize(m_launch.get_stream()));
+    NVBENCH_CUDA_CALL_RESET_ERROR(cudaStreamSynchronize(m_launch.get_stream()));
   }
 
   KernelLauncher &m_kernel_launcher;

diff --git a/testing/CMakeLists.txt b/testing/CMakeLists.txt
@@ -16,6 +16,7 @@ set(test_srcs
   named_values.cu
   option_parser.cu
   range.cu
+  reset_error.cu
   ring_buffer.cu
   runner.cu
   state.cu

diff --git a/testing/reset_error.cu b/testing/reset_error.cu
@@ -0,0 +1,29 @@
+#include <nvbench/cuda_call.cuh>
+
+#include "test_asserts.cuh"
+
+namespace
+{
+    __global__ void multiply5(const int32_t* __restrict__ a, int32_t* __restrict__ b)
+    {
+      const auto id = blockIdx.x * blockDim.x + threadIdx.x;
+      b[id] = 5 * a[id];
+    }
+}
+
+int main()
+{ 
+  multiply5<<<256, 256>>>(nullptr, nullptr);
+
+  try
+  {
+    NVBENCH_CUDA_CALL_RESET_ERROR(cudaStreamSynchronize(0));
+    return 1;
+  }
+  catch (const std::runtime_error &)
+  {
+    ASSERT(cudaGetLastError() == cudaError_t::cudaSuccess);
+  }
+
+  return 0;
+}