================
@@ -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:
> Presumably we _do_ want to fail here. We should at least call
> __builtin_abort() or something.
This fails in fact. It's a "trick" to trigger a compilation failure on a
constantly evaluated context: the code ends up calling a non-constexpr function
from a constantly evaluated context when the assertion fails.
The new thing in this libstdc++ version is that even when assertions are
disabled, on a constexpr context, we still have assert checks.
> Also, __glibcxx_assert_fail() appears to have variants with multiple
> parameters, too:
>
> https://github.com/gcc-mirror/gcc/blob/4cff0434e8bf6683988482a7e47f8459c06f2c05/libstdc%2B%2B-v3/src/c%2B%2B11/assert_fail.cc#L33
>
> It looks like it's version-dependent, and we may need to provide both
> variants, otherwise the issue will still be there for other libstdc++
> versions.
Thanks for pointing this out. I'll add the non-constexpr version too.
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
