Fix thrust::raw_reference_cast for tuple_of_iterator_references and simplify thrust::generate

thrust/testing/raw_reference_cast.cu

@@ -0,0 +1,48 @@
+#include <thrust/detail/raw_reference_cast.h>
+#include <thrust/device_vector.h>
+#include <thrust/iterator/zip_iterator.h>
+
+#include <unittest/unittest.h>
+
+void TestRawReferenceCast()
+{
+  using ::cuda::std::is_same_v;
+
+  {
+    [[maybe_unused]] int i        = 42;
+    [[maybe_unused]] const int ci = 42;
+    static_assert(is_same_v<decltype(thrust::raw_reference_cast(i)), int&>);
+    static_assert(is_same_v<decltype(thrust::raw_reference_cast(ci)), const int&>);
+  }
+  {
+    [[maybe_unused]] thrust::host_vector<int> vec(1);
+    static_assert(is_same_v<decltype(thrust::raw_reference_cast(*vec.begin())), int&>);
+    static_assert(is_same_v<decltype(thrust::raw_reference_cast(*vec.cbegin())), const int&>);
+
+    [[maybe_unused]] auto zip = thrust::make_zip_iterator(vec.begin(), vec.begin());
+    static_assert(
+      is_same_v<decltype(thrust::raw_reference_cast(*zip)), thrust::detail::tuple_of_iterator_references<int&, int&>>);
+
+    [[maybe_unused]] auto zip2 = thrust::make_zip_iterator(zip, zip);
+    static_assert(
+      is_same_v<decltype(thrust::raw_reference_cast(*zip2)),
+                thrust::detail::tuple_of_iterator_references<thrust::detail::tuple_of_iterator_references<int&, int&>,
+                                                             thrust::detail::tuple_of_iterator_references<int&, int&>>>);
+  }
+  {
+    [[maybe_unused]] thrust::device_vector<int> vec(1);
+    static_assert(is_same_v<decltype(thrust::raw_reference_cast(*vec.begin())), int&>);
+    static_assert(is_same_v<decltype(thrust::raw_reference_cast(*vec.cbegin())), const int&>);
+
+    [[maybe_unused]] auto zip = thrust::make_zip_iterator(vec.begin(), vec.begin());
+    static_assert(
+      is_same_v<decltype(thrust::raw_reference_cast(*zip)), thrust::detail::tuple_of_iterator_references<int&, int&>>);
+
+    [[maybe_unused]] auto zip2 = thrust::make_zip_iterator(zip, zip);
+    static_assert(
+      is_same_v<decltype(thrust::raw_reference_cast(*zip2)),
+                thrust::detail::tuple_of_iterator_references<thrust::detail::tuple_of_iterator_references<int&, int&>,
+                                                             thrust::detail::tuple_of_iterator_references<int&, int&>>>);
+  }
+}
+DECLARE_UNITTEST(TestRawReferenceCast);

thrust/thrust/detail/internal_functional.h

@@ -112,79 +112,6 @@ struct tuple_not_binary_predicate
   mutable Predicate pred;
 };
 
-template <typename Generator>
-struct host_generate_functor
-{
-  using result_type = void;
-
-  _CCCL_EXEC_CHECK_DISABLE
-  _CCCL_HOST_DEVICE host_generate_functor(Generator g)
-      : gen(g)
-  {}
-
-  // operator() does not take an lvalue reference because some iterators
-  // produce temporary proxy references when dereferenced. for example,
-  // consider the temporary tuple of references produced by zip_iterator.
-  // such temporaries cannot bind to an lvalue reference.
-  //
-  // to WAR this, accept a const reference (which is bindable to a temporary),
-  // and const_cast in the implementation.
-  //
-  // XXX change to an rvalue reference upon c++0x (which either a named variable
-  //     or temporary can bind to)
-  template <typename T>
-  _CCCL_HOST void operator()(const T& x)
-  {
-    // we have to be naughty and const_cast this to get it to work
-    T& lvalue = const_cast<T&>(x);
-
-    // this assigns correctly whether x is a true reference or proxy
-    lvalue = gen();
-  }
-
-  Generator gen;
-};
-
-template <typename Generator>
-struct device_generate_functor
-{
-  using result_type = void;
-
-  _CCCL_EXEC_CHECK_DISABLE
-  _CCCL_HOST_DEVICE device_generate_functor(Generator g)
-      : gen(g)
-  {}
-
-  // operator() does not take an lvalue reference because some iterators
-  // produce temporary proxy references when dereferenced. for example,
-  // consider the temporary tuple of references produced by zip_iterator.
-  // such temporaries cannot bind to an lvalue reference.
-  //
-  // to WAR this, accept a const reference (which is bindable to a temporary),
-  // and const_cast in the implementation.
-  //
-  // XXX change to an rvalue reference upon c++0x (which either a named variable
-  //     or temporary can bind to)
-  template <typename T>
-  _CCCL_HOST_DEVICE void operator()(const T& x)
-  {
-    // we have to be naughty and const_cast this to get it to work
-    T& lvalue = const_cast<T&>(x);
-
-    // this assigns correctly whether x is a true reference or proxy
-    lvalue = gen();
-  }
-
-  Generator gen;
-};
-
-template <typename System, typename Generator>
-struct generate_functor
-    : thrust::detail::eval_if<::cuda::std::is_convertible<System, thrust::host_system_tag>::value,
-                              thrust::detail::identity_<host_generate_functor<Generator>>,
-                              thrust::detail::identity_<device_generate_functor<Generator>>>
-{};
-
 template <typename T>
 struct is_non_const_reference
     : ::cuda::std::_And<thrust::detail::not_<::cuda::std::is_const<T>>,