[c2h] Make catch macros work on device #8928
There are no files selected for viewing
| Original file line number | Diff line number | Diff line change |
|---|---|---|
| @@ -0,0 +1,171 @@ | ||
| // SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. | ||
| // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception | ||
|
|
||
| #pragma once | ||
|
|
||
| #include <cuda/std/detail/__config> | ||
|
|
||
| #include <nv/target> | ||
|
|
||
| #include <catch2/catch_message.hpp> | ||
| #include <catch2/catch_template_test_macros.hpp> | ||
| #include <catch2/catch_test_macros.hpp> | ||
| #include <catch2/matchers/catch_matchers.hpp> | ||
|
|
||
| // 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. | ||
|
|
||
| // We must pass the COND as a cstring parameter, because it might contain the '%' character that would break the printf | ||
|
Jacobfaib marked this conversation as resolved.
|
||
| // 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", \ | ||
| COND, \ | ||
| blockIdx.x, \ | ||
| blockIdx.y, \ | ||
| blockIdx.z, \ | ||
| threadIdx.x, \ | ||
| threadIdx.y, \ | ||
| threadIdx.z) | ||
|
|
||
| // <catch2/catch2_test_macros.hpp> | ||
|
|
||
| #define REQUIRE(...) \ | ||
| NV_IF_ELSE_TARGET(NV_IS_HOST, (CATCH_REQUIRE(__VA_ARGS__);), ({ \ | ||
|
Contributor
There was a problem hiding this comment. Don't you need to handle clang-coda here?
Contributor
Author
There was a problem hiding this comment. It seems to be working correctly in most cases, for other problematic cases I added |
||
| 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(...) 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(...) \ | ||
|
Contributor
There was a problem hiding this comment. Similarly, we should consider removing
Contributor
There was a problem hiding this comment. I think there could be value in having more information in the log from some external testing like CI or QA, so you might not have to reproduce the issue yourself, while REQUIRE would not give you enough information to figure out the failure. I wouldn't say its common and I can see using CHECK too often lead to a more messy log with multiple lines of repeated same failure for example. So its very case by case, but I think there is space for both of these macros
Contributor
Author
There was a problem hiding this comment. Agreed! But maybe someone else would have a valid use case for that
Contributor
There was a problem hiding this comment.
Ehhh, overall I think this ends up being rare in practice. I see it similarly to C++ compiler error messages. Yes, once in a blue moon seeing the 30 different overloads of More often than not, CHECK(container.size() >= n);
// a lot of other unrelated code...
CHECK(container[n] == foo); // oopsAnd now you are left trying parse the crash log instead.
Contributor
There was a problem hiding this comment. I use |
||
| 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(...) 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__) | ||
| #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_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__) | ||
|
|
||
| #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__) | ||
|
|
||
| // <catch2/catch_message.hpp> | ||
|
|
||
| #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__) | ||
|
|
||
| // <catch2/catch_template_test_macros.hpp> | ||
|
|
||
| #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__) | ||
|
|
||
| // <catch2/matchers/catch_matchers.hpp> | ||
|
|
||
| #define REQUIRE_THROWS_WITH(...) CATCH_REQUIRE_THROWS_WITH(__VA_ARGS__) | ||
| #define REQUIRE_THROWS_MATCHES(...) CATCH_REQUIRE_THROWS_MATCHES(__VA_ARGS__) | ||
|
Contributor
There was a problem hiding this comment. In my humble (and correct!) opinion, Perhaps we can take this time to fix up some defaults and just delete the other macros.
Contributor
Author
There was a problem hiding this comment. We definitely can! Let's move this discussion to slack/new issue. This PR really only aims to provide host + device working alternatives
Contributor
There was a problem hiding this comment. I would strongly advise to stay with the exact semantics of Catch2. If If we change the meaning of macros between C2H and Catch2, it will create a lot of confusion for maintainers. |
||
| #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__) | ||
|
|
||
| // extensions | ||
|
|
||
| // 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) | ||
|
|
||
| // 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_CUDART(...) CHECK((__VA_ARGS__) == cudaSuccess) | ||
| Original file line number | Diff line number | Diff line change |
|---|---|---|
|
|
@@ -12,7 +12,7 @@ | |
|
|
||
| #include <cstddef> | ||
|
|
||
| #include <c2h/catch2_test_helper.h> | ||
|
|
||
| namespace cudax = cuda::experimental; | ||
|
|
||
|
|
||
Uh oh!
There was an error while loading. Please reload this page.