Reapply "[CUDA][HIP] Add a __device__ version of std::__glibcxx_assert_fail() (llvm#144886) (llvm#3189)

Modifications to reapply the commit:
* Add noexcept only after C++11 on __glibcxx_assert_fail
* Remove vararg version of __glibcxx_assert_fail

And doc CP.

Issue
[SWDEV-518041](https://ontrack-internal.amd.com/browse/SWDEV-518041)
& doc task
[SWDEV-538485](https://ontrack-internal.amd.com/browse/SWDEV-538485)

---------

Co-authored-by: Juan Manuel Martinez Caamaño <[email protected]>

Author: jmmartinez

clang/docs/ReleaseNotes.rst

@@ -1274,6 +1274,8 @@ CUDA/HIP Language Changes
 - Fixed a bug about overriding a constexpr pure-virtual member function with a non-constexpr virtual member function which causes compilation failure when including standard C++ header `format`.
 - Added initial support for version 3 of the compressed offload bundle format, which uses 64-bit fields for Total File Size and Uncompressed Binary Size. This enables support for files larger than 4GB. The support is currently experimental and can be enabled by setting the environment variable `COMPRESSED_BUNDLE_FORMAT_VERSION=3`.
 
+* Provide a __device__ version of std::__glibcxx_assert_fail() in a header wrapper.
+
 CUDA Support
 ^^^^^^^^^^^^
 - Clang now supports CUDA SDK up to 12.6

clang/lib/Headers/CMakeLists.txt

@@ -333,6 +333,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

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