Allow mutation through a transform_iterator

But only if the transform iterator's base iterator returns a true l-value reference (and not a proxy reference).

Author: bernhardmgruber

thrust/testing/transform_iterator.cu

@@ -108,3 +108,33 @@ void TestTransformIteratorNonCopyable()
 }
 
 DECLARE_UNITTEST(TestTransformIteratorNonCopyable);
+
+struct foo
+{
+  int x, y;
+};
+
+struct access_x
+{
+  _CCCL_HOST_DEVICE int& operator()(foo& f) const noexcept
+  {
+    return f.x;
+  }
+};
+
+void TestTransformIteratorAsDestination()
+{
+  constexpr auto n = 10;
+  thrust::host_vector<int> src(n, 1234);
+  thrust::host_vector<foo> dst(n, foo{1, 2});
+
+  thrust::copy(src.begin(), src.end(), thrust::make_transform_iterator(dst.begin(), access_x{}));
+
+  for (const auto& f : dst)
+  {
+    ASSERT_EQUAL(f.x, 1234);
+    ASSERT_EQUAL(f.y, 2);
+  }
+}
+
+DECLARE_UNITTEST(TestTransformIteratorAsDestination);

thrust/thrust/detail/reference.h

@@ -308,7 +308,7 @@ class reference
   pointer const ptr;
 
   // `thrust::detail::is_wrapped_reference` is a trait that indicates whether
-  // a type is a fancy reference. It detects such types by loooking for a
+  // a type is a fancy reference. It detects such types by looking for a
   // nested `wrapped_reference_hint` type.
   struct wrapped_reference_hint
   {};

thrust/thrust/iterator/detail/transform_iterator.inl

@@ -45,10 +45,14 @@ template <class UnaryFunc, class Iterator, class Reference, class Value>
 struct transform_iterator_base
 {
 private:
+  using unary_func_input_t =
+    ::cuda::std::_If<::cuda::std::is_lvalue_reference<iterator_reference_t<Iterator>>::value,
+                     iterator_reference_t<Iterator>,
+                     const iterator_value_t<Iterator>&>;
+
   // By default, dereferencing the iterator yields the same as the function.
-  using reference = typename thrust::detail::ia_dflt_help<
-    Reference,
-    thrust::detail::result_of_adaptable_function<UnaryFunc(typename thrust::iterator_value<Iterator>::type)>>::type;
+  using reference = typename thrust::detail::
+    ia_dflt_help<Reference, thrust::detail::result_of_adaptable_function<UnaryFunc(unary_func_input_t)>>::type;
 
   // To get the default for Value: remove cvref on the result type.
   using value_type = typename thrust::detail::ia_dflt_help<Value, thrust::remove_cvref<reference>>::type;

thrust/thrust/iterator/transform_iterator.h

@@ -298,8 +298,25 @@ class transform_iterator
   // See goo.gl/LELTNp
   THRUST_DISABLE_MSVC_WARNING_BEGIN(4172)
 
-  _CCCL_EXEC_CHECK_DISABLE
   _CCCL_HOST_DEVICE typename super_t::reference dereference() const
+  {
+    // TODO(bgruber): I am not sure this is the correct check here. There is also the trait
+    // thrust::detail::is_wrapped_reference that sounds fitting. Only allowing to pass through l-value references
+    // strikes me as more conservative though.
+    // TODO(bgruber): use an if constexpr in C++17
+    return dereference_impl(::cuda::std::is_lvalue_reference<iterator_reference_t<Iterator>>{});
+  }
+
+  _CCCL_EXEC_CHECK_DISABLE
+  _CCCL_HOST_DEVICE
+  typename super_t::reference dereference_impl(::cuda::std::true_type /* iterator returns a T& */) const
+  {
+    return m_f(*this->base());
+  }
+
+  _CCCL_EXEC_CHECK_DISABLE
+  _CCCL_HOST_DEVICE
+  typename super_t::reference dereference_impl(::cuda::std::false_type /* iterator returns a proxy ref */) const
   {
     // Create a temporary to allow iterators with wrapped references to
     // convert to their value type before calling m_f. Note that this