From 21aff40f4813870f5abd7b031a1f62cfbe2e1c70 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Juan=20Manuel=20Martinez=20Caama=C3=B1o?= Date: Thu, 19 Jun 2025 09:13:26 +0200 Subject: [PATCH 1/3] Reapply "[CUDA][HIP] Add a __device__ version of std::__glibcxx_assert_fail() (#136133)" This reverts commit 5bfb9bb3a0002de9007d4bc04c2b0e300a72c52d. --- clang/lib/Headers/CMakeLists.txt | 1 + .../Headers/cuda_wrappers/bits/c++config.h | 51 +++++++++++++++++++ 2 files changed, 52 insertions(+) create mode 100644 clang/lib/Headers/cuda_wrappers/bits/c++config.h diff --git a/clang/lib/Headers/CMakeLists.txt b/clang/lib/Headers/CMakeLists.txt index c1c9d2e8c7b79..c96d209c1fc0c 100644 --- a/clang/lib/Headers/CMakeLists.txt +++ b/clang/lib/Headers/CMakeLists.txt @@ -341,6 +341,7 @@ set(cuda_wrapper_files ) set(cuda_wrapper_bits_files + cuda_wrappers/bits/c++config.h cuda_wrappers/bits/shared_ptr_base.h cuda_wrappers/bits/basic_string.h cuda_wrappers/bits/basic_string.tcc diff --git a/clang/lib/Headers/cuda_wrappers/bits/c++config.h b/clang/lib/Headers/cuda_wrappers/bits/c++config.h new file mode 100644 index 0000000000000..eafa13a9cc640 --- /dev/null +++ b/clang/lib/Headers/cuda_wrappers/bits/c++config.h @@ -0,0 +1,51 @@ +// libstdc++ uses the non-constexpr function std::__glibcxx_assert_fail() +// to trigger compilation errors when the __glibcxx_assert(cond) macro +// is used in a constexpr context. +// Compilation fails when using code from the libstdc++ (such as std::array) on +// device code, since these assertions invoke a non-constexpr host function from +// device code. +// +// To work around this issue, we declare our own device version of the function + +#ifndef __CLANG_CUDA_WRAPPERS_BITS_CPP_CONFIG +#define __CLANG_CUDA_WRAPPERS_BITS_CPP_CONFIG + +#include_next + +#ifdef _LIBCPP_BEGIN_NAMESPACE_STD +_LIBCPP_BEGIN_NAMESPACE_STD +#else +namespace std { +#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION +_GLIBCXX_BEGIN_NAMESPACE_VERSION +#endif + +#ifdef _GLIBCXX_VERBOSE_ASSERT +__attribute__((device, noreturn)) inline void +__glibcxx_assert_fail(const char *file, int line, const char *function, + const char *condition) noexcept { + if (file && function && condition) + __builtin_printf("%s:%d: %s: Assertion '%s' failed.\n", file, line, + function, condition); + else if (function) + __builtin_printf("%s: Undefined behavior detected.\n", function); + __builtin_abort(); +} +#endif + +#endif +__attribute__((device, noreturn, __always_inline__, + __visibility__("default"))) inline void +__glibcxx_assert_fail(...) noexcept { + __builtin_abort(); +} +#ifdef _LIBCPP_END_NAMESPACE_STD +_LIBCPP_END_NAMESPACE_STD +#else +#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION +_GLIBCXX_END_NAMESPACE_VERSION +#endif +} // namespace std +#endif + +#endif From 316644d6e9149a714a01c5cf6860fae16176db0b Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Juan=20Manuel=20Martinez=20Caama=C3=B1o?= Date: Thu, 19 Jun 2025 09:15:50 +0200 Subject: [PATCH 2/3] [CUDA][HIP] Add noexcept only after C++11 on __glibcxx_assert_fail --- clang/lib/Headers/cuda_wrappers/bits/c++config.h | 14 ++++++++++++-- 1 file changed, 12 insertions(+), 2 deletions(-) diff --git a/clang/lib/Headers/cuda_wrappers/bits/c++config.h b/clang/lib/Headers/cuda_wrappers/bits/c++config.h index eafa13a9cc640..98773ff868c3e 100644 --- a/clang/lib/Headers/cuda_wrappers/bits/c++config.h +++ b/clang/lib/Headers/cuda_wrappers/bits/c++config.h @@ -20,10 +20,17 @@ namespace std { _GLIBCXX_BEGIN_NAMESPACE_VERSION #endif +#pragma push_macro("CUDA_NOEXCEPT") +#if __cplusplus >= 201103L +#define CUDA_NOEXCEPT noexcept +#else +#define CUDA_NOEXCEPT +#endif + #ifdef _GLIBCXX_VERBOSE_ASSERT __attribute__((device, noreturn)) inline void __glibcxx_assert_fail(const char *file, int line, const char *function, - const char *condition) noexcept { + const char *condition) CUDA_NOEXCEPT { if (file && function && condition) __builtin_printf("%s:%d: %s: Assertion '%s' failed.\n", file, line, function, condition); @@ -36,9 +43,12 @@ __glibcxx_assert_fail(const char *file, int line, const char *function, #endif __attribute__((device, noreturn, __always_inline__, __visibility__("default"))) inline void -__glibcxx_assert_fail(...) noexcept { +__glibcxx_assert_fail(...) CUDA_NOEXCEPT { __builtin_abort(); } + +#pragma pop_macro("CUDA_NOEXCEPT") + #ifdef _LIBCPP_END_NAMESPACE_STD _LIBCPP_END_NAMESPACE_STD #else From 71d1ca33e843c8b9e01f1a29a697691e50a68d6c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Juan=20Manuel=20Martinez=20Caama=C3=B1o?= Date: Thu, 19 Jun 2025 10:36:19 +0200 Subject: [PATCH 3/3] [CUDA][HIP] Fix error 'CUDA device code does not support variadic functions' --- clang/lib/Headers/cuda_wrappers/bits/c++config.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/lib/Headers/cuda_wrappers/bits/c++config.h b/clang/lib/Headers/cuda_wrappers/bits/c++config.h index 98773ff868c3e..27083253181d2 100644 --- a/clang/lib/Headers/cuda_wrappers/bits/c++config.h +++ b/clang/lib/Headers/cuda_wrappers/bits/c++config.h @@ -27,23 +27,23 @@ _GLIBCXX_BEGIN_NAMESPACE_VERSION #define CUDA_NOEXCEPT #endif -#ifdef _GLIBCXX_VERBOSE_ASSERT __attribute__((device, noreturn)) inline void __glibcxx_assert_fail(const char *file, int line, const char *function, const char *condition) CUDA_NOEXCEPT { +#ifdef _GLIBCXX_VERBOSE_ASSERT if (file && function && condition) __builtin_printf("%s:%d: %s: Assertion '%s' failed.\n", file, line, function, condition); else if (function) __builtin_printf("%s: Undefined behavior detected.\n", function); +#endif __builtin_abort(); } -#endif #endif __attribute__((device, noreturn, __always_inline__, __visibility__("default"))) inline void -__glibcxx_assert_fail(...) CUDA_NOEXCEPT { +__glibcxx_assert_fail() CUDA_NOEXCEPT { __builtin_abort(); }