From 9f8818b419ce1938862e4a9b9f549c05aaef749c Mon Sep 17 00:00:00 2001 From: Sergey Pavlov Date: Thu, 30 May 2024 13:42:01 +0400 Subject: [PATCH] Moved cudaGetLastError() call to NVBENCH_CUDA_CALL macro --- nvbench/cuda_call.cuh | 17 ++--------------- nvbench/detail/measure_cold.cuh | 2 +- nvbench/detail/measure_cupti.cuh | 2 +- nvbench/detail/measure_hot.cuh | 2 +- testing/reset_error.cu | 7 ++++--- 5 files changed, 9 insertions(+), 21 deletions(-) diff --git a/nvbench/cuda_call.cuh b/nvbench/cuda_call.cuh index 737ca1c2..5b2ae362 100644 --- a/nvbench/cuda_call.cuh +++ b/nvbench/cuda_call.cuh @@ -24,12 +24,14 @@ #include /// 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 index ecbcfb3b..2b0183f5 100644 --- 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 index bf193469..ec7b2120 100644 --- 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 index 60ce9761..c9175830 100644 --- 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 index e3a28ef8..8fece930 100644 --- 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,8 +18,8 @@ int main() try { - NVBENCH_CUDA_CALL_RESET_ERROR(cudaStreamSynchronize(0)); - return 1; + NVBENCH_CUDA_CALL(cudaStreamSynchronize(0)); + ASSERT(false); } catch (const std::runtime_error &) { @@ -26,4 +27,4 @@ int main() } return 0; -} \ No newline at end of file +}