[clang] 8ec0552 - Reapply "[CUDA][HIP] Add a __device__ version of std::__glibcxx_assert_fail() (#144886)

[via cfe-commits]

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