From 66498203d99b085842e631cf833002d9202a9a4d Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Wed, 21 Feb 2024 09:51:35 +0100 Subject: [PATCH] Implement `ranges::ref_view` --- .../include/cuda/std/__ranges/ref_view.h | 125 +++++++ libcudacxx/include/cuda/std/ranges | 1 + .../range.ref.view/borrowing.compile.pass.cpp | 30 ++ .../range.ref.view/range.ref.view.pass.cpp | 318 ++++++++++++++++++ 4 files changed, 474 insertions(+) create mode 100644 libcudacxx/include/cuda/std/__ranges/ref_view.h create mode 100644 libcudacxx/test/libcudacxx/std/ranges/range.adaptors/range.all/range.ref.view/borrowing.compile.pass.cpp create mode 100644 libcudacxx/test/libcudacxx/std/ranges/range.adaptors/range.all/range.ref.view/range.ref.view.pass.cpp diff --git a/libcudacxx/include/cuda/std/__ranges/ref_view.h b/libcudacxx/include/cuda/std/__ranges/ref_view.h new file mode 100644 index 00000000000..4376f566438 --- /dev/null +++ b/libcudacxx/include/cuda/std/__ranges/ref_view.h @@ -0,0 +1,125 @@ +// -*- C++ -*- +//===----------------------------------------------------------------------===// +// +// Part of the LLVM Project, 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) 2025 NVIDIA CORPORATION & AFFILIATES +// +//===----------------------------------------------------------------------===// +#ifndef _LIBCUDACXX___RANGES_REF_VIEW_H +#define _LIBCUDACXX___RANGES_REF_VIEW_H + +#include + +#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 + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +_LIBCUDACXX_BEGIN_NAMESPACE_RANGES + +#if _CCCL_STD_VER >= 2017 && !_CCCL_COMPILER(MSVC2017) + +_LIBCUDACXX_BEGIN_NAMESPACE_RANGES_ABI + +template +struct __conversion_tester +{ + _LIBCUDACXX_HIDE_FROM_ABI static void __fun(_Range&); + static void __fun(_Range&&) = delete; +}; + +template +_CCCL_CONCEPT __convertible_to_lvalue = + _CCCL_REQUIRES_EXPR((_Tp, _Range))((__conversion_tester<_Range>::__fun(declval<_Tp>()))); + +# if _CCCL_STD_VER >= 2020 + +template + requires is_object_v<_Range> +# else // ^^^ C++20 ^^^ / vvv C++17 vvv +template , int> = 0, enable_if_t, int> = 0> +# endif // _CCCL_STD_VER <= 2017 +class ref_view : public view_interface> +{ + _Range* __range_; + +public: + _CCCL_TEMPLATE(class _Tp) + _CCCL_REQUIRES(__different_from<_Tp, ref_view> _CCCL_AND convertible_to<_Tp, _Range&> _CCCL_AND + __convertible_to_lvalue<_Tp, _Range>) + _LIBCUDACXX_HIDE_FROM_ABI constexpr ref_view(_Tp&& __t) + : view_interface>() + , __range_(_CUDA_VSTD::addressof(static_cast<_Range&>(_CUDA_VSTD::forward<_Tp>(__t)))) + {} + + _LIBCUDACXX_HIDE_FROM_ABI constexpr _Range& base() const + { + return *__range_; + } + + _LIBCUDACXX_HIDE_FROM_ABI constexpr iterator_t<_Range> begin() const + { + return _CUDA_VRANGES::begin(*__range_); + } + _LIBCUDACXX_HIDE_FROM_ABI constexpr sentinel_t<_Range> end() const + { + return _CUDA_VRANGES::end(*__range_); + } + + _CCCL_TEMPLATE(class _Range2 = _Range) + _CCCL_REQUIRES(invocable<_CUDA_VRANGES::__empty::__fn, const _Range2&>) + _LIBCUDACXX_HIDE_FROM_ABI constexpr bool empty() const + { + return _CUDA_VRANGES::empty(*__range_); + } + + _CCCL_TEMPLATE(class _Range2 = _Range) + _CCCL_REQUIRES(sized_range<_Range2>) + _LIBCUDACXX_HIDE_FROM_ABI constexpr auto size() const + { + return _CUDA_VRANGES::size(*__range_); + } + + _CCCL_TEMPLATE(class _Range2 = _Range) + _CCCL_REQUIRES(contiguous_range<_Range2>) + _LIBCUDACXX_HIDE_FROM_ABI constexpr auto data() const + { + return _CUDA_VRANGES::data(*__range_); + } +}; + +template +_CCCL_HOST_DEVICE ref_view(_Range&) -> ref_view<_Range>; + +_LIBCUDACXX_END_NAMESPACE_RANGES_ABI + +template +_CCCL_INLINE_VAR constexpr bool enable_borrowed_range> = true; + +#endif // _CCCL_STD_VER >= 2017 && !_CCCL_COMPILER_MSVC_2017 + +_LIBCUDACXX_END_NAMESPACE_RANGES + +#endif // _LIBCUDACXX___RANGES_REF_VIEW_H diff --git a/libcudacxx/include/cuda/std/ranges b/libcudacxx/include/cuda/std/ranges index 03f1438f456..3e46bb6f375 100644 --- a/libcudacxx/include/cuda/std/ranges +++ b/libcudacxx/include/cuda/std/ranges @@ -34,6 +34,7 @@ _CCCL_DIAG_SUPPRESS_MSVC(4848) #include #include #include +#include #include #include #include diff --git a/libcudacxx/test/libcudacxx/std/ranges/range.adaptors/range.all/range.ref.view/borrowing.compile.pass.cpp b/libcudacxx/test/libcudacxx/std/ranges/range.adaptors/range.all/range.ref.view/borrowing.compile.pass.cpp new file mode 100644 index 00000000000..9ac2e7847d9 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/ranges/range.adaptors/range.all/range.ref.view/borrowing.compile.pass.cpp @@ -0,0 +1,30 @@ +//===----------------------------------------------------------------------===// +// +// 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) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11, c++14 +// UNSUPPORTED: msvc-19.16 + +// template +// inline constexpr bool enable_borrowed_range> = true; + +#include + +#include "test_range.h" + +static_assert(cuda::std::ranges::borrowed_range>); +static_assert(cuda::std::ranges::borrowed_range>); +#if _LIBCUDACXX_HAS_RANGES +static_assert(cuda::std::ranges::borrowed_range>); +#endif + +int main(int, char**) +{ + return 0; +} diff --git a/libcudacxx/test/libcudacxx/std/ranges/range.adaptors/range.all/range.ref.view/range.ref.view.pass.cpp b/libcudacxx/test/libcudacxx/std/ranges/range.adaptors/range.all/range.ref.view/range.ref.view.pass.cpp new file mode 100644 index 00000000000..1d86cb9e2dd --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/ranges/range.adaptors/range.all/range.ref.view/range.ref.view.pass.cpp @@ -0,0 +1,318 @@ +//===----------------------------------------------------------------------===// +// +// 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) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// UNSUPPORTED: c++03, c++11, c++14 +// UNSUPPORTED: msvc-19.16 + +// template +// requires is_object_v +// class ref_view; + +#include +#include + +#include "test_iterators.h" +#include "test_macros.h" + +__device__ int globalBuff[8]; + +#if TEST_STD_VER >= 2020 +template +concept ValidRefView = requires { typename cuda::std::ranges::ref_view; }; +#else // ^^^ C++20 ^^^ / vvv C++17 vvv +template +constexpr bool ValidRefView = false; +template +constexpr bool ValidRefView>> = true; +#endif // TEST_STD_VER <= 2017 + +struct Range +{ + int start = 0; + __host__ __device__ friend constexpr int* begin(Range const& range) + { + return globalBuff + range.start; + } + __host__ __device__ friend constexpr int* end(Range const&) + { + return globalBuff + 8; + } + __host__ __device__ friend constexpr int* begin(Range& range) + { + return globalBuff + range.start; + } + __host__ __device__ friend constexpr int* end(Range&) + { + return globalBuff + 8; + } +}; + +struct BeginOnly +{ + __host__ __device__ friend int* begin(BeginOnly const&); + __host__ __device__ friend int* begin(BeginOnly&); +}; + +static_assert(ValidRefView); +static_assert(!ValidRefView); +static_assert(!ValidRefView); +static_assert(ValidRefView); + +static_assert(cuda::std::derived_from, + cuda::std::ranges::view_interface>>); + +struct RangeConvertible +{ + __host__ __device__ operator Range&(); +}; + +struct RValueRangeConvertible +{ + __host__ __device__ operator Range&&(); +}; + +static_assert(cuda::std::is_constructible_v, Range&>); +static_assert(cuda::std::is_constructible_v, RangeConvertible>); +static_assert(!cuda::std::is_constructible_v, RValueRangeConvertible>); + +struct ConstConvertibleToLValueAndRValue +{ + __host__ __device__ operator Range&() const; + __host__ __device__ operator Range&&() const; +}; +static_assert(cuda::std::is_convertible_v>); +static_assert(!cuda::std::is_convertible_v>); +static_assert(!cuda::std::is_convertible_v>); + +struct ForwardRange +{ + __host__ __device__ constexpr forward_iterator begin() const + { + return forward_iterator(globalBuff); + } + __host__ __device__ constexpr forward_iterator end() const + { + return forward_iterator(globalBuff + 8); + } +}; + +struct Cpp17InputRange +{ + struct sentinel + { + __host__ __device__ friend constexpr bool operator==(sentinel, cpp17_input_iterator iter) + { + return base(iter) == globalBuff + 8; + } +#if TEST_STD_VER <= 2017 + __host__ __device__ friend constexpr bool operator==(cpp17_input_iterator iter, sentinel) + { + return base(iter) == globalBuff + 8; + } + __host__ __device__ friend constexpr bool operator!=(sentinel, cpp17_input_iterator iter) + { + return base(iter) != globalBuff + 8; + } + __host__ __device__ friend constexpr bool operator!=(cpp17_input_iterator iter, sentinel) + { + return base(iter) != globalBuff + 8; + } +#endif // TEST_STD_VER <= 2017 + __host__ __device__ friend constexpr cuda::std::ptrdiff_t operator-(sentinel, cpp17_input_iterator) + { + return -8; + } + __host__ __device__ friend constexpr cuda::std::ptrdiff_t operator-(cpp17_input_iterator, sentinel) + { + return 8; + } + }; + + __host__ __device__ constexpr cpp17_input_iterator begin() const + { + return cpp17_input_iterator(globalBuff); + } + __host__ __device__ constexpr sentinel end() const + { + return {}; + } +}; + +struct Cpp20InputRange +{ + struct sentinel + { + __host__ __device__ friend constexpr bool operator==(sentinel, const cpp20_input_iterator& iter) + { + return base(iter) == globalBuff + 8; + } +#if TEST_STD_VER <= 2017 + __host__ __device__ friend constexpr bool operator==(const cpp20_input_iterator& iter, sentinel) + { + return base(iter) == globalBuff + 8; + } + __host__ __device__ friend constexpr bool operator!=(sentinel, const cpp20_input_iterator& iter) + { + return base(iter) != globalBuff + 8; + } + __host__ __device__ friend constexpr bool operator!=(const cpp20_input_iterator& iter, sentinel) + { + return base(iter) != globalBuff + 8; + } +#endif // TEST_STD_VER <= 2017 + __host__ __device__ friend constexpr cuda::std::ptrdiff_t operator-(sentinel, const cpp20_input_iterator&) + { + return -8; + } + }; + + __host__ __device__ constexpr cpp20_input_iterator begin() const + { + return cpp20_input_iterator(globalBuff); + } + __host__ __device__ constexpr sentinel end() const + { + return {}; + } +}; +template <> +inline constexpr bool cuda::std::ranges::enable_borrowed_range = true; + +#if TEST_STD_VER >= 2020 +template +concept EmptyIsInvocable = requires(cuda::std::ranges::ref_view view) { view.empty(); }; + +template +concept SizeIsInvocable = requires(cuda::std::ranges::ref_view view) { view.size(); }; + +template +concept DataIsInvocable = requires(cuda::std::ranges::ref_view view) { view.data(); }; +#else // ^^^ C++20 ^^^ / vvv C++17 vvv +template +_CCCL_CONCEPT_FRAGMENT(EmptyIsInvocable_, requires(cuda::std::ranges::ref_view view)((view.empty()))); +template +_CCCL_CONCEPT EmptyIsInvocable = _CCCL_FRAGMENT(EmptyIsInvocable_, R); + +template +_CCCL_CONCEPT_FRAGMENT(SizeIsInvocable_, requires(cuda::std::ranges::ref_view view)((view.size()))); +template +_CCCL_CONCEPT SizeIsInvocable = _CCCL_FRAGMENT(SizeIsInvocable_, R); + +template +_CCCL_CONCEPT_FRAGMENT(DataIsInvocable_, requires(cuda::std::ranges::ref_view view)((view.data()))); +template +_CCCL_CONCEPT DataIsInvocable = _CCCL_FRAGMENT(DataIsInvocable_, R); +#endif // TEST_STD_VER <= 2017 + +// Testing ctad. +static_assert(cuda::std::same_as())), + cuda::std::ranges::ref_view>); + +__host__ __device__ constexpr bool test() +{ + { + // ref_view::base + Range range{}; + cuda::std::ranges::ref_view view{range}; + assert(view.begin() == globalBuff); + view.base() = Range{2}; + assert(view.begin() == globalBuff + 2); + } + + { + // ref_view::begin + Range range1{}; + cuda::std::ranges::ref_view view1 = range1; + assert(view1.begin() == globalBuff); + + ForwardRange range2{}; + cuda::std::ranges::ref_view view2 = range2; + assert(base(view2.begin()) == globalBuff); + + Cpp17InputRange range3{}; + cuda::std::ranges::ref_view view3 = range3; + assert(base(view3.begin()) == globalBuff); + + Cpp20InputRange range4{}; + cuda::std::ranges::ref_view view4 = range4; + assert(base(view4.begin()) == globalBuff); + } + + { + // ref_view::end + Range range1{}; + cuda::std::ranges::ref_view view1 = range1; + assert(view1.end() == globalBuff + 8); + + ForwardRange range2{}; + cuda::std::ranges::ref_view view2 = range2; + assert(base(view2.end()) == globalBuff + 8); + + Cpp17InputRange range3{}; + cuda::std::ranges::ref_view view3 = range3; + assert(view3.end() == cpp17_input_iterator(globalBuff + 8)); + + Cpp20InputRange range4{}; + cuda::std::ranges::ref_view view4 = range4; + assert(view4.end() == cpp20_input_iterator(globalBuff + 8)); + } + + { + // ref_view::empty + Range range{8}; + cuda::std::ranges::ref_view view1 = range; + assert(view1.empty()); + + ForwardRange range2{}; + cuda::std::ranges::ref_view view2 = range2; + assert(!view2.empty()); + + static_assert(!EmptyIsInvocable); + static_assert(!EmptyIsInvocable); + } + + { + // ref_view::size + Range range1{8}; + cuda::std::ranges::ref_view view1 = range1; + assert(view1.size() == 0); + + Range range2{2}; + cuda::std::ranges::ref_view view2 = range2; + assert(view2.size() == 6); + + static_assert(!SizeIsInvocable); + } + + { + // ref_view::data + Range range1{}; + cuda::std::ranges::ref_view view1 = range1; + assert(view1.data() == globalBuff); + + Range range2{2}; + cuda::std::ranges::ref_view view2 = range2; + assert(view2.data() == globalBuff + 2); + + static_assert(!DataIsInvocable); + } + + return true; +} + +int main(int, char**) +{ + test(); +#if defined(_LIBCUDACXX_ADDRESSOF) + static_assert(test()); +#endif // _LIBCUDACXX_ADDRESSOF + + return 0; +}