Emit diagnostics for device lambdas

thrust/detail/type_traits.h

@@ -1,5 +1,5 @@
 /*
- *  Copyright 2008-2018 NVIDIA Corporation
+ *  Copyright 2008-2022 NVIDIA Corporation
  *
  *  Licensed under the Apache License, Version 2.0 (the "License");
  *  you may not use this file except in compliance with the License.
@@ -24,9 +24,9 @@
 
 #include <thrust/detail/config.h>
 
-#if THRUST_CPP_DIALECT >= 2011
-#  include <type_traits>
-#endif
+#include <cuda/std/type_traits>
+
+#include <type_traits>
 
 THRUST_NAMESPACE_BEGIN
 
@@ -47,7 +47,6 @@ namespace detail
      // We don't want to switch to std::integral_constant, because we want access
      // to the C++14 operator(), but we'd like standard traits to interoperate
      // with our version when tag dispatching.
-     #if THRUST_CPP_DIALECT >= 2011
      integral_constant() = default;
 
      integral_constant(integral_constant const&) = default;
@@ -56,7 +55,6 @@ namespace detail
 
      constexpr __host__ __device__
      integral_constant(std::integral_constant<T, v>) noexcept {}
-     #endif
 
      constexpr __host__ __device__ operator value_type() const noexcept { return value; }
      constexpr __host__ __device__ value_type operator()() const noexcept { return value; }
@@ -715,6 +713,20 @@ template<typename T>
   {
   };
 
+template <typename Invokable, typename... Args>
+using invoke_result_t =
+#if THRUST_CPP_DIALECT < 2017
+  typename cuda::std::result_of<Invokable(Args...)>::type;
+#else // 2017+
+  cuda::std::invoke_result_t<Invokable, Args...>;
+#endif
+
+template <class F, class... Us> 
+struct invoke_result
+{
+  using type = invoke_result_t<F, Us...>;
+};
+
 } // end detail
 
 using detail::integral_constant;

thrust/detail/type_traits/result_of_adaptable_function.h

@@ -29,7 +29,6 @@ namespace detail
 // Sets `type` to the result of the specified Signature invocation. If the
 // callable defines a `result_type` alias member, that type is used instead.
 // Use invoke_result / result_of when FuncType::result_type is not defined.
-#if THRUST_CPP_DIALECT >= 2017
 template <typename Signature, typename Enable = void>
 struct result_of_adaptable_function
 {
@@ -39,16 +38,12 @@ struct result_of_adaptable_function
   template <typename F, typename...Args>
   struct impl<F(Args...)>
   {
-    using type = std::invoke_result_t<F, Args...>;
+    using type = invoke_result_t<F, Args...>;
   };
 
 public:
   using type = typename impl<Signature>::type;
 };
-#else // < C++17
-template <typename Signature, typename Enable = void>
-struct result_of_adaptable_function : std::result_of<Signature> {};
-#endif // < C++17
 
 // specialization for invocations which define result_type
 template <typename Functor, typename... ArgTypes>

thrust/iterator/detail/transform_input_output_iterator.inl

@@ -17,8 +17,8 @@
 #pragma once
 
 #include <thrust/detail/config.h>
-
 #include <thrust/iterator/iterator_adaptor.h>
+#include <thrust/detail/type_traits.h>
 
 THRUST_NAMESPACE_BEGIN
 
@@ -35,12 +35,7 @@ template <typename InputFunction, typename OutputFunction, typename Iterator>
 {
   using iterator_value_type = typename thrust::iterator_value<Iterator>::type;
 
-  // std::result_of is deprecated in 2017, replace with std::invoke_result
-#if THRUST_CPP_DIALECT < 2017
-  using Value = typename std::result_of<InputFunction(iterator_value_type)>::type;
-#else
-  using Value = std::invoke_result_t<InputFunction, iterator_value_type>;
-#endif
+  using Value = invoke_result_t<InputFunction, iterator_value_type>;
 
   public:
     __host__ __device__
@@ -93,12 +88,7 @@ public:
     <
         transform_input_output_iterator<InputFunction, OutputFunction, Iterator>
       , Iterator
-    // std::result_of is deprecated in 2017, replace with std::invoke_result
-#if THRUST_CPP_DIALECT < 2017
-      , typename std::result_of<InputFunction(iterator_value_type)>::type
-#else
-      , std::invoke_result_t<InputFunction, iterator_value_type>
-#endif
+      , detail::invoke_result_t<InputFunction, iterator_value_type>
       , thrust::use_default
       , thrust::use_default
       , transform_input_output_iterator_proxy<InputFunction, OutputFunction, Iterator>

thrust/optional.h

@@ -15,6 +15,7 @@
 
 #include <thrust/detail/config.h>
 #include <thrust/detail/cpp11_required.h>
+#include <thrust/detail/type_traits.h>
 
 #if THRUST_CPP_DIALECT >= 2011
 
@@ -255,22 +256,6 @@ constexpr auto invoke(Fn &&f, Args &&... args)
 {
   return std::forward<Fn>(f)(std::forward<Args>(args)...);
 }
-
-// std::invoke_result from C++17
-template <class F, class, class... Us> struct invoke_result_impl;
-
-template <class F, class... Us>
-struct invoke_result_impl<
-    F, decltype(detail::invoke(std::declval<F>(), std::declval<Us>()...), void()),
-    Us...> {
-  using type = decltype(detail::invoke(std::declval<F>(), std::declval<Us>()...));
-};
-
-template <class F, class... Us>
-using invoke_result = invoke_result_impl<F, void, Us...>;
-
-template <class F, class... Us>
-using invoke_result_t = typename invoke_result<F, Us...>::type;
 #endif
 
 // std::void_t from C++17

thrust/system/cuda/detail/transform_scan.h

@@ -30,8 +30,9 @@
 
 #if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
 #include <iterator>
-#include <thrust/system/cuda/detail/scan.h>
+#include <thrust/detail/type_traits.h>
 #include <thrust/distance.h>
+#include <thrust/system/cuda/detail/scan.h>
 
 THRUST_NAMESPACE_BEGIN
 
@@ -52,12 +53,7 @@ transform_inclusive_scan(execution_policy<Derived> &policy,
 {
   // Use the transformed input iterator's value type per https://wg21.link/P0571
   using input_type = typename thrust::iterator_value<InputIt>::type;
-#if THRUST_CPP_DIALECT < 2017
-  using result_type = typename std::result_of<TransformOp(input_type)>::type;
-#else
-  using result_type = std::invoke_result_t<TransformOp, input_type>;
-#endif
-
+  using result_type = thrust::detail::invoke_result_t<TransformOp, input_type>;
   using value_type = typename std::remove_reference<result_type>::type;
 
   typedef typename iterator_traits<InputIt>::difference_type size_type;

thrust/system/detail/generic/transform_scan.inl

@@ -49,11 +49,7 @@ __host__ __device__
 {
   // Use the input iterator's value type per https://wg21.link/P0571
   using InputType = typename thrust::iterator_value<InputIterator>::type;
-#if THRUST_CPP_DIALECT < 2017
-  using ResultType = typename std::result_of<UnaryFunction(InputType)>::type;
-#else
-  using ResultType = std::invoke_result_t<UnaryFunction, InputType>;
-#endif
+  using ResultType = thrust::detail::invoke_result_t<UnaryFunction, InputType>;
   using ValueType = typename std::remove_reference<ResultType>::type;
 
   thrust::transform_iterator<UnaryFunction, InputIterator, ValueType> _first(first, unary_op);