https://github.com/Artem-B created
https://github.com/llvm/llvm-project/pull/148918
Since #116709 more libc++ code relies on std::declval() and it broke some CUDA
compilations.
The new wrapper adds GPU-side overloads for the declval() helper functions
which allows it to continue working when used from CUDA sources.
>From ea1949d13608ac948ab34d1eeb073decdd11e2a3 Mon Sep 17 00:00:00 2001
From: Artem Belevich <t...@google.com>
Date: Tue, 15 Jul 2025 11:10:40 -0700
Subject: [PATCH] [CUDA] add wrapper header for libc++'s __utlility/declval.h
Since #116709 more libc++ code relies on std::declval() and it broke some CUDA
compilations.
The new wrapper adds GPU-side overloads for the declval() helper functions
which allows it to continue working when used from CUDA sources.
---
clang/lib/Headers/CMakeLists.txt | 22 +++++++++++--
.../Headers/cuda_wrappers/__utility/declval.h | 31 +++++++++++++++++++
2 files changed, 50 insertions(+), 3 deletions(-)
create mode 100644 clang/lib/Headers/cuda_wrappers/__utility/declval.h
diff --git a/clang/lib/Headers/CMakeLists.txt b/clang/lib/Headers/CMakeLists.txt
index c96d209c1fc0c..b4618fe4a46da 100644
--- a/clang/lib/Headers/CMakeLists.txt
+++ b/clang/lib/Headers/CMakeLists.txt
@@ -347,6 +347,10 @@ set(cuda_wrapper_bits_files
cuda_wrappers/bits/basic_string.tcc
)
+set(cuda_wrapper_utility_files
+ cuda_wrappers/__utility/declval.h
+)
+
set(ppc_wrapper_files
ppc_wrappers/mmintrin.h
ppc_wrappers/xmmintrin.h
@@ -443,8 +447,9 @@ endfunction(clang_generate_header)
# Copy header files from the source directory to the build directory
foreach( f ${files} ${cuda_wrapper_files} ${cuda_wrapper_bits_files}
- ${ppc_wrapper_files} ${openmp_wrapper_files} ${zos_wrapper_files}
${hlsl_files}
- ${llvm_libc_wrapper_files} ${llvm_offload_wrapper_files})
+ ${cuda_wrapper_utility_files} ${ppc_wrapper_files}
${openmp_wrapper_files}
+ ${zos_wrapper_files} ${hlsl_files} ${llvm_libc_wrapper_files}
+ ${llvm_offload_wrapper_files})
copy_header_to_output_dir(${CMAKE_CURRENT_SOURCE_DIR} ${f})
endforeach( f )
@@ -553,7 +558,7 @@ add_header_target("arm-common-resource-headers"
"${arm_common_files};${arm_commo
# Architecture/platform specific targets
add_header_target("arm-resource-headers"
"${arm_only_files};${arm_only_generated_files}")
add_header_target("aarch64-resource-headers"
"${aarch64_only_files};${aarch64_only_generated_files}")
-add_header_target("cuda-resource-headers"
"${cuda_files};${cuda_wrapper_files};${cuda_wrapper_bits_files}")
+add_header_target("cuda-resource-headers"
"${cuda_files};${cuda_wrapper_files};${cuda_wrapper_bits_files};${cuda_wrapper_utility_files}")
add_header_target("hexagon-resource-headers" "${hexagon_files}")
add_header_target("hip-resource-headers" "${hip_files}")
add_header_target("loongarch-resource-headers" "${loongarch_files}")
@@ -600,6 +605,11 @@ install(
DESTINATION ${header_install_dir}/cuda_wrappers/bits
COMPONENT clang-resource-headers)
+install(
+ FILES ${cuda_wrapper_utility_files}
+ DESTINATION ${header_install_dir}/cuda_wrappers/__utility
+ COMPONENT clang-resource-headers)
+
install(
FILES ${ppc_wrapper_files}
DESTINATION ${header_install_dir}/ppc_wrappers
@@ -663,6 +673,12 @@ install(
EXCLUDE_FROM_ALL
COMPONENT cuda-resource-headers)
+install(
+ FILES ${cuda_wrapper_utility_files}
+ DESTINATION ${header_install_dir}/cuda_wrappers/__utility
+ EXCLUDE_FROM_ALL
+ COMPONENT cuda-resource-headers)
+
install(
FILES ${cuda_files}
DESTINATION ${header_install_dir}
diff --git a/clang/lib/Headers/cuda_wrappers/__utility/declval.h
b/clang/lib/Headers/cuda_wrappers/__utility/declval.h
new file mode 100644
index 0000000000000..b16311e849fa4
--- /dev/null
+++ b/clang/lib/Headers/cuda_wrappers/__utility/declval.h
@@ -0,0 +1,31 @@
+#ifndef __CUDA_WRAPPERS_UTILITY_DECLVAL_H__
+#define __CUDA_WRAPPERS_UTILITY_DECLVAL_H__
+
+#include_next <__utility/declval.h>
+
+// The stuff below is the exact copy of the <__utility/declval.h>,
+// but with __device__ attribute applied to the functions, so it works on a
GPU.
+
+_LIBCPP_BEGIN_NAMESPACE_STD
+
+// Suppress deprecation notice for volatile-qualified return type resulting
+// from volatile-qualified types _Tp.
+_LIBCPP_SUPPRESS_DEPRECATED_PUSH
+template <class _Tp>
+__attribute__((device))
+_Tp&& __declval(int);
+template <class _Tp>
+__attribute__((device))
+_Tp __declval(long);
+_LIBCPP_SUPPRESS_DEPRECATED_POP
+
+template <class _Tp>
+__attribute__((device))
+_LIBCPP_HIDE_FROM_ABI decltype(std::__declval<_Tp>(0)) declval() _NOEXCEPT {
+ static_assert(!__is_same(_Tp, _Tp),
+ "std::declval can only be used in an unevaluated context. "
+ "It's likely that your current usage is trying to extract a
value from the function.");
+}
+
+_LIBCPP_END_NAMESPACE_STD
+#endif // __CUDA_WRAPPERS_UTILITY_DECLVAL_H__
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
