Add hidden friends for access and reverse access

libcudacxx/include/cuda/__cccl_config

@@ -17,6 +17,7 @@
 #include <cuda/std/__cccl/dialect.h>
 #include <cuda/std/__cccl/execution_space.h>
 #include <cuda/std/__cccl/ptx_isa.h>
+#include <cuda/std/__cccl/sequence_access.h>
 #include <cuda/std/__cccl/system_header.h>
 #include <cuda/std/__cccl/version.h>
 #include <cuda/std/__cccl/visibility.h>

libcudacxx/include/cuda/std/__cccl/sequence_access.h

@@ -0,0 +1,77 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of libcu++, the C++ Standard Library for your entire system,
+// under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CCCL_SEQUENCE_ACCESS_H
+#define __CCCL_SEQUENCE_ACCESS_H
+
+#include <cuda/std/__cccl/compiler.h>
+#include <cuda/std/__cccl/system_header.h>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+// We need to define hidden friends for {cr,r,}{begin,end} of our containers as we will otherwise encounter ambigouities
+#define _CCCL_SYNTHESIZE_SEQUENCE_ACCESS(_ClassName, _ConstIter)                                \
+  friend iterator begin(_ClassName& __sequence) noexcept(noexcept(__sequence.begin()))          \
+  {                                                                                             \
+    return __sequence.begin();                                                                  \
+  }                                                                                             \
+  friend _ConstIter begin(const _ClassName& __sequence) noexcept(noexcept(__sequence.begin()))  \
+  {                                                                                             \
+    return __sequence.begin();                                                                  \
+  }                                                                                             \
+  friend iterator end(_ClassName& __sequence) noexcept(noexcept(__sequence.end()))              \
+  {                                                                                             \
+    return __sequence.end();                                                                    \
+  }                                                                                             \
+  friend _ConstIter end(const _ClassName& __sequence) noexcept(noexcept(__sequence.end()))      \
+  {                                                                                             \
+    return __sequence.end();                                                                    \
+  }                                                                                             \
+  friend _ConstIter cbegin(const _ClassName& __sequence) noexcept(noexcept(__sequence.begin())) \
+  {                                                                                             \
+    return __sequence.begin();                                                                  \
+  }                                                                                             \
+  friend _ConstIter cend(const _ClassName& __sequence) noexcept(noexcept(__sequence.end()))     \
+  {                                                                                             \
+    return __sequence.end();                                                                    \
+  }
+#define _CCCL_SYNTHESIZE_SEQUENCE_REVERSE_ACCESS(_ClassName, _ConstIter)                          \
+  friend reverse_iterator rbegin(_ClassName& __sequence) noexcept(noexcept(__sequence.rbegin()))  \
+  {                                                                                               \
+    return __sequence.rbegin();                                                                   \
+  }                                                                                               \
+  friend _ConstIter rbegin(const _ClassName& __sequence) noexcept(noexcept(__sequence.rbegin()))  \
+  {                                                                                               \
+    return __sequence.rbegin();                                                                   \
+  }                                                                                               \
+  friend reverse_iterator rend(_ClassName& __sequence) noexcept(noexcept(__sequence.rend()))      \
+  {                                                                                               \
+    return __sequence.rend();                                                                     \
+  }                                                                                               \
+  friend _ConstIter rend(const _ClassName& __sequence) noexcept(noexcept(__sequence.rend()))      \
+  {                                                                                               \
+    return __sequence.rend();                                                                     \
+  }                                                                                               \
+  friend _ConstIter crbegin(const _ClassName& __sequence) noexcept(noexcept(__sequence.rbegin())) \
+  {                                                                                               \
+    return __sequence.rbegin();                                                                   \
+  }                                                                                               \
+  friend _ConstIter crend(const _ClassName& __sequence) noexcept(noexcept(__sequence.rend()))     \
+  {                                                                                               \
+    return __sequence.rend();                                                                     \
+  }
+
+#endif // __CCCL_SEQUENCE_ACCESS_H

libcudacxx/include/cuda/std/__ranges/subrange.h

@@ -316,7 +316,7 @@ class _LIBCUDACXX_TEMPLATE_VIS subrange : public view_interface<subrange<_Iter,
 
   _LIBCUDACXX_TEMPLATE(class _It = _Iter)
   _LIBCUDACXX_REQUIRES(copyable<_It>)
-  _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr _It begin() const
+  _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr _It begin() const
   {
     return __begin_;
   }
@@ -328,7 +328,7 @@ class _LIBCUDACXX_TEMPLATE_VIS subrange : public view_interface<subrange<_Iter,
     return _CUDA_VSTD::move(__begin_);
   }
 
-  _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr _Sent end() const
+  _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr _Sent end() const
   {
     return __end_;
   }

libcudacxx/include/cuda/std/__ranges/view_interface.h

@@ -33,6 +33,7 @@
 #include <cuda/std/__type_traits/is_class.h>
 #include <cuda/std/__type_traits/is_reference.h>
 #include <cuda/std/__type_traits/remove_cv.h>
+#include <cuda/std/__utility/declval.h>
 #include <cuda/std/detail/libcxx/include/__assert>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_RANGES
@@ -177,6 +178,42 @@ class view_interface
   {
     return _CUDA_VRANGES::begin(__derived())[__index];
   }
+
+  template <class _D2 = _Derived, class _Ret = decltype(_CUDA_VRANGES::begin(_CUDA_VSTD::declval<_D2&>()))>
+  _CCCL_NODISCARD_FRIEND constexpr _Ret begin(_Derived& __range)
+  {
+    return _CUDA_VRANGES::begin(__range);
+  }
+
+  template <class _D2 = _Derived, class _Ret = decltype(_CUDA_VRANGES::begin(_CUDA_VSTD::declval<const _D2&>()))>
+  _CCCL_NODISCARD_FRIEND constexpr _Ret begin(const _Derived& __range)
+  {
+    return _CUDA_VRANGES::begin(__range);
+  }
+
+  template <class _D2 = _Derived, class _Ret = decltype(_CUDA_VRANGES::cbegin(_CUDA_VSTD::declval<const _D2&>()))>
+  _CCCL_NODISCARD_FRIEND constexpr auto cbegin(const _Derived& __range)
+  {
+    return _CUDA_VRANGES::cbegin(__range);
+  }
+
+  template <class _D2 = _Derived, class _Ret = decltype(_CUDA_VRANGES::end(_CUDA_VSTD::declval<_D2&>()))>
+  _CCCL_NODISCARD_FRIEND constexpr auto end(_Derived& __range)
+  {
+    return _CUDA_VRANGES::end(__range);
+  }
+
+  template <class _D2 = _Derived, class _Ret = decltype(_CUDA_VRANGES::end(_CUDA_VSTD::declval<const _D2&>()))>
+  _CCCL_NODISCARD_FRIEND constexpr auto end(const _Derived& __range)
+  {
+    return _CUDA_VRANGES::end(__range);
+  }
+
+  template <class _D2 = _Derived, class _Ret = decltype(_CUDA_VRANGES::cend(_CUDA_VSTD::declval<const _D2&>()))>
+  _CCCL_NODISCARD_FRIEND constexpr auto cend(const _Derived& __range)
+  {
+    return _CUDA_VRANGES::cend(__range);
+  }
 };
 
 _LIBCUDACXX_END_NAMESPACE_RANGES_ABI

libcudacxx/include/cuda/std/detail/libcxx/include/array

@@ -323,6 +323,9 @@ struct _LIBCUDACXX_TEMPLATE_VIS array
   {
     return __elems_;
   }
+
+  _CCCL_SYNTHESIZE_SEQUENCE_ACCESS(array, const_iterator);
+  _CCCL_SYNTHESIZE_SEQUENCE_REVERSE_ACCESS(array, const_reverse_iterator);
 };
 
 template <class _Tp>
@@ -498,6 +501,9 @@ struct _LIBCUDACXX_TEMPLATE_VIS array<_Tp, 0>
     return *data();
   }
   _CCCL_DIAG_POP
+
+  _CCCL_SYNTHESIZE_SEQUENCE_ACCESS(array, const_iterator);
+  _CCCL_SYNTHESIZE_SEQUENCE_REVERSE_ACCESS(array, const_reverse_iterator);
 };
 
 #if _CCCL_STD_VER >= 2017

libcudacxx/include/cuda/std/detail/libcxx/include/span

@@ -503,6 +503,9 @@ public:
     return span<byte, _Extent * sizeof(element_type)>{reinterpret_cast<byte*>(data()), size_bytes()};
   }
 
+  _CCCL_SYNTHESIZE_SEQUENCE_ACCESS(span, iterator);
+  _CCCL_SYNTHESIZE_SEQUENCE_REVERSE_ACCESS(span, reverse_iterator);
+
 private:
   pointer __data_;
 };
@@ -730,6 +733,9 @@ public:
     return {reinterpret_cast<byte*>(data()), size_bytes()};
   }
 
+  _CCCL_SYNTHESIZE_SEQUENCE_ACCESS(span, iterator);
+  _CCCL_SYNTHESIZE_SEQUENCE_REVERSE_ACCESS(span, reverse_iterator);
+
 private:
   pointer __data_;
   size_type __size_;

thrust/thrust/detail/contiguous_storage.h

@@ -140,6 +140,8 @@ class contiguous_storage
   // on move assignment
   _CCCL_HOST_DEVICE contiguous_storage& operator=(contiguous_storage&& other);
 
+  _CCCL_SYNTHESIZE_SEQUENCE_ACCESS(contiguous_storage, const_iterator);
+
 private:
   // XXX we could inherit from this to take advantage of empty base class optimization
   allocator_type m_allocator;

thrust/thrust/detail/range/head_flags.h

@@ -122,6 +122,8 @@ class head_flags_with_init
     return *(begin() + i);
   }
 
+  _CCCL_SYNTHESIZE_SEQUENCE_ACCESS(head_flags_with_init, iterator);
+
 private:
   iterator m_begin, m_end;
 };
@@ -202,6 +204,8 @@ class head_flags
     return *(begin() + i);
   }
 
+  _CCCL_SYNTHESIZE_SEQUENCE_ACCESS(head_flags, iterator);
+
 private:
   iterator m_begin, m_end;
 };

thrust/thrust/detail/range/tail_flags.h

@@ -107,6 +107,8 @@ class tail_flags
     return *(begin() + i);
   }
 
+  _CCCL_SYNTHESIZE_SEQUENCE_ACCESS(tail_flags, iterator);
+
 private:
   iterator m_begin, m_end;
 };

thrust/thrust/detail/trivial_sequence.h

@@ -65,10 +65,20 @@ struct _trivial_sequence<Iterator, DerivedPolicy, thrust::detail::true_type>
     return first;
   }
 
+  _CCCL_HOST_DEVICE friend iterator_type begin(_trivial_sequence& sequence)
+  {
+    return sequence.first;
+  }
+
   _CCCL_HOST_DEVICE iterator_type end()
   {
     return last;
   }
+
+  _CCCL_HOST_DEVICE friend iterator_type end(_trivial_sequence& sequence)
+  {
+    return sequence.last;
+  }
 };
 
 // non-trivial case
@@ -89,10 +99,20 @@ struct _trivial_sequence<Iterator, DerivedPolicy, thrust::detail::false_type>
     return buffer.begin();
   }
 
+  _CCCL_HOST_DEVICE friend iterator_type begin(_trivial_sequence& sequence)
+  {
+    return sequence.begin();
+  }
+
   _CCCL_HOST_DEVICE iterator_type end()
   {
     return buffer.end();
   }
+
+  _CCCL_HOST_DEVICE friend iterator_type end(_trivial_sequence& sequence)
+  {
+    return sequence.end();
+  }
 };
 
 template <typename Iterator, typename DerivedPolicy>

thrust/thrust/detail/vector_base.h

@@ -472,6 +472,9 @@ class vector_base
    */
   allocator_type get_allocator(void) const;
 
+  _CCCL_SYNTHESIZE_SEQUENCE_ACCESS(vector_base, const_iterator);
+  _CCCL_SYNTHESIZE_SEQUENCE_REVERSE_ACCESS(vector_base, const_reverse_iterator);
+
 protected:
   // Our storage
   storage_type m_storage;