Author: Juan Manuel Martinez Caamaño Date: 2025-06-24T09:13:13+02:00 New Revision: 8ec0552a7f1f50986dda6d13eae310d121d7e3ba
URL: https://github.com/llvm/llvm-project/commit/8ec0552a7f1f50986dda6d13eae310d121d7e3ba DIFF: https://github.com/llvm/llvm-project/commit/8ec0552a7f1f50986dda6d13eae310d121d7e3ba.diff LOG: Reapply "[CUDA][HIP] Add a __device__ version of std::__glibcxx_assert_fail() (#144886) Modifications to reapply the commit: * Add noexcept only after C++11 on __glibcxx_assert_fail * Remove vararg version of __glibcxx_assert_fail Added: clang/lib/Headers/cuda_wrappers/bits/c++config.h Modified: clang/lib/Headers/CMakeLists.txt Removed: ################################################################################ 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..27083253181d2 --- /dev/null +++ b/clang/lib/Headers/cuda_wrappers/bits/c++config.h @@ -0,0 +1,61 @@ +// 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 <bits/c++config.h> + +#ifdef _LIBCPP_BEGIN_NAMESPACE_STD +_LIBCPP_BEGIN_NAMESPACE_STD +#else +namespace std { +#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION +_GLIBCXX_BEGIN_NAMESPACE_VERSION +#endif + +#pragma push_macro("CUDA_NOEXCEPT") +#if __cplusplus >= 201103L +#define CUDA_NOEXCEPT noexcept +#else +#define CUDA_NOEXCEPT +#endif + +__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 +__attribute__((device, noreturn, __always_inline__, + __visibility__("default"))) inline void +__glibcxx_assert_fail() CUDA_NOEXCEPT { + __builtin_abort(); +} + +#pragma pop_macro("CUDA_NOEXCEPT") + +#ifdef _LIBCPP_END_NAMESPACE_STD +_LIBCPP_END_NAMESPACE_STD +#else +#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION +_GLIBCXX_END_NAMESPACE_VERSION +#endif +} // namespace std +#endif + +#endif _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
