[CUDA] add wrapper header for libc++'s __utlility/declval.h

Since llvm#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.

Author: Artem-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}

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__