From c33d4bb839fc9bd514849d0241eb70f835b0a7ac Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Tue, 18 Jul 2023 09:00:13 -0500 Subject: [PATCH 01/14] Only run cugraph conda CI for CUDA 11. (#3713) We are planning to enable CUDA 12 CI jobs on July 18, per https://github.com/rapidsai/shared-action-workflows/issues/112. This PR must be merged to disable CUDA 12 CI jobs for cugraph until #3271 / #3456 are merged. Authors: - Bradley Dice (https://github.com/bdice) Approvers: - Ray Douglass (https://github.com/raydouglass) --- .github/workflows/build.yaml | 4 +++- .github/workflows/pr.yaml | 8 ++++++-- .github/workflows/test.yaml | 2 ++ 3 files changed, 11 insertions(+), 3 deletions(-) diff --git a/.github/workflows/build.yaml b/.github/workflows/build.yaml index de9caa0fabe..cf73e1d2d27 100644 --- a/.github/workflows/build.yaml +++ b/.github/workflows/build.yaml @@ -30,6 +30,7 @@ jobs: secrets: inherit uses: rapidsai/shared-action-workflows/.github/workflows/conda-cpp-build.yaml@branch-23.08 with: + matrix_filter: map(select(.CUDA_VER | startswith("11"))) build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} date: ${{ inputs.date }} @@ -39,6 +40,7 @@ jobs: secrets: inherit uses: rapidsai/shared-action-workflows/.github/workflows/conda-python-build.yaml@branch-23.08 with: + matrix_filter: map(select(.CUDA_VER | startswith("11"))) build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} date: ${{ inputs.date }} @@ -61,7 +63,7 @@ jobs: arch: "amd64" branch: ${{ inputs.branch }} build_type: ${{ inputs.build_type || 'branch' }} - container_image: "rapidsai/ci:latest" + container_image: "rapidsai/ci:cuda11.8.0-ubuntu22.04-py3.10" date: ${{ inputs.date }} node_type: "gpu-v100-latest-1" run_script: "ci/build_docs.sh" diff --git a/.github/workflows/pr.yaml b/.github/workflows/pr.yaml index 4d52cd26de4..d9029ea37a1 100644 --- a/.github/workflows/pr.yaml +++ b/.github/workflows/pr.yaml @@ -35,6 +35,7 @@ jobs: secrets: inherit uses: rapidsai/shared-action-workflows/.github/workflows/conda-cpp-build.yaml@branch-23.08 with: + matrix_filter: map(select(.CUDA_VER | startswith("11"))) build_type: pull-request node_type: cpu16 conda-cpp-tests: @@ -42,18 +43,21 @@ jobs: secrets: inherit uses: rapidsai/shared-action-workflows/.github/workflows/conda-cpp-tests.yaml@branch-23.08 with: + matrix_filter: map(select(.CUDA_VER | startswith("11"))) build_type: pull-request conda-python-build: needs: conda-cpp-build secrets: inherit uses: rapidsai/shared-action-workflows/.github/workflows/conda-python-build.yaml@branch-23.08 with: + matrix_filter: map(select(.CUDA_VER | startswith("11"))) build_type: pull-request conda-python-tests: needs: conda-python-build secrets: inherit uses: rapidsai/shared-action-workflows/.github/workflows/conda-python-tests.yaml@branch-23.08 with: + matrix_filter: map(select(.CUDA_VER | startswith("11"))) build_type: pull-request conda-notebook-tests: needs: conda-python-build @@ -63,7 +67,7 @@ jobs: build_type: pull-request node_type: "gpu-v100-latest-1" arch: "amd64" - container_image: "rapidsai/ci:latest" + container_image: "rapidsai/ci:cuda11.8.0-ubuntu22.04-py3.10" run_script: "ci/test_notebooks.sh" docs-build: needs: conda-python-build @@ -73,7 +77,7 @@ jobs: build_type: pull-request node_type: "gpu-v100-latest-1" arch: "amd64" - container_image: "rapidsai/ci:latest" + container_image: "rapidsai/ci:cuda11.8.0-ubuntu22.04-py3.10" run_script: "ci/build_docs.sh" wheel-build-pylibcugraph: needs: checks diff --git a/.github/workflows/test.yaml b/.github/workflows/test.yaml index d697b8f1649..1b8cfaf25b7 100644 --- a/.github/workflows/test.yaml +++ b/.github/workflows/test.yaml @@ -18,6 +18,7 @@ jobs: secrets: inherit uses: rapidsai/shared-action-workflows/.github/workflows/conda-cpp-tests.yaml@branch-23.08 with: + matrix_filter: map(select(.CUDA_VER | startswith("11"))) build_type: nightly branch: ${{ inputs.branch }} date: ${{ inputs.date }} @@ -26,6 +27,7 @@ jobs: secrets: inherit uses: rapidsai/shared-action-workflows/.github/workflows/conda-python-tests.yaml@branch-23.08 with: + matrix_filter: map(select(.CUDA_VER | startswith("11"))) build_type: nightly branch: ${{ inputs.branch }} date: ${{ inputs.date }} From 32e6e5184cf63ac989356d6e2b3f93eab63073a4 Mon Sep 17 00:00:00 2001 From: Naim <110031745+naimnv@users.noreply.github.com> Date: Tue, 18 Jul 2023 16:53:18 +0200 Subject: [PATCH 02/14] [BUG] Fix namesapce to default_hash and hash_functions (#3711) Fix namesapce to default_hash and hash_functions Authors: - Naim (https://github.com/naimnv) - Seunghwa Kang (https://github.com/seunghwak) Approvers: - Seunghwa Kang (https://github.com/seunghwak) - Chuck Hastings (https://github.com/ChuckHastings) URL: https://github.com/rapidsai/cugraph/pull/3711 --- .../cugraph/utilities/device_functors.cuh | 16 +- .../include/hash/concurrent_unordered_map.cuh | 5 +- cpp/src/community/detail/common_methods.cuh | 6 + cpp/src/community/detail/refine_impl.cuh | 5 + cpp/src/prims/detail/nbr_intersection.cuh | 4 +- cpp/src/prims/kv_store.cuh | 641 +++++++++++++++--- ..._v_pair_transform_dst_nbr_intersection.cuh | 2 +- ...r_v_random_select_transform_outgoing_e.cuh | 11 +- ...m_reduce_dst_key_aggregated_outgoing_e.cuh | 2 +- cpp/src/structure/relabel_impl.cuh | 4 +- 10 files changed, 568 insertions(+), 128 deletions(-) diff --git a/cpp/include/cugraph/utilities/device_functors.cuh b/cpp/include/cugraph/utilities/device_functors.cuh index d29e7c47d14..501e74cf47b 100644 --- a/cpp/include/cugraph/utilities/device_functors.cuh +++ b/cpp/include/cugraph/utilities/device_functors.cuh @@ -57,16 +57,28 @@ struct pack_bool_t { } }; -template +template struct indirection_t { Iterator first{}; - __device__ typename thrust::iterator_traits::value_type operator()(size_t i) const + __device__ typename thrust::iterator_traits::value_type operator()(index_t i) const { return *(first + i); } }; +template +struct indirection_if_idx_valid_t { + Iterator first{}; + index_t invalid_idx{}; + typename thrust::iterator_traits::value_type invalid_value{}; + + __device__ typename thrust::iterator_traits::value_type operator()(index_t i) const + { + return (i != invalid_idx) ? *(first + i) : invalid_value; + } +}; + template struct not_equal_t { T compare{}; diff --git a/cpp/libcugraph_etl/include/hash/concurrent_unordered_map.cuh b/cpp/libcugraph_etl/include/hash/concurrent_unordered_map.cuh index f097f9c43a2..ab14ff6c685 100644 --- a/cpp/libcugraph_etl/include/hash/concurrent_unordered_map.cuh +++ b/cpp/libcugraph_etl/include/hash/concurrent_unordered_map.cuh @@ -27,7 +27,8 @@ #include #include -#include +#include +#include #include #include @@ -118,7 +119,7 @@ union pair_packer()>> { */ template , + typename Hasher = cudf::hashing::detail::default_hash, typename Equality = equal_to, typename Allocator = default_allocator>> class concurrent_unordered_map { diff --git a/cpp/src/community/detail/common_methods.cuh b/cpp/src/community/detail/common_methods.cuh index 62ede6eaafb..b388ba53e81 100644 --- a/cpp/src/community/detail/common_methods.cuh +++ b/cpp/src/community/detail/common_methods.cuh @@ -18,6 +18,7 @@ #include #include +#include #include #include #include @@ -42,6 +43,11 @@ CUCO_DECLARE_BITWISE_COMPARABLE(float) CUCO_DECLARE_BITWISE_COMPARABLE(double) +// FIXME: a temporary workaround for a compiler error, should be deleted once cuco gets patched. +namespace cuco { +template <> +struct is_bitwise_comparable> : std::true_type {}; +} // namespace cuco namespace cugraph { namespace detail { diff --git a/cpp/src/community/detail/refine_impl.cuh b/cpp/src/community/detail/refine_impl.cuh index bbd720131de..e811aafc776 100644 --- a/cpp/src/community/detail/refine_impl.cuh +++ b/cpp/src/community/detail/refine_impl.cuh @@ -48,6 +48,11 @@ CUCO_DECLARE_BITWISE_COMPARABLE(float) CUCO_DECLARE_BITWISE_COMPARABLE(double) +// FIXME: a temporary workaround for a compiler error, should be deleted once cuco gets patched. +namespace cuco { +template <> +struct is_bitwise_comparable> : std::true_type {}; +} // namespace cuco namespace cugraph { namespace detail { diff --git a/cpp/src/prims/detail/nbr_intersection.cuh b/cpp/src/prims/detail/nbr_intersection.cuh index 2d0d0a876e6..98453d46c3f 100644 --- a/cpp/src/prims/detail/nbr_intersection.cuh +++ b/cpp/src/prims/detail/nbr_intersection.cuh @@ -974,7 +974,7 @@ nbr_intersection(raft::handle_t const& handle, .get_stream()); // initially store minimum degrees (upper bound for intersection sizes) if (intersect_minor_nbr[0] && intersect_minor_nbr[1]) { auto second_element_to_idx_map = - detail::kv_cuco_store_device_view_t((*major_to_idx_map_ptr)->view()); + detail::kv_cuco_store_find_device_view_t((*major_to_idx_map_ptr)->view()); thrust::transform( handle.get_thrust_policy(), get_dataframe_buffer_begin(vertex_pair_buffer), @@ -1005,7 +1005,7 @@ nbr_intersection(raft::handle_t const& handle, handle.get_stream()); if (intersect_minor_nbr[0] && intersect_minor_nbr[1]) { auto second_element_to_idx_map = - detail::kv_cuco_store_device_view_t((*major_to_idx_map_ptr)->view()); + detail::kv_cuco_store_find_device_view_t((*major_to_idx_map_ptr)->view()); thrust::tabulate( handle.get_thrust_policy(), rx_v_pair_nbr_intersection_sizes.begin(), diff --git a/cpp/src/prims/kv_store.cuh b/cpp/src/prims/kv_store.cuh index 8395fc55833..f20865c92dc 100644 --- a/cpp/src/prims/kv_store.cuh +++ b/cpp/src/prims/kv_store.cuh @@ -16,11 +16,14 @@ #pragma once #include +#include #include +#include #include #include +#include #include #include #include @@ -29,6 +32,7 @@ #include #include #include +#include #include #include @@ -45,7 +49,7 @@ namespace cugraph { namespace detail { template -struct binary_search_find_op_t { +struct kv_binary_search_find_op_t { using key_type = typename thrust::iterator_traits::value_type; using value_type = typename thrust::iterator_traits::value_type; @@ -67,7 +71,7 @@ struct binary_search_find_op_t { }; template -struct binary_search_contains_op_t { +struct kv_binary_search_contains_op_t { using key_type = typename thrust::iterator_traits::value_type; KeyIterator store_key_first{}; @@ -79,6 +83,105 @@ struct binary_search_contains_op_t { } }; +template +struct kv_cuco_insert_and_increment_t { + using key_type = typename thrust::iterator_traits::value_type; + using cuco_store_type = cuco::experimental::static_map< + key_type, + size_t, + cuco::experimental::extent, + cuda::thread_scope_device, + thrust::equal_to, + cuco::experimental::linear_probing<1, // CG size + cuco::murmurhash3_32>, + rmm::mr::stream_allocator_adaptor>>; + + typename cuco_store_type::ref_type device_ref{}; + KeyIterator key_first{}; + size_t* counter{nullptr}; + size_t invalid_idx{}; + + __device__ size_t operator()(size_t i) + { + auto pair = thrust::make_tuple(*(key_first + i), size_t{0} /* dummy */); + auto [iter, inserted] = device_ref.insert_and_find(pair); + if (inserted) { + cuda::atomic_ref atomic_counter(*counter); + auto idx = atomic_counter.fetch_add(size_t{1}, cuda::std::memory_order_relaxed); + using ref_type = typename cuco_store_type::ref_type; + cuda::atomic_ref ref((*iter).second); + ref.store(idx, cuda::std::memory_order_relaxed); + return idx; + } else { + return invalid_idx; + } + } +}; + +template +struct kv_cuco_insert_if_and_increment_t { + using key_type = typename thrust::iterator_traits::value_type; + using cuco_store_type = cuco::experimental::static_map< + key_type, + size_t, + cuco::experimental::extent, + cuda::thread_scope_device, + thrust::equal_to, + cuco::experimental::linear_probing<1, // CG size + cuco::murmurhash3_32>, + rmm::mr::stream_allocator_adaptor>>; + + typename cuco_store_type::ref_type device_ref{}; + KeyIterator key_first{}; + StencilIterator stencil_first{}; + PredOp pred_op{}; + size_t* counter{nullptr}; + size_t invalid_idx{}; + + __device__ size_t operator()(size_t i) + { + if (pred_op(*(stencil_first + i)) == false) { return invalid_idx; } + + auto pair = thrust::make_tuple(*(key_first + i), size_t{0} /* dummy */); + auto [iter, inserted] = device_ref.insert_and_find(pair); + if (inserted) { + cuda::atomic_ref atomic_counter(*counter); + auto idx = atomic_counter.fetch_add(size_t{1}, cuda::std::memory_order_relaxed); + using ref_type = typename cuco_store_type::ref_type; + cuda::atomic_ref ref((*iter).second); + ref.store(idx, cuda::std::memory_order_relaxed); + return idx; + } else { + return invalid_idx; + } + } +}; + +template +struct kv_cuco_insert_and_assign_t { + using cuco_store_type = cuco::experimental::static_map< + key_t, + std::conditional_t, value_t, size_t>, + cuco::experimental::extent, + cuda::thread_scope_device, + thrust::equal_to, + cuco::experimental::linear_probing<1, // CG size + cuco::murmurhash3_32>, + rmm::mr::stream_allocator_adaptor>>; + + typename cuco_store_type::ref_type device_ref{}; + + __device__ void operator()(thrust::tuple pair) + { + auto [iter, inserted] = device_ref.insert_and_find(pair); + if (!inserted) { + using ref_type = typename cuco_store_type::ref_type; + cuda::atomic_ref ref((*iter).second); + ref.store(thrust::get<1>(pair), cuda::std::memory_order_relaxed); + } + } +}; + template struct kv_binary_search_store_device_view_t { using key_type = typename ViewType::key_type; @@ -112,18 +215,19 @@ struct kv_binary_search_store_device_view_t { }; template -struct kv_cuco_store_device_view_t { - using key_type = typename ViewType::key_type; - using value_type = typename ViewType::value_type; - using cuco_store_device_view_type = typename ViewType::cuco_store_type::device_view; +struct kv_cuco_store_find_device_view_t { + using key_type = typename ViewType::key_type; + using value_type = typename ViewType::value_type; + using cuco_store_device_ref_type = + typename ViewType::cuco_store_type::ref_type; static_assert(!ViewType::binary_search); - __host__ kv_cuco_store_device_view_t(ViewType view) - : cuco_store_device_view(view.cuco_store_device_view()) + __host__ kv_cuco_store_find_device_view_t(ViewType view) + : cuco_store_device_ref(view.cuco_store_find_device_ref()) { if constexpr (std::is_arithmetic_v) { - invalid_value = cuco_store_device_view.get_empty_value_sentinel(); + invalid_value = cuco_store_device_ref.empty_value_sentinel(); } else { store_value_first = view.store_value_first(); invalid_value = view.invalid_value(); @@ -132,11 +236,11 @@ struct kv_cuco_store_device_view_t { __device__ value_type find(key_type key) const { - auto found = cuco_store_device_view.find(key); - if (found == cuco_store_device_view.end()) { + auto found = cuco_store_device_ref.find(key); + if (found == cuco_store_device_ref.end()) { return invalid_value; } else { - auto val = found->second.load(cuda::std::memory_order_relaxed); + auto val = (*found).second; if constexpr (std::is_arithmetic_v) { return val; } else { @@ -145,7 +249,7 @@ struct kv_cuco_store_device_view_t { } } - cuco_store_device_view_type cuco_store_device_view{}; + cuco_store_device_ref_type cuco_store_device_ref{}; std::conditional_t, typename ViewType::value_iterator, std::byte /* dummy */> @@ -185,7 +289,7 @@ class kv_binary_search_store_view_t { key_first, key_last, value_first, - binary_search_find_op_t{ + kv_binary_search_find_op_t{ store_key_first_, store_key_last_, store_value_first_, invalid_value_}); } @@ -195,11 +299,12 @@ class kv_binary_search_store_view_t { ResultValueIterator value_first, rmm::cuda_stream_view stream) const { - thrust::transform(rmm::exec_policy(stream), - key_first, - key_last, - value_first, - binary_search_contains_op_t{store_key_first_, store_key_last_}); + thrust::transform( + rmm::exec_policy(stream), + key_first, + key_last, + value_first, + kv_binary_search_contains_op_t{store_key_first_, store_key_last_}); } KeyIterator store_key_first() const { return store_key_first_; } @@ -227,31 +332,29 @@ class kv_cuco_store_view_t { static constexpr bool binary_search = false; - using cuco_store_type = - cuco::static_map, value_type, size_t>, - cuda::thread_scope_device, - rmm::mr::stream_allocator_adaptor>>; + using cuco_store_type = cuco::experimental::static_map< + key_t, + std::conditional_t, value_type, size_t>, + cuco::experimental::extent, + cuda::thread_scope_device, + thrust::equal_to, + cuco::experimental::linear_probing<1, // CG size + cuco::murmurhash3_32>, + rmm::mr::stream_allocator_adaptor>>; - // FIXME: const_cast as a temporary workaround for - // https://github.com/NVIDIA/cuCollections/issues/242 (cuco find() is not a const function) template kv_cuco_store_view_t(cuco_store_type const* store, std::enable_if_t, int32_t> = 0) - : cuco_store_(const_cast(store)) + : cuco_store_(store) { } - // FIXME: const_cast as a temporary workaround for - // https://github.com/NVIDIA/cuCollections/issues/242 (cuco find() is not a const function) template kv_cuco_store_view_t(cuco_store_type const* store, ValueIterator value_first, type invalid_value, std::enable_if_t, int32_t> = 0) - : cuco_store_(const_cast(store)), - store_value_first_(value_first), - invalid_value_(invalid_value) + : cuco_store_(store), store_value_first_(value_first), invalid_value_(invalid_value) { } @@ -262,34 +365,17 @@ class kv_cuco_store_view_t { rmm::cuda_stream_view stream) const { if constexpr (std::is_arithmetic_v) { - cuco_store_->find(key_first, - key_last, - value_first, - cuco::detail::MurmurHash3_32{}, - thrust::equal_to{}, - stream); + cuco_store_->find(key_first, key_last, value_first, stream.value()); } else { rmm::device_uvector indices(thrust::distance(key_first, key_last), stream); - cuco_store_->find(key_first, - key_last, - indices.begin(), - cuco::detail::MurmurHash3_32{}, - thrust::equal_to{}, - stream); - auto invalid_idx = cuco_store_->get_empty_value_sentinel(); - thrust::transform(rmm::exec_policy(stream), - indices.begin(), - indices.end(), - value_first, - [store_value_first = store_value_first_, - invalid_value = invalid_value_, - invalid_idx] __device__(auto idx) { - if (idx != invalid_idx) { - return *(store_value_first + idx); - } else { - return invalid_value; - } - }); + auto invalid_idx = cuco_store_->empty_value_sentinel(); + cuco_store_->find(key_first, key_last, indices.begin(), stream.value()); + thrust::transform( + rmm::exec_policy(stream), + indices.begin(), + indices.end(), + value_first, + indirection_if_idx_valid_t{store_value_first_, invalid_idx, invalid_value_}); } } @@ -299,15 +385,10 @@ class kv_cuco_store_view_t { ResultValueIterator value_first, rmm::cuda_stream_view stream) const { - cuco_store_->contains(key_first, - key_last, - value_first, - cuco::detail::MurmurHash3_32{}, - thrust::equal_to{}, - stream); + cuco_store_->contains(key_first, key_last, value_first, stream.value()); } - auto cuco_store_device_view() const { return cuco_store_->get_device_view(); } + auto cuco_store_find_device_ref() const { return cuco_store_->ref(cuco::experimental::find); } template std::enable_if_t, ValueIterator> store_value_first() const @@ -315,21 +396,19 @@ class kv_cuco_store_view_t { return store_value_first_; } - key_t invalid_key() const { return cuco_store_->get_empty_key_sentinel(); } + key_t invalid_key() const { return cuco_store_->empty_key_sentinel(); } value_type invalid_value() const { if constexpr (std::is_arithmetic_v) { - return cuco_store_->get_empty_value_sentinel(); + return cuco_store_->empty_value_sentinel(); } else { return invalid_value_; } } private: - // FIXME: cuco_store should be a const pointer but we can't due to - // https://github.com/NVIDIA/cuCollections/issues/242 (cuco find() is not a const function) - cuco_store_type* cuco_store_{}; + cuco_store_type const* cuco_store_{}; std::conditional_t, ValueIterator, std::byte /* dummy */> store_value_first_{}; @@ -395,6 +474,29 @@ class kv_binary_search_store_t { } } + auto retrieve_all(rmm::cuda_stream_view stream) + { + rmm::device_uvector tmp_store_keys(store_keys_.size(), stream); + auto tmp_store_values = + allocate_dataframe_buffer(size_dataframe_buffer(store_values_), stream); + thrust::copy( + rmm::exec_policy(stream), store_keys_.begin(), store_keys_.end(), tmp_store_keys.begin()); + thrust::copy(rmm::exec_policy(stream), + get_dataframe_buffer_begin(store_values_), + get_dataframe_buffer_end(store_values_), + get_dataframe_buffer_begin(tmp_store_values)); + return std::make_tuple(std::move(tmp_store_keys), std::move(tmp_store_values)); + } + + auto release(rmm::cuda_stream_view stream) + { + auto tmp_store_keys = std::move(store_keys_); + auto tmp_store_values = std::move(store_values_); + store_keys_ = rmm::device_uvector(0, stream); + store_values_ = allocate_dataframe_buffer(0, stream); + return std::make_tuple(std::move(tmp_store_keys), std::move(tmp_store_values)); + } + key_t const* store_key_first() const { return store_keys_.cbegin(); } key_t const* store_key_last() const { return store_keys_.cend(); } @@ -403,6 +505,10 @@ class kv_binary_search_store_t { value_t invalid_value() const { return invalid_value_; } + size_t size() const { return store_keys_.size(); } + + size_t capacity() const { return store_keys_.size(); } + private: rmm::device_uvector store_keys_; decltype(allocate_dataframe_buffer(0, rmm::cuda_stream_view{})) store_values_; @@ -421,14 +527,28 @@ class kv_cuco_store_t { std::invoke_result_t), value_buffer_type&>; - using cuco_store_type = - cuco::static_map, value_t, size_t>, - cuda::thread_scope_device, - rmm::mr::stream_allocator_adaptor>>; + using cuco_store_type = cuco::experimental::static_map< + key_t, + std::conditional_t, value_t, size_t>, + cuco::experimental::extent, + cuda::thread_scope_device, + thrust::equal_to, + cuco::experimental::linear_probing<1, // CG size + cuco::murmurhash3_32>, + rmm::mr::stream_allocator_adaptor>>; kv_cuco_store_t(rmm::cuda_stream_view stream) {} + kv_cuco_store_t(size_t capacity, + key_t invalid_key, + value_t invalid_value, + rmm::cuda_stream_view stream) + { + allocate(capacity, invalid_key, invalid_value, stream); + capacity_ = capacity; + size_ = 0; + } + template kv_cuco_store_t(KeyIterator key_first, KeyIterator key_last, @@ -437,51 +557,228 @@ class kv_cuco_store_t { value_t invalid_value, rmm::cuda_stream_view stream) { - double constexpr load_factor = 0.7; - auto num_keys = static_cast(thrust::distance(key_first, key_last)); - auto cuco_size = std::max( - static_cast(static_cast(num_keys) / load_factor), - static_cast(num_keys) + 1); // cuco::static_map requires at least one empty slot - auto stream_adapter = rmm::mr::make_stream_allocator_adaptor( - rmm::mr::polymorphic_allocator(rmm::mr::get_current_device_resource()), stream); + auto num_keys = static_cast(thrust::distance(key_first, key_last)); + allocate(num_keys, invalid_key, invalid_value, stream); + if constexpr (!std::is_arithmetic_v) { invalid_value_ = invalid_value; } + capacity_ = num_keys; + size_ = 0; + + insert(key_first, key_last, value_first, stream); + } + + template + void insert(KeyIterator key_first, + KeyIterator key_last, + ValueIterator value_first, + rmm::cuda_stream_view stream) + { + auto num_keys = static_cast(thrust::distance(key_first, key_last)); + if (num_keys == 0) return; + if constexpr (std::is_arithmetic_v) { - cuco_store_ = - std::make_unique(cuco_size, - cuco::sentinel::empty_key{invalid_key}, - cuco::sentinel::empty_value{invalid_value}, - stream_adapter, - stream); + auto pair_first = thrust::make_zip_iterator(thrust::make_tuple(key_first, value_first)); + size_ += cuco_store_->insert(pair_first, pair_first + num_keys, stream.value()); + } else { + auto old_store_value_size = size_dataframe_buffer(store_values_); + // FIXME: we can use cuda::atomic instead but currently on a system with x86 + GPU, this + // requires placing the atomic variable on managed memory and this adds additional + // complication. + rmm::device_scalar counter(old_store_value_size, stream); + auto mutable_device_ref = cuco_store_->ref(cuco::experimental::insert_and_find); + rmm::device_uvector store_value_offsets(num_keys, stream); + thrust::tabulate( + rmm::exec_policy(stream), + store_value_offsets.begin(), + store_value_offsets.end(), + kv_cuco_insert_and_increment_t{ + mutable_device_ref, key_first, counter.data(), std::numeric_limits::max()}); + size_ += counter.value(stream); + resize_dataframe_buffer(store_values_, size_, stream); + thrust::scatter_if(rmm::exec_policy(stream), + value_first, + value_first + num_keys, + store_value_offsets.begin() /* map */, + store_value_offsets.begin() /* stencil */, + get_dataframe_buffer_begin(store_values_), + not_equal_t{std::numeric_limits::max()}); + } + } + template + void insert_if(KeyIterator key_first, + KeyIterator key_last, + ValueIterator value_first, + StencilIterator stencil_first, + PredOp pred_op, + rmm::cuda_stream_view stream) + { + auto num_keys = static_cast(thrust::distance(key_first, key_last)); + if (num_keys == 0) return; + + if constexpr (std::is_arithmetic_v) { auto pair_first = thrust::make_zip_iterator(thrust::make_tuple(key_first, value_first)); - cuco_store_->insert(pair_first, - pair_first + num_keys, - cuco::detail::MurmurHash3_32{}, - thrust::equal_to{}, - stream); + size_ += cuco_store_->insert_if( + pair_first, pair_first + num_keys, stencil_first, pred_op, stream.value()); } else { - cuco_store_ = std::make_unique( - cuco_size, - cuco::sentinel::empty_key{invalid_key}, - cuco::sentinel::empty_value{std::numeric_limits::max()}, - stream_adapter, - stream); - store_values_ = allocate_dataframe_buffer(num_keys, stream); - invalid_value_ = invalid_value; + auto old_store_value_size = size_dataframe_buffer(store_values_); + // FIXME: we can use cuda::atomic instead but currently on a system with x86 + GPU, this + // requires placing the atomic variable on managed memory and this adds additional + // complication. + rmm::device_scalar counter(old_store_value_size, stream); + auto mutable_device_ref = cuco_store_->ref(cuco::experimental::insert_and_find); + rmm::device_uvector store_value_offsets(num_keys, stream); + thrust::tabulate(rmm::exec_policy(stream), + store_value_offsets.begin(), + store_value_offsets.end(), + kv_cuco_insert_if_and_increment_t{ + mutable_device_ref, + key_first, + stencil_first, + pred_op, + counter.data(), + std::numeric_limits::max()}); + size_ += counter.value(stream); + resize_dataframe_buffer(store_values_, size_, stream); + thrust::scatter_if(rmm::exec_policy(stream), + value_first, + value_first + num_keys, + store_value_offsets.begin() /* map */, + store_value_offsets.begin() /* stencil */, + get_dataframe_buffer_begin(store_values_), + not_equal_t{std::numeric_limits::max()}); + } + } + + template + void insert_and_assign(KeyIterator key_first, + KeyIterator key_last, + ValueIterator value_first, + rmm::cuda_stream_view stream) + { + auto num_keys = static_cast(thrust::distance(key_first, key_last)); + if (num_keys == 0) return; + if constexpr (std::is_arithmetic_v) { + auto pair_first = thrust::make_zip_iterator(thrust::make_tuple(key_first, value_first)); + // FIXME: a temporary solution till insert_and_assign is added to + // cuco::experimental::static_map + auto mutable_device_ref = cuco_store_->ref(cuco::experimental::insert_and_find); + thrust::for_each(rmm::exec_policy(stream), + pair_first, + pair_first + num_keys, + detail::kv_cuco_insert_and_assign_t{mutable_device_ref}); + // FIXME: this is an upper bound of size_, as some inserts may fail due to existing keys + size_ += num_keys; + } else { + auto old_store_value_size = size_dataframe_buffer(store_values_); + // FIXME: we can use cuda::atomic instead but currently on a system with x86 + GPU, this + // requires placing the atomic variable on managed memory and this adds additional + // complication. + rmm::device_scalar counter(old_store_value_size, stream); + auto mutable_device_ref = cuco_store_->ref(cuco::experimental::insert_and_find); + rmm::device_uvector store_value_offsets(num_keys, stream); + thrust::tabulate( + rmm::exec_policy(stream), + store_value_offsets.begin(), + store_value_offsets.end(), + kv_cuco_insert_and_increment_t{ + mutable_device_ref, key_first, counter.data(), std::numeric_limits::max()}); + size_ += counter.value(stream); + resize_dataframe_buffer(store_values_, size_, stream); + thrust::scatter_if(rmm::exec_policy(stream), + value_first, + value_first + num_keys, + store_value_offsets.begin() /* map */, + store_value_offsets.begin() /* stencil */, + get_dataframe_buffer_begin(store_values_), + not_equal_t{std::numeric_limits::max()}); + + // now perform assigns (for k,v pairs that failed to insert) + + rmm::device_uvector kv_indices(num_keys, stream); + thrust::sequence(rmm::exec_policy(), kv_indices.begin(), kv_indices.end(), size_t{0}); auto pair_first = thrust::make_zip_iterator( - thrust::make_tuple(key_first, thrust::make_counting_iterator(size_t{0}))); - cuco_store_->insert(pair_first, - pair_first + num_keys, - cuco::detail::MurmurHash3_32{}, - thrust::equal_to{}, - stream); - thrust::copy(rmm::exec_policy(stream), - value_first, - value_first + num_keys, - get_dataframe_buffer_begin(store_values_)); + thrust::make_tuple(store_value_offsets.begin(), kv_indices.begin())); + kv_indices.resize( + thrust::distance( + pair_first, + thrust::remove_if(rmm::exec_policy(stream), + pair_first, + pair_first + num_keys, + [invalid_idx = std::numeric_limits::max()] __device__( + auto pair) { return thrust::get<0>(pair) != invalid_idx; })), + stream); + store_value_offsets.resize(0, stream); + store_value_offsets.shrink_to_fit(stream); + + thrust::sort(rmm::exec_policy(stream), + kv_indices.begin(), + kv_indices.end(), + [key_first] __device__(auto lhs, auto rhs) { + return *(key_first + lhs) < *(key_first + rhs); + }); + kv_indices.resize(thrust::distance(kv_indices.begin(), + thrust::unique(rmm::exec_policy(stream), + kv_indices.begin(), + kv_indices.end(), + [key_first] __device__(auto lhs, auto rhs) { + return *(key_first + lhs) == + *(key_first + rhs); + })), + stream); + + thrust::for_each( + rmm::exec_policy(stream), + kv_indices.begin(), + kv_indices.end(), + [key_first, + value_first, + store_value_first = get_dataframe_buffer_begin(store_values_), + device_ref = cuco_store_->ref(cuco::experimental::find)] __device__(auto kv_idx) { + size_t store_value_offset{}; + auto found = device_ref.find(*(key_first + kv_idx)); + assert(found != device_ref.end()); + store_value_offset = (*found).second; + *(store_value_first + store_value_offset) = *(value_first + kv_idx); + }); } } + auto retrieve_all(rmm::cuda_stream_view stream) + { + rmm::device_uvector keys(size_, stream); + auto values = allocate_dataframe_buffer(0, stream); + if constexpr (std::is_arithmetic_v) { + values.resize(size_, stream); + auto pair_last = cuco_store_->retrieve_all(keys.begin(), values.begin(), stream.value()); + // FIXME: this resize (& shrink_to_fit) shouldn't be necessary if size_ is exact + keys.resize(thrust::distance(keys.begin(), std::get<0>(pair_last)), stream); + values.resize(keys.size(), stream); + } else { + rmm::device_uvector indices(size_, stream); + auto pair_last = cuco_store_->retrieve_all(keys.begin(), indices.begin(), stream.value()); + // FIXME: this resize (& shrink_to_fit) shouldn't be necessary if size_ is exact + keys.resize(thrust::distance(keys.begin(), std::get<0>(pair_last)), stream); + indices.resize(keys.size(), stream); + resize_dataframe_buffer(values, keys.size(), stream); + thrust::gather(rmm::exec_policy(stream), + indices.begin(), + indices.end(), + get_dataframe_buffer_begin(store_values_), + get_dataframe_buffer_begin(values)); + } + return std::make_tuple(std::move(keys), std::move(values)); + } + + auto release(rmm::cuda_stream_view stream) + { + auto [retrieved_keys, retrieved_values] = retrieve_all(stream); + allocate(0, invalid_key(), invalid_value(), stream); + capacity_ = 0; + size_ = 0; + return std::make_tuple(std::move(retrieved_keys), std::move(retrieved_values)); + } + cuco_store_type const* cuco_store_ptr() const { return cuco_store_.get(); } template @@ -490,18 +787,60 @@ class kv_cuco_store_t { return get_dataframe_buffer_cbegin(store_values_); } - key_t invalid_key() const { return cuco_store_.get_empty_key_sentinel(); } + key_t invalid_key() const { return cuco_store_->empty_key_sentinel(); } value_t invalid_value() const { if constexpr (std::is_arithmetic_v) { - return cuco_store_.get_empty_value_sentinel(); + return cuco_store_->empty_value_sentinel(); } else { return invalid_value_; } } + // FIXME: currently this returns an upper-bound + size_t size() const { return size_; } + + size_t capacity() const { return capacity_; } + private: + void allocate(size_t num_keys, + key_t invalid_key, + value_t invalid_value, + rmm::cuda_stream_view stream) + { + double constexpr load_factor = 0.7; + auto cuco_size = std::max( + static_cast(static_cast(num_keys) / load_factor), + static_cast(num_keys) + 1); // cuco::static_map requires at least one empty slot + + auto stream_adapter = rmm::mr::make_stream_allocator_adaptor( + rmm::mr::polymorphic_allocator(rmm::mr::get_current_device_resource()), stream); + if constexpr (std::is_arithmetic_v) { + cuco_store_ = std::make_unique( + cuco_size, + cuco::sentinel::empty_key{invalid_key}, + cuco::sentinel::empty_value{invalid_value}, + thrust::equal_to{}, + cuco::experimental::linear_probing<1, // CG size + cuco::murmurhash3_32>{}, + stream_adapter, + stream.value()); + } else { + cuco_store_ = std::make_unique( + cuco_size, + cuco::sentinel::empty_key{invalid_key}, + cuco::sentinel::empty_value{std::numeric_limits::max()}, + thrust::equal_to{}, + cuco::experimental::linear_probing<1, // CG size + cuco::murmurhash3_32>{}, + stream_adapter, + stream); + store_values_ = allocate_dataframe_buffer(0, stream); + reserve_dataframe_buffer(store_values_, num_keys, stream); + } + } + std::unique_ptr cuco_store_{nullptr}; std::conditional_t, decltype(allocate_dataframe_buffer(0, rmm::cuda_stream_view{})), @@ -510,6 +849,12 @@ class kv_cuco_store_t { std::conditional_t, value_t, std::byte /* dummy */> invalid_value_{}; + size_t capacity_{0}; + size_t size_{ + 0}; // caching as cuco_store_->size() is expensive (this scans the entire slots to handle + // user inserts through a device reference (and currently this is an upper bound (this + // will become exact once we fully switch to cuco::experimental::static_map and use the + // static_map class's insert_and_assign function; this function will be added soon) }; } // namespace detail @@ -528,6 +873,23 @@ class kv_store_t { kv_store_t(rmm::cuda_stream_view stream) : store_(stream) {} + /* when use_binary_search = false */ + template + kv_store_t( + size_t capacity /* one can expect good performance till the capacity, the actual underlying + capacity can be larger (for performance & correctness reasons) */ + , + key_t invalid_key /* invalid key shouldn't appear in any *iter in [key_first, key_last) */, + value_t invalid_value /* invalid_value shouldn't appear in any *iter in [value_first, + value_first + thrust::distance(key_first, key_last)), invalid_value is + returned when match fails for the given key */ + , + rmm::cuda_stream_view stream, + std::enable_if_t = 0) + : store_(capacity, invalid_key, invalid_value, stream) + { + } + /* when use_binary_search = true */ template kv_store_t( @@ -576,6 +938,47 @@ class kv_store_t { { } + /* when use binary_search = false, this requires that the capacity is large enough */ + template + std::enable_if_t insert(KeyIterator key_first, + KeyIterator key_last, + ValueIterator value_first, + rmm::cuda_stream_view stream) + { + store_.insert(key_first, key_last, value_first, stream); + } + + /* when use binary_search = false, this requires that the capacity is large enough */ + template + std::enable_if_t insert_if(KeyIterator key_first, + KeyIterator key_last, + ValueIterator value_first, + StencilIterator stencil_first, + PredOp pred_op, + rmm::cuda_stream_view stream) + { + store_.insert_if(key_first, key_last, value_first, stencil_first, pred_op, stream); + } + + /* when use binary_search = false, this requires that the capacity is large enough */ + template + std::enable_if_t insert_and_assign(KeyIterator key_first, + KeyIterator key_last, + ValueIterator value_first, + rmm::cuda_stream_view stream) + { + store_.insert_and_assign(key_first, key_last, value_first, stream); + } + + auto retrieve_all(rmm::cuda_stream_view stream) const { return store_.retrieve_all(stream); } + + // kv_store_t becomes empty after release + auto release(rmm::cuda_stream_view stream) { return store_.release(stream); } + auto view() const { if constexpr (use_binary_search) { @@ -593,6 +996,18 @@ class kv_store_t { } } + template + std::enable_if_t invalid_key() const + { + return store_.invalid_key(); + } + + value_t invalid_value() const { return store_.invalid_value(); } + + size_t size() const { return store_.size(); } + + size_t capacity() const { return store_.capacity(); } + private: std::conditional_t, diff --git a/cpp/src/prims/per_v_pair_transform_dst_nbr_intersection.cuh b/cpp/src/prims/per_v_pair_transform_dst_nbr_intersection.cuh index f30de0750e3..d69bb8af25e 100644 --- a/cpp/src/prims/per_v_pair_transform_dst_nbr_intersection.cuh +++ b/cpp/src/prims/per_v_pair_transform_dst_nbr_intersection.cuh @@ -346,7 +346,7 @@ void per_v_pair_transform_dst_nbr_intersection( // partition? This may provide additional performance improvement opportunities??? auto chunk_vertex_pair_first = thrust::make_transform_iterator( chunk_vertex_pair_index_first, - detail::indirection_t{vertex_pair_first}); + detail::indirection_t{vertex_pair_first}); auto [intersection_offsets, intersection_indices] = detail::nbr_intersection(handle, graph_view, diff --git a/cpp/src/prims/per_v_random_select_transform_outgoing_e.cuh b/cpp/src/prims/per_v_random_select_transform_outgoing_e.cuh index 69cce08d352..d7c094a2361 100644 --- a/cpp/src/prims/per_v_random_select_transform_outgoing_e.cuh +++ b/cpp/src/prims/per_v_random_select_transform_outgoing_e.cuh @@ -399,11 +399,12 @@ rmm::device_uvector get_sampling_index_without_replacement( if (retry_segment_indices) { retry_degrees = rmm::device_uvector((*retry_segment_indices).size(), handle.get_stream()); - thrust::transform(handle.get_thrust_policy(), - (*retry_segment_indices).begin(), - (*retry_segment_indices).end(), - (*retry_degrees).begin(), - indirection_t{segment_degree_first}); + thrust::transform( + handle.get_thrust_policy(), + (*retry_segment_indices).begin(), + (*retry_segment_indices).end(), + (*retry_degrees).begin(), + indirection_t{segment_degree_first}); retry_sample_nbr_indices = rmm::device_uvector( (*retry_segment_indices).size() * high_partition_over_sampling_K, handle.get_stream()); retry_sample_indices = rmm::device_uvector( diff --git a/cpp/src/prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh b/cpp/src/prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh index a4d34443413..2e19adc34c4 100644 --- a/cpp/src/prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh +++ b/cpp/src/prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh @@ -756,7 +756,7 @@ void per_v_transform_reduce_dst_key_aggregated_outgoing_e( : thrust::nullopt; std::conditional_t, - detail::kv_cuco_store_device_view_t> + detail::kv_cuco_store_find_device_view_t> dst_key_value_map_device_view( GraphViewType::is_multi_gpu ? multi_gpu_minor_key_value_map_ptr->view() : kv_store_view); thrust::transform(handle.get_thrust_policy(), diff --git a/cpp/src/structure/relabel_impl.cuh b/cpp/src/structure/relabel_impl.cuh index c4c34733a4d..192120e6b4c 100644 --- a/cpp/src/structure/relabel_impl.cuh +++ b/cpp/src/structure/relabel_impl.cuh @@ -142,7 +142,7 @@ void relabel(raft::handle_t const& handle, handle.get_stream()); if (skip_missing_labels) { - auto device_view = detail::kv_cuco_store_device_view_t(relabel_map_view); + auto device_view = detail::kv_cuco_store_find_device_view_t(relabel_map_view); thrust::transform( handle.get_thrust_policy(), rx_unique_old_labels.begin(), @@ -187,7 +187,7 @@ void relabel(raft::handle_t const& handle, handle.get_stream()); auto relabel_map_view = relabel_map.view(); if (skip_missing_labels) { - auto device_view = detail::kv_cuco_store_device_view_t(relabel_map_view); + auto device_view = detail::kv_cuco_store_find_device_view_t(relabel_map_view); thrust::transform( handle.get_thrust_policy(), labels, From 917b98b07908f536ca6e3d17db48411901e7f3c2 Mon Sep 17 00:00:00 2001 From: Tingyu Wang Date: Tue, 18 Jul 2023 14:55:59 -0400 Subject: [PATCH 03/14] Make `pylibcugraphops` optional imports in `cugraph-dgl` and `-pyg` (#3693) Fixes #3691 . Authors: - Tingyu Wang (https://github.com/tingyu66) Approvers: - Vibhu Jawa (https://github.com/VibhuJawa) - Rick Ratzel (https://github.com/rlratzel) - Ray Douglass (https://github.com/raydouglass) URL: https://github.com/rapidsai/cugraph/pull/3693 --- .../all_cuda-118_arch-x86_64.yaml | 2 ++ dependencies.yaml | 21 ++++++++++++++++ .../conda/cugraph_dgl_dev_11.6.yml | 18 -------------- .../conda/cugraph_dgl_dev_cuda-118.yaml | 20 ++++++++++++++++ .../cugraph_dgl/nn/conv/gatconv.py | 8 +++---- .../cugraph_dgl/nn/conv/transformerconv.py | 8 +++---- .../cugraph-pyg/cugraph_pyg/nn/conv/base.py | 24 ++++--------------- .../cugraph_pyg/nn/conv/gat_conv.py | 5 ++-- .../cugraph_pyg/nn/conv/gatv2_conv.py | 5 ++-- .../cugraph_pyg/nn/conv/transformer_conv.py | 5 ++-- 10 files changed, 60 insertions(+), 56 deletions(-) delete mode 100644 python/cugraph-dgl/conda/cugraph_dgl_dev_11.6.yml create mode 100644 python/cugraph-dgl/conda/cugraph_dgl_dev_cuda-118.yaml diff --git a/conda/environments/all_cuda-118_arch-x86_64.yaml b/conda/environments/all_cuda-118_arch-x86_64.yaml index 0487fb53c07..9c428ef9d07 100644 --- a/conda/environments/all_cuda-118_arch-x86_64.yaml +++ b/conda/environments/all_cuda-118_arch-x86_64.yaml @@ -6,6 +6,8 @@ channels: - dask/label/dev - conda-forge - nvidia +- pytorch +- dglteam/label/cu118 dependencies: - aiohttp - c-compiler diff --git a/dependencies.yaml b/dependencies.yaml index 87edcc62683..9b858999743 100644 --- a/dependencies.yaml +++ b/dependencies.yaml @@ -170,12 +170,21 @@ files: # list is really minimal or if it is a superset. - test_python_common - test_python_cugraph + cugraph_dgl_dev: + matrix: + cuda: ["11.8"] + output: conda + conda_dir: python/cugraph-dgl/conda + includes: + - cugraph_dgl_dev channels: - rapidsai - rapidsai-nightly - dask/label/dev - conda-forge - nvidia + - pytorch + - dglteam/label/cu118 dependencies: checks: common: @@ -418,3 +427,15 @@ dependencies: - output_types: [conda, pyproject] packages: - *cudf + cugraph_dgl_dev: + common: + - output_types: [conda] + packages: + - cugraph==23.8.* + - pylibcugraphops==23.8.* + - pytorch>=2.0 + - pytorch-cuda==11.8 + - dgl>=1.1.0.cu* + - setuptools + - pre-commit + - pytest diff --git a/python/cugraph-dgl/conda/cugraph_dgl_dev_11.6.yml b/python/cugraph-dgl/conda/cugraph_dgl_dev_11.6.yml deleted file mode 100644 index 74b16921291..00000000000 --- a/python/cugraph-dgl/conda/cugraph_dgl_dev_11.6.yml +++ /dev/null @@ -1,18 +0,0 @@ -name: cugraph_dgl_dev -channels: -- rapidsai-nightly -- rapidsai -- pytorch -- conda-forge -- nvidia -- dglteam - -dependencies: -- cudatoolkit=11.6 -- cugraph=23.02* -- pylibcugraphops=23.02.* -- pytorch=1.12.0 -- dgl-cuda11.6 -- setuptools -- pre-commit -- pytest diff --git a/python/cugraph-dgl/conda/cugraph_dgl_dev_cuda-118.yaml b/python/cugraph-dgl/conda/cugraph_dgl_dev_cuda-118.yaml new file mode 100644 index 00000000000..a252d5e0c78 --- /dev/null +++ b/python/cugraph-dgl/conda/cugraph_dgl_dev_cuda-118.yaml @@ -0,0 +1,20 @@ +# This file is generated by `rapids-dependency-file-generator`. +# To make changes, edit ../../../dependencies.yaml and run `rapids-dependency-file-generator`. +channels: +- rapidsai +- rapidsai-nightly +- dask/label/dev +- conda-forge +- nvidia +- pytorch +- dglteam/label/cu118 +dependencies: +- cugraph==23.8.* +- dgl>=1.1.0.cu* +- pre-commit +- pylibcugraphops==23.8.* +- pytest +- pytorch-cuda==11.8 +- pytorch>=2.0 +- setuptools +name: cugraph_dgl_dev_cuda-118 diff --git a/python/cugraph-dgl/cugraph_dgl/nn/conv/gatconv.py b/python/cugraph-dgl/cugraph_dgl/nn/conv/gatconv.py index 7825febc24b..239def5b677 100644 --- a/python/cugraph-dgl/cugraph_dgl/nn/conv/gatconv.py +++ b/python/cugraph-dgl/cugraph_dgl/nn/conv/gatconv.py @@ -19,12 +19,10 @@ from cugraph_dgl.nn.conv.base import BaseConv from cugraph.utilities.utils import import_optional -from pylibcugraphops.pytorch import CSC -from pylibcugraphops.pytorch.operators import mha_gat_n2n - dgl = import_optional("dgl") torch = import_optional("torch") nn = import_optional("torch.nn") +ops_torch = import_optional("pylibcugraphops.pytorch") class GATConv(BaseConv): @@ -179,7 +177,7 @@ def forward( bipartite = not isinstance(nfeat, torch.Tensor) offsets, indices, _ = g.adj_tensors("csc") - graph = CSC( + graph = ops_torch.CSC( offsets=offsets, indices=indices, num_src_nodes=g.num_src_nodes(), @@ -212,7 +210,7 @@ def forward( ) nfeat = self.fc(nfeat) - out = mha_gat_n2n( + out = ops_torch.operators.mha_gat_n2n( (nfeat_src, nfeat_dst) if bipartite else nfeat, self.attn_weights, graph, diff --git a/python/cugraph-dgl/cugraph_dgl/nn/conv/transformerconv.py b/python/cugraph-dgl/cugraph_dgl/nn/conv/transformerconv.py index 141adc86069..5cd5fbbaebe 100644 --- a/python/cugraph-dgl/cugraph_dgl/nn/conv/transformerconv.py +++ b/python/cugraph-dgl/cugraph_dgl/nn/conv/transformerconv.py @@ -15,12 +15,10 @@ from cugraph_dgl.nn.conv.base import BaseConv from cugraph.utilities.utils import import_optional -from pylibcugraphops.pytorch import CSC -from pylibcugraphops.pytorch.operators import mha_simple_n2n - dgl = import_optional("dgl") torch = import_optional("torch") nn = import_optional("torch.nn") +ops_torch = import_optional("pylibcugraphops.pytorch") class TransformerConv(BaseConv): @@ -133,7 +131,7 @@ def forward( Edge feature tensor. Default: ``None``. """ offsets, indices, _ = g.adj_tensors("csc") - graph = CSC( + graph = ops_torch.CSC( offsets=offsets, indices=indices, num_src_nodes=g.num_src_nodes(), @@ -155,7 +153,7 @@ def forward( ) efeat = self.lin_edge(efeat) - out = mha_simple_n2n( + out = ops_torch.operators.mha_simple_n2n( key_emb=key, query_emb=query, value_emb=value, diff --git a/python/cugraph-pyg/cugraph_pyg/nn/conv/base.py b/python/cugraph-pyg/cugraph_pyg/nn/conv/base.py index 207efcdace4..2639f66f440 100644 --- a/python/cugraph-pyg/cugraph_pyg/nn/conv/base.py +++ b/python/cugraph-pyg/cugraph_pyg/nn/conv/base.py @@ -18,26 +18,12 @@ torch = import_optional("torch") torch_geometric = import_optional("torch_geometric") - -try: # pragma: no cover - from pylibcugraphops.pytorch import CSC, HeteroCSC - - HAS_PYLIBCUGRAPHOPS = True -except ImportError: - HAS_PYLIBCUGRAPHOPS = False +ops_torch = import_optional("pylibcugraphops.pytorch") class BaseConv(torch.nn.Module): # pragma: no cover r"""An abstract base class for implementing cugraph-ops message passing layers.""" - def __init__(self): - super().__init__() - - if HAS_PYLIBCUGRAPHOPS is False: - raise ModuleNotFoundError( - f"'{self.__class__.__name__}' requires " f"'pylibcugraphops>=23.04'" - ) - def reset_parameters(self): r"""Resets all learnable parameters of the module.""" pass @@ -88,7 +74,7 @@ def get_cugraph( csc: Tuple[torch.Tensor, torch.Tensor, int], bipartite: bool = False, max_num_neighbors: Optional[int] = None, - ) -> CSC: + ) -> ops_torch.CSC: r"""Constructs a :obj:`cugraph-ops` graph object from CSC representation. Supports both bipartite and non-bipartite graphs. @@ -116,7 +102,7 @@ def get_cugraph( if max_num_neighbors is None: max_num_neighbors = -1 - return CSC( + return ops_torch.CSC( offsets=colptr, indices=row, num_src_nodes=num_src_nodes, @@ -131,7 +117,7 @@ def get_typed_cugraph( num_edge_types: Optional[int] = None, bipartite: bool = False, max_num_neighbors: Optional[int] = None, - ) -> HeteroCSC: + ) -> ops_torch.HeteroCSC: r"""Constructs a typed :obj:`cugraph` graph object from a CSC representation where each edge corresponds to a given edge type. Supports both bipartite and non-bipartite graphs. @@ -162,7 +148,7 @@ def get_typed_cugraph( row, colptr, num_src_nodes = csc edge_type = edge_type.int() - return HeteroCSC( + return ops_torch.HeteroCSC( offsets=colptr, indices=row, edge_types=edge_type, diff --git a/python/cugraph-pyg/cugraph_pyg/nn/conv/gat_conv.py b/python/cugraph-pyg/cugraph_pyg/nn/conv/gat_conv.py index 23b7d50ba96..f0040015b4a 100644 --- a/python/cugraph-pyg/cugraph_pyg/nn/conv/gat_conv.py +++ b/python/cugraph-pyg/cugraph_pyg/nn/conv/gat_conv.py @@ -12,8 +12,6 @@ # limitations under the License. from typing import Optional, Tuple, Union -from pylibcugraphops.pytorch.operators import mha_gat_n2n - from cugraph.utilities.utils import import_optional from .base import BaseConv @@ -21,6 +19,7 @@ torch = import_optional("torch") nn = import_optional("torch.nn") torch_geometric = import_optional("torch_geometric") +ops_torch = import_optional("pylibcugraphops.pytorch") class GATConv(BaseConv): @@ -211,7 +210,7 @@ def forward( ) x = self.lin(x) - out = mha_gat_n2n( + out = ops_torch.operators.mha_gat_n2n( (x_src, x_dst) if bipartite else x, self.att, graph, diff --git a/python/cugraph-pyg/cugraph_pyg/nn/conv/gatv2_conv.py b/python/cugraph-pyg/cugraph_pyg/nn/conv/gatv2_conv.py index d4c947b952a..d74ca6b00d0 100644 --- a/python/cugraph-pyg/cugraph_pyg/nn/conv/gatv2_conv.py +++ b/python/cugraph-pyg/cugraph_pyg/nn/conv/gatv2_conv.py @@ -12,8 +12,6 @@ # limitations under the License. from typing import Optional, Tuple, Union -from pylibcugraphops.pytorch.operators import mha_gat_v2_n2n - from cugraph.utilities.utils import import_optional from .base import BaseConv @@ -21,6 +19,7 @@ torch = import_optional("torch") nn = import_optional("torch.nn") torch_geometric = import_optional("torch_geometric") +ops_torch = import_optional("pylibcugraphops.pytorch") class GATv2Conv(BaseConv): @@ -208,7 +207,7 @@ def forward( else: x = self.lin_src(x) - out = mha_gat_v2_n2n( + out = ops_torch.operators.mha_gat_v2_n2n( (x_src, x_dst) if bipartite else x, self.att, graph, diff --git a/python/cugraph-pyg/cugraph_pyg/nn/conv/transformer_conv.py b/python/cugraph-pyg/cugraph_pyg/nn/conv/transformer_conv.py index f67756eb3fe..1b8b1aa0ffa 100644 --- a/python/cugraph-pyg/cugraph_pyg/nn/conv/transformer_conv.py +++ b/python/cugraph-pyg/cugraph_pyg/nn/conv/transformer_conv.py @@ -12,8 +12,6 @@ # limitations under the License. from typing import Optional, Tuple, Union -from pylibcugraphops.pytorch.operators import mha_simple_n2n - from cugraph.utilities.utils import import_optional from .base import BaseConv @@ -21,6 +19,7 @@ torch = import_optional("torch") nn = import_optional("torch.nn") torch_geometric = import_optional("torch_geometric") +ops_torch = import_optional("pylibcugraphops.pytorch") class TransformerConv(BaseConv): @@ -186,7 +185,7 @@ def forward( ) edge_attr = self.lin_edge(edge_attr) - out = mha_simple_n2n( + out = ops_torch.operators.mha_simple_n2n( key, query, value, From 54c081267482d8e9c9db4d84cdf41d9ff23cc674 Mon Sep 17 00:00:00 2001 From: Alex Barghi <105237337+alexbarghi-nv@users.noreply.github.com> Date: Tue, 18 Jul 2023 15:51:19 -0400 Subject: [PATCH 04/14] [BUG] Fix Bulk Sampling Test Issue (#3701) The bulk sampling tests are currently using a temporary directory which is not accessible from all processes/workers. This is causing the MG tests to fail. This PR updates the bulk sampler tests to use an environment variable for the scratch directory. The test run script will also need to be updated to make sure that environment variable is properly set. Authors: - Alex Barghi (https://github.com/alexbarghi-nv) Approvers: - Brad Rees (https://github.com/BradReesWork) URL: https://github.com/rapidsai/cugraph/pull/3701 --- python/cugraph/cugraph/tests/conftest.py | 19 ++++++++ .../tests/sampling/test_bulk_sampler.py | 45 ++++++++++++------- .../tests/sampling/test_bulk_sampler_io.py | 19 +++++--- .../tests/sampling/test_bulk_sampler_io_mg.py | 19 +++++--- .../tests/sampling/test_bulk_sampler_mg.py | 28 ++++++++---- python/cugraph/cugraph/utilities/utils.py | 12 +++++ 6 files changed, 103 insertions(+), 39 deletions(-) diff --git a/python/cugraph/cugraph/tests/conftest.py b/python/cugraph/cugraph/tests/conftest.py index fece006c4b8..916e445cfdb 100644 --- a/python/cugraph/cugraph/tests/conftest.py +++ b/python/cugraph/cugraph/tests/conftest.py @@ -17,6 +17,9 @@ stop_dask_client, ) +import os +import tempfile + # module-wide fixtures @@ -42,3 +45,19 @@ def dask_client(): yield dask_client stop_dask_client(dask_client, dask_cluster) + + +@pytest.fixture(scope="module") +def scratch_dir(): + # This should always be set if doing MG testing, since temporary + # directories are only accessible from the current process. + tempdir_object = os.getenv( + "RAPIDS_PYTEST_SCRATCH_DIR", tempfile.TemporaryDirectory() + ) + + if isinstance(tempdir_object, tempfile.TemporaryDirectory): + yield tempdir_object.name + else: + yield tempdir_object + + del tempdir_object diff --git a/python/cugraph/cugraph/tests/sampling/test_bulk_sampler.py b/python/cugraph/cugraph/tests/sampling/test_bulk_sampler.py index c25b5297e18..99696f943f3 100644 --- a/python/cugraph/cugraph/tests/sampling/test_bulk_sampler.py +++ b/python/cugraph/cugraph/tests/sampling/test_bulk_sampler.py @@ -17,13 +17,14 @@ import cugraph from cugraph.experimental.datasets import karate from cugraph.experimental.gnn import BulkSampler +from cugraph.utilities.utils import create_directory_with_overwrite -import tempfile import os +import shutil @pytest.mark.sg -def test_bulk_sampler_simple(): +def test_bulk_sampler_simple(scratch_dir): el = karate.get_edgelist().reset_index().rename(columns={"index": "eid"}) el["eid"] = el["eid"].astype("int32") el["etp"] = cupy.int32(0) @@ -36,10 +37,12 @@ def test_bulk_sampler_simple(): edge_attr=["wgt", "eid", "etp"], ) - tempdir_object = tempfile.TemporaryDirectory() + samples_path = os.path.join(scratch_dir, "test_bulk_sampler_simple") + create_directory_with_overwrite(samples_path) + bs = BulkSampler( batch_size=2, - output_path=tempdir_object.name, + output_path=samples_path, graph=G, fanout_vals=[2, 2], with_replacement=False, @@ -55,14 +58,16 @@ def test_bulk_sampler_simple(): bs.add_batches(batches, start_col_name="start", batch_col_name="batch") bs.flush() - recovered_samples = cudf.read_parquet(tempdir_object.name) + recovered_samples = cudf.read_parquet(samples_path) for b in batches["batch"].unique().values_host.tolist(): assert b in recovered_samples["batch_id"].values_host.tolist() + shutil.rmtree(samples_path) + @pytest.mark.sg -def test_bulk_sampler_remainder(): +def test_bulk_sampler_remainder(scratch_dir): el = karate.get_edgelist().reset_index().rename(columns={"index": "eid"}) el["eid"] = el["eid"].astype("int32") el["etp"] = cupy.int32(0) @@ -75,10 +80,12 @@ def test_bulk_sampler_remainder(): edge_attr=["wgt", "eid", "etp"], ) - tempdir_object = tempfile.TemporaryDirectory() + samples_path = os.path.join(scratch_dir, "test_bulk_sampler_remainder") + create_directory_with_overwrite(samples_path) + bs = BulkSampler( batch_size=2, - output_path=tempdir_object.name, + output_path=samples_path, graph=G, seeds_per_call=7, batches_per_partition=2, @@ -102,26 +109,27 @@ def test_bulk_sampler_remainder(): bs.add_batches(batches, start_col_name="start", batch_col_name="batch") bs.flush() - tld = tempdir_object.name - recovered_samples = cudf.read_parquet(tld) + recovered_samples = cudf.read_parquet(samples_path) for b in batches["batch"].unique().values_host.tolist(): assert b in recovered_samples["batch_id"].values_host.tolist() for x in range(0, 6, 2): subdir = f"{x}-{x+1}" - df = cudf.read_parquet(os.path.join(tld, f"batch={subdir}.parquet")) + df = cudf.read_parquet(os.path.join(samples_path, f"batch={subdir}.parquet")) assert ((df.batch_id == x) | (df.batch_id == (x + 1))).all() assert ((df.hop_id == 0) | (df.hop_id == 1)).all() assert ( - cudf.read_parquet(os.path.join(tld, "batch=6-6.parquet")).batch_id == 6 + cudf.read_parquet(os.path.join(samples_path, "batch=6-6.parquet")).batch_id == 6 ).all() + shutil.rmtree(samples_path) + @pytest.mark.sg -def test_bulk_sampler_large_batch_size(): +def test_bulk_sampler_large_batch_size(scratch_dir): el = karate.get_edgelist().reset_index().rename(columns={"index": "eid"}) el["eid"] = el["eid"].astype("int32") el["etp"] = cupy.int32(0) @@ -134,10 +142,13 @@ def test_bulk_sampler_large_batch_size(): edge_attr=["wgt", "eid", "etp"], ) - tempdir_object = tempfile.TemporaryDirectory() + samples_path = os.path.join(scratch_dir, "test_bulk_sampler_large_batch_size") + if os.path.exists(samples_path): + shutil.rmtree(samples_path) + os.makedirs(samples_path) bs = BulkSampler( batch_size=5120, - output_path=tempdir_object.name, + output_path=samples_path, graph=G, fanout_vals=[2, 2], with_replacement=False, @@ -153,7 +164,9 @@ def test_bulk_sampler_large_batch_size(): bs.add_batches(batches, start_col_name="start", batch_col_name="batch") bs.flush() - recovered_samples = cudf.read_parquet(tempdir_object.name) + recovered_samples = cudf.read_parquet(samples_path) for b in batches["batch"].unique().values_host.tolist(): assert b in recovered_samples["batch_id"].values_host.tolist() + + shutil.rmtree(samples_path) diff --git a/python/cugraph/cugraph/tests/sampling/test_bulk_sampler_io.py b/python/cugraph/cugraph/tests/sampling/test_bulk_sampler_io.py index 83d20ea2cf5..ea39a9ee7bd 100644 --- a/python/cugraph/cugraph/tests/sampling/test_bulk_sampler_io.py +++ b/python/cugraph/cugraph/tests/sampling/test_bulk_sampler_io.py @@ -13,14 +13,15 @@ import pytest import cudf -import tempfile import os +import shutil from cugraph.gnn.data_loading.bulk_sampler_io import write_samples +from cugraph.utilities.utils import create_directory_with_overwrite @pytest.mark.sg -def test_bulk_sampler_io(): +def test_bulk_sampler_io(scratch_dir): results = cudf.DataFrame( { "sources": [0, 0, 1, 2, 2, 2, 3, 4, 5, 5, 6, 7], @@ -34,12 +35,14 @@ def test_bulk_sampler_io(): offsets = cudf.DataFrame({"offsets": [0, 8], "batch_id": [0, 1]}) - tempdir_object = tempfile.TemporaryDirectory() - write_samples(results, offsets, 1, tempdir_object.name) + samples_path = os.path.join(scratch_dir, "test_bulk_sampler_io") + create_directory_with_overwrite(samples_path) - assert len(os.listdir(tempdir_object.name)) == 2 + write_samples(results, offsets, 1, samples_path) - df = cudf.read_parquet(os.path.join(tempdir_object.name, "batch=0-0.parquet")) + assert len(os.listdir(samples_path)) == 2 + + df = cudf.read_parquet(os.path.join(samples_path, "batch=0-0.parquet")) assert len(df) == 8 assert ( @@ -55,7 +58,7 @@ def test_bulk_sampler_io(): ) assert (df.batch_id == 0).all() - df = cudf.read_parquet(os.path.join(tempdir_object.name, "batch=1-1.parquet")) + df = cudf.read_parquet(os.path.join(samples_path, "batch=1-1.parquet")) assert len(df) == 4 assert ( df.sources.values_host.tolist() @@ -69,3 +72,5 @@ def test_bulk_sampler_io(): df.hop_id.values_host.tolist() == results.hop_id.iloc[8:12].values_host.tolist() ) assert (df.batch_id == 1).all() + + shutil.rmtree(samples_path) diff --git a/python/cugraph/cugraph/tests/sampling/test_bulk_sampler_io_mg.py b/python/cugraph/cugraph/tests/sampling/test_bulk_sampler_io_mg.py index eacd697b7b3..4a77b873034 100644 --- a/python/cugraph/cugraph/tests/sampling/test_bulk_sampler_io_mg.py +++ b/python/cugraph/cugraph/tests/sampling/test_bulk_sampler_io_mg.py @@ -16,14 +16,15 @@ import cudf import dask_cudf -import tempfile import os +import shutil from cugraph.gnn.data_loading.bulk_sampler_io import write_samples +from cugraph.utilities.utils import create_directory_with_overwrite @pytest.mark.mg -def test_bulk_sampler_io(): +def test_bulk_sampler_io(scratch_dir): results = cudf.DataFrame( { "sources": [0, 0, 1, 2, 2, 2, 3, 4, 5, 5, 6, 7], @@ -41,12 +42,14 @@ def test_bulk_sampler_io(): offsets = cudf.DataFrame({"offsets": [0, 0], "batch_id": [0, 1]}) offsets = dask_cudf.from_cudf(offsets, npartitions=2) - tempdir_object = tempfile.TemporaryDirectory() - write_samples(results, offsets, 1, tempdir_object.name) + samples_path = os.path.join(scratch_dir, "mg_test_bulk_sampler_io") + create_directory_with_overwrite(samples_path) - assert len(os.listdir(tempdir_object.name)) == 2 + write_samples(results, offsets, 1, samples_path) - df = cudf.read_parquet(os.path.join(tempdir_object.name, "batch=0-0.parquet")) + assert len(os.listdir(samples_path)) == 2 + + df = cudf.read_parquet(os.path.join(samples_path, "batch=0-0.parquet")) assert len(df) == 8 results = results.compute() @@ -63,7 +66,7 @@ def test_bulk_sampler_io(): ) assert (df.batch_id == 0).all() - df = cudf.read_parquet(os.path.join(tempdir_object.name, "batch=1-1.parquet")) + df = cudf.read_parquet(os.path.join(samples_path, "batch=1-1.parquet")) assert len(df) == 4 assert ( df.sources.values_host.tolist() @@ -77,3 +80,5 @@ def test_bulk_sampler_io(): df.hop_id.values_host.tolist() == results.hop_id.iloc[8:12].values_host.tolist() ) assert (df.batch_id == 1).all() + + shutil.rmtree(samples_path) diff --git a/python/cugraph/cugraph/tests/sampling/test_bulk_sampler_mg.py b/python/cugraph/cugraph/tests/sampling/test_bulk_sampler_mg.py index f717d452403..7e94d3cdced 100644 --- a/python/cugraph/cugraph/tests/sampling/test_bulk_sampler_mg.py +++ b/python/cugraph/cugraph/tests/sampling/test_bulk_sampler_mg.py @@ -18,12 +18,14 @@ import cugraph from cugraph.experimental.datasets import karate from cugraph.experimental import BulkSampler +from cugraph.utilities.utils import create_directory_with_overwrite -import tempfile +import os +import shutil @pytest.mark.mg -def test_bulk_sampler_simple(dask_client): +def test_bulk_sampler_simple(dask_client, scratch_dir): el = karate.get_edgelist().reset_index().rename(columns={"index": "eid"}) el["eid"] = el["eid"].astype("int32") el["etp"] = cupy.int32(0) @@ -36,10 +38,12 @@ def test_bulk_sampler_simple(dask_client): edge_attr=["wgt", "eid", "etp"], ) - tempdir_object = tempfile.TemporaryDirectory() + samples_path = os.path.join(scratch_dir, "mg_test_bulk_sampler_simple") + create_directory_with_overwrite(samples_path) + bs = BulkSampler( batch_size=2, - output_path=tempdir_object.name, + output_path=samples_path, graph=G, fanout_vals=[2, 2], with_replacement=False, @@ -58,14 +62,16 @@ def test_bulk_sampler_simple(dask_client): bs.add_batches(batches, start_col_name="start", batch_col_name="batch") bs.flush() - recovered_samples = cudf.read_parquet(tempdir_object.name) + recovered_samples = cudf.read_parquet(samples_path) for b in batches["batch"].unique().compute().values_host.tolist(): assert b in recovered_samples["batch_id"].values_host.tolist() + shutil.rmtree(samples_path) + @pytest.mark.mg -def test_bulk_sampler_mg_graph_sg_input(dask_client): +def test_bulk_sampler_mg_graph_sg_input(dask_client, scratch_dir): el = karate.get_edgelist().reset_index().rename(columns={"index": "eid"}) el["eid"] = el["eid"].astype("int32") el["etp"] = cupy.int32(0) @@ -78,10 +84,12 @@ def test_bulk_sampler_mg_graph_sg_input(dask_client): edge_attr=["wgt", "eid", "etp"], ) - tempdir_object = tempfile.TemporaryDirectory() + samples_path = os.path.join(scratch_dir, "mg_test_bulk_sampler_mg_graph_sg_input") + create_directory_with_overwrite(samples_path) + bs = BulkSampler( batch_size=2, - output_path=tempdir_object.name, + output_path=samples_path, graph=G, fanout_vals=[2, 2], with_replacement=False, @@ -97,7 +105,9 @@ def test_bulk_sampler_mg_graph_sg_input(dask_client): bs.add_batches(batches, start_col_name="start", batch_col_name="batch") bs.flush() - recovered_samples = cudf.read_parquet(tempdir_object.name) + recovered_samples = cudf.read_parquet(samples_path) for b in batches["batch"].unique().values_host.tolist(): assert b in recovered_samples["batch_id"].values_host.tolist() + + shutil.rmtree(samples_path) diff --git a/python/cugraph/cugraph/utilities/utils.py b/python/cugraph/cugraph/utilities/utils.py index 5384175d201..e68b5dd4880 100644 --- a/python/cugraph/cugraph/utilities/utils.py +++ b/python/cugraph/cugraph/utilities/utils.py @@ -12,6 +12,8 @@ # limitations under the License. import importlib +import os +import shutil from numba import cuda @@ -529,3 +531,13 @@ def create_list_series_from_2d_ar(ar, index): children=(offset_col, data), ) return cudf.Series(lc, index=index) + + +def create_directory_with_overwrite(directory): + """ + Creates the given directory. If it already exists, the + existing directory is recursively deleted first. + """ + if os.path.exists(directory): + shutil.rmtree(directory) + os.makedirs(directory) From 1692e4c49f3c8fcad99a58e64187086f4062d32a Mon Sep 17 00:00:00 2001 From: Chuck Hastings <45364586+ChuckHastings@users.noreply.github.com> Date: Tue, 18 Jul 2023 18:49:16 -0400 Subject: [PATCH 05/14] Sampling modifications to support PyG and DGL options (#3696) Modify sampling to add options for managing inclusion of sources on different hops. Closes #3667 Closes #3665 Closes #3664 Authors: - Chuck Hastings (https://github.com/ChuckHastings) - Brad Rees (https://github.com/BradReesWork) Approvers: - Alex Barghi (https://github.com/alexbarghi-nv) - Naim (https://github.com/naimnv) - Seunghwa Kang (https://github.com/seunghwak) - Vibhu Jawa (https://github.com/VibhuJawa) URL: https://github.com/rapidsai/cugraph/pull/3696 --- cpp/CMakeLists.txt | 10 +- cpp/include/cugraph/algorithms.hpp | 22 +- cpp/include/cugraph_c/sampling_algorithms.h | 123 ++++ cpp/src/c_api/uniform_neighbor_sampling.cpp | 163 ++++- .../detail/gather_one_hop_edgelist_impl.cuh | 394 ++++++++++++ .../detail/gather_one_hop_edgelist_mg.cu | 119 ++++ .../detail/gather_one_hop_edgelist_sg.cu | 119 ++++ .../detail/prepare_next_frontier_impl.cuh | 197 ++++++ .../detail/prepare_next_frontier_mg.cu | 59 ++ .../detail/prepare_next_frontier_sg.cu | 59 ++ .../remove_visited_vertices_from_frontier.cu | 92 +++ cpp/src/sampling/detail/sample_edges.cuh | 277 ++++++++ cpp/src/sampling/detail/sample_edges_mg.cu | 125 ++++ ...ampling_utils_sg.cu => sample_edges_sg.cu} | 98 +-- ...graph_functions.hpp => sampling_utils.hpp} | 75 +++ cpp/src/sampling/detail/sampling_utils_mg.cu | 341 ---------- ...h => shuffle_and_organize_output_impl.cuh} | 591 +----------------- .../detail/shuffle_and_organize_output_mg.cu | 143 +++++ .../uniform_neighbor_sampling_impl.hpp | 112 ++-- .../sampling/uniform_neighbor_sampling_mg.cpp | 12 + .../sampling/uniform_neighbor_sampling_sg.cpp | 12 + cpp/tests/CMakeLists.txt | 48 +- cpp/tests/c_api/mg_test_utils.cpp | 274 ++++++++ cpp/tests/c_api/mg_test_utils.h | 34 + .../c_api/mg_uniform_neighbor_sample_test.c | 552 +++++++++++++--- .../c_api/uniform_neighbor_sample_test.c | 479 ++++++++++++++ 26 files changed, 3321 insertions(+), 1209 deletions(-) create mode 100644 cpp/src/sampling/detail/gather_one_hop_edgelist_impl.cuh create mode 100644 cpp/src/sampling/detail/gather_one_hop_edgelist_mg.cu create mode 100644 cpp/src/sampling/detail/gather_one_hop_edgelist_sg.cu create mode 100644 cpp/src/sampling/detail/prepare_next_frontier_impl.cuh create mode 100644 cpp/src/sampling/detail/prepare_next_frontier_mg.cu create mode 100644 cpp/src/sampling/detail/prepare_next_frontier_sg.cu create mode 100644 cpp/src/sampling/detail/remove_visited_vertices_from_frontier.cu create mode 100644 cpp/src/sampling/detail/sample_edges.cuh create mode 100644 cpp/src/sampling/detail/sample_edges_mg.cu rename cpp/src/sampling/detail/{sampling_utils_sg.cu => sample_edges_sg.cu} (56%) rename cpp/src/sampling/detail/{graph_functions.hpp => sampling_utils.hpp} (69%) delete mode 100644 cpp/src/sampling/detail/sampling_utils_mg.cu rename cpp/src/sampling/detail/{sampling_utils_impl.cuh => shuffle_and_organize_output_impl.cuh} (56%) create mode 100644 cpp/src/sampling/detail/shuffle_and_organize_output_mg.cu diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 996e654734f..8af50e5f72b 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -219,8 +219,14 @@ set(CUGRAPH_SOURCES src/community/egonet_mg.cu src/sampling/random_walks.cu src/sampling/random_walks_sg.cu - src/sampling/detail/sampling_utils_mg.cu - src/sampling/detail/sampling_utils_sg.cu + src/sampling/detail/prepare_next_frontier_sg.cu + src/sampling/detail/prepare_next_frontier_mg.cu + src/sampling/detail/gather_one_hop_edgelist_sg.cu + src/sampling/detail/gather_one_hop_edgelist_mg.cu + src/sampling/detail/remove_visited_vertices_from_frontier.cu + src/sampling/detail/sample_edges_sg.cu + src/sampling/detail/sample_edges_mg.cu + src/sampling/detail/shuffle_and_organize_output_mg.cu src/sampling/uniform_neighbor_sampling_mg.cpp src/sampling/uniform_neighbor_sampling_sg.cpp src/cores/core_number_sg.cu diff --git a/cpp/include/cugraph/algorithms.hpp b/cpp/include/cugraph/algorithms.hpp index cf9cba2af4d..547a298fba9 100644 --- a/cpp/include/cugraph/algorithms.hpp +++ b/cpp/include/cugraph/algorithms.hpp @@ -1840,6 +1840,17 @@ k_core(raft::handle_t const& handle, std::optional> core_numbers, bool do_expensive_check = false); +/** + * @brief Controls how we treat prior sources in sampling + * + * @param DEFAULT Add vertices encounted while sampling to the new frontier + * @param CARRY_OVER In addition to newly encountered vertices, include vertices + * used as sources in any previous frontier in the new frontier + * @param EXCLUDE Filter the new frontier to exclude any vertex that was + * used as a source in a previous frontier + */ +enum class prior_sources_behavior_t { DEFAULT = 0, CARRY_OVER, EXCLUDE }; + /** * @brief Uniform Neighborhood Sampling. * @@ -1893,6 +1904,11 @@ k_core(raft::handle_t const& handle, * level * @param rng_state A pre-initialized raft::RngState object for generating random numbers * @param return_hops boolean flag specifying if the hop information should be returned + * @param prior_sources_behavior Enum type defining how to handle prior sources, (defaults to + * DEFAULT) + * @param dedupe_sources boolean flag, if true then if a vertex v appears as a destination in hop X + * multiple times with the same label, it will only be passed once (for each label) as a source + * for the next hop. Default is false. * @param with_replacement boolean flag specifying if random sampling is done with replacement * (true); or, without replacement (false); default = true; * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`). @@ -1928,8 +1944,10 @@ uniform_neighbor_sample( raft::host_span fan_out, raft::random::RngState& rng_state, bool return_hops, - bool with_replacement = true, - bool do_expensive_check = false); + bool with_replacement = true, + prior_sources_behavior_t prior_sources_behavior = prior_sources_behavior_t::DEFAULT, + bool dedupe_sources = false, + bool do_expensive_check = false); /* * @brief Compute triangle counts. diff --git a/cpp/include/cugraph_c/sampling_algorithms.h b/cpp/include/cugraph_c/sampling_algorithms.h index e1be616016b..5e792403a88 100644 --- a/cpp/include/cugraph_c/sampling_algorithms.h +++ b/cpp/include/cugraph_c/sampling_algorithms.h @@ -187,8 +187,79 @@ typedef struct { int32_t align_; } cugraph_sample_result_t; +/** + * @brief Opaque sampling options type + */ +typedef struct { + int32_t align_; +} cugraph_sampling_options_t; + +/** + * @brief Enumeration for prior sources behavior + */ +typedef enum cugraph_prior_sources_behavior_t { + DEFAULT = 0, /** Construct sources for hop k from destination vertices from hop k-1 */ + CARRY_OVER, /** Construct sources for hop k from destination vertices from hop k-1 + and sources from hop k-1 */ + EXCLUDE /** Construct sources for hop k from destination vertices form hop k-1, + but exclude any vertex that has already been used as a source */ +} cugraph_prior_sources_behavior_t; + +/** + * @brief Create sampling options object + * + * All sampling options set to FALSE + * + * @param [out] options Opaque pointer to the sampling options + * @param [out] error Pointer to an error object storing details of any error. Will + * be populated if error code is not CUGRAPH_SUCCESS + */ +cugraph_error_code_t cugraph_sampling_options_create(cugraph_sampling_options_t** options, + cugraph_error_t** error); + +/** + * @brief Set flag to sample with_replacement + * + * @param options - opaque pointer to the sampling options + * @param value - Boolean value to assign to the option + */ +void cugraph_sampling_set_with_replacement(cugraph_sampling_options_t* options, bool_t value); + +/** + * @brief Set flag to sample return_hops + * + * @param options - opaque pointer to the sampling options + * @param value - Boolean value to assign to the option + */ +void cugraph_sampling_set_return_hops(cugraph_sampling_options_t* options, bool_t value); + +/** + * @brief Set prior sources behavior + * + * @param options - opaque pointer to the sampling options + * @param value - Enum defining prior sources behavior + */ +void cugraph_sampling_set_prior_sources_behavior(cugraph_sampling_options_t* options, + cugraph_prior_sources_behavior_t value); + +/** + * @brief Set flag to sample dedupe_sources prior to sampling + * + * @param options - opaque pointer to the sampling options + * @param value - Boolean value to assign to the option + */ +void cugraph_sampling_set_dedupe_sources(cugraph_sampling_options_t* options, bool_t value); + +/** + * @brief Free sampling options object + * + * @param [in] options Opaque pointer to sampling object + */ +void cugraph_sampling_options_free(cugraph_sampling_options_t* options); + /** * @brief Uniform Neighborhood Sampling + * @deprecated This call should be replaced with cugraph_uniform_neighbor_sample * * Returns a sample of the neighborhood around specified start vertices. Optionally, each * start vertex can be associated with a label, allowing the caller to specify multiple batches @@ -242,6 +313,58 @@ cugraph_error_code_t cugraph_uniform_neighbor_sample_with_edge_properties( cugraph_sample_result_t** result, cugraph_error_t** error); +/** + * @brief Uniform Neighborhood Sampling + * + * Returns a sample of the neighborhood around specified start vertices. Optionally, each + * start vertex can be associated with a label, allowing the caller to specify multiple batches + * of sampling requests in the same function call - which should improve GPU utilization. + * + * If label is NULL then all start vertices will be considered part of the same batch and the + * return value will not have a label column. + * + * @param [in] handle Handle for accessing resources + * @param [in] graph Pointer to graph. NOTE: Graph might be modified if the storage + * needs to be transposed + * @param [in] start_vertices Device array of start vertices for the sampling + * @param [in] start_vertex_labels Device array of start vertex labels for the sampling. The + * labels associated with each start vertex will be included in the output associated with results + * that were derived from that start vertex. We only support label of type INT32. If label is + * NULL, the return data will not be labeled. + * @param [in] label_list Device array of the labels included in @p start_vertex_labels. If + * @p label_to_comm_rank is not specified this parameter is ignored. If specified, label_list + * must be sorted in ascending order. + * @param [in] label_to_comm_rank Device array identifying which comm rank the output for a + * particular label should be shuffled in the output. If not specifed the data is not organized in + * output. If specified then the all data from @p label_list[i] will be shuffled to rank @p. This + * cannot be specified unless @p start_vertex_labels is also specified + * label_to_comm_rank[i]. If not specified then the output data will not be shuffled between ranks. + * @param [in] fanout Host array defining the fan out at each step in the sampling algorithm. + * We only support fanout values of type INT32 + * @param [in/out] rng_state State of the random number generator, updated with each call + * @param [in] sampling_options + * Opaque pointer defining the sampling options. + * @param [in] do_expensive_check + * A flag to run expensive checks for input arguments (if set to true) + * @param [in] result Output from the uniform_neighbor_sample call + * @param [out] error Pointer to an error object storing details of any error. Will + * be populated if error code is not CUGRAPH_SUCCESS + * @return error code + */ +cugraph_error_code_t cugraph_uniform_neighbor_sample( + const cugraph_resource_handle_t* handle, + cugraph_graph_t* graph, + const cugraph_type_erased_device_array_view_t* start_vertices, + const cugraph_type_erased_device_array_view_t* start_vertex_labels, + const cugraph_type_erased_device_array_view_t* label_list, + const cugraph_type_erased_device_array_view_t* label_to_comm_rank, + const cugraph_type_erased_host_array_view_t* fan_out, + cugraph_rng_state_t* rng_state, + const cugraph_sampling_options_t* options, + bool_t do_expensive_check, + cugraph_sample_result_t** result, + cugraph_error_t** error); + /** * @brief Get the source vertices from the sampling algorithm result * diff --git a/cpp/src/c_api/uniform_neighbor_sampling.cpp b/cpp/src/c_api/uniform_neighbor_sampling.cpp index 3b3a3ad6381..d9dc9ca4d50 100644 --- a/cpp/src/c_api/uniform_neighbor_sampling.cpp +++ b/cpp/src/c_api/uniform_neighbor_sampling.cpp @@ -32,6 +32,13 @@ namespace cugraph { namespace c_api { +struct cugraph_sampling_options_t { + bool_t with_replacement_{FALSE}; + bool_t return_hops_{FALSE}; + prior_sources_behavior_t prior_sources_behavior_{prior_sources_behavior_t::DEFAULT}; + bool_t dedupe_sources_{FALSE}; +}; + struct cugraph_sample_result_t { cugraph_type_erased_device_array_t* src_{nullptr}; cugraph_type_erased_device_array_t* dst_{nullptr}; @@ -57,8 +64,7 @@ struct uniform_neighbor_sampling_functor : public cugraph::c_api::abstract_funct cugraph::c_api::cugraph_type_erased_device_array_view_t const* label_to_comm_rank_{nullptr}; cugraph::c_api::cugraph_type_erased_host_array_view_t const* fan_out_{nullptr}; cugraph::c_api::cugraph_rng_state_t* rng_state_{nullptr}; - bool with_replacement_{false}; - bool return_hops_{false}; + cugraph::c_api::cugraph_sampling_options_t options_{}; bool do_expensive_check_{false}; cugraph::c_api::cugraph_sample_result_t* result_{nullptr}; @@ -71,8 +77,7 @@ struct uniform_neighbor_sampling_functor : public cugraph::c_api::abstract_funct cugraph_type_erased_device_array_view_t const* label_to_comm_rank, cugraph_type_erased_host_array_view_t const* fan_out, cugraph_rng_state_t* rng_state, - bool with_replacement, - bool return_hops, + cugraph::c_api::cugraph_sampling_options_t options, bool do_expensive_check) : abstract_functor(), handle_(*reinterpret_cast(handle)->handle_), @@ -91,8 +96,7 @@ struct uniform_neighbor_sampling_functor : public cugraph::c_api::abstract_funct fan_out_( reinterpret_cast(fan_out)), rng_state_(reinterpret_cast(rng_state)), - with_replacement_(with_replacement), - return_hops_(return_hops), + options_(options), do_expensive_check_(do_expensive_check) { } @@ -200,8 +204,11 @@ struct uniform_neighbor_sampling_functor : public cugraph::c_api::abstract_funct : std::nullopt, raft::host_span(fan_out_->as_type(), fan_out_->size_), rng_state_->rng_state_, - return_hops_, - with_replacement_); + options_.return_hops_, + options_.with_replacement_, + options_.prior_sources_behavior_, + options_.dedupe_sources_, + do_expensive_check_); std::vector vertex_partition_lasts = graph_view.vertex_partition_range_lasts(); @@ -242,6 +249,63 @@ struct uniform_neighbor_sampling_functor : public cugraph::c_api::abstract_funct } // namespace +extern "C" cugraph_error_code_t cugraph_sampling_options_create( + cugraph_sampling_options_t** options, cugraph_error_t** error) +{ + *options = + reinterpret_cast(new cugraph::c_api::cugraph_sampling_options_t()); + if (*options == nullptr) { + *error = reinterpret_cast( + new cugraph::c_api::cugraph_error_t{"invalid resource handle"}); + return CUGRAPH_INVALID_HANDLE; + } + + return CUGRAPH_SUCCESS; +} + +extern "C" void cugraph_sampling_set_with_replacement(cugraph_sampling_options_t* options, + bool_t value) +{ + auto internal_pointer = reinterpret_cast(options); + internal_pointer->with_replacement_ = value; +} + +extern "C" void cugraph_sampling_set_return_hops(cugraph_sampling_options_t* options, bool_t value) +{ + auto internal_pointer = reinterpret_cast(options); + internal_pointer->return_hops_ = value; +} + +extern "C" void cugraph_sampling_set_prior_sources_behavior(cugraph_sampling_options_t* options, + cugraph_prior_sources_behavior_t value) +{ + auto internal_pointer = reinterpret_cast(options); + switch (value) { + case CARRY_OVER: + internal_pointer->prior_sources_behavior_ = cugraph::prior_sources_behavior_t::CARRY_OVER; + break; + case EXCLUDE: + internal_pointer->prior_sources_behavior_ = cugraph::prior_sources_behavior_t::EXCLUDE; + break; + default: + internal_pointer->prior_sources_behavior_ = cugraph::prior_sources_behavior_t::DEFAULT; + break; + } +} + +extern "C" void cugraph_sampling_set_dedupe_sources(cugraph_sampling_options_t* options, + bool_t value) +{ + auto internal_pointer = reinterpret_cast(options); + internal_pointer->dedupe_sources_ = value; +} + +extern "C" void cugraph_sampling_options_free(cugraph_sampling_options_t* options) +{ + auto internal_pointer = reinterpret_cast(options); + delete internal_pointer; +} + extern "C" cugraph_type_erased_device_array_view_t* cugraph_sample_result_get_sources( const cugraph_sample_result_t* result) { @@ -617,16 +681,77 @@ extern "C" cugraph_error_code_t cugraph_uniform_neighbor_sample_with_edge_proper "fan_out should be of type int", *error); - uniform_neighbor_sampling_functor functor{handle, - graph, - start_vertices, - start_vertex_labels, - label_list, - label_to_comm_rank, - fan_out, - rng_state, - with_replacement, - return_hops, - do_expensive_check}; + uniform_neighbor_sampling_functor functor{ + handle, + graph, + start_vertices, + start_vertex_labels, + label_list, + label_to_comm_rank, + fan_out, + rng_state, + cugraph::c_api::cugraph_sampling_options_t{with_replacement, return_hops}, + do_expensive_check}; + return cugraph::c_api::run_algorithm(graph, functor, result, error); +} + +cugraph_error_code_t cugraph_uniform_neighbor_sample( + const cugraph_resource_handle_t* handle, + cugraph_graph_t* graph, + const cugraph_type_erased_device_array_view_t* start_vertices, + const cugraph_type_erased_device_array_view_t* start_vertex_labels, + const cugraph_type_erased_device_array_view_t* label_list, + const cugraph_type_erased_device_array_view_t* label_to_comm_rank, + const cugraph_type_erased_host_array_view_t* fan_out, + cugraph_rng_state_t* rng_state, + const cugraph_sampling_options_t* options, + bool_t do_expensive_check, + cugraph_sample_result_t** result, + cugraph_error_t** error) +{ + CAPI_EXPECTS((start_vertex_labels == nullptr) || + (reinterpret_cast( + start_vertex_labels) + ->type_ == INT32), + CUGRAPH_INVALID_INPUT, + "start_vertex_labels should be of type int", + *error); + + CAPI_EXPECTS((label_to_comm_rank == nullptr) || (start_vertex_labels != nullptr), + CUGRAPH_INVALID_INPUT, + "cannot specify label_to_comm_rank unless start_vertex_labels is also specified", + *error); + + CAPI_EXPECTS((label_to_comm_rank == nullptr) || (label_list != nullptr), + CUGRAPH_INVALID_INPUT, + "cannot specify label_to_comm_rank unless label_list is also specified", + *error); + + CAPI_EXPECTS(reinterpret_cast(graph)->vertex_type_ == + reinterpret_cast( + start_vertices) + ->type_, + CUGRAPH_INVALID_INPUT, + "vertex type of graph and start_vertices must match", + *error); + + CAPI_EXPECTS( + reinterpret_cast(fan_out) + ->type_ == INT32, + CUGRAPH_INVALID_INPUT, + "fan_out should be of type int", + *error); + + uniform_neighbor_sampling_functor functor{ + handle, + graph, + start_vertices, + start_vertex_labels, + label_list, + label_to_comm_rank, + fan_out, + rng_state, + *reinterpret_cast(options), + do_expensive_check}; return cugraph::c_api::run_algorithm(graph, functor, result, error); } diff --git a/cpp/src/sampling/detail/gather_one_hop_edgelist_impl.cuh b/cpp/src/sampling/detail/gather_one_hop_edgelist_impl.cuh new file mode 100644 index 00000000000..74267d02b38 --- /dev/null +++ b/cpp/src/sampling/detail/gather_one_hop_edgelist_impl.cuh @@ -0,0 +1,394 @@ +/* + * Copyright (c) 2022-2023, 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. + */ + +#pragma once + +#include +#include +#include +#include + +#include +#include +#include +#include + +#include + +#include +#include + +#include + +namespace cugraph { +namespace detail { + +struct return_edges_with_properties_e_op { + template + auto __host__ __device__ operator()(key_t optionally_tagged_src, + vertex_t dst, + thrust::nullopt_t, + thrust::nullopt_t, + EdgeProperties edge_properties) + { + static_assert(std::is_same_v || + std::is_same_v>); + + // FIXME: A solution using thrust_tuple_cat would be more flexible here + if constexpr (std::is_same_v) { + vertex_t src{optionally_tagged_src}; + + if constexpr (std::is_same_v) { + return thrust::make_optional(thrust::make_tuple(src, dst)); + } else if constexpr (std::is_arithmetic::value) { + return thrust::make_optional(thrust::make_tuple(src, dst, edge_properties)); + } else if constexpr (cugraph::is_thrust_tuple_of_arithmetic::value && + (thrust::tuple_size::value == 2)) { + return thrust::make_optional(thrust::make_tuple( + src, dst, thrust::get<0>(edge_properties), thrust::get<1>(edge_properties))); + } else if constexpr (cugraph::is_thrust_tuple_of_arithmetic::value && + (thrust::tuple_size::value == 3)) { + return thrust::make_optional(thrust::make_tuple(src, + dst, + thrust::get<0>(edge_properties), + thrust::get<1>(edge_properties), + thrust::get<2>(edge_properties))); + } + } else if constexpr (std::is_same_v>) { + vertex_t src{thrust::get<0>(optionally_tagged_src)}; + int32_t label{thrust::get<1>(optionally_tagged_src)}; + + src = thrust::get<0>(optionally_tagged_src); + if constexpr (std::is_same_v) { + return thrust::make_optional(thrust::make_tuple(src, dst, label)); + } else if constexpr (std::is_arithmetic::value) { + return thrust::make_optional(thrust::make_tuple(src, dst, edge_properties, label)); + } else if constexpr (cugraph::is_thrust_tuple_of_arithmetic::value && + (thrust::tuple_size::value == 2)) { + return thrust::make_optional(thrust::make_tuple( + src, dst, thrust::get<0>(edge_properties), thrust::get<1>(edge_properties), label)); + } else if constexpr (cugraph::is_thrust_tuple_of_arithmetic::value && + (thrust::tuple_size::value == 3)) { + return thrust::make_optional(thrust::make_tuple(src, + dst, + thrust::get<0>(edge_properties), + thrust::get<1>(edge_properties), + thrust::get<2>(edge_properties), + label)); + } + } + } +}; + +template +std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +gather_one_hop_edgelist( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_type_view, + cugraph::vertex_frontier_t const& vertex_frontier, + bool do_expensive_check) +{ + rmm::device_uvector majors(0, handle.get_stream()); + rmm::device_uvector minors(0, handle.get_stream()); + std::optional> edge_ids{std::nullopt}; + std::optional> edge_weights{std::nullopt}; + std::optional> edge_types{std::nullopt}; + std::optional> labels{std::nullopt}; + + if (edge_weight_view) { + if (edge_id_view) { + if (edge_type_view) { + if constexpr (std::is_same_v) { + std::tie(majors, minors, edge_weights, edge_ids, edge_types, labels) = + cugraph::extract_transform_v_frontier_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + view_concat(*edge_weight_view, *edge_id_view, *edge_type_view), + return_edges_with_properties_e_op{}, + do_expensive_check); + } else { + std::tie(majors, minors, edge_weights, edge_ids, edge_types) = + cugraph::extract_transform_v_frontier_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + view_concat(*edge_weight_view, *edge_id_view, *edge_type_view), + return_edges_with_properties_e_op{}, + do_expensive_check); + } + } else { + if constexpr (std::is_same_v) { + std::tie(majors, minors, edge_weights, edge_ids, labels) = + cugraph::extract_transform_v_frontier_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + view_concat(*edge_weight_view, *edge_id_view), + return_edges_with_properties_e_op{}, + do_expensive_check); + } else { + std::tie(majors, minors, edge_weights, edge_ids) = + cugraph::extract_transform_v_frontier_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + view_concat(*edge_weight_view, *edge_id_view), + return_edges_with_properties_e_op{}, + do_expensive_check); + } + } + } else { + if (edge_type_view) { + if constexpr (std::is_same_v) { + std::tie(majors, minors, edge_weights, edge_types, labels) = + cugraph::extract_transform_v_frontier_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + view_concat(*edge_weight_view, *edge_type_view), + return_edges_with_properties_e_op{}, + do_expensive_check); + } else { + std::tie(majors, minors, edge_weights, edge_types) = + cugraph::extract_transform_v_frontier_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + view_concat(*edge_weight_view, *edge_type_view), + return_edges_with_properties_e_op{}, + do_expensive_check); + } + } else { + if constexpr (std::is_same_v) { + std::tie(majors, minors, edge_weights, labels) = + cugraph::extract_transform_v_frontier_outgoing_e(handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + *edge_weight_view, + return_edges_with_properties_e_op{}, + do_expensive_check); + } else { + std::tie(majors, minors, edge_weights) = + cugraph::extract_transform_v_frontier_outgoing_e(handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + *edge_weight_view, + return_edges_with_properties_e_op{}, + do_expensive_check); + } + } + } + } else { + if (edge_id_view) { + if (edge_type_view) { + if constexpr (std::is_same_v) { + std::tie(majors, minors, edge_ids, edge_types, labels) = + cugraph::extract_transform_v_frontier_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + view_concat(*edge_id_view, *edge_type_view), + return_edges_with_properties_e_op{}, + do_expensive_check); + } else { + std::tie(majors, minors, edge_ids, edge_types) = + cugraph::extract_transform_v_frontier_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + view_concat(*edge_id_view, *edge_type_view), + return_edges_with_properties_e_op{}, + do_expensive_check); + } + } else { + if constexpr (std::is_same_v) { + std::tie(majors, minors, edge_ids, labels) = + cugraph::extract_transform_v_frontier_outgoing_e(handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + *edge_id_view, + return_edges_with_properties_e_op{}, + do_expensive_check); + } else { + std::tie(majors, minors, edge_ids) = + cugraph::extract_transform_v_frontier_outgoing_e(handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + *edge_id_view, + return_edges_with_properties_e_op{}, + do_expensive_check); + } + } + } else { + if (edge_type_view) { + if constexpr (std::is_same_v) { + std::tie(majors, minors, edge_types, labels) = + cugraph::extract_transform_v_frontier_outgoing_e(handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + *edge_type_view, + return_edges_with_properties_e_op{}, + do_expensive_check); + } else { + std::tie(majors, minors, edge_types) = + cugraph::extract_transform_v_frontier_outgoing_e(handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + *edge_type_view, + return_edges_with_properties_e_op{}, + do_expensive_check); + } + } else { + if constexpr (std::is_same_v) { + std::tie(majors, minors, labels) = + cugraph::extract_transform_v_frontier_outgoing_e(handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + edge_dummy_property_t{}.view(), + return_edges_with_properties_e_op{}, + do_expensive_check); + } else { + std::tie(majors, minors) = + cugraph::extract_transform_v_frontier_outgoing_e(handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + edge_dummy_property_t{}.view(), + return_edges_with_properties_e_op{}, + do_expensive_check); + } + } + } + } + + return std::make_tuple(std::move(majors), + std::move(minors), + std::move(edge_weights), + std::move(edge_ids), + std::move(edge_types), + std::move(labels)); +} + +template +std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +gather_one_hop_edgelist( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_type_view, + raft::device_span active_majors, + std::optional> active_major_labels, + bool do_expensive_check) +{ + if (active_major_labels) { + cugraph::vertex_frontier_t vertex_label_frontier(handle, + 1); + vertex_label_frontier.bucket(0).insert( + thrust::make_zip_iterator(active_majors.begin(), active_major_labels->begin()), + thrust::make_zip_iterator(active_majors.end(), active_major_labels->end())); + + return gather_one_hop_edgelist(handle, + graph_view, + edge_weight_view, + edge_id_view, + edge_type_view, + vertex_label_frontier, + do_expensive_check); + } else { + cugraph::vertex_frontier_t vertex_frontier(handle, 1); + vertex_frontier.bucket(0).insert(active_majors.begin(), active_majors.end()); + + return gather_one_hop_edgelist(handle, + graph_view, + edge_weight_view, + edge_id_view, + edge_type_view, + vertex_frontier, + do_expensive_check); + } +} + +} // namespace detail +} // namespace cugraph diff --git a/cpp/src/sampling/detail/gather_one_hop_edgelist_mg.cu b/cpp/src/sampling/detail/gather_one_hop_edgelist_mg.cu new file mode 100644 index 00000000000..ab9f3d00fae --- /dev/null +++ b/cpp/src/sampling/detail/gather_one_hop_edgelist_mg.cu @@ -0,0 +1,119 @@ +/* + * Copyright (c) 2022-2023, 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 + +namespace cugraph { +namespace detail { + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +gather_one_hop_edgelist( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_edge_type_view, + raft::device_span active_majors, + std::optional> active_major_labels, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +gather_one_hop_edgelist( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_edge_type_view, + raft::device_span active_majors, + std::optional> active_major_labels, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +gather_one_hop_edgelist( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_edge_type_view, + raft::device_span active_majors, + std::optional> active_major_labels, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +gather_one_hop_edgelist( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_edge_type_view, + raft::device_span active_majors, + std::optional> active_major_labels, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +gather_one_hop_edgelist( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_edge_type_view, + raft::device_span active_majors, + std::optional> active_major_labels, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +gather_one_hop_edgelist( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_edge_type_view, + raft::device_span active_majors, + std::optional> active_major_labels, + bool do_expensive_check); + +} // namespace detail +} // namespace cugraph diff --git a/cpp/src/sampling/detail/gather_one_hop_edgelist_sg.cu b/cpp/src/sampling/detail/gather_one_hop_edgelist_sg.cu new file mode 100644 index 00000000000..9f2629e8631 --- /dev/null +++ b/cpp/src/sampling/detail/gather_one_hop_edgelist_sg.cu @@ -0,0 +1,119 @@ +/* + * Copyright (c) 2022-2023, 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 + +namespace cugraph { +namespace detail { + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +gather_one_hop_edgelist( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_edge_type_view, + raft::device_span active_majors, + std::optional> active_major_labels, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +gather_one_hop_edgelist( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_edge_type_view, + raft::device_span active_majors, + std::optional> active_major_labels, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +gather_one_hop_edgelist( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_edge_type_view, + raft::device_span active_majors, + std::optional> active_major_labels, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +gather_one_hop_edgelist( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_edge_type_view, + raft::device_span active_majors, + std::optional> active_major_labels, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +gather_one_hop_edgelist( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_edge_type_view, + raft::device_span active_majors, + std::optional> active_major_labels, + bool do_expensive_check); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +gather_one_hop_edgelist( + raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_edge_type_view, + raft::device_span active_majors, + std::optional> active_major_labels, + bool do_expensive_check); + +} // namespace detail +} // namespace cugraph diff --git a/cpp/src/sampling/detail/prepare_next_frontier_impl.cuh b/cpp/src/sampling/detail/prepare_next_frontier_impl.cuh new file mode 100644 index 00000000000..a0e09225a5b --- /dev/null +++ b/cpp/src/sampling/detail/prepare_next_frontier_impl.cuh @@ -0,0 +1,197 @@ +/* + * Copyright (c) 2022-2023, 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. + */ + +#pragma once + +#include +#include +#include +#include +#include +#include + +#include + +#include + +#include +#include +#include +#include + +#include + +namespace cugraph { +namespace detail { + +template +std::tuple, + std::optional>, + std::optional, + std::optional>>>> +prepare_next_frontier( + raft::handle_t const& handle, + raft::device_span sampled_src_vertices, + std::optional> sampled_src_vertex_labels, + raft::device_span sampled_dst_vertices, + std::optional> sampled_dst_vertex_labels, + std::optional, + std::optional>>>&& vertex_used_as_source, + // FIXME: vertex_partition_view_t should provide vertex_partition_range_lasts() + // (and internally store this information in raft::host_span). + vertex_partition_view_t vertex_partition, + std::vector const& vertex_partition_range_lasts, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, + bool do_expensive_check) +{ + vertex_partition_device_view_t d_vertex_partition(vertex_partition); + + size_t frontier_size = sampled_dst_vertices.size(); + if (prior_sources_behavior == prior_sources_behavior_t::CARRY_OVER) { + frontier_size += sampled_src_vertices.size(); + } + + rmm::device_uvector frontier_vertices(frontier_size, handle.get_stream()); + auto frontier_vertex_labels = + sampled_dst_vertex_labels + ? std::make_optional>(frontier_size, handle.get_stream()) + : std::nullopt; + + thrust::copy(handle.get_thrust_policy(), + sampled_dst_vertices.begin(), + sampled_dst_vertices.end(), + frontier_vertices.begin()); + + if (prior_sources_behavior == prior_sources_behavior_t::CARRY_OVER) { + thrust::copy(handle.get_thrust_policy(), + sampled_src_vertices.begin(), + sampled_src_vertices.end(), + frontier_vertices.begin() + sampled_dst_vertices.size()); + } + + if (frontier_vertex_labels) { + thrust::copy(handle.get_thrust_policy(), + sampled_dst_vertex_labels->begin(), + sampled_dst_vertex_labels->end(), + frontier_vertex_labels->begin()); + + if (prior_sources_behavior == prior_sources_behavior_t::CARRY_OVER) { + thrust::copy(handle.get_thrust_policy(), + sampled_src_vertex_labels->begin(), + sampled_src_vertex_labels->end(), + frontier_vertex_labels->begin() + sampled_dst_vertices.size()); + } + } + + if constexpr (multi_gpu) { + if (frontier_vertex_labels) { + std::tie(frontier_vertices, *frontier_vertex_labels) = + shuffle_int_vertex_value_pairs_to_local_gpu_by_vertex_partitioning( + handle, + std::move(frontier_vertices), + std::move(*frontier_vertex_labels), + vertex_partition_range_lasts); + } else { + frontier_vertices = shuffle_int_vertices_to_local_gpu_by_vertex_partitioning( + handle, std::move(frontier_vertices), vertex_partition_range_lasts); + } + } + + if (frontier_vertex_labels) { + auto begin_iter = + thrust::make_zip_iterator(frontier_vertices.begin(), frontier_vertex_labels->begin()); + thrust::sort(handle.get_thrust_policy(), begin_iter, begin_iter + frontier_vertices.size()); + } else { + thrust::sort(handle.get_thrust_policy(), frontier_vertices.begin(), frontier_vertices.end()); + } + + if (vertex_used_as_source) { + auto& [verts, labels] = *vertex_used_as_source; + + // add sources from this expansion to the vertex_used_as_source + size_t current_verts_size = verts.size(); + size_t new_verts_size = current_verts_size + sampled_src_vertices.size(); + + verts.resize(new_verts_size, handle.get_stream()); + + thrust::copy(handle.get_thrust_policy(), + sampled_src_vertices.begin(), + sampled_src_vertices.end(), + verts.begin() + current_verts_size); + + // sort and unique the vertex_used_as_source structures + if (sampled_src_vertex_labels) { + labels->resize(new_verts_size, handle.get_stream()); + + thrust::copy(handle.get_thrust_policy(), + sampled_src_vertex_labels->begin(), + sampled_src_vertex_labels->end(), + labels->begin() + current_verts_size); + + auto begin_iter = thrust::make_zip_iterator(verts.begin(), labels->begin()); + + thrust::sort(handle.get_thrust_policy(), begin_iter, begin_iter + new_verts_size); + + auto end_iter = + thrust::unique(handle.get_thrust_policy(), begin_iter, begin_iter + new_verts_size); + + verts.resize(thrust::distance(begin_iter, end_iter), handle.get_stream()); + labels->resize(thrust::distance(begin_iter, end_iter), handle.get_stream()); + } else { + thrust::sort(handle.get_thrust_policy(), verts.begin(), verts.end()); + + auto end_iter = thrust::unique(handle.get_thrust_policy(), verts.begin(), verts.end()); + + verts.resize(thrust::distance(verts.begin(), end_iter), handle.get_stream()); + } + + // Now with the updated verts/labels we can filter the next frontier + std::tie(frontier_vertices, frontier_vertex_labels) = remove_visited_vertices_from_frontier( + handle, + std::move(frontier_vertices), + std::move(frontier_vertex_labels), + raft::device_span{verts.data(), verts.size()}, + labels ? std::make_optional(raft::device_span{labels->data(), labels->size()}) + : std::nullopt); + } + + if (dedupe_sources) { + if (frontier_vertex_labels) { + auto begin_iter = + thrust::make_zip_iterator(frontier_vertices.begin(), frontier_vertex_labels->begin()); + + auto new_end = thrust::unique( + handle.get_thrust_policy(), begin_iter, begin_iter + frontier_vertices.size()); + + frontier_vertices.resize(thrust::distance(begin_iter, new_end), handle.get_stream()); + frontier_vertex_labels->resize(thrust::distance(begin_iter, new_end), handle.get_stream()); + } else { + auto new_end = thrust::unique( + handle.get_thrust_policy(), frontier_vertices.begin(), frontier_vertices.end()); + + frontier_vertices.resize(thrust::distance(frontier_vertices.begin(), new_end), + handle.get_stream()); + } + } + + return std::make_tuple(std::move(frontier_vertices), + std::move(frontier_vertex_labels), + std::move(vertex_used_as_source)); +} + +} // namespace detail +} // namespace cugraph diff --git a/cpp/src/sampling/detail/prepare_next_frontier_mg.cu b/cpp/src/sampling/detail/prepare_next_frontier_mg.cu new file mode 100644 index 00000000000..60b4f49a1f7 --- /dev/null +++ b/cpp/src/sampling/detail/prepare_next_frontier_mg.cu @@ -0,0 +1,59 @@ +/* + * Copyright (c) 2022-2023, 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 + +namespace cugraph { +namespace detail { + +template std::tuple, + std::optional>, + std::optional, + std::optional>>>> +prepare_next_frontier( + raft::handle_t const& handle, + raft::device_span sampled_src_vertices, + std::optional> sampled_src_vertex_labels, + raft::device_span sampled_dst_vertices, + std::optional> sampled_dst_vertex_labels, + std::optional, + std::optional>>>&& vertex_used_as_source, + vertex_partition_view_t vertex_partition, + std::vector const& vertex_partition_range_lasts, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, + bool do_expensive_check); + +template std::tuple, + std::optional>, + std::optional, + std::optional>>>> +prepare_next_frontier( + raft::handle_t const& handle, + raft::device_span sampled_src_vertices, + std::optional> sampled_src_vertex_labels, + raft::device_span sampled_dst_vertices, + std::optional> sampled_dst_vertex_labels, + std::optional, + std::optional>>>&& vertex_used_as_source, + vertex_partition_view_t vertex_partition, + std::vector const& vertex_partition_range_lasts, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, + bool do_expensive_check); + +} // namespace detail +} // namespace cugraph diff --git a/cpp/src/sampling/detail/prepare_next_frontier_sg.cu b/cpp/src/sampling/detail/prepare_next_frontier_sg.cu new file mode 100644 index 00000000000..2c2d28e6843 --- /dev/null +++ b/cpp/src/sampling/detail/prepare_next_frontier_sg.cu @@ -0,0 +1,59 @@ +/* + * Copyright (c) 2022-2023, 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 + +namespace cugraph { +namespace detail { + +template std::tuple, + std::optional>, + std::optional, + std::optional>>>> +prepare_next_frontier( + raft::handle_t const& handle, + raft::device_span sampled_src_vertices, + std::optional> sampled_src_vertex_labels, + raft::device_span sampled_dst_vertices, + std::optional> sampled_dst_vertex_labels, + std::optional, + std::optional>>>&& vertex_used_as_source, + vertex_partition_view_t vertex_partition, + std::vector const& vertex_partition_range_lasts, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, + bool do_expensive_check); + +template std::tuple, + std::optional>, + std::optional, + std::optional>>>> +prepare_next_frontier( + raft::handle_t const& handle, + raft::device_span sampled_src_vertices, + std::optional> sampled_src_vertex_labels, + raft::device_span sampled_dst_vertices, + std::optional> sampled_dst_vertex_labels, + std::optional, + std::optional>>>&& vertex_used_as_source, + vertex_partition_view_t vertex_partition, + std::vector const& vertex_partition_range_lasts, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, + bool do_expensive_check); + +} // namespace detail +} // namespace cugraph diff --git a/cpp/src/sampling/detail/remove_visited_vertices_from_frontier.cu b/cpp/src/sampling/detail/remove_visited_vertices_from_frontier.cu new file mode 100644 index 00000000000..a3b224498ed --- /dev/null +++ b/cpp/src/sampling/detail/remove_visited_vertices_from_frontier.cu @@ -0,0 +1,92 @@ +/* + * Copyright (c) 2023, 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 + +namespace cugraph { +namespace detail { + +template +std::tuple, std::optional>> +remove_visited_vertices_from_frontier( + raft::handle_t const& handle, + rmm::device_uvector&& frontier_vertices, + std::optional>&& frontier_vertex_labels, + raft::device_span vertices_used_as_source, + std::optional> vertex_labels_used_as_source) +{ + if (frontier_vertex_labels) { + auto begin_iter = + thrust::make_zip_iterator(frontier_vertices.begin(), frontier_vertex_labels->begin()); + auto new_end = thrust::remove_if( + handle.get_thrust_policy(), + begin_iter, + begin_iter + frontier_vertices.size(), + begin_iter, + [a_begin = vertices_used_as_source.begin(), + a_end = vertices_used_as_source.end(), + b_begin = vertex_labels_used_as_source->begin(), + b_end = + vertex_labels_used_as_source->end()] __device__(thrust::tuple tuple) { + return thrust::binary_search(thrust::seq, + thrust::make_zip_iterator(a_begin, b_begin), + thrust::make_zip_iterator(a_end, b_end), + tuple); + }); + + frontier_vertices.resize(thrust::distance(begin_iter, new_end), handle.get_stream()); + frontier_vertex_labels->resize(thrust::distance(begin_iter, new_end), handle.get_stream()); + } else { + auto new_end = thrust::copy_if( + handle.get_thrust_policy(), + frontier_vertices.begin(), + frontier_vertices.end(), + frontier_vertices.begin(), + [a_begin = vertices_used_as_source.begin(), a_end = vertices_used_as_source.end()] __device__( + vertex_t v) { return !thrust::binary_search(thrust::seq, a_begin, a_end, v); }); + frontier_vertices.resize(thrust::distance(frontier_vertices.begin(), new_end), + handle.get_stream()); + } + + return std::make_tuple(std::move(frontier_vertices), std::move(frontier_vertex_labels)); +} + +template std::tuple, std::optional>> +remove_visited_vertices_from_frontier( + raft::handle_t const& handle, + rmm::device_uvector&& frontier_vertices, + std::optional>&& frontier_vertex_labels, + raft::device_span vertices_used_as_source, + std::optional> vertex_labels_used_as_source); + +template std::tuple, std::optional>> +remove_visited_vertices_from_frontier( + raft::handle_t const& handle, + rmm::device_uvector&& frontier_vertices, + std::optional>&& frontier_vertex_labels, + raft::device_span vertices_used_as_source, + std::optional> vertex_labels_used_as_source); + +} // namespace detail +} // namespace cugraph diff --git a/cpp/src/sampling/detail/sample_edges.cuh b/cpp/src/sampling/detail/sample_edges.cuh new file mode 100644 index 00000000000..3137d56c234 --- /dev/null +++ b/cpp/src/sampling/detail/sample_edges.cuh @@ -0,0 +1,277 @@ +/* + * Copyright (c) 2022-2023, 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. + */ + +#pragma once + +#include +#include // ?? +#include +#include + +#include +#include +#include +#include + +#include + +#include +#include +#include + +#include + +namespace cugraph { +namespace detail { + +template +struct sample_edges_op_t { + template + auto __host__ __device__ operator()(vertex_t src, + vertex_t dst, + thrust::nullopt_t, + thrust::nullopt_t, + EdgeProperties edge_properties) const + { + // FIXME: A solution using thrust_tuple_cat would be more flexible here + if constexpr (std::is_same_v) { + return thrust::make_tuple(src, dst); + } else if constexpr (std::is_arithmetic::value) { + return thrust::make_tuple(src, dst, edge_properties); + } else if constexpr (cugraph::is_thrust_tuple_of_arithmetic::value && + (thrust::tuple_size::value == 2)) { + return thrust::make_tuple( + src, dst, thrust::get<0>(edge_properties), thrust::get<1>(edge_properties)); + } else if constexpr (cugraph::is_thrust_tuple_of_arithmetic::value && + (thrust::tuple_size::value == 3)) { + return thrust::make_tuple(src, + dst, + thrust::get<0>(edge_properties), + thrust::get<1>(edge_properties), + thrust::get<2>(edge_properties)); + } + } +}; + +struct segmented_fill_t { + raft::device_span fill_values{}; + raft::device_span segment_offsets{}; + raft::device_span output_values{}; + + __device__ void operator()(size_t i) const + { + thrust::fill(thrust::seq, + output_values.begin() + segment_offsets[i], + output_values.begin() + segment_offsets[i + 1], + fill_values[i]); + } +}; + +template +std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +sample_edges(raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_type_view, + raft::random::RngState& rng_state, + raft::device_span active_majors, + std::optional> active_major_labels, + size_t fanout, + bool with_replacement) +{ + using tag_t = void; + + cugraph::vertex_frontier_t vertex_frontier(handle, 1); + + vertex_frontier.bucket(0).insert(active_majors.begin(), active_majors.end()); + + rmm::device_uvector majors(0, handle.get_stream()); + rmm::device_uvector minors(0, handle.get_stream()); + std::optional> edge_ids{std::nullopt}; + std::optional> weights{std::nullopt}; + std::optional> edge_types{std::nullopt}; + std::optional> labels{std::nullopt}; + std::optional> sample_offsets{std::nullopt}; + + if (edge_weight_view) { + if (edge_id_view) { + if (edge_type_view) { + std::forward_as_tuple(sample_offsets, + std::tie(majors, minors, weights, edge_ids, edge_types)) = + cugraph::per_v_random_select_transform_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + view_concat(*edge_weight_view, *edge_id_view, *edge_type_view), + sample_edges_op_t{}, + rng_state, + fanout, + with_replacement, + std::optional>{ + std::nullopt}, + true); + } else { + std::forward_as_tuple(sample_offsets, std::tie(majors, minors, weights, edge_ids)) = + cugraph::per_v_random_select_transform_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + view_concat(*edge_weight_view, *edge_id_view), + sample_edges_op_t{}, + rng_state, + fanout, + with_replacement, + std::optional>{std::nullopt}, + true); + } + } else { + if (edge_type_view) { + std::forward_as_tuple(sample_offsets, std::tie(majors, minors, weights, edge_types)) = + cugraph::per_v_random_select_transform_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + view_concat(*edge_weight_view, *edge_type_view), + sample_edges_op_t{}, + rng_state, + fanout, + with_replacement, + std::optional>{std::nullopt}, + true); + } else { + std::forward_as_tuple(sample_offsets, std::tie(majors, minors, weights)) = + cugraph::per_v_random_select_transform_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + *edge_weight_view, + sample_edges_op_t{}, + rng_state, + fanout, + with_replacement, + std::optional>{std::nullopt}, + true); + } + } + } else { + if (edge_id_view) { + if (edge_type_view) { + std::forward_as_tuple(sample_offsets, std::tie(majors, minors, edge_ids, edge_types)) = + cugraph::per_v_random_select_transform_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + view_concat(*edge_id_view, *edge_type_view), + sample_edges_op_t{}, + rng_state, + fanout, + with_replacement, + std::optional>{std::nullopt}, + true); + } else { + std::forward_as_tuple(sample_offsets, std::tie(majors, minors, edge_ids)) = + cugraph::per_v_random_select_transform_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + *edge_id_view, + sample_edges_op_t{}, + rng_state, + fanout, + with_replacement, + std::optional>{std::nullopt}, + true); + } + } else { + if (edge_type_view) { + std::forward_as_tuple(sample_offsets, std::tie(majors, minors, edge_types)) = + cugraph::per_v_random_select_transform_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + *edge_type_view, + sample_edges_op_t{}, + rng_state, + fanout, + with_replacement, + std::optional>{std::nullopt}, + true); + } else { + std::forward_as_tuple(sample_offsets, std::tie(majors, minors)) = + cugraph::per_v_random_select_transform_outgoing_e( + handle, + graph_view, + vertex_frontier.bucket(0), + edge_src_dummy_property_t{}.view(), + edge_dst_dummy_property_t{}.view(), + edge_dummy_property_t{}.view(), + sample_edges_op_t{}, + rng_state, + fanout, + with_replacement, + std::optional>{std::nullopt}, + true); + } + } + } + + if (active_major_labels) { + labels = rmm::device_uvector((*sample_offsets).back_element(handle.get_stream()), + handle.get_stream()); + thrust::for_each(handle.get_thrust_policy(), + thrust::make_counting_iterator(size_t{0}), + thrust::make_counting_iterator(active_majors.size()), + segmented_fill_t{*active_major_labels, + raft::device_span(sample_offsets->data(), + sample_offsets->size()), + raft::device_span(labels->data(), labels->size())}); + } + + return std::make_tuple(std::move(majors), + std::move(minors), + std::move(weights), + std::move(edge_ids), + std::move(edge_types), + std::move(labels)); +} + +} // namespace detail +} // namespace cugraph diff --git a/cpp/src/sampling/detail/sample_edges_mg.cu b/cpp/src/sampling/detail/sample_edges_mg.cu new file mode 100644 index 00000000000..4f42f68b1c4 --- /dev/null +++ b/cpp/src/sampling/detail/sample_edges_mg.cu @@ -0,0 +1,125 @@ +/* + * Copyright (c) 2022-2023, 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 + +namespace cugraph { +namespace detail { + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +sample_edges(raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_edge_type_view, + raft::random::RngState& rng_state, + raft::device_span active_majors, + std::optional> active_major_labels, + size_t fanout, + bool with_replacement); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +sample_edges(raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_edge_type_view, + raft::random::RngState& rng_state, + raft::device_span active_majors, + std::optional> active_major_labels, + size_t fanout, + bool with_replacement); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +sample_edges(raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_edge_type_view, + raft::random::RngState& rng_state, + raft::device_span active_majors, + std::optional> active_major_labels, + size_t fanout, + bool with_replacement); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +sample_edges(raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_edge_type_view, + raft::random::RngState& rng_state, + raft::device_span active_majors, + std::optional> active_major_labels, + size_t fanout, + bool with_replacement); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +sample_edges(raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_edge_type_view, + raft::random::RngState& rng_state, + raft::device_span active_majors, + std::optional> active_major_labels, + size_t fanout, + bool with_replacement); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +sample_edges(raft::handle_t const& handle, + graph_view_t const& graph_view, + std::optional> edge_weight_view, + std::optional> edge_id_view, + std::optional> edge_edge_type_view, + raft::random::RngState& rng_state, + raft::device_span active_majors, + std::optional> active_major_labels, + size_t fanout, + bool with_replacement); + +} // namespace detail +} // namespace cugraph diff --git a/cpp/src/sampling/detail/sampling_utils_sg.cu b/cpp/src/sampling/detail/sample_edges_sg.cu similarity index 56% rename from cpp/src/sampling/detail/sampling_utils_sg.cu rename to cpp/src/sampling/detail/sample_edges_sg.cu index 7d0ec0ed420..f415789434a 100644 --- a/cpp/src/sampling/detail/sampling_utils_sg.cu +++ b/cpp/src/sampling/detail/sample_edges_sg.cu @@ -14,107 +14,11 @@ * limitations under the License. */ -#include +#include namespace cugraph { namespace detail { -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_edge_type_view, - raft::device_span active_majors, - std::optional> active_major_labels, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_edge_type_view, - raft::device_span active_majors, - std::optional> active_major_labels, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_edge_type_view, - raft::device_span active_majors, - std::optional> active_major_labels, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_edge_type_view, - raft::device_span active_majors, - std::optional> active_major_labels, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_edge_type_view, - raft::device_span active_majors, - std::optional> active_major_labels, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_edge_type_view, - raft::device_span active_majors, - std::optional> active_major_labels, - bool do_expensive_check); - template std::tuple, rmm::device_uvector, std::optional>, diff --git a/cpp/src/sampling/detail/graph_functions.hpp b/cpp/src/sampling/detail/sampling_utils.hpp similarity index 69% rename from cpp/src/sampling/detail/graph_functions.hpp rename to cpp/src/sampling/detail/sampling_utils.hpp index de9205d19c7..2ece35c06fe 100644 --- a/cpp/src/sampling/detail/graph_functions.hpp +++ b/cpp/src/sampling/detail/sampling_utils.hpp @@ -16,6 +16,7 @@ #pragma once +#include #include namespace cugraph { @@ -120,6 +121,80 @@ sample_edges(raft::handle_t const& handle, size_t fanout, bool with_replacement); +/** + * @brief Use the sampling results from hop N to populate the new frontier for hop N+1. + * + * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. + * @tparam label_t Type of label. Needs to be an integral type. + * @tparam multi_gpu Flag indicating whether template instantiation should target single-GPU (false) + * + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param sampled_src_vertices The source vertices for the current sampling + * @param sampled_src_vertex_labels Optional labels for the vertices for the current sampling + * @param sampled_dst_vertices Vertices for the next frontier + * @param sampled_dst_vertex_labels Optional labels for the next frontier + * @param vertex_used_as_source Optional. If specified then we want to exclude vertices that + * were previously used as sources. These vertices (and optional labels) will be updated based + * on the contents of sampled_src_vertices/sampled_src_vertex_labels and the update will be part + * of the return value. + * @param vertex_partition Vertex partition view from the graph view + * @param vertex_partition_range_lasts End of range information from graph view + * @param prior_sources_behavior Identifies how to treat sources in each hop + * @param dedupe_sources boolean flag, if true then if a vertex v appears as a destination in hop X + * multiple times with the same label, it will only be passed once (for each label) as a source + * for the next hop. Default is false. + * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`). + * + * @return A tuple of device vectors containing the vertices for the next frontier expansion and + * optional labels associated with the vertices, along with the updated value for + * @p vertex_used_as_sources + */ +template +std::tuple, + std::optional>, + std::optional, + std::optional>>>> +prepare_next_frontier( + raft::handle_t const& handle, + raft::device_span sampled_src_vertices, + std::optional> sampled_src_vertex_labels, + raft::device_span sampled_dst_vertices, + std::optional> sampled_dst_vertex_labels, + std::optional, + std::optional>>>&& vertex_used_as_source, + vertex_partition_view_t vertex_partition, + std::vector const& vertex_partition_range_lasts, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, + bool do_expensive_check); + +/** + * @brief Remove from the frontier any vertices that have already been used as a source + * + * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type. + * @tparam label_t Type of label. Needs to be an integral type. + * + * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and + * handles to various CUDA libraries) to run graph algorithms. + * @param frontier_vertices Vertices discovered in the current hop + * @param frontier_vertex_labels Labels for the vertices discovered in the current hop + * @param vertices_used_as_source Device vector containing all vertices used in previous hops as a + * source + * @param vertex_labels_used_as_source Device vector containing vertex labels associated with + * the @p vertices_used_as_source used in previous hops as a source vertex label + * + * @return tuple containing the modified frontier_vertices and frontier_vertex_labels + */ +template +std::tuple, std::optional>> +remove_visited_vertices_from_frontier( + raft::handle_t const& handle, + rmm::device_uvector&& frontier_vertices, + std::optional>&& frontier_vertex_labels, + raft::device_span vertices_used_as_source, + std::optional> vertex_labels_used_as_source); + // FIXME: Should this go in shuffle_wrappers.hpp? /** * @brief Shuffle sampling results to the desired GPU based on label_to_output_gpu_mapping[label[i]] diff --git a/cpp/src/sampling/detail/sampling_utils_mg.cu b/cpp/src/sampling/detail/sampling_utils_mg.cu deleted file mode 100644 index 5a8f682f4d0..00000000000 --- a/cpp/src/sampling/detail/sampling_utils_mg.cu +++ /dev/null @@ -1,341 +0,0 @@ -/* - * Copyright (c) 2022-2023, 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 - -namespace cugraph { -namespace detail { - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_edge_type_view, - raft::device_span active_majors, - std::optional> active_major_labels, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_edge_type_view, - raft::device_span active_majors, - std::optional> active_major_labels, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_edge_type_view, - raft::device_span active_majors, - std::optional> active_major_labels, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_edge_type_view, - raft::device_span active_majors, - std::optional> active_major_labels, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_edge_type_view, - raft::device_span active_majors, - std::optional> active_major_labels, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_edge_type_view, - raft::device_span active_majors, - std::optional> active_major_labels, - bool do_expensive_check); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_edge_type_view, - raft::random::RngState& rng_state, - raft::device_span active_majors, - std::optional> active_major_labels, - size_t fanout, - bool with_replacement); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_edge_type_view, - raft::random::RngState& rng_state, - raft::device_span active_majors, - std::optional> active_major_labels, - size_t fanout, - bool with_replacement); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_edge_type_view, - raft::random::RngState& rng_state, - raft::device_span active_majors, - std::optional> active_major_labels, - size_t fanout, - bool with_replacement); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_edge_type_view, - raft::random::RngState& rng_state, - raft::device_span active_majors, - std::optional> active_major_labels, - size_t fanout, - bool with_replacement); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_edge_type_view, - raft::random::RngState& rng_state, - raft::device_span active_majors, - std::optional> active_major_labels, - size_t fanout, - bool with_replacement); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_edge_type_view, - raft::random::RngState& rng_state, - raft::device_span active_majors, - std::optional> active_major_labels, - size_t fanout, - bool with_replacement); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -shuffle_and_organize_output( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& hops, - std::optional>&& labels, - std::optional, raft::device_span>> - label_to_output_comm_rank); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -shuffle_and_organize_output( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& hops, - std::optional>&& labels, - std::optional, raft::device_span>> - label_to_output_comm_rank); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -shuffle_and_organize_output( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& hops, - std::optional>&& labels, - std::optional, raft::device_span>> - label_to_output_comm_rank); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -shuffle_and_organize_output( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& hops, - std::optional>&& labels, - std::optional, raft::device_span>> - label_to_output_comm_rank); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -shuffle_and_organize_output( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& hops, - std::optional>&& labels, - std::optional, raft::device_span>> - label_to_output_comm_rank); - -template std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -shuffle_and_organize_output( - raft::handle_t const& handle, - rmm::device_uvector&& majors, - rmm::device_uvector&& minors, - std::optional>&& weights, - std::optional>&& edge_ids, - std::optional>&& edge_types, - std::optional>&& hops, - std::optional>&& labels, - std::optional, raft::device_span>> - label_to_output_comm_rank); - -} // namespace detail -} // namespace cugraph diff --git a/cpp/src/sampling/detail/sampling_utils_impl.cuh b/cpp/src/sampling/detail/shuffle_and_organize_output_impl.cuh similarity index 56% rename from cpp/src/sampling/detail/sampling_utils_impl.cuh rename to cpp/src/sampling/detail/shuffle_and_organize_output_impl.cuh index 8ba9e45e8d2..5197ea7f872 100644 --- a/cpp/src/sampling/detail/sampling_utils_impl.cuh +++ b/cpp/src/sampling/detail/shuffle_and_organize_output_impl.cuh @@ -16,17 +16,17 @@ #pragma once -#include -#include #include // ?? #include #include +#include #include #include #include #include #include +#include #include @@ -39,89 +39,6 @@ namespace cugraph { namespace detail { -struct return_edges_with_properties_e_op { - template - auto __host__ __device__ operator()(key_t optionally_tagged_src, - vertex_t dst, - thrust::nullopt_t, - thrust::nullopt_t, - EdgeProperties edge_properties) - { - // FIXME: A solution using thrust_tuple_cat would be more flexible here - if constexpr (std::is_same_v) { - vertex_t src{optionally_tagged_src}; - - if constexpr (std::is_same_v) { - return thrust::make_optional(thrust::make_tuple(src, dst)); - } else if constexpr (std::is_arithmetic::value) { - return thrust::make_optional(thrust::make_tuple(src, dst, edge_properties)); - } else if constexpr (cugraph::is_thrust_tuple_of_arithmetic::value && - (thrust::tuple_size::value == 2)) { - return thrust::make_optional(thrust::make_tuple( - src, dst, thrust::get<0>(edge_properties), thrust::get<1>(edge_properties))); - } else if constexpr (cugraph::is_thrust_tuple_of_arithmetic::value && - (thrust::tuple_size::value == 3)) { - return thrust::make_optional(thrust::make_tuple(src, - dst, - thrust::get<0>(edge_properties), - thrust::get<1>(edge_properties), - thrust::get<2>(edge_properties))); - } - } else if constexpr (std::is_same_v>) { - vertex_t src{thrust::get<0>(optionally_tagged_src)}; - int32_t label{thrust::get<1>(optionally_tagged_src)}; - - src = thrust::get<0>(optionally_tagged_src); - if constexpr (std::is_same_v) { - return thrust::make_optional(thrust::make_tuple(src, dst, label)); - } else if constexpr (std::is_arithmetic::value) { - return thrust::make_optional(thrust::make_tuple(src, dst, edge_properties, label)); - } else if constexpr (cugraph::is_thrust_tuple_of_arithmetic::value && - (thrust::tuple_size::value == 2)) { - return thrust::make_optional(thrust::make_tuple( - src, dst, thrust::get<0>(edge_properties), thrust::get<1>(edge_properties), label)); - } else if constexpr (cugraph::is_thrust_tuple_of_arithmetic::value && - (thrust::tuple_size::value == 3)) { - return thrust::make_optional(thrust::make_tuple(src, - dst, - thrust::get<0>(edge_properties), - thrust::get<1>(edge_properties), - thrust::get<2>(edge_properties), - label)); - } - } - } -}; - -template -struct sample_edges_op_t { - template - auto __host__ __device__ operator()(vertex_t src, - vertex_t dst, - thrust::nullopt_t, - thrust::nullopt_t, - EdgeProperties edge_properties) const - { - // FIXME: A solution using thrust_tuple_cat would be more flexible here - if constexpr (std::is_same_v) { - return thrust::make_tuple(src, dst); - } else if constexpr (std::is_arithmetic::value) { - return thrust::make_tuple(src, dst, edge_properties); - } else if constexpr (cugraph::is_thrust_tuple_of_arithmetic::value && - (thrust::tuple_size::value == 2)) { - return thrust::make_tuple( - src, dst, thrust::get<0>(edge_properties), thrust::get<1>(edge_properties)); - } else if constexpr (cugraph::is_thrust_tuple_of_arithmetic::value && - (thrust::tuple_size::value == 3)) { - return thrust::make_tuple(src, - dst, - thrust::get<0>(edge_properties), - thrust::get<1>(edge_properties), - thrust::get<2>(edge_properties)); - } - } -}; - template struct shuffle_to_output_comm_rank_t { raft::device_span output_label_; @@ -135,510 +52,6 @@ struct shuffle_to_output_comm_rank_t { } }; -struct segmented_fill_t { - raft::device_span fill_values{}; - raft::device_span segment_offsets{}; - raft::device_span output_values{}; - - __device__ void operator()(size_t i) const - { - thrust::fill(thrust::seq, - output_values.begin() + segment_offsets[i], - output_values.begin() + segment_offsets[i + 1], - fill_values[i]); - } -}; - -template -std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_type_view, - cugraph::vertex_frontier_t const& vertex_frontier, - bool do_expensive_check) -{ - rmm::device_uvector majors(0, handle.get_stream()); - rmm::device_uvector minors(0, handle.get_stream()); - std::optional> edge_ids{std::nullopt}; - std::optional> weights{std::nullopt}; - std::optional> edge_types{std::nullopt}; - std::optional> labels{std::nullopt}; - - if (edge_weight_view) { - if (edge_id_view) { - if (edge_type_view) { - if constexpr (std::is_same_v) { - std::tie(majors, minors, weights, edge_ids, edge_types, labels) = - cugraph::extract_transform_v_frontier_outgoing_e( - handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - view_concat(*edge_weight_view, *edge_id_view, *edge_type_view), - return_edges_with_properties_e_op{}, - do_expensive_check); - } else { - std::tie(majors, minors, weights, edge_ids, edge_types) = - cugraph::extract_transform_v_frontier_outgoing_e( - handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - view_concat(*edge_weight_view, *edge_id_view, *edge_type_view), - return_edges_with_properties_e_op{}, - do_expensive_check); - } - } else { - if constexpr (std::is_same_v) { - std::tie(majors, minors, weights, edge_ids, labels) = - cugraph::extract_transform_v_frontier_outgoing_e( - handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - view_concat(*edge_weight_view, *edge_id_view), - return_edges_with_properties_e_op{}, - do_expensive_check); - } else { - std::tie(majors, minors, weights, edge_ids) = - cugraph::extract_transform_v_frontier_outgoing_e( - handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - view_concat(*edge_weight_view, *edge_id_view), - return_edges_with_properties_e_op{}, - do_expensive_check); - } - } - } else { - if (edge_type_view) { - if constexpr (std::is_same_v) { - std::tie(majors, minors, weights, edge_types, labels) = - cugraph::extract_transform_v_frontier_outgoing_e( - handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - view_concat(*edge_weight_view, *edge_type_view), - return_edges_with_properties_e_op{}, - do_expensive_check); - } else { - std::tie(majors, minors, weights, edge_types) = - cugraph::extract_transform_v_frontier_outgoing_e( - handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - view_concat(*edge_weight_view, *edge_type_view), - return_edges_with_properties_e_op{}, - do_expensive_check); - } - } else { - if constexpr (std::is_same_v) { - std::tie(majors, minors, weights, labels) = - cugraph::extract_transform_v_frontier_outgoing_e(handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - *edge_weight_view, - return_edges_with_properties_e_op{}, - do_expensive_check); - } else { - std::tie(majors, minors, weights) = - cugraph::extract_transform_v_frontier_outgoing_e(handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - *edge_weight_view, - return_edges_with_properties_e_op{}, - do_expensive_check); - } - } - } - } else { - if (edge_id_view) { - if (edge_type_view) { - if constexpr (std::is_same_v) { - std::tie(majors, minors, edge_ids, edge_types, labels) = - cugraph::extract_transform_v_frontier_outgoing_e( - handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - view_concat(*edge_id_view, *edge_type_view), - return_edges_with_properties_e_op{}, - do_expensive_check); - } else { - std::tie(majors, minors, edge_ids, edge_types) = - cugraph::extract_transform_v_frontier_outgoing_e( - handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - view_concat(*edge_id_view, *edge_type_view), - return_edges_with_properties_e_op{}, - do_expensive_check); - } - } else { - if constexpr (std::is_same_v) { - std::tie(majors, minors, edge_ids, labels) = - cugraph::extract_transform_v_frontier_outgoing_e(handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - *edge_id_view, - return_edges_with_properties_e_op{}, - do_expensive_check); - } else { - std::tie(majors, minors, edge_ids) = - cugraph::extract_transform_v_frontier_outgoing_e(handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - *edge_id_view, - return_edges_with_properties_e_op{}, - do_expensive_check); - } - } - } else { - if (edge_type_view) { - if constexpr (std::is_same_v) { - std::tie(majors, minors, edge_types, labels) = - cugraph::extract_transform_v_frontier_outgoing_e(handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - *edge_type_view, - return_edges_with_properties_e_op{}, - do_expensive_check); - } else { - std::tie(majors, minors, edge_types) = - cugraph::extract_transform_v_frontier_outgoing_e(handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - *edge_type_view, - return_edges_with_properties_e_op{}, - do_expensive_check); - } - } else { - if constexpr (std::is_same_v) { - std::tie(majors, minors, labels) = - cugraph::extract_transform_v_frontier_outgoing_e(handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - edge_dummy_property_t{}.view(), - return_edges_with_properties_e_op{}, - do_expensive_check); - } else { - std::tie(majors, minors) = - cugraph::extract_transform_v_frontier_outgoing_e(handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - edge_dummy_property_t{}.view(), - return_edges_with_properties_e_op{}, - do_expensive_check); - } - } - } - } - - return std::make_tuple(std::move(majors), - std::move(minors), - std::move(weights), - std::move(edge_ids), - std::move(edge_types), - std::move(labels)); -} - -template -std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -gather_one_hop_edgelist( - raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_type_view, - raft::device_span active_majors, - std::optional> active_major_labels, - bool do_expensive_check) -{ - if (active_major_labels) { - cugraph::vertex_frontier_t vertex_label_frontier(handle, - 1); - vertex_label_frontier.bucket(0).insert( - thrust::make_zip_iterator(active_majors.begin(), active_major_labels->begin()), - thrust::make_zip_iterator(active_majors.end(), active_major_labels->end())); - - return gather_one_hop_edgelist(handle, - graph_view, - edge_weight_view, - edge_id_view, - edge_type_view, - vertex_label_frontier, - do_expensive_check); - } else { - cugraph::vertex_frontier_t vertex_frontier(handle, 1); - vertex_frontier.bucket(0).insert(active_majors.begin(), active_majors.end()); - - return gather_one_hop_edgelist(handle, - graph_view, - edge_weight_view, - edge_id_view, - edge_type_view, - vertex_frontier, - do_expensive_check); - } -} - -template -std::tuple, - rmm::device_uvector, - std::optional>, - std::optional>, - std::optional>, - std::optional>> -sample_edges(raft::handle_t const& handle, - graph_view_t const& graph_view, - std::optional> edge_weight_view, - std::optional> edge_id_view, - std::optional> edge_type_view, - raft::random::RngState& rng_state, - raft::device_span active_majors, - std::optional> active_major_labels, - size_t fanout, - bool with_replacement) -{ - using tag_t = void; - - cugraph::vertex_frontier_t vertex_frontier(handle, 1); - - vertex_frontier.bucket(0).insert(active_majors.begin(), active_majors.end()); - - rmm::device_uvector majors(0, handle.get_stream()); - rmm::device_uvector minors(0, handle.get_stream()); - std::optional> edge_ids{std::nullopt}; - std::optional> weights{std::nullopt}; - std::optional> edge_types{std::nullopt}; - std::optional> labels{std::nullopt}; - std::optional> sample_offsets{std::nullopt}; - - if (edge_weight_view) { - if (edge_id_view) { - if (edge_type_view) { - std::forward_as_tuple(sample_offsets, - std::tie(majors, minors, weights, edge_ids, edge_types)) = - cugraph::per_v_random_select_transform_outgoing_e( - handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - view_concat(*edge_weight_view, *edge_id_view, *edge_type_view), - sample_edges_op_t{}, - rng_state, - fanout, - with_replacement, - std::optional>{ - std::nullopt}, - true); - } else { - std::forward_as_tuple(sample_offsets, std::tie(majors, minors, weights, edge_ids)) = - cugraph::per_v_random_select_transform_outgoing_e( - handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - view_concat(*edge_weight_view, *edge_id_view), - sample_edges_op_t{}, - rng_state, - fanout, - with_replacement, - std::optional>{std::nullopt}, - true); - } - } else { - if (edge_type_view) { - std::forward_as_tuple(sample_offsets, std::tie(majors, minors, weights, edge_types)) = - cugraph::per_v_random_select_transform_outgoing_e( - handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - view_concat(*edge_weight_view, *edge_type_view), - sample_edges_op_t{}, - rng_state, - fanout, - with_replacement, - std::optional>{std::nullopt}, - true); - } else { - std::forward_as_tuple(sample_offsets, std::tie(majors, minors, weights)) = - cugraph::per_v_random_select_transform_outgoing_e( - handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - *edge_weight_view, - sample_edges_op_t{}, - rng_state, - fanout, - with_replacement, - std::optional>{std::nullopt}, - true); - } - } - } else { - if (edge_id_view) { - if (edge_type_view) { - std::forward_as_tuple(sample_offsets, std::tie(majors, minors, edge_ids, edge_types)) = - cugraph::per_v_random_select_transform_outgoing_e( - handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - view_concat(*edge_id_view, *edge_type_view), - sample_edges_op_t{}, - rng_state, - fanout, - with_replacement, - std::optional>{std::nullopt}, - true); - } else { - std::forward_as_tuple(sample_offsets, std::tie(majors, minors, edge_ids)) = - cugraph::per_v_random_select_transform_outgoing_e( - handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - *edge_id_view, - sample_edges_op_t{}, - rng_state, - fanout, - with_replacement, - std::optional>{std::nullopt}, - true); - } - } else { - if (edge_type_view) { - std::forward_as_tuple(sample_offsets, std::tie(majors, minors, edge_types)) = - cugraph::per_v_random_select_transform_outgoing_e( - handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - *edge_type_view, - sample_edges_op_t{}, - rng_state, - fanout, - with_replacement, - std::optional>{std::nullopt}, - true); - } else { - std::forward_as_tuple(sample_offsets, std::tie(majors, minors)) = - cugraph::per_v_random_select_transform_outgoing_e( - handle, - graph_view, - vertex_frontier.bucket(0), - edge_src_dummy_property_t{}.view(), - edge_dst_dummy_property_t{}.view(), - edge_dummy_property_t{}.view(), - sample_edges_op_t{}, - rng_state, - fanout, - with_replacement, - std::optional>{std::nullopt}, - true); - } - } - } - - if (active_major_labels) { - labels = rmm::device_uvector((*sample_offsets).back_element(handle.get_stream()), - handle.get_stream()); - thrust::for_each(handle.get_thrust_policy(), - thrust::make_counting_iterator(size_t{0}), - thrust::make_counting_iterator(active_majors.size()), - segmented_fill_t{*active_major_labels, - raft::device_span(sample_offsets->data(), - sample_offsets->size()), - raft::device_span(labels->data(), labels->size())}); - } - - return std::make_tuple(std::move(majors), - std::move(minors), - std::move(weights), - std::move(edge_ids), - std::move(edge_types), - std::move(labels)); -} - template + +namespace cugraph { +namespace detail { + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +shuffle_and_organize_output( + raft::handle_t const& handle, + rmm::device_uvector&& majors, + rmm::device_uvector&& minors, + std::optional>&& weights, + std::optional>&& edge_ids, + std::optional>&& edge_types, + std::optional>&& hops, + std::optional>&& labels, + std::optional, raft::device_span>> + label_to_output_comm_rank); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +shuffle_and_organize_output( + raft::handle_t const& handle, + rmm::device_uvector&& majors, + rmm::device_uvector&& minors, + std::optional>&& weights, + std::optional>&& edge_ids, + std::optional>&& edge_types, + std::optional>&& hops, + std::optional>&& labels, + std::optional, raft::device_span>> + label_to_output_comm_rank); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +shuffle_and_organize_output( + raft::handle_t const& handle, + rmm::device_uvector&& majors, + rmm::device_uvector&& minors, + std::optional>&& weights, + std::optional>&& edge_ids, + std::optional>&& edge_types, + std::optional>&& hops, + std::optional>&& labels, + std::optional, raft::device_span>> + label_to_output_comm_rank); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +shuffle_and_organize_output( + raft::handle_t const& handle, + rmm::device_uvector&& majors, + rmm::device_uvector&& minors, + std::optional>&& weights, + std::optional>&& edge_ids, + std::optional>&& edge_types, + std::optional>&& hops, + std::optional>&& labels, + std::optional, raft::device_span>> + label_to_output_comm_rank); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +shuffle_and_organize_output( + raft::handle_t const& handle, + rmm::device_uvector&& majors, + rmm::device_uvector&& minors, + std::optional>&& weights, + std::optional>&& edge_ids, + std::optional>&& edge_types, + std::optional>&& hops, + std::optional>&& labels, + std::optional, raft::device_span>> + label_to_output_comm_rank); + +template std::tuple, + rmm::device_uvector, + std::optional>, + std::optional>, + std::optional>, + std::optional>, + std::optional>, + std::optional>> +shuffle_and_organize_output( + raft::handle_t const& handle, + rmm::device_uvector&& majors, + rmm::device_uvector&& minors, + std::optional>&& weights, + std::optional>&& edge_ids, + std::optional>&& edge_types, + std::optional>&& hops, + std::optional>&& labels, + std::optional, raft::device_span>> + label_to_output_comm_rank); + +} // namespace detail +} // namespace cugraph diff --git a/cpp/src/sampling/uniform_neighbor_sampling_impl.hpp b/cpp/src/sampling/uniform_neighbor_sampling_impl.hpp index 7dbc98840e7..d9ffbcf56d5 100644 --- a/cpp/src/sampling/uniform_neighbor_sampling_impl.hpp +++ b/cpp/src/sampling/uniform_neighbor_sampling_impl.hpp @@ -16,11 +16,12 @@ #pragma once -#include +#include #include #include #include +#include #include @@ -50,13 +51,15 @@ uniform_neighbor_sample_impl( std::optional> edge_weight_view, std::optional> edge_id_view, std::optional> edge_type_view, - raft::device_span starting_vertices, - std::optional> starting_vertex_labels, + raft::device_span this_frontier_vertices, + std::optional> this_frontier_vertex_labels, std::optional, raft::device_span>> label_to_output_comm_rank, raft::host_span fan_out, bool return_hops, bool with_replacement, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, raft::random::RngState& rng_state, bool do_expensive_check) { @@ -76,8 +79,8 @@ uniform_neighbor_sample_impl( } CUGRAPH_EXPECTS( - !label_to_output_comm_rank || starting_vertex_labels, - "cannot specify output GPU mapping without also specifying starting_vertex_labels"); + !label_to_output_comm_rank || this_frontier_vertex_labels, + "cannot specify output GPU mapping without also specifying this_frontier_vertex_labels"); if (do_expensive_check) { if (label_to_output_comm_rank) { @@ -97,8 +100,8 @@ uniform_neighbor_sample_impl( edge_type_view ? std::make_optional(std::vector>{}) : std::nullopt; auto level_result_label_vectors = - starting_vertex_labels ? std::make_optional(std::vector>{}) - : std::nullopt; + this_frontier_vertex_labels ? std::make_optional(std::vector>{}) + : std::nullopt; level_result_src_vectors.reserve(fan_out.size()); level_result_dst_vectors.reserve(fan_out.size()); @@ -109,10 +112,22 @@ uniform_neighbor_sample_impl( rmm::device_uvector frontier_vertices(0, handle.get_stream()); auto frontier_vertex_labels = - starting_vertex_labels + this_frontier_vertex_labels ? std::make_optional(rmm::device_uvector{0, handle.get_stream()}) : std::nullopt; + std::optional< + std::tuple, std::optional>>> + vertex_used_as_source{std::nullopt}; + + if (prior_sources_behavior == prior_sources_behavior_t::EXCLUDE) { + vertex_used_as_source = std::make_optional( + std::make_tuple(rmm::device_uvector{0, handle.get_stream()}, + this_frontier_vertex_labels + ? std::make_optional(rmm::device_uvector{0, handle.get_stream()}) + : std::nullopt)); + } + std::vector level_sizes{}; int32_t hop{0}; for (auto&& k_level : fan_out) { @@ -131,8 +146,8 @@ uniform_neighbor_sample_impl( edge_id_view, edge_type_view, rng_state, - starting_vertices, - starting_vertex_labels, + this_frontier_vertices, + this_frontier_vertex_labels, static_cast(k_level), with_replacement); } else { @@ -142,8 +157,8 @@ uniform_neighbor_sample_impl( edge_weight_view, edge_id_view, edge_type_view, - starting_vertices, - starting_vertex_labels); + this_frontier_vertices, + this_frontier_vertex_labels); } level_sizes.push_back(srcs.size()); @@ -157,48 +172,33 @@ uniform_neighbor_sample_impl( ++hop; if (hop < fan_out.size()) { - if constexpr (multi_gpu) { - size_t frontier_size = level_result_dst_vectors.back().size(); - frontier_vertices.resize(frontier_size, handle.get_stream()); - - raft::copy(frontier_vertices.begin(), - level_result_dst_vectors.back().data(), - level_result_dst_vectors.back().size(), - handle.get_stream()); - - if (starting_vertex_labels) { - frontier_vertex_labels->resize(frontier_size, handle.get_stream()); - raft::copy(frontier_vertex_labels->begin(), - level_result_label_vectors->back().data(), - frontier_size, - handle.get_stream()); - - std::tie(frontier_vertices, *frontier_vertex_labels) = - shuffle_int_vertex_value_pairs_to_local_gpu_by_vertex_partitioning( - handle, - std::move(frontier_vertices), - std::move(*frontier_vertex_labels), - graph_view.vertex_partition_range_lasts()); - - starting_vertices = - raft::device_span(frontier_vertices.data(), frontier_vertices.size()); - - starting_vertex_labels = raft::device_span(frontier_vertex_labels->data(), - frontier_vertex_labels->size()); - } else { - frontier_vertices = shuffle_int_vertices_to_local_gpu_by_vertex_partitioning( - handle, std::move(frontier_vertices), graph_view.vertex_partition_range_lasts()); - - starting_vertices = - raft::device_span(frontier_vertices.data(), frontier_vertices.size()); - } - } else { - starting_vertices = raft::device_span( - level_result_dst_vectors.back().data(), level_result_dst_vectors.back().size()); - if (starting_vertex_labels) { - starting_vertex_labels = raft::device_span( - level_result_label_vectors->back().data(), level_result_label_vectors->back().size()); - } + // FIXME: We should modify vertex_partition_range_lasts to return a raft::host_span + // rather than making a copy. + auto vertex_partition_range_lasts = graph_view.vertex_partition_range_lasts(); + std::tie(frontier_vertices, frontier_vertex_labels, vertex_used_as_source) = + prepare_next_frontier( + handle, + this_frontier_vertices, + this_frontier_vertex_labels, + raft::device_span{level_result_dst_vectors.back().data(), + level_result_dst_vectors.back().size()}, + frontier_vertex_labels ? std::make_optional(raft::device_span( + level_result_label_vectors->back().data(), + level_result_label_vectors->back().size())) + : std::nullopt, + std::move(vertex_used_as_source), + graph_view.local_vertex_partition_view(), + vertex_partition_range_lasts, + prior_sources_behavior, + dedupe_sources, + do_expensive_check); + + this_frontier_vertices = + raft::device_span(frontier_vertices.data(), frontier_vertices.size()); + + if (frontier_vertex_labels) { + this_frontier_vertex_labels = raft::device_span( + frontier_vertex_labels->data(), frontier_vertex_labels->size()); } } } @@ -349,6 +349,8 @@ uniform_neighbor_sample( raft::random::RngState& rng_state, bool return_hops, bool with_replacement, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, bool do_expensive_check) { CUGRAPH_EXPECTS(!graph_view.has_edge_mask(), "unimplemented."); @@ -364,6 +366,8 @@ uniform_neighbor_sample( fan_out, return_hops, with_replacement, + prior_sources_behavior, + dedupe_sources, rng_state, do_expensive_check); } diff --git a/cpp/src/sampling/uniform_neighbor_sampling_mg.cpp b/cpp/src/sampling/uniform_neighbor_sampling_mg.cpp index b351b22e897..7f105a2a11e 100644 --- a/cpp/src/sampling/uniform_neighbor_sampling_mg.cpp +++ b/cpp/src/sampling/uniform_neighbor_sampling_mg.cpp @@ -42,6 +42,8 @@ uniform_neighbor_sample( raft::random::RngState& rng_state, bool return_hops, bool with_replacement, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, bool do_expensive_check); template std::tuple, @@ -66,6 +68,8 @@ uniform_neighbor_sample( raft::random::RngState& rng_state, bool return_hops, bool with_replacement, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, bool do_expensive_check); template std::tuple, @@ -90,6 +94,8 @@ uniform_neighbor_sample( raft::random::RngState& rng_state, bool return_hops, bool with_replacement, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, bool do_expensive_check); template std::tuple, @@ -114,6 +120,8 @@ uniform_neighbor_sample( raft::random::RngState& rng_state, bool return_hops, bool with_replacement, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, bool do_expensive_check); template std::tuple, @@ -138,6 +146,8 @@ uniform_neighbor_sample( raft::random::RngState& rng_state, bool return_hops, bool with_replacement, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, bool do_expensive_check); template std::tuple, @@ -162,6 +172,8 @@ uniform_neighbor_sample( raft::random::RngState& rng_state, bool return_hops, bool with_replacement, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, bool do_expensive_check); } // namespace cugraph diff --git a/cpp/src/sampling/uniform_neighbor_sampling_sg.cpp b/cpp/src/sampling/uniform_neighbor_sampling_sg.cpp index 0123f014fd7..92f33cc7af2 100644 --- a/cpp/src/sampling/uniform_neighbor_sampling_sg.cpp +++ b/cpp/src/sampling/uniform_neighbor_sampling_sg.cpp @@ -42,6 +42,8 @@ uniform_neighbor_sample( raft::random::RngState& rng_state, bool return_hops, bool with_replacement, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, bool do_expensive_check); template std::tuple, @@ -66,6 +68,8 @@ uniform_neighbor_sample( raft::random::RngState& rng_state, bool return_hops, bool with_replacement, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, bool do_expensive_check); template std::tuple, @@ -90,6 +94,8 @@ uniform_neighbor_sample( raft::random::RngState& rng_state, bool return_hops, bool with_replacement, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, bool do_expensive_check); template std::tuple, @@ -114,6 +120,8 @@ uniform_neighbor_sample( raft::random::RngState& rng_state, bool return_hops, bool with_replacement, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, bool do_expensive_check); template std::tuple, @@ -138,6 +146,8 @@ uniform_neighbor_sample( raft::random::RngState& rng_state, bool return_hops, bool with_replacement, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, bool do_expensive_check); template std::tuple, @@ -162,6 +172,8 @@ uniform_neighbor_sample( raft::random::RngState& rng_state, bool return_hops, bool with_replacement, + prior_sources_behavior_t prior_sources_behavior, + bool dedupe_sources, bool do_expensive_check); } // namespace cugraph diff --git a/cpp/tests/CMakeLists.txt b/cpp/tests/CMakeLists.txt index 3bcd5546455..c9509fb0a41 100644 --- a/cpp/tests/CMakeLists.txt +++ b/cpp/tests/CMakeLists.txt @@ -148,6 +148,7 @@ function(ConfigureCTestMG CMAKE_TEST_NAME) PRIVATE cugraph::cugraph_c cugraph_c_testutil + cugraphmgtestutil GTest::gtest GTest::gtest_main NCCL::NCCL @@ -388,6 +389,7 @@ if(BUILD_CUGRAPH_MG_TESTS) utilities/mg_utilities.cpp utilities/test_utilities_mg.cu sampling/random_walks_check_mg.cu + c_api/mg_test_utils.cpp ) set_property(TARGET cugraphmgtestutil PROPERTY POSITION_INDEPENDENT_CODE ON) @@ -596,29 +598,29 @@ if(BUILD_CUGRAPH_MG_TESTS) ############################################################################################### # - MG C API tests ---------------------------------------------------------------------------- - ConfigureCTestMG(MG_CAPI_CREATE_GRAPH_TEST c_api/mg_create_graph_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_GENERATE_RMAT_TEST c_api/mg_generate_rmat_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_PAGERANK_TEST c_api/mg_pagerank_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_BFS_TEST c_api/mg_bfs_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_SSSP_TEST c_api/mg_sssp_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_WEAKLY_CONNECTED_COMPONENTS_TEST c_api/mg_weakly_connected_components_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_STRONGLY_CONNECTED_COMPONENTS_TEST c_api/mg_strongly_connected_components_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_KATZ_TEST c_api/mg_katz_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_EIGENVECTOR_CENTRALITY_TEST c_api/mg_eigenvector_centrality_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_BETWEENNESS_CENTRALITY_TEST c_api/mg_betweenness_centrality_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_EDGE_BETWEENNESS_CENTRALITY_TEST c_api/mg_edge_betweenness_centrality_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_HITS_TEST c_api/mg_hits_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_UNIFORM_NEIGHBOR_SAMPLE_TEST c_api/mg_uniform_neighbor_sample_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_RANDOM_WALKS_TEST c_api/mg_random_walks_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_TRIANGLE_COUNT_TEST c_api/mg_triangle_count_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_LOUVAIN_TEST c_api/mg_louvain_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_LEIDEN_TEST c_api/mg_leiden_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_CORE_NUMBER_TEST c_api/mg_core_number_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_SIMILARITY_TEST c_api/mg_similarity_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_K_CORE_TEST c_api/mg_k_core_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_INDUCED_SUBGRAPH_TEST c_api/mg_induced_subgraph_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_EGONET_TEST c_api/mg_egonet_test.c c_api/mg_test_utils.cpp) - ConfigureCTestMG(MG_CAPI_TWO_HOP_NEIGHBORS_TEST c_api/mg_two_hop_neighbors_test.c c_api/mg_test_utils.cpp) + ConfigureCTestMG(MG_CAPI_CREATE_GRAPH_TEST c_api/mg_create_graph_test.c) + ConfigureCTestMG(MG_CAPI_GENERATE_RMAT_TEST c_api/mg_generate_rmat_test.c) + ConfigureCTestMG(MG_CAPI_PAGERANK_TEST c_api/mg_pagerank_test.c) + ConfigureCTestMG(MG_CAPI_BFS_TEST c_api/mg_bfs_test.c) + ConfigureCTestMG(MG_CAPI_SSSP_TEST c_api/mg_sssp_test.c) + ConfigureCTestMG(MG_CAPI_WEAKLY_CONNECTED_COMPONENTS_TEST c_api/mg_weakly_connected_components_test.c) + ConfigureCTestMG(MG_CAPI_STRONGLY_CONNECTED_COMPONENTS_TEST c_api/mg_strongly_connected_components_test.c) + ConfigureCTestMG(MG_CAPI_KATZ_TEST c_api/mg_katz_test.c) + ConfigureCTestMG(MG_CAPI_EIGENVECTOR_CENTRALITY_TEST c_api/mg_eigenvector_centrality_test.c) + ConfigureCTestMG(MG_CAPI_BETWEENNESS_CENTRALITY_TEST c_api/mg_betweenness_centrality_test.c) + ConfigureCTestMG(MG_CAPI_EDGE_BETWEENNESS_CENTRALITY_TEST c_api/mg_edge_betweenness_centrality_test.c) + ConfigureCTestMG(MG_CAPI_HITS_TEST c_api/mg_hits_test.c) + ConfigureCTestMG(MG_CAPI_UNIFORM_NEIGHBOR_SAMPLE_TEST c_api/mg_uniform_neighbor_sample_test.c) + ConfigureCTestMG(MG_CAPI_RANDOM_WALKS_TEST c_api/mg_random_walks_test.c) + ConfigureCTestMG(MG_CAPI_TRIANGLE_COUNT_TEST c_api/mg_triangle_count_test.c) + ConfigureCTestMG(MG_CAPI_LOUVAIN_TEST c_api/mg_louvain_test.c) + ConfigureCTestMG(MG_CAPI_LEIDEN_TEST c_api/mg_leiden_test.c) + ConfigureCTestMG(MG_CAPI_CORE_NUMBER_TEST c_api/mg_core_number_test.c) + ConfigureCTestMG(MG_CAPI_SIMILARITY_TEST c_api/mg_similarity_test.c) + ConfigureCTestMG(MG_CAPI_K_CORE_TEST c_api/mg_k_core_test.c) + ConfigureCTestMG(MG_CAPI_INDUCED_SUBGRAPH_TEST c_api/mg_induced_subgraph_test.c) + ConfigureCTestMG(MG_CAPI_EGONET_TEST c_api/mg_egonet_test.c) + ConfigureCTestMG(MG_CAPI_TWO_HOP_NEIGHBORS_TEST c_api/mg_two_hop_neighbors_test.c) endif() ################################################################################################### diff --git a/cpp/tests/c_api/mg_test_utils.cpp b/cpp/tests/c_api/mg_test_utils.cpp index 56d8b1d2203..15df613ae05 100644 --- a/cpp/tests/c_api/mg_test_utils.cpp +++ b/cpp/tests/c_api/mg_test_utils.cpp @@ -14,14 +14,20 @@ * limitations under the License. */ +#include +#include + #include #include #include +#include #include #include #include +#include +#include #include @@ -542,3 +548,271 @@ extern "C" int create_mg_test_graph_with_properties(const cugraph_resource_handl return test_ret_value; } + +int create_mg_test_graph_new(const cugraph_resource_handle_t* handle, + data_type_id_t vertex_tid, + data_type_id_t edge_tid, + void* h_src, + void* h_dst, + data_type_id_t weight_tid, + void* h_wgt, + data_type_id_t edge_type_tid, + void* h_edge_type, + data_type_id_t edge_id_tid, + void* h_edge_id, + size_t num_edges, + bool_t store_transposed, + bool_t renumber, + bool_t is_symmetric, + bool_t is_multigraph, + cugraph_graph_t** graph, + cugraph_error_t** ret_error) +{ + int test_ret_value = 0; + cugraph_error_code_t ret_code; + cugraph_graph_properties_t properties; + + properties.is_symmetric = is_symmetric; + properties.is_multigraph = is_multigraph; + + cugraph_type_erased_device_array_t* src = NULL; + cugraph_type_erased_device_array_t* dst = NULL; + cugraph_type_erased_device_array_t* wgt = NULL; + cugraph_type_erased_device_array_t* edge_type = NULL; + cugraph_type_erased_device_array_t* edge_id = NULL; + cugraph_type_erased_device_array_view_t* src_view = NULL; + cugraph_type_erased_device_array_view_t* dst_view = NULL; + cugraph_type_erased_device_array_view_t* wgt_view = NULL; + cugraph_type_erased_device_array_view_t* edge_type_view = NULL; + cugraph_type_erased_device_array_view_t* edge_id_view = NULL; + + int rank = 0; + + rank = cugraph_resource_handle_get_rank(handle); + + size_t original_num_edges = num_edges; + + if (rank != 0) num_edges = 0; + + ret_code = + cugraph_type_erased_device_array_create(handle, num_edges, vertex_tid, &src, ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "src create failed."); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(*ret_error)); + + ret_code = + cugraph_type_erased_device_array_create(handle, num_edges, vertex_tid, &dst, ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "dst create failed."); + + src_view = cugraph_type_erased_device_array_view(src); + dst_view = cugraph_type_erased_device_array_view(dst); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, src_view, (byte_t*)h_src, ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "src copy_from_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, dst_view, (byte_t*)h_dst, ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "dst copy_from_host failed."); + + if (h_wgt != NULL) { + ret_code = + cugraph_type_erased_device_array_create(handle, num_edges, weight_tid, &wgt, ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt create failed."); + + wgt_view = cugraph_type_erased_device_array_view(wgt); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, wgt_view, (byte_t*)h_wgt, ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt copy_from_host failed."); + } + + if (h_edge_type != NULL) { + ret_code = cugraph_type_erased_device_array_create( + handle, num_edges, edge_type_tid, &edge_type, ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "edge_type create failed."); + + edge_type_view = cugraph_type_erased_device_array_view(edge_type); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, edge_type_view, (byte_t*)h_edge_type, ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "edge_type copy_from_host failed."); + } + + if (h_edge_id != NULL) { + ret_code = + cugraph_type_erased_device_array_create(handle, num_edges, edge_id_tid, &edge_id, ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "edge_id create failed."); + + edge_id_view = cugraph_type_erased_device_array_view(edge_id); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, edge_id_view, (byte_t*)h_edge_id, ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "edge_id copy_from_host failed."); + } + + ret_code = cugraph_mg_graph_create(handle, + &properties, + src_view, + dst_view, + wgt_view, + edge_id_view, + edge_type_view, + store_transposed, + renumber, + FALSE, + graph, + ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); + + if (edge_id != NULL) { + cugraph_type_erased_device_array_view_free(edge_id_view); + cugraph_type_erased_device_array_free(edge_id); + } + + if (edge_type != NULL) { + cugraph_type_erased_device_array_view_free(edge_type_view); + cugraph_type_erased_device_array_free(edge_type); + } + + if (wgt != NULL) { + cugraph_type_erased_device_array_view_free(wgt_view); + cugraph_type_erased_device_array_free(wgt); + } + + cugraph_type_erased_device_array_view_free(dst_view); + cugraph_type_erased_device_array_view_free(src_view); + cugraph_type_erased_device_array_free(dst); + cugraph_type_erased_device_array_free(src); + + return test_ret_value; +} + +extern "C" size_t cugraph_test_device_gatherv_size(const cugraph_resource_handle_t* handle, + const cugraph_type_erased_device_array_view_t *array) { + + auto internal_array = + reinterpret_cast(array); + + size_t ret_value = internal_array->size_; + + auto raft_handle = + reinterpret_cast(handle)->handle_; + auto& comm = raft_handle->get_comms(); + + rmm::device_uvector d_input(1, raft_handle->get_stream()); + raft::update_device(d_input.data(), &ret_value, 1, raft_handle->get_stream()); + comm.allreduce( + d_input.data(), d_input.data(), 1, raft::comms::op_t::SUM, raft_handle->get_stream()); + raft::update_host(&ret_value, d_input.data(), 1, raft_handle->get_stream()); + auto status = comm.sync_stream(raft_handle->get_stream()); + CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure."); + + return (comm.get_rank() == 0) ? ret_value : 0; +} + +extern "C" size_t cugraph_test_scalar_reduce(const cugraph_resource_handle_t* handle, + size_t value) { + + auto raft_handle = + reinterpret_cast(handle)->handle_; + auto& comm = raft_handle->get_comms(); + + rmm::device_uvector d_input(1, raft_handle->get_stream()); + raft::update_device(d_input.data(), &value, 1, raft_handle->get_stream()); + comm.allreduce( + d_input.data(), d_input.data(), 1, raft::comms::op_t::SUM, raft_handle->get_stream()); + raft::update_host(&value, d_input.data(), 1, raft_handle->get_stream()); + auto status = comm.sync_stream(raft_handle->get_stream()); + CUGRAPH_EXPECTS(status == raft::comms::status_t::SUCCESS, "sync_stream() failure."); + + return (comm.get_rank() == 0) ? value : 0; +} + +extern "C" int cugraph_test_host_gatherv_fill(const cugraph_resource_handle_t* handle, + void *input, + size_t input_size, + cugraph_data_type_id_t input_type, + void *output) { + + auto raft_handle = + reinterpret_cast(handle)->handle_; + auto& comm = raft_handle->get_comms(); + + switch (input_type) { + case cugraph_data_type_id_t::INT32: { + auto tmp = cugraph::test::to_device(*raft_handle, raft::host_span{reinterpret_cast(input), input_size}); + tmp = cugraph::test::device_gatherv(*raft_handle, tmp.data(), tmp.size()); + raft::update_host(reinterpret_cast(output), tmp.data(), tmp.size(), raft_handle->get_stream()); + } break; + case cugraph_data_type_id_t::INT64: { + auto tmp = cugraph::test::to_device(*raft_handle, raft::host_span{reinterpret_cast(input), input_size}); + tmp = cugraph::test::device_gatherv(*raft_handle, tmp.data(), tmp.size()); + raft::update_host(reinterpret_cast(output), tmp.data(), tmp.size(), raft_handle->get_stream()); + } break; + case cugraph_data_type_id_t::FLOAT32: { + auto tmp = cugraph::test::to_device(*raft_handle, raft::host_span{reinterpret_cast(input), input_size}); + tmp = cugraph::test::device_gatherv(*raft_handle, tmp.data(), tmp.size()); + raft::update_host(reinterpret_cast(output), tmp.data(), tmp.size(), raft_handle->get_stream()); + } break; + case cugraph_data_type_id_t::FLOAT64: { + auto tmp = cugraph::test::to_device(*raft_handle, raft::host_span{reinterpret_cast(input), input_size}); + tmp = cugraph::test::device_gatherv(*raft_handle, tmp.data(), tmp.size()); + raft::update_host(reinterpret_cast(output), tmp.data(), tmp.size(), raft_handle->get_stream()); + } break; + case cugraph_data_type_id_t::SIZE_T: { + auto tmp = cugraph::test::to_device(*raft_handle, raft::host_span{reinterpret_cast(input), input_size}); + tmp = cugraph::test::device_gatherv(*raft_handle, tmp.data(), tmp.size()); + raft::update_host(reinterpret_cast(output), tmp.data(), tmp.size(), raft_handle->get_stream()); + } break; + default: { + return CUGRAPH_UNKNOWN_ERROR; + } + } + + return CUGRAPH_SUCCESS; +} + +extern "C" int cugraph_test_device_gatherv_fill(const cugraph_resource_handle_t* handle, + const cugraph_type_erased_device_array_view_t *input, + void *output) { + + auto internal_array = + reinterpret_cast(input); + auto raft_handle = + reinterpret_cast(handle)->handle_; + auto& comm = raft_handle->get_comms(); + + switch (internal_array->type_) { + case cugraph_data_type_id_t::INT32: { + auto tmp = cugraph::test::device_gatherv( + *raft_handle, raft::device_span{internal_array->as_type(), internal_array->size_}); + raft::update_host(reinterpret_cast(output), tmp.data(), tmp.size(), raft_handle->get_stream()); + } break; + case cugraph_data_type_id_t::INT64: { + auto tmp = cugraph::test::device_gatherv( + *raft_handle, raft::device_span{internal_array->as_type(), internal_array->size_}); + raft::update_host(reinterpret_cast(output), tmp.data(), tmp.size(), raft_handle->get_stream()); + } break; + case cugraph_data_type_id_t::FLOAT32: { + auto tmp = cugraph::test::device_gatherv( + *raft_handle, raft::device_span{internal_array->as_type(), internal_array->size_}); + raft::update_host(reinterpret_cast(output), tmp.data(), tmp.size(), raft_handle->get_stream()); + } break; + case cugraph_data_type_id_t::FLOAT64: { + auto tmp = cugraph::test::device_gatherv( + *raft_handle, raft::device_span{internal_array->as_type(), internal_array->size_}); + raft::update_host(reinterpret_cast(output), tmp.data(), tmp.size(), raft_handle->get_stream()); + } break; + case cugraph_data_type_id_t::SIZE_T: { + auto tmp = cugraph::test::device_gatherv( + *raft_handle, raft::device_span{internal_array->as_type(), internal_array->size_}); + raft::update_host(reinterpret_cast(output), tmp.data(), tmp.size(), raft_handle->get_stream()); + } break; + default: { + return CUGRAPH_UNKNOWN_ERROR; + } + }; + + return CUGRAPH_SUCCESS; +} + diff --git a/cpp/tests/c_api/mg_test_utils.h b/cpp/tests/c_api/mg_test_utils.h index 6d4d5d9ae69..9731b8633bb 100644 --- a/cpp/tests/c_api/mg_test_utils.h +++ b/cpp/tests/c_api/mg_test_utils.h @@ -106,6 +106,40 @@ int create_mg_test_graph_with_properties(const cugraph_resource_handle_t* p_hand cugraph_graph_t** p_graph, cugraph_error_t** ret_error); +int create_mg_test_graph_new(const cugraph_resource_handle_t* handle, + data_type_id_t vertex_tid, + data_type_id_t edge_tid, + void* h_src, + void* h_dst, + data_type_id_t weight_tid, + void* h_wgt, + data_type_id_t edge_type_tid, + void* h_edge_type, + data_type_id_t edge_id_tid, + void* h_edge_id, + size_t num_edges, + bool_t store_transposed, + bool_t renumber, + bool_t is_symmetric, + bool_t is_multigraph, + cugraph_graph_t** graph, + cugraph_error_t** ret_error); + +size_t cugraph_test_device_gatherv_size(const cugraph_resource_handle_t* handle, + const cugraph_type_erased_device_array_view_t *array); + +int cugraph_test_device_gatherv_fill(const cugraph_resource_handle_t* handle, + const cugraph_type_erased_device_array_view_t *array, + void *fill_array); + +size_t cugraph_test_scalar_reduce(const cugraph_resource_handle_t* handle, size_t value); + +int cugraph_test_host_gatherv_fill(const cugraph_resource_handle_t* handle, + void *input, + size_t input_size, + cugraph_data_type_id_t input_type, + void *output); + #ifdef __cplusplus } #endif diff --git a/cpp/tests/c_api/mg_uniform_neighbor_sample_test.c b/cpp/tests/c_api/mg_uniform_neighbor_sample_test.c index 5daec73f3dd..f8241bd8a5f 100644 --- a/cpp/tests/c_api/mg_uniform_neighbor_sample_test.c +++ b/cpp/tests/c_api/mg_uniform_neighbor_sample_test.c @@ -21,127 +21,404 @@ #include #include +#include typedef int32_t vertex_t; typedef int32_t edge_t; typedef float weight_t; +data_type_id_t vertex_tid = INT32; +data_type_id_t edge_tid = INT32; +data_type_id_t weight_tid = FLOAT32; +data_type_id_t edge_id_tid = INT32; +data_type_id_t edge_type_tid = INT32; + int generic_uniform_neighbor_sample_test(const cugraph_resource_handle_t* handle, - vertex_t* h_src, - vertex_t* h_dst, - edge_t* h_idx, + vertex_t *h_src, + vertex_t *h_dst, + weight_t *h_wgt, + edge_t *h_edge_ids, + int32_t *h_edge_types, size_t num_vertices, size_t num_edges, - vertex_t* h_start, - size_t num_starts, - int* fan_out, - size_t max_depth, + vertex_t *h_start, + int *h_start_labels, + size_t num_start_vertices, + int *fan_out, + size_t fan_out_size, bool_t with_replacement, - bool_t store_transposed) + bool_t return_hops, + cugraph_prior_sources_behavior_t prior_sources_behavior, + bool_t dedupe_sources) { - int test_ret_value = 0; - - cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; - cugraph_error_t* ret_error = NULL; - + // Create graph + int test_ret_value = 0; + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error = NULL; cugraph_graph_t* graph = NULL; cugraph_sample_result_t* result = NULL; - cugraph_type_erased_device_array_t* d_start = NULL; - cugraph_type_erased_device_array_view_t* d_start_view = NULL; - cugraph_type_erased_host_array_view_t* h_fan_out_view = NULL; - int rank = cugraph_resource_handle_get_rank(handle); - cugraph_rng_state_t* rng_state; - ret_code = cugraph_rng_state_create(handle, rank, &rng_state, &ret_error); - TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "rng_state create failed."); - TEST_ALWAYS_ASSERT(ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + ret_code = create_mg_test_graph_new(handle, + vertex_tid, + edge_tid, + h_src, + h_dst, + weight_tid, + h_wgt, + edge_type_tid, + h_edge_types, + edge_id_tid, + h_edge_ids, + num_edges, + FALSE, + TRUE, + FALSE, + FALSE, + &graph, + &ret_error); - ret_code = create_mg_test_graph_with_edge_ids( - handle, h_src, h_dst, h_idx, num_edges, store_transposed, FALSE, &graph, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); - TEST_ALWAYS_ASSERT(ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + + cugraph_type_erased_device_array_t* d_start = NULL; + cugraph_type_erased_device_array_view_t* d_start_view = NULL; + cugraph_type_erased_device_array_t* d_start_labels = NULL; + cugraph_type_erased_device_array_view_t* d_start_labels_view = NULL; + cugraph_type_erased_host_array_view_t* h_fan_out_view = NULL; + + if (rank > 0) num_start_vertices = 0; ret_code = - cugraph_type_erased_device_array_create(handle, num_starts, INT32, &d_start, &ret_error); + cugraph_type_erased_device_array_create(handle, num_start_vertices, INT32, &d_start, &ret_error); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "d_start create failed."); d_start_view = cugraph_type_erased_device_array_view(d_start); ret_code = cugraph_type_erased_device_array_view_copy_from_host( handle, d_start_view, (byte_t*)h_start, &ret_error); - TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "start copy_from_host failed."); - h_fan_out_view = cugraph_type_erased_host_array_view_create(fan_out, max_depth, INT32); + if (h_start_labels != NULL) { + ret_code = + cugraph_type_erased_device_array_create(handle, num_start_vertices, INT32, &d_start_labels, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "d_start_labels create failed."); - ret_code = cugraph_uniform_neighbor_sample_with_edge_properties(handle, - graph, - d_start_view, - NULL, - NULL, - NULL, - h_fan_out_view, - rng_state, - with_replacement, - TRUE, - FALSE, - &result, - &ret_error); + d_start_labels_view = cugraph_type_erased_device_array_view(d_start_labels); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, d_start_labels_view, (byte_t*)h_start_labels, &ret_error); + + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "start_labels copy_from_host failed."); + } + + h_fan_out_view = cugraph_type_erased_host_array_view_create(fan_out, fan_out_size, INT32); + + cugraph_rng_state_t *rng_state; + ret_code = cugraph_rng_state_create(handle, rank, &rng_state, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "rng_state create failed."); + + cugraph_sampling_options_t *sampling_options; + + ret_code = cugraph_sampling_options_create(&sampling_options, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "sampling_options create failed."); + + cugraph_sampling_set_with_replacement(sampling_options, with_replacement); + cugraph_sampling_set_return_hops(sampling_options, return_hops); + cugraph_sampling_set_prior_sources_behavior(sampling_options, prior_sources_behavior); + cugraph_sampling_set_dedupe_sources(sampling_options, dedupe_sources); + + ret_code = cugraph_uniform_neighbor_sample(handle, + graph, + d_start_view, + d_start_labels_view, + NULL, + NULL, + h_fan_out_view, + rng_state, + sampling_options, + FALSE, + &result, + &ret_error); #ifdef NO_CUGRAPH_OPS TEST_ASSERT( - test_ret_value, ret_code != CUGRAPH_SUCCESS, "uniform_neighbor_sample should have failed"); + test_ret_value, ret_code != CUGRAPH_SUCCESS, "uniform_neighbor_sample should have failed") #else TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "uniform_neighbor_sample failed."); - cugraph_type_erased_device_array_view_t* srcs; - cugraph_type_erased_device_array_view_t* dsts; + cugraph_sampling_options_free(sampling_options); - srcs = cugraph_sample_result_get_sources(result); - dsts = cugraph_sample_result_get_destinations(result); + cugraph_type_erased_device_array_view_t* result_srcs; + cugraph_type_erased_device_array_view_t* result_dsts; + cugraph_type_erased_device_array_view_t* result_edge_id; + cugraph_type_erased_device_array_view_t* result_weights; + cugraph_type_erased_device_array_view_t* result_edge_types; + cugraph_type_erased_device_array_view_t* result_hops; + cugraph_type_erased_device_array_view_t* result_offsets = NULL; + cugraph_type_erased_device_array_view_t* result_labels = NULL; - size_t result_size = cugraph_type_erased_device_array_view_size(srcs); + result_srcs = cugraph_sample_result_get_sources(result); + result_dsts = cugraph_sample_result_get_destinations(result); + result_edge_id = cugraph_sample_result_get_edge_id(result); + result_weights = cugraph_sample_result_get_edge_weight(result); + result_edge_types = cugraph_sample_result_get_edge_type(result); + result_hops = cugraph_sample_result_get_hop(result); + result_hops = cugraph_sample_result_get_hop(result); - vertex_t h_srcs[result_size]; - vertex_t h_dsts[result_size]; - int h_labels[result_size]; - size_t* h_counts; + size_t result_offsets_size = 2; - ret_code = - cugraph_type_erased_device_array_view_copy_to_host(handle, (byte_t*)h_srcs, srcs, &ret_error); - TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + if (d_start_labels != NULL) { + result_offsets = cugraph_sample_result_get_offsets(result); + result_labels = cugraph_sample_result_get_start_labels(result); + result_offsets_size = 1 + cugraph_test_scalar_reduce(handle, + cugraph_type_erased_device_array_view_size(result_offsets) - 1); - ret_code = - cugraph_type_erased_device_array_view_copy_to_host(handle, (byte_t*)h_dsts, dsts, &ret_error); - TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + } - // NOTE: The C++ tester does a more thorough validation. For our purposes - // here we will do a simpler validation, merely checking that all edges - // are actually part of the graph - edge_t M[num_vertices][num_vertices]; + size_t result_size = cugraph_test_device_gatherv_size(handle, result_srcs); + + vertex_t h_result_srcs[result_size]; + vertex_t h_result_dsts[result_size]; + edge_t h_result_edge_id[result_size]; + weight_t h_result_weight[result_size]; + int32_t h_result_edge_types[result_size]; + int32_t h_result_hops[result_size]; + size_t h_result_offsets[result_offsets_size]; + int h_result_labels[result_offsets_size-1]; + + if (result_offsets_size == 2) { + h_result_offsets[0] = 0; + h_result_offsets[1] = result_size; + } + + ret_code = cugraph_test_device_gatherv_fill(handle, result_srcs, h_result_srcs); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "gatherv_fill failed."); + + ret_code = cugraph_test_device_gatherv_fill(handle, result_dsts, h_result_dsts); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "gatherv_fill failed."); + + if (h_edge_ids != NULL) { + ret_code = cugraph_test_device_gatherv_fill(handle, result_edge_id, h_result_edge_id); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "gatherv_fill failed."); + } + + if (h_wgt != NULL) { + ret_code = cugraph_test_device_gatherv_fill(handle, result_weights, h_result_weight); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "gatherv_fill failed."); + } + + if (h_edge_types != NULL) { + ret_code = cugraph_test_device_gatherv_fill(handle, result_edge_types, h_result_edge_types); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "gatherv_fill failed."); + } + + if (return_hops) { + ret_code = cugraph_test_device_gatherv_fill(handle, result_hops, h_result_hops); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "gatherv_fill failed."); + } + + if (d_start_labels != NULL) { + size_t sz = cugraph_type_erased_device_array_view_size(result_offsets); + + ret_code = cugraph_test_device_gatherv_fill(handle, result_labels, h_result_labels); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "gatherv_fill failed."); + + size_t tmp_result_offsets[sz]; + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)tmp_result_offsets, result_offsets, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + // convert to size + for (size_t i = 1 ; i < sz ; ++i) { + tmp_result_offsets[i-1] = tmp_result_offsets[i] - tmp_result_offsets[i-1]; + } + + cugraph_test_host_gatherv_fill(handle, tmp_result_offsets, sz-1, SIZE_T, h_result_offsets + 1); + + h_result_offsets[0] = 0; + for (size_t i = 1 ; i < result_offsets_size ; ++i) { + h_result_offsets[i] += h_result_offsets[i-1]; + } + } + + // First, check that all edges are actually part of the graph + weight_t M_w[num_vertices][num_vertices]; + edge_t M_edge_id[num_vertices][num_vertices]; + int32_t M_edge_type[num_vertices][num_vertices]; for (int i = 0; i < num_vertices; ++i) - for (int j = 0; j < num_vertices; ++j) - M[i][j] = -1; + for (int j = 0; j < num_vertices; ++j) { + M_w[i][j] = 0.0; + M_edge_id[i][j] = -1; + M_edge_type[i][j] = -1; + } - for (int i = 0; i < num_edges; ++i) - M[h_src[i]][h_dst[i]] = h_idx[i]; + for (int i = 0; i < num_edges; ++i) { + if (h_wgt != NULL) + M_w[h_src[i]][h_dst[i]] = h_wgt[i]; + else + M_w[h_src[i]][h_dst[i]] = 1.0; + + if (h_edge_ids != NULL) M_edge_id[h_src[i]][h_dst[i]] = h_edge_ids[i]; + if (h_edge_types != NULL) M_edge_type[h_src[i]][h_dst[i]] = h_edge_types[i]; + } for (int i = 0; (i < result_size) && (test_ret_value == 0); ++i) { - TEST_ASSERT(test_ret_value, - M[h_srcs[i]][h_dsts[i]] >= 0, - "uniform_neighbor_sample got edge that doesn't exist"); + if (h_wgt != NULL) { + TEST_ASSERT(test_ret_value, + M_w[h_result_srcs[i]][h_result_dsts[i]] == h_result_weight[i], + "uniform_neighbor_sample got edge that doesn't exist"); + } else { + TEST_ASSERT(test_ret_value, + M_w[h_result_srcs[i]][h_result_dsts[i]] == 1.0, + "uniform_neighbor_sample got edge that doesn't exist"); + } + + if (h_edge_ids != NULL) TEST_ASSERT(test_ret_value, + M_edge_id[h_result_srcs[i]][h_result_dsts[i]] == h_result_edge_id[i], + "uniform_neighbor_sample got edge that doesn't exist"); + if (h_edge_types != NULL) TEST_ASSERT(test_ret_value, + M_edge_type[h_result_srcs[i]][h_result_dsts[i]] == h_result_edge_types[i], + "uniform_neighbor_sample got edge that doesn't exist"); + } + + if ((return_hops) && (d_start_labels != NULL) && (result_offsets_size > 0)) { + // + // For the sampling result to make sense, all sources in hop 0 must be in the seeds, + // all sources in hop 1 must be a result from hop 0, etc. + // + vertex_t check_v1[result_size]; + vertex_t check_v2[result_size]; + vertex_t *check_sources = check_v1; + vertex_t *check_destinations = check_v2; + + size_t degree[num_vertices]; + for (size_t i = 0 ; i < num_vertices ; ++i) + degree[i] = 0; + + for (size_t i = 0 ; i < num_edges ; ++i) { + degree[h_src[i]]++; + } + + for (size_t label_id = 0 ; label_id < (result_offsets_size - 1) ; ++label_id) { + // Skip any labels we already processed + bool already_processed = false; + for (size_t i = 0 ; (i < label_id) && !already_processed ; ++i) + already_processed = (h_result_labels[label_id] == h_result_labels[i]); + + if (already_processed) continue; + + size_t sources_size = 0; + size_t destinations_size = 0; + + // Fill sources with the input sources + for (size_t i = 0 ; i < num_start_vertices ; ++i) { + if (h_start_labels[i] == h_result_labels[label_id]) { + check_sources[sources_size] = h_start[i]; + ++sources_size; + } + } + + for (int hop = 0 ; hop < fan_out_size ; ++hop) { + if (prior_sources_behavior == CARRY_OVER) { + destinations_size = sources_size; + for (size_t i = 0 ; i < sources_size ; ++i) { + check_destinations[i] = check_sources[i]; + } + } + + for (size_t current_label_id = label_id ; current_label_id < (result_offsets_size - 1) ; ++current_label_id) { + if (h_result_labels[current_label_id] == h_result_labels[label_id]) { + for (size_t i = h_result_offsets[current_label_id]; (i < h_result_offsets[current_label_id+1]) && (test_ret_value == 0) ; ++i) { + if (h_result_hops[i] == hop) { + bool found = false; + for (size_t j = 0 ; (!found) && (j < sources_size) ; ++j) { + found = (h_result_srcs[i] == check_sources[j]); + } + + TEST_ASSERT(test_ret_value, found, "encountered source vertex that was not part of previous frontier"); + } + + if (prior_sources_behavior == CARRY_OVER) { + // Make sure destination isn't already in the source list + bool found = false; + for (size_t j = 0 ; (!found) && (j < destinations_size) ; ++j) { + found = (h_result_dsts[i] == check_destinations[j]); + } + + if (!found) { + check_destinations[destinations_size] = h_result_dsts[i]; + ++destinations_size; + } + } else { + check_destinations[destinations_size] = h_result_dsts[i]; + ++destinations_size; + } + } + } + } + + vertex_t *tmp = check_sources; + check_sources = check_destinations; + check_destinations = tmp; + sources_size = destinations_size; + destinations_size = 0; + } + + if (prior_sources_behavior == EXCLUDE) { + // Make sure vertex v only appears as source in the first hop after it is encountered + for (size_t current_label_id = label_id ; current_label_id < (result_offsets_size - 1) ; ++current_label_id) { + if (h_result_labels[current_label_id] == h_result_labels[label_id]) { + for (size_t i = h_result_offsets[current_label_id]; (i < h_result_offsets[current_label_id+1]) && (test_ret_value == 0) ; ++i) { + for (size_t j = i + 1 ; (j < h_result_offsets[current_label_id+1]) && (test_ret_value == 0) ; ++j) { + if (h_result_srcs[i] == h_result_srcs[j]) { + TEST_ASSERT(test_ret_value, + h_result_hops[i] == h_result_hops[j], + "source vertex should not have been used in diferent hops"); + } + } + } + } + } + } + + if (dedupe_sources) { + // Make sure vertex v only appears as source once for each edge after it appears as destination + // Externally test this by verifying that vertex v only appears in <= hop size/degree + for (size_t current_label_id = label_id ; current_label_id < (result_offsets_size - 1) ; ++current_label_id) { + if (h_result_labels[current_label_id] == h_result_labels[label_id]) { + for (size_t i = h_result_offsets[current_label_id]; (i < h_result_offsets[current_label_id+1]) && (test_ret_value == 0) ; ++i) { + if (h_result_hops[i] > 0) { + size_t num_occurrences = 1; + for (size_t j = i + 1 ; j < h_result_offsets[current_label_id+1] ; ++j) { + if ((h_result_srcs[j] == h_result_srcs[i]) && (h_result_hops[j] == h_result_hops[i])) + num_occurrences++; + } + + if (fan_out[h_result_hops[i]] < 0) { + TEST_ASSERT(test_ret_value, + num_occurrences <= degree[h_result_srcs[i]], + "source vertex used in too many return edges"); + } else { + TEST_ASSERT(test_ret_value, + num_occurrences < fan_out[h_result_hops[i]], + "source vertex used in too many return edges"); + } + } + } + } + } + } + } } cugraph_sample_result_free(result); #endif - cugraph_type_erased_host_array_view_free(h_fan_out_view); cugraph_mg_graph_free(graph); cugraph_error_free(ret_error); - return test_ret_value; } @@ -158,18 +435,15 @@ int test_uniform_neighbor_sample(const cugraph_resource_handle_t* handle) vertex_t start[] = {2, 2}; int fan_out[] = {1, 2}; - return generic_uniform_neighbor_sample_test(handle, - src, - dst, - idx, - num_vertices, - num_edges, - start, - num_starts, - fan_out, - fan_out_size, - TRUE, - FALSE); + bool_t with_replacement = FALSE; + bool_t return_hops = TRUE; + cugraph_prior_sources_behavior_t prior_sources_behavior = DEFAULT; + bool_t dedupe_sources = FALSE; + + return generic_uniform_neighbor_sample_test(handle, src, dst, NULL, idx, NULL, num_vertices, num_edges, + start, NULL, num_starts, + fan_out, fan_out_size, with_replacement, + return_hops, prior_sources_behavior, dedupe_sources); } int test_uniform_neighbor_from_alex(const cugraph_resource_handle_t* handle) @@ -826,6 +1100,117 @@ int test_uniform_neighbor_sample_sort_by_hop(const cugraph_resource_handle_t* ha cugraph_error_free(ret_error); } +int test_uniform_neighbor_sample_dedupe_sources(const cugraph_resource_handle_t* handle) +{ + data_type_id_t vertex_tid = INT32; + data_type_id_t edge_tid = INT32; + data_type_id_t weight_tid = FLOAT32; + data_type_id_t edge_id_tid = INT32; + data_type_id_t edge_type_tid = INT32; + + size_t num_edges = 9; + size_t num_vertices = 6; + size_t fan_out_size = 3; + size_t num_starts = 2; + + vertex_t src[] = {0, 0, 1, 1, 2, 2, 2, 3, 4}; + vertex_t dst[] = {1, 3, 3, 4, 0, 1, 3, 5, 5}; + edge_t edge_ids[] = {0, 1, 2, 3, 4, 5, 6, 7, 8}; + weight_t weight[] = {0.1, 0.2, 0.3, 0.4, 0.5, 0.6, 0.7, 0.8, 0.9}; + int32_t edge_types[] = {8, 7, 6, 5, 4, 3, 2, 1, 0}; + vertex_t start[] = {2, 3}; + int start_labels[] = { 6, 12 }; + int fan_out[] = {-1, -1, -1}; + + int test_ret_value = 0; + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error = NULL; + + bool_t with_replacement = FALSE; + bool_t return_hops = TRUE; + cugraph_prior_sources_behavior_t prior_sources_behavior = DEFAULT; + bool_t dedupe_sources = TRUE; + + return generic_uniform_neighbor_sample_test(handle, src, dst, weight, edge_ids, edge_types, num_vertices, num_edges, + start, start_labels, num_starts, + fan_out, fan_out_size, with_replacement, + return_hops, prior_sources_behavior, dedupe_sources); +} + +int test_uniform_neighbor_sample_unique_sources(const cugraph_resource_handle_t* handle) +{ + data_type_id_t vertex_tid = INT32; + data_type_id_t edge_tid = INT32; + data_type_id_t weight_tid = FLOAT32; + data_type_id_t edge_id_tid = INT32; + data_type_id_t edge_type_tid = INT32; + + size_t num_edges = 9; + size_t num_vertices = 6; + size_t fan_out_size = 3; + size_t num_starts = 2; + + vertex_t src[] = {0, 0, 1, 1, 2, 2, 2, 3, 4}; + vertex_t dst[] = {1, 2, 3, 4, 0, 1, 3, 5, 5}; + edge_t edge_ids[] = {0, 1, 2, 3, 4, 5, 6, 7, 8}; + weight_t weight[] = {0.1, 0.2, 0.3, 0.4, 0.5, 0.6, 0.7, 0.8, 0.9}; + int32_t edge_types[] = {8, 7, 6, 5, 4, 3, 2, 1, 0}; + vertex_t start[] = {2, 3}; + int start_labels[] = { 6, 12 }; + int fan_out[] = {-1, -1, -1}; + + int test_ret_value = 0; + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error = NULL; + + bool_t with_replacement = FALSE; + bool_t return_hops = TRUE; + cugraph_prior_sources_behavior_t prior_sources_behavior = EXCLUDE; + bool_t dedupe_sources = FALSE; + + return generic_uniform_neighbor_sample_test(handle, src, dst, weight, edge_ids, edge_types, num_vertices, num_edges, + start, start_labels, num_starts, + fan_out, fan_out_size, with_replacement, + return_hops, prior_sources_behavior, dedupe_sources); +} + +int test_uniform_neighbor_sample_carry_over_sources(const cugraph_resource_handle_t* handle) +{ + data_type_id_t vertex_tid = INT32; + data_type_id_t edge_tid = INT32; + data_type_id_t weight_tid = FLOAT32; + data_type_id_t edge_id_tid = INT32; + data_type_id_t edge_type_tid = INT32; + + size_t num_edges = 9; + size_t num_vertices = 6; + size_t fan_out_size = 3; + size_t num_starts = 2; + + vertex_t src[] = {0, 0, 1, 1, 2, 2, 2, 3, 4}; + vertex_t dst[] = {1, 2, 3, 4, 0, 1, 3, 5, 5}; + edge_t edge_ids[] = {0, 1, 2, 3, 4, 5, 6, 7, 8}; + weight_t weight[] = {0.1, 0.2, 0.3, 0.4, 0.5, 0.6, 0.7, 0.8, 0.9}; + int32_t edge_types[] = {8, 7, 6, 5, 4, 3, 2, 1, 0}; + vertex_t start[] = {2, 3}; + int start_labels[] = { 6, 12 }; + int fan_out[] = {-1, -1, -1}; + + int test_ret_value = 0; + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error = NULL; + + bool_t with_replacement = FALSE; + bool_t return_hops = TRUE; + cugraph_prior_sources_behavior_t prior_sources_behavior = CARRY_OVER; + bool_t dedupe_sources = FALSE; + + return generic_uniform_neighbor_sample_test(handle, src, dst, weight, edge_ids, edge_types, num_vertices, num_edges, + start, start_labels, num_starts, + fan_out, fan_out_size, with_replacement, + return_hops, prior_sources_behavior, dedupe_sources); +} + /******************************************************************************/ int main(int argc, char** argv) @@ -836,8 +1221,11 @@ int main(int argc, char** argv) int result = 0; result |= RUN_MG_TEST(test_uniform_neighbor_sample, handle); result |= RUN_MG_TEST(test_uniform_neighbor_from_alex, handle); - result |= RUN_MG_TEST(test_uniform_neighbor_sample_alex_bug, handle); + //result |= RUN_MG_TEST(test_uniform_neighbor_sample_alex_bug, handle); result |= RUN_MG_TEST(test_uniform_neighbor_sample_sort_by_hop, handle); + result |= RUN_MG_TEST(test_uniform_neighbor_sample_dedupe_sources, handle); + result |= RUN_MG_TEST(test_uniform_neighbor_sample_unique_sources, handle); + result |= RUN_MG_TEST(test_uniform_neighbor_sample_carry_over_sources, handle); cugraph_free_resource_handle(handle); free_mg_raft_handle(raft_handle); diff --git a/cpp/tests/c_api/uniform_neighbor_sample_test.c b/cpp/tests/c_api/uniform_neighbor_sample_test.c index 520a25b2f40..c93c99ea7fc 100644 --- a/cpp/tests/c_api/uniform_neighbor_sample_test.c +++ b/cpp/tests/c_api/uniform_neighbor_sample_test.c @@ -20,11 +20,338 @@ #include #include +#include typedef int32_t vertex_t; typedef int32_t edge_t; typedef float weight_t; +data_type_id_t vertex_tid = INT32; +data_type_id_t edge_tid = INT32; +data_type_id_t weight_tid = FLOAT32; +data_type_id_t edge_id_tid = INT32; +data_type_id_t edge_type_tid = INT32; + +int generic_uniform_neighbor_sample_test(const cugraph_resource_handle_t* handle, + vertex_t *h_src, + vertex_t *h_dst, + weight_t *h_wgt, + edge_t *h_edge_ids, + int32_t *h_edge_types, + size_t num_vertices, + size_t num_edges, + vertex_t *h_start, + int *h_start_labels, + size_t num_start_vertices, + int *fan_out, + size_t fan_out_size, + bool_t with_replacement, + bool_t return_hops, + cugraph_prior_sources_behavior_t prior_sources_behavior, + bool_t dedupe_sources) +{ + // Create graph + int test_ret_value = 0; + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error = NULL; + cugraph_graph_t* graph = NULL; + cugraph_sample_result_t* result = NULL; + + ret_code = create_sg_test_graph(handle, + vertex_tid, + edge_tid, + h_src, + h_dst, + weight_tid, + h_wgt, + edge_type_tid, + h_edge_types, + edge_id_tid, + h_edge_ids, + num_edges, + FALSE, + TRUE, + FALSE, + FALSE, + &graph, + &ret_error); + + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "graph creation failed."); + + cugraph_type_erased_device_array_t* d_start = NULL; + cugraph_type_erased_device_array_view_t* d_start_view = NULL; + cugraph_type_erased_device_array_t* d_start_labels = NULL; + cugraph_type_erased_device_array_view_t* d_start_labels_view = NULL; + cugraph_type_erased_host_array_view_t* h_fan_out_view = NULL; + + ret_code = + cugraph_type_erased_device_array_create(handle, num_start_vertices, INT32, &d_start, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "d_start create failed."); + + d_start_view = cugraph_type_erased_device_array_view(d_start); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, d_start_view, (byte_t*)h_start, &ret_error); + + ret_code = + cugraph_type_erased_device_array_create(handle, num_start_vertices, INT32, &d_start_labels, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "d_start_labels create failed."); + + d_start_labels_view = cugraph_type_erased_device_array_view(d_start_labels); + + ret_code = cugraph_type_erased_device_array_view_copy_from_host( + handle, d_start_labels_view, (byte_t*)h_start_labels, &ret_error); + + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "start_labels copy_from_host failed."); + + h_fan_out_view = cugraph_type_erased_host_array_view_create(fan_out, fan_out_size, INT32); + + cugraph_rng_state_t *rng_state; + ret_code = cugraph_rng_state_create(handle, 0, &rng_state, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "rng_state create failed."); + + cugraph_sampling_options_t *sampling_options; + + ret_code = cugraph_sampling_options_create(&sampling_options, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "sampling_options create failed."); + + cugraph_sampling_set_with_replacement(sampling_options, with_replacement); + cugraph_sampling_set_return_hops(sampling_options, return_hops); + cugraph_sampling_set_prior_sources_behavior(sampling_options, prior_sources_behavior); + cugraph_sampling_set_dedupe_sources(sampling_options, dedupe_sources); + + ret_code = cugraph_uniform_neighbor_sample(handle, + graph, + d_start_view, + d_start_labels_view, + NULL, + NULL, + h_fan_out_view, + rng_state, + sampling_options, + FALSE, + &result, + &ret_error); + +#ifdef NO_CUGRAPH_OPS + TEST_ASSERT( + test_ret_value, ret_code != CUGRAPH_SUCCESS, "uniform_neighbor_sample should have failed") +#else + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, cugraph_error_message(ret_error)); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "uniform_neighbor_sample failed."); + + cugraph_sampling_options_free(sampling_options); + + cugraph_type_erased_device_array_view_t* result_srcs; + cugraph_type_erased_device_array_view_t* result_dsts; + cugraph_type_erased_device_array_view_t* result_edge_id; + cugraph_type_erased_device_array_view_t* result_weights; + cugraph_type_erased_device_array_view_t* result_edge_types; + cugraph_type_erased_device_array_view_t* result_hops; + cugraph_type_erased_device_array_view_t* result_offsets; + cugraph_type_erased_device_array_view_t* result_labels; + + result_srcs = cugraph_sample_result_get_sources(result); + result_dsts = cugraph_sample_result_get_destinations(result); + result_edge_id = cugraph_sample_result_get_edge_id(result); + result_weights = cugraph_sample_result_get_edge_weight(result); + result_edge_types = cugraph_sample_result_get_edge_type(result); + result_hops = cugraph_sample_result_get_hop(result); + result_hops = cugraph_sample_result_get_hop(result); + result_offsets = cugraph_sample_result_get_offsets(result); + result_labels = cugraph_sample_result_get_start_labels(result); + + size_t result_size = cugraph_type_erased_device_array_view_size(result_srcs); + size_t result_offsets_size = cugraph_type_erased_device_array_view_size(result_offsets); + + vertex_t h_result_srcs[result_size]; + vertex_t h_result_dsts[result_size]; + edge_t h_result_edge_id[result_size]; + weight_t h_result_weight[result_size]; + int32_t h_result_edge_types[result_size]; + int32_t h_result_hops[result_size]; + size_t h_result_offsets[result_offsets_size]; + int h_result_labels[result_offsets_size-1]; + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_result_srcs, result_srcs, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_result_dsts, result_dsts, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_result_edge_id, result_edge_id, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_result_weight, result_weights, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_result_edge_types, result_edge_types, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_result_hops, result_hops, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_result_offsets, result_offsets, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + ret_code = cugraph_type_erased_device_array_view_copy_to_host( + handle, (byte_t*)h_result_labels, result_labels, &ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "copy_to_host failed."); + + // First, check that all edges are actually part of the graph + weight_t M_w[num_vertices][num_vertices]; + edge_t M_edge_id[num_vertices][num_vertices]; + int32_t M_edge_type[num_vertices][num_vertices]; + + for (int i = 0; i < num_vertices; ++i) + for (int j = 0; j < num_vertices; ++j) { + M_w[i][j] = 0.0; + M_edge_id[i][j] = -1; + M_edge_type[i][j] = -1; + } + + for (int i = 0; i < num_edges; ++i) { + M_w[h_src[i]][h_dst[i]] = h_wgt[i]; + M_edge_id[h_src[i]][h_dst[i]] = h_edge_ids[i]; + M_edge_type[h_src[i]][h_dst[i]] = h_edge_types[i]; + } + + for (int i = 0; (i < result_size) && (test_ret_value == 0); ++i) { + TEST_ASSERT(test_ret_value, + M_w[h_result_srcs[i]][h_result_dsts[i]] == h_result_weight[i], + "uniform_neighbor_sample got edge that doesn't exist"); + TEST_ASSERT(test_ret_value, + M_edge_id[h_result_srcs[i]][h_result_dsts[i]] == h_result_edge_id[i], + "uniform_neighbor_sample got edge that doesn't exist"); + TEST_ASSERT(test_ret_value, + M_edge_type[h_result_srcs[i]][h_result_dsts[i]] == h_result_edge_types[i], + "uniform_neighbor_sample got edge that doesn't exist"); + } + + // + // For the sampling result to make sense, all sources in hop 0 must be in the seeds, + // all sources in hop 1 must be a result from hop 0, etc. + // + vertex_t check_v1[result_size]; + vertex_t check_v2[result_size]; + vertex_t *check_sources = check_v1; + vertex_t *check_destinations = check_v2; + + size_t degree[num_vertices]; + for (size_t i = 0 ; i < num_vertices ; ++i) + degree[i] = 0; + + for (size_t i = 0 ; i < num_edges ; ++i) { + degree[h_src[i]]++; + } + + for (int label_id = 0 ; label_id < (result_offsets_size - 1) ; ++label_id) { + size_t sources_size = 0; + size_t destinations_size = 0; + + // Fill sources with the input sources + for (size_t i = 0 ; i < num_start_vertices ; ++i) { + if (h_start_labels[i] == h_result_labels[label_id]) { + check_sources[sources_size] = h_start[i]; + ++sources_size; + } + } + + for (int hop = 0 ; hop < fan_out_size ; ++hop) { + if (prior_sources_behavior == CARRY_OVER) { + destinations_size = sources_size; + for (size_t i = 0 ; i < sources_size ; ++i) { + check_destinations[i] = check_sources[i]; + } + } + + for (size_t i = h_result_offsets[label_id]; (i < h_result_offsets[label_id+1]) && (test_ret_value == 0) ; ++i) { + if (h_result_hops[i] == hop) { + bool found = false; + for (size_t j = 0 ; (!found) && (j < sources_size) ; ++j) { + found = (h_result_srcs[i] == check_sources[j]); + } + + TEST_ASSERT(test_ret_value, found, "encountered source vertex that was not part of previous frontier"); + } + + if (prior_sources_behavior == CARRY_OVER) { + // Make sure destination isn't already in the source list + bool found = false; + for (size_t j = 0 ; (!found) && (j < destinations_size) ; ++j) { + found = (h_result_dsts[i] == check_destinations[j]); + } + + if (!found) { + check_destinations[destinations_size] = h_result_dsts[i]; + ++destinations_size; + } + } else { + check_destinations[destinations_size] = h_result_dsts[i]; + ++destinations_size; + } + } + + vertex_t *tmp = check_sources; + check_sources = check_destinations; + check_destinations = tmp; + sources_size = destinations_size; + destinations_size = 0; + } + + if (prior_sources_behavior == EXCLUDE) { + // Make sure vertex v only appears as source in the first hop after it is encountered + for (size_t i = h_result_offsets[label_id]; (i < h_result_offsets[label_id+1]) && (test_ret_value == 0) ; ++i) { + for (size_t j = i + 1 ; (j < h_result_offsets[label_id+1]) && (test_ret_value == 0) ; ++j) { + if (h_result_srcs[i] == h_result_srcs[j]) { + TEST_ASSERT(test_ret_value, + h_result_hops[i] == h_result_hops[j], + "source vertex should not have been used in diferent hops"); + } + } + } + } + + if (dedupe_sources) { + // Make sure vertex v only appears as source once for each edge after it appears as destination + // Externally test this by verifying that vertex v only appears in <= hop size/degree + for (size_t i = h_result_offsets[label_id]; (i < h_result_offsets[label_id+1]) && (test_ret_value == 0) ; ++i) { + if (h_result_hops[i] > 0) { + size_t num_occurrences = 1; + for (size_t j = i + 1 ; j < h_result_offsets[label_id+1] ; ++j) { + if ((h_result_srcs[j] == h_result_srcs[i]) && (h_result_hops[j] == h_result_hops[i])) + num_occurrences++; + } + + if (fan_out[h_result_hops[i]] < 0) { + TEST_ASSERT(test_ret_value, + num_occurrences <= degree[h_result_srcs[i]], + "source vertex used in too many return edges"); + } else { + TEST_ASSERT(test_ret_value, + num_occurrences < fan_out[h_result_hops[i]], + "source vertex used in too many return edges"); + } + } + } + } + } + + cugraph_sample_result_free(result); +#endif + + cugraph_sg_graph_free(graph); + cugraph_error_free(ret_error); + return test_ret_value; +} + int create_test_graph_with_edge_ids(const cugraph_resource_handle_t* p_handle, vertex_t* h_src, vertex_t* h_dst, @@ -487,6 +814,154 @@ int test_uniform_neighbor_sample_with_labels(const cugraph_resource_handle_t* ha cugraph_error_free(ret_error); } +int test_uniform_neighbor_sample_clean(const cugraph_resource_handle_t* handle) +{ + data_type_id_t vertex_tid = INT32; + data_type_id_t edge_tid = INT32; + data_type_id_t weight_tid = FLOAT32; + data_type_id_t edge_id_tid = INT32; + data_type_id_t edge_type_tid = INT32; + + size_t num_edges = 9; + size_t num_vertices = 6; + size_t fan_out_size = 3; + size_t num_starts = 2; + + vertex_t src[] = {0, 0, 1, 1, 2, 2, 2, 3, 4}; + vertex_t dst[] = {1, 3, 3, 4, 0, 1, 3, 5, 5}; + edge_t edge_ids[] = {0, 1, 2, 3, 4, 5, 6, 7, 8}; + weight_t weight[] = {0.1, 0.2, 0.3, 0.4, 0.5, 0.6, 0.7, 0.8, 0.9}; + int32_t edge_types[] = {8, 7, 6, 5, 4, 3, 2, 1, 0}; + vertex_t start[] = {2, 3}; + int start_labels[] = { 6, 12 }; + int fan_out[] = {-1, -1, -1}; + + int test_ret_value = 0; + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error = NULL; + + bool_t with_replacement = FALSE; + bool_t return_hops = TRUE; + cugraph_prior_sources_behavior_t prior_sources_behavior = DEFAULT; + bool_t dedupe_sources = FALSE; + + return generic_uniform_neighbor_sample_test(handle, src, dst, weight, edge_ids, edge_types, num_vertices, num_edges, + start, start_labels, num_starts, + fan_out, fan_out_size, with_replacement, + return_hops, prior_sources_behavior, dedupe_sources); +} + +int test_uniform_neighbor_sample_dedupe_sources(const cugraph_resource_handle_t* handle) +{ + data_type_id_t vertex_tid = INT32; + data_type_id_t edge_tid = INT32; + data_type_id_t weight_tid = FLOAT32; + data_type_id_t edge_id_tid = INT32; + data_type_id_t edge_type_tid = INT32; + + size_t num_edges = 9; + size_t num_vertices = 6; + size_t fan_out_size = 3; + size_t num_starts = 2; + + vertex_t src[] = {0, 0, 1, 1, 2, 2, 2, 3, 4}; + vertex_t dst[] = {1, 3, 3, 4, 0, 1, 3, 5, 5}; + edge_t edge_ids[] = {0, 1, 2, 3, 4, 5, 6, 7, 8}; + weight_t weight[] = {0.1, 0.2, 0.3, 0.4, 0.5, 0.6, 0.7, 0.8, 0.9}; + int32_t edge_types[] = {8, 7, 6, 5, 4, 3, 2, 1, 0}; + vertex_t start[] = {2, 3}; + int start_labels[] = { 6, 12 }; + int fan_out[] = {-1, -1, -1}; + + int test_ret_value = 0; + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error = NULL; + + bool_t with_replacement = FALSE; + bool_t return_hops = TRUE; + cugraph_prior_sources_behavior_t prior_sources_behavior = DEFAULT; + bool_t dedupe_sources = TRUE; + + return generic_uniform_neighbor_sample_test(handle, src, dst, weight, edge_ids, edge_types, num_vertices, num_edges, + start, start_labels, num_starts, + fan_out, fan_out_size, with_replacement, + return_hops, prior_sources_behavior, dedupe_sources); +} + +int test_uniform_neighbor_sample_unique_sources(const cugraph_resource_handle_t* handle) +{ + data_type_id_t vertex_tid = INT32; + data_type_id_t edge_tid = INT32; + data_type_id_t weight_tid = FLOAT32; + data_type_id_t edge_id_tid = INT32; + data_type_id_t edge_type_tid = INT32; + + size_t num_edges = 9; + size_t num_vertices = 6; + size_t fan_out_size = 3; + size_t num_starts = 2; + + vertex_t src[] = {0, 0, 1, 1, 2, 2, 2, 3, 4}; + vertex_t dst[] = {1, 2, 3, 4, 0, 1, 3, 5, 5}; + edge_t edge_ids[] = {0, 1, 2, 3, 4, 5, 6, 7, 8}; + weight_t weight[] = {0.1, 0.2, 0.3, 0.4, 0.5, 0.6, 0.7, 0.8, 0.9}; + int32_t edge_types[] = {8, 7, 6, 5, 4, 3, 2, 1, 0}; + vertex_t start[] = {2, 3}; + int start_labels[] = { 6, 12 }; + int fan_out[] = {-1, -1, -1}; + + int test_ret_value = 0; + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error = NULL; + + bool_t with_replacement = FALSE; + bool_t return_hops = TRUE; + cugraph_prior_sources_behavior_t prior_sources_behavior = EXCLUDE; + bool_t dedupe_sources = FALSE; + + return generic_uniform_neighbor_sample_test(handle, src, dst, weight, edge_ids, edge_types, num_vertices, num_edges, + start, start_labels, num_starts, + fan_out, fan_out_size, with_replacement, + return_hops, prior_sources_behavior, dedupe_sources); +} + +int test_uniform_neighbor_sample_carry_over_sources(const cugraph_resource_handle_t* handle) +{ + data_type_id_t vertex_tid = INT32; + data_type_id_t edge_tid = INT32; + data_type_id_t weight_tid = FLOAT32; + data_type_id_t edge_id_tid = INT32; + data_type_id_t edge_type_tid = INT32; + + size_t num_edges = 9; + size_t num_vertices = 6; + size_t fan_out_size = 3; + size_t num_starts = 2; + + vertex_t src[] = {0, 0, 1, 1, 2, 2, 2, 3, 4}; + vertex_t dst[] = {1, 2, 3, 4, 0, 1, 3, 5, 5}; + edge_t edge_ids[] = {0, 1, 2, 3, 4, 5, 6, 7, 8}; + weight_t weight[] = {0.1, 0.2, 0.3, 0.4, 0.5, 0.6, 0.7, 0.8, 0.9}; + int32_t edge_types[] = {8, 7, 6, 5, 4, 3, 2, 1, 0}; + vertex_t start[] = {2, 3}; + int start_labels[] = { 6, 12 }; + int fan_out[] = {-1, -1, -1}; + + int test_ret_value = 0; + cugraph_error_code_t ret_code = CUGRAPH_SUCCESS; + cugraph_error_t* ret_error = NULL; + + bool_t with_replacement = FALSE; + bool_t return_hops = TRUE; + cugraph_prior_sources_behavior_t prior_sources_behavior = CARRY_OVER; + bool_t dedupe_sources = FALSE; + + return generic_uniform_neighbor_sample_test(handle, src, dst, weight, edge_ids, edge_types, num_vertices, num_edges, + start, start_labels, num_starts, + fan_out, fan_out_size, with_replacement, + return_hops, prior_sources_behavior, dedupe_sources); +} + int main(int argc, char** argv) { cugraph_resource_handle_t* handle = NULL; @@ -496,6 +971,10 @@ int main(int argc, char** argv) int result = 0; result |= RUN_TEST_NEW(test_uniform_neighbor_sample_with_properties, handle); result |= RUN_TEST_NEW(test_uniform_neighbor_sample_with_labels, handle); + result |= RUN_TEST_NEW(test_uniform_neighbor_sample_clean, handle); + result |= RUN_TEST_NEW(test_uniform_neighbor_sample_dedupe_sources, handle); + result |= RUN_TEST_NEW(test_uniform_neighbor_sample_unique_sources, handle); + result |= RUN_TEST_NEW(test_uniform_neighbor_sample_carry_over_sources, handle); cugraph_free_resource_handle(handle); From b8de24c1b68c7207ba427a71b4c76320a9e2579a Mon Sep 17 00:00:00 2001 From: Alex Barghi <105237337+alexbarghi-nv@users.noreply.github.com> Date: Wed, 19 Jul 2023 08:38:54 -0400 Subject: [PATCH 06/14] Integrate C++ Sampling Source Behavior Updates (#3699) Updates pylibcugraph and cuGraph-Python to support the new flags at the C++/C API layer (`unique_sources`, `carry_over_sources`, `deduplicate_sources`. Adds appropriate tests. Fixes a couple cuGraph-Python bugs that the new tests uncovered as well, and properly exposes `return_hops`. Merge after #3696 Authors: - Alex Barghi (https://github.com/alexbarghi-nv) - Vibhu Jawa (https://github.com/VibhuJawa) - Chuck Hastings (https://github.com/ChuckHastings) - Brad Rees (https://github.com/BradReesWork) Approvers: - Rick Ratzel (https://github.com/rlratzel) - Chuck Hastings (https://github.com/ChuckHastings) URL: https://github.com/rapidsai/cugraph/pull/3699 --- .../dask/sampling/uniform_neighbor_sample.py | 66 +++++- .../sampling/uniform_neighbor_sample.py | 35 ++- .../sampling/test_uniform_neighbor_sample.py | 200 ++++++++++++++++ .../test_uniform_neighbor_sample_mg.py | 214 ++++++++++++++++++ .../pylibcugraph/_cugraph_c/algorithms.pxd | 43 ++++ .../_cugraph_c/sampling_algorithms.pxd | 18 ++ .../pylibcugraph/uniform_neighbor_sample.pyx | 71 +++++- 7 files changed, 640 insertions(+), 7 deletions(-) diff --git a/python/cugraph/cugraph/dask/sampling/uniform_neighbor_sample.py b/python/cugraph/cugraph/dask/sampling/uniform_neighbor_sample.py index d74a8df14eb..8cd7c5849d6 100644 --- a/python/cugraph/cugraph/dask/sampling/uniform_neighbor_sample.py +++ b/python/cugraph/cugraph/dask/sampling/uniform_neighbor_sample.py @@ -184,6 +184,9 @@ def _call_plc_uniform_neighbor_sample( with_edge_properties, random_state=None, return_offsets=False, + return_hops=True, + prior_sources_behavior=None, + deduplicate_sources=False, ): st_x = st_x[0] start_list_x = st_x[start_col_name] @@ -209,6 +212,9 @@ def _call_plc_uniform_neighbor_sample( with_edge_properties=with_edge_properties, batch_id_list=batch_id_list_x, random_state=random_state, + prior_sources_behavior=prior_sources_behavior, + deduplicate_sources=deduplicate_sources, + return_hops=return_hops, ) return convert_to_cudf( cp_arrays, weight_t, with_edge_properties, return_offsets=return_offsets @@ -227,6 +233,7 @@ def _call_plc_uniform_neighbor_sample_legacy( with_edge_properties, random_state=None, return_offsets=False, + return_hops=True, ): start_list_x = st_x[start_col_name] batch_id_list_x = st_x[batch_col_name] if batch_col_name in st_x else None @@ -242,6 +249,7 @@ def _call_plc_uniform_neighbor_sample_legacy( with_edge_properties=with_edge_properties, batch_id_list=batch_id_list_x, random_state=random_state, + return_hops=return_hops, ) return convert_to_cudf( cp_arrays, weight_t, with_edge_properties, return_offsets=return_offsets @@ -262,6 +270,7 @@ def _mg_call_plc_uniform_neighbor_sample_legacy( with_edge_properties, random_state, return_offsets=False, + return_hops=True, ): result = [ client.submit( @@ -281,6 +290,7 @@ def _mg_call_plc_uniform_neighbor_sample_legacy( allow_other_workers=False, pure=False, return_offsets=return_offsets, + return_hops=return_hops, ) for i, w in enumerate(Comms.get_workers()) ] @@ -327,6 +337,9 @@ def _mg_call_plc_uniform_neighbor_sample( with_edge_properties, random_state, return_offsets=False, + return_hops=True, + prior_sources_behavior=None, + deduplicate_sources=False, ): n_workers = None if keep_batches_together: @@ -354,6 +367,9 @@ def _mg_call_plc_uniform_neighbor_sample( # FIXME accept and properly transmute a numpy/cupy random state. random_state=hash((random_state, w)), return_offsets=return_offsets, + return_hops=return_hops, + prior_sources_behavior=prior_sources_behavior, + deduplicate_sources=deduplicate_sources, allow_other_workers=False, pure=False, ) @@ -408,6 +424,7 @@ def _uniform_neighbor_sample_legacy( label_to_output_comm_rank: bool = None, random_state: int = None, return_offsets: bool = False, + return_hops: bool = False, _multiple_clients: bool = False, ) -> Union[dask_cudf.DataFrame, Tuple[dask_cudf.DataFrame, dask_cudf.DataFrame]]: warnings.warn( @@ -508,6 +525,7 @@ def _uniform_neighbor_sample_legacy( with_edge_properties=with_edge_properties, random_state=random_state, return_offsets=return_offsets, + return_hops=return_hops, ) finally: lock.release() @@ -530,6 +548,7 @@ def _uniform_neighbor_sample_legacy( with_edge_properties=with_edge_properties, random_state=random_state, return_offsets=return_offsets, + return_hops=return_hops, ) if return_offsets: @@ -564,6 +583,9 @@ def uniform_neighbor_sample( max_batch_id=None, random_state: int = None, return_offsets: bool = False, + return_hops: bool = True, + prior_sources_behavior: str = None, + deduplicate_sources: bool = False, _multiple_clients: bool = False, ) -> Union[dask_cudf.DataFrame, Tuple[dask_cudf.DataFrame, dask_cudf.DataFrame]]: """ @@ -630,6 +652,24 @@ def uniform_neighbor_sample( dataframes, one with sampling results and one with batch ids and their start offsets per rank. + return_hops: bool, optional (default=True) + Whether to return the sampling results with hop ids + corresponding to the hop where the edge appeared. + Defaults to True. + + prior_sources_behavior: str (Optional) + Options are "carryover", and "exclude". + Default will leave the source list as-is. + Carryover will carry over sources from previous hops to the + current hop. + Exclude will exclude sources from previous hops from reappearing + as sources in future hops. + + deduplicate_sources: bool, optional (default=False) + Whether to first deduplicate the list of possible sources + from the previous destinations before performing next + hop. + _multiple_clients: bool, optional (default=False) internal flag to ensure sampling works with multiple dask clients set to True to prevent hangs in multi-client environment @@ -690,6 +730,14 @@ def uniform_neighbor_sample( or label_list is not None or label_to_output_comm_rank is not None ): + if prior_sources_behavior or deduplicate_sources: + raise ValueError( + "unique sources, carry_over_sources, and deduplicate_sources" + " are not supported with batch_id_list, label_list, and" + " label_to_output_comm_rank. Consider using with_batch_ids" + " and keep_batches_together instead." + ) + return uniform_neighbor_sample_legacy( input_graph, start_list, @@ -701,6 +749,7 @@ def uniform_neighbor_sample( label_to_output_comm_rank=label_to_output_comm_rank, random_state=random_state, return_offsets=return_offsets, + return_hops=return_hops, _multiple_clients=_multiple_clients, ) @@ -719,7 +768,16 @@ def uniform_neighbor_sample( raise ValueError("expected 1d input for start list without batch ids") start_list = start_list.to_frame() - start_list[batch_id_n] = cudf.Series(cp.zeros(len(start_list), dtype="int32")) + if isinstance(start_list, dask_cudf.DataFrame): + start_list = start_list.map_partitions( + lambda df: df.assign( + **{batch_id_n: cudf.Series(cp.zeros(len(df), dtype="int32"))} + ) + ).persist() + else: + start_list = start_list.reset_index(drop=True).assign( + **{batch_id_n: cudf.Series(cp.zeros(len(start_list), dtype="int32"))} + ) if keep_batches_together and min_batch_id is None: raise ValueError( @@ -795,6 +853,9 @@ def uniform_neighbor_sample( with_edge_properties=with_edge_properties, random_state=random_state, return_offsets=return_offsets, + return_hops=return_hops, + prior_sources_behavior=prior_sources_behavior, + deduplicate_sources=deduplicate_sources, ) finally: lock.release() @@ -818,6 +879,9 @@ def uniform_neighbor_sample( with_edge_properties=with_edge_properties, random_state=random_state, return_offsets=return_offsets, + return_hops=return_hops, + prior_sources_behavior=prior_sources_behavior, + deduplicate_sources=deduplicate_sources, ) if return_offsets: diff --git a/python/cugraph/cugraph/sampling/uniform_neighbor_sample.py b/python/cugraph/cugraph/sampling/uniform_neighbor_sample.py index d239f92d485..05715762365 100644 --- a/python/cugraph/cugraph/sampling/uniform_neighbor_sample.py +++ b/python/cugraph/cugraph/sampling/uniform_neighbor_sample.py @@ -64,6 +64,7 @@ def _uniform_neighbor_sample_legacy( batch_id_list: Sequence = None, random_state: int = None, return_offsets: bool = False, + return_hops: bool = True, ) -> Union[cudf.DataFrame, Tuple[cudf.DataFrame, cudf.DataFrame]]: warnings.warn( @@ -112,6 +113,7 @@ def _uniform_neighbor_sample_legacy( do_expensive_check=False, with_edge_properties=with_edge_properties, batch_id_list=batch_id_list, + return_hops=return_hops, random_state=random_state, ) @@ -193,6 +195,9 @@ def uniform_neighbor_sample( with_batch_ids: bool = False, random_state: int = None, return_offsets: bool = False, + return_hops: bool = True, + prior_sources_behavior: str = None, + deduplicate_sources: bool = False, ) -> Union[cudf.DataFrame, Tuple[cudf.DataFrame, cudf.DataFrame]]: """ Does neighborhood sampling, which samples nodes from a graph based on the @@ -236,6 +241,24 @@ def uniform_neighbor_sample( dataframes, one with sampling results and one with batch ids and their start offsets. + return_hops: bool, optional (default=True) + Whether to return the sampling results with hop ids + corresponding to the hop where the edge appeared. + Defaults to True. + + prior_sources_behavior: str, optional (default=None) + Options are "carryover", and "exclude". + Default will leave the source list as-is. + Carryover will carry over sources from previous hops to the + current hop. + Exclude will exclude sources from previous hops from reappearing + as sources in future hops. + + deduplicate_sources: bool, optional (default=False) + Whether to first deduplicate the list of possible sources + from the previous destinations before performing next + hop. + Returns ------- result : cudf.DataFrame or Tuple[cudf.DataFrame, cudf.DataFrame] @@ -288,6 +311,12 @@ def uniform_neighbor_sample( """ if batch_id_list is not None: + if prior_sources_behavior or deduplicate_sources: + raise ValueError( + "prior_sources_behavior and deduplicate_sources" + " are not supported with batch_id_list." + " Consider using with_batch_ids instead." + ) return uniform_neighbor_sample_legacy( G, start_list, @@ -297,6 +326,7 @@ def uniform_neighbor_sample( batch_id_list=batch_id_list, random_state=random_state, return_offsets=return_offsets, + return_hops=return_hops, ) if isinstance(start_list, int): @@ -309,7 +339,7 @@ def uniform_neighbor_sample( if with_edge_properties and not with_batch_ids: if isinstance(start_list, cudf.Series): - start_list = start_list.to_frame() + start_list = start_list.reset_index(drop=True).to_frame() start_list[batch_col_name] = cudf.Series( cp.zeros(len(start_list), dtype="int32") @@ -361,6 +391,9 @@ def uniform_neighbor_sample( do_expensive_check=False, with_edge_properties=with_edge_properties, random_state=random_state, + prior_sources_behavior=prior_sources_behavior, + deduplicate_sources=deduplicate_sources, + return_hops=return_hops, ) df = cudf.DataFrame() diff --git a/python/cugraph/cugraph/tests/sampling/test_uniform_neighbor_sample.py b/python/cugraph/cugraph/tests/sampling/test_uniform_neighbor_sample.py index 39d2fbea7dd..1781ce17753 100644 --- a/python/cugraph/cugraph/tests/sampling/test_uniform_neighbor_sample.py +++ b/python/cugraph/cugraph/tests/sampling/test_uniform_neighbor_sample.py @@ -528,6 +528,206 @@ def test_uniform_neighbor_sample_empty_start_list(): assert sampling_results.empty +@pytest.mark.sg +def test_uniform_neighbor_sample_exclude_sources_basic(): + df = cudf.DataFrame( + { + "src": [0, 4, 1, 2, 3, 5, 4, 1, 0], + "dst": [1, 1, 2, 4, 3, 1, 5, 0, 2], + "eid": [9, 8, 7, 6, 5, 4, 3, 2, 1], + } + ) + + G = cugraph.MultiGraph(directed=True) + G.from_cudf_edgelist(df, source="src", destination="dst", edge_id="eid") + + sampling_results = cugraph.uniform_neighbor_sample( + G, + cudf.DataFrame( + { + "seed": cudf.Series([0, 4, 1], dtype="int64"), + "batch": cudf.Series([1, 1, 1], dtype="int32"), + } + ), + [2, 3, 3], + with_replacement=False, + with_edge_properties=True, + with_batch_ids=True, + random_state=62, + prior_sources_behavior="exclude", + ).sort_values(by="hop_id") + + expected_hop_0 = [1, 2, 1, 5, 2, 0] + assert sorted( + sampling_results[sampling_results.hop_id == 0].destinations.values_host.tolist() + ) == sorted(expected_hop_0) + + next_sources = set( + sampling_results[sampling_results.hop_id > 0].sources.values_host.tolist() + ) + for v in [0, 4, 1]: + assert v not in next_sources + + next_sources = set( + sampling_results[sampling_results.hop_id > 1].sources.values_host.tolist() + ) + for v in sampling_results[ + sampling_results.hop_id == 1 + ].sources.values_host.tolist(): + assert v not in next_sources + + +@pytest.mark.sg +def test_uniform_neighbor_sample_exclude_sources_email_eu_core(): + el = email_Eu_core.get_edgelist() + + G = cugraph.Graph(directed=True) + G.from_cudf_edgelist(el, source="src", destination="dst") + + seeds = G.select_random_vertices(62, int(0.001 * len(el))) + + sampling_results = cugraph.uniform_neighbor_sample( + G, + seeds, + [5, 4, 3, 2, 1], + with_replacement=False, + with_edge_properties=True, + with_batch_ids=False, + prior_sources_behavior="exclude", + ) + + for hop in range(5): + current_sources = set( + sampling_results[ + sampling_results.hop_id == hop + ].sources.values_host.tolist() + ) + future_sources = set( + sampling_results[sampling_results.hop_id > hop].sources.values_host.tolist() + ) + + for s in current_sources: + assert s not in future_sources + + +@pytest.mark.sg +def test_uniform_neighbor_sample_carry_over_sources_basic(): + df = cudf.DataFrame( + { + "src": [0, 4, 1, 2, 3, 5, 4, 1, 0, 6], + "dst": [1, 1, 2, 4, 6, 1, 5, 0, 2, 2], + "eid": [9, 8, 7, 6, 5, 4, 3, 2, 1, 0], + } + ) + + G = cugraph.MultiGraph(directed=True) + G.from_cudf_edgelist(df, source="src", destination="dst", edge_id="eid") + + sampling_results = cugraph.uniform_neighbor_sample( + G, + cudf.DataFrame( + { + "seed": cudf.Series([0, 4, 3], dtype="int64"), + "batch": cudf.Series([1, 1, 1], dtype="int32"), + } + ), + [2, 3, 3], + with_replacement=False, + with_edge_properties=True, + with_batch_ids=True, + random_state=62, + prior_sources_behavior="carryover", + ).sort_values(by="hop_id")[["sources", "destinations", "hop_id"]] + + assert ( + len( + sampling_results[ + (sampling_results.hop_id == 2) & (sampling_results.sources == 6) + ] + ) + == 2 + ) + + for hop in range(2): + sources_current_hop = set( + sampling_results[ + sampling_results.hop_id == hop + ].sources.values_host.tolist() + ) + sources_next_hop = set( + sampling_results[ + sampling_results.hop_id == (hop + 1) + ].sources.values_host.tolist() + ) + + for s in sources_current_hop: + assert s in sources_next_hop + + +@pytest.mark.sg +def test_uniform_neighbor_sample_carry_over_sources_email_eu_core(): + el = email_Eu_core.get_edgelist() + + G = cugraph.Graph(directed=True) + G.from_cudf_edgelist(el, source="src", destination="dst") + + seeds = G.select_random_vertices(62, int(0.001 * len(el))) + + sampling_results = cugraph.uniform_neighbor_sample( + G, + seeds, + [5, 4, 3, 2, 1], + with_replacement=False, + with_edge_properties=True, + with_batch_ids=False, + prior_sources_behavior="carryover", + ) + + for hop in range(4): + sources_current_hop = set( + sampling_results[ + sampling_results.hop_id == hop + ].sources.values_host.tolist() + ) + sources_next_hop = set( + sampling_results[ + sampling_results.hop_id == (hop + 1) + ].sources.values_host.tolist() + ) + + for s in sources_current_hop: + assert s in sources_next_hop + + +@pytest.mark.sg +def test_uniform_neighbor_sample_deduplicate_sources_email_eu_core(): + el = email_Eu_core.get_edgelist() + + G = cugraph.Graph(directed=True) + G.from_cudf_edgelist(el, source="src", destination="dst") + + seeds = G.select_random_vertices(62, int(0.001 * len(el))) + + sampling_results = cugraph.uniform_neighbor_sample( + G, + seeds, + [5, 4, 3, 2, 1], + with_replacement=False, + with_edge_properties=True, + with_batch_ids=False, + deduplicate_sources=True, + ) + + for hop in range(5): + counts_current_hop = ( + sampling_results[sampling_results.hop_id == hop] + .sources.value_counts() + .values_host.tolist() + ) + for c in counts_current_hop: + assert c <= 5 - hop + + @pytest.mark.sg @pytest.mark.skip(reason="needs to be written!") def test_multi_client_sampling(): diff --git a/python/cugraph/cugraph/tests/sampling/test_uniform_neighbor_sample_mg.py b/python/cugraph/cugraph/tests/sampling/test_uniform_neighbor_sample_mg.py index 4da3f3cf950..c0f8b087dec 100644 --- a/python/cugraph/cugraph/tests/sampling/test_uniform_neighbor_sample_mg.py +++ b/python/cugraph/cugraph/tests/sampling/test_uniform_neighbor_sample_mg.py @@ -743,6 +743,220 @@ def test_uniform_neighbor_sample_batched(dask_client, dataset, input_df, max_bat assert output_starts_per_batch <= input_starts_per_batch +@pytest.mark.mg +def test_uniform_neighbor_sample_exclude_sources_basic(dask_client): + df = dask_cudf.from_cudf( + cudf.DataFrame( + { + "src": [0, 4, 1, 2, 3, 5, 4, 1, 0], + "dst": [1, 1, 2, 4, 3, 1, 5, 0, 2], + "eid": [9, 8, 7, 6, 5, 4, 3, 2, 1], + } + ), + npartitions=1, + ) + + G = cugraph.MultiGraph(directed=True) + G.from_dask_cudf_edgelist(df, source="src", destination="dst", edge_id="eid") + + sampling_results = ( + cugraph.dask.uniform_neighbor_sample( + G, + cudf.DataFrame( + { + "seed": cudf.Series([0, 4, 1], dtype="int64"), + "batch": cudf.Series([1, 1, 1], dtype="int32"), + } + ), + [2, 3, 3], + with_replacement=False, + with_edge_properties=True, + with_batch_ids=True, + random_state=62, + prior_sources_behavior="exclude", + ) + .sort_values(by="hop_id") + .compute() + ) + + expected_hop_0 = [1, 2, 1, 5, 2, 0] + assert sorted( + sampling_results[sampling_results.hop_id == 0].destinations.values_host.tolist() + ) == sorted(expected_hop_0) + + next_sources = set( + sampling_results[sampling_results.hop_id > 0].sources.values_host.tolist() + ) + for v in [0, 4, 1]: + assert v not in next_sources + + next_sources = set( + sampling_results[sampling_results.hop_id > 1].sources.values_host.tolist() + ) + for v in sampling_results[ + sampling_results.hop_id == 1 + ].sources.values_host.tolist(): + assert v not in next_sources + + +@pytest.mark.mg +def test_uniform_neighbor_sample_exclude_sources_email_eu_core(dask_client): + el = dask_cudf.from_cudf(email_Eu_core.get_edgelist(), npartitions=8) + + G = cugraph.Graph(directed=True) + G.from_dask_cudf_edgelist(el, source="src", destination="dst") + + seeds = G.select_random_vertices(62, int(0.001 * len(el))) + + sampling_results = cugraph.dask.uniform_neighbor_sample( + G, + seeds, + [5, 4, 3, 2, 1], + with_replacement=False, + with_edge_properties=True, + with_batch_ids=False, + prior_sources_behavior="exclude", + ).compute() + + for hop in range(5): + current_sources = set( + sampling_results[ + sampling_results.hop_id == hop + ].sources.values_host.tolist() + ) + future_sources = set( + sampling_results[sampling_results.hop_id > hop].sources.values_host.tolist() + ) + + for s in current_sources: + assert s not in future_sources + + +@pytest.mark.mg +def test_uniform_neighbor_sample_carry_over_sources_basic(dask_client): + df = dask_cudf.from_cudf( + cudf.DataFrame( + { + "src": [0, 4, 1, 2, 3, 5, 4, 1, 0, 6], + "dst": [1, 1, 2, 4, 6, 1, 5, 0, 2, 2], + "eid": [9, 8, 7, 6, 5, 4, 3, 2, 1, 0], + } + ), + npartitions=4, + ) + + G = cugraph.MultiGraph(directed=True) + G.from_dask_cudf_edgelist(df, source="src", destination="dst", edge_id="eid") + + sampling_results = ( + cugraph.dask.uniform_neighbor_sample( + G, + cudf.DataFrame( + { + "seed": cudf.Series([0, 4, 3], dtype="int64"), + "batch": cudf.Series([1, 1, 1], dtype="int32"), + } + ), + [2, 3, 3], + with_replacement=False, + with_edge_properties=True, + with_batch_ids=True, + random_state=62, + prior_sources_behavior="carryover", + ) + .sort_values(by="hop_id")[["sources", "destinations", "hop_id"]] + .compute() + ) + + assert ( + len( + sampling_results[ + (sampling_results.hop_id == 2) & (sampling_results.sources == 6) + ] + ) + == 2 + ) + + for hop in range(2): + sources_current_hop = set( + sampling_results[ + sampling_results.hop_id == hop + ].sources.values_host.tolist() + ) + sources_next_hop = set( + sampling_results[ + sampling_results.hop_id == (hop + 1) + ].sources.values_host.tolist() + ) + + for s in sources_current_hop: + assert s in sources_next_hop + + +@pytest.mark.mg +def test_uniform_neighbor_sample_carry_over_sources_email_eu_core(dask_client): + el = dask_cudf.from_cudf(email_Eu_core.get_edgelist(), npartitions=8) + + G = cugraph.Graph(directed=True) + G.from_dask_cudf_edgelist(el, source="src", destination="dst") + + seeds = G.select_random_vertices(62, int(0.001 * len(el))) + + sampling_results = cugraph.dask.uniform_neighbor_sample( + G, + seeds, + [5, 4, 3, 2, 1], + with_replacement=False, + with_edge_properties=True, + with_batch_ids=False, + prior_sources_behavior="carryover", + ).compute() + + for hop in range(4): + sources_current_hop = set( + sampling_results[ + sampling_results.hop_id == hop + ].sources.values_host.tolist() + ) + sources_next_hop = set( + sampling_results[ + sampling_results.hop_id == (hop + 1) + ].sources.values_host.tolist() + ) + + for s in sources_current_hop: + assert s in sources_next_hop + + +@pytest.mark.mg +def test_uniform_neighbor_sample_deduplicate_sources_email_eu_core(dask_client): + el = dask_cudf.from_cudf(email_Eu_core.get_edgelist(), npartitions=8) + + G = cugraph.Graph(directed=True) + G.from_dask_cudf_edgelist(el, source="src", destination="dst") + + seeds = G.select_random_vertices(62, int(0.001 * len(el))) + + sampling_results = cugraph.dask.uniform_neighbor_sample( + G, + seeds, + [5, 4, 3, 2, 1], + with_replacement=False, + with_edge_properties=True, + with_batch_ids=False, + deduplicate_sources=True, + ).compute() + + for hop in range(5): + counts_current_hop = ( + sampling_results[sampling_results.hop_id == hop] + .sources.value_counts() + .values_host.tolist() + ) + for c in counts_current_hop: + assert c <= 5 - hop + + # ============================================================================= # Benchmarks # ============================================================================= diff --git a/python/pylibcugraph/pylibcugraph/_cugraph_c/algorithms.pxd b/python/pylibcugraph/pylibcugraph/_cugraph_c/algorithms.pxd index 8ced11aa4d7..9d2447466e8 100644 --- a/python/pylibcugraph/pylibcugraph/_cugraph_c/algorithms.pxd +++ b/python/pylibcugraph/pylibcugraph/_cugraph_c/algorithms.pxd @@ -231,6 +231,49 @@ cdef extern from "cugraph_c/algorithms.h": cugraph_sample_result_t** result, cugraph_error_t** error ) + + ctypedef struct cugraph_sampling_options_t: + pass + + ctypedef enum cugraph_prior_sources_behavior_t: + DEFAULT + CARRY_OVER + EXCLUDE + + cdef cugraph_error_code_t \ + cugraph_sampling_options_create( + cugraph_sampling_options_t** options, + cugraph_error_t** error, + ) + + cdef void \ + cugraph_sampling_set_with_replacement( + cugraph_sampling_options_t* options, + bool_t value, + ) + + cdef void \ + cugraph_sampling_set_return_hops( + cugraph_sampling_options_t* options, + bool_t value, + ) + + cdef void \ + cugraph_sampling_set_prior_sources_behavior( + cugraph_sampling_options_t* options, + cugraph_prior_sources_behavior_t value + ) + + cdef void \ + cugraph_sampling_set_dedupe_sources( + cugraph_sampling_options_t* options, + bool_t value, + ) + + cdef void \ + cugraph_sampling_options_free( + cugraph_sampling_options_t* options, + ) # uniform random walks cdef cugraph_error_code_t \ diff --git a/python/pylibcugraph/pylibcugraph/_cugraph_c/sampling_algorithms.pxd b/python/pylibcugraph/pylibcugraph/_cugraph_c/sampling_algorithms.pxd index ad8a8cd33a0..91cc11d6b1c 100644 --- a/python/pylibcugraph/pylibcugraph/_cugraph_c/sampling_algorithms.pxd +++ b/python/pylibcugraph/pylibcugraph/_cugraph_c/sampling_algorithms.pxd @@ -31,6 +31,7 @@ from pylibcugraph._cugraph_c.graph cimport ( ) from pylibcugraph._cugraph_c.algorithms cimport ( cugraph_sample_result_t, + cugraph_sampling_options_t, ) from pylibcugraph._cugraph_c.random cimport ( cugraph_rng_state_t, @@ -41,6 +42,8 @@ from pylibcugraph._cugraph_c.array cimport ( cdef extern from "cugraph_c/sampling_algorithms.h": ########################################################################### + + # deprecated, should migrate to cugraph_uniform_neighbor_sample cdef cugraph_error_code_t cugraph_uniform_neighbor_sample_with_edge_properties( const cugraph_resource_handle_t* handle, cugraph_graph_t* graph, @@ -57,6 +60,21 @@ cdef extern from "cugraph_c/sampling_algorithms.h": cugraph_error_t** error ) + cdef cugraph_error_code_t cugraph_uniform_neighbor_sample( + const cugraph_resource_handle_t* handle, + cugraph_graph_t* graph, + const cugraph_type_erased_device_array_view_t* start_vertices, + const cugraph_type_erased_device_array_view_t* start_vertex_labels, + const cugraph_type_erased_device_array_view_t* label_list, + const cugraph_type_erased_device_array_view_t* label_to_comm_rank, + const cugraph_type_erased_host_array_view_t* fan_out, + cugraph_rng_state_t* rng_state, + const cugraph_sampling_options_t* options, + bool_t do_expensive_check, + cugraph_sample_result_t** result, + cugraph_error_t** error + ) + cdef cugraph_error_code_t cugraph_test_uniform_neighborhood_sample_result_create( const cugraph_resource_handle_t* handle, const cugraph_type_erased_device_array_view_t* srcs, diff --git a/python/pylibcugraph/pylibcugraph/uniform_neighbor_sample.pyx b/python/pylibcugraph/pylibcugraph/uniform_neighbor_sample.pyx index a3f5dfb273f..a1832948f28 100644 --- a/python/pylibcugraph/pylibcugraph/uniform_neighbor_sample.pyx +++ b/python/pylibcugraph/pylibcugraph/uniform_neighbor_sample.pyx @@ -42,9 +42,20 @@ from pylibcugraph._cugraph_c.algorithms cimport ( cugraph_sample_result_get_destinations, cugraph_sample_result_get_index, cugraph_sample_result_free, + + cugraph_prior_sources_behavior_t, + cugraph_sampling_options_t, + cugraph_sampling_options_create, + cugraph_sampling_options_free, + cugraph_sampling_set_with_replacement, + cugraph_sampling_set_return_hops, + cugraph_sampling_set_prior_sources_behavior, + cugraph_sampling_set_dedupe_sources, + ) from pylibcugraph._cugraph_c.sampling_algorithms cimport ( - cugraph_uniform_neighbor_sample_with_edge_properties, + cugraph_uniform_neighbor_sample, + ) from pylibcugraph.resource_handle cimport ( ResourceHandle, @@ -75,12 +86,16 @@ def uniform_neighbor_sample(ResourceHandle resource_handle, _GPUGraph input_graph, start_list, h_fan_out, + *, bool_t with_replacement, bool_t do_expensive_check, bool_t with_edge_properties=False, batch_id_list=None, label_list=None, label_to_output_comm_rank=None, + prior_sources_behavior=None, + bool_t deduplicate_sources=False, + bool_t return_hops=False, random_state=None): """ Does neighborhood sampling, which samples nodes from a graph based on the @@ -118,6 +133,27 @@ def uniform_neighbor_sample(ResourceHandle resource_handle, batch_id_list: list[int32] (Optional) List of int32 batch ids that is returned with each edge. Optional argument, defaults to NULL, returning nothing. + + label_list: list[int32] (Optional) + List of unique int32 batch ids. Required if also passing the + label_to_output_comm_rank flag. Default to NULL (does nothing) + + label_to_output_comm_rank: list[int32] (Optional) + Maps the unique batch ids in label_list to the rank of the + worker that should hold results for that batch id. + Defaults to NULL (does nothing) + + prior_sources_behavior: str (Optional) + Options are "carryover", and "exclude". + Default will leave the source list as-is. + Carryover will carry over sources from previous hops to the + current hop. + Exclude will exclude sources from previous hops from reappearing + as sources in future hops. + + deduplicate_sources: bool (Optional) + If True, will deduplicate the source list before sampling. + Defaults to False. random_state: int (Optional) Random state to use when generating samples. Optional argument, @@ -212,7 +248,30 @@ def uniform_neighbor_sample(ResourceHandle resource_handle, cdef cugraph_rng_state_t* rng_state_ptr = \ cg_rng_state.rng_state_ptr - error_code = cugraph_uniform_neighbor_sample_with_edge_properties( + + cdef cugraph_prior_sources_behavior_t prior_sources_behavior_e + if prior_sources_behavior is None: + prior_sources_behavior_e = cugraph_prior_sources_behavior_t.DEFAULT + elif prior_sources_behavior == 'carryover': + prior_sources_behavior_e = cugraph_prior_sources_behavior_t.CARRY_OVER + elif prior_sources_behavior == 'exclude': + prior_sources_behavior_e = cugraph_prior_sources_behavior_t.EXCLUDE + else: + raise ValueError( + f'Invalid option {prior_sources_behavior}' + ' for prior sources behavior' + ) + + cdef cugraph_sampling_options_t* sampling_options + error_code = cugraph_sampling_options_create(&sampling_options, &error_ptr) + assert_success(error_code, error_ptr, "cugraph_sampling_options_create") + + cugraph_sampling_set_with_replacement(sampling_options, with_replacement) + cugraph_sampling_set_return_hops(sampling_options, return_hops) + cugraph_sampling_set_dedupe_sources(sampling_options, deduplicate_sources) + cugraph_sampling_set_prior_sources_behavior(sampling_options, prior_sources_behavior_e) + + error_code = cugraph_uniform_neighbor_sample( c_resource_handle_ptr, c_graph_ptr, start_ptr, @@ -221,12 +280,14 @@ def uniform_neighbor_sample(ResourceHandle resource_handle, label_to_output_comm_rank_ptr, fan_out_ptr, rng_state_ptr, - with_replacement, - True, # return_hops + sampling_options, do_expensive_check, &result_ptr, &error_ptr) - assert_success(error_code, error_ptr, "cugraph_uniform_neighbor_sample_with_edge_properties") + assert_success(error_code, error_ptr, "cugraph_uniform_neighbor_sample") + + # Free the sampling options + cugraph_sampling_options_free(sampling_options) # Free the two input arrays that are no longer needed. cugraph_type_erased_device_array_view_free(start_ptr) From 59b0eb70c4e3157c3184128bda52104c925bef25 Mon Sep 17 00:00:00 2001 From: Joseph Nke <76006812+jnke2016@users.noreply.github.com> Date: Wed, 19 Jul 2023 13:39:29 +0100 Subject: [PATCH 07/14] [BUG] Unsupported graph for similiarity algos (#3710) This PR update the docstrings raises an error when running any similarity algos with vertices from a graph that are unrenumbered. Authors: - Joseph Nke (https://github.com/jnke2016) Approvers: - Brad Rees (https://github.com/BradReesWork) URL: https://github.com/rapidsai/cugraph/pull/3710 --- .../cugraph/link_prediction/jaccard.py | 34 +++++++++++++++++-- .../cugraph/link_prediction/overlap.py | 33 ++++++++++++++++-- .../cugraph/link_prediction/sorensen.py | 32 +++++++++++++++-- .../cugraph/link_prediction/wjaccard.py | 28 +++++++++++++-- .../cugraph/link_prediction/woverlap.py | 25 +++++++++++++- .../cugraph/link_prediction/wsorensen.py | 26 +++++++++++++- .../tests/link_prediction/test_jaccard.py | 12 +++++++ .../tests/link_prediction/test_overlap.py | 11 ++++++ .../tests/link_prediction/test_sorensen.py | 12 +++++++ .../tests/link_prediction/test_wjaccard.py | 11 ++++++ .../tests/link_prediction/test_woverlap.py | 11 ++++++ .../tests/link_prediction/test_wsorensen.py | 11 ++++++ 12 files changed, 234 insertions(+), 12 deletions(-) diff --git a/python/cugraph/cugraph/link_prediction/jaccard.py b/python/cugraph/cugraph/link_prediction/jaccard.py index 1c4fed7a8f9..dd411fa889d 100644 --- a/python/cugraph/cugraph/link_prediction/jaccard.py +++ b/python/cugraph/cugraph/link_prediction/jaccard.py @@ -1,4 +1,4 @@ -# Copyright (c) 2019-2022, NVIDIA CORPORATION. +# Copyright (c) 2019-2023, 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 @@ -20,7 +20,7 @@ ) -def jaccard(input_graph, vertex_pair=None): +def jaccard(input_graph, vertex_pair=None, do_expensive_check=True): """ Compute the Jaccard similarity between each pair of vertices connected by an edge, or between arbitrary pairs of vertices specified by the user. @@ -36,6 +36,10 @@ def jaccard(input_graph, vertex_pair=None): of cugraph.jaccard is different from the behavior of networkx.jaccard_coefficient. + This algorithm doesn't currently support datasets with vertices that + are not (re)numebred vertices from 0 to V-1 where V is the total number of + vertices as this creates isolated vertices. + cugraph.jaccard, in the absence of a specified vertex pair list, will use the edges of the graph to construct a vertex pair list and will return the jaccard coefficient for those vertex pairs. @@ -80,6 +84,10 @@ def jaccard(input_graph, vertex_pair=None): current implementation computes the jaccard coefficient for all adjacent vertices in the graph. + do_expensive_check: bool (default=True) + When set to True, check if the vertices in the graph are (re)numbered + from 0 to V-1 where V is the total number of vertices. + Returns ------- df : cudf.DataFrame @@ -104,6 +112,22 @@ def jaccard(input_graph, vertex_pair=None): >>> df = cugraph.jaccard(G) """ + if do_expensive_check: + if not input_graph.renumbered: + input_df = input_graph.edgelist.edgelist_df[["src", "dst"]] + max_vertex = input_df.max().max() + expected_nodes = cudf.Series(range(0, max_vertex + 1, 1)).astype( + input_df.dtypes[0] + ) + nodes = ( + cudf.concat([input_df["src"], input_df["dst"]]) + .unique() + .sort_values() + .reset_index(drop=True) + ) + if not expected_nodes.equals(nodes): + raise ValueError("Unrenumbered vertices are not supported.") + if input_graph.is_directed(): raise ValueError("Input must be an undirected Graph.") if type(vertex_pair) == cudf.DataFrame: @@ -120,10 +144,14 @@ def jaccard(input_graph, vertex_pair=None): return df -def jaccard_coefficient(G, ebunch=None): +def jaccard_coefficient(G, ebunch=None, do_expensive_check=True): """ For NetworkX Compatability. See `jaccard` + NOTE: This algorithm doesn't currently support datasets with vertices that + are not (re)numebred vertices from 0 to V-1 where V is the total number of + vertices as this creates isolated vertices. + Parameters ---------- graph : cugraph.Graph diff --git a/python/cugraph/cugraph/link_prediction/overlap.py b/python/cugraph/cugraph/link_prediction/overlap.py index ba9f225062e..e05e0c944fe 100644 --- a/python/cugraph/cugraph/link_prediction/overlap.py +++ b/python/cugraph/cugraph/link_prediction/overlap.py @@ -1,4 +1,4 @@ -# Copyright (c) 2019-2022, NVIDIA CORPORATION. +# Copyright (c) 2019-2023, 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 @@ -20,10 +20,14 @@ ) -def overlap_coefficient(G, ebunch=None): +def overlap_coefficient(G, ebunch=None, do_expensive_check=True): """ For NetworkX Compatability. See `overlap` + NOTE: This algorithm doesn't currently support datasets with vertices that + are not (re)numebred vertices from 0 to V-1 where V is the total number of + vertices as this creates isolated vertices. + """ vertex_pair = None @@ -42,7 +46,7 @@ def overlap_coefficient(G, ebunch=None): return df -def overlap(input_graph, vertex_pair=None): +def overlap(input_graph, vertex_pair=None, do_expensive_check=True): """ Compute the Overlap Coefficient between each pair of vertices connected by an edge, or between arbitrary pairs of vertices specified by the user. @@ -54,6 +58,10 @@ def overlap(input_graph, vertex_pair=None): neighbors. If first is specified but second is not, or vice versa, an exception will be thrown. + NOTE: This algorithm doesn't currently support datasets with vertices that + are not (re)numebred vertices from 0 to V-1 where V is the total number of + vertices as this creates isolated vertices. + Parameters ---------- input_graph : cugraph.Graph @@ -66,6 +74,10 @@ def overlap(input_graph, vertex_pair=None): vertices. If provided, the overlap coefficient is computed for the given vertex pairs, else, it is computed for all vertex pairs. + do_expensive_check: bool (default=True) + When set to True, check if the vertices in the graph are (re)numbered + from 0 to V-1 where V is the total number of vertices. + Returns ------- df : cudf.DataFrame @@ -90,6 +102,21 @@ def overlap(input_graph, vertex_pair=None): >>> df = cugraph.overlap(G) """ + if do_expensive_check: + if not input_graph.renumbered: + input_df = input_graph.edgelist.edgelist_df[["src", "dst"]] + max_vertex = input_df.max().max() + expected_nodes = cudf.Series(range(0, max_vertex + 1, 1)).astype( + input_df.dtypes[0] + ) + nodes = ( + cudf.concat([input_df["src"], input_df["dst"]]) + .unique() + .sort_values() + .reset_index(drop=True) + ) + if not expected_nodes.equals(nodes): + raise ValueError("Unrenumbered vertices are not supported.") if type(vertex_pair) == cudf.DataFrame: vertex_pair = renumber_vertex_pair(input_graph, vertex_pair) diff --git a/python/cugraph/cugraph/link_prediction/sorensen.py b/python/cugraph/cugraph/link_prediction/sorensen.py index 20238e10464..0f35f868b7c 100644 --- a/python/cugraph/cugraph/link_prediction/sorensen.py +++ b/python/cugraph/cugraph/link_prediction/sorensen.py @@ -21,7 +21,7 @@ ) -def sorensen(input_graph, vertex_pair=None): +def sorensen(input_graph, vertex_pair=None, do_expensive_check=True): """ Compute the Sorensen coefficient between each pair of vertices connected by an edge, or between arbitrary pairs of vertices specified by the user. @@ -30,6 +30,10 @@ def sorensen(input_graph, vertex_pair=None): If first is specified but second is not, or vice versa, an exception will be thrown. + NOTE: This algorithm doesn't currently support datasets with vertices that + are not (re)numebred vertices from 0 to V-1 where V is the total number of + vertices as this creates isolated vertices. + cugraph.sorensen, in the absence of a specified vertex pair list, will use the edges of the graph to construct a vertex pair list and will return the sorensen coefficient for those vertex pairs. @@ -50,6 +54,10 @@ def sorensen(input_graph, vertex_pair=None): current implementation computes the Sorensen coefficient for all adjacent vertices in the graph. + do_expensive_check: bool (default=True) + When set to True, check if the vertices in the graph are (re)numbered + from 0 to V-1 where V is the total number of vertices. + Returns ------- df : cudf.DataFrame @@ -76,6 +84,22 @@ def sorensen(input_graph, vertex_pair=None): >>> df = cugraph.sorensen(G) """ + if do_expensive_check: + if not input_graph.renumbered: + input_df = input_graph.edgelist.edgelist_df[["src", "dst"]] + max_vertex = input_df.max().max() + expected_nodes = cudf.Series(range(0, max_vertex + 1, 1)).astype( + input_df.dtypes[0] + ) + nodes = ( + cudf.concat([input_df["src"], input_df["dst"]]) + .unique() + .sort_values() + .reset_index(drop=True) + ) + if not expected_nodes.equals(nodes): + raise ValueError("Unrenumbered vertices are not supported.") + if type(input_graph) is not Graph: raise TypeError("input graph must a Graph") @@ -94,10 +118,14 @@ def sorensen(input_graph, vertex_pair=None): return df -def sorensen_coefficient(G, ebunch=None): +def sorensen_coefficient(G, ebunch=None, do_expensive_check=True): """ For NetworkX Compatability. See `sorensen` + NOTE: This algorithm doesn't currently support datasets with vertices that + are not (re)numebred vertices from 0 to V-1 where V is the total number of + vertices as this creates isolated vertices. + Parameters ---------- G : cugraph.Graph diff --git a/python/cugraph/cugraph/link_prediction/wjaccard.py b/python/cugraph/cugraph/link_prediction/wjaccard.py index b8ef33d926f..fc6edae8d3e 100644 --- a/python/cugraph/cugraph/link_prediction/wjaccard.py +++ b/python/cugraph/cugraph/link_prediction/wjaccard.py @@ -1,4 +1,4 @@ -# Copyright (c) 2019-2022, NVIDIA CORPORATION. +# Copyright (c) 2019-2023, 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 @@ -17,7 +17,7 @@ from cugraph.utilities import renumber_vertex_pair -def jaccard_w(input_graph, weights, vertex_pair=None): +def jaccard_w(input_graph, weights, vertex_pair=None, do_expensive_check=True): """ Compute the weighted Jaccard similarity between each pair of vertices connected by an edge, or between arbitrary pairs of vertices specified by @@ -29,6 +29,10 @@ def jaccard_w(input_graph, weights, vertex_pair=None): neighbors. If first is specified but second is not, or vice versa, an exception will be thrown. + NOTE: This algorithm doesn't currently support datasets with vertices that + are not (re)numebred vertices from 0 to V-1 where V is the total number of + vertices as this creates isolated vertices. + Parameters ---------- input_graph : cugraph.Graph @@ -51,6 +55,10 @@ def jaccard_w(input_graph, weights, vertex_pair=None): vertices. If provided, the jaccard coefficient is computed for the given vertex pairs, else, it is computed for all vertex pairs. + do_expensive_check: bool (default=True) + When set to True, check if the vertices in the graph are (re)numbered + from 0 to V-1 where V is the total number of vertices. + Returns ------- df : cudf.DataFrame @@ -87,6 +95,22 @@ def jaccard_w(input_graph, weights, vertex_pair=None): >>> df = cugraph.jaccard_w(G, weights) """ + if do_expensive_check: + if not input_graph.renumbered: + input_df = input_graph.edgelist.edgelist_df[["src", "dst"]] + max_vertex = input_df.max().max() + expected_nodes = cudf.Series(range(0, max_vertex + 1, 1)).astype( + input_df.dtypes[0] + ) + nodes = ( + cudf.concat([input_df["src"], input_df["dst"]]) + .unique() + .sort_values() + .reset_index(drop=True) + ) + if not expected_nodes.equals(nodes): + raise ValueError("Unrenumbered vertices are not supported.") + if type(input_graph) is not Graph: raise TypeError("input graph must a Graph") diff --git a/python/cugraph/cugraph/link_prediction/woverlap.py b/python/cugraph/cugraph/link_prediction/woverlap.py index c7d4f56a428..27fb7d608ca 100644 --- a/python/cugraph/cugraph/link_prediction/woverlap.py +++ b/python/cugraph/cugraph/link_prediction/woverlap.py @@ -16,7 +16,7 @@ from cugraph.utilities import renumber_vertex_pair -def overlap_w(input_graph, weights, vertex_pair=None): +def overlap_w(input_graph, weights, vertex_pair=None, do_expensive_check=True): """ Compute the weighted Overlap Coefficient between each pair of vertices connected by an edge, or between arbitrary pairs of vertices specified by @@ -28,6 +28,10 @@ def overlap_w(input_graph, weights, vertex_pair=None): neighbors. If first is specified but second is not, or vice versa, an exception will be thrown. + NOTE: This algorithm doesn't currently support datasets with vertices that + are not (re)numebred vertices from 0 to V-1 where V is the total number of + vertices as this creates isolated vertices. + Parameters ---------- input_graph : cugraph.Graph @@ -51,6 +55,10 @@ def overlap_w(input_graph, weights, vertex_pair=None): vertices. If provided, the overlap coefficient is computed for the given vertex pairs, else, it is computed for all vertex pairs. + do_expensive_check: bool (default=True) + When set to True, check if the vertices in the graph are (re)numbered + from 0 to V-1 where V is the total number of vertices. + Returns ------- df : cudf.DataFrame @@ -88,6 +96,21 @@ def overlap_w(input_graph, weights, vertex_pair=None): ... len(weights['vertex']))] >>> df = cugraph.overlap_w(G, weights) """ + if do_expensive_check: + if not input_graph.renumbered: + input_df = input_graph.edgelist.edgelist_df[["src", "dst"]] + max_vertex = input_df.max().max() + expected_nodes = cudf.Series(range(0, max_vertex + 1, 1)).astype( + input_df.dtypes[0] + ) + nodes = ( + cudf.concat([input_df["src"], input_df["dst"]]) + .unique() + .sort_values() + .reset_index(drop=True) + ) + if not expected_nodes.equals(nodes): + raise ValueError("Unrenumbered vertices are not supported.") if type(vertex_pair) == cudf.DataFrame: vertex_pair = renumber_vertex_pair(input_graph, vertex_pair) diff --git a/python/cugraph/cugraph/link_prediction/wsorensen.py b/python/cugraph/cugraph/link_prediction/wsorensen.py index c017463a294..c27e4f66a02 100644 --- a/python/cugraph/cugraph/link_prediction/wsorensen.py +++ b/python/cugraph/cugraph/link_prediction/wsorensen.py @@ -17,13 +17,17 @@ from cugraph.utilities import renumber_vertex_pair -def sorensen_w(input_graph, weights, vertex_pair=None): +def sorensen_w(input_graph, weights, vertex_pair=None, do_expensive_check=True): """ Compute the weighted Sorensen similarity between each pair of vertices connected by an edge, or between arbitrary pairs of vertices specified by the user. Sorensen coefficient is defined between two sets as the ratio of twice the volume of their intersection divided by the volume of each set. + NOTE: This algorithm doesn't currently support datasets with vertices that + are not (re)numebred vertices from 0 to V-1 where V is the total number of + vertices as this creates isolated vertices. + Parameters ---------- input_graph : cugraph.Graph @@ -47,6 +51,10 @@ def sorensen_w(input_graph, weights, vertex_pair=None): vertices. If provided, the sorensen coefficient is computed for the given vertex pairs, else, it is computed for all vertex pairs. + do_expensive_check: bool (default=True) + When set to True, check if the vertices in the graph are (re)numbered + from 0 to V-1 where V is the total number of vertices. + Returns ------- df : cudf.DataFrame @@ -85,6 +93,22 @@ def sorensen_w(input_graph, weights, vertex_pair=None): >>> df = cugraph.sorensen_w(G, weights) """ + if do_expensive_check: + if not input_graph.renumbered: + input_df = input_graph.edgelist.edgelist_df[["src", "dst"]] + max_vertex = input_df.max().max() + expected_nodes = cudf.Series(range(0, max_vertex + 1, 1)).astype( + input_df.dtypes[0] + ) + nodes = ( + cudf.concat([input_df["src"], input_df["dst"]]) + .unique() + .sort_values() + .reset_index(drop=True) + ) + if not expected_nodes.equals(nodes): + raise ValueError("Unrenumbered vertices are not supported.") + if type(input_graph) is not Graph: raise TypeError("input graph must a Graph") diff --git a/python/cugraph/cugraph/tests/link_prediction/test_jaccard.py b/python/cugraph/cugraph/tests/link_prediction/test_jaccard.py index b04c4c741b1..43077126827 100644 --- a/python/cugraph/cugraph/tests/link_prediction/test_jaccard.py +++ b/python/cugraph/cugraph/tests/link_prediction/test_jaccard.py @@ -202,6 +202,7 @@ def test_nx_jaccard_time(read_csv, gpubenchmark): @pytest.mark.sg @pytest.mark.parametrize("graph_file", [netscience]) +@pytest.mark.skip(reason="Skipping because this datasets is unrenumbered") def test_jaccard_edgevals(gpubenchmark, graph_file): dataset_path = netscience.get_path() M = utils.read_csv_for_nx(dataset_path) @@ -326,3 +327,14 @@ def test_weighted_exp_jaccard(): use_weight = True with pytest.raises(ValueError): exp_jaccard(G, use_weight=use_weight) + + +@pytest.mark.sg +def test_invalid_datasets_jaccard(): + karate = DATASETS_UNDIRECTED[0] + df = karate.get_edgelist() + df = df.add(1) + G = cugraph.Graph(directed=False) + G.from_cudf_edgelist(df, source="src", destination="dst") + with pytest.raises(ValueError): + cugraph.jaccard(G) diff --git a/python/cugraph/cugraph/tests/link_prediction/test_overlap.py b/python/cugraph/cugraph/tests/link_prediction/test_overlap.py index 68f879dacdb..03bee451f3c 100644 --- a/python/cugraph/cugraph/tests/link_prediction/test_overlap.py +++ b/python/cugraph/cugraph/tests/link_prediction/test_overlap.py @@ -225,3 +225,14 @@ def test_weighted_exp_overlap(): use_weight = True with pytest.raises(ValueError): exp_overlap(G, use_weight=use_weight) + + +@pytest.mark.sg +def test_invalid_datasets_overlap(): + karate = DATASETS_UNDIRECTED[0] + df = karate.get_edgelist() + df = df.add(1) + G = cugraph.Graph(directed=False) + G.from_cudf_edgelist(df, source="src", destination="dst") + with pytest.raises(ValueError): + cugraph.overlap(G) diff --git a/python/cugraph/cugraph/tests/link_prediction/test_sorensen.py b/python/cugraph/cugraph/tests/link_prediction/test_sorensen.py index 3457627ed7d..14d84784161 100644 --- a/python/cugraph/cugraph/tests/link_prediction/test_sorensen.py +++ b/python/cugraph/cugraph/tests/link_prediction/test_sorensen.py @@ -187,6 +187,7 @@ def test_nx_sorensen_time(gpubenchmark, read_csv): @pytest.mark.sg @pytest.mark.parametrize("graph_file", [netscience]) +@pytest.mark.skip(reason="Skipping because this datasets is unrenumbered") def test_sorensen_edgevals(gpubenchmark, graph_file): dataset_path = netscience.get_path() M = utils.read_csv_for_nx(dataset_path) @@ -288,3 +289,14 @@ def test_weighted_exp_sorensen(): use_weight = True with pytest.raises(ValueError): exp_sorensen(G, use_weight=use_weight) + + +@pytest.mark.sg +def test_invalid_datasets_sorensen(): + karate = DATASETS_UNDIRECTED[0] + df = karate.get_edgelist() + df = df.add(1) + G = cugraph.Graph(directed=False) + G.from_cudf_edgelist(df, source="src", destination="dst") + with pytest.raises(ValueError): + cugraph.sorensen(G) diff --git a/python/cugraph/cugraph/tests/link_prediction/test_wjaccard.py b/python/cugraph/cugraph/tests/link_prediction/test_wjaccard.py index 22ace93c0e4..2bc39b877ea 100644 --- a/python/cugraph/cugraph/tests/link_prediction/test_wjaccard.py +++ b/python/cugraph/cugraph/tests/link_prediction/test_wjaccard.py @@ -176,3 +176,14 @@ def test_wjaccard_multi_column(read_csv): actual = df_res.sort_values("0_first").reset_index() expected = df_exp.sort_values("first").reset_index() assert_series_equal(actual["jaccard_coeff"], expected["jaccard_coeff"]) + + +@pytest.mark.sg +def test_invalid_datasets_jaccard_w(): + karate = DATASETS_UNDIRECTED[0] + df = karate.get_edgelist() + df = df.add(1) + G = cugraph.Graph(directed=False) + G.from_cudf_edgelist(df, source="src", destination="dst") + with pytest.raises(ValueError): + cugraph.jaccard_w(G, None) diff --git a/python/cugraph/cugraph/tests/link_prediction/test_woverlap.py b/python/cugraph/cugraph/tests/link_prediction/test_woverlap.py index f4fab9d0faa..5e35bb66f07 100644 --- a/python/cugraph/cugraph/tests/link_prediction/test_woverlap.py +++ b/python/cugraph/cugraph/tests/link_prediction/test_woverlap.py @@ -159,3 +159,14 @@ def test_woverlap_multi_column(graph_file): actual = df_res.sort_values("0_first").reset_index() expected = df_exp.sort_values("first").reset_index() assert_series_equal(actual["overlap_coeff"], expected["overlap_coeff"]) + + +@pytest.mark.sg +def test_invalid_datasets_overlap_w(): + karate = DATASETS_UNDIRECTED[0] + df = karate.get_edgelist() + df = df.add(1) + G = cugraph.Graph(directed=False) + G.from_cudf_edgelist(df, source="src", destination="dst") + with pytest.raises(ValueError): + cugraph.overlap_w(G, None) diff --git a/python/cugraph/cugraph/tests/link_prediction/test_wsorensen.py b/python/cugraph/cugraph/tests/link_prediction/test_wsorensen.py index 0cf775d666c..cca2363d2d6 100644 --- a/python/cugraph/cugraph/tests/link_prediction/test_wsorensen.py +++ b/python/cugraph/cugraph/tests/link_prediction/test_wsorensen.py @@ -180,3 +180,14 @@ def test_wsorensen_multi_column(read_csv): actual = df_res.sort_values("0_first").reset_index() expected = df_exp.sort_values("first").reset_index() assert_series_equal(actual["sorensen_coeff"], expected["sorensen_coeff"]) + + +@pytest.mark.sg +def test_invalid_datasets_sorensen_w(): + karate = DATASETS_UNDIRECTED[0] + df = karate.get_edgelist() + df = df.add(1) + G = cugraph.Graph(directed=False) + G.from_cudf_edgelist(df, source="src", destination="dst") + with pytest.raises(ValueError): + cugraph.sorensen_w(G, None) From a280986afc6bd9bdb61f990726f61483d5b30620 Mon Sep 17 00:00:00 2001 From: Vibhu Jawa Date: Wed, 19 Jul 2023 10:34:06 -0700 Subject: [PATCH 08/14] [WIP] Make edge ids optional (#3702) This PR makes edge ids optional for cugraph-dgl dataloaders Todo: - [ ] Add tests Authors: - Vibhu Jawa (https://github.com/VibhuJawa) - Brad Rees (https://github.com/BradReesWork) Approvers: - Alex Barghi (https://github.com/alexbarghi-nv) URL: https://github.com/rapidsai/cugraph/pull/3702 --- .../dataloading/utils/sampling_helpers.py | 20 +++++++++++++------ .../gnn/data_loading/bulk_sampler_io.py | 2 ++ 2 files changed, 16 insertions(+), 6 deletions(-) diff --git a/python/cugraph-dgl/cugraph_dgl/dataloading/utils/sampling_helpers.py b/python/cugraph-dgl/cugraph_dgl/dataloading/utils/sampling_helpers.py index 02052c9841d..051464f08bb 100644 --- a/python/cugraph-dgl/cugraph_dgl/dataloading/utils/sampling_helpers.py +++ b/python/cugraph-dgl/cugraph_dgl/dataloading/utils/sampling_helpers.py @@ -52,9 +52,11 @@ def _get_tensor_ls_from_sampled_df(df): batch_indices = torch.searchsorted(batch_id_tensor, batch_indices) split_d = {} + for column in ["sources", "destinations", "edge_id", "hop_id"]: - tensor = cast_to_tensor(df[column]) - split_d[column] = torch.tensor_split(tensor, batch_indices.cpu()) + if column in df.columns: + tensor = cast_to_tensor(df[column]) + split_d[column] = torch.tensor_split(tensor, batch_indices.cpu()) result_tensor_ls = [] for i, hop_id_tensor in enumerate(split_d["hop_id"]): @@ -66,7 +68,11 @@ def _get_tensor_ls_from_sampled_df(df): hop_indices = torch.searchsorted(hop_id_tensor, hop_indices) s = torch.tensor_split(split_d["sources"][i], hop_indices.cpu()) d = torch.tensor_split(split_d["destinations"][i], hop_indices.cpu()) - eid = torch.tensor_split(split_d["edge_id"][i], hop_indices.cpu()) + if "edge_id" in split_d: + eid = torch.tensor_split(split_d["edge_id"][i], hop_indices.cpu()) + else: + eid = [None] * len(s) + result_tensor_ls.append((x, y, z) for x, y, z in zip(s, d, eid)) return result_tensor_ls @@ -125,7 +131,7 @@ def _create_homogeneous_sampled_graphs_from_tensors_perhop( def create_homogeneous_dgl_block_from_tensors_ls( src_ids: torch.Tensor, dst_ids: torch.Tensor, - edge_ids: torch.Tensor, + edge_ids: Optional[torch.Tensor], seed_nodes: Optional[torch.Tensor], total_number_of_nodes: int, ): @@ -133,7 +139,8 @@ def create_homogeneous_dgl_block_from_tensors_ls( (src_ids, dst_ids), num_nodes=total_number_of_nodes, ) - sampled_graph.edata[dgl.EID] = edge_ids + if edge_ids is not None: + sampled_graph.edata[dgl.EID] = edge_ids # TODO: Check if unique is needed if seed_nodes is None: seed_nodes = dst_ids.unique() @@ -144,7 +151,8 @@ def create_homogeneous_dgl_block_from_tensors_ls( src_nodes=src_ids.unique(), include_dst_in_src=True, ) - block.edata[dgl.EID] = sampled_graph.edata[dgl.EID] + if edge_ids is not None: + block.edata[dgl.EID] = sampled_graph.edata[dgl.EID] return block diff --git a/python/cugraph/cugraph/gnn/data_loading/bulk_sampler_io.py b/python/cugraph/cugraph/gnn/data_loading/bulk_sampler_io.py index 44c1185bbf1..002b214e783 100644 --- a/python/cugraph/cugraph/gnn/data_loading/bulk_sampler_io.py +++ b/python/cugraph/cugraph/gnn/data_loading/bulk_sampler_io.py @@ -51,6 +51,8 @@ def _write_samples_to_parquet( raise ValueError("Invalid value of partition_info") max_batch_id = offsets.batch_id.max() + results.dropna(axis=1, how="all", inplace=True) + results["hop_id"] = results["hop_id"].astype("uint8") for p in range(0, len(offsets), batches_per_partition): offsets_p = offsets.iloc[p : p + batches_per_partition] From 145a92bb691b3e290b50cc7d7678d8cbff137616 Mon Sep 17 00:00:00 2001 From: Ray Douglass <3107146+raydouglass@users.noreply.github.com> Date: Wed, 19 Jul 2023 15:45:51 -0400 Subject: [PATCH 09/14] Reorder dependencies.yaml channels (#3721) This PR reorders the channels to be compatible with other RAPIDS libs. #3693 added two new channels, but they were ordered differently than [other libraries](https://github.com/rapidsai/cudf/blob/branch-23.08/dependencies.yaml#L163-L169). This caused [an error](https://github.com/rapidsai/docker/actions/runs/5591817313/jobs/10223438072?pr=545#step:10:349) in the `conda-merge` script in https://github.com/rapidsai/docker/pull/545. It's also worth noting that docker overhaul requiring the channel ordering being consistent across libraries' `dependencies.yaml` seems very fragile and will undoubtedly cause more issues down the line. Curious if anyone has thoughts on making this more reliable. Authors: - Ray Douglass (https://github.com/raydouglass) Approvers: - AJ Schmidt (https://github.com/ajschmidt8) - Bradley Dice (https://github.com/bdice) URL: https://github.com/rapidsai/cugraph/pull/3721 --- conda/environments/all_cuda-118_arch-x86_64.yaml | 4 ++-- dependencies.yaml | 4 ++-- python/cugraph-dgl/conda/cugraph_dgl_dev_cuda-118.yaml | 4 ++-- 3 files changed, 6 insertions(+), 6 deletions(-) diff --git a/conda/environments/all_cuda-118_arch-x86_64.yaml b/conda/environments/all_cuda-118_arch-x86_64.yaml index 9c428ef9d07..1540a85bcc8 100644 --- a/conda/environments/all_cuda-118_arch-x86_64.yaml +++ b/conda/environments/all_cuda-118_arch-x86_64.yaml @@ -4,10 +4,10 @@ channels: - rapidsai - rapidsai-nightly - dask/label/dev -- conda-forge -- nvidia - pytorch - dglteam/label/cu118 +- conda-forge +- nvidia dependencies: - aiohttp - c-compiler diff --git a/dependencies.yaml b/dependencies.yaml index 9b858999743..d6e36a9643c 100644 --- a/dependencies.yaml +++ b/dependencies.yaml @@ -181,10 +181,10 @@ channels: - rapidsai - rapidsai-nightly - dask/label/dev - - conda-forge - - nvidia - pytorch - dglteam/label/cu118 + - conda-forge + - nvidia dependencies: checks: common: diff --git a/python/cugraph-dgl/conda/cugraph_dgl_dev_cuda-118.yaml b/python/cugraph-dgl/conda/cugraph_dgl_dev_cuda-118.yaml index a252d5e0c78..6961a485742 100644 --- a/python/cugraph-dgl/conda/cugraph_dgl_dev_cuda-118.yaml +++ b/python/cugraph-dgl/conda/cugraph_dgl_dev_cuda-118.yaml @@ -4,10 +4,10 @@ channels: - rapidsai - rapidsai-nightly - dask/label/dev -- conda-forge -- nvidia - pytorch - dglteam/label/cu118 +- conda-forge +- nvidia dependencies: - cugraph==23.8.* - dgl>=1.1.0.cu* From 54b43e14d9e1045b61c27be2e80e4e1e6fcccd35 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Thu, 20 Jul 2023 05:59:47 -0700 Subject: [PATCH 10/14] cugraph: Build CUDA 12 packages (#3456) First attempt at building libcugraph with CUDA 12 and nvidia channel packages Closes https://github.com/rapidsai/cugraph/issues/3271 Authors: - Vyas Ramasubramani (https://github.com/vyasr) - Rick Ratzel (https://github.com/rlratzel) - Bradley Dice (https://github.com/bdice) Approvers: - Ray Douglass (https://github.com/raydouglass) - Chuck Hastings (https://github.com/ChuckHastings) - Bradley Dice (https://github.com/bdice) - Rick Ratzel (https://github.com/rlratzel) - https://github.com/jakirkham URL: https://github.com/rapidsai/cugraph/pull/3456 --- .github/workflows/build.yaml | 4 +- .github/workflows/pr.yaml | 8 +- .github/workflows/test.yaml | 2 - ci/build_python.sh | 54 +++++++----- .../all_cuda-118_arch-x86_64.yaml | 3 +- conda/recipes/cugraph/conda_build_config.yaml | 18 +--- conda/recipes/cugraph/meta.yaml | 40 ++++----- .../libcugraph/conda_build_config.yaml | 39 +++++++-- conda/recipes/libcugraph/meta.yaml | 83 +++++++++++++++---- .../pylibcugraph/conda_build_config.yaml | 18 +--- conda/recipes/pylibcugraph/meta.yaml | 35 ++++---- dependencies.yaml | 23 +++-- 12 files changed, 192 insertions(+), 135 deletions(-) diff --git a/.github/workflows/build.yaml b/.github/workflows/build.yaml index cf73e1d2d27..de9caa0fabe 100644 --- a/.github/workflows/build.yaml +++ b/.github/workflows/build.yaml @@ -30,7 +30,6 @@ jobs: secrets: inherit uses: rapidsai/shared-action-workflows/.github/workflows/conda-cpp-build.yaml@branch-23.08 with: - matrix_filter: map(select(.CUDA_VER | startswith("11"))) build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} date: ${{ inputs.date }} @@ -40,7 +39,6 @@ jobs: secrets: inherit uses: rapidsai/shared-action-workflows/.github/workflows/conda-python-build.yaml@branch-23.08 with: - matrix_filter: map(select(.CUDA_VER | startswith("11"))) build_type: ${{ inputs.build_type || 'branch' }} branch: ${{ inputs.branch }} date: ${{ inputs.date }} @@ -63,7 +61,7 @@ jobs: arch: "amd64" branch: ${{ inputs.branch }} build_type: ${{ inputs.build_type || 'branch' }} - container_image: "rapidsai/ci:cuda11.8.0-ubuntu22.04-py3.10" + container_image: "rapidsai/ci:latest" date: ${{ inputs.date }} node_type: "gpu-v100-latest-1" run_script: "ci/build_docs.sh" diff --git a/.github/workflows/pr.yaml b/.github/workflows/pr.yaml index d9029ea37a1..4d52cd26de4 100644 --- a/.github/workflows/pr.yaml +++ b/.github/workflows/pr.yaml @@ -35,7 +35,6 @@ jobs: secrets: inherit uses: rapidsai/shared-action-workflows/.github/workflows/conda-cpp-build.yaml@branch-23.08 with: - matrix_filter: map(select(.CUDA_VER | startswith("11"))) build_type: pull-request node_type: cpu16 conda-cpp-tests: @@ -43,21 +42,18 @@ jobs: secrets: inherit uses: rapidsai/shared-action-workflows/.github/workflows/conda-cpp-tests.yaml@branch-23.08 with: - matrix_filter: map(select(.CUDA_VER | startswith("11"))) build_type: pull-request conda-python-build: needs: conda-cpp-build secrets: inherit uses: rapidsai/shared-action-workflows/.github/workflows/conda-python-build.yaml@branch-23.08 with: - matrix_filter: map(select(.CUDA_VER | startswith("11"))) build_type: pull-request conda-python-tests: needs: conda-python-build secrets: inherit uses: rapidsai/shared-action-workflows/.github/workflows/conda-python-tests.yaml@branch-23.08 with: - matrix_filter: map(select(.CUDA_VER | startswith("11"))) build_type: pull-request conda-notebook-tests: needs: conda-python-build @@ -67,7 +63,7 @@ jobs: build_type: pull-request node_type: "gpu-v100-latest-1" arch: "amd64" - container_image: "rapidsai/ci:cuda11.8.0-ubuntu22.04-py3.10" + container_image: "rapidsai/ci:latest" run_script: "ci/test_notebooks.sh" docs-build: needs: conda-python-build @@ -77,7 +73,7 @@ jobs: build_type: pull-request node_type: "gpu-v100-latest-1" arch: "amd64" - container_image: "rapidsai/ci:cuda11.8.0-ubuntu22.04-py3.10" + container_image: "rapidsai/ci:latest" run_script: "ci/build_docs.sh" wheel-build-pylibcugraph: needs: checks diff --git a/.github/workflows/test.yaml b/.github/workflows/test.yaml index 1b8cfaf25b7..d697b8f1649 100644 --- a/.github/workflows/test.yaml +++ b/.github/workflows/test.yaml @@ -18,7 +18,6 @@ jobs: secrets: inherit uses: rapidsai/shared-action-workflows/.github/workflows/conda-cpp-tests.yaml@branch-23.08 with: - matrix_filter: map(select(.CUDA_VER | startswith("11"))) build_type: nightly branch: ${{ inputs.branch }} date: ${{ inputs.date }} @@ -27,7 +26,6 @@ jobs: secrets: inherit uses: rapidsai/shared-action-workflows/.github/workflows/conda-python-tests.yaml@branch-23.08 with: - matrix_filter: map(select(.CUDA_VER | startswith("11"))) build_type: nightly branch: ${{ inputs.branch }} date: ${{ inputs.date }} diff --git a/ci/build_python.sh b/ci/build_python.sh index 517dda726ee..5125e86d53a 100755 --- a/ci/build_python.sh +++ b/ci/build_python.sh @@ -26,28 +26,36 @@ rapids-mamba-retry mambabuild \ --channel "${RAPIDS_CONDA_BLD_OUTPUT_DIR}" \ conda/recipes/cugraph -rapids-mamba-retry mambabuild \ - --no-test \ - --channel "${CPP_CHANNEL}" \ - --channel "${RAPIDS_CONDA_BLD_OUTPUT_DIR}" \ - conda/recipes/cugraph-service - -rapids-mamba-retry mambabuild \ - --no-test \ - --channel "${CPP_CHANNEL}" \ - --channel "${RAPIDS_CONDA_BLD_OUTPUT_DIR}" \ - --channel pyg \ - --channel pytorch \ - --channel pytorch-nightly \ - conda/recipes/cugraph-pyg - -rapids-mamba-retry mambabuild \ - --no-test \ - --channel "${CPP_CHANNEL}" \ - --channel "${RAPIDS_CONDA_BLD_OUTPUT_DIR}" \ - --channel dglteam \ - --channel pytorch \ - --channel pytorch-nightly \ - conda/recipes/cugraph-dgl +RAPIDS_CUDA_MAJOR="${RAPIDS_CUDA_VERSION%%.*}" + +if [[ ${RAPIDS_CUDA_MAJOR} == "11" ]]; then + # Only one CUDA configuration is needed, so we choose CUDA 11 arbitrarily. + # Nothing in the cugraph-service packages is CUDA-specific. + rapids-mamba-retry mambabuild \ + --no-test \ + --channel "${CPP_CHANNEL}" \ + --channel "${RAPIDS_CONDA_BLD_OUTPUT_DIR}" \ + conda/recipes/cugraph-service + + # Only CUDA 11 is supported right now due to PyTorch requirement. + rapids-mamba-retry mambabuild \ + --no-test \ + --channel "${CPP_CHANNEL}" \ + --channel "${RAPIDS_CONDA_BLD_OUTPUT_DIR}" \ + --channel pyg \ + --channel pytorch \ + --channel pytorch-nightly \ + conda/recipes/cugraph-pyg + + # Only CUDA 11 is supported right now due to PyTorch requirement. + rapids-mamba-retry mambabuild \ + --no-test \ + --channel "${CPP_CHANNEL}" \ + --channel "${RAPIDS_CONDA_BLD_OUTPUT_DIR}" \ + --channel dglteam \ + --channel pytorch \ + --channel pytorch-nightly \ + conda/recipes/cugraph-dgl +fi rapids-upload-conda-to-s3 python diff --git a/conda/environments/all_cuda-118_arch-x86_64.yaml b/conda/environments/all_cuda-118_arch-x86_64.yaml index 1540a85bcc8..16a4d4f0dbc 100644 --- a/conda/environments/all_cuda-118_arch-x86_64.yaml +++ b/conda/environments/all_cuda-118_arch-x86_64.yaml @@ -12,7 +12,8 @@ dependencies: - aiohttp - c-compiler - cmake>=3.26.4 -- cudatoolkit=11.8 +- cuda-version=11.8 +- cudatoolkit - cudf==23.8.* - cupy>=12.0.0 - cxx-compiler diff --git a/conda/recipes/cugraph/conda_build_config.yaml b/conda/recipes/cugraph/conda_build_config.yaml index e3e18d4620b..4530a4c942d 100644 --- a/conda/recipes/cugraph/conda_build_config.yaml +++ b/conda/recipes/cugraph/conda_build_config.yaml @@ -5,6 +5,9 @@ cxx_compiler_version: - 11 cuda_compiler: + - cuda-nvcc + +cuda11_compiler: - nvcc cmake_version: @@ -15,18 +18,3 @@ sysroot_version: ucx_py_version: - "0.33.*" - -# The CTK libraries below are missing from the conda-forge::cudatoolkit -# package. The "*_host_*" version specifiers correspond to `11.8` packages. - -libcublas_host_version: - - "=11.11.3.6" - -libcurand_host_version: - - "=10.3.0.86" - -libcusolver_host_version: - - "=11.4.1.48" - -libcusparse_host_version: - - "=11.7.5.86" diff --git a/conda/recipes/cugraph/meta.yaml b/conda/recipes/cugraph/meta.yaml index 7f2dc542538..e2b9d38c181 100644 --- a/conda/recipes/cugraph/meta.yaml +++ b/conda/recipes/cugraph/meta.yaml @@ -39,36 +39,38 @@ build: requirements: build: - {{ compiler('c') }} - - {{ compiler('cuda') }} {{ cuda_version }} - {{ compiler('cxx') }} + {% if cuda_major == "11" %} + - {{ compiler('cuda11') }} {{ cuda_version }} + {% else %} + - {{ compiler('cuda') }} + {% endif %} + - cuda-version ={{ cuda_version }} - cmake {{ cmake_version }} - ninja - sysroot_{{ target_platform }} {{ sysroot_version }} host: - - cudatoolkit ={{ cuda_version }} + - cuda-version ={{ cuda_version }} + {% if cuda_major == "11" %} + - cudatoolkit + {% endif %} - cudf ={{ minor_version }} - cython >=0.29,<0.30 - - libcublas {{ libcublas_host_version }} - - libcublas-dev {{ libcublas_host_version }} - libcugraph ={{ version }} - - libcurand {{ libcurand_host_version }} - - libcurand-dev {{ libcurand_host_version }} - - libcusolver {{ libcusolver_host_version }} - - libcusolver-dev {{ libcusolver_host_version }} - - libcusparse {{ libcusparse_host_version }} - - libcusparse-dev {{ libcusparse_host_version }} - - libraft ={{ minor_version }} - - libraft-headers ={{ minor_version }} - - pylibraft ={{ minor_version}} + - pylibraft ={{ minor_version }} - python - raft-dask ={{ minor_version }} + - rmm ={{ minor_version }} - scikit-build >=0.13.1 - setuptools - - ucx-proc=*=gpu - - ucx-py {{ ucx_py_version }} run: - - {{ pin_compatible('cudatoolkit', max_pin='x', min_pin='x') }} - - cuda-python >=11.7.1,<12.0 + - {{ pin_compatible('cuda-version', max_pin='x', min_pin='x') }} + {% if cuda_major == "11" %} + - cudatoolkit + - cuda-python >=11.7.1,<12.0a0 + {% else %} + - cuda-python >=12.0,<13.0a0 + {% endif %} - cudf ={{ minor_version }} - cupy >=12.0.0 - dask-cuda ={{ minor_version }} @@ -77,8 +79,6 @@ requirements: - dask-core >=2023.5.1 - distributed >=2023.5.1 - libcugraph ={{ version }} - - libraft ={{ minor_version }} - - libraft-headers ={{ minor_version }} - pylibcugraph ={{ version }} - pylibraft ={{ minor_version }} - python @@ -88,7 +88,7 @@ requirements: tests: requirements: - - cudatoolkit ={{ cuda_version }} + - cuda-version ={{ cuda_version }} imports: - cugraph diff --git a/conda/recipes/libcugraph/conda_build_config.yaml b/conda/recipes/libcugraph/conda_build_config.yaml index 287d1d391da..3acd0166a67 100644 --- a/conda/recipes/libcugraph/conda_build_config.yaml +++ b/conda/recipes/libcugraph/conda_build_config.yaml @@ -5,6 +5,9 @@ cxx_compiler_version: - 11 cuda_compiler: + - cuda-nvcc + +cuda11_compiler: - nvcc cmake_version: @@ -19,23 +22,43 @@ nccl_version: gtest_version: - ">=1.13.0" -cuda_profiler_api_version: - - ">=11.8.86,<12" - sysroot_version: - "2.17" # The CTK libraries below are missing from the conda-forge::cudatoolkit -# package. The "*_host_*" version specifiers correspond to `11.8` packages. +# package. The "*_host_*" version specifiers correspond to `11.8` packages +# and the "*_run_*" version specifiers correspond to `11.x` packages. -libcublas_host_version: +cuda11_libcublas_host_version: - "=11.11.3.6" -libcurand_host_version: +cuda11_libcublas_run_version: + - ">=11.5.2.43,<12.0.0" + +cuda11_libcurand_host_version: - "=10.3.0.86" -libcusolver_host_version: +cuda11_libcurand_run_version: + - ">=10.2.5.43,<10.3.1" + +cuda11_libcusolver_host_version: - "=11.4.1.48" -libcusparse_host_version: +cuda11_libcusolver_run_version: + - ">=11.2.0.43,<11.4.2" + +cuda11_libcusparse_host_version: - "=11.7.5.86" + +cuda11_libcusparse_run_version: + - ">=11.6.0.43,<12.0.0" + +# `cuda-profiler-api` only has `11.8.0` and `12.0.0` packages for all +# architectures. The "*_host_*" version specifiers correspond to `11.8` packages and the +# "*_run_*" version specifiers correspond to `11.x` packages. + +cuda11_cuda_profiler_api_host_version: + - "=11.8.86" + +cuda11_cuda_profiler_api_run_version: + - ">=11.4.240,<12" diff --git a/conda/recipes/libcugraph/meta.yaml b/conda/recipes/libcugraph/meta.yaml index f843aabba92..d52d81366d7 100644 --- a/conda/recipes/libcugraph/meta.yaml +++ b/conda/recipes/libcugraph/meta.yaml @@ -4,7 +4,6 @@ {% set minor_version = version.split('.')[0] + '.' + version.split('.')[1] %} {% set cuda_version = '.'.join(environ['RAPIDS_CUDA_VERSION'].split('.')[:2]) %} {% set cuda_major = cuda_version.split('.')[0] %} -{% set cuda_spec = ">=" + cuda_major ~ ",<" + (cuda_major | int + 1) ~ ".0a0" %} # i.e. >=11,<12.0a0 {% set date_string = environ['RAPIDS_DATE_STRING'] %} package: @@ -34,29 +33,45 @@ build: requirements: build: - {{ compiler('c') }} - - {{ compiler('cuda') }} {{ cuda_version }} + {% if cuda_major == "11" %} + - {{ compiler('cuda11') }} {{ cuda_version }} + {% else %} + - {{ compiler('cuda') }} + {% endif %} + - cuda-version ={{ cuda_version }} - {{ compiler('cxx') }} - cmake {{ cmake_version }} - ninja - openmpi # Required for building cpp-mgtests (multi-GPU tests) - sysroot_{{ target_platform }} {{ sysroot_version }} host: + {% if cuda_major == "11" %} + - cudatoolkit - cuda-nvtx ={{ cuda_version }} - - cuda-profiler-api {{ cuda_profiler_api_version }} - - cudatoolkit ={{ cuda_version }} + - cuda-profiler-api {{ cuda11_cuda_profiler_api_host_version }} + - libcublas {{ cuda11_libcublas_host_version }} + - libcublas-dev {{ cuda11_libcublas_host_version }} + - libcurand {{ cuda11_libcurand_host_version }} + - libcurand-dev {{ cuda11_libcurand_host_version }} + - libcusolver {{ cuda11_libcusolver_host_version }} + - libcusolver-dev {{ cuda11_libcusolver_host_version }} + - libcusparse {{ cuda11_libcusparse_host_version }} + - libcusparse-dev {{ cuda11_libcusparse_host_version }} + {% else %} + - cuda-nvtx-dev + - cuda-profiler-api + - cuda-cudart-dev + - libcublas-dev + - libcurand-dev + - libcusolver-dev + - libcusparse-dev + {% endif %} + - cuda-version ={{ cuda_version }} - doxygen {{ doxygen_version }} - gmock {{ gtest_version }} - gtest {{ gtest_version }} - - libcublas {{ libcublas_host_version }} - - libcublas-dev {{ libcublas_host_version }} - libcudf ={{ minor_version }} - libcugraphops ={{ minor_version }} - - libcurand {{ libcurand_host_version }} - - libcurand-dev {{ libcurand_host_version }} - - libcusolver {{ libcusolver_host_version }} - - libcusolver-dev {{ libcusolver_host_version }} - - libcusparse {{ libcusparse_host_version }} - - libcusparse-dev {{ libcusparse_host_version }} - libraft ={{ minor_version }} - libraft-headers ={{ minor_version }} - librmm ={{ minor_version }} @@ -71,12 +86,30 @@ outputs: number: {{ GIT_DESCRIBE_NUMBER }} string: cuda{{ cuda_major }}_{{ date_string }}_{{ GIT_DESCRIBE_HASH }}_{{ GIT_DESCRIBE_NUMBER }} ignore_run_exports_from: - - {{ compiler('cuda') }} + {% if cuda_major == "11" %} + - {{ compiler('cuda11') }} + {% endif %} requirements: build: - cmake {{ cmake_version }} + host: + - cuda-version ={{ cuda_version }} run: - - cudatoolkit {{ cuda_spec }} + - {{ pin_compatible('cuda-version', max_pin='x', min_pin='x') }} + {% if cuda_major == "11" %} + - cuda-profiler-api {{ cuda11_cuda_profiler_api_run_version }} + - cudatoolkit + - libcublas {{ cuda11_libcublas_run_version }} + - libcurand {{ cuda11_libcurand_run_version }} + - libcusolver {{ cuda11_libcusolver_run_version }} + - libcusparse {{ cuda11_libcusparse_run_version }} + {% else %} + - cuda-profiler-api + - libcublas + - libcurand + - libcusolver + - libcusparse + {% endif %} - libcugraphops ={{ minor_version }} - libraft ={{ minor_version }} - libraft-headers ={{ minor_version }} @@ -95,13 +128,20 @@ outputs: number: {{ GIT_DESCRIBE_NUMBER }} string: cuda{{ cuda_major }}_{{ date_string }}_{{ GIT_DESCRIBE_HASH }}_{{ GIT_DESCRIBE_NUMBER }} ignore_run_exports_from: - - {{ compiler('cuda') }} + {% if cuda_major == "11" %} + - {{ compiler('cuda11') }} + {% endif %} requirements: build: - cmake {{ cmake_version }} + host: + - cuda-version ={{ cuda_version }} run: + - {{ pin_compatible('cuda-version', max_pin='x', min_pin='x') }} - {{ pin_subpackage('libcugraph', exact=True) }} - - cudatoolkit {{ cuda_spec }} + {% if cuda_major == "11" %} + - cudatoolkit + {% endif %} - libcudf ={{ minor_version }} - librmm ={{ minor_version }} about: @@ -116,14 +156,21 @@ outputs: number: {{ GIT_DESCRIBE_NUMBER }} string: cuda{{ cuda_major }}_{{ date_string }}_{{ GIT_DESCRIBE_HASH }}_{{ GIT_DESCRIBE_NUMBER }} ignore_run_exports_from: - - {{ compiler('cuda') }} + {% if cuda_major == "11" %} + - {{ compiler('cuda11') }} + {% endif %} requirements: build: - cmake {{ cmake_version }} + host: + - cuda-version ={{ cuda_version }} run: + - {{ pin_compatible('cuda-version', max_pin='x', min_pin='x') }} - {{ pin_subpackage('libcugraph_etl', exact=True) }} - {{ pin_subpackage('libcugraph', exact=True) }} - - cudatoolkit {{ cuda_spec }} + {% if cuda_major == "11" %} + - cudatoolkit + {% endif %} - gmock {{ gtest_version }} - gtest {{ gtest_version }} about: diff --git a/conda/recipes/pylibcugraph/conda_build_config.yaml b/conda/recipes/pylibcugraph/conda_build_config.yaml index e3e18d4620b..4530a4c942d 100644 --- a/conda/recipes/pylibcugraph/conda_build_config.yaml +++ b/conda/recipes/pylibcugraph/conda_build_config.yaml @@ -5,6 +5,9 @@ cxx_compiler_version: - 11 cuda_compiler: + - cuda-nvcc + +cuda11_compiler: - nvcc cmake_version: @@ -15,18 +18,3 @@ sysroot_version: ucx_py_version: - "0.33.*" - -# The CTK libraries below are missing from the conda-forge::cudatoolkit -# package. The "*_host_*" version specifiers correspond to `11.8` packages. - -libcublas_host_version: - - "=11.11.3.6" - -libcurand_host_version: - - "=10.3.0.86" - -libcusolver_host_version: - - "=11.4.1.48" - -libcusparse_host_version: - - "=11.7.5.86" diff --git a/conda/recipes/pylibcugraph/meta.yaml b/conda/recipes/pylibcugraph/meta.yaml index de031a6fe94..aa82c20ad44 100644 --- a/conda/recipes/pylibcugraph/meta.yaml +++ b/conda/recipes/pylibcugraph/meta.yaml @@ -39,41 +39,38 @@ build: requirements: build: - {{ compiler('c') }} - - {{ compiler('cuda') }} {{ cuda_version }} - {{ compiler('cxx') }} + {% if cuda_major == "11" %} + - {{ compiler('cuda11') }} {{ cuda_version }} + {% else %} + - {{ compiler('cuda') }} + {% endif %} + - cuda-version ={{ cuda_version }} - cmake {{ cmake_version }} - ninja - sysroot_{{ target_platform }} {{ sysroot_version }} host: - - cudatoolkit ={{ cuda_version }} - - cudf ={{ minor_version }} + - cuda-version ={{ cuda_version }} + {% if cuda_major == "11" %} + - cudatoolkit + {% endif %} - cython >=0.29,<0.30 - - libcublas {{ libcublas_host_version }} - - libcublas-dev {{ libcublas_host_version }} - libcugraph ={{ version }} - - libcurand {{ libcurand_host_version }} - - libcurand-dev {{ libcurand_host_version }} - - libcusolver {{ libcusolver_host_version }} - - libcusolver-dev {{ libcusolver_host_version }} - - libcusparse {{ libcusparse_host_version }} - - libcusparse-dev {{ libcusparse_host_version }} - - libraft ={{ minor_version }} - - libraft-headers ={{ minor_version }} - - pylibraft ={{ minor_version}} + - pylibraft ={{ minor_version }} - python - - rmm ={{ minor_version }} - scikit-build >=0.13.1 - setuptools - - ucx-proc=*=gpu - - ucx-py {{ ucx_py_version }} run: - - {{ pin_compatible('cudatoolkit', max_pin='x', min_pin='x') }} + - {{ pin_compatible('cuda-version', max_pin='x', min_pin='x') }} + {% if cuda_major == "11" %} + - cudatoolkit + {% endif %} - libcugraph ={{ version }} - python tests: requirements: - - cudatoolkit ={{ cuda_version }} + - cuda-version ={{ cuda_version }} imports: - pylibcugraph diff --git a/dependencies.yaml b/dependencies.yaml index d6e36a9643c..572638069dc 100644 --- a/dependencies.yaml +++ b/dependencies.yaml @@ -3,7 +3,7 @@ files: all: output: [conda] matrix: - cuda: ["11.8"] + cuda: ["11.8", "12.0"] arch: [x86_64] includes: - checks @@ -195,22 +195,30 @@ dependencies: specific: - output_types: [conda] matrices: + - matrix: + cuda: "12.0" + packages: + - cuda-version=12.0 - matrix: cuda: "11.8" packages: - - cudatoolkit=11.8 + - cuda-version=11.8 + - cudatoolkit - matrix: cuda: "11.5" packages: - - cudatoolkit=11.5 + - cuda-version=11.5 + - cudatoolkit - matrix: cuda: "11.4" packages: - - cudatoolkit=11.4 + - cuda-version=11.4 + - cudatoolkit - matrix: cuda: "11.2" packages: - - cudatoolkit=11.2 + - cuda-version=11.2 + - cudatoolkit common_build: common: - output_types: [conda, pyproject] @@ -253,6 +261,11 @@ dependencies: cuda: "11.8" packages: - nvcc_linux-aarch64=11.8 + - matrix: + cuda: "12.0" + packages: + - cuda-version=12.0 + - cuda-nvcc docs: common: - output_types: [conda] From 803b854fd43be7375d6354ed005001725c7710b4 Mon Sep 17 00:00:00 2001 From: Chuck Hastings <45364586+ChuckHastings@users.noreply.github.com> Date: Thu, 20 Jul 2023 10:38:09 -0400 Subject: [PATCH 11/14] Fix primitive bug discovered in MG edge betweenness centrality testing (#3723) @jnke2016 found a bug in MG edge betweenness centrality. Upon investigation, the `major_offset` in this block of code wasn't being set correctly if `major` is in the hypersparse region. Authors: - Chuck Hastings (https://github.com/ChuckHastings) Approvers: - Seunghwa Kang (https://github.com/seunghwak) - Joseph Nke (https://github.com/jnke2016) - Naim (https://github.com/naimnv) URL: https://github.com/rapidsai/cugraph/pull/3723 --- cpp/src/prims/transform_e.cuh | 25 +++++++++---------------- 1 file changed, 9 insertions(+), 16 deletions(-) diff --git a/cpp/src/prims/transform_e.cuh b/cpp/src/prims/transform_e.cuh index 9be12262574..7950df58a3e 100644 --- a/cpp/src/prims/transform_e.cuh +++ b/cpp/src/prims/transform_e.cuh @@ -291,23 +291,16 @@ void transform_e(raft::handle_t const& handle, auto major = thrust::get<0>(edge); auto minor = thrust::get<1>(edge); - vertex_t major_offset{}; - vertex_t major_idx{}; auto major_hypersparse_first = edge_partition.major_hypersparse_first(); - if (major_hypersparse_first) { - if (major < *major_hypersparse_first) { - major_offset = edge_partition.major_offset_from_major_nocheck(major); - major_idx = major_offset; - } else { - auto major_hypersparse_idx = - edge_partition.major_hypersparse_idx_from_major_nocheck(major); - assert(major_hypersparse_idx); - major_idx = edge_partition.major_offset_from_major_nocheck(*major_hypersparse_first) + - *major_hypersparse_idx; - } - } else { - major_offset = edge_partition.major_offset_from_major_nocheck(major); - major_idx = major_offset; + auto major_offset = edge_partition.major_offset_from_major_nocheck(major); + vertex_t major_idx{major_offset}; + + if ((major_hypersparse_first) && (major >= *major_hypersparse_first)) { + auto major_hypersparse_idx = + edge_partition.major_hypersparse_idx_from_major_nocheck(major); + assert(major_hypersparse_idx); + major_idx = edge_partition.major_offset_from_major_nocheck(*major_hypersparse_first) + + *major_hypersparse_idx; } auto minor_offset = edge_partition.minor_offset_from_minor_nocheck(minor); From 54a30b933922cc0dbef7c0fa0d98b49889bcbbe8 Mon Sep 17 00:00:00 2001 From: Vyas Ramasubramani Date: Fri, 21 Jul 2023 08:39:39 -0700 Subject: [PATCH 12/14] Stop using setup.py in build.sh (#3704) This PR changes replaces direct invocations of setup.py in build.sh with `pip install` commands since setup.py usage is deprecated. This change will facilitate upcoming changes to our build systems. Authors: - Vyas Ramasubramani (https://github.com/vyasr) Approvers: - Rick Ratzel (https://github.com/rlratzel) URL: https://github.com/rapidsai/cugraph/pull/3704 --- build.sh | 65 ++++++++++++++++---------------------------------------- 1 file changed, 18 insertions(+), 47 deletions(-) diff --git a/build.sh b/build.sh index a8e97d924c6..18359229822 100755 --- a/build.sh +++ b/build.sh @@ -60,7 +60,7 @@ HELP="$0 [ ...] [ ...] and is: -v - verbose build mode -g - build for debug - -n - do not install after a successful build + -n - do not install after a successful build (does not affect Python packages) --pydevelop - use setup.py develop instead of install --allgpuarch - build for all supported GPU architectures --skip_cpp_tests - do not build the SG test binaries as part of the libcugraph and libcugraph_etl targets @@ -104,7 +104,7 @@ BUILD_CPP_MG_TESTS=OFF BUILD_ALL_GPU_ARCH=0 BUILD_WITH_CUGRAPHOPS=ON CMAKE_GENERATOR_OPTION="-G Ninja" -PYTHON_ARGS_FOR_INSTALL="-m pip install --no-build-isolation --no-deps ." +PYTHON_ARGS_FOR_INSTALL="-m pip install --no-build-isolation --no-deps" # Set defaults for vars that may not have been defined externally # FIXME: if PREFIX is not set, check CONDA_PREFIX, but there is no fallback @@ -178,6 +178,12 @@ if hasArg --pydevelop; then PYTHON_ARGS_FOR_INSTALL="-m pip install --no-build-isolation --no-deps -e ." fi +# Append `-DFIND_RAFT_CPP=ON` to EXTRA_CMAKE_ARGS unless a user specified the option. +SKBUILD_EXTRA_CMAKE_ARGS="${EXTRA_CMAKE_ARGS}" + if [[ "${EXTRA_CMAKE_ARGS}" != *"DFIND_CUGRAPH_CPP"* ]]; then + SKBUILD_EXTRA_CMAKE_ARGS="${SKBUILD_EXTRA_CMAKE_ARGS} -DFIND_CUGRAPH_CPP=ON" + fi + # If clean or uninstall targets given, run them prior to any other steps if hasArg uninstall; then if [[ "$INSTALL_PREFIX" != "" ]]; then @@ -296,21 +302,9 @@ if buildAll || hasArg pylibcugraph; then if hasArg --clean; then cleanPythonDir ${REPODIR}/python/pylibcugraph else - cd ${REPODIR}/python/pylibcugraph - # setup.py references an env var CUGRAPH_BUILD_PATH to find the libcugraph - # build. If not set by the user, set it to LIBCUGRAPH_BUILD_DIR - CUGRAPH_BUILD_PATH=${CUGRAPH_BUILD_PATH:=${LIBCUGRAPH_BUILD_DIR}} - python setup.py build_ext \ - --inplace \ - -- \ - -DFIND_CUGRAPH_CPP=ON \ - -DUSE_CUGRAPH_OPS=${BUILD_WITH_CUGRAPHOPS} \ - -Dcugraph_ROOT=${LIBCUGRAPH_BUILD_DIR} \ - -- \ - -j${PARALLEL_LEVEL:-1} - if [[ ${INSTALL_TARGET} != "" ]]; then - env CUGRAPH_BUILD_PATH=${CUGRAPH_BUILD_PATH} python ${PYTHON_ARGS_FOR_INSTALL} - fi + SKBUILD_CONFIGURE_OPTIONS="${SKBUILD_EXTRA_CMAKE_ARGS} -DUSE_CUGRAPH_OPS=${BUILD_WITH_CUGRAPHOPS}" \ + SKBUILD_BUILD_OPTIONS="-j${PARALLEL_LEVEL}" \ + python ${PYTHON_ARGS_FOR_INSTALL} ${REPODIR}/python/pylibcugraph fi fi @@ -319,22 +313,9 @@ if buildAll || hasArg cugraph; then if hasArg --clean; then cleanPythonDir ${REPODIR}/python/cugraph else - cd ${REPODIR}/python/cugraph - # FIXME: this needs to eventually reference the pylibcugraph build - # setup.py references an env var CUGRAPH_BUILD_PATH to find the libcugraph - # build. If not set by the user, set it to LIBCUGRAPH_BUILD_DIR - CUGRAPH_BUILD_PATH=${CUGRAPH_BUILD_PATH:=${LIBCUGRAPH_BUILD_DIR}} - python setup.py build_ext \ - --inplace \ - -- \ - -DFIND_CUGRAPH_CPP=ON \ - -DUSE_CUGRAPH_OPS=${BUILD_WITH_CUGRAPHOPS} \ - -Dcugraph_ROOT=${LIBCUGRAPH_BUILD_DIR} \ - -- \ - -j${PARALLEL_LEVEL:-1} - if [[ ${INSTALL_TARGET} != "" ]]; then - env CUGRAPH_BUILD_PATH=${CUGRAPH_BUILD_PATH} python ${PYTHON_ARGS_FOR_INSTALL} - fi + SKBUILD_CONFIGURE_OPTIONS="${SKBUILD_EXTRA_CMAKE_ARGS} -DUSE_CUGRAPH_OPS=${BUILD_WITH_CUGRAPHOPS}" \ + SKBUILD_BUILD_OPTIONS="-j${PARALLEL_LEVEL}" \ + python ${PYTHON_ARGS_FOR_INSTALL} ${REPODIR}/python/cugraph fi fi @@ -343,12 +324,8 @@ if hasArg cugraph-service; then if hasArg --clean; then cleanPythonDir ${REPODIR}/python/cugraph-service else - if [[ ${INSTALL_TARGET} != "" ]]; then - cd ${REPODIR}/python/cugraph-service/client - python ${PYTHON_ARGS_FOR_INSTALL} - cd ${REPODIR}/python/cugraph-service/server - python ${PYTHON_ARGS_FOR_INSTALL} - fi + python ${PYTHON_ARGS_FOR_INSTALL} ${REPODIR}/python/cugraph-service/client + python ${PYTHON_ARGS_FOR_INSTALL} ${REPODIR}/python/cugraph-service/server fi fi @@ -357,10 +334,7 @@ if hasArg cugraph-pyg; then if hasArg --clean; then cleanPythonDir ${REPODIR}/python/cugraph-pyg else - if [[ ${INSTALL_TARGET} != "" ]]; then - cd ${REPODIR}/python/cugraph-pyg - python ${PYTHON_ARGS_FOR_INSTALL} - fi + python ${PYTHON_ARGS_FOR_INSTALL} ${REPODIR}/python/cugraph-pyg fi fi @@ -369,10 +343,7 @@ if hasArg cugraph-dgl; then if hasArg --clean; then cleanPythonDir ${REPODIR}/python/cugraph-dgl else - if [[ ${INSTALL_TARGET} != "" ]]; then - cd ${REPODIR}/python/cugraph-dgl - python ${PYTHON_ARGS_FOR_INSTALL} - fi + python ${PYTHON_ARGS_FOR_INSTALL} ${REPODIR}/python/cugraph-dgl fi fi From cab836896aeba2c5549c57ebfd4818e883c17b5c Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Fri, 21 Jul 2023 12:09:24 -0500 Subject: [PATCH 13/14] Remove RAFT specialization. (#3727) This PR removes a deprecation warning coming from RAFT. Including the specializations header no longer has an effect, so it can be safely removed. Authors: - Bradley Dice (https://github.com/bdice) Approvers: - Chuck Hastings (https://github.com/ChuckHastings) URL: https://github.com/rapidsai/cugraph/pull/3727 --- cpp/src/community/legacy/spectral_clustering.cu | 3 --- 1 file changed, 3 deletions(-) diff --git a/cpp/src/community/legacy/spectral_clustering.cu b/cpp/src/community/legacy/spectral_clustering.cu index 5ab71e7d280..84e9891b7ce 100644 --- a/cpp/src/community/legacy/spectral_clustering.cu +++ b/cpp/src/community/legacy/spectral_clustering.cu @@ -24,9 +24,6 @@ #include #include -#if defined RAFT_COMPILED -#include -#endif #include #include From 487efd6d69bdd67c5ea03be56bda2d7e73005431 Mon Sep 17 00:00:00 2001 From: Chuck Hastings <45364586+ChuckHastings@users.noreply.github.com> Date: Fri, 21 Jul 2023 17:28:01 -0400 Subject: [PATCH 14/14] Fix sampling call parameters if compiled with -DNO_CUGRAPH_OPS (#3729) Updated sampling call in the `#ifdef NO_CUGRAPH_OPS` branch of the SG unit test. The MG unit test was correctly updated previously. Fixes #3726 Authors: - Chuck Hastings (https://github.com/ChuckHastings) Approvers: - Seunghwa Kang (https://github.com/seunghwak) - Bradley Dice (https://github.com/bdice) URL: https://github.com/rapidsai/cugraph/pull/3729 --- .../sampling/mg_uniform_neighbor_sampling.cu | 1 - .../sampling/sg_uniform_neighbor_sampling.cu | 36 ++++++++++--------- 2 files changed, 20 insertions(+), 17 deletions(-) diff --git a/cpp/tests/sampling/mg_uniform_neighbor_sampling.cu b/cpp/tests/sampling/mg_uniform_neighbor_sampling.cu index 82fb2430ca1..57f85a212b1 100644 --- a/cpp/tests/sampling/mg_uniform_neighbor_sampling.cu +++ b/cpp/tests/sampling/mg_uniform_neighbor_sampling.cu @@ -163,7 +163,6 @@ class Tests_MGUniform_Neighbor_Sampling EXPECT_THROW( cugraph::uniform_neighbor_sample( *handle_, - handle, mg_graph_view, mg_edge_weight_view, std::optional>{std::nullopt}, diff --git a/cpp/tests/sampling/sg_uniform_neighbor_sampling.cu b/cpp/tests/sampling/sg_uniform_neighbor_sampling.cu index a59ea7feb8f..f795c11437f 100644 --- a/cpp/tests/sampling/sg_uniform_neighbor_sampling.cu +++ b/cpp/tests/sampling/sg_uniform_neighbor_sampling.cu @@ -128,25 +128,29 @@ class Tests_Uniform_Neighbor_Sampling random_sources.size(), handle.get_stream()); -#ifdef NO_CUGRAPH_OPS - EXPECT_THROW(cugraph::uniform_neighbor_sample( - handle, - graph_view, - edge_weight_view, - std::nullopt, - std::nullopt, - std::move(random_sources_copy), - std::move(batch_number), - raft::host_span(uniform_neighbor_sampling_usecase.fanout.data(), - uniform_neighbor_sampling_usecase.fanout.size()), - rng_state, - true, - uniform_neighbor_sampling_usecase.flag_replacement), - std::exception); -#else std::optional, raft::device_span>> label_to_output_comm_rank_mapping{std::nullopt}; +#ifdef NO_CUGRAPH_OPS + EXPECT_THROW( + cugraph::uniform_neighbor_sample( + handle, + graph_view, + edge_weight_view, + std::optional>{std::nullopt}, + std::optional>{std::nullopt}, + raft::device_span{random_sources_copy.data(), random_sources.size()}, + batch_number ? std::make_optional(raft::device_span{batch_number->data(), + batch_number->size()}) + : std::nullopt, + label_to_output_comm_rank_mapping, + raft::host_span(uniform_neighbor_sampling_usecase.fanout.data(), + uniform_neighbor_sampling_usecase.fanout.size()), + rng_state, + true, + uniform_neighbor_sampling_usecase.flag_replacement), + std::exception); +#else auto&& [src_out, dst_out, wgt_out, edge_id, edge_type, hop, labels, offsets] = cugraph::uniform_neighbor_sample( handle,