https://github.com/jmmartinez updated
https://github.com/llvm/llvm-project/pull/136133
From b433add1b7c2027b89cdd9be11680c0a9a59e807 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Juan=20Manuel=20Martinez=20Caama=C3=B1o?= <juama...@amd.com>
Date: Thu, 17 Apr 2025 13:41:55 +0200
Subject: [PATCH] [CUDA][HIP] Add a __device__ version of
std::__glibcxx_assert_fail()
libstdc++ 15 uses the non-constexpr function
std::__glibcxx_assert_fail() to trigger compilation errors when the
__glibcxx_assert(cond) macro is used in a constantly evaluated 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.
This patch proposes a cuda wrapper header "bits/c++config.h" which adds a
__device__ version of std::__glibcxx_assert_fail().
---
clang/lib/Headers/CMakeLists.txt | 1 +
.../Headers/cuda_wrappers/bits/c++config.h | 37 +++++++++++++++++++
2 files changed, 38 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 acf49e40c447e..54395e053dbc4 100644
--- a/clang/lib/Headers/CMakeLists.txt
+++ b/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
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..61c7e4a2043d1
--- /dev/null
+++ b/clang/lib/Headers/cuda_wrappers/bits/c++config.h
@@ -0,0 +1,37 @@
+// 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
+#endif
+__attribute__((device))
+__attribute__((__always_inline__, __visibility__("default"))) inline void
+__glibcxx_assert_fail(...) {
+ __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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
