From 073d5766685fd8bb5b80c813eed14cda8abd87ce Mon Sep 17 00:00:00 2001 From: Jiaming Yuan Date: Tue, 14 Jul 2020 17:05:13 -0400 Subject: [PATCH 01/14] Implement mdarray and integrate mdspan. --- cpp/CMakeLists.txt | 6 +- cpp/cmake/thirdparty/get_mdspan.cmake | 15 + cpp/include/raft/mdarray.h | 502 ++++++++++++++++++++++++++ cpp/test/CMakeLists.txt | 1 + cpp/test/mdarray.cu | 289 +++++++++++++++ 5 files changed, 811 insertions(+), 2 deletions(-) create mode 100644 cpp/cmake/thirdparty/get_mdspan.cmake create mode 100644 cpp/include/raft/mdarray.h create mode 100644 cpp/test/mdarray.cu diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index f3a0f2d554..4951597f5b 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -111,6 +111,7 @@ include(cmake/thirdparty/get_rmm.cmake) include(cmake/thirdparty/get_cuco.cmake) include(cmake/thirdparty/get_libcudacxx.cmake) include(cmake/thirdparty/get_faiss.cmake) +include(cmake/thirdparty/get_mdspan.cmake) if(BUILD_TESTS) include(cmake/thirdparty/get_gtest.cmake) @@ -137,7 +138,8 @@ target_link_libraries(raft INTERFACE CUDA::cusparse $<$:CUDA::nvToolsExt> rmm::rmm - cuco::cuco) + cuco::cuco + std::mdspan) target_compile_definitions(raft INTERFACE $<$:NVTX_ENABLED>) target_compile_features(raft INTERFACE cxx_std_17 $) @@ -248,7 +250,7 @@ install(DIRECTORY include/raft/ # Temporary install of raft.hpp while the file is removed install(FILES include/raft.hpp - DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/raft) + DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/raft) ############################################################################## # - install export ----------------------------------------------------------- diff --git a/cpp/cmake/thirdparty/get_mdspan.cmake b/cpp/cmake/thirdparty/get_mdspan.cmake new file mode 100644 index 0000000000..0830205c9d --- /dev/null +++ b/cpp/cmake/thirdparty/get_mdspan.cmake @@ -0,0 +1,15 @@ +function(find_and_configure_mdspan VERSION) + rapids_cpm_find( + mdspan ${VERSION} + GLOBAL_TARGETS std::mdspan + BUILD_EXPORT_SET raft-exports + INSTALL_EXPORT_SET raft-exports + CPM_ARGS + GIT_REPOSITORY https://github.com/trivialfis/mdspan + GIT_TAG 0193f075e977cc5f3c957425fd899e53d598f524 + OPTIONS "MDSPAN_ENABLE_CUDA ON" + "MDSPAN_CXX_STANDARD ON" + ) +endfunction() + +find_and_configure_mdspan(0.2.0) diff --git a/cpp/include/raft/mdarray.h b/cpp/include/raft/mdarray.h new file mode 100644 index 0000000000..ce4f030876 --- /dev/null +++ b/cpp/include/raft/mdarray.h @@ -0,0 +1,502 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#include +#include +#include +#include +#include + +namespace raft { +namespace detail { +/** + * @brief A simplified version of thrust::device_reference with support for CUDA stream. + */ +template +class device_reference { + public: + using value_type = typename std::remove_cv_t; + using pointer = thrust::device_ptr; + using const_pointer = thrust::device_ptr; + + private: + std::conditional_t::value, const_pointer, pointer> ptr_; + rmm::cuda_stream_view stream_; + + public: + device_reference(thrust::device_ptr ptr, rmm::cuda_stream_view stream) + : ptr_{ptr}, stream_{stream} + { + } + + operator value_type() const // NOLINT + { + auto* raw = ptr_.get(); + value_type v{}; + update_host(&v, raw, 1, stream_); + return v; + } + auto operator=(T const& other) -> device_reference& + { + auto* raw = ptr_.get(); + update_device(raw, &other, 1, stream_); + return *this; + } +}; + +/** + * @brief A thin wrapper over rmm::device_uvector for implementing the mdarray container policy. + * + */ +template +class device_uvector { + rmm::device_uvector data_; + + public: + using value_type = T; + using size_type = std::size_t; + + using reference = device_reference; + using const_reference = device_reference; + + using pointer = value_type*; + using const_pointer = value_type const*; + + using iterator = pointer; + using const_iterator = const_pointer; + + public: + ~device_uvector() = default; + device_uvector(device_uvector&&) noexcept = default; + device_uvector(device_uvector const& that) : data_{that.data_, that.data_.stream()} {} + + auto operator=(device_uvector const& that) -> device_uvector& + { + data_ = rmm::device_uvector{that.data_, that.data_.stream()}; + return *this; + } + auto operator=(device_uvector&& that) noexcept -> device_uvector& = default; + + /** + * @brief Default ctor is deleted as it doesn't accept stream. + */ + device_uvector() = delete; + /** + * @brief Ctor that accepts a size, stream and an optional mr. + */ + explicit device_uvector( + std::size_t size, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) + : data_{size, stream, mr} + { + } + /** + * @brief Index operator that returns a proxy to the actual data. + */ + template + auto operator[](Index i) noexcept -> reference + { + return device_reference{thrust::device_ptr{data_.data() + i}, data_.stream()}; + } + /** + * @brief Index operator that returns a proxy to the actual data. + */ + template + auto operator[](Index i) const noexcept + { + return device_reference{thrust::device_ptr{data_.data() + i}, data_.stream()}; + } + + [[nodiscard]] auto data() noexcept -> pointer { return data_.data(); } + [[nodiscard]] auto data() const noexcept -> const_pointer { return data_.data(); } +}; + +/** + * @brief A container policy for device mdarray. + */ +template +class device_uvector_policy { + rmm::cuda_stream_view stream_; + + public: + using element_type = ElementType; + using container_type = device_uvector; + // FIXME(jiamingy): allocator type is not supported by rmm::device_uvector + using pointer = typename container_type::pointer; + using const_pointer = typename container_type::const_pointer; + using reference = device_reference; + using const_reference = device_reference; + + using accessor_policy = std::experimental::default_accessor; + using const_accessor_policy = std::experimental::default_accessor; + + public: + auto create(size_t n) -> container_type { return container_type(n, stream_); } + + device_uvector_policy() = delete; + explicit device_uvector_policy(rmm::cuda_stream_view stream) noexcept( + std::is_nothrow_copy_constructible_v) + : stream_{stream} + { + } + + [[nodiscard]] constexpr auto access(container_type& c, size_t n) const noexcept -> reference + { + return c[n]; + } + [[nodiscard]] constexpr auto access(container_type const& c, size_t n) const noexcept + -> const_reference + { + return c[n]; + } + + [[nodiscard]] auto make_accessor_policy() noexcept { return accessor_policy{}; } + [[nodiscard]] auto make_accessor_policy() const noexcept { return const_accessor_policy{}; } +}; + +/** + * @brief A container policy for host mdarray. + */ +template > +class host_vector_policy { + public: + using element_type = ElementType; + using container_type = std::vector; + using allocator_type = typename container_type::allocator_type; + using pointer = typename container_type::pointer; + using const_pointer = typename container_type::const_pointer; + using reference = element_type&; + using const_reference = element_type const&; + using accessor_policy = std::experimental::default_accessor; + using const_accessor_policy = std::experimental::default_accessor; + + public: + auto create(size_t n) -> container_type { return container_type(n); } + + constexpr host_vector_policy() noexcept(std::is_nothrow_default_constructible_v) = + default; + explicit constexpr host_vector_policy(rmm::cuda_stream_view) noexcept( + std::is_nothrow_default_constructible_v) + : host_vector_policy() + { + } + + [[nodiscard]] constexpr auto access(container_type& c, size_t n) const noexcept -> reference + { + return c[n]; + } + [[nodiscard]] constexpr auto access(container_type const& c, size_t n) const noexcept + -> const_reference + { + return c[n]; + } + + [[nodiscard]] auto make_accessor_policy() noexcept { return accessor_policy{}; } + [[nodiscard]] auto make_accessor_policy() const noexcept { return const_accessor_policy{}; } +}; + +/** + * @brief A mixin to distinguish host and device memory. + */ +template +struct accessor_mixin : public AccessorPolicy { + using accessor_type = AccessorPolicy; + using is_host_type = std::conditional_t; + // make sure the explicit ctor can fall through + using AccessorPolicy::AccessorPolicy; + accessor_mixin(AccessorPolicy const& that) : AccessorPolicy{that} {} // NOLINT +}; + +template +using host_accessor = accessor_mixin; + +template +using device_accessor = accessor_mixin; +} // namespace detail + +namespace stdex = std::experimental; + +/** + * @brief stdex::mdspan with device tag to avoid accessing incorrect memory location. + */ +template > +using device_mdspan = + stdex::mdspan>; + +/** + * @brief stdex::mdspan with host tag to avoid accessing incorrect memory location. + */ +template > +using host_mdspan = + stdex::mdspan>; + +/** + * @brief Modified from the c++ mdarray proposal + * + * https://isocpp.org/files/papers/D1684R0.html + * + * mdarray is a container type for mdspan with similar template arguments. However there + * are some inconsistencies in between them. We have made some modificiations to fit our + * needs, which are listed below. + * + * - Layout policy is different, the mdarray in raft uses `stdex::extent` directly just + * like `mdspan`, while the `mdarray` in the reference implementation uses varidic + * template. + * + * - Most of the constructors from the reference implementation is removed to make sure + * CUDA stream is honorred. + * + * - unique_size is not implemented, which is still working in progress in the proposal + * + * - For container policy, we adopt the alternative approach documented in the proposal + * [sec 2.4.3], which requires an additional make_accessor method for it to be used in + * mdspan. The container policy reference implementation has multiple `access` methods + * that accommodate needs for both mdarray and mdspan. This is more difficult for us + * since the policy might contain states that are unwanted inside a CUDA kernel. Also, + * on host we return a proxy to the actual value as `device_ref` so different access + * methods will have different return type, which is less desirable. + * + * - For the above reasons, copying from other mdarray with different policy type is also + * removed. + */ +template +class mdarray { + static_assert(!std::is_const::value, + "Element type for container must not be const."); + + public: + using extents_type = Extents; + using layout_type = LayoutPolicy; + using mapping_type = typename layout_type::template mapping; + using element_type = ElementType; + + using value_type = std::remove_cv_t; + using index_type = std::size_t; + using difference_type = std::ptrdiff_t; + // Naming: ref impl: container_policy_type, proposal: container_policy + using container_policy_type = ContainerPolicy; + using container_type = typename container_policy_type::container_type; + + using pointer = typename container_policy_type::pointer; + using const_pointer = typename container_policy_type::const_pointer; + using reference = typename container_policy_type::reference; + using const_reference = typename container_policy_type::const_reference; + + private: + template , + typename container_policy_type::const_accessor_policy, + typename container_policy_type::accessor_policy>> + using view_type_impl = + std::conditional_t, + device_mdspan>; + + public: + /** + * \brief the mdspan type returned by view method. + */ + using view_type = view_type_impl; + using const_view_type = view_type_impl; + + public: + constexpr mdarray() noexcept(std::is_nothrow_default_constructible_v) + : cp_{rmm::cuda_stream_default}, c_{cp_.create(0)} {}; + constexpr mdarray(mdarray const&) noexcept(std::is_nothrow_copy_constructible_v) = + default; + constexpr mdarray(mdarray&&) noexcept(std::is_nothrow_move_constructible::value) = + default; + + constexpr auto operator =(mdarray const&) noexcept( + std::is_nothrow_copy_assignable::value) -> mdarray& = default; + constexpr auto operator =(mdarray&&) noexcept( + std::is_nothrow_move_assignable::value) -> mdarray& = default; + + ~mdarray() noexcept(std::is_nothrow_destructible::value) = default; + +#ifndef RAFT_MDARRAY_CTOR_CONSTEXPR +#if !(__CUDACC_VER_MAJOR__ == 11 && __CUDACC_VER_MINOR__ <= 2) +// 11.0: +// Error: Internal Compiler Error (codegen): "there was an error in verifying the lgenfe output!" +// +// 11.2: +// Call parameter type does not match function signature! +// i8** null +// i8* %call14 = call i32 null(void (i8*)* null, i8* null, i8** null), !dbg !1060 +// : parse Invalid record (Producer: 'LLVM7.0.1' Reader: 'LLVM 7.0.1') +#define RAFT_MDARRAY_CTOR_CONSTEXPR constexpr +#else +#define RAFT_MDARRAY_CTOR_CONSTEXPR +#endif // !(__CUDACC_VER_MAJOR__ == 11 && __CUDACC_VER_MINOR__ <= 2) +#endif // RAFT_MDARRAY_CTOR_CONSTEXPR + + /** + * @brief The only constructor that can create storage, this is to make sure CUDA stream is being + * used. + */ + RAFT_MDARRAY_CTOR_CONSTEXPR mdarray(mapping_type const& m, container_policy_type const& cp) + : cp_(cp), map_(m), c_(cp_.create(map_.required_span_size())) + { + } + RAFT_MDARRAY_CTOR_CONSTEXPR mdarray(mapping_type const& m, container_policy_type& cp) + : cp_(cp), map_(m), c_(cp_.create(map_.required_span_size())) + { + } + +#undef RAFT_MDARRAY_CTOR_CONSTEXPR + + /** + * @brief Get a mdspan that can be passed down to CUDA kernels. + */ + auto view() noexcept { return view_type(c_.data(), map_, cp_.make_accessor_policy()); } + /** + * @brief Get a mdspan that can be passed down to CUDA kernels. + */ + auto view() const noexcept + { + return const_view_type(c_.data(), map_, cp_.make_accessor_policy()); + } + + [[nodiscard]] constexpr auto size() const noexcept -> index_type { return this->view().size(); } + + [[nodiscard]] auto data() noexcept -> pointer { return c_.data(); } + [[nodiscard]] constexpr auto data() const noexcept -> const_pointer { return c_.data(); } + + /** + * @brief Indexing operator, use it sparingly since it triggers a device<->host copy. + */ + template + auto operator()(IndexType&&... indices) + -> std::enable_if_t && ...) && + std::is_constructible_v && + std::is_constructible_v, + /* device policy is not default constructible due to requirement for CUDA + stream. */ + /* std::is_default_constructible_v */ + reference> + { + return cp_.access(c_, map_(std::forward(indices)...)); + } + + /** + * @brief Indexing operator, use it sparingly since it triggers a device<->host copy. + */ + template + auto operator()(IndexType&&... indices) const + -> std::enable_if_t && ...) && + std::is_constructible_v && + std::is_constructible::value, + /* device policy is not default constructible due to requirement for CUDA + stream. */ + /* std::is_default_constructible_v */ + const_reference> + { + return cp_.access(c_, map_(std::forward(indices)...)); + } + + // basic_mdarray observers of the domain multidimensional index space (also in basic_mdspan) + [[nodiscard]] MDSPAN_INLINE_FUNCTION static constexpr auto rank() noexcept -> index_type + { + return extents_type::rank(); + } + [[nodiscard]] MDSPAN_INLINE_FUNCTION static constexpr auto rank_dynamic() noexcept -> index_type + { + return extents_type::rank_dynamic(); + } + [[nodiscard]] MDSPAN_INLINE_FUNCTION static constexpr auto static_extent(size_t r) noexcept + -> index_type + { + return extents_type::static_extent(r); + } + [[nodiscard]] MDSPAN_INLINE_FUNCTION constexpr auto extents() const noexcept -> extents_type + { + return map_.extents(); + } + /** + * @brief the extent of rank r + */ + [[nodiscard]] MDSPAN_INLINE_FUNCTION constexpr auto extent(size_t r) const noexcept -> index_type + { + return map_.extents().extent(r); + } + // mapping + [[nodiscard]] MDSPAN_INLINE_FUNCTION constexpr auto mapping() const noexcept -> mapping_type + { + return map_; + } + [[nodiscard]] MDSPAN_INLINE_FUNCTION constexpr auto is_unique() const noexcept -> bool + { + return map_.is_unique(); + } + [[nodiscard]] MDSPAN_INLINE_FUNCTION constexpr auto is_contiguous() const noexcept -> bool + { + return map_.is_contiguous(); + } + [[nodiscard]] MDSPAN_INLINE_FUNCTION constexpr auto is_strided() const noexcept -> bool + { + return map_.is_strided(); + } + [[nodiscard]] MDSPAN_INLINE_FUNCTION constexpr auto stride(size_t r) const -> index_type + { + return map_.stride(r); + } + + [[nodiscard]] MDSPAN_INLINE_FUNCTION static constexpr auto is_always_unique() noexcept -> bool + { + return mapping_type::is_always_unique(); + } + [[nodiscard]] MDSPAN_INLINE_FUNCTION static constexpr auto is_always_contiguous() noexcept -> bool + { + return mapping_type::is_always_contiguous(); + } + [[nodiscard]] MDSPAN_INLINE_FUNCTION static constexpr auto is_always_strided() noexcept -> bool + { + return mapping_type::is_always_strided(); + } + + private: + template + friend class mdarray; + + private: + container_policy_type cp_; + mapping_type map_; + container_type c_; +}; + +template > +using host_mdarray = + mdarray>; + +template > +using device_mdarray = + mdarray>; +} // namespace raft diff --git a/cpp/test/CMakeLists.txt b/cpp/test/CMakeLists.txt index 07f04ad2ab..7749efe624 100644 --- a/cpp/test/CMakeLists.txt +++ b/cpp/test/CMakeLists.txt @@ -36,6 +36,7 @@ add_executable(test_raft test/eigen_solvers.cu test/handle.cpp test/integer_utils.cpp + test/mdarray.cu test/nvtx.cpp test/pow2_utils.cu test/label/label.cu diff --git a/cpp/test/mdarray.cu b/cpp/test/mdarray.cu new file mode 100644 index 0000000000..d548e97960 --- /dev/null +++ b/cpp/test/mdarray.cu @@ -0,0 +1,289 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace { +namespace stdex = std::experimental; +void check_status(int32_t* d_status, rmm::cuda_stream_view stream) +{ + stream.synchronize(); + int32_t h_status{1}; + raft::update_host(&h_status, d_status, 1, stream); + ASSERT_EQ(h_status, 0); +} + +// just simple integration test, main tests are in mdspan ref implementation. +void test_mdspan() +{ + auto stream = rmm::cuda_stream_default; + rmm::device_uvector a{16ul, stream}; + thrust::sequence(rmm::exec_policy(stream), a.begin(), a.end()); + stdex::mdspan> span{ + a.data(), 4, 4}; + thrust::device_vector status(1, 0); + auto p_status = status.data().get(); + thrust::for_each_n( + rmm::exec_policy(stream), thrust::make_counting_iterator(0ul), 4, [=] __device__(size_t i) { + auto v = span(0, i); + if (v != i) { raft::myAtomicAdd(p_status, 1); } + auto k = stdex::submdspan(span, 0, stdex::full_extent); + if (k(i) != i) { raft::myAtomicAdd(p_status, 1); } + }); + check_status(p_status, stream); +} +} // namespace + +TEST(MDSpan, Basic) { test_mdspan(); } + +namespace raft { +void test_uvector_policy() +{ + auto s = rmm::cuda_stream{}; + detail::device_uvector dvec(10, s); + auto a = dvec[2]; + a = 3; + float c = a; + ASSERT_EQ(c, 3); +} + +TEST(MDArray, Policy) { test_uvector_policy(); } + +void test_mdarray_basic() +{ + using matrix_extent = stdex::extents; + auto s = rmm::cuda_stream_default; + { + /** + * device policy + */ + stdex::layout_right::mapping layout{matrix_extent{4, 4}}; + using mdarray_t = device_mdarray; + auto policy = mdarray_t::container_policy_type{s}; + static_assert(std::is_same_v>); + device_mdarray array{layout, policy}; + + array(0, 3) = 1; + ASSERT_EQ(array(0, 3), 1); + // non-const access + auto d_view = array.view(); + static_assert(!decltype(d_view)::accessor_type::is_host_type::value); + + thrust::device_vector status(1, 0); + auto p_status = status.data().get(); + thrust::for_each_n(rmm::exec_policy(s), + thrust::make_counting_iterator(0ul), + 1, + [d_view, p_status] __device__(auto i) { + if (d_view(0, 3) != 1) { myAtomicAdd(p_status, 1); } + d_view(0, 2) = 3; + if (d_view(0, 2) != 3) { myAtomicAdd(p_status, 1); } + }); + check_status(p_status, s); + + // const ref access + auto const& arr = array; + ASSERT_EQ(arr(0, 3), 1); + auto const_d_view = arr.view(); + thrust::for_each_n(rmm::exec_policy(s), + thrust::make_counting_iterator(0ul), + 1, + [const_d_view, p_status] __device__(auto i) { + if (const_d_view(0, 3) != 1) { myAtomicAdd(p_status, 1); } + }); + check_status(p_status, s); + + // utilities + static_assert(array.rank_dynamic() == 2); + static_assert(array.rank() == 2); + static_assert(array.is_unique()); + static_assert(array.is_contiguous()); + static_assert(array.is_strided()); + + static_assert(!std::is_nothrow_default_constructible::value); // cuda stream + static_assert(std::is_nothrow_move_constructible::value); + static_assert(std::is_nothrow_move_assignable::value); + } + { + /** + * host policy + */ + using mdarray_t = host_mdarray; + mdarray_t::container_policy_type policy; + static_assert( + std::is_same_v>); + stdex::layout_right::mapping layout{matrix_extent{4, 4}}; + host_mdarray array{layout, policy}; + + array(0, 3) = 1; + ASSERT_EQ(array(0, 3), 1); + auto h_view = array.view(); + static_assert(decltype(h_view)::accessor_type::is_host_type::value); + thrust::for_each_n(thrust::host, thrust::make_counting_iterator(0ul), 1, [h_view](auto i) { + ASSERT_EQ(h_view(0, 3), 1); + }); + + static_assert(std::is_nothrow_default_constructible::value); + static_assert(std::is_nothrow_move_constructible::value); + static_assert(std::is_nothrow_move_assignable::value); + } + { + /** + * static extent + */ + using static_extent = stdex::extents<16, 16>; + stdex::layout_right::mapping layout{static_extent{}}; + using mdarray_t = device_mdarray; + mdarray_t::container_policy_type policy{s}; + device_mdarray array{layout, policy}; + + static_assert(array.rank_dynamic() == 0); + static_assert(array.rank() == 2); + static_assert(array.is_unique()); + static_assert(array.is_contiguous()); + static_assert(array.is_strided()); + + array(0, 3) = 1; + ASSERT_EQ(array(0, 3), 1); + + auto const& ref = array; + ASSERT_EQ(ref(0, 3), 1); + } +} + +TEST(MDArray, Basic) { test_mdarray_basic(); } + +template +void test_mdarray_copy_move(ThrustPolicy exec, PolicyFn make_policy) +{ + using matrix_extent = stdex::extents; + stdex::layout_right::mapping layout{matrix_extent{4, 4}}; + + using mdarray_t = BasicMDarray; + using policy_t = typename mdarray_t::container_policy_type; + auto policy = make_policy(); + + mdarray_t arr_origin{layout, policy}; + thrust::sequence(exec, arr_origin.data(), arr_origin.data() + arr_origin.size()); + + auto check_eq = [](auto const& l, auto const& r) { + ASSERT_EQ(l.extents(), r.extents()); + for (size_t i = 0; i < l.view().extent(0); ++i) { + for (size_t j = 0; j < l.view().extent(1); ++j) { + ASSERT_EQ(l(i, j), r(i, j)); + } + } + }; + + { + // copy ctor + auto policy = make_policy(); + mdarray_t arr{layout, policy}; + thrust::sequence(exec, arr.data(), arr.data() + arr.size()); + mdarray_t arr_copy_construct{arr}; + check_eq(arr, arr_copy_construct); + + auto const& ref = arr; + mdarray_t arr_copy_construct_1{ref}; + check_eq(ref, arr_copy_construct_1); + } + + { + // copy assign + auto policy = make_policy(); + mdarray_t arr{layout, policy}; + thrust::sequence(exec, arr.data(), arr.data() + arr.size()); + mdarray_t arr_copy_assign{layout, policy}; + arr_copy_assign = arr; + check_eq(arr, arr_copy_assign); + + auto const& ref = arr; + mdarray_t arr_copy_assign_1{layout, policy}; + arr_copy_assign_1 = ref; + check_eq(ref, arr_copy_assign_1); + } + + { + // move ctor + auto policy = make_policy(); + mdarray_t arr{layout, policy}; + thrust::sequence(exec, arr.data(), arr.data() + arr.size()); + mdarray_t arr_move_construct{std::move(arr)}; + ASSERT_EQ(arr.data(), nullptr); + check_eq(arr_origin, arr_move_construct); + } + + { + // move assign + auto policy = make_policy(); + mdarray_t arr{layout, policy}; + thrust::sequence(exec, arr.data(), arr.data() + arr.size()); + mdarray_t arr_move_assign{layout, policy}; + arr_move_assign = std::move(arr); + ASSERT_EQ(arr.data(), nullptr); + check_eq(arr_origin, arr_move_assign); + } +} + +TEST(MDArray, CopyMove) +{ + using matrix_extent = stdex::extents; + using d_matrix_t = device_mdarray; + using policy_t = typename d_matrix_t::container_policy_type; + auto s = rmm::cuda_stream_default; + test_mdarray_copy_move(rmm::exec_policy(s), [s]() { return policy_t{s}; }); + + using h_matrix_t = host_mdarray; + test_mdarray_copy_move(thrust::host, + []() { return detail::host_vector_policy{}; }); + + { + d_matrix_t arr; + auto s = rmm::cuda_stream(); + policy_t policy{s}; + matrix_extent extents{3, 3}; + d_matrix_t::layout_type::mapping layout{extents}; + d_matrix_t non_dft{layout, policy}; + + arr = non_dft; + ASSERT_NE(arr.data(), non_dft.data()); + ASSERT_EQ(arr.extent(0), non_dft.extent(0)); + } + { + h_matrix_t arr; + using h_policy_t = typename h_matrix_t::container_policy_type; + h_policy_t policy{s}; + matrix_extent extents{3, 3}; + h_matrix_t::layout_type::mapping layout{extents}; + h_matrix_t non_dft{layout, policy}; + + arr = non_dft; + ASSERT_NE(arr.data(), non_dft.data()); + ASSERT_EQ(arr.extent(0), non_dft.extent(0)); + } +} +} // namespace raft From 04c5b3918dade6ee3786c1aa144de11baf486dcb Mon Sep 17 00:00:00 2001 From: fis Date: Tue, 25 Jan 2022 12:04:05 +0800 Subject: [PATCH 02/14] Move into detail. --- cpp/include/raft/detail/mdarray.hpp | 228 +++++++++++++++++++++++++++ cpp/include/raft/mdarray.h | 231 ++-------------------------- 2 files changed, 238 insertions(+), 221 deletions(-) create mode 100644 cpp/include/raft/detail/mdarray.hpp diff --git a/cpp/include/raft/detail/mdarray.hpp b/cpp/include/raft/detail/mdarray.hpp new file mode 100644 index 0000000000..8f7166ea16 --- /dev/null +++ b/cpp/include/raft/detail/mdarray.hpp @@ -0,0 +1,228 @@ +/* + * Copyright (c) 2022, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ +#include +#include +#include +#include + +namespace raft::detail { +/** + * @brief A simplified version of thrust::device_reference with support for CUDA stream. + */ +template +class device_reference { + public: + using value_type = typename std::remove_cv_t; + using pointer = thrust::device_ptr; + using const_pointer = thrust::device_ptr; + + private: + std::conditional_t::value, const_pointer, pointer> ptr_; + rmm::cuda_stream_view stream_; + + public: + device_reference(thrust::device_ptr ptr, rmm::cuda_stream_view stream) + : ptr_{ptr}, stream_{stream} + { + } + + operator value_type() const // NOLINT + { + auto* raw = ptr_.get(); + value_type v{}; + update_host(&v, raw, 1, stream_); + return v; + } + auto operator=(T const& other) -> device_reference& + { + auto* raw = ptr_.get(); + update_device(raw, &other, 1, stream_); + return *this; + } +}; + +/** + * @brief A thin wrapper over rmm::device_uvector for implementing the mdarray container policy. + * + */ +template +class device_uvector { + rmm::device_uvector data_; + + public: + using value_type = T; + using size_type = std::size_t; + + using reference = device_reference; + using const_reference = device_reference; + + using pointer = value_type*; + using const_pointer = value_type const*; + + using iterator = pointer; + using const_iterator = const_pointer; + + public: + ~device_uvector() = default; + device_uvector(device_uvector&&) noexcept = default; + device_uvector(device_uvector const& that) : data_{that.data_, that.data_.stream()} {} + + auto operator=(device_uvector const& that) -> device_uvector& + { + data_ = rmm::device_uvector{that.data_, that.data_.stream()}; + return *this; + } + auto operator=(device_uvector&& that) noexcept -> device_uvector& = default; + + /** + * @brief Default ctor is deleted as it doesn't accept stream. + */ + device_uvector() = delete; + /** + * @brief Ctor that accepts a size, stream and an optional mr. + */ + explicit device_uvector( + std::size_t size, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) + : data_{size, stream, mr} + { + } + /** + * @brief Index operator that returns a proxy to the actual data. + */ + template + auto operator[](Index i) noexcept -> reference + { + return device_reference{thrust::device_ptr{data_.data() + i}, data_.stream()}; + } + /** + * @brief Index operator that returns a proxy to the actual data. + */ + template + auto operator[](Index i) const noexcept + { + return device_reference{thrust::device_ptr{data_.data() + i}, data_.stream()}; + } + + [[nodiscard]] auto data() noexcept -> pointer { return data_.data(); } + [[nodiscard]] auto data() const noexcept -> const_pointer { return data_.data(); } +}; + +/** + * @brief A container policy for device mdarray. + */ +template +class device_uvector_policy { + rmm::cuda_stream_view stream_; + + public: + using element_type = ElementType; + using container_type = device_uvector; + // FIXME(jiamingy): allocator type is not supported by rmm::device_uvector + using pointer = typename container_type::pointer; + using const_pointer = typename container_type::const_pointer; + using reference = device_reference; + using const_reference = device_reference; + + using accessor_policy = std::experimental::default_accessor; + using const_accessor_policy = std::experimental::default_accessor; + + public: + auto create(size_t n) -> container_type { return container_type(n, stream_); } + + device_uvector_policy() = delete; + explicit device_uvector_policy(rmm::cuda_stream_view stream) noexcept( + std::is_nothrow_copy_constructible_v) + : stream_{stream} + { + } + + [[nodiscard]] constexpr auto access(container_type& c, size_t n) const noexcept -> reference + { + return c[n]; + } + [[nodiscard]] constexpr auto access(container_type const& c, size_t n) const noexcept + -> const_reference + { + return c[n]; + } + + [[nodiscard]] auto make_accessor_policy() noexcept { return accessor_policy{}; } + [[nodiscard]] auto make_accessor_policy() const noexcept { return const_accessor_policy{}; } +}; + +/** + * @brief A container policy for host mdarray. + */ +template > +class host_vector_policy { + public: + using element_type = ElementType; + using container_type = std::vector; + using allocator_type = typename container_type::allocator_type; + using pointer = typename container_type::pointer; + using const_pointer = typename container_type::const_pointer; + using reference = element_type&; + using const_reference = element_type const&; + using accessor_policy = std::experimental::default_accessor; + using const_accessor_policy = std::experimental::default_accessor; + + public: + auto create(size_t n) -> container_type { return container_type(n); } + + constexpr host_vector_policy() noexcept(std::is_nothrow_default_constructible_v) = + default; + explicit constexpr host_vector_policy(rmm::cuda_stream_view) noexcept( + std::is_nothrow_default_constructible_v) + : host_vector_policy() + { + } + + [[nodiscard]] constexpr auto access(container_type& c, size_t n) const noexcept -> reference + { + return c[n]; + } + [[nodiscard]] constexpr auto access(container_type const& c, size_t n) const noexcept + -> const_reference + { + return c[n]; + } + + [[nodiscard]] auto make_accessor_policy() noexcept { return accessor_policy{}; } + [[nodiscard]] auto make_accessor_policy() const noexcept { return const_accessor_policy{}; } +}; + +/** + * @brief A mixin to distinguish host and device memory. + */ +template +struct accessor_mixin : public AccessorPolicy { + using accessor_type = AccessorPolicy; + using is_host_type = std::conditional_t; + // make sure the explicit ctor can fall through + using AccessorPolicy::AccessorPolicy; + accessor_mixin(AccessorPolicy const& that) : AccessorPolicy{that} {} // NOLINT +}; + +template +using host_accessor = accessor_mixin; + +template +using device_accessor = accessor_mixin; + +namespace stdex = std::experimental; +} // namespace raft::detail diff --git a/cpp/include/raft/mdarray.h b/cpp/include/raft/mdarray.h index ce4f030876..ec4378e6f6 100644 --- a/cpp/include/raft/mdarray.h +++ b/cpp/include/raft/mdarray.h @@ -14,240 +14,29 @@ * limitations under the License. */ #include -#include +#include #include -#include -#include namespace raft { -namespace detail { -/** - * @brief A simplified version of thrust::device_reference with support for CUDA stream. - */ -template -class device_reference { - public: - using value_type = typename std::remove_cv_t; - using pointer = thrust::device_ptr; - using const_pointer = thrust::device_ptr; - - private: - std::conditional_t::value, const_pointer, pointer> ptr_; - rmm::cuda_stream_view stream_; - - public: - device_reference(thrust::device_ptr ptr, rmm::cuda_stream_view stream) - : ptr_{ptr}, stream_{stream} - { - } - - operator value_type() const // NOLINT - { - auto* raw = ptr_.get(); - value_type v{}; - update_host(&v, raw, 1, stream_); - return v; - } - auto operator=(T const& other) -> device_reference& - { - auto* raw = ptr_.get(); - update_device(raw, &other, 1, stream_); - return *this; - } -}; - -/** - * @brief A thin wrapper over rmm::device_uvector for implementing the mdarray container policy. - * - */ -template -class device_uvector { - rmm::device_uvector data_; - - public: - using value_type = T; - using size_type = std::size_t; - - using reference = device_reference; - using const_reference = device_reference; - - using pointer = value_type*; - using const_pointer = value_type const*; - - using iterator = pointer; - using const_iterator = const_pointer; - - public: - ~device_uvector() = default; - device_uvector(device_uvector&&) noexcept = default; - device_uvector(device_uvector const& that) : data_{that.data_, that.data_.stream()} {} - - auto operator=(device_uvector const& that) -> device_uvector& - { - data_ = rmm::device_uvector{that.data_, that.data_.stream()}; - return *this; - } - auto operator=(device_uvector&& that) noexcept -> device_uvector& = default; - - /** - * @brief Default ctor is deleted as it doesn't accept stream. - */ - device_uvector() = delete; - /** - * @brief Ctor that accepts a size, stream and an optional mr. - */ - explicit device_uvector( - std::size_t size, - rmm::cuda_stream_view stream, - rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource()) - : data_{size, stream, mr} - { - } - /** - * @brief Index operator that returns a proxy to the actual data. - */ - template - auto operator[](Index i) noexcept -> reference - { - return device_reference{thrust::device_ptr{data_.data() + i}, data_.stream()}; - } - /** - * @brief Index operator that returns a proxy to the actual data. - */ - template - auto operator[](Index i) const noexcept - { - return device_reference{thrust::device_ptr{data_.data() + i}, data_.stream()}; - } - - [[nodiscard]] auto data() noexcept -> pointer { return data_.data(); } - [[nodiscard]] auto data() const noexcept -> const_pointer { return data_.data(); } -}; - -/** - * @brief A container policy for device mdarray. - */ -template -class device_uvector_policy { - rmm::cuda_stream_view stream_; - - public: - using element_type = ElementType; - using container_type = device_uvector; - // FIXME(jiamingy): allocator type is not supported by rmm::device_uvector - using pointer = typename container_type::pointer; - using const_pointer = typename container_type::const_pointer; - using reference = device_reference; - using const_reference = device_reference; - - using accessor_policy = std::experimental::default_accessor; - using const_accessor_policy = std::experimental::default_accessor; - - public: - auto create(size_t n) -> container_type { return container_type(n, stream_); } - - device_uvector_policy() = delete; - explicit device_uvector_policy(rmm::cuda_stream_view stream) noexcept( - std::is_nothrow_copy_constructible_v) - : stream_{stream} - { - } - - [[nodiscard]] constexpr auto access(container_type& c, size_t n) const noexcept -> reference - { - return c[n]; - } - [[nodiscard]] constexpr auto access(container_type const& c, size_t n) const noexcept - -> const_reference - { - return c[n]; - } - - [[nodiscard]] auto make_accessor_policy() noexcept { return accessor_policy{}; } - [[nodiscard]] auto make_accessor_policy() const noexcept { return const_accessor_policy{}; } -}; - -/** - * @brief A container policy for host mdarray. - */ -template > -class host_vector_policy { - public: - using element_type = ElementType; - using container_type = std::vector; - using allocator_type = typename container_type::allocator_type; - using pointer = typename container_type::pointer; - using const_pointer = typename container_type::const_pointer; - using reference = element_type&; - using const_reference = element_type const&; - using accessor_policy = std::experimental::default_accessor; - using const_accessor_policy = std::experimental::default_accessor; - - public: - auto create(size_t n) -> container_type { return container_type(n); } - - constexpr host_vector_policy() noexcept(std::is_nothrow_default_constructible_v) = - default; - explicit constexpr host_vector_policy(rmm::cuda_stream_view) noexcept( - std::is_nothrow_default_constructible_v) - : host_vector_policy() - { - } - - [[nodiscard]] constexpr auto access(container_type& c, size_t n) const noexcept -> reference - { - return c[n]; - } - [[nodiscard]] constexpr auto access(container_type const& c, size_t n) const noexcept - -> const_reference - { - return c[n]; - } - - [[nodiscard]] auto make_accessor_policy() noexcept { return accessor_policy{}; } - [[nodiscard]] auto make_accessor_policy() const noexcept { return const_accessor_policy{}; } -}; - -/** - * @brief A mixin to distinguish host and device memory. - */ -template -struct accessor_mixin : public AccessorPolicy { - using accessor_type = AccessorPolicy; - using is_host_type = std::conditional_t; - // make sure the explicit ctor can fall through - using AccessorPolicy::AccessorPolicy; - accessor_mixin(AccessorPolicy const& that) : AccessorPolicy{that} {} // NOLINT -}; - -template -using host_accessor = accessor_mixin; - -template -using device_accessor = accessor_mixin; -} // namespace detail - -namespace stdex = std::experimental; - /** * @brief stdex::mdspan with device tag to avoid accessing incorrect memory location. */ template > -using device_mdspan = - stdex::mdspan>; + class LayoutPolicy = detail::stdex::layout_right, + class AccessorPolicy = detail::stdex::default_accessor> +using device_mdspan = detail::stdex:: + mdspan>; /** * @brief stdex::mdspan with host tag to avoid accessing incorrect memory location. */ template > + class LayoutPolicy = detail::stdex::layout_right, + class AccessorPolicy = detail::stdex::default_accessor> using host_mdspan = - stdex::mdspan>; + detail::stdex::mdspan>; /** * @brief Modified from the c++ mdarray proposal @@ -488,14 +277,14 @@ class mdarray { template > using host_mdarray = mdarray>; template > using device_mdarray = mdarray>; From 7b69dd0be2e39bf0695af7ca9d11a587fd080bc1 Mon Sep 17 00:00:00 2001 From: fis Date: Tue, 25 Jan 2022 12:07:05 +0800 Subject: [PATCH 03/14] Rename header to hpp. --- cpp/include/raft/{mdarray.h => mdarray.hpp} | 0 cpp/test/mdarray.cu | 2 +- 2 files changed, 1 insertion(+), 1 deletion(-) rename cpp/include/raft/{mdarray.h => mdarray.hpp} (100%) diff --git a/cpp/include/raft/mdarray.h b/cpp/include/raft/mdarray.hpp similarity index 100% rename from cpp/include/raft/mdarray.h rename to cpp/include/raft/mdarray.hpp diff --git a/cpp/test/mdarray.cu b/cpp/test/mdarray.cu index d548e97960..b3efc8237a 100644 --- a/cpp/test/mdarray.cu +++ b/cpp/test/mdarray.cu @@ -17,7 +17,7 @@ #include #include #include -#include +#include #include #include #include From 1b01e688c00418b2c3b2be82274118c8c9431201 Mon Sep 17 00:00:00 2001 From: fis Date: Tue, 25 Jan 2022 12:35:52 +0800 Subject: [PATCH 04/14] Factory methods. --- cpp/include/raft/detail/mdarray.hpp | 3 + cpp/include/raft/mdarray.hpp | 121 ++++++++++++++++++++++++++++ cpp/test/mdarray.cu | 56 +++++++++++++ 3 files changed, 180 insertions(+) diff --git a/cpp/include/raft/detail/mdarray.hpp b/cpp/include/raft/detail/mdarray.hpp index 8f7166ea16..aa99de4a9a 100644 --- a/cpp/include/raft/detail/mdarray.hpp +++ b/cpp/include/raft/detail/mdarray.hpp @@ -225,4 +225,7 @@ template using device_accessor = accessor_mixin; namespace stdex = std::experimental; + +using vector_extent_t = stdex::extents; +using matrix_extent_t = stdex::extents; } // namespace raft::detail diff --git a/cpp/include/raft/mdarray.hpp b/cpp/include/raft/mdarray.hpp index ec4378e6f6..a6672ab449 100644 --- a/cpp/include/raft/mdarray.hpp +++ b/cpp/include/raft/mdarray.hpp @@ -288,4 +288,125 @@ template > using device_mdarray = mdarray>; + +template +using host_vector_t = host_mdarray; + +template +using device_vector_t = device_mdarray; + +/** + * @brief Shorthand for c-contiguous host matrix. + */ +template +using host_matrix_t = host_mdarray; +/** + * @brief Shorthand for c-contiguous device matrix. + */ +template +using device_matrix_t = device_mdarray; + +template +using host_vector_view_t = host_mdspan; + +template +using device_vector_view_t = device_mdspan; + +/** + * @brief Shorthand for c-contiguous host matrix view. + */ +template +using host_matrix_view_t = host_mdspan; +/** + * @brief Shorthand for c-contiguous device matrix view. + */ +template +using device_matrix_view_t = device_mdspan; + +/** + * @brief Create a 2-dim c-contiguous mdspan instance for host pointer. + */ +template +auto make_host_matrix_view(ElementType* ptr, size_t n_rows, size_t n_cols) +{ + detail::matrix_extent_t extents{n_rows, n_cols}; + return host_matrix_view_t{ptr, extents}; +} +/** + * @brief Create a 2-dim c-contiguous mdspan instance for device pointer. + */ +template +auto make_device_matrix_view(ElementType* ptr, size_t n_rows, size_t n_cols) +{ + detail::matrix_extent_t extents{n_rows, n_cols}; + return device_matrix_view_t{ptr, extents}; +} + +/** + * @brief Create a 1-dim mdspan instance for host pointer. + */ +template +auto make_host_vector_view(ElementType* ptr, size_t n) +{ + detail::vector_extent_t extents{n}; + return host_matrix_view_t{ptr, extents}; +} + +/** + * @brief Create a 1-dim mdspan instance for device pointer. + */ +template +auto make_device_vector_view(ElementType* ptr, size_t n) +{ + detail::vector_extent_t extents{n}; + return device_matrix_view_t{ptr, extents}; +} + +/** + * @brief Create a 2-dim c-contiguous host mdarray. + */ +template +auto make_host_matrix(size_t n_rows, size_t n_cols) +{ + detail::matrix_extent_t extents{n_rows, n_cols}; + using policy_t = typename host_matrix_t::container_policy_type; + policy_t policy; + return host_matrix_t{extents, policy}; +} + +/** + * @brief Create a 2-dim c-contiguous device mdarray. + */ +template +auto make_device_matrix(size_t n_rows, size_t n_cols, rmm::cuda_stream_view stream) +{ + detail::matrix_extent_t extents{n_rows, n_cols}; + using policy_t = typename device_matrix_t::container_policy_type; + policy_t policy{stream}; + return device_matrix_t{extents, policy}; +} + +/** + * @brief Create a 1-dim host mdarray. + */ +template +auto make_host_vector(size_t n) +{ + detail::vector_extent_t extents{n}; + using policy_t = typename host_vector_t::container_policy_type; + policy_t policy; + return host_vector_t{extents, policy}; +} + +/** + * @brief Create a 1-dim device mdarray. + */ +template +auto make_device_vector(size_t n, rmm::cuda_stream_view stream) +{ + detail::vector_extent_t extents{n}; + using policy_t = typename device_vector_t::container_policy_type; + policy_t policy{stream}; + return device_vector_t{extents, policy}; +} } // namespace raft diff --git a/cpp/test/mdarray.cu b/cpp/test/mdarray.cu index b3efc8237a..9ad3f9a65e 100644 --- a/cpp/test/mdarray.cu +++ b/cpp/test/mdarray.cu @@ -286,4 +286,60 @@ TEST(MDArray, CopyMove) ASSERT_EQ(arr.extent(0), non_dft.extent(0)); } } + +TEST(MDArray, Factory) +{ + size_t n{100}; + rmm::device_uvector d_vec(n, rmm::cuda_stream_default); + { + auto d_matrix = make_device_matrix_view(d_vec.data(), d_vec.size() / 2, 2); + ASSERT_EQ(d_matrix.extent(0), n / 2); + ASSERT_EQ(d_matrix.extent(1), 2); + ASSERT_EQ(d_matrix.data(), d_vec.data()); + } + { + auto const& vec_ref = d_vec; + auto d_matrix = make_device_matrix_view(vec_ref.data(), d_vec.size() / 2, 2); + ASSERT_EQ(d_matrix.extent(0), n / 2); + ASSERT_EQ(d_matrix.extent(1), 2); + ASSERT_EQ(d_matrix.data(), d_vec.data()); + } + + std::vector h_vec(n); + { + auto h_matrix = make_host_matrix_view(h_vec.data(), h_vec.size() / 2, 2); + ASSERT_EQ(h_matrix.extent(0), n / 2); + ASSERT_EQ(h_matrix.extent(1), 2); + ASSERT_EQ(h_matrix.data(), h_vec.data()); + h_matrix(0, 0) = 13; + ASSERT_EQ(h_matrix(0, 0), 13); + } + { + auto const& vec_ref = h_vec; + auto h_matrix = make_host_matrix_view(vec_ref.data(), d_vec.size() / 2, 2); + ASSERT_EQ(h_matrix.extent(0), n / 2); + ASSERT_EQ(h_matrix.extent(1), 2); + ASSERT_EQ(h_matrix.data(), h_vec.data()); + // const, cannot assign + // h_matrix(0, 0) = 13; + ASSERT_EQ(h_matrix(0, 0), 13); + } + + { + // host mdarray + auto h_matrix = make_host_matrix(n, n); + ASSERT_EQ(h_matrix.extent(0), n); + ASSERT_EQ(h_matrix.extent(1), n); + + auto h_vec = make_host_vector(n); + } + { + // device mdarray + auto d_matrix = make_device_matrix(n, n, rmm::cuda_stream_default); + ASSERT_EQ(d_matrix.extent(0), n); + ASSERT_EQ(d_matrix.extent(1), n); + + auto d_vec = make_device_vector(n, rmm::cuda_stream_default); + } +} } // namespace raft From 511fafa773ca86b4f1ba858497a5af62daddb23b Mon Sep 17 00:00:00 2001 From: fis Date: Wed, 26 Jan 2022 18:06:09 +0800 Subject: [PATCH 05/14] pragma --- cpp/include/raft/detail/mdarray.hpp | 1 + cpp/include/raft/mdarray.hpp | 1 + 2 files changed, 2 insertions(+) diff --git a/cpp/include/raft/detail/mdarray.hpp b/cpp/include/raft/detail/mdarray.hpp index aa99de4a9a..61d15e105a 100644 --- a/cpp/include/raft/detail/mdarray.hpp +++ b/cpp/include/raft/detail/mdarray.hpp @@ -13,6 +13,7 @@ * See the License for the specific language governing permissions and * limitations under the License. */ +#pragma once #include #include #include diff --git a/cpp/include/raft/mdarray.hpp b/cpp/include/raft/mdarray.hpp index a6672ab449..b41adceaba 100644 --- a/cpp/include/raft/mdarray.hpp +++ b/cpp/include/raft/mdarray.hpp @@ -13,6 +13,7 @@ * See the License for the specific language governing permissions and * limitations under the License. */ +#pragma once #include #include #include From debcaba44444209dca92da1481a9953feb779226 Mon Sep 17 00:00:00 2001 From: fis Date: Wed, 26 Jan 2022 18:11:37 +0800 Subject: [PATCH 06/14] Remove postfix in naming. --- cpp/include/raft/detail/mdarray.hpp | 4 +- cpp/include/raft/mdarray.hpp | 68 +++++++++++++++++------------ 2 files changed, 42 insertions(+), 30 deletions(-) diff --git a/cpp/include/raft/detail/mdarray.hpp b/cpp/include/raft/detail/mdarray.hpp index 61d15e105a..0d6009584a 100644 --- a/cpp/include/raft/detail/mdarray.hpp +++ b/cpp/include/raft/detail/mdarray.hpp @@ -227,6 +227,6 @@ using device_accessor = accessor_mixin; namespace stdex = std::experimental; -using vector_extent_t = stdex::extents; -using matrix_extent_t = stdex::extents; +using vector_extent = stdex::extents; +using matrix_extent = stdex::extents; } // namespace raft::detail diff --git a/cpp/include/raft/mdarray.hpp b/cpp/include/raft/mdarray.hpp index b41adceaba..672c715a44 100644 --- a/cpp/include/raft/mdarray.hpp +++ b/cpp/include/raft/mdarray.hpp @@ -276,6 +276,9 @@ class mdarray { container_type c_; }; +/** + * @brief mdarray with host container policy + */ template >; +/** + * @brief mdarray with device container policy + */ template >; +/** + * @brief Shorthand for 1-dim host mdarray. + */ template -using host_vector_t = host_mdarray; +using host_vector = host_mdarray; +/** + * @brief Shorthand for 1-dim device mdarray. + */ template -using device_vector_t = device_mdarray; +using device_vector = device_mdarray; /** * @brief Shorthand for c-contiguous host matrix. */ template -using host_matrix_t = host_mdarray; +using host_matrix = host_mdarray; /** * @brief Shorthand for c-contiguous device matrix. */ template -using device_matrix_t = device_mdarray; +using device_matrix = device_mdarray; template -using host_vector_view_t = host_mdspan; +using host_vector_view = host_mdspan; template -using device_vector_view_t = device_mdspan; +using device_vector_view = device_mdspan; /** * @brief Shorthand for c-contiguous host matrix view. */ template -using host_matrix_view_t = host_mdspan; +using host_matrix_view = host_mdspan; /** * @brief Shorthand for c-contiguous device matrix view. */ template -using device_matrix_view_t = device_mdspan; +using device_matrix_view = device_mdspan; /** * @brief Create a 2-dim c-contiguous mdspan instance for host pointer. @@ -330,8 +342,8 @@ using device_matrix_view_t = device_mdspan template auto make_host_matrix_view(ElementType* ptr, size_t n_rows, size_t n_cols) { - detail::matrix_extent_t extents{n_rows, n_cols}; - return host_matrix_view_t{ptr, extents}; + detail::matrix_extent extents{n_rows, n_cols}; + return host_matrix_view{ptr, extents}; } /** * @brief Create a 2-dim c-contiguous mdspan instance for device pointer. @@ -339,8 +351,8 @@ auto make_host_matrix_view(ElementType* ptr, size_t n_rows, size_t n_cols) template auto make_device_matrix_view(ElementType* ptr, size_t n_rows, size_t n_cols) { - detail::matrix_extent_t extents{n_rows, n_cols}; - return device_matrix_view_t{ptr, extents}; + detail::matrix_extent extents{n_rows, n_cols}; + return device_matrix_view{ptr, extents}; } /** @@ -349,8 +361,8 @@ auto make_device_matrix_view(ElementType* ptr, size_t n_rows, size_t n_cols) template auto make_host_vector_view(ElementType* ptr, size_t n) { - detail::vector_extent_t extents{n}; - return host_matrix_view_t{ptr, extents}; + detail::vector_extent extents{n}; + return host_matrix_view{ptr, extents}; } /** @@ -359,8 +371,8 @@ auto make_host_vector_view(ElementType* ptr, size_t n) template auto make_device_vector_view(ElementType* ptr, size_t n) { - detail::vector_extent_t extents{n}; - return device_matrix_view_t{ptr, extents}; + detail::vector_extent extents{n}; + return device_matrix_view{ptr, extents}; } /** @@ -369,10 +381,10 @@ auto make_device_vector_view(ElementType* ptr, size_t n) template auto make_host_matrix(size_t n_rows, size_t n_cols) { - detail::matrix_extent_t extents{n_rows, n_cols}; - using policy_t = typename host_matrix_t::container_policy_type; + detail::matrix_extent extents{n_rows, n_cols}; + using policy_t = typename host_matrix::container_policy_type; policy_t policy; - return host_matrix_t{extents, policy}; + return host_matrix{extents, policy}; } /** @@ -381,10 +393,10 @@ auto make_host_matrix(size_t n_rows, size_t n_cols) template auto make_device_matrix(size_t n_rows, size_t n_cols, rmm::cuda_stream_view stream) { - detail::matrix_extent_t extents{n_rows, n_cols}; - using policy_t = typename device_matrix_t::container_policy_type; + detail::matrix_extent extents{n_rows, n_cols}; + using policy_t = typename device_matrix::container_policy_type; policy_t policy{stream}; - return device_matrix_t{extents, policy}; + return device_matrix{extents, policy}; } /** @@ -393,10 +405,10 @@ auto make_device_matrix(size_t n_rows, size_t n_cols, rmm::cuda_stream_view stre template auto make_host_vector(size_t n) { - detail::vector_extent_t extents{n}; - using policy_t = typename host_vector_t::container_policy_type; + detail::vector_extent extents{n}; + using policy_t = typename host_vector::container_policy_type; policy_t policy; - return host_vector_t{extents, policy}; + return host_vector{extents, policy}; } /** @@ -405,9 +417,9 @@ auto make_host_vector(size_t n) template auto make_device_vector(size_t n, rmm::cuda_stream_view stream) { - detail::vector_extent_t extents{n}; - using policy_t = typename device_vector_t::container_policy_type; + detail::vector_extent extents{n}; + using policy_t = typename device_vector::container_policy_type; policy_t policy{stream}; - return device_vector_t{extents, policy}; + return device_vector{extents, policy}; } } // namespace raft From 297f9d4f943b52c98e124771430be660c41f44c6 Mon Sep 17 00:00:00 2001 From: fis Date: Wed, 26 Jan 2022 18:14:26 +0800 Subject: [PATCH 07/14] checks. --- cpp/test/mdarray.cu | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/cpp/test/mdarray.cu b/cpp/test/mdarray.cu index 9ad3f9a65e..b55f2aa847 100644 --- a/cpp/test/mdarray.cu +++ b/cpp/test/mdarray.cu @@ -330,16 +330,22 @@ TEST(MDArray, Factory) auto h_matrix = make_host_matrix(n, n); ASSERT_EQ(h_matrix.extent(0), n); ASSERT_EQ(h_matrix.extent(1), n); + static_assert(h_matrix.rank() == 2); auto h_vec = make_host_vector(n); + static_assert(h_vec.rank() == 1); + ASSERT_EQ(h_vec.extent(0), n); } { // device mdarray auto d_matrix = make_device_matrix(n, n, rmm::cuda_stream_default); ASSERT_EQ(d_matrix.extent(0), n); ASSERT_EQ(d_matrix.extent(1), n); + static_assert(d_matrix.rank() == 2); auto d_vec = make_device_vector(n, rmm::cuda_stream_default); + static_assert(d_vec.rank() == 1); + ASSERT_EQ(d_vec.extent(0), n); } } } // namespace raft From 07be70415dc18929f6906295b05e6bb3009e7f0f Mon Sep 17 00:00:00 2001 From: fis Date: Fri, 28 Jan 2022 17:32:33 +0800 Subject: [PATCH 08/14] Optional matrix layout. --- cpp/include/raft/mdarray.hpp | 43 +++++++++++++++++++++--------------- cpp/test/mdarray.cu | 34 ++++++++++++++++++++++++++++ 2 files changed, 59 insertions(+), 18 deletions(-) diff --git a/cpp/include/raft/mdarray.hpp b/cpp/include/raft/mdarray.hpp index 672c715a44..00a89bdaa5 100644 --- a/cpp/include/raft/mdarray.hpp +++ b/cpp/include/raft/mdarray.hpp @@ -117,9 +117,9 @@ class mdarray { constexpr mdarray(mdarray&&) noexcept(std::is_nothrow_move_constructible::value) = default; - constexpr auto operator =(mdarray const&) noexcept( + constexpr auto operator=(mdarray const&) noexcept( std::is_nothrow_copy_assignable::value) -> mdarray& = default; - constexpr auto operator =(mdarray&&) noexcept( + constexpr auto operator=(mdarray&&) noexcept( std::is_nothrow_move_assignable::value) -> mdarray& = default; ~mdarray() noexcept(std::is_nothrow_destructible::value) = default; @@ -311,48 +311,55 @@ using device_vector = device_mdarray; /** * @brief Shorthand for c-contiguous host matrix. */ -template -using host_matrix = host_mdarray; +template +using host_matrix = host_mdarray; + /** * @brief Shorthand for c-contiguous device matrix. */ -template -using device_matrix = device_mdarray; +template +using device_matrix = device_mdarray; +/** + * @brief Shorthand for 1-dim host mdspan. + */ template using host_vector_view = host_mdspan; +/** + * @brief Shorthand for 1-dim device mdspan. + */ template using device_vector_view = device_mdspan; /** * @brief Shorthand for c-contiguous host matrix view. */ -template -using host_matrix_view = host_mdspan; +template +using host_matrix_view = host_mdspan; /** * @brief Shorthand for c-contiguous device matrix view. */ -template -using device_matrix_view = device_mdspan; +template +using device_matrix_view = device_mdspan; /** * @brief Create a 2-dim c-contiguous mdspan instance for host pointer. */ -template +template auto make_host_matrix_view(ElementType* ptr, size_t n_rows, size_t n_cols) { detail::matrix_extent extents{n_rows, n_cols}; - return host_matrix_view{ptr, extents}; + return host_matrix_view{ptr, extents}; } /** * @brief Create a 2-dim c-contiguous mdspan instance for device pointer. */ -template +template auto make_device_matrix_view(ElementType* ptr, size_t n_rows, size_t n_cols) { detail::matrix_extent extents{n_rows, n_cols}; - return device_matrix_view{ptr, extents}; + return device_matrix_view{ptr, extents}; } /** @@ -378,25 +385,25 @@ auto make_device_vector_view(ElementType* ptr, size_t n) /** * @brief Create a 2-dim c-contiguous host mdarray. */ -template +template auto make_host_matrix(size_t n_rows, size_t n_cols) { detail::matrix_extent extents{n_rows, n_cols}; using policy_t = typename host_matrix::container_policy_type; policy_t policy; - return host_matrix{extents, policy}; + return host_matrix{extents, policy}; } /** * @brief Create a 2-dim c-contiguous device mdarray. */ -template +template auto make_device_matrix(size_t n_rows, size_t n_cols, rmm::cuda_stream_view stream) { detail::matrix_extent extents{n_rows, n_cols}; using policy_t = typename device_matrix::container_policy_type; policy_t policy{stream}; - return device_matrix{extents, policy}; + return device_matrix{extents, policy}; } /** diff --git a/cpp/test/mdarray.cu b/cpp/test/mdarray.cu index b55f2aa847..2dfd87bc9c 100644 --- a/cpp/test/mdarray.cu +++ b/cpp/test/mdarray.cu @@ -348,4 +348,38 @@ TEST(MDArray, Factory) ASSERT_EQ(d_vec.extent(0), n); } } + +namespace { +template +void check_matrix_layout(device_matrix_view in) +{ + static_assert(in.rank() == 2); + static_assert(in.is_contiguous()); + + bool constexpr kIsCContiguous = std::is_same_v; + bool constexpr kIsFContiguous = std::is_same_v; + // only 1 of them is true + static_assert(kIsCContiguous || kIsFContiguous); + static_assert(!(kIsCContiguous && kIsFContiguous)); +} +} // anonymous namespace + +TEST(MDArray, FuncArg) +{ + { + auto d_matrix = make_device_matrix(10, 10, rmm::cuda_stream_default); + check_matrix_layout(d_matrix.view()); + } + { + auto d_matrix = make_device_matrix(10, 10, rmm::cuda_stream_default); + check_matrix_layout(d_matrix.view()); + + // FIXME(jiamingy): The slice has a default accessor instead of accessor_mixin, due to + // the hardcoded policy in submdspan implementation. We need to have a rewritten + // version of submdspan for implementing padding. + // auto slice = + // stdex::submdspan(d_matrix.view(), std::make_tuple(2ul, 4ul), std::make_tuple(2ul, 5ul)); + // check_matrix_layout(slice); + } +} } // namespace raft From e4fdf3286c5f292fd7b42f4c2dd87198720f0506 Mon Sep 17 00:00:00 2001 From: fis Date: Fri, 28 Jan 2022 17:37:51 +0800 Subject: [PATCH 09/14] clang format. --- cpp/include/raft/mdarray.hpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/include/raft/mdarray.hpp b/cpp/include/raft/mdarray.hpp index 00a89bdaa5..d89c304ed6 100644 --- a/cpp/include/raft/mdarray.hpp +++ b/cpp/include/raft/mdarray.hpp @@ -117,9 +117,9 @@ class mdarray { constexpr mdarray(mdarray&&) noexcept(std::is_nothrow_move_constructible::value) = default; - constexpr auto operator=(mdarray const&) noexcept( + constexpr auto operator =(mdarray const&) noexcept( std::is_nothrow_copy_assignable::value) -> mdarray& = default; - constexpr auto operator=(mdarray&&) noexcept( + constexpr auto operator =(mdarray&&) noexcept( std::is_nothrow_move_assignable::value) -> mdarray& = default; ~mdarray() noexcept(std::is_nothrow_destructible::value) = default; From c28a8b39a288f35e2f27aede8c110a006ef12f86 Mon Sep 17 00:00:00 2001 From: fis Date: Sat, 29 Jan 2022 18:22:15 +0800 Subject: [PATCH 10/14] Add alias for c/f contiguous. --- cpp/include/raft/mdarray.hpp | 9 +++++++++ cpp/test/mdarray.cu | 4 ++-- 2 files changed, 11 insertions(+), 2 deletions(-) diff --git a/cpp/include/raft/mdarray.hpp b/cpp/include/raft/mdarray.hpp index d89c304ed6..cbbb3b2489 100644 --- a/cpp/include/raft/mdarray.hpp +++ b/cpp/include/raft/mdarray.hpp @@ -19,6 +19,15 @@ #include namespace raft { +/** + * @\brief C-Contiguous layout for mdarray and mdspan. Implies row-major and contiguous memory. + */ +using layout_c_contiguous = detail::stdex::layout_right; +/** + * @\brief F-Contiguous layout for mdarray and mdspan. Implies column-major and contiguous memory. + */ +using layout_f_contiguous = detail::stdex::layout_left; + /** * @brief stdex::mdspan with device tag to avoid accessing incorrect memory location. */ diff --git a/cpp/test/mdarray.cu b/cpp/test/mdarray.cu index 2dfd87bc9c..9a1b996ef1 100644 --- a/cpp/test/mdarray.cu +++ b/cpp/test/mdarray.cu @@ -356,8 +356,8 @@ void check_matrix_layout(device_matrix_view in) static_assert(in.rank() == 2); static_assert(in.is_contiguous()); - bool constexpr kIsCContiguous = std::is_same_v; - bool constexpr kIsFContiguous = std::is_same_v; + bool constexpr kIsCContiguous = std::is_same_v; + bool constexpr kIsFContiguous = std::is_same_v; // only 1 of them is true static_assert(kIsCContiguous || kIsFContiguous); static_assert(!(kIsCContiguous && kIsFContiguous)); From db9ec8a54b8246eadd638772e73730aa682a2434 Mon Sep 17 00:00:00 2001 From: fis Date: Sat, 29 Jan 2022 18:59:28 +0800 Subject: [PATCH 11/14] scalar & rename. --- cpp/include/raft/detail/mdarray.hpp | 1 + cpp/include/raft/mdarray.hpp | 134 ++++++++++++++++++++++------ cpp/test/mdarray.cu | 59 +++++++++--- 3 files changed, 157 insertions(+), 37 deletions(-) diff --git a/cpp/include/raft/detail/mdarray.hpp b/cpp/include/raft/detail/mdarray.hpp index 0d6009584a..2cc87c409a 100644 --- a/cpp/include/raft/detail/mdarray.hpp +++ b/cpp/include/raft/detail/mdarray.hpp @@ -229,4 +229,5 @@ namespace stdex = std::experimental; using vector_extent = stdex::extents; using matrix_extent = stdex::extents; +using scalar_extent = stdex::extents<1>; } // namespace raft::detail diff --git a/cpp/include/raft/mdarray.hpp b/cpp/include/raft/mdarray.hpp index cbbb3b2489..81808772b1 100644 --- a/cpp/include/raft/mdarray.hpp +++ b/cpp/include/raft/mdarray.hpp @@ -23,6 +23,7 @@ namespace raft { * @\brief C-Contiguous layout for mdarray and mdspan. Implies row-major and contiguous memory. */ using layout_c_contiguous = detail::stdex::layout_right; + /** * @\brief F-Contiguous layout for mdarray and mdspan. Implies column-major and contiguous memory. */ @@ -31,20 +32,20 @@ using layout_f_contiguous = detail::stdex::layout_left; /** * @brief stdex::mdspan with device tag to avoid accessing incorrect memory location. */ -template > +template > using device_mdspan = detail::stdex:: mdspan>; /** * @brief stdex::mdspan with host tag to avoid accessing incorrect memory location. */ -template > +template > using host_mdspan = detail::stdex::mdspan>; @@ -77,7 +78,7 @@ using host_mdspan = * - For the above reasons, copying from other mdarray with different policy type is also * removed. */ -template +template class mdarray { static_assert(!std::is_const::value, "Element type for container must not be const."); @@ -288,23 +289,39 @@ class mdarray { /** * @brief mdarray with host container policy */ -template > +template > using host_mdarray = mdarray>; /** * @brief mdarray with device container policy */ -template > +template > using device_mdarray = mdarray>; +/** + * @brief Shorthand for 0-dim host mdarray (scalar). + * + * Underlying storage is std::vector. + */ +template +using host_scalar = host_mdarray; + +/** + * @brief Shorthand for 0-dim host mdarray (scalar). + * + * Similar to rmm::device_scalar, underying storage is rmm::device_uvector. + */ +template +using device_scalar = device_mdarray; + /** * @brief Shorthand for 1-dim host mdarray. */ @@ -320,15 +337,27 @@ using device_vector = device_mdarray; /** * @brief Shorthand for c-contiguous host matrix. */ -template +template using host_matrix = host_mdarray; /** * @brief Shorthand for c-contiguous device matrix. */ -template +template using device_matrix = device_mdarray; +/** + * @brief Shorthand for 0-dim host mdspan (scalar). + */ +template +using host_scalar_view = host_mdspan; + +/** + * @brief Shorthand for 0-dim host mdspan (scalar). + */ +template +using device_scalar_view = device_mdspan; + /** * @brief Shorthand for 1-dim host mdspan. */ @@ -344,18 +373,38 @@ using device_vector_view = device_mdspan; /** * @brief Shorthand for c-contiguous host matrix view. */ -template +template using host_matrix_view = host_mdspan; /** * @brief Shorthand for c-contiguous device matrix view. */ -template +template using device_matrix_view = device_mdspan; +/** + * @brief Create a 0-dim (scalar) mdspan instance for host value. + */ +template +auto make_host_scalar_view(ElementType* ptr) +{ + detail::scalar_extent extents; + return host_scalar_view{ptr, extents}; +} + +/** + * @brief Create a 0-dim (scalar) mdspan instance for device value. + */ +template +auto make_device_scalar_view(ElementType* ptr) +{ + detail::scalar_extent extents; + return device_scalar_view{ptr, extents}; +} + /** * @brief Create a 2-dim c-contiguous mdspan instance for host pointer. */ -template +template auto make_host_matrix_view(ElementType* ptr, size_t n_rows, size_t n_cols) { detail::matrix_extent extents{n_rows, n_cols}; @@ -364,7 +413,7 @@ auto make_host_matrix_view(ElementType* ptr, size_t n_rows, size_t n_cols) /** * @brief Create a 2-dim c-contiguous mdspan instance for device pointer. */ -template +template auto make_device_matrix_view(ElementType* ptr, size_t n_rows, size_t n_cols) { detail::matrix_extent extents{n_rows, n_cols}; @@ -394,7 +443,7 @@ auto make_device_vector_view(ElementType* ptr, size_t n) /** * @brief Create a 2-dim c-contiguous host mdarray. */ -template +template auto make_host_matrix(size_t n_rows, size_t n_cols) { detail::matrix_extent extents{n_rows, n_cols}; @@ -406,7 +455,7 @@ auto make_host_matrix(size_t n_rows, size_t n_cols) /** * @brief Create a 2-dim c-contiguous device mdarray. */ -template +template auto make_device_matrix(size_t n_rows, size_t n_cols, rmm::cuda_stream_view stream) { detail::matrix_extent extents{n_rows, n_cols}; @@ -415,6 +464,41 @@ auto make_device_matrix(size_t n_rows, size_t n_cols, rmm::cuda_stream_view stre return device_matrix{extents, policy}; } +/** + * @brief Create a host scalar from v. + * + * Underlying storage is std::vector. + */ +template +auto make_host_scalar(ElementType const& v) +{ + // FIXME(jiamingy): We can optimize this by using std::array as container policy, which + // requires some more compile time dispatching. This is enabled in the ref impl but + // hasn't been ported here yet. + detail::scalar_extent extents; + using policy_t = typename host_scalar::container_policy_type; + policy_t policy; + auto scalar = host_scalar{extents, policy}; + scalar(0) = v; + return scalar; +} + +/** + * @brief Create a device scalar from v. + * + * Similar to rmm::device_scalar, underying storage is rmm::device_uvector. + */ +template +auto make_device_scalar(ElementType const& v, rmm::cuda_stream_view stream) +{ + detail::scalar_extent extents; + using policy_t = typename device_scalar::container_policy_type; + policy_t policy{stream}; + auto scalar = device_scalar{extents, policy}; + scalar(0) = v; + return scalar; +} + /** * @brief Create a 1-dim host mdarray. */ diff --git a/cpp/test/mdarray.cu b/cpp/test/mdarray.cu index 9a1b996ef1..60860f90f4 100644 --- a/cpp/test/mdarray.cu +++ b/cpp/test/mdarray.cu @@ -80,12 +80,12 @@ void test_mdarray_basic() /** * device policy */ - stdex::layout_right::mapping layout{matrix_extent{4, 4}}; - using mdarray_t = device_mdarray; + layout_c_contiguous::mapping layout{matrix_extent{4, 4}}; + using mdarray_t = device_mdarray; auto policy = mdarray_t::container_policy_type{s}; static_assert(std::is_same_v>); - device_mdarray array{layout, policy}; + device_mdarray array{layout, policy}; array(0, 3) = 1; ASSERT_EQ(array(0, 3), 1); @@ -132,12 +132,12 @@ void test_mdarray_basic() /** * host policy */ - using mdarray_t = host_mdarray; + using mdarray_t = host_mdarray; mdarray_t::container_policy_type policy; static_assert( std::is_same_v>); - stdex::layout_right::mapping layout{matrix_extent{4, 4}}; - host_mdarray array{layout, policy}; + layout_c_contiguous::mapping layout{matrix_extent{4, 4}}; + host_mdarray array{layout, policy}; array(0, 3) = 1; ASSERT_EQ(array(0, 3), 1); @@ -156,10 +156,10 @@ void test_mdarray_basic() * static extent */ using static_extent = stdex::extents<16, 16>; - stdex::layout_right::mapping layout{static_extent{}}; - using mdarray_t = device_mdarray; + layout_c_contiguous::mapping layout{static_extent{}}; + using mdarray_t = device_mdarray; mdarray_t::container_policy_type policy{s}; - device_mdarray array{layout, policy}; + device_mdarray array{layout, policy}; static_assert(array.rank_dynamic() == 0); static_assert(array.rank() == 2); @@ -181,7 +181,7 @@ template void test_mdarray_copy_move(ThrustPolicy exec, PolicyFn make_policy) { using matrix_extent = stdex::extents; - stdex::layout_right::mapping layout{matrix_extent{4, 4}}; + layout_c_contiguous::mapping layout{matrix_extent{4, 4}}; using mdarray_t = BasicMDarray; using policy_t = typename mdarray_t::container_policy_type; @@ -287,7 +287,8 @@ TEST(MDArray, CopyMove) } } -TEST(MDArray, Factory) +namespace { +void test_factory_methods() { size_t n{100}; rmm::device_uvector d_vec(n, rmm::cuda_stream_default); @@ -347,7 +348,40 @@ TEST(MDArray, Factory) static_assert(d_vec.rank() == 1); ASSERT_EQ(d_vec.extent(0), n); } + + { + // device scalar + auto d_scalar = make_device_scalar(17.0, rmm::cuda_stream_default); + static_assert(d_scalar.rank() == 1); + static_assert(d_scalar.rank_dynamic() == 0); + ASSERT_EQ(d_scalar(0), 17.0); + + auto view = d_scalar.view(); + thrust::device_vector status(1, 0); + auto p_status = status.data().get(); + thrust::for_each_n(rmm::exec_policy(rmm::cuda_stream_default), + thrust::make_counting_iterator(0), + 1, + [=] __device__(auto i) { + if (view(i) != 17.0) { myAtomicAdd(p_status, 1); } + }); + check_status(p_status, rmm::cuda_stream_default); + } + { + // host scalar + auto h_scalar = make_host_scalar(17.0); + static_assert(h_scalar.rank() == 1); + static_assert(h_scalar.rank_dynamic() == 0); + ASSERT_EQ(h_scalar(0), 17.0); + ASSERT_EQ(h_scalar.view()(0), 17.0); + + auto view = make_host_scalar_view(h_scalar.data()); + ASSERT_EQ(view(0), 17.0); + } } +} // anonymous namespace + +TEST(MDArray, Factory) { test_factory_methods(); } namespace { template @@ -371,7 +405,8 @@ TEST(MDArray, FuncArg) check_matrix_layout(d_matrix.view()); } { - auto d_matrix = make_device_matrix(10, 10, rmm::cuda_stream_default); + auto d_matrix = + make_device_matrix(10, 10, rmm::cuda_stream_default); check_matrix_layout(d_matrix.view()); // FIXME(jiamingy): The slice has a default accessor instead of accessor_mixin, due to From c02fbaf8c9022d48e5cabe741839c521e9e4f2ec Mon Sep 17 00:00:00 2001 From: fis Date: Tue, 15 Feb 2022 03:20:47 +0800 Subject: [PATCH 12/14] Update cmake module. --- cpp/cmake/thirdparty/get_mdspan.cmake | 4 ++-- cpp/test/CMakeLists.txt | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/cpp/cmake/thirdparty/get_mdspan.cmake b/cpp/cmake/thirdparty/get_mdspan.cmake index 0830205c9d..c88d4e6857 100644 --- a/cpp/cmake/thirdparty/get_mdspan.cmake +++ b/cpp/cmake/thirdparty/get_mdspan.cmake @@ -5,8 +5,8 @@ function(find_and_configure_mdspan VERSION) BUILD_EXPORT_SET raft-exports INSTALL_EXPORT_SET raft-exports CPM_ARGS - GIT_REPOSITORY https://github.com/trivialfis/mdspan - GIT_TAG 0193f075e977cc5f3c957425fd899e53d598f524 + GIT_REPOSITORY https://github.com/rapidsai/mdspan.git + GIT_TAG b3042485358d2ee168ae2b486c98c2c61ec5aec1 OPTIONS "MDSPAN_ENABLE_CUDA ON" "MDSPAN_CXX_STANDARD ON" ) diff --git a/cpp/test/CMakeLists.txt b/cpp/test/CMakeLists.txt index 7749efe624..221eef2a79 100644 --- a/cpp/test/CMakeLists.txt +++ b/cpp/test/CMakeLists.txt @@ -36,7 +36,6 @@ add_executable(test_raft test/eigen_solvers.cu test/handle.cpp test/integer_utils.cpp - test/mdarray.cu test/nvtx.cpp test/pow2_utils.cu test/label/label.cu @@ -65,6 +64,7 @@ add_executable(test_raft test/matrix/math.cu test/matrix/matrix.cu test/matrix/linewise_op.cu + test/mdarray.cu test/mr/device/buffer.cpp test/mr/host/buffer.cpp test/mst.cu From 96f3484ff0ec02d5466d505ebea7ae0d5d917a4b Mon Sep 17 00:00:00 2001 From: fis Date: Tue, 15 Feb 2022 03:58:12 +0800 Subject: [PATCH 13/14] Add license. --- cpp/include/raft/mdarray.hpp | 7 +++++ thirdparty/LICENSES/mdarray.license | 42 +++++++++++++++++++++++++++++ 2 files changed, 49 insertions(+) create mode 100644 thirdparty/LICENSES/mdarray.license diff --git a/cpp/include/raft/mdarray.hpp b/cpp/include/raft/mdarray.hpp index 81808772b1..44ca526c16 100644 --- a/cpp/include/raft/mdarray.hpp +++ b/cpp/include/raft/mdarray.hpp @@ -1,3 +1,10 @@ +/* + * Copyright (2019) Sandia Corporation + * + * The source code is licensed under the 3-clause BSD license found in the LICENSE file + * thirdparty/LICENSES/mdarray.license + */ + /* * Copyright (c) 2022, NVIDIA CORPORATION. * diff --git a/thirdparty/LICENSES/mdarray.license b/thirdparty/LICENSES/mdarray.license new file mode 100644 index 0000000000..e636b86032 --- /dev/null +++ b/thirdparty/LICENSES/mdarray.license @@ -0,0 +1,42 @@ +/* +//@HEADER +// ************************************************************************ +// +// Kokkos v. 2.0 +// Copyright (2019) Sandia Corporation +// +// Under the terms of Contract DE-AC04-94AL85000 with Sandia Corporation, +// the U.S. Government retains certain rights in this software. +// +// Redistribution and use in source and binary forms, with or without +// modification, are permitted provided that the following conditions are +// met: +// +// 1. Redistributions of source code must retain the above copyright +// notice, this list of conditions and the following disclaimer. +// +// 2. Redistributions in binary form must reproduce the above copyright +// notice, this list of conditions and the following disclaimer in the +// documentation and/or other materials provided with the distribution. +// +// 3. Neither the name of the Corporation nor the names of the +// contributors may be used to endorse or promote products derived from +// this software without specific prior written permission. +// +// THIS SOFTWARE IS PROVIDED BY SANDIA CORPORATION "AS IS" AND ANY +// EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR +// PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL SANDIA CORPORATION OR THE +// CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, +// EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, +// PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR +// PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF +// LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING +// NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS +// SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. +// +// Questions? Contact Christian R. Trott (crtrott@sandia.gov) +// +// ************************************************************************ +//@HEADER +*/ \ No newline at end of file From 94dc6099c98e3e6099291b1bfbe431ef81c1adb1 Mon Sep 17 00:00:00 2001 From: "Corey J. Nolet" Date: Thu, 17 Feb 2022 15:56:54 -0500 Subject: [PATCH 14/14] Adding copyright to detail/mdarray.hpp --- cpp/include/raft/detail/mdarray.hpp | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/cpp/include/raft/detail/mdarray.hpp b/cpp/include/raft/detail/mdarray.hpp index 2cc87c409a..9f0f275eaa 100644 --- a/cpp/include/raft/detail/mdarray.hpp +++ b/cpp/include/raft/detail/mdarray.hpp @@ -1,3 +1,10 @@ +/* + * Copyright (2019) Sandia Corporation + * + * The source code is licensed under the 3-clause BSD license found in the LICENSE file + * thirdparty/LICENSES/mdarray.license + */ + /* * Copyright (c) 2022, NVIDIA CORPORATION. *