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

[via cfe-commits]

Author: Artem Belevich
Date: 2025-07-15T13:56:22-07:00
New Revision: 7f2bcd97cf4afc94e4badb2f6acd229003ef5912

URL: https://github.com/llvm/llvm-project/commit/7f2bcd97cf4afc94e4badb2f6acd229003ef5912
DIFF: https://github.com/llvm/llvm-project/commit/7f2bcd97cf4afc94e4badb2f6acd229003ef5912.diff

LOG: [CUDA] add wrapper header for libc++'s __utlility/declval.h (#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.

Added: 
    clang/lib/Headers/cuda_wrappers/__utility/declval.h

Modified: 
    clang/lib/Headers/CMakeLists.txt

Removed: 
    


################################################################################
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..a8434e3601939
--- /dev/null
+++ b/clang/lib/Headers/cuda_wrappers/__utility/declval.h
@@ -0,0 +1,28 @@
+#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__