From 0aa807d24ecb11daac19fc35e9ab22300c83a3a5 Mon Sep 17 00:00:00 2001 From: Josh Romero Date: Mon, 6 Jul 2026 11:52:12 -0700 Subject: [PATCH] Add some missing CUDA launch checks to benchmark and examples. Signed-off-by: Josh Romero --- benchmark/benchmark.cu | 2 ++ examples/cc/basic_usage/basic_usage.cu | 10 ++++++++++ examples/cc/basic_usage/basic_usage_autotune.cu | 10 ++++++++++ 3 files changed, 22 insertions(+) diff --git a/benchmark/benchmark.cu b/benchmark/benchmark.cu index 51d22d7..77d4559 100644 --- a/benchmark/benchmark.cu +++ b/benchmark/benchmark.cu @@ -593,10 +593,12 @@ int main(int argc, char** argv) { // Note: excluding scaling from timing #ifdef R2C scale<<<(pinfo_x_r.size + 1024 - 1) / 1024, 1024>>>(output_r, 1.0 / fftsize, pinfo_x_r); + CHECK_CUDA_LAUNCH_EXIT(); if (out_of_place) std::swap(input, output); if (out_of_place) std::swap(input_r, output_r); #else scale<<<(pinfo_x_c.size + 1024 - 1) / 1024, 1024>>>(output, 1.0 / fftsize, pinfo_x_c); + CHECK_CUDA_LAUNCH_EXIT(); if (out_of_place) std::swap(input, output); #endif diff --git a/examples/cc/basic_usage/basic_usage.cu b/examples/cc/basic_usage/basic_usage.cu index 63249d9..d84fe8a 100644 --- a/examples/cc/basic_usage/basic_usage.cu +++ b/examples/cc/basic_usage/basic_usage.cu @@ -48,6 +48,15 @@ } \ } while (false) +#define CHECK_CUDA_LAUNCH_EXIT() \ + do { \ + cudaError_t err = cudaGetLastError(); \ + if (cudaSuccess != err) { \ + fprintf(stderr, "%s:%d CUDA error. (%s)\n", __FILE__, __LINE__, cudaGetErrorString(err)); \ + exit(EXIT_FAILURE); \ + } \ + } while (false) + #define CHECK_MPI_EXIT(call) \ { \ int err = call; \ @@ -223,6 +232,7 @@ int main(int argc, char** argv) { int threads_per_block = 256; int nblocks = (pinfo_x.size + threads_per_block - 1) / threads_per_block; initialize_pencil<<>>(data_d, pinfo_x); + CHECK_CUDA_LAUNCH_EXIT(); // Allocating cuDecomp workspace diff --git a/examples/cc/basic_usage/basic_usage_autotune.cu b/examples/cc/basic_usage/basic_usage_autotune.cu index 3f76d14..1da7f75 100644 --- a/examples/cc/basic_usage/basic_usage_autotune.cu +++ b/examples/cc/basic_usage/basic_usage_autotune.cu @@ -48,6 +48,15 @@ } \ } while (false) +#define CHECK_CUDA_LAUNCH_EXIT() \ + do { \ + cudaError_t err = cudaGetLastError(); \ + if (cudaSuccess != err) { \ + fprintf(stderr, "%s:%d CUDA error. (%s)\n", __FILE__, __LINE__, cudaGetErrorString(err)); \ + exit(EXIT_FAILURE); \ + } \ + } while (false) + #define CHECK_MPI_EXIT(call) \ { \ int err = call; \ @@ -207,6 +216,7 @@ int main(int argc, char** argv) { int threads_per_block = 256; int nblocks = (pinfo_x.size + threads_per_block - 1) / threads_per_block; initialize_pencil<<>>(data_d, pinfo_x); + CHECK_CUDA_LAUNCH_EXIT(); // Allocating cuDecomp workspace