Emit diagnostics for device lambdas

Author: gevtushenko

thrust/detail/type_traits/result_of_adaptable_function.h

@@ -20,6 +20,8 @@
 #include <thrust/detail/type_traits.h>
 #include <thrust/detail/type_traits/function_traits.h>
 
+#include <cub/detail/type_traits.cuh>
+
 #include <type_traits>
 
 THRUST_NAMESPACE_BEGIN
@@ -29,7 +31,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 +40,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 = CUB_NS_QUALIFIER::detail::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,9 +17,10 @@
 #pragma once
 
 #include <thrust/detail/config.h>
-
 #include <thrust/iterator/iterator_adaptor.h>
 
+#include <cub/detail/type_traits.cuh>
+
 THRUST_NAMESPACE_BEGIN
 
 template <typename InputFunction, typename OutputFunction, typename Iterator>
@@ -35,12 +36,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 = cub::detail::invoke_result_t<InputFunction, iterator_value_type>;
 
   public:
     __host__ __device__
@@ -93,12 +89,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
+      , cub::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

@@ -16,6 +16,8 @@
 #include <thrust/detail/config.h>
 #include <thrust/detail/cpp11_required.h>
 
+#include <cub/detail/type_traits.cuh>
+
 #if THRUST_CPP_DIALECT >= 2011
 
 #include <thrust/addressof.h>
@@ -256,19 +258,12 @@ 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> 
+struct invoke_result
+{
+  using type = CUB_NS_QUALIFIER::detail::invoke_result_t<F, 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

thrust/system/cuda/detail/transform_scan.h

@@ -33,6 +33,8 @@
 #include <thrust/system/cuda/detail/scan.h>
 #include <thrust/distance.h>
 
+#include <cub/detail/type_traits.cuh>
+
 THRUST_NAMESPACE_BEGIN
 
 namespace cuda_cub {
@@ -52,12 +54,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 = CUB_NS_QUALIFIER::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

@@ -25,6 +25,8 @@
 #include <thrust/detail/type_traits/function_traits.h>
 #include <thrust/detail/type_traits/iterator/is_output_iterator.h>
 
+#include <cub/detail/type_traits.cuh>
+
 THRUST_NAMESPACE_BEGIN
 namespace system
 {
@@ -49,11 +51,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 = CUB_NS_QUALIFIER::detail::invoke_result_t<UnaryFunction, InputType>;
   using ValueType = typename std::remove_reference<ResultType>::type;
 
   thrust::transform_iterator<UnaryFunction, InputIterator, ValueType> _first(first, unary_op);