Skip to content

Navigation Menu

Sign in
Appearance settings

Search code, repositories, users, issues, pull requests...

Provide feedback

We read every piece of feedback, and take your input very seriously.

Saved searches

Use saved searches to filter your results more quickly

Appearance settings

Commit 66d1103

Browse filesBrowse files
committed
[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().
1 parent 181872f commit 66d1103
Copy full SHA for 66d1103

File tree

2 files changed

+40
-0
lines changed
Filter options

2 files changed

+40
-0
lines changed

‎clang/lib/Headers/CMakeLists.txt

Copy file name to clipboardExpand all lines: clang/lib/Headers/CMakeLists.txt
+1Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -333,6 +333,7 @@ set(cuda_wrapper_files
333333
)
334334

335335
set(cuda_wrapper_bits_files
336+
cuda_wrappers/bits/c++config.h
336337
cuda_wrappers/bits/shared_ptr_base.h
337338
cuda_wrappers/bits/basic_string.h
338339
cuda_wrappers/bits/basic_string.tcc
+39Lines changed: 39 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,39 @@
1+
// libstdc++ uses the non-constexpr function std::__glibcxx_assert_fail()
2+
// to trigger compilation errors when the __glibcxx_assert(cond) macro
3+
// is used in a constexpr context.
4+
// Compilation fails when using code from the libstdc++ (such as std::array) on
5+
// device code, since these assertions invoke a non-constexpr host function from
6+
// device code.
7+
//
8+
// To work around this issue, we declare our own device version of the function
9+
10+
#ifndef __CLANG_CUDA_WRAPPERS_BITS_CPP_CONFIG
11+
#define __CLANG_CUDA_WRAPPERS_BITS_CPP_CONFIG
12+
13+
#include_next <bits/c++config.h>
14+
15+
#if _GLIBCXX_HAVE_IS_CONSTANT_EVALUATED
16+
17+
#ifdef _LIBCPP_BEGIN_NAMESPACE_STD
18+
_LIBCPP_BEGIN_NAMESPACE_STD
19+
#else
20+
namespace std {
21+
#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
22+
_GLIBCXX_BEGIN_NAMESPACE_VERSION
23+
#endif
24+
#endif
25+
__device__
26+
__attribute__((__always_inline__, __visibility__("default"))) inline void
27+
__glibcxx_assert_fail() {}
28+
#ifdef _LIBCPP_END_NAMESPACE_STD
29+
_LIBCPP_END_NAMESPACE_STD
30+
#else
31+
#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
32+
_GLIBCXX_END_NAMESPACE_VERSION
33+
#endif
34+
} // namespace std
35+
#endif
36+
37+
#endif
38+
39+
#endif

0 commit comments

Comments
0 (0)
Morty Proxy This is a proxified and sanitized view of the page, visit original site.