From d11755ca1b1e18a41750d649f50a0c380bbeb1fd Mon Sep 17 00:00:00 2001 From: David Bayer Date: Tue, 12 May 2026 19:23:23 +0200 Subject: [PATCH 01/10] [c2h] Make catch macros work on device --- c2h/CMakeLists.txt | 1 + c2h/include/c2h/catch2_test_helper.h | 9 +- c2h/include/c2h/catch2_test_macros.h | 145 ++++++++++++++++++ cub/test/insert_nested_NVTX_range_guard.h | 2 +- cudax/test/common/testing.cuh | 51 +----- cudax/test/utility/optionally_static.cu | 2 +- .../libcudacxx/cuda/ccclrt/common/testing.cuh | 43 ++---- 7 files changed, 168 insertions(+), 85 deletions(-) create mode 100644 c2h/include/c2h/catch2_test_macros.h diff --git a/c2h/CMakeLists.txt b/c2h/CMakeLists.txt index 6e0367b1e36..0f9a437a6fd 100644 --- a/c2h/CMakeLists.txt +++ b/c2h/CMakeLists.txt @@ -36,6 +36,7 @@ add_library( generators_uniform_offsets.cu generators_vector.cu ) +target_compile_definitions(cccl.c2h PUBLIC CATCH_CONFIG_PREFIX_ALL) target_include_directories(cccl.c2h PUBLIC "${C2H_SOURCE_DIR}/include") target_link_libraries( cccl.c2h diff --git a/c2h/include/c2h/catch2_test_helper.h b/c2h/include/c2h/catch2_test_helper.h index ff62c885970..a32a54106d4 100644 --- a/c2h/include/c2h/catch2_test_helper.h +++ b/c2h/include/c2h/catch2_test_helper.h @@ -19,6 +19,7 @@ #include #include +#include #include #include #include @@ -595,26 +596,26 @@ class nvtx_fixture #define C2H_TEST_IMPL(ID, NAME, TAG, ...) \ using C2H_TEST_CONCAT(types_, ID) = c2h::cartesian_product<__VA_ARGS__>; \ - TEMPLATE_LIST_TEST_CASE_METHOD(::detail::nvtx_fixture, C2H_TEST_NAME(NAME), TAG, C2H_TEST_CONCAT(types_, ID)) + CATCH_TEMPLATE_LIST_TEST_CASE_METHOD(::detail::nvtx_fixture, C2H_TEST_NAME(NAME), TAG, C2H_TEST_CONCAT(types_, ID)) #define C2H_TEST(NAME, TAG, ...) C2H_TEST_IMPL(__LINE__, NAME, TAG, __VA_ARGS__) #define C2H_TEST_WITH_FIXTURE_IMPL(ID, FIXTURE, NAME, TAG, ...) \ using C2H_TEST_CONCAT(types_, ID) = c2h::cartesian_product<__VA_ARGS__>; \ - TEMPLATE_LIST_TEST_CASE_METHOD(FIXTURE, C2H_TEST_NAME(NAME), TAG, C2H_TEST_CONCAT(types_, ID)) + CATCH_TEMPLATE_LIST_TEST_CASE_METHOD(FIXTURE, C2H_TEST_NAME(NAME), TAG, C2H_TEST_CONCAT(types_, ID)) #define C2H_TEST_WITH_FIXTURE(FIXTURE, NAME, TAG, ...) \ C2H_TEST_WITH_FIXTURE_IMPL(__LINE__, FIXTURE, NAME, TAG, __VA_ARGS__) #define C2H_TEST_LIST_IMPL(ID, NAME, TAG, ...) \ using C2H_TEST_CONCAT(types_, ID) = c2h::type_list<__VA_ARGS__>; \ - TEMPLATE_LIST_TEST_CASE_METHOD(::detail::nvtx_fixture, C2H_TEST_NAME(NAME), TAG, C2H_TEST_CONCAT(types_, ID)) + CATCH_TEMPLATE_LIST_TEST_CASE_METHOD(::detail::nvtx_fixture, C2H_TEST_NAME(NAME), TAG, C2H_TEST_CONCAT(types_, ID)) #define C2H_TEST_LIST(NAME, TAG, ...) C2H_TEST_LIST_IMPL(__LINE__, NAME, TAG, __VA_ARGS__) #define C2H_TEST_LIST_WITH_FIXTURE_IMPL(ID, FIXTURE, NAME, TAG, ...) \ using C2H_TEST_CONCAT(types_, ID) = c2h::type_list<__VA_ARGS__>; \ - TEMPLATE_LIST_TEST_CASE_METHOD(FIXTURE, C2H_TEST_NAME(NAME), TAG, C2H_TEST_CONCAT(types_, ID)) + CATCH_TEMPLATE_LIST_TEST_CASE_METHOD(FIXTURE, C2H_TEST_NAME(NAME), TAG, C2H_TEST_CONCAT(types_, ID)) #define C2H_TEST_LIST_WITH_FIXTURE(FIXTURE, NAME, TAG, ...) \ C2H_TEST_LIST_WITH_FIXTURE_IMPL(__LINE__, FIXTURE, NAME, TAG, __VA_ARGS__) diff --git a/c2h/include/c2h/catch2_test_macros.h b/c2h/include/c2h/catch2_test_macros.h new file mode 100644 index 00000000000..7db47f143a2 --- /dev/null +++ b/c2h/include/c2h/catch2_test_macros.h @@ -0,0 +1,145 @@ +// SPDX-FileCopyrightText: Copyright (c) 2011-2022, NVIDIA CORPORATION. All rights reserved. +// SPDX-License-Identifier: BSD-3-Clause + +#pragma once + +#include + +#include + +#include +#include +#include + +// This file implements Catch2's test macros that work both in host and device code. We globally define the +// CATCH_CONFIG_PREFIX_ALL macro to force Catch2 to prepend it's macros with CATCH_ prefix. That allows us to implement +// the non-prefixed versions ourselves. +// +// In host code, we just use the CATCH_-prefixed variant, in device code we implement the functionality, so it +// corresponds the desired functionality. +// +// Only a subset of the Catch2's macro are provided. If needed, feel free to extend the support. Host-only macros can +// be determined by missing NV_IF_ELSE_TARGET wrapper and immediate dispatch to CATCH_-prefixed variant. + +#define C2H_INTERNAL_DEVICE_PRINTF(FMT, ...) \ + ::printf(__FILE__ ":" _CCCL_TO_STRING(__LINE__) ":\n " FMT "\n block [%u, %u, %u], thread [%u, %u, %u]\n\n", \ + __VA_ARGS__, \ + blockIdx.x, \ + blockIdx.y, \ + blockIdx.z, \ + threadIdx.x, \ + threadIdx.y, \ + threadIdx.z) + +#define C2H_INTERNAL_DEVICE_TEST_PRINT(KIND, ...) C2H_INTERNAL_DEVICE_PRINTF(KIND "(%s) failed", "" #__VA_ARGS__) + +#define REQUIRE(...) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE(__VA_ARGS__);), ({ \ + if (!(__VA_ARGS__)) \ + { \ + C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE", __VA_ARGS__); \ + } \ + ::__trap(); \ + })) +#define REQUIRE_FALSE(...) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_FALSE(__VA_ARGS__);), ({ \ + if (__VA_ARGS__) \ + { \ + C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE_FALSE", __VA_ARGS__); \ + } \ + ::__trap(); \ + })) + +#define REQUIRE_THROWS(...) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_THROWS(__VA_ARGS__);), ({ \ + (void) __VA_ARGS__; \ + C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE_THROWS", __VA_ARGS__); \ + ::__trap(); \ + })) +#define REQUIRE_THROWS_AS(EXPR, TYPE) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_THROWS_AS(EXPR, TYPE);), ({ \ + (void) EXPR; \ + C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE_THROWS_AS", EXPR, TYPE); \ + ::__trap(); \ + })) +#define REQUIRE_NOTHROW(...) NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_NOTHROW(__VA_ARGS__);), (__VA_ARGS__;)) + +#define CHECK(...) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK(__VA_ARGS__);), ({ \ + if (!(__VA_ARGS__)) \ + { \ + C2H_INTERNAL_DEVICE_TEST_PRINT("CHECK", __VA_ARGS__); \ + } \ + ::__trap(); \ + })) +#define CHECK_FALSE(...) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_FALSE(__VA_ARGS__);), ({ \ + if (__VA_ARGS__) \ + { \ + C2H_INTERNAL_DEVICE_TEST_PRINT("CHECK_FALSE", __VA_ARGS__); \ + } \ + ::__trap(); \ + })) +#define CHECKED_IF(...) CATCH_CHECKED_IF(__VA_ARGS__) +#define CHECKED_ELSE(...) CATCH_CHECKED_ELSE(__VA_ARGS__) +#define CHECK_NOFAIL(...) CATCH_CHECK_NOFAIL(__VA_ARGS__) + +#define CHECK_THROWS(...) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_THROWS(__VA_ARGS__);), ({ \ + (void) __VA_ARGS__; \ + C2H_INTERNAL_DEVICE_TEST_PRINT("CHECK_THROWS", __VA_ARGS__); \ + })) +#define CHECK_THROWS_AS(EXPR, TYPE) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_THROWS_AS(EXPR, TYPE);), ({ \ + (void) EXPR; \ + C2H_INTERNAL_DEVICE_TEST_PRINT("CHECK_THROWS_AS", EXPR, TYPE); \ + })) +#define CHECK_NOTHROW(...) NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_NOTHROW(__VA_ARGS__);), (EXPR;)) + +#define TEST_CASE(...) CATCH_TEST_CASE(__VA_ARGS__) +#define TEST_CASE_METHOD(...) CATCH_TEST_CASE_METHOD(__VA_ARGS__) +#define METHOD_AS_TEST_CASE(...) CATCH_METHOD_AS_TEST_CASE(__VA_ARGS__) +#define REGISTER_TEST_CASE(...) CATCH_REGISTER_TEST_CASE(__VA_ARGS__) +#define SECTION(...) CATCH_SECTION(__VA_ARGS__) +#define DYNAMIC_SECTION(...) CATCH_DYNAMIC_SECTION(__VA_ARGS__) +#define FAIL(...) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_FAIL(__VA_ARGS__);), ({ C2H_INTERNAL_DEVICE_PRINTF("FAIL: %s", __VA_ARGS__); })) +#define FAIL_CHECK(...) CATCH_FAIL_CHECK(__VA_ARGS__) +#define SUCCEED(...) CATCH_SUCCEED(__VA_ARGS__) +#define SKIP(...) CATCH_SKIP(__VA_ARGS__) + +#define STATIC_REQUIRE(...) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_STATIC_REQUIRE(__VA_ARGS__);), (static_assert(__VA_ARGS__, #__VA_ARGS__);)) +#define STATIC_REQUIRE_FALSE(...) \ + NV_IF_ELSE_TARGET( \ + NV_IS_HOST, (CATCH_STATIC_REQUIRE_FALSE(__VA_ARGS__);), (static_assert(!(__VA_ARGS__), "!(" #__VA_ARGS__ ")");)) +#define STATIC_CHECK(...) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_STATIC_CHECK(__VA_ARGS__);), (static_assert(__VA_ARGS__, #__VA_ARGS__);)) +#define STATIC_CHECK_FALSE(...) \ + NV_IF_ELSE_TARGET( \ + NV_IS_HOST, (CATCH_STATIC_CHECK_FALSE(__VA_ARGS__);), (static_assert(!(__VA_ARGS__), "!(" #__VA_ARGS__ ")");)) + +#define SCENARIO(...) CATCH_SCENARIO(__VA_ARGS__) +#define SCENARIO_METHOD(...) CATCH_SCENARIO_METHOD(__VA_ARGS__) +#define GIVEN(...) CATCH_GIVEN(__VA_ARGS__) +#define AND_GIVEN(...) CATCH_AND_GIVEN(__VA_ARGS__) +#define WHEN(...) CATCH_WHEN(__VA_ARGS__) +#define AND_WHEN(...) CATCH_AND_WHEN(__VA_ARGS__) +#define THEN(...) CATCH_THEN(__VA_ARGS__) +#define AND_THEN(...) CATCH_AND_THEN(__VA_ARGS__) + +#define INFO(...) CATCH_INFO(__VA_ARGS__) +#define UNSCOPED_INFO(...) CATCH_UNSCOPED_INFO(__VA_ARGS__) +#define WARN(...) CATCH_WARN(__VA_ARGS__) +#define CAPTURE(...) CATCH_CAPTURE(__VA_ARGS__) + +#define TEMPLATE_TEST_CASE(...) CATCH_TEMPLATE_TEST_CASE(__VA_ARGS__) +#define TEMPLATE_TEST_CASE_SIG(...) CATCH_TEMPLATE_TEST_CASE_SIG(__VA_ARGS__) +#define TEMPLATE_TEST_CASE_METHOD(...) CATCH_TEMPLATE_TEST_CASE_METHOD(__VA_ARGS__) +#define TEMPLATE_TEST_CASE_METHOD_SIG(...) CATCH_TEMPLATE_TEST_CASE_METHOD_SIG(__VA_ARGS__) +#define TEMPLATE_PRODUCT_TEST_CASE(...) CATCH_TEMPLATE_PRODUCT_TEST_CASE(__VA_ARGS__) +#define TEMPLATE_PRODUCT_TEST_CASE_SIG(...) CATCH_TEMPLATE_PRODUCT_TEST_CASE_SIG(__VA_ARGS__) +#define TEMPLATE_PRODUCT_TEST_CASE_METHOD(...) CATCH_TEMPLATE_PRODUCT_TEST_CASE_METHOD(__VA_ARGS__) +#define TEMPLATE_PRODUCT_TEST_CASE_METHOD_SIG(...) CATCH_TEMPLATE_PRODUCT_TEST_CASE_METHOD_SIG(__VA_ARGS__) +#define TEMPLATE_LIST_TEST_CASE(...) CATCH_TEMPLATE_LIST_TEST_CASE(__VA_ARGS__) +#define TEMPLATE_LIST_TEST_CASE_METHOD(...) CATCH_TEMPLATE_LIST_TEST_CASE_METHOD(__VA_ARGS__) diff --git a/cub/test/insert_nested_NVTX_range_guard.h b/cub/test/insert_nested_NVTX_range_guard.h index 1337666815b..ff1738610a1 100644 --- a/cub/test/insert_nested_NVTX_range_guard.h +++ b/cub/test/insert_nested_NVTX_range_guard.h @@ -8,7 +8,7 @@ #include #include -#include +#include inline thread_local const char* current_nvtx_range_name = nullptr; diff --git a/cudax/test/common/testing.cuh b/cudax/test/common/testing.cuh index 3d834cf7fee..7e550050fd9 100644 --- a/cudax/test/common/testing.cuh +++ b/cudax/test/common/testing.cuh @@ -32,56 +32,13 @@ namespace cudax_async = cuda::experimental::execution; // NOLINT: misc-unused-al #define CUDART(call) REQUIRE((call) == cudaSuccess) -// Unlike nvcc, when clang parses CUDA, both the host and device sections are present (see -// https://llvm.org/docs/CompileCudaWithLLVM.html#compilation-models). The upshot here is that -// certain calls of CUDAX_REQUIRE() from host code will try to use device calls because -// NV_IS_DEVICE is true. -#if _CCCL_CUDA_COMPILER(CLANG) -__host__ -#endif - __device__ inline void - cudax_require_impl(bool condition, - [[maybe_unused]] const char* condition_text, - [[maybe_unused]] const char* filename, - [[maybe_unused]] unsigned int linenum, - [[maybe_unused]] const char* funcname) -{ - if (!condition) - { -#if !_CCCL_CUDA_COMPILER(CLANG) - // TODO do warp aggregate prints for easier readability? - printf("%s:%u: %s: block: [%d,%d,%d], thread: [%d,%d,%d] Condition `%s` failed.\n", - filename, - linenum, - funcname, - blockIdx.x, - blockIdx.y, - blockIdx.z, - threadIdx.x, - threadIdx.y, - threadIdx.z, - condition_text); -#endif - ::cuda::std::terminate(); - } -} - -#define CUDAX_REQUIRE(condition) \ - NV_IF_ELSE_TARGET(NV_IS_DEVICE, \ - (cudax_require_impl(condition, #condition, __FILE__, __LINE__, __PRETTY_FUNCTION__);), \ - (REQUIRE(condition);)) +#define CUDAX_REQUIRE(condition) REQUIRE(condition) -#define CUDAX_CHECK(condition) \ - NV_IF_ELSE_TARGET(NV_IS_DEVICE, \ - (cudax_require_impl(condition, #condition, __FILE__, __LINE__, __PRETTY_FUNCTION__);), \ - (CHECK(condition);)) +#define CUDAX_CHECK(condition) CHECK(condition) -#define CUDAX_FAIL(message) /* */ \ - NV_IF_ELSE_TARGET(NV_IS_DEVICE, /* */ \ - (cudax_require_impl(false, message, __FILE__, __LINE__, __PRETTY_FUNCTION__);), \ - (FAIL(message);)) +#define CUDAX_FAIL(message) FAIL(message) -#define CUDAX_CHECK_FALSE(condition) CUDAX_CHECK(!(condition)) +#define CUDAX_CHECK_FALSE(condition) CHECK_FALSE(condition) __host__ __device__ constexpr bool operator==(const dim3& lhs, const dim3& rhs) noexcept { diff --git a/cudax/test/utility/optionally_static.cu b/cudax/test/utility/optionally_static.cu index 4a8ea7b789e..e25b33d206b 100644 --- a/cudax/test/utility/optionally_static.cu +++ b/cudax/test/utility/optionally_static.cu @@ -12,7 +12,7 @@ #include -#include +#include namespace cudax = cuda::experimental; diff --git a/libcudacxx/test/libcudacxx/cuda/ccclrt/common/testing.cuh b/libcudacxx/test/libcudacxx/cuda/ccclrt/common/testing.cuh index 1f7f54ab84c..c27433a82c2 100644 --- a/libcudacxx/test/libcudacxx/cuda/ccclrt/common/testing.cuh +++ b/libcudacxx/test/libcudacxx/cuda/ccclrt/common/testing.cuh @@ -26,40 +26,19 @@ #define CUDART(call) REQUIRE((call) == cudaSuccess) -// There is a problem with clang-cuda and nv/target, but we don't need the device side macros yet, -// disable them for now -#if _CCCL_CUDA_COMPILER(CLANG) -# define CCCLRT_REQUIRE(condition) REQUIRE(condition) -# define CCCLRT_CHECK(condition) CHECK(condition) -# define CCCLRT_FAIL(message) FAIL(message) -# define CCCLRT_CHECK_FALSE(condition) CCCLRT_CHECK(!(condition)) - -#else // _CCCL_CUDA_COMPILER(CLANG) -# define CCCLRT_REQUIRE(condition) \ - NV_IF_ELSE_TARGET(NV_IS_DEVICE, \ - (ccclrt_require_impl(condition, #condition, __FILE__, __LINE__, __PRETTY_FUNCTION__);), \ - (REQUIRE(condition);)) - -# define CCCLRT_CHECK(condition) \ - NV_IF_ELSE_TARGET(NV_IS_DEVICE, \ - (ccclrt_require_impl(condition, #condition, __FILE__, __LINE__, __PRETTY_FUNCTION__);), \ - (CHECK(condition);)) - -# define CCCLRT_FAIL(message) /* */ \ - NV_IF_ELSE_TARGET(NV_IS_DEVICE, /* */ \ - (ccclrt_require_impl(false, message, __FILE__, __LINE__, __PRETTY_FUNCTION__);), \ - (FAIL(message);)) - -# define CCCLRT_CHECK_FALSE(condition) CCCLRT_CHECK(!(condition)) -#endif // _CCCL_CUDA_COMPILER(CLANG) +#define CCCLRT_REQUIRE(condition) REQUIRE(condition) + +#define CCCLRT_CHECK(condition) CHECK(condition) + +#define CCCLRT_FAIL(message) FAIL(message) + +#define CCCLRT_CHECK_FALSE(condition) CHECK_FALSE(condition) // Explicit device side require macros for clang-cuda -#define CCCLRT_REQUIRE_DEVICE(condition) \ - ccclrt_require_impl(condition, #condition, __FILE__, __LINE__, __PRETTY_FUNCTION__); -#define CCCLRT_CHECK_DEVICE(condition) \ - ccclrt_require_impl(condition, #condition, __FILE__, __LINE__, __PRETTY_FUNCTION__); -#define CCCLRT_FAIL_DEVICE(message) ccclrt_require_impl(false, message, __FILE__, __LINE__, __PRETTY_FUNCTION__); -#define CCCLRT_CHECK_FALSE_DEVICE(condition) CCCLRT_CHECK_DEVICE(!(condition)) +#define CCCLRT_REQUIRE_DEVICE(condition) REQUIRE(condition) +#define CCCLRT_CHECK_DEVICE(condition) CHECK(condition) +#define CCCLRT_FAIL_DEVICE(message) FAIL(message) +#define CCCLRT_CHECK_FALSE_DEVICE(condition) CHECK_FALSE(condition) TEST_FUNC constexpr bool operator==(const dim3& lhs, const dim3& rhs) noexcept { From 32bfdbf5e00d9a90ca4684857716cc2d410b3989 Mon Sep 17 00:00:00 2001 From: David Bayer Date: Tue, 12 May 2026 19:55:08 +0200 Subject: [PATCH 02/10] update --- c2h/include/c2h/catch2_test_macros.h | 32 +++++++++++++++++++++------- 1 file changed, 24 insertions(+), 8 deletions(-) diff --git a/c2h/include/c2h/catch2_test_macros.h b/c2h/include/c2h/catch2_test_macros.h index 7db47f143a2..01c375065f6 100644 --- a/c2h/include/c2h/catch2_test_macros.h +++ b/c2h/include/c2h/catch2_test_macros.h @@ -10,6 +10,7 @@ #include #include #include +#include // This file implements Catch2's test macros that work both in host and device code. We globally define the // CATCH_CONFIG_PREFIX_ALL macro to force Catch2 to prepend it's macros with CATCH_ prefix. That allows us to implement @@ -33,6 +34,8 @@ #define C2H_INTERNAL_DEVICE_TEST_PRINT(KIND, ...) C2H_INTERNAL_DEVICE_PRINTF(KIND "(%s) failed", "" #__VA_ARGS__) +// + #define REQUIRE(...) \ NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE(__VA_ARGS__);), ({ \ if (!(__VA_ARGS__)) \ @@ -52,13 +55,13 @@ #define REQUIRE_THROWS(...) \ NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_THROWS(__VA_ARGS__);), ({ \ - (void) __VA_ARGS__; \ + __VA_ARGS__; \ C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE_THROWS", __VA_ARGS__); \ ::__trap(); \ })) #define REQUIRE_THROWS_AS(EXPR, TYPE) \ - NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_THROWS_AS(EXPR, TYPE);), ({ \ - (void) EXPR; \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_THROWS_AS((EXPR), TYPE);), ({ \ + EXPR; \ C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE_THROWS_AS", EXPR, TYPE); \ ::__trap(); \ })) @@ -86,15 +89,15 @@ #define CHECK_THROWS(...) \ NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_THROWS(__VA_ARGS__);), ({ \ - (void) __VA_ARGS__; \ + __VA_ARGS__; \ C2H_INTERNAL_DEVICE_TEST_PRINT("CHECK_THROWS", __VA_ARGS__); \ })) #define CHECK_THROWS_AS(EXPR, TYPE) \ - NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_THROWS_AS(EXPR, TYPE);), ({ \ - (void) EXPR; \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_THROWS_AS((EXPR), TYPE);), ({ \ + EXPR; \ C2H_INTERNAL_DEVICE_TEST_PRINT("CHECK_THROWS_AS", EXPR, TYPE); \ })) -#define CHECK_NOTHROW(...) NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_NOTHROW(__VA_ARGS__);), (EXPR;)) +#define CHECK_NOTHROW(...) NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_NOTHROW(__VA_ARGS__);), (__VA_ARGS__;)) #define TEST_CASE(...) CATCH_TEST_CASE(__VA_ARGS__) #define TEST_CASE_METHOD(...) CATCH_TEST_CASE_METHOD(__VA_ARGS__) @@ -103,7 +106,7 @@ #define SECTION(...) CATCH_SECTION(__VA_ARGS__) #define DYNAMIC_SECTION(...) CATCH_DYNAMIC_SECTION(__VA_ARGS__) #define FAIL(...) \ - NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_FAIL(__VA_ARGS__);), ({ C2H_INTERNAL_DEVICE_PRINTF("FAIL: %s", __VA_ARGS__); })) + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_FAIL(__VA_ARGS__);), ({ C2H_INTERNAL_DEVICE_TEST_PRINT("FAIL", __VA_ARGS__); })) #define FAIL_CHECK(...) CATCH_FAIL_CHECK(__VA_ARGS__) #define SUCCEED(...) CATCH_SUCCEED(__VA_ARGS__) #define SKIP(...) CATCH_SKIP(__VA_ARGS__) @@ -128,11 +131,15 @@ #define THEN(...) CATCH_THEN(__VA_ARGS__) #define AND_THEN(...) CATCH_AND_THEN(__VA_ARGS__) +// + #define INFO(...) CATCH_INFO(__VA_ARGS__) #define UNSCOPED_INFO(...) CATCH_UNSCOPED_INFO(__VA_ARGS__) #define WARN(...) CATCH_WARN(__VA_ARGS__) #define CAPTURE(...) CATCH_CAPTURE(__VA_ARGS__) +// + #define TEMPLATE_TEST_CASE(...) CATCH_TEMPLATE_TEST_CASE(__VA_ARGS__) #define TEMPLATE_TEST_CASE_SIG(...) CATCH_TEMPLATE_TEST_CASE_SIG(__VA_ARGS__) #define TEMPLATE_TEST_CASE_METHOD(...) CATCH_TEMPLATE_TEST_CASE_METHOD(__VA_ARGS__) @@ -143,3 +150,12 @@ #define TEMPLATE_PRODUCT_TEST_CASE_METHOD_SIG(...) CATCH_TEMPLATE_PRODUCT_TEST_CASE_METHOD_SIG(__VA_ARGS__) #define TEMPLATE_LIST_TEST_CASE(...) CATCH_TEMPLATE_LIST_TEST_CASE(__VA_ARGS__) #define TEMPLATE_LIST_TEST_CASE_METHOD(...) CATCH_TEMPLATE_LIST_TEST_CASE_METHOD(__VA_ARGS__) + +// + +#define REQUIRE_THROWS_WITH(...) CATCH_REQUIRE_THROWS_WITH(__VA_ARGS__) +#define REQUIRE_THROWS_MATCHES(...) CATCH_REQUIRE_THROWS_MATCHES(__VA_ARGS__) +#define CHECK_THROWS_WITH(...) CATCH_CHECK_THROWS_WITH(__VA_ARGS__) +#define CHECK_THROWS_MATCHES(...) CATCH_CHECK_THROWS_MATCHES(__VA_ARGS__) +#define CHECK_THAT(...) CATCH_CHECK_THAT(__VA_ARGS__) +#define REQUIRE_THAT(...) CATCH_REQUIRE_THAT(__VA_ARGS__) From 7e3c5974a243279cac19ee188e94b3e789256ba7 Mon Sep 17 00:00:00 2001 From: David Bayer Date: Tue, 12 May 2026 20:09:55 +0200 Subject: [PATCH 03/10] fixes --- c2h/include/c2h/catch2_test_macros.h | 7 +++++++ cub/test/catch2_test_device_radix_sort_custom.cu | 4 ++-- .../test/libcudacxx/cuda/containers/buffer/access.cu | 2 +- 3 files changed, 10 insertions(+), 3 deletions(-) diff --git a/c2h/include/c2h/catch2_test_macros.h b/c2h/include/c2h/catch2_test_macros.h index 01c375065f6..7d888087aa4 100644 --- a/c2h/include/c2h/catch2_test_macros.h +++ b/c2h/include/c2h/catch2_test_macros.h @@ -159,3 +159,10 @@ #define CHECK_THROWS_MATCHES(...) CATCH_CHECK_THROWS_MATCHES(__VA_ARGS__) #define CHECK_THAT(...) CATCH_CHECK_THAT(__VA_ARGS__) #define REQUIRE_THAT(...) CATCH_REQUIRE_THAT(__VA_ARGS__) + +// extensions + +#define REQUIRE_CUDA(...) REQUIRE((__VA_ARGS__) == CUDA_SUCCESS) +#define REQUIRE_CUDART(...) REQUIRE((__VA_ARGS__) == cudaSuccess) +#define CHECK_CUDA(...) CHECK((__VA_ARGS__) == CUDA_SUCCESS) +#define CHECK_CUDART(...) CHECK((__VA_ARGS__) == cudaSuccess) diff --git a/cub/test/catch2_test_device_radix_sort_custom.cu b/cub/test/catch2_test_device_radix_sort_custom.cu index 9a2e1cb2f78..5e2b21c4f08 100644 --- a/cub/test/catch2_test_device_radix_sort_custom.cu +++ b/cub/test/catch2_test_device_radix_sort_custom.cu @@ -499,12 +499,12 @@ struct decomposer_t }; // example-end custom-type -static __host__ std::ostream& operator<<(std::ostream& os, const custom_t& self) +__host__ std::ostream& operator<<(std::ostream& os, const custom_t& self) { return os << "{ " << self.f << ", " << self.lli << " }"; } -static __host__ __device__ bool operator==(const custom_t& lhs, const custom_t& rhs) +__host__ __device__ bool operator==(const custom_t& lhs, const custom_t& rhs) { return lhs.f == rhs.f && lhs.lli == rhs.lli; } diff --git a/libcudacxx/test/libcudacxx/cuda/containers/buffer/access.cu b/libcudacxx/test/libcudacxx/cuda/containers/buffer/access.cu index 36db24bb697..a5e6bcbd063 100644 --- a/libcudacxx/test/libcudacxx/cuda/containers/buffer/access.cu +++ b/libcudacxx/test/libcudacxx/cuda/containers/buffer/access.cu @@ -121,7 +121,7 @@ C2H_CCCLRT_TEST("cuda::buffer access and stream", "[container][buffer]", test_ty static_assert(cuda::std::is_same_v); CCCLRT_CHECK(v2 == T(4)); - CHECK_THROWS_AS((buf.at(4)), std::out_of_range); + CHECK_THROWS_AS(((void) buf.at(4)), std::out_of_range); } } From 6b3cfcd91bd8eaa5eefc78b293b8bf56fcd7b2d5 Mon Sep 17 00:00:00 2001 From: David Bayer Date: Tue, 12 May 2026 20:12:39 +0200 Subject: [PATCH 04/10] simplify printing --- c2h/include/c2h/catch2_test_macros.h | 22 +++++++++++----------- 1 file changed, 11 insertions(+), 11 deletions(-) diff --git a/c2h/include/c2h/catch2_test_macros.h b/c2h/include/c2h/catch2_test_macros.h index 7d888087aa4..8133d943576 100644 --- a/c2h/include/c2h/catch2_test_macros.h +++ b/c2h/include/c2h/catch2_test_macros.h @@ -22,17 +22,17 @@ // Only a subset of the Catch2's macro are provided. If needed, feel free to extend the support. Host-only macros can // be determined by missing NV_IF_ELSE_TARGET wrapper and immediate dispatch to CATCH_-prefixed variant. -#define C2H_INTERNAL_DEVICE_PRINTF(FMT, ...) \ - ::printf(__FILE__ ":" _CCCL_TO_STRING(__LINE__) ":\n " FMT "\n block [%u, %u, %u], thread [%u, %u, %u]\n\n", \ - __VA_ARGS__, \ - blockIdx.x, \ - blockIdx.y, \ - blockIdx.z, \ - threadIdx.x, \ - threadIdx.y, \ - threadIdx.z) - -#define C2H_INTERNAL_DEVICE_TEST_PRINT(KIND, ...) C2H_INTERNAL_DEVICE_PRINTF(KIND "(%s) failed", "" #__VA_ARGS__) +#define C2H_INTERNAL_DEVICE_TEST_PRINT(KIND, ...) \ + ::printf( \ + __FILE__ \ + ":" _CCCL_TO_STRING(__LINE__) ":\n " KIND "(%s) failed\n block [%u, %u, %u], thread [%u, %u, %u]\n\n", \ + #__VA_ARGS__, \ + blockIdx.x, \ + blockIdx.y, \ + blockIdx.z, \ + threadIdx.x, \ + threadIdx.y, \ + threadIdx.z) // From 3dab09dbc42bafce11a0c6f9fbc340f8628f1079 Mon Sep 17 00:00:00 2001 From: David Bayer Date: Tue, 12 May 2026 20:53:37 +0200 Subject: [PATCH 05/10] fixes --- c2h/include/c2h/catch2_test_macros.h | 10 ++++++++++ cub/test/catch2_test_device_topk_api.cu | 8 ++++---- .../test/libcudacxx/cuda/ccclrt/common/testing.cuh | 2 +- 3 files changed, 15 insertions(+), 5 deletions(-) diff --git a/c2h/include/c2h/catch2_test_macros.h b/c2h/include/c2h/catch2_test_macros.h index 8133d943576..b3f86995ced 100644 --- a/c2h/include/c2h/catch2_test_macros.h +++ b/c2h/include/c2h/catch2_test_macros.h @@ -162,6 +162,16 @@ // extensions +#define REQUIRE_DEVICE(...) \ + do \ + { \ + if (!(__VA_ARGS__)) \ + { \ + C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE", __VA_ARGS__); \ + } \ + ::__trap(); \ + } while (false) + #define REQUIRE_CUDA(...) REQUIRE((__VA_ARGS__) == CUDA_SUCCESS) #define REQUIRE_CUDART(...) REQUIRE((__VA_ARGS__) == cudaSuccess) #define CHECK_CUDA(...) CHECK((__VA_ARGS__) == CUDA_SUCCESS) diff --git a/cub/test/catch2_test_device_topk_api.cu b/cub/test/catch2_test_device_topk_api.cu index 8b7d7a52a70..8e8b229172b 100644 --- a/cub/test/catch2_test_device_topk_api.cu +++ b/cub/test/catch2_test_device_topk_api.cu @@ -226,22 +226,22 @@ struct decomposer_t }; // example-end topk-custom-type -static __host__ std::ostream& operator<<(std::ostream& os, const custom_t& self) +__host__ std::ostream& operator<<(std::ostream& os, const custom_t& self) { return os << "{ " << self.f << ", " << self.lli << " }"; } -static __host__ __device__ bool operator==(const custom_t& lhs, const custom_t& rhs) +__host__ __device__ bool operator==(const custom_t& lhs, const custom_t& rhs) { return lhs.f == rhs.f && lhs.lli == rhs.lli; } -static __host__ __device__ bool operator<(const custom_t& lhs, const custom_t& rhs) +__host__ __device__ bool operator<(const custom_t& lhs, const custom_t& rhs) { return lhs.lli == rhs.lli ? lhs.f < rhs.f : lhs.lli < rhs.lli; } -static __host__ __device__ bool operator>(const custom_t& lhs, const custom_t& rhs) +__host__ __device__ bool operator>(const custom_t& lhs, const custom_t& rhs) { return rhs < lhs; } diff --git a/libcudacxx/test/libcudacxx/cuda/ccclrt/common/testing.cuh b/libcudacxx/test/libcudacxx/cuda/ccclrt/common/testing.cuh index c27433a82c2..52468fa155d 100644 --- a/libcudacxx/test/libcudacxx/cuda/ccclrt/common/testing.cuh +++ b/libcudacxx/test/libcudacxx/cuda/ccclrt/common/testing.cuh @@ -35,7 +35,7 @@ #define CCCLRT_CHECK_FALSE(condition) CHECK_FALSE(condition) // Explicit device side require macros for clang-cuda -#define CCCLRT_REQUIRE_DEVICE(condition) REQUIRE(condition) +#define CCCLRT_REQUIRE_DEVICE(condition) REQUIRE_DEVICE(condition) #define CCCLRT_CHECK_DEVICE(condition) CHECK(condition) #define CCCLRT_FAIL_DEVICE(message) FAIL(message) #define CCCLRT_CHECK_FALSE_DEVICE(condition) CHECK_FALSE(condition) From 39c07c5c66fd6031ed46d830ee1ba8c7feab865a Mon Sep 17 00:00:00 2001 From: David Bayer Date: Tue, 12 May 2026 22:15:19 +0200 Subject: [PATCH 06/10] more fixes --- c2h/include/c2h/catch2_test_macros.h | 10 +++++----- cudax/test/cufile/driver_register.cu | 2 +- 2 files changed, 6 insertions(+), 6 deletions(-) diff --git a/c2h/include/c2h/catch2_test_macros.h b/c2h/include/c2h/catch2_test_macros.h index b3f86995ced..bb2dc26a4da 100644 --- a/c2h/include/c2h/catch2_test_macros.h +++ b/c2h/include/c2h/catch2_test_macros.h @@ -41,16 +41,16 @@ if (!(__VA_ARGS__)) \ { \ C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE", __VA_ARGS__); \ + ::__trap(); \ } \ - ::__trap(); \ })) #define REQUIRE_FALSE(...) \ NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_FALSE(__VA_ARGS__);), ({ \ if (__VA_ARGS__) \ { \ C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE_FALSE", __VA_ARGS__); \ + ::__trap(); \ } \ - ::__trap(); \ })) #define REQUIRE_THROWS(...) \ @@ -72,16 +72,16 @@ if (!(__VA_ARGS__)) \ { \ C2H_INTERNAL_DEVICE_TEST_PRINT("CHECK", __VA_ARGS__); \ + ::__trap(); \ } \ - ::__trap(); \ })) #define CHECK_FALSE(...) \ NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_FALSE(__VA_ARGS__);), ({ \ if (__VA_ARGS__) \ { \ C2H_INTERNAL_DEVICE_TEST_PRINT("CHECK_FALSE", __VA_ARGS__); \ + ::__trap(); \ } \ - ::__trap(); \ })) #define CHECKED_IF(...) CATCH_CHECKED_IF(__VA_ARGS__) #define CHECKED_ELSE(...) CATCH_CHECKED_ELSE(__VA_ARGS__) @@ -168,8 +168,8 @@ if (!(__VA_ARGS__)) \ { \ C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE", __VA_ARGS__); \ + ::__trap(); \ } \ - ::__trap(); \ } while (false) #define REQUIRE_CUDA(...) REQUIRE((__VA_ARGS__) == CUDA_SUCCESS) diff --git a/cudax/test/cufile/driver_register.cu b/cudax/test/cufile/driver_register.cu index 5cd30afdcfa..d943df8f026 100644 --- a/cudax/test/cufile/driver_register.cu +++ b/cudax/test/cufile/driver_register.cu @@ -47,7 +47,7 @@ void test_register_native_handle() CUDAX_REQUIRE(file.get() != nullptr); // 4. Reregistering the same file handle should result in an cufile_error. - CHECK_THROWS_AS(cudax::cufile_driver.register_native_handle(fd), cudax::cufile_error); + CHECK_THROWS_AS((void) cudax::cufile_driver.register_native_handle(fd), cudax::cufile_error); // 5. Deregister the cuFile handles. Can be called multiple times. cudax::cufile_driver.deregister_native_handle(file); From 1265a5c0003f21e1c31d9631c7a9dbdd06c0f261 Mon Sep 17 00:00:00 2001 From: David Bayer Date: Tue, 12 May 2026 22:35:33 +0200 Subject: [PATCH 07/10] improve doc strings --- c2h/include/c2h/catch2_test_macros.h | 126 ++++++++++++++------------- 1 file changed, 67 insertions(+), 59 deletions(-) diff --git a/c2h/include/c2h/catch2_test_macros.h b/c2h/include/c2h/catch2_test_macros.h index bb2dc26a4da..0ea56bf0003 100644 --- a/c2h/include/c2h/catch2_test_macros.h +++ b/c2h/include/c2h/catch2_test_macros.h @@ -22,11 +22,13 @@ // Only a subset of the Catch2's macro are provided. If needed, feel free to extend the support. Host-only macros can // be determined by missing NV_IF_ELSE_TARGET wrapper and immediate dispatch to CATCH_-prefixed variant. -#define C2H_INTERNAL_DEVICE_TEST_PRINT(KIND, ...) \ +// We must pass the COND as a cstring parameter, because it might contain the '%' character that would break the printf +// formatting. +#define C2H_INTERNAL_DEVICE_TEST_PRINT(KIND, COND) \ ::printf( \ __FILE__ \ ":" _CCCL_TO_STRING(__LINE__) ":\n " KIND "(%s) failed\n block [%u, %u, %u], thread [%u, %u, %u]\n\n", \ - #__VA_ARGS__, \ + COND, \ blockIdx.x, \ blockIdx.y, \ blockIdx.z, \ @@ -36,66 +38,66 @@ // -#define REQUIRE(...) \ - NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE(__VA_ARGS__);), ({ \ - if (!(__VA_ARGS__)) \ - { \ - C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE", __VA_ARGS__); \ - ::__trap(); \ - } \ +#define REQUIRE(...) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE(__VA_ARGS__);), ({ \ + if (!(__VA_ARGS__)) \ + { \ + C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE", #__VA_ARGS__); \ + ::__trap(); \ + } \ })) -#define REQUIRE_FALSE(...) \ - NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_FALSE(__VA_ARGS__);), ({ \ - if (__VA_ARGS__) \ - { \ - C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE_FALSE", __VA_ARGS__); \ - ::__trap(); \ - } \ +#define REQUIRE_FALSE(...) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_FALSE(__VA_ARGS__);), ({ \ + if (__VA_ARGS__) \ + { \ + C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE_FALSE", #__VA_ARGS__); \ + ::__trap(); \ + } \ })) -#define REQUIRE_THROWS(...) \ - NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_THROWS(__VA_ARGS__);), ({ \ - __VA_ARGS__; \ - C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE_THROWS", __VA_ARGS__); \ - ::__trap(); \ +#define REQUIRE_THROWS(...) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_THROWS(__VA_ARGS__);), ({ \ + __VA_ARGS__; \ + C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE_THROWS", #__VA_ARGS__); \ + ::__trap(); \ })) -#define REQUIRE_THROWS_AS(EXPR, TYPE) \ - NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_THROWS_AS((EXPR), TYPE);), ({ \ - EXPR; \ - C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE_THROWS_AS", EXPR, TYPE); \ - ::__trap(); \ +#define REQUIRE_THROWS_AS(EXPR, TYPE) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_THROWS_AS((EXPR), TYPE);), ({ \ + EXPR; \ + C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE_THROWS_AS", #EXPR #TYPE); \ + ::__trap(); \ })) #define REQUIRE_NOTHROW(...) NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_NOTHROW(__VA_ARGS__);), (__VA_ARGS__;)) -#define CHECK(...) \ - NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK(__VA_ARGS__);), ({ \ - if (!(__VA_ARGS__)) \ - { \ - C2H_INTERNAL_DEVICE_TEST_PRINT("CHECK", __VA_ARGS__); \ - ::__trap(); \ - } \ +#define CHECK(...) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK(__VA_ARGS__);), ({ \ + if (!(__VA_ARGS__)) \ + { \ + C2H_INTERNAL_DEVICE_TEST_PRINT("CHECK", #__VA_ARGS__); \ + ::__trap(); \ + } \ })) -#define CHECK_FALSE(...) \ - NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_FALSE(__VA_ARGS__);), ({ \ - if (__VA_ARGS__) \ - { \ - C2H_INTERNAL_DEVICE_TEST_PRINT("CHECK_FALSE", __VA_ARGS__); \ - ::__trap(); \ - } \ +#define CHECK_FALSE(...) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_FALSE(__VA_ARGS__);), ({ \ + if (__VA_ARGS__) \ + { \ + C2H_INTERNAL_DEVICE_TEST_PRINT("CHECK_FALSE", #__VA_ARGS__); \ + ::__trap(); \ + } \ })) #define CHECKED_IF(...) CATCH_CHECKED_IF(__VA_ARGS__) #define CHECKED_ELSE(...) CATCH_CHECKED_ELSE(__VA_ARGS__) #define CHECK_NOFAIL(...) CATCH_CHECK_NOFAIL(__VA_ARGS__) -#define CHECK_THROWS(...) \ - NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_THROWS(__VA_ARGS__);), ({ \ - __VA_ARGS__; \ - C2H_INTERNAL_DEVICE_TEST_PRINT("CHECK_THROWS", __VA_ARGS__); \ +#define CHECK_THROWS(...) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_THROWS(__VA_ARGS__);), ({ \ + __VA_ARGS__; \ + C2H_INTERNAL_DEVICE_TEST_PRINT("CHECK_THROWS", #__VA_ARGS__); \ })) -#define CHECK_THROWS_AS(EXPR, TYPE) \ - NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_THROWS_AS((EXPR), TYPE);), ({ \ - EXPR; \ - C2H_INTERNAL_DEVICE_TEST_PRINT("CHECK_THROWS_AS", EXPR, TYPE); \ +#define CHECK_THROWS_AS(EXPR, TYPE) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_THROWS_AS((EXPR), TYPE);), ({ \ + EXPR; \ + C2H_INTERNAL_DEVICE_TEST_PRINT("CHECK_THROWS_AS", #EXPR #TYPE); \ })) #define CHECK_NOTHROW(...) NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_NOTHROW(__VA_ARGS__);), (__VA_ARGS__;)) @@ -106,7 +108,7 @@ #define SECTION(...) CATCH_SECTION(__VA_ARGS__) #define DYNAMIC_SECTION(...) CATCH_DYNAMIC_SECTION(__VA_ARGS__) #define FAIL(...) \ - NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_FAIL(__VA_ARGS__);), ({ C2H_INTERNAL_DEVICE_TEST_PRINT("FAIL", __VA_ARGS__); })) + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_FAIL(__VA_ARGS__);), ({ C2H_INTERNAL_DEVICE_TEST_PRINT("FAIL", #__VA_ARGS__); })) #define FAIL_CHECK(...) CATCH_FAIL_CHECK(__VA_ARGS__) #define SUCCEED(...) CATCH_SUCCEED(__VA_ARGS__) #define SKIP(...) CATCH_SKIP(__VA_ARGS__) @@ -162,17 +164,23 @@ // extensions -#define REQUIRE_DEVICE(...) \ - do \ - { \ - if (!(__VA_ARGS__)) \ - { \ - C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE", __VA_ARGS__); \ - ::__trap(); \ - } \ +// Sometimes clang-cuda has problems with REQUIRE(...) when used in __device__ function - it tries to instantiate the +// host path. This is related to clang-cuda's compilation trajectory. For these cases, we provide REQUIRE_DEVICE(...) as +// a fallback. +#define REQUIRE_DEVICE(...) \ + do \ + { \ + if (!(__VA_ARGS__)) \ + { \ + C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE", #__VA_ARGS__); \ + ::__trap(); \ + } \ } while (false) -#define REQUIRE_CUDA(...) REQUIRE((__VA_ARGS__) == CUDA_SUCCESS) +// Macros to require/check success of a CUDA Driver call. +#define REQUIRE_CUDA(...) REQUIRE((__VA_ARGS__) == CUDA_SUCCESS) +#define CHECK_CUDA(...) CHECK((__VA_ARGS__) == CUDA_SUCCESS) + +// Macros to require/check success of a CUDA Runtime call. #define REQUIRE_CUDART(...) REQUIRE((__VA_ARGS__) == cudaSuccess) -#define CHECK_CUDA(...) CHECK((__VA_ARGS__) == CUDA_SUCCESS) #define CHECK_CUDART(...) CHECK((__VA_ARGS__) == cudaSuccess) From 969064931861c7865eaa8a5c6dfce2db72a8310f Mon Sep 17 00:00:00 2001 From: David Bayer Date: Thu, 14 May 2026 20:36:41 +0200 Subject: [PATCH 08/10] fix clang tidy --- .../cuda/memory_resource/resources/shared_memory_pools.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/libcudacxx/test/libcudacxx/cuda/memory_resource/resources/shared_memory_pools.cu b/libcudacxx/test/libcudacxx/cuda/memory_resource/resources/shared_memory_pools.cu index 6d9afa5cab8..a83df8c77cd 100644 --- a/libcudacxx/test/libcudacxx/cuda/memory_resource/resources/shared_memory_pools.cu +++ b/libcudacxx/test/libcudacxx/cuda/memory_resource/resources/shared_memory_pools.cu @@ -168,7 +168,7 @@ C2H_CCCLRT_TEST_LIST("shared_memory_pool comparison", "[memory_resource]", SHARE SECTION("Copies are equal") { - shared_pool copy(pool1); + shared_pool copy(pool1); // NOLINT(performance-unnecessary-copy-initialization) CHECK(pool1 == copy); } From 694ebaf250dadd003fb76632a91b51cebd78544a Mon Sep 17 00:00:00 2001 From: David Bayer Date: Thu, 14 May 2026 22:35:17 +0200 Subject: [PATCH 09/10] fixes --- c2h/include/c2h/catch2_test_macros.h | 27 ++++++++++----------------- 1 file changed, 10 insertions(+), 17 deletions(-) diff --git a/c2h/include/c2h/catch2_test_macros.h b/c2h/include/c2h/catch2_test_macros.h index 0ea56bf0003..cd6cd528226 100644 --- a/c2h/include/c2h/catch2_test_macros.h +++ b/c2h/include/c2h/catch2_test_macros.h @@ -1,5 +1,5 @@ -// SPDX-FileCopyrightText: Copyright (c) 2011-2022, NVIDIA CORPORATION. All rights reserved. -// SPDX-License-Identifier: BSD-3-Clause +// SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception #pragma once @@ -55,19 +55,9 @@ } \ })) -#define REQUIRE_THROWS(...) \ - NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_THROWS(__VA_ARGS__);), ({ \ - __VA_ARGS__; \ - C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE_THROWS", #__VA_ARGS__); \ - ::__trap(); \ - })) -#define REQUIRE_THROWS_AS(EXPR, TYPE) \ - NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_THROWS_AS((EXPR), TYPE);), ({ \ - EXPR; \ - C2H_INTERNAL_DEVICE_TEST_PRINT("REQUIRE_THROWS_AS", #EXPR #TYPE); \ - ::__trap(); \ - })) -#define REQUIRE_NOTHROW(...) NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_NOTHROW(__VA_ARGS__);), (__VA_ARGS__;)) +#define REQUIRE_THROWS(...) CATCH_REQUIRE_THROWS(__VA_ARGS__) +#define REQUIRE_THROWS_AS(...) CATCH_REQUIRE_THROWS_AS(__VA_ARGS__) +#define REQUIRE_NOTHROW(...) NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE_NOTHROW(__VA_ARGS__);), (__VA_ARGS__;)) #define CHECK(...) \ NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK(__VA_ARGS__);), ({ \ @@ -107,8 +97,11 @@ #define REGISTER_TEST_CASE(...) CATCH_REGISTER_TEST_CASE(__VA_ARGS__) #define SECTION(...) CATCH_SECTION(__VA_ARGS__) #define DYNAMIC_SECTION(...) CATCH_DYNAMIC_SECTION(__VA_ARGS__) -#define FAIL(...) \ - NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_FAIL(__VA_ARGS__);), ({ C2H_INTERNAL_DEVICE_TEST_PRINT("FAIL", #__VA_ARGS__); })) +#define FAIL(...) \ + NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_FAIL(__VA_ARGS__);), ({ \ + C2H_INTERNAL_DEVICE_TEST_PRINT("FAIL", #__VA_ARGS__); \ + ::__trap(); \ + })) #define FAIL_CHECK(...) CATCH_FAIL_CHECK(__VA_ARGS__) #define SUCCEED(...) CATCH_SUCCEED(__VA_ARGS__) #define SKIP(...) CATCH_SKIP(__VA_ARGS__) From 490c5312a563f7433e56da3ad75a6e63e0bba754 Mon Sep 17 00:00:00 2001 From: David Bayer Date: Thu, 14 May 2026 22:39:27 +0200 Subject: [PATCH 10/10] fixes --- c2h/include/c2h/catch2_test_macros.h | 14 +++----------- 1 file changed, 3 insertions(+), 11 deletions(-) diff --git a/c2h/include/c2h/catch2_test_macros.h b/c2h/include/c2h/catch2_test_macros.h index cd6cd528226..8de20d3330e 100644 --- a/c2h/include/c2h/catch2_test_macros.h +++ b/c2h/include/c2h/catch2_test_macros.h @@ -79,17 +79,9 @@ #define CHECKED_ELSE(...) CATCH_CHECKED_ELSE(__VA_ARGS__) #define CHECK_NOFAIL(...) CATCH_CHECK_NOFAIL(__VA_ARGS__) -#define CHECK_THROWS(...) \ - NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_THROWS(__VA_ARGS__);), ({ \ - __VA_ARGS__; \ - C2H_INTERNAL_DEVICE_TEST_PRINT("CHECK_THROWS", #__VA_ARGS__); \ - })) -#define CHECK_THROWS_AS(EXPR, TYPE) \ - NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_THROWS_AS((EXPR), TYPE);), ({ \ - EXPR; \ - C2H_INTERNAL_DEVICE_TEST_PRINT("CHECK_THROWS_AS", #EXPR #TYPE); \ - })) -#define CHECK_NOTHROW(...) NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_NOTHROW(__VA_ARGS__);), (__VA_ARGS__;)) +#define CHECK_THROWS(...) CATCH_CHECK_THROWS(__VA_ARGS__) +#define CHECK_THROWS_AS(...) CATCH_CHECK_THROWS_AS(__VA_ARGS__) +#define CHECK_NOTHROW(...) NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_CHECK_NOTHROW(__VA_ARGS__);), (__VA_ARGS__;)) #define TEST_CASE(...) CATCH_TEST_CASE(__VA_ARGS__) #define TEST_CASE_METHOD(...) CATCH_TEST_CASE_METHOD(__VA_ARGS__)