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 b98abc3

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 b98abc3
Copy full SHA for b98abc3

File tree

2 files changed

+36
-0
lines changed
Filter options

2 files changed

+36
-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
+35Lines changed: 35 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,35 @@
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+
#ifdef _LIBCPP_BEGIN_NAMESPACE_STD
16+
_LIBCPP_BEGIN_NAMESPACE_STD
17+
#else
18+
namespace std {
19+
#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
20+
_GLIBCXX_BEGIN_NAMESPACE_VERSION
21+
#endif
22+
#endif
23+
__device__
24+
__attribute__((__always_inline__, __visibility__("default"))) inline void
25+
__glibcxx_assert_fail() {}
26+
#ifdef _LIBCPP_END_NAMESPACE_STD
27+
_LIBCPP_END_NAMESPACE_STD
28+
#else
29+
#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
30+
_GLIBCXX_END_NAMESPACE_VERSION
31+
#endif
32+
} // namespace std
33+
#endif
34+
35+
#endif

0 commit comments

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