llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Juan Manuel Martinez Caamaño (jmmartinez) <details> <summary>Changes</summary> 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(). --- Full diff: https://github.com/llvm/llvm-project/pull/136133.diff 2 Files Affected: - (modified) clang/lib/Headers/CMakeLists.txt (+1) - (added) clang/lib/Headers/cuda_wrappers/bits/c++config.h (+39) ``````````diff 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..583e595f7f529 --- /dev/null +++ b/clang/lib/Headers/cuda_wrappers/bits/c++config.h @@ -0,0 +1,39 @@ +// 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> + +#if _GLIBCXX_HAVE_IS_CONSTANT_EVALUATED + +#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() {} +#ifdef _LIBCPP_END_NAMESPACE_STD +_LIBCPP_END_NAMESPACE_STD +#else +#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION +_GLIBCXX_END_NAMESPACE_VERSION +#endif +} // namespace std +#endif + +#endif + +#endif `````````` </details> 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
