Add example loading offset via transform_iterator and extend doc

Author: bernhardmgruber

thrust/testing/cuda/offset_iterator.cu

@@ -1,7 +1,9 @@
 #include <cub/util_type.cuh>
 
 #include <thrust/distance.h>
+#include <thrust/functional.h>
 #include <thrust/iterator/offset_iterator.h>
+#include <thrust/iterator/transform_iterator.h>
 
 #include <cuda/std/iterator>
 
@@ -131,3 +133,18 @@ void TestOffsetIteratorIndirectValue()
   ASSERT_EQUAL(*iter, 2);
 }
 DECLARE_VECTOR_UNITTEST(TestOffsetIteratorIndirectValue);
+
+// this is not strictly a CUDA test, but it uses CUB
+template <typename Vector>
+void TestOffsetIteratorIndirectValueFancyIterator()
+{
+  using thrust::placeholders::_1;
+
+  Vector v{0, 1, 2, 3, 4, 5, 6, 7, 8};
+  thrust::device_vector<typename Vector::difference_type> offsets{2};
+  auto it = thrust::make_transform_iterator(offsets.begin(), _1 * 3);
+  thrust::offset_iterator iter(v.begin(), cub::FutureValue{it});
+
+  ASSERT_EQUAL(*iter, 6);
+}
+DECLARE_VECTOR_UNITTEST(TestOffsetIteratorIndirectValueFancyIterator);

thrust/thrust/iterator/offset_iterator.h

@@ -53,6 +53,33 @@ THRUST_NAMESPACE_BEGIN
 //!   return 0;
 //! }
 //! \endcode
+//!
+//! Combining a \p offset_iterator with \p cub::FutureValue can also be used to retrieve the offset from an iterator.
+//! However, such an \p offset_iterator cannot be moved anymore, since changes to the offset can not be written back.
+//!
+//! \code
+//! #include <thrust/iterator/offset_iterator.h>
+//! #include <thrust/fill.h>
+//! #include <thrust/functional.h>
+//! #include <thrust/device_vector.h>
+//! #include <cub/util_type.h>
+//!
+//! int main()
+//! {
+//!   using thrust::placeholders::_1;
+//!   thrust::device_vector<int> data{1, 2, 3, 4};
+//!
+//!   thrust::device_vector<ptrdiff> offsets{1}; // offset is only available on device
+//!   auto offset = cub::FutureValue{thrust::make_transform_iterator(offsets.begin(), _1 * 2)};
+//!   thrust::offset_iterator iter(v.begin(), offset); // load and transform offset upon access
+//!   // iter is at position 2 (= 1 * 2) in data, and would return 3 in device code
+//!
+//!   return 0;
+//! }
+//! \endcode
+//!
+//! In the above example, the offset is loaded from a device vector, transformed by a \p transform_iterator, and then
+//! applied to the underlying iterator, when the \p offset_iterator is accessed.
 template <typename Iterator, typename Offset = typename ::cuda::std::iterator_traits<Iterator>::difference_type>
 class offset_iterator : public iterator_adaptor<offset_iterator<Iterator, Offset>, Iterator>
 {