[clang] [CUDA] add wrapper header for libc++'s __utlility/declval.h (PR #148918)

[Artem Belevich via cfe-commits]

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 <tra at 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__