================
@@ -0,0 +1,35 @@
+// 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
+__device__
+    __attribute__((__always_inline__, __visibility__("default"))) inline void
+    __glibcxx_assert_fail() {}
----------------
jmmartinez wrote:

Maybe this could help to understand the intention of the empty 
`__glibcxx_assert_fail` function: https://godbolt.org/z/YfY6nec1K

This snippet tries to reflect what would happen when runtime assertions are 
disabled; but constexpr assertions are enabled.
- On a non-constexpr context the whole assertion code gets optimized out since 
std::is_constant_evaluated is always false
- On a constexpr context, when we uncomment the constexpr from the declaration 
of Elt, the code ends up calling the "non_constexpr" function from a constexpr 
context; which raises an error.
- If we change the index to be inbounds, the code compiles.

https://github.com/llvm/llvm-project/pull/136133
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

Reply via email to