Address review feedback

Author: miscco

libcudacxx/include/cuda/__iterator/transform_iterator.h

@@ -88,7 +88,7 @@ inline constexpr bool __transform_iterator_nothrow_subscript<_Fn, _Iter, true> =
 //! needed while saving both memory capacity and bandwidth.
 //!
 //! The following code snippet demonstrates how to create a \p transform_iterator which represents the result of
-//! \c sqrtf applied to the contents of a \p device_vector.
+//! \c sqrtf applied to the contents of a \p thrust::device_vector.
 //!
 //! @code
 //! #include <cuda/iterator>
@@ -105,11 +105,7 @@ inline constexpr bool __transform_iterator_nothrow_subscript<_Fn, _Iter, true> =
 //!
 //! int main()
 //! {
-//!   thrust::device_vector<float> v(4);
-//!   v[0] = 1.0f;
-//!   v[1] = 4.0f;
-//!   v[2] = 9.0f;
-//!   v[3] = 16.0f;
+//!   thrust::device_vector<float> v{1.0f, 4.0f, 9.0f, 16.0f};
 //!
 //!   using FloatIterator = thrust::device_vector<float>::iterator;
 //!
@@ -126,8 +122,8 @@ inline constexpr bool __transform_iterator_nothrow_subscript<_Fn, _Iter, true> =
 //! @endcode
 //!
 //! This next example demonstrates how to use a \p transform_iterator with the \p thrust::reduce function to compute the
-//! sum of squares of a sequence. We will create temporary \p transform_iterators with the \p make_transform_iterator
-//! function in order to avoid explicitly specifying their type:
+//! sum of squares of a sequence. We will create temporary \p transform_iterators utilising class template argument
+//! deduction avoid explicitly specifying their type:
 //!
 //! @code
 //! #include <cuda/iterator>
@@ -151,62 +147,21 @@ inline constexpr bool __transform_iterator_nothrow_subscript<_Fn, _Iter, true> =
 //!   v[0] = 1.0f;
 //!   v[1] = 2.0f;
 //!   v[2] = 3.0f;
-//!   v[3] = 4.0f;
-//!
-//!   float sum_of_squares =
-//!    thrust::reduce(cuda::make_transform_iterator(v.begin(), square{}),
-//!                   cuda::make_transform_iterator(v.end(),   square{}));
+//!   thrust::device_vector<float> v{1.0f, 2.0f, 3.0f, 4.0f};
+//!   thrust::reduce(cuda::transform_iterator{v.begin(), square{}},
+//!                  cuda::transform_iterator{v.end(),   square{}});
 //!
 //!   std::cout << "sum of squares: " << sum_of_squares << std::endl;
 //!   return 0;
 //! }
 //! @endcode
-//!
-//! The following example illustrates how to use the third template argument to explicitly specify the return type of
-//! the function.
-//!
-//! @code
-//! #include <cuda/iterator>
-//! #include <thrust/device_vector.h>
-//!
-//! struct square_root
-//! {
-//!   __host__ __device__
-//!   float operator()(float x) const
-//!   {
-//!     return sqrtf(x);
-//!   }
-//! };
-//!
-//! int main()
-//! {
-//!   thrust::device_vector<float> v(4);
-//!   v[0] = 1.0f;
-//!   v[1] = 4.0f;
-//!   v[2] = 9.0f;
-//!   v[3] = 16.0f;
-//!
-//!   using FloatIterator = thrust::device_vector<float>::iterator;
-//!
-//!   // note: float result_type is specified explicitly
-//!   cuda::transform_iterator iter(v.begin(), square_root{});
-//!
-//!   *iter;   // returns 1.0f
-//!   iter[0]; // returns 1.0f;
-//!   iter[1]; // returns 2.0f;
-//!   iter[2]; // returns 3.0f;
-//!   iter[3]; // returns 4.0f;
-//!
-//!   // iter[4] is an out-of-bounds error
-//! }
-//! @endcode
 template <class _Iter, class _Fn>
 class transform_iterator : public __transform_iterator_category_base<_Iter, _Fn>
 {
   static_assert(_CUDA_VSTD::is_object_v<_Fn>, "cuda::transform_iterator requires that _Fn is a function object");
   static_assert(_CUDA_VSTD::regular_invocable<_Fn&, _CUDA_VSTD::iter_reference_t<_Iter>>,
                 "cuda::transform_iterator requires that _Fn is invocable with iter_reference_t<_Iter>");
-  static_assert(_CUDA_VSTD::regular_invocable<_Fn&, _CUDA_VSTD::iter_reference_t<_Iter>>,
+  static_assert(_CUDA_VSTD::__can_reference<_CUDA_VSTD::invoke_result_t<_Fn&, _CUDA_VSTD::iter_reference_t<_Iter>>>,
                 "cuda::transform_iterator requires that the return type of _Fn is referenceable");
 
 public:
@@ -341,10 +296,7 @@ class transform_iterator : public __transform_iterator_category_base<_Iter, _Fn>
     return _CUDA_VSTD::invoke(const_cast<_Fn&>(*__func_), __current_[__n]);
   }
 
-  //! @brief Compares two \c transform_iterator \p __lhs and \p __rhs for equality
-  //! @param __lhs The left \c transform_iterator
-  //! @param __rhs The right \c transform_iterator
-  //! @return true if the stored iterators compares equal
+  //! @brief Compares two \c transform_iterator for equality, directly comparing the stored iterators
   template <class _Iter2 = _Iter>
   [[nodiscard]] _LIBCUDACXX_HIDE_FROM_ABI friend constexpr auto
   operator==(const transform_iterator& __lhs, const transform_iterator& __rhs) noexcept(
@@ -355,10 +307,7 @@ class transform_iterator : public __transform_iterator_category_base<_Iter, _Fn>
   }
 
 #if _CCCL_STD_VER <= 2017
-  //! @brief Compares two \c transform_iterator \p __lhs and \p __rhs for inequality
-  //! @param __lhs The left \c transform_iterator
-  //! @param __rhs The right \c transform_iterator
-  //! @return true if the stored iterators do not compare equal
+  //! @brief Compares two \c transform_iterator for inequality, directly comparing the stored iterators
   template <class _Iter2 = _Iter>
   [[nodiscard]] _LIBCUDACXX_HIDE_FROM_ABI friend constexpr auto
   operator!=(const transform_iterator& __lhs, const transform_iterator& __rhs) noexcept(
@@ -369,10 +318,7 @@ class transform_iterator : public __transform_iterator_category_base<_Iter, _Fn>
   }
 #endif // _CCCL_STD_VER <= 2017
 
-  //! @brief Compares two \c transform_iterator \p __lhs and \p __rhs for less than
-  //! @param __lhs The left \c transform_iterator
-  //! @param __rhs The right \c transform_iterator
-  //! @return true if the stored iterators compares less than
+  //! @brief Compares two \c transform_iterator for less then, directly comparing the stored iterators
   template <class _Iter2 = _Iter>
   [[nodiscard]] _LIBCUDACXX_HIDE_FROM_ABI friend constexpr auto
   operator<(const transform_iterator& __lhs, const transform_iterator& __rhs) noexcept(
@@ -382,10 +328,7 @@ class transform_iterator : public __transform_iterator_category_base<_Iter, _Fn>
     return __lhs.__current_ < __rhs.__current_;
   }
 
-  //! @brief Compares two \c transform_iterator \p __lhs and \p __rhs for greater than
-  //! @param __lhs The left \c transform_iterator
-  //! @param __rhs The right \c transform_iterator
-  //! @return true if the stored iterators compares greater than
+  //! @brief Compares two \c transform_iterator for greater then, directly comparing the stored iterators
   template <class _Iter2 = _Iter>
   [[nodiscard]] _LIBCUDACXX_HIDE_FROM_ABI friend constexpr auto
   operator>(const transform_iterator& __lhs, const transform_iterator& __rhs) noexcept(
@@ -395,10 +338,7 @@ class transform_iterator : public __transform_iterator_category_base<_Iter, _Fn>
     return __lhs.__current_ > __rhs.__current_;
   }
 
-  //! @brief Compares two \c transform_iterator \p __lhs and \p __rhs for less equal
-  //! @param __lhs The left \c transform_iterator
-  //! @param __rhs The right \c transform_iterator
-  //! @return true if the stored iterators compares less equal
+  //! @brief Compares two \c transform_iterator for less equal, directly comparing the stored iterators
   template <class _Iter2 = _Iter>
   [[nodiscard]] _LIBCUDACXX_HIDE_FROM_ABI friend constexpr auto
   operator<=(const transform_iterator& __lhs, const transform_iterator& __rhs) noexcept(
@@ -408,10 +348,7 @@ class transform_iterator : public __transform_iterator_category_base<_Iter, _Fn>
     return __lhs.__current_ <= __rhs.__current_;
   }
 
-  //! @brief Compares two \c transform_iterator \p __lhs and \p __rhs for greater equal
-  //! @param __lhs The left \c transform_iterator
-  //! @param __rhs The right \c transform_iterator
-  //! @return true if the stored iterators compares greater equal
+  //! @brief Compares two \c transform_iterator for greater equal, directly comparing the stored iterators
   template <class _Iter2 = _Iter>
   [[nodiscard]] _LIBCUDACXX_HIDE_FROM_ABI friend constexpr auto
   operator>=(const transform_iterator& __lhs, const transform_iterator& __rhs) noexcept(
@@ -422,10 +359,7 @@ class transform_iterator : public __transform_iterator_category_base<_Iter, _Fn>
   }
 
 #if _LIBCUDACXX_HAS_SPACESHIP_OPERATOR()
-  //! @brief Three-way-compares two \c transform_iterator \p __lhs and \p __rhs
-  //! @param __lhs The left \c transform_iterator
-  //! @param __rhs The right \c transform_iterator
-  //! @return The three-way ordering of stored iterators
+  //! @brief Three-way-compares two \c transform_iterator, directly three-way-comparing the stored iterators
   template <class _Iter2 = _Iter>
   [[nodiscard]] _LIBCUDACXX_HIDE_FROM_ABI friend constexpr auto
   operator<=>(const transform_iterator& __lhs, const transform_iterator& __rhs) noexcept(

libcudacxx/test/libcudacxx/cuda/iterators/transform_iterator/ctor.pass.cpp

@@ -21,12 +21,12 @@
 struct NoDefaultInitIter
 {
   int* ptr_;
-  typedef cuda::std::random_access_iterator_tag iterator_category;
-  typedef int value_type;
-  typedef cuda::std::ptrdiff_t difference_type;
-  typedef int* pointer;
-  typedef int& reference;
-  typedef NoDefaultInitIter self;
+  using iterator_category = cuda::std::random_access_iterator_tag;
+  using value_type        = int;
+  using difference_type   = cuda::std::ptrdiff_t;
+  using pointer           = int*;
+  using reference         = int&;
+  using self              = NoDefaultInitIter;
 
   __host__ __device__ constexpr NoDefaultInitIter(int* ptr)
       : ptr_(ptr)

libcudacxx/test/libcudacxx/cuda/iterators/transform_iterator/subscript.pass.cpp

@@ -12,6 +12,7 @@
 
 #include <cuda/iterator>
 #include <cuda/std/cassert>
+#include <cuda/std/utility>
 
 #include "test_iterators.h"
 #include "test_macros.h"
@@ -84,13 +85,25 @@ __host__ __device__ constexpr bool test()
   test<random_access_iterator<int*>>();
   test<int*>();
 
+  { // Enjusre that we can assign through projections
+    using pair    = cuda::std::pair<int, int>;
+    pair buffer[] = {{0, -1}, {1, -1}, {2, -1}, {3, -1}, {4, -1}, {5, -1}, {6, -1}, {7, -1}};
+    cuda::transform_iterator iter{buffer, &cuda::std::pair<int, int>::second};
+    assert(iter[4] == -1);
+    iter[4] = 42;
+    assert(iter[4] == 42);
+
+    const pair expected{4, 42};
+    assert(buffer[4] == expected);
+  }
+
   return true;
 }
 
 int main(int, char**)
 {
   test();
-  static_assert(test(), "");
+  // static_assert(test(), "");
 
   return 0;
 }