Extend cub::FutureValue (#4075)

Co-authored-by: Michael Schellenberger Costa <miscco@nvidia.com>

Author: bernhardmgruber

cub/cub/util_type.cuh

@@ -262,10 +262,12 @@ struct FutureValue
 {
   using value_type    = T;
   using iterator_type = IterT;
+
   explicit _CCCL_HOST_DEVICE _CCCL_FORCEINLINE FutureValue(IterT iter)
       : m_iter(iter)
   {}
-  _CCCL_HOST_DEVICE _CCCL_FORCEINLINE operator T()
+
+  _CCCL_HOST_DEVICE _CCCL_FORCEINLINE operator T() const noexcept
   {
     return *m_iter;
   }
@@ -274,6 +276,9 @@ private:
   IterT m_iter;
 };
 
+template <typename IterT>
+FutureValue(IterT) -> FutureValue<detail::it_value_t<IterT>, IterT>;
+
 namespace detail
 {
 

cub/test/catch2_test_util_type.cu

@@ -107,3 +107,23 @@ C2H_TEST("Test CustomHalf", "[util][type]")
   CHECK(cuda::std::numeric_limits<half_t>::max() == half_t::max());
   CHECK(cuda::std::numeric_limits<half_t>::lowest() == half_t::lowest());
 }
+
+C2H_TEST("Test FutureValue", "[util][type]")
+{
+  // read
+  int value;
+  cub::FutureValue<int> fv{&value};
+  value = 42;
+  CHECK(fv == 42);
+  value = 43;
+  CHECK(fv == 43);
+
+  // CTAD
+  cub::FutureValue fv2{&value};
+  STATIC_REQUIRE(::cuda::std::is_same_v<decltype(fv2), cub::FutureValue<int, int*>>);
+
+  c2h::device_vector<int> v(0);
+  cub::FutureValue fv3{v.begin()};
+  STATIC_REQUIRE(
+    ::cuda::std::is_same_v<decltype(fv3), cub::FutureValue<int, typename c2h::device_vector<int>::iterator>>);
+}