diff --git a/cpp/include/cugraph/graph_view.hpp b/cpp/include/cugraph/graph_view.hpp index f30a8b7e2af..d79d4635c54 100644 --- a/cpp/include/cugraph/graph_view.hpp +++ b/cpp/include/cugraph/graph_view.hpp @@ -268,7 +268,11 @@ class graph_base_t { properties_(properties){}; vertex_t number_of_vertices() const { return number_of_vertices_; } - edge_t number_of_edges() const { return number_of_edges_; } + edge_t number_of_edges() const + { + CUGRAPH_EXPECTS(!(this->has_edge_mask()), "unimplemented."); + return number_of_edges_; + } template std::enable_if_t::value, bool> is_valid_vertex(vertex_type v) const @@ -285,6 +289,20 @@ class graph_base_t { bool is_symmetric() const { return properties_.is_symmetric; } bool is_multigraph() const { return properties_.is_multigraph; } + void attach_edge_mask(edge_property_view_t edge_mask_view) + { + edge_mask_view_ = edge_mask_view; + } + + void clear_edge_mask() { edge_mask_view_ = std::nullopt; } + + bool has_edge_mask() const { return edge_mask_view_.has_value(); } + + std::optional> edge_mask_view() const + { + return edge_mask_view_; + } + protected: raft::handle_t const* handle_ptr() const { return handle_ptr_; }; graph_properties_t graph_properties() const { return properties_; } @@ -296,6 +314,8 @@ class graph_base_t { edge_t number_of_edges_{0}; graph_properties_t properties_{}; + + std::optional> edge_mask_view_{std::nullopt}; }; } // namespace detail @@ -731,20 +751,6 @@ class graph_view_t edge_mask_view) - { - edge_mask_view_ = edge_mask_view; - } - - void clear_edge_mask() { edge_mask_view_ = std::nullopt; } - - bool has_edge_mask() const { return edge_mask_view_.has_value(); } - - std::optional> edge_mask_view() const - { - return edge_mask_view_; - } - private: std::vector edge_partition_offsets_{}; std::vector edge_partition_indices_{}; @@ -790,8 +796,6 @@ class graph_view_t>, std::optional /* dummy */> local_sorted_unique_edge_dst_vertex_partition_offsets_{std::nullopt}; - - std::optional> edge_mask_view_{std::nullopt}; }; // single-GPU version @@ -1012,28 +1016,12 @@ class graph_view_t edge_mask_view) - { - edge_mask_view_ = edge_mask_view; - } - - void clear_edge_mask() { edge_mask_view_ = std::nullopt; } - - bool has_edge_mask() const { return edge_mask_view_.has_value(); } - - std::optional> edge_mask_view() const - { - return edge_mask_view_; - } - private: edge_t const* offsets_{nullptr}; vertex_t const* indices_{nullptr}; // segment offsets based on vertex degree, relevant only if vertex IDs are renumbered std::optional> segment_offsets_{std::nullopt}; - - std::optional> edge_mask_view_{std::nullopt}; }; } // namespace cugraph diff --git a/cpp/include/cugraph/mtmg/detail/device_shared_wrapper.hpp b/cpp/include/cugraph/mtmg/detail/device_shared_wrapper.hpp index c4cacb401af..3e4b2513a8d 100644 --- a/cpp/include/cugraph/mtmg/detail/device_shared_wrapper.hpp +++ b/cpp/include/cugraph/mtmg/detail/device_shared_wrapper.hpp @@ -57,10 +57,10 @@ class device_shared_wrapper_t { { std::lock_guard lock(lock_); - auto pos = objects_.find(handle.get_local_rank()); + auto pos = objects_.find(handle.get_rank()); CUGRAPH_EXPECTS(pos == objects_.end(), "Cannot overwrite wrapped object"); - objects_.insert(std::make_pair(handle.get_local_rank(), std::move(obj))); + objects_.insert(std::make_pair(handle.get_rank(), std::move(obj))); } /** @@ -90,7 +90,7 @@ class device_shared_wrapper_t { { std::lock_guard lock(lock_); - auto pos = objects_.find(handle.get_local_rank()); + auto pos = objects_.find(handle.get_rank()); CUGRAPH_EXPECTS(pos != objects_.end(), "Uninitialized wrapped object"); return pos->second; @@ -106,7 +106,7 @@ class device_shared_wrapper_t { { std::lock_guard lock(lock_); - auto pos = objects_.find(handle.get_local_rank()); + auto pos = objects_.find(handle.get_rank()); CUGRAPH_EXPECTS(pos != objects_.end(), "Uninitialized wrapped object"); diff --git a/cpp/include/cugraph/mtmg/handle.hpp b/cpp/include/cugraph/mtmg/handle.hpp index 6223de1781d..0b02091a3cc 100644 --- a/cpp/include/cugraph/mtmg/handle.hpp +++ b/cpp/include/cugraph/mtmg/handle.hpp @@ -32,18 +32,19 @@ namespace mtmg { * */ class handle_t { + handle_t(handle_t const&) = delete; + handle_t operator=(handle_t const&) = delete; + public: /** * @brief Constructor * * @param raft_handle Raft handle for the resources * @param thread_rank Rank for this thread + * @param device_id Device id for the device this handle operates on */ - handle_t(raft::handle_t const& raft_handle, int thread_rank, size_t device_id) - : raft_handle_(raft_handle), - thread_rank_(thread_rank), - local_rank_(raft_handle.get_comms().get_rank()), // FIXME: update for multi-node - device_id_(device_id) + handle_t(raft::handle_t const& raft_handle, int thread_rank, rmm::cuda_device_id device_id) + : raft_handle_(raft_handle), thread_rank_(thread_rank), device_id_raii_(device_id) { } @@ -118,18 +119,10 @@ class handle_t { */ int get_rank() const { return raft_handle_.get_comms().get_rank(); } - /** - * @brief Get local gpu rank - * - * @return local gpu rank - */ - int get_local_rank() const { return local_rank_; } - private: raft::handle_t const& raft_handle_; int thread_rank_; - int local_rank_; - size_t device_id_; + rmm::cuda_set_device_raii device_id_raii_; }; } // namespace mtmg diff --git a/cpp/include/cugraph/mtmg/instance_manager.hpp b/cpp/include/cugraph/mtmg/instance_manager.hpp index f819a5a0abe..f60063c4101 100644 --- a/cpp/include/cugraph/mtmg/instance_manager.hpp +++ b/cpp/include/cugraph/mtmg/instance_manager.hpp @@ -47,15 +47,10 @@ class instance_manager_t { ~instance_manager_t() { - int current_device{}; - RAFT_CUDA_TRY(cudaGetDevice(¤t_device)); - for (size_t i = 0; i < nccl_comms_.size(); ++i) { - RAFT_CUDA_TRY(cudaSetDevice(device_ids_[i].value())); + rmm::cuda_set_device_raii local_set_device(device_ids_[i]); RAFT_NCCL_TRY(ncclCommDestroy(*nccl_comms_[i])); } - - RAFT_CUDA_TRY(cudaSetDevice(current_device)); } /** @@ -75,8 +70,7 @@ class instance_manager_t { int gpu_id = local_id % raft_handle_.size(); int thread_id = local_id / raft_handle_.size(); - RAFT_CUDA_TRY(cudaSetDevice(device_ids_[gpu_id].value())); - return handle_t(*raft_handle_[gpu_id], thread_id, static_cast(gpu_id)); + return handle_t(*raft_handle_[gpu_id], thread_id, device_ids_[gpu_id]); } /** diff --git a/cpp/include/cugraph/mtmg/resource_manager.hpp b/cpp/include/cugraph/mtmg/resource_manager.hpp index 127944cf7ba..bc312c9ae77 100644 --- a/cpp/include/cugraph/mtmg/resource_manager.hpp +++ b/cpp/include/cugraph/mtmg/resource_manager.hpp @@ -89,7 +89,7 @@ class resource_manager_t { local_rank_map_.insert(std::pair(global_rank, local_device_id)); - RAFT_CUDA_TRY(cudaSetDevice(local_device_id.value())); + rmm::cuda_set_device_raii local_set_device(local_device_id); // FIXME: There is a bug in the cuda_memory_resource that results in a Hang. // using the pool resource as a work-around. @@ -182,14 +182,12 @@ class resource_manager_t { --gpu_row_comm_size; } - int current_device{}; - RAFT_CUDA_TRY(cudaGetDevice(¤t_device)); RAFT_NCCL_TRY(ncclGroupStart()); for (size_t i = 0; i < local_ranks_to_include.size(); ++i) { int rank = local_ranks_to_include[i]; auto pos = local_rank_map_.find(rank); - RAFT_CUDA_TRY(cudaSetDevice(pos->second.value())); + rmm::cuda_set_device_raii local_set_device(pos->second); nccl_comms.push_back(std::make_unique()); handles.push_back( @@ -204,7 +202,6 @@ class resource_manager_t { handles[i].get(), *nccl_comms[i], ranks_to_include.size(), rank); } RAFT_NCCL_TRY(ncclGroupEnd()); - RAFT_CUDA_TRY(cudaSetDevice(current_device)); std::vector running_threads; @@ -217,9 +214,7 @@ class resource_manager_t { &device_ids, &nccl_comms, &handles]() { - int rank = local_ranks_to_include[idx]; - RAFT_CUDA_TRY(cudaSetDevice(device_ids[idx].value())); - + rmm::cuda_set_device_raii local_set_device(device_ids[idx]); cugraph::partition_manager::init_subcomm(*handles[idx], gpu_row_comm_size); }); } diff --git a/cpp/include/cugraph/utilities/misc_utils.cuh b/cpp/include/cugraph/utilities/misc_utils.cuh index a62e8ce85ec..28e2853727f 100644 --- a/cpp/include/cugraph/utilities/misc_utils.cuh +++ b/cpp/include/cugraph/utilities/misc_utils.cuh @@ -19,6 +19,7 @@ #include #include +#include #include #include #include diff --git a/cpp/src/c_api/graph_mg.cpp b/cpp/src/c_api/graph_mg.cpp index 5413949e3a3..326022a3fa9 100644 --- a/cpp/src/c_api/graph_mg.cpp +++ b/cpp/src/c_api/graph_mg.cpp @@ -523,9 +523,9 @@ extern "C" cugraph_error_code_t cugraph_mg_graph_create( NULL, &src, &dst, - &weights, - &edge_ids, - &edge_type_ids, + (weights == nullptr) ? nullptr : &weights, + (edge_ids == nullptr) ? nullptr : &edge_ids, + (edge_type_ids == nullptr) ? nullptr : &edge_type_ids, store_transposed, 1, FALSE, diff --git a/cpp/src/prims/count_if_e.cuh b/cpp/src/prims/count_if_e.cuh index f6e4bc9bead..9cff4f5eceb 100644 --- a/cpp/src/prims/count_if_e.cuh +++ b/cpp/src/prims/count_if_e.cuh @@ -74,8 +74,6 @@ typename GraphViewType::edge_type count_if_e(raft::handle_t const& handle, using vertex_t = typename GraphViewType::vertex_type; using edge_t = typename GraphViewType::edge_type; - CUGRAPH_EXPECTS(!graph_view.has_edge_mask(), "unimplemented."); - if (do_expensive_check) { // currently, nothing to do } diff --git a/cpp/src/prims/fill_edge_property.cuh b/cpp/src/prims/fill_edge_property.cuh index d446944b65b..e6875576044 100644 --- a/cpp/src/prims/fill_edge_property.cuh +++ b/cpp/src/prims/fill_edge_property.cuh @@ -15,6 +15,7 @@ */ #pragma once +#include #include #include #include @@ -23,6 +24,7 @@ #include #include +#include #include @@ -38,21 +40,78 @@ void fill_edge_property(raft::handle_t const& handle, { static_assert(std::is_same_v); + using edge_t = typename GraphViewType::edge_type; + + auto edge_mask_view = graph_view.edge_mask_view(); + auto value_firsts = edge_property_output.value_firsts(); auto edge_counts = edge_property_output.edge_counts(); for (size_t i = 0; i < graph_view.number_of_local_edge_partitions(); ++i) { + auto edge_partition_e_mask = + edge_mask_view + ? thrust::make_optional< + detail::edge_partition_edge_property_device_view_t>( + *edge_mask_view, i) + : thrust::nullopt; + if constexpr (cugraph::has_packed_bool_element< std::remove_reference_t, T>()) { static_assert(std::is_arithmetic_v, "unimplemented for thrust::tuple types."); auto packed_input = input ? packed_bool_full_mask() : packed_bool_empty_mask(); - thrust::fill_n(handle.get_thrust_policy(), - value_firsts[i], - packed_bool_size(static_cast(edge_counts[i])), - packed_input); + auto rem = edge_counts[i] % packed_bools_per_word(); + if (edge_partition_e_mask) { + auto input_first = + thrust::make_zip_iterator(value_firsts[i], (*edge_partition_e_mask).value_first()); + thrust::transform(handle.get_thrust_policy(), + input_first, + input_first + packed_bool_size(static_cast(edge_counts[i] - rem)), + value_firsts[i], + [packed_input] __device__(thrust::tuple pair) { + auto old_value = thrust::get<0>(pair); + auto mask = thrust::get<1>(pair); + return (old_value & ~mask) | (packed_input & mask); + }); + if (rem > 0) { + thrust::transform( + handle.get_thrust_policy(), + input_first + packed_bool_size(static_cast(edge_counts[i] - rem)), + input_first + packed_bool_size(static_cast(edge_counts[i])), + value_firsts[i] + packed_bool_size(static_cast(edge_counts[i] - rem)), + [packed_input, rem] __device__(thrust::tuple pair) { + auto old_value = thrust::get<0>(pair); + auto mask = thrust::get<1>(pair); + return ((old_value & ~mask) | (packed_input & mask)) & packed_bool_partial_mask(rem); + }); + } + } else { + thrust::fill_n(handle.get_thrust_policy(), + value_firsts[i], + packed_bool_size(static_cast(edge_counts[i] - rem)), + packed_input); + if (rem > 0) { + thrust::fill_n( + handle.get_thrust_policy(), + value_firsts[i] + packed_bool_size(static_cast(edge_counts[i] - rem)), + 1, + packed_input & packed_bool_partial_mask(rem)); + } + } } else { - thrust::fill_n( - handle.get_thrust_policy(), value_firsts[i], static_cast(edge_counts[i]), input); + if (edge_partition_e_mask) { + thrust::transform_if(handle.get_thrust_policy(), + thrust::make_constant_iterator(input), + thrust::make_constant_iterator(input) + edge_counts[i], + thrust::make_counting_iterator(edge_t{0}), + value_firsts[i], + thrust::identity{}, + [edge_partition_e_mask = *edge_partition_e_mask] __device__(edge_t i) { + return edge_partition_e_mask.get(i); + }); + } else { + thrust::fill_n( + handle.get_thrust_policy(), value_firsts[i], static_cast(edge_counts[i]), input); + } } } } @@ -79,8 +138,6 @@ void fill_edge_property(raft::handle_t const& handle, edge_property_t& edge_property_output, bool do_expensive_check = false) { - CUGRAPH_EXPECTS(!graph_view.has_edge_mask(), "unimplemented."); - if (do_expensive_check) { // currently, nothing to do } diff --git a/cpp/src/prims/transform_e.cuh b/cpp/src/prims/transform_e.cuh index edacdc8a970..c6623621d24 100644 --- a/cpp/src/prims/transform_e.cuh +++ b/cpp/src/prims/transform_e.cuh @@ -16,10 +16,12 @@ #pragma once #include +#include #include #include #include #include +#include #include #include @@ -44,6 +46,7 @@ template __global__ void transform_e_packed_bool( @@ -53,6 +56,7 @@ __global__ void transform_e_packed_bool( EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, EdgePartitionEdgeValueInputWrapper edge_partition_e_value_input, + thrust::optional edge_partition_e_mask, EdgePartitionEdgeValueOutputWrapper edge_partition_e_value_output, EdgeOp e_op) { @@ -68,11 +72,14 @@ __global__ void transform_e_packed_bool( auto num_edges = edge_partition.number_of_edges(); while (idx < static_cast(packed_bool_size(num_edges))) { + auto edge_mask = packed_bool_full_mask(); + if (edge_partition_e_mask) { edge_mask = *((*edge_partition_e_mask).value_first() + idx); } + auto local_edge_idx = idx * static_cast(packed_bools_per_word()) + static_cast(lane_id); - uint32_t mask{0}; int predicate{0}; - if (local_edge_idx < num_edges) { + + if ((local_edge_idx < num_edges) && (edge_mask & packed_bool_mask(lane_id))) { auto major_idx = edge_partition.major_idx_from_local_edge_idx_nocheck(local_edge_idx); auto major = edge_partition.major_from_major_idx_nocheck(major_idx); auto major_offset = edge_partition.major_offset_from_major_nocheck(major); @@ -91,8 +98,15 @@ __global__ void transform_e_packed_bool( ? int{1} : int{0}; } - mask = __ballot_sync(uint32_t{0xffffffff}, predicate); - if (lane_id == 0) { *(edge_partition_e_value_output.value_first() + idx) = mask; } + uint32_t new_val = __ballot_sync(uint32_t{0xffffffff}, predicate); + if (lane_id == 0) { + if (edge_mask == packed_bool_full_mask()) { + *(edge_partition_e_value_output.value_first() + idx) = new_val; + } else { + auto old_val = *(edge_partition_e_value_output.value_first() + idx); + *(edge_partition_e_value_output.value_first() + idx) = (old_val & ~edge_mask) | new_val; + } + } idx += static_cast(gridDim.x * (blockDim.x / raft::warp_size())); } @@ -178,12 +192,18 @@ void transform_e(raft::handle_t const& handle, typename EdgeValueOutputWrapper::value_iterator, typename EdgeValueOutputWrapper::value_type>; - CUGRAPH_EXPECTS(!graph_view.has_edge_mask(), "unimplemented."); + auto edge_mask_view = graph_view.edge_mask_view(); for (size_t i = 0; i < graph_view.number_of_local_edge_partitions(); ++i) { auto edge_partition = edge_partition_device_view_t( graph_view.local_edge_partition_view(i)); + auto edge_partition_e_mask = + edge_mask_view + ? thrust::make_optional< + detail::edge_partition_edge_property_device_view_t>( + *edge_mask_view, i) + : thrust::nullopt; edge_partition_src_input_device_view_t edge_partition_src_value_input{}; edge_partition_dst_input_device_view_t edge_partition_dst_value_input{}; @@ -214,35 +234,40 @@ void transform_e(raft::handle_t const& handle, edge_partition_src_value_input, edge_partition_dst_value_input, edge_partition_e_value_input, + edge_partition_e_mask, edge_partition_e_value_output, e_op); } } else { - thrust::transform( + thrust::for_each( handle.get_thrust_policy(), thrust::make_counting_iterator(edge_t{0}), thrust::make_counting_iterator(num_edges), - edge_partition_e_value_output.value_first(), [e_op, edge_partition, edge_partition_src_value_input, edge_partition_dst_value_input, - edge_partition_e_value_input] __device__(edge_t i) { - auto major_idx = edge_partition.major_idx_from_local_edge_idx_nocheck(i); - auto major = edge_partition.major_from_major_idx_nocheck(major_idx); - auto major_offset = edge_partition.major_offset_from_major_nocheck(major); - auto minor = *(edge_partition.indices() + i); - auto minor_offset = edge_partition.minor_offset_from_minor_nocheck(minor); - - auto src = GraphViewType::is_storage_transposed ? minor : major; - auto dst = GraphViewType::is_storage_transposed ? major : minor; - auto src_offset = GraphViewType::is_storage_transposed ? minor_offset : major_offset; - auto dst_offset = GraphViewType::is_storage_transposed ? major_offset : minor_offset; - return e_op(src, - dst, - edge_partition_src_value_input.get(src_offset), - edge_partition_dst_value_input.get(dst_offset), - edge_partition_e_value_input.get(i)); + edge_partition_e_value_input, + edge_partition_e_mask, + edge_partition_e_value_output] __device__(edge_t i) { + if (!edge_partition_e_mask || (*edge_partition_e_mask).get(i)) { + auto major_idx = edge_partition.major_idx_from_local_edge_idx_nocheck(i); + auto major = edge_partition.major_from_major_idx_nocheck(major_idx); + auto major_offset = edge_partition.major_offset_from_major_nocheck(major); + auto minor = *(edge_partition.indices() + i); + auto minor_offset = edge_partition.minor_offset_from_minor_nocheck(minor); + + auto src = GraphViewType::is_storage_transposed ? minor : major; + auto dst = GraphViewType::is_storage_transposed ? major : minor; + auto src_offset = GraphViewType::is_storage_transposed ? minor_offset : major_offset; + auto dst_offset = GraphViewType::is_storage_transposed ? major_offset : minor_offset; + auto e_op_result = e_op(src, + dst, + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), + edge_partition_e_value_input.get(i)); + edge_partition_e_value_output.set(i, e_op_result); + } }); } } @@ -336,14 +361,12 @@ void transform_e(raft::handle_t const& handle, typename EdgeValueOutputWrapper::value_iterator, typename EdgeValueOutputWrapper::value_type>; - CUGRAPH_EXPECTS(!graph_view.has_edge_mask(), "unimplemented."); - auto major_first = GraphViewType::is_storage_transposed ? edge_list.dst_begin() : edge_list.src_begin(); auto minor_first = GraphViewType::is_storage_transposed ? edge_list.src_begin() : edge_list.dst_begin(); - auto edge_first = thrust::make_zip_iterator(thrust::make_tuple(major_first, minor_first)); + auto edge_first = thrust::make_zip_iterator(major_first, minor_first); if (do_expensive_check) { CUGRAPH_EXPECTS( @@ -382,10 +405,18 @@ void transform_e(raft::handle_t const& handle, edge_partition_offsets.back() = edge_list.size(); } + auto edge_mask_view = graph_view.edge_mask_view(); + for (size_t i = 0; i < graph_view.number_of_local_edge_partitions(); ++i) { auto edge_partition = edge_partition_device_view_t( graph_view.local_edge_partition_view(i)); + auto edge_partition_e_mask = + edge_mask_view + ? thrust::make_optional< + detail::edge_partition_edge_property_device_view_t>( + *edge_mask_view, i) + : thrust::nullopt; if (do_expensive_check) { CUGRAPH_EXPECTS( @@ -393,7 +424,8 @@ void transform_e(raft::handle_t const& handle, handle.get_thrust_policy(), edge_first + edge_partition_offsets[i], edge_first + edge_partition_offsets[i + 1], - [edge_partition] __device__(thrust::tuple edge) { + [edge_partition, + edge_partition_e_mask] __device__(thrust::tuple edge) { auto major = thrust::get<0>(edge); auto minor = thrust::get<1>(edge); vertex_t major_idx{}; @@ -416,8 +448,19 @@ void transform_e(raft::handle_t const& handle, edge_t edge_offset{}; edge_t local_degree{}; thrust::tie(indices, edge_offset, local_degree) = edge_partition.local_edges(major_idx); - auto it = thrust::lower_bound(thrust::seq, indices, indices + local_degree, minor); - return *it != minor; + auto lower_it = + thrust::lower_bound(thrust::seq, indices, indices + local_degree, minor); + if (*lower_it != minor) { return true; } + if (edge_partition_e_mask) { + auto upper_it = + thrust::upper_bound(thrust::seq, lower_it, indices + local_degree, minor); + if (detail::count_set_bits((*edge_partition_e_mask).value_first(), + edge_offset + thrust::distance(indices, lower_it), + thrust::distance(lower_it, upper_it)) == 0) { + return true; + } + } + return false; }) == 0, "Invalid input arguments: edge_list contains edges that do not exist in the input graph."); } @@ -446,6 +489,7 @@ void transform_e(raft::handle_t const& handle, edge_partition_src_value_input, edge_partition_dst_value_input, edge_partition_e_value_input, + edge_partition_e_mask, edge_partition_e_value_output] __device__(thrust::tuple edge) { auto major = thrust::get<0>(edge); auto minor = thrust::get<1>(edge); @@ -469,7 +513,7 @@ void transform_e(raft::handle_t const& handle, edge_t local_degree{}; thrust::tie(indices, edge_offset, local_degree) = edge_partition.local_edges(major_idx); auto lower_it = thrust::lower_bound(thrust::seq, indices, indices + local_degree, minor); - auto upper_it = thrust::upper_bound(thrust::seq, indices, indices + local_degree, minor); + auto upper_it = thrust::upper_bound(thrust::seq, lower_it, indices + local_degree, minor); auto src = GraphViewType::is_storage_transposed ? minor : major; auto dst = GraphViewType::is_storage_transposed ? major : minor; @@ -478,14 +522,17 @@ void transform_e(raft::handle_t const& handle, for (auto it = lower_it; it != upper_it; ++it) { assert(*it == minor); - auto e_op_result = - e_op(src, - dst, - edge_partition_src_value_input.get(src_offset), - edge_partition_dst_value_input.get(dst_offset), - edge_partition_e_value_input.get(edge_offset + thrust::distance(indices, it))); - edge_partition_e_value_output.set(edge_offset + thrust::distance(indices, it), - e_op_result); + if (!edge_partition_e_mask || + ((*edge_partition_e_mask).get(edge_offset + thrust::distance(indices, it)))) { + auto e_op_result = + e_op(src, + dst, + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), + edge_partition_e_value_input.get(edge_offset + thrust::distance(indices, it))); + edge_partition_e_value_output.set(edge_offset + thrust::distance(indices, it), + e_op_result); + } } }); } diff --git a/cpp/src/prims/transform_reduce_e.cuh b/cpp/src/prims/transform_reduce_e.cuh index 9c23f3fca18..483ab64dcd9 100644 --- a/cpp/src/prims/transform_reduce_e.cuh +++ b/cpp/src/prims/transform_reduce_e.cuh @@ -56,6 +56,7 @@ template __global__ void transform_reduce_e_hypersparse( @@ -65,6 +66,7 @@ __global__ void transform_reduce_e_hypersparse( EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, EdgePartitionEdgeValueInputWrapper edge_partition_e_value_input, + thrust::optional edge_partition_e_mask, ResultIterator result_iter /* size 1 */, EdgeOp e_op) { @@ -101,24 +103,31 @@ __global__ void transform_reduce_e_hypersparse( &edge_partition_src_value_input, &edge_partition_dst_value_input, &edge_partition_e_value_input, + &edge_partition_e_mask, &e_op, major, indices, edge_offset] __device__(auto i) { - auto major_offset = edge_partition.major_offset_from_major_nocheck(major); - auto minor = indices[i]; - auto minor_offset = edge_partition.minor_offset_from_minor_nocheck(minor); - auto src = GraphViewType::is_storage_transposed ? minor : major; - auto dst = GraphViewType::is_storage_transposed ? major : minor; - auto src_offset = - GraphViewType::is_storage_transposed ? minor_offset : static_cast(major_offset); - auto dst_offset = - GraphViewType::is_storage_transposed ? static_cast(major_offset) : minor_offset; - return e_op(src, - dst, - edge_partition_src_value_input.get(src_offset), - edge_partition_dst_value_input.get(dst_offset), - edge_partition_e_value_input.get(edge_offset + i)); + if (!edge_partition_e_mask || (*edge_partition_e_mask).get(edge_offset + i)) { + auto major_offset = edge_partition.major_offset_from_major_nocheck(major); + auto minor = indices[i]; + auto minor_offset = edge_partition.minor_offset_from_minor_nocheck(minor); + auto src = GraphViewType::is_storage_transposed ? minor : major; + auto dst = GraphViewType::is_storage_transposed ? major : minor; + auto src_offset = GraphViewType::is_storage_transposed + ? minor_offset + : static_cast(major_offset); + auto dst_offset = GraphViewType::is_storage_transposed + ? static_cast(major_offset) + : minor_offset; + return e_op(src, + dst, + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), + edge_partition_e_value_input.get(edge_offset + i)); + } else { + return e_op_result_t{}; + } }, e_op_result_t{}, edge_property_add); @@ -135,6 +144,7 @@ template __global__ void transform_reduce_e_low_degree( @@ -146,6 +156,7 @@ __global__ void transform_reduce_e_low_degree( EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, EdgePartitionEdgeValueInputWrapper edge_partition_e_value_input, + thrust::optional edge_partition_e_mask, ResultIterator result_iter /* size 1 */, EdgeOp e_op) { @@ -177,27 +188,34 @@ __global__ void transform_reduce_e_low_degree( &edge_partition_src_value_input, &edge_partition_dst_value_input, &edge_partition_e_value_input, + &edge_partition_e_mask, &e_op, major_offset, indices, edge_offset] __device__(auto i) { - auto minor = indices[i]; - auto minor_offset = edge_partition.minor_offset_from_minor_nocheck(minor); - auto src = GraphViewType::is_storage_transposed - ? minor - : edge_partition.major_from_major_offset_nocheck(major_offset); - auto dst = GraphViewType::is_storage_transposed - ? edge_partition.major_from_major_offset_nocheck(major_offset) - : minor; - auto src_offset = - GraphViewType::is_storage_transposed ? minor_offset : static_cast(major_offset); - auto dst_offset = - GraphViewType::is_storage_transposed ? static_cast(major_offset) : minor_offset; - return e_op(src, - dst, - edge_partition_src_value_input.get(src_offset), - edge_partition_dst_value_input.get(dst_offset), - edge_partition_e_value_input.get(edge_offset + i)); + if (!edge_partition_e_mask || (*edge_partition_e_mask).get(edge_offset + i)) { + auto minor = indices[i]; + auto minor_offset = edge_partition.minor_offset_from_minor_nocheck(minor); + auto src = GraphViewType::is_storage_transposed + ? minor + : edge_partition.major_from_major_offset_nocheck(major_offset); + auto dst = GraphViewType::is_storage_transposed + ? edge_partition.major_from_major_offset_nocheck(major_offset) + : minor; + auto src_offset = GraphViewType::is_storage_transposed + ? minor_offset + : static_cast(major_offset); + auto dst_offset = GraphViewType::is_storage_transposed + ? static_cast(major_offset) + : minor_offset; + return e_op(src, + dst, + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), + edge_partition_e_value_input.get(edge_offset + i)); + } else { + return e_op_result_t{}; + } }, e_op_result_t{}, edge_property_add); @@ -214,6 +232,7 @@ template __global__ void transform_reduce_e_mid_degree( @@ -225,6 +244,7 @@ __global__ void transform_reduce_e_mid_degree( EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, EdgePartitionEdgeValueInputWrapper edge_partition_e_value_input, + thrust::optional edge_partition_e_mask, ResultIterator result_iter /* size 1 */, EdgeOp e_op) { @@ -250,24 +270,26 @@ __global__ void transform_reduce_e_mid_degree( edge_t local_degree{}; thrust::tie(indices, edge_offset, local_degree) = edge_partition.local_edges(major_offset); for (edge_t i = lane_id; i < local_degree; i += raft::warp_size()) { - auto minor = indices[i]; - auto minor_offset = edge_partition.minor_offset_from_minor_nocheck(minor); - auto src = GraphViewType::is_storage_transposed - ? minor - : edge_partition.major_from_major_offset_nocheck(major_offset); - auto dst = GraphViewType::is_storage_transposed - ? edge_partition.major_from_major_offset_nocheck(major_offset) - : minor; - auto src_offset = - GraphViewType::is_storage_transposed ? minor_offset : static_cast(major_offset); - auto dst_offset = - GraphViewType::is_storage_transposed ? static_cast(major_offset) : minor_offset; - auto e_op_result = e_op(src, - dst, - edge_partition_src_value_input.get(src_offset), - edge_partition_dst_value_input.get(dst_offset), - edge_partition_e_value_input.get(edge_offset + i)); - e_op_result_sum = edge_property_add(e_op_result_sum, e_op_result); + if (!edge_partition_e_mask || (*edge_partition_e_mask).get(edge_offset + i)) { + auto minor = indices[i]; + auto minor_offset = edge_partition.minor_offset_from_minor_nocheck(minor); + auto src = GraphViewType::is_storage_transposed + ? minor + : edge_partition.major_from_major_offset_nocheck(major_offset); + auto dst = GraphViewType::is_storage_transposed + ? edge_partition.major_from_major_offset_nocheck(major_offset) + : minor; + auto src_offset = + GraphViewType::is_storage_transposed ? minor_offset : static_cast(major_offset); + auto dst_offset = + GraphViewType::is_storage_transposed ? static_cast(major_offset) : minor_offset; + auto e_op_result = e_op(src, + dst, + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), + edge_partition_e_value_input.get(edge_offset + i)); + e_op_result_sum = edge_property_add(e_op_result_sum, e_op_result); + } } idx += gridDim.x * (blockDim.x / raft::warp_size()); } @@ -280,6 +302,7 @@ template __global__ void transform_reduce_e_high_degree( @@ -291,6 +314,7 @@ __global__ void transform_reduce_e_high_degree( EdgePartitionSrcValueInputWrapper edge_partition_src_value_input, EdgePartitionDstValueInputWrapper edge_partition_dst_value_input, EdgePartitionEdgeValueInputWrapper edge_partition_e_value_input, + thrust::optional edge_partition_e_mask, ResultIterator result_iter /* size 1 */, EdgeOp e_op) { @@ -313,24 +337,26 @@ __global__ void transform_reduce_e_high_degree( edge_t local_degree{}; thrust::tie(indices, edge_offset, local_degree) = edge_partition.local_edges(major_offset); for (edge_t i = threadIdx.x; i < local_degree; i += blockDim.x) { - auto minor = indices[i]; - auto minor_offset = edge_partition.minor_offset_from_minor_nocheck(minor); - auto src = GraphViewType::is_storage_transposed - ? minor - : edge_partition.major_from_major_offset_nocheck(major_offset); - auto dst = GraphViewType::is_storage_transposed - ? edge_partition.major_from_major_offset_nocheck(major_offset) - : minor; - auto src_offset = - GraphViewType::is_storage_transposed ? minor_offset : static_cast(major_offset); - auto dst_offset = - GraphViewType::is_storage_transposed ? static_cast(major_offset) : minor_offset; - auto e_op_result = e_op(src, - dst, - edge_partition_src_value_input.get(src_offset), - edge_partition_dst_value_input.get(dst_offset), - edge_partition_e_value_input.get(edge_offset + i)); - e_op_result_sum = edge_property_add(e_op_result_sum, e_op_result); + if (!edge_partition_e_mask || (*edge_partition_e_mask).get(edge_offset + i)) { + auto minor = indices[i]; + auto minor_offset = edge_partition.minor_offset_from_minor_nocheck(minor); + auto src = GraphViewType::is_storage_transposed + ? minor + : edge_partition.major_from_major_offset_nocheck(major_offset); + auto dst = GraphViewType::is_storage_transposed + ? edge_partition.major_from_major_offset_nocheck(major_offset) + : minor; + auto src_offset = + GraphViewType::is_storage_transposed ? minor_offset : static_cast(major_offset); + auto dst_offset = + GraphViewType::is_storage_transposed ? static_cast(major_offset) : minor_offset; + auto e_op_result = e_op(src, + dst, + edge_partition_src_value_input.get(src_offset), + edge_partition_dst_value_input.get(dst_offset), + edge_partition_e_value_input.get(edge_offset + i)); + e_op_result_sum = edge_property_add(e_op_result_sum, e_op_result); + } } idx += gridDim.x; } @@ -417,8 +443,6 @@ T transform_reduce_e(raft::handle_t const& handle, typename EdgeValueInputWrapper::value_iterator, typename EdgeValueInputWrapper::value_type>>; - CUGRAPH_EXPECTS(!graph_view.has_edge_mask(), "unimplemented."); - if (do_expensive_check) { // currently, nothing to do } @@ -431,10 +455,18 @@ T transform_reduce_e(raft::handle_t const& handle, get_dataframe_buffer_begin(result_buffer) + 1, T{}); + auto edge_mask_view = graph_view.edge_mask_view(); + for (size_t i = 0; i < graph_view.number_of_local_edge_partitions(); ++i) { auto edge_partition = edge_partition_device_view_t( graph_view.local_edge_partition_view(i)); + auto edge_partition_e_mask = + edge_mask_view + ? thrust::make_optional< + detail::edge_partition_edge_property_device_view_t>( + *edge_mask_view, i) + : thrust::nullopt; edge_partition_src_input_device_view_t edge_partition_src_value_input{}; edge_partition_dst_input_device_view_t edge_partition_dst_value_input{}; @@ -467,6 +499,7 @@ T transform_reduce_e(raft::handle_t const& handle, edge_partition_src_value_input, edge_partition_dst_value_input, edge_partition_e_value_input, + edge_partition_e_mask, get_dataframe_buffer_begin(result_buffer), e_op); } @@ -482,6 +515,7 @@ T transform_reduce_e(raft::handle_t const& handle, edge_partition_src_value_input, edge_partition_dst_value_input, edge_partition_e_value_input, + edge_partition_e_mask, get_dataframe_buffer_begin(result_buffer), e_op); } @@ -497,6 +531,7 @@ T transform_reduce_e(raft::handle_t const& handle, edge_partition_src_value_input, edge_partition_dst_value_input, edge_partition_e_value_input, + edge_partition_e_mask, get_dataframe_buffer_begin(result_buffer), e_op); } @@ -510,6 +545,7 @@ T transform_reduce_e(raft::handle_t const& handle, edge_partition_src_value_input, edge_partition_dst_value_input, edge_partition_e_value_input, + edge_partition_e_mask, get_dataframe_buffer_begin(result_buffer), e_op); } @@ -527,6 +563,7 @@ T transform_reduce_e(raft::handle_t const& handle, edge_partition_src_value_input, edge_partition_dst_value_input, edge_partition_e_value_input, + edge_partition_e_mask, get_dataframe_buffer_begin(result_buffer), e_op); } @@ -601,8 +638,6 @@ auto transform_reduce_e(raft::handle_t const& handle, edge_op_result_type::type; static_assert(!std::is_same_v); - CUGRAPH_EXPECTS(!graph_view.has_edge_mask(), "unimplemented."); - if (do_expensive_check) { // currently, nothing to do } diff --git a/cpp/src/prims/update_edge_src_dst_property.cuh b/cpp/src/prims/update_edge_src_dst_property.cuh index 2d72a075ca5..b8621e122c6 100644 --- a/cpp/src/prims/update_edge_src_dst_property.cuh +++ b/cpp/src/prims/update_edge_src_dst_property.cuh @@ -866,8 +866,6 @@ void update_edge_src_property( edge_src_property_output, bool do_expensive_check = false) { - CUGRAPH_EXPECTS(!graph_view.has_edge_mask(), "unimplemented."); - if (do_expensive_check) { // currently, nothing to do } @@ -917,8 +915,6 @@ void update_edge_src_property( edge_src_property_output, bool do_expensive_check = false) { - CUGRAPH_EXPECTS(!graph_view.has_edge_mask(), "unimplemented."); - if (do_expensive_check) { auto num_invalids = thrust::count_if( handle.get_thrust_policy(), @@ -985,8 +981,6 @@ void update_edge_dst_property( edge_dst_property_output, bool do_expensive_check = false) { - CUGRAPH_EXPECTS(!graph_view.has_edge_mask(), "unimplemented."); - if (do_expensive_check) { // currently, nothing to do } @@ -1036,8 +1030,6 @@ void update_edge_dst_property( edge_dst_property_output, bool do_expensive_check = false) { - CUGRAPH_EXPECTS(!graph_view.has_edge_mask(), "unimplemented."); - if (do_expensive_check) { auto num_invalids = thrust::count_if( handle.get_thrust_policy(), diff --git a/cpp/src/structure/detail/structure_utils.cuh b/cpp/src/structure/detail/structure_utils.cuh index c49b62e4543..7630d5855a0 100644 --- a/cpp/src/structure/detail/structure_utils.cuh +++ b/cpp/src/structure/detail/structure_utils.cuh @@ -524,6 +524,7 @@ std::tuple> mark_entries(raft::handle_t co return word; }); + // FIXME: use detail::count_set_bits size_t bit_count = thrust::transform_reduce( handle.get_thrust_policy(), marked_entries.begin(), diff --git a/cpp/src/structure/graph_view_impl.cuh b/cpp/src/structure/graph_view_impl.cuh index 64a8a3212b3..37a553dcdbd 100644 --- a/cpp/src/structure/graph_view_impl.cuh +++ b/cpp/src/structure/graph_view_impl.cuh @@ -548,7 +548,7 @@ graph_view_tpartition_, this->edge_partition_segment_offsets_); } else { - CUGRAPH_EXPECTS(!has_edge_mask(), "unimplemented."); + CUGRAPH_EXPECTS(!(this->has_edge_mask()), "unimplemented."); return compute_minor_degrees(handle, *this); } } @@ -566,7 +566,7 @@ graph_view_tlocal_vertex_partition_range_size()); } else { - CUGRAPH_EXPECTS(!has_edge_mask(), "unimplemented."); + CUGRAPH_EXPECTS(!(this->has_edge_mask()), "unimplemented."); return compute_minor_degrees(handle, *this); } } @@ -577,7 +577,7 @@ graph_view_thas_edge_mask()), "unimplemented."); return compute_minor_degrees(handle, *this); } else { return compute_major_degrees(handle, @@ -598,7 +598,7 @@ graph_view_thas_edge_mask()), "unimplemented."); return compute_minor_degrees(handle, *this); } else { return compute_major_degrees( @@ -614,7 +614,7 @@ template >:: compute_max_in_degree(raft::handle_t const& handle) const { - CUGRAPH_EXPECTS(!has_edge_mask(), "unimplemented."); + CUGRAPH_EXPECTS(!(this->has_edge_mask()), "unimplemented."); auto in_degrees = compute_in_degrees(handle); auto it = thrust::max_element(handle.get_thrust_policy(), in_degrees.begin(), in_degrees.end()); @@ -632,7 +632,7 @@ template >:: compute_max_in_degree(raft::handle_t const& handle) const { - CUGRAPH_EXPECTS(!has_edge_mask(), "unimplemented."); + CUGRAPH_EXPECTS(!(this->has_edge_mask()), "unimplemented."); auto in_degrees = compute_in_degrees(handle); auto it = thrust::max_element(handle.get_thrust_policy(), in_degrees.begin(), in_degrees.end()); @@ -646,7 +646,7 @@ template >:: compute_max_out_degree(raft::handle_t const& handle) const { - CUGRAPH_EXPECTS(!has_edge_mask(), "unimplemented."); + CUGRAPH_EXPECTS(!(this->has_edge_mask()), "unimplemented."); auto out_degrees = compute_out_degrees(handle); auto it = thrust::max_element(handle.get_thrust_policy(), out_degrees.begin(), out_degrees.end()); @@ -664,7 +664,7 @@ template >:: compute_max_out_degree(raft::handle_t const& handle) const { - CUGRAPH_EXPECTS(!has_edge_mask(), "unimplemented."); + CUGRAPH_EXPECTS(!(this->has_edge_mask()), "unimplemented."); auto out_degrees = compute_out_degrees(handle); auto it = thrust::max_element(handle.get_thrust_policy(), out_degrees.begin(), out_degrees.end()); @@ -678,7 +678,7 @@ template >:: count_self_loops(raft::handle_t const& handle) const { - CUGRAPH_EXPECTS(!has_edge_mask(), "unimplemented."); + CUGRAPH_EXPECTS(!(this->has_edge_mask()), "unimplemented."); return count_if_e( handle, @@ -693,7 +693,7 @@ template >:: count_self_loops(raft::handle_t const& handle) const { - CUGRAPH_EXPECTS(!has_edge_mask(), "unimplemented."); + CUGRAPH_EXPECTS(!(this->has_edge_mask()), "unimplemented."); return count_if_e( handle, @@ -708,7 +708,7 @@ template >:: count_multi_edges(raft::handle_t const& handle) const { - CUGRAPH_EXPECTS(!has_edge_mask(), "unimplemented."); + CUGRAPH_EXPECTS(!(this->has_edge_mask()), "unimplemented."); if (!this->is_multigraph()) { return edge_t{0}; } @@ -728,7 +728,7 @@ template >:: count_multi_edges(raft::handle_t const& handle) const { - CUGRAPH_EXPECTS(!has_edge_mask(), "unimplemented."); + CUGRAPH_EXPECTS(!(this->has_edge_mask()), "unimplemented."); if (!this->is_multigraph()) { return edge_t{0}; } diff --git a/cpp/tests/c_api/mg_test_utils.cpp b/cpp/tests/c_api/mg_test_utils.cpp index 15df613ae05..6eec436e77d 100644 --- a/cpp/tests/c_api/mg_test_utils.cpp +++ b/cpp/tests/c_api/mg_test_utils.cpp @@ -158,30 +158,22 @@ extern "C" int create_mg_test_graph(const cugraph_resource_handle_t* handle, rank = cugraph_resource_handle_get_rank(handle); - if (rank == 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)); + size_t original_num_edges = num_edges; - 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."); + if (rank != 0) num_edges = 0; - 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."); - } else { - ret_code = cugraph_type_erased_device_array_create(handle, 0, 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, &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, 0, vertex_tid, &dst, ret_error); - TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "dst create failed."); + 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."); - ret_code = cugraph_type_erased_device_array_create(handle, 0, weight_tid, &wgt, ret_error); - TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt create failed."); - } + 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."); src_view = cugraph_type_erased_device_array_view(src); dst_view = cugraph_type_erased_device_array_view(dst); @@ -207,7 +199,7 @@ extern "C" int create_mg_test_graph(const cugraph_resource_handle_t* handle, NULL, NULL, store_transposed, - num_edges, + original_num_edges, // UNUSED FALSE, p_graph, ret_error); @@ -260,30 +252,22 @@ extern "C" int create_mg_test_graph_double(const cugraph_resource_handle_t* hand rank = cugraph_resource_handle_get_rank(handle); - if (rank == 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)); + size_t original_num_edges = num_edges; - 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."); + if (rank != 0) num_edges = 0; - 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."); - } else { - ret_code = cugraph_type_erased_device_array_create(handle, 0, 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, &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, 0, vertex_tid, &dst, ret_error); - TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "dst create failed."); + 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."); - ret_code = cugraph_type_erased_device_array_create(handle, 0, weight_tid, &wgt, ret_error); - TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt create failed."); - } + 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."); src_view = cugraph_type_erased_device_array_view(src); dst_view = cugraph_type_erased_device_array_view(dst); @@ -309,7 +293,7 @@ extern "C" int create_mg_test_graph_double(const cugraph_resource_handle_t* hand NULL, NULL, store_transposed, - num_edges, + original_num_edges, // UNUSED FALSE, p_graph, ret_error); @@ -357,30 +341,22 @@ extern "C" int create_mg_test_graph_with_edge_ids(const cugraph_resource_handle_ rank = cugraph_resource_handle_get_rank(handle); - if (rank == 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)); + size_t original_num_edges = num_edges; - 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."); + if (rank != 0) num_edges = 0; - ret_code = - cugraph_type_erased_device_array_create(handle, num_edges, edge_tid, &idx, ret_error); - TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "idx create failed."); - } else { - ret_code = cugraph_type_erased_device_array_create(handle, 0, 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, &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, 0, vertex_tid, &dst, ret_error); - TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "dst create failed."); + 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."); - ret_code = cugraph_type_erased_device_array_create(handle, 0, edge_tid, &idx, ret_error); - TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "wgt create failed."); - } + ret_code = + cugraph_type_erased_device_array_create(handle, num_edges, edge_tid, &idx, ret_error); + TEST_ASSERT(test_ret_value, ret_code == CUGRAPH_SUCCESS, "idx create failed."); src_view = cugraph_type_erased_device_array_view(src); dst_view = cugraph_type_erased_device_array_view(dst); @@ -406,7 +382,7 @@ extern "C" int create_mg_test_graph_with_edge_ids(const cugraph_resource_handle_ idx_view, NULL, store_transposed, - num_edges, + original_num_edges, // UNUSED FALSE, p_graph, ret_error); @@ -464,7 +440,7 @@ extern "C" int create_mg_test_graph_with_properties(const cugraph_resource_handl size_t original_num_edges = num_edges; - if (rank == 0) num_edges = 0; + if (rank != 0) num_edges = 0; ret_code = cugraph_type_erased_device_array_create(handle, num_edges, vertex_tid, &src, ret_error); @@ -529,7 +505,7 @@ extern "C" int create_mg_test_graph_with_properties(const cugraph_resource_handl idx_view, type_view, store_transposed, - original_num_edges, + original_num_edges, // UNUSED FALSE, p_graph, ret_error); @@ -593,7 +569,7 @@ int create_mg_test_graph_new(const cugraph_resource_handle_t* 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."); @@ -658,7 +634,7 @@ int create_mg_test_graph_new(const cugraph_resource_handle_t* handle, edge_id_view, edge_type_view, store_transposed, - renumber, + original_num_edges, // UNUSED FALSE, graph, ret_error); diff --git a/cpp/tests/mtmg/threaded_test.cu b/cpp/tests/mtmg/threaded_test.cu index bc4d8cfef6a..1a6a17eaa18 100644 --- a/cpp/tests/mtmg/threaded_test.cu +++ b/cpp/tests/mtmg/threaded_test.cu @@ -155,10 +155,25 @@ class Tests_Multithreaded input_usecase.template construct_edgelist( handle, multithreaded_usecase.test_weighted, false, false); + rmm::device_uvector d_unique_vertices(2 * d_src_v.size(), handle.get_stream()); + thrust::copy( + handle.get_thrust_policy(), d_src_v.begin(), d_src_v.end(), d_unique_vertices.begin()); + thrust::copy(handle.get_thrust_policy(), + d_dst_v.begin(), + d_dst_v.end(), + d_unique_vertices.begin() + d_src_v.size()); + thrust::sort(handle.get_thrust_policy(), d_unique_vertices.begin(), d_unique_vertices.end()); + + d_unique_vertices.resize(thrust::distance(d_unique_vertices.begin(), + thrust::unique(handle.get_thrust_policy(), + d_unique_vertices.begin(), + d_unique_vertices.end())), + handle.get_stream()); + auto h_src_v = cugraph::test::to_host(handle, d_src_v); auto h_dst_v = cugraph::test::to_host(handle, d_dst_v); auto h_weights_v = cugraph::test::to_host(handle, d_weights_v); - auto unique_vertices = cugraph::test::to_host(handle, d_vertices_v); + auto unique_vertices = cugraph::test::to_host(handle, d_unique_vertices); // Load edgelist from different threads. We'll use more threads than GPUs here for (int i = 0; i < num_threads; ++i) { @@ -293,13 +308,13 @@ class Tests_Multithreaded num_threads]() { auto thread_handle = instance_manager->get_handle(); - auto number_of_vertices = unique_vertices->size(); + auto number_of_vertices = unique_vertices.size(); std::vector my_vertex_list; my_vertex_list.reserve((number_of_vertices + num_threads - 1) / num_threads); for (size_t j = i; j < number_of_vertices; j += num_threads) { - my_vertex_list.push_back((*unique_vertices)[j]); + my_vertex_list.push_back(unique_vertices[j]); } rmm::device_uvector d_my_vertex_list(my_vertex_list.size(), diff --git a/cpp/tests/prims/mg_count_if_e.cu b/cpp/tests/prims/mg_count_if_e.cu index 449aa728d87..03bf8ae0ae5 100644 --- a/cpp/tests/prims/mg_count_if_e.cu +++ b/cpp/tests/prims/mg_count_if_e.cu @@ -53,8 +53,9 @@ #include struct Prims_Usecase { - bool check_correctness{true}; bool test_weighted{false}; + bool edge_masking{false}; + bool check_correctness{true}; }; template @@ -102,6 +103,13 @@ class Tests_MGCountIfE auto mg_graph_view = mg_graph.view(); + std::optional> edge_mask{std::nullopt}; + if (prims_usecase.edge_masking) { + edge_mask = + cugraph::test::generate::edge_property(*handle_, mg_graph_view, 2); + mg_graph_view.attach_edge_mask((*edge_mask).view()); + } + // 2. run MG count_if_e const int hash_bin_count = 5; @@ -148,19 +156,19 @@ class Tests_MGCountIfE (*mg_renumber_map).size()), false); - auto sg_graph_view = sg_graph.view(); + if (handle_->get_comms().get_rank() == 0) { + auto sg_graph_view = sg_graph.view(); - auto sg_vertex_prop = cugraph::test::generate::vertex_property( - *handle_, - thrust::make_counting_iterator(sg_graph_view.local_vertex_partition_range_first()), - thrust::make_counting_iterator(sg_graph_view.local_vertex_partition_range_last()), - hash_bin_count); - auto sg_src_prop = cugraph::test::generate::src_property( - *handle_, sg_graph_view, sg_vertex_prop); - auto sg_dst_prop = cugraph::test::generate::dst_property( - *handle_, sg_graph_view, sg_vertex_prop); + auto sg_vertex_prop = cugraph::test::generate::vertex_property( + *handle_, + thrust::make_counting_iterator(sg_graph_view.local_vertex_partition_range_first()), + thrust::make_counting_iterator(sg_graph_view.local_vertex_partition_range_last()), + hash_bin_count); + auto sg_src_prop = cugraph::test::generate::src_property( + *handle_, sg_graph_view, sg_vertex_prop); + auto sg_dst_prop = cugraph::test::generate::dst_property( + *handle_, sg_graph_view, sg_vertex_prop); - if (handle_->get_comms().get_rank() == 0) { auto expected_result = count_if_e( *handle_, sg_graph_view, @@ -312,7 +320,10 @@ INSTANTIATE_TEST_SUITE_P( file_test, Tests_MGCountIfE_File, ::testing::Combine( - ::testing::Values(Prims_Usecase{true}), + ::testing::Values(Prims_Usecase{false, false, true}, + Prims_Usecase{false, true, true}, + Prims_Usecase{true, false, true}, + Prims_Usecase{true, true, true}), ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx"), cugraph::test::File_Usecase("test/datasets/web-Google.mtx"), cugraph::test::File_Usecase("test/datasets/ljournal-2008.mtx"), @@ -320,7 +331,10 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P(rmat_small_test, Tests_MGCountIfE_Rmat, - ::testing::Combine(::testing::Values(Prims_Usecase{true}), + ::testing::Combine(::testing::Values(Prims_Usecase{false, false, true}, + Prims_Usecase{false, true, true}, + Prims_Usecase{true, false, true}, + Prims_Usecase{true, true, true}), ::testing::Values(cugraph::test::Rmat_Usecase( 10, 16, 0.57, 0.19, 0.19, 0, false, false)))); @@ -332,7 +346,10 @@ INSTANTIATE_TEST_SUITE_P( factor (to avoid running same benchmarks more than once) */ Tests_MGCountIfE_Rmat, ::testing::Combine( - ::testing::Values(Prims_Usecase{false}), + ::testing::Values(Prims_Usecase{false, false, false}, + Prims_Usecase{false, true, false}, + Prims_Usecase{true, false, false}, + Prims_Usecase{true, true, false}), ::testing::Values(cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false)))); CUGRAPH_MG_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_intersection.cu b/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_intersection.cu index a3edb1f6372..ac73c446d89 100644 --- a/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_intersection.cu +++ b/cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_intersection.cu @@ -14,6 +14,8 @@ * limitations under the License. */ +#include "property_generator.cuh" + #include #include #include @@ -116,29 +118,8 @@ class Tests_MGPerVPairTransformDstNbrIntersection std::optional> edge_mask{std::nullopt}; if (prims_usecase.edge_masking) { - cugraph::edge_src_property_t edge_src_renumber_map( - *handle_, mg_graph_view); - cugraph::edge_dst_property_t edge_dst_renumber_map( - *handle_, mg_graph_view); - cugraph::update_edge_src_property( - *handle_, mg_graph_view, (*mg_renumber_map).begin(), edge_src_renumber_map); - cugraph::update_edge_dst_property( - *handle_, mg_graph_view, (*mg_renumber_map).begin(), edge_dst_renumber_map); - - edge_mask = cugraph::edge_property_t(*handle_, mg_graph_view); - - cugraph::transform_e( - *handle_, - mg_graph_view, - edge_src_renumber_map.view(), - edge_dst_renumber_map.view(), - cugraph::edge_dummy_property_t{}.view(), - [] __device__(auto src, auto dst, auto src_property, auto dst_property, thrust::nullopt_t) { - return ((src_property % 2 == 0) && (dst_property % 2 == 0)) - ? false - : true; // mask out the edges with even unrenumbered src & dst vertex IDs - }, - (*edge_mask).mutable_view()); + edge_mask = + cugraph::test::generate::edge_property(*handle_, mg_graph_view, 2); mg_graph_view.attach_edge_mask((*edge_mask).view()); } @@ -257,42 +238,6 @@ class Tests_MGPerVPairTransformDstNbrIntersection if (handle_->get_comms().get_rank() == 0) { auto sg_graph_view = sg_graph.view(); - if (prims_usecase.edge_masking) { - rmm::device_uvector srcs(0, handle_->get_stream()); - rmm::device_uvector dsts(0, handle_->get_stream()); - std::tie(srcs, dsts, std::ignore, std::ignore) = - cugraph::decompress_to_edgelist( - *handle_, sg_graph_view, std::nullopt, std::nullopt, std::nullopt); - auto edge_first = thrust::make_zip_iterator(srcs.begin(), dsts.begin()); - srcs.resize(thrust::distance(edge_first, - thrust::remove_if(handle_->get_thrust_policy(), - edge_first, - edge_first + srcs.size(), - [] __device__(auto pair) { - return (thrust::get<0>(pair) % 2 == 0) && - (thrust::get<1>(pair) % 2 == 0); - })), - handle_->get_stream()); - dsts.resize(srcs.size(), handle_->get_stream()); - rmm::device_uvector vertices(sg_graph_view.number_of_vertices(), - handle_->get_stream()); - thrust::sequence( - handle_->get_thrust_policy(), vertices.begin(), vertices.end(), vertex_t{0}); - std::tie(sg_graph, std::ignore, std::ignore, std::ignore, std::ignore) = cugraph:: - create_graph_from_edgelist( - *handle_, - std::move(vertices), - std::move(srcs), - std::move(dsts), - std::nullopt, - std::nullopt, - std::nullopt, - cugraph::graph_properties_t{sg_graph_view.is_symmetric(), - sg_graph_view.is_multigraph()}, - false); - sg_graph_view = sg_graph.view(); - } - auto sg_result_buffer = cugraph::allocate_dataframe_buffer>( cugraph::size_dataframe_buffer(mg_aggregate_vertex_pair_buffer), handle_->get_stream()); auto sg_out_degrees = sg_graph_view.compute_out_degrees(*handle_); diff --git a/cpp/tests/prims/mg_per_v_random_select_transform_outgoing_e.cu b/cpp/tests/prims/mg_per_v_random_select_transform_outgoing_e.cu index eb6a8fd5cb6..2b9e9aafa3f 100644 --- a/cpp/tests/prims/mg_per_v_random_select_transform_outgoing_e.cu +++ b/cpp/tests/prims/mg_per_v_random_select_transform_outgoing_e.cu @@ -324,8 +324,9 @@ class Tests_MGPerVRandomSelectTransformOutgoingE with_replacement = prims_usecase.with_replacement, invalid_value = invalid_value ? thrust::make_optional(*invalid_value) : thrust::nullopt, - property_transform = cugraph::test::detail::property_transform{ - hash_bin_count}] __device__(size_t i) { + property_transform = + cugraph::test::detail::vertex_property_transform{ + hash_bin_count}] __device__(size_t i) { auto v = *(frontier_vertex_first + i); // check sample_offsets diff --git a/cpp/tests/prims/mg_transform_e.cu b/cpp/tests/prims/mg_transform_e.cu index 24deaad810a..e9be80f1f7d 100644 --- a/cpp/tests/prims/mg_transform_e.cu +++ b/cpp/tests/prims/mg_transform_e.cu @@ -52,6 +52,7 @@ struct Prims_Usecase { bool use_edgelist{false}; + bool edge_masking{false}; bool check_correctness{true}; }; @@ -100,6 +101,13 @@ class Tests_MGTransformE auto mg_graph_view = mg_graph.view(); + std::optional> edge_mask{std::nullopt}; + if (prims_usecase.edge_masking) { + edge_mask = + cugraph::test::generate::edge_property(*handle_, mg_graph_view, 2); + mg_graph_view.attach_edge_mask((*edge_mask).view()); + } + // 2. run MG transform_e const int hash_bin_count = 5; @@ -439,7 +447,10 @@ INSTANTIATE_TEST_SUITE_P( file_test, Tests_MGTransformE_File, ::testing::Combine( - ::testing::Values(Prims_Usecase{false, true}, Prims_Usecase{true, true}), + ::testing::Values(Prims_Usecase{false, false, true}, + Prims_Usecase{false, true, true}, + Prims_Usecase{true, false, true}, + Prims_Usecase{true, true, true}), ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx"), cugraph::test::File_Usecase("test/datasets/web-Google.mtx"), cugraph::test::File_Usecase("test/datasets/ljournal-2008.mtx"), @@ -447,8 +458,10 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P(rmat_small_test, Tests_MGTransformE_Rmat, - ::testing::Combine(::testing::Values(Prims_Usecase{false, true}, - Prims_Usecase{true, true}), + ::testing::Combine(::testing::Values(Prims_Usecase{false, false, true}, + Prims_Usecase{false, true, true}, + Prims_Usecase{true, false, true}, + Prims_Usecase{true, true, true}), ::testing::Values(cugraph::test::Rmat_Usecase( 10, 16, 0.57, 0.19, 0.19, 0, false, false)))); @@ -460,7 +473,10 @@ INSTANTIATE_TEST_SUITE_P( factor (to avoid running same benchmarks more than once) */ Tests_MGTransformE_Rmat, ::testing::Combine( - ::testing::Values(Prims_Usecase{false, false}, Prims_Usecase{true, false}), + ::testing::Values(Prims_Usecase{false, false, false}, + Prims_Usecase{false, true, false}, + Prims_Usecase{true, false, false}, + Prims_Usecase{true, true, false}), ::testing::Values(cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false)))); CUGRAPH_MG_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/prims/mg_transform_reduce_e.cu b/cpp/tests/prims/mg_transform_reduce_e.cu index 79aa3da54df..c4ae11ab7c9 100644 --- a/cpp/tests/prims/mg_transform_reduce_e.cu +++ b/cpp/tests/prims/mg_transform_reduce_e.cu @@ -91,8 +91,9 @@ struct result_compare> { }; struct Prims_Usecase { - bool check_correctness{true}; bool test_weighted{false}; + bool edge_masking{false}; + bool check_correctness{true}; }; template @@ -141,6 +142,13 @@ class Tests_MGTransformReduceE auto mg_graph_view = mg_graph.view(); + std::optional> edge_mask{std::nullopt}; + if (prims_usecase.edge_masking) { + edge_mask = + cugraph::test::generate::edge_property(*handle_, mg_graph_view, 2); + mg_graph_view.attach_edge_mask((*edge_mask).view()); + } + // 2. run MG transform reduce const int hash_bin_count = 5; @@ -365,7 +373,10 @@ INSTANTIATE_TEST_SUITE_P( file_test, Tests_MGTransformReduceE_File, ::testing::Combine( - ::testing::Values(Prims_Usecase{true}), + ::testing::Values(Prims_Usecase{false, false, true}, + Prims_Usecase{false, true, true}, + Prims_Usecase{true, false, true}, + Prims_Usecase{true, true, true}), ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx"), cugraph::test::File_Usecase("test/datasets/web-Google.mtx"), cugraph::test::File_Usecase("test/datasets/ljournal-2008.mtx"), @@ -373,7 +384,10 @@ INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P(rmat_small_test, Tests_MGTransformReduceE_Rmat, - ::testing::Combine(::testing::Values(Prims_Usecase{true}), + ::testing::Combine(::testing::Values(Prims_Usecase{false, false, true}, + Prims_Usecase{false, true, true}, + Prims_Usecase{true, false, true}, + Prims_Usecase{true, true, true}), ::testing::Values(cugraph::test::Rmat_Usecase( 10, 16, 0.57, 0.19, 0.19, 0, false, false)))); @@ -385,7 +399,10 @@ INSTANTIATE_TEST_SUITE_P( factor (to avoid running same benchmarks more than once) */ Tests_MGTransformReduceE_Rmat, ::testing::Combine( - ::testing::Values(Prims_Usecase{false}), + ::testing::Values(Prims_Usecase{false, false, false}, + Prims_Usecase{false, true, false}, + Prims_Usecase{true, false, false}, + Prims_Usecase{true, true, false}), ::testing::Values(cugraph::test::Rmat_Usecase(20, 32, 0.57, 0.19, 0.19, 0, false, false)))); CUGRAPH_MG_TEST_PROGRAM_MAIN() diff --git a/cpp/tests/prims/property_generator.cuh b/cpp/tests/prims/property_generator.cuh index e7264cd276f..680455eda79 100644 --- a/cpp/tests/prims/property_generator.cuh +++ b/cpp/tests/prims/property_generator.cuh @@ -15,6 +15,7 @@ */ #pragma once +#include #include #include @@ -61,7 +62,7 @@ __host__ __device__ auto make_property_value(T val) } template -struct property_transform { +struct vertex_property_transform { int32_t mod{}; constexpr __device__ property_t operator()(vertex_t v) const @@ -73,6 +74,20 @@ struct property_transform { } }; +template +struct edge_property_transform { + int32_t mod{}; + + constexpr __device__ property_t operator()( + vertex_t src, vertex_t dst, thrust::nullopt_t, thrust::nullopt_t, thrust::nullopt_t) const + { + static_assert(cugraph::is_thrust_tuple_of_arithmetic::value || + std::is_arithmetic_v); + cuco::detail::MurmurHash3_32 hash_func{}; + return make_property_value(hash_func(src + dst) % mod); + } +}; + } // namespace detail template @@ -96,7 +111,7 @@ struct generate { labels.begin(), labels.end(), cugraph::get_dataframe_buffer_begin(data), - detail::property_transform{hash_bin_count}); + detail::vertex_property_transform{hash_bin_count}); return data; } @@ -111,7 +126,7 @@ struct generate { begin, end, cugraph::get_dataframe_buffer_begin(data), - detail::property_transform{hash_bin_count}); + detail::vertex_property_transform{hash_bin_count}); return data; } @@ -138,6 +153,22 @@ struct generate { handle, graph_view, cugraph::get_dataframe_buffer_begin(property), output_property); return output_property; } + + template + static auto edge_property(raft::handle_t const& handle, + graph_view_type const& graph_view, + int32_t hash_bin_count) + { + auto output_property = cugraph::edge_property_t(handle, graph_view); + cugraph::transform_e(handle, + graph_view, + cugraph::edge_src_dummy_property_t{}.view(), + cugraph::edge_dst_dummy_property_t{}.view(), + cugraph::edge_dummy_property_t{}.view(), + detail::edge_property_transform{hash_bin_count}, + output_property.mutable_view()); + return output_property; + } }; } // namespace test diff --git a/cpp/tests/structure/mg_select_random_vertices_test.cpp b/cpp/tests/structure/mg_select_random_vertices_test.cpp index 79c50301922..e49e1ebcb99 100644 --- a/cpp/tests/structure/mg_select_random_vertices_test.cpp +++ b/cpp/tests/structure/mg_select_random_vertices_test.cpp @@ -90,7 +90,7 @@ class Tests_MGSelectRandomVertices std::iota( h_given_set.begin(), h_given_set.end(), mg_graph_view.local_vertex_partition_range_first()); std::shuffle(h_given_set.begin(), h_given_set.end(), std::mt19937{std::random_device{}()}); - h_given_set.resize(std::rand() % mg_graph_view.local_vertex_partition_range_size() + 1); + h_given_set.resize(std::rand() % (mg_graph_view.local_vertex_partition_range_size() + 1)); // Compute size of the distributed vertex set int num_of_elements_in_given_set = static_cast(h_given_set.size()); @@ -105,7 +105,7 @@ class Tests_MGSelectRandomVertices size_t select_count = num_of_elements_in_given_set > select_random_vertices_usecase.select_count ? select_random_vertices_usecase.select_count - : std::rand() % num_of_elements_in_given_set + 1; + : std::rand() % (num_of_elements_in_given_set + 1); for (int idx = 0; idx < with_replacement_flags.size(); idx++) { bool with_replacement = with_replacement_flags[idx]; diff --git a/docs/cugraph/source/installation/source_build.md b/docs/cugraph/source/installation/source_build.md index f5ee0741da6..1a129d45295 100644 --- a/docs/cugraph/source/installation/source_build.md +++ b/docs/cugraph/source/installation/source_build.md @@ -1,53 +1,46 @@ # Building from Source -The following instructions are for users wishing to build cuGraph from source code. These instructions are tested on supported distributions of Linux, CUDA, and Python - See [RAPIDS Getting Started](https://rapids.ai/start.html) for list of supported environments. Other operating systems _might be_ compatible, but are not currently tested. - -The cuGraph package include both a C/C++ CUDA portion and a python portion. Both libraries need to be installed in order for cuGraph to operate correctly. +These instructions are tested on supported versions/distributions of Linux, +CUDA, and Python - See [RAPIDS Getting Started](https://rapids.ai/start.html) +for the list of supported environments. Other environments _might be_ +compatible, but are not currently tested. ## Prerequisites -__Compiler:__ +__Compilers:__ * `gcc` version 9.3+ -* `nvcc` version 11.0+ -* `cmake` version 3.20.1+ +* `nvcc` version 11.5+ __CUDA:__ -* CUDA 11.0+ +* CUDA 11.2+ * NVIDIA driver 450.80.02+ * Pascal architecture or better -You can obtain CUDA from [https://developer.nvidia.com/cuda-downloads](https://developer.nvidia.com/cuda-downloads). - -__Packages:__ -* `cmake` version 3.20.1+ -* `libcugraphops` (version matching source branch version, eg. `23.10`) - -You can obtain `libcugraphops` using `conda`/`mamba` from the `nvidia` channel, or using `pip` with the `--extra-index-url=https://pypi.nvidia.com` option. See the [RAPIDS docs](https://docs.rapids.ai/install#environment) for more details. - -## Building cuGraph -To install cuGraph from source, ensure the dependencies are met. +Further details and download links for these prerequisites are available on the +[RAPIDS System Requirements page](https://docs.rapids.ai/install#system-req). +## Setting up the development environment -### Clone Repo and Configure Conda Environment -__GIT clone a version of the repository__ - - ```bash - # Set the localtion to cuGraph in an environment variable CUGRAPH_HOME - export CUGRAPH_HOME=$(pwd)/cugraph - - # Download the cuGraph repo - if you have a folked version, use that path here instead - git clone https://github.com/rapidsai/cugraph.git $CUGRAPH_HOME +### Clone the repository: +```bash +CUGRAPH_HOME=$(pwd)/cugraph +git clone https://github.com/rapidsai/cugraph.git $CUGRAPH_HOME +cd $CUGRAPH_HOME +``` - cd $CUGRAPH_HOME - ``` +### Create the conda environment -__Create the conda development environment__ +Using conda is the easiest way to install both the build and runtime +dependencies for cugraph. While it is possible to build and run cugraph without +conda, the required packages occasionally change, making it difficult to +document here. The best way to see the current dependencies needed for a build +and run environment is to examine the list of packages in the [conda +environment YAML +files](https://github.com/rapidsai/cugraph/blob/main/conda/environments). ```bash -# create the conda environment (assuming in base `cugraph` directory) - # for CUDA 11.x -conda env create --name cugraph_dev --file conda/environments/all_cuda-118_arch-x86_64.yaml +conda env create --name cugraph_dev --file $CUGRAPH_HOME/conda/environments/all_cuda-118_arch-x86_64.yaml # activate the environment conda activate cugraph_dev @@ -56,101 +49,53 @@ conda activate cugraph_dev conda deactivate ``` - - The environment can be updated as development includes/changes the dependencies. To do so, run: - +The environment can be updated as cugraph adds/removes/updates its dependencies. To do so, run: ```bash - -# Where XXX is the CUDA 11 version -conda env update --name cugraph_dev --file conda/environments/cugraph_dev_cuda11.XXX.yml - +# for CUDA 11.x +conda env update --name cugraph_dev --file $CUGRAPH_HOME/conda/environments/all_cuda-118_arch-x86_64.yaml conda activate cugraph_dev ``` +### Build and Install -### Build and Install Using the `build.sh` Script -Using the `build.sh` script make compiling and installing cuGraph a breeze. To build and install, simply do: +#### Build and install using `build.sh` +Using the `build.sh` script, located in the `$CUGRAPH_HOME` directory, is the +recommended way to build and install the cugraph libraries. By default, +`build.sh` will build and install a predefined set of targets +(packages/libraries), but can also accept a list of targets to build. -```bash -$ cd $CUGRAPH_HOME -$ ./build.sh clean -$ ./build.sh libcugraph -$ ./build.sh cugraph -``` +For example, to build only the cugraph C++ library (`libcugraph`) and the +high-level python library (`cugraph`) without building the C++ test binaries, +run this command: -There are several other options available on the build script for advanced users. -`build.sh` options: ```bash -build.sh [ ...] [ ...] - where is: - clean - remove all existing build artifacts and configuration (start over) - uninstall - uninstall libcugraph and cugraph from a prior build/install (see also -n) - libcugraph - build libcugraph.so and SG test binaries - libcugraph_etl - build libcugraph_etl.so and SG test binaries - pylibcugraph - build the pylibcugraph Python package - cugraph - build the cugraph Python package - nx-cugraph - build the nx-cugraph Python package - cugraph-service - build the cugraph-service_client and cugraph-service_server Python package - cpp-mgtests - build libcugraph and libcugraph_etl MG tests. Builds MPI communicator, adding MPI as a dependency. - cugraph-dgl - build the cugraph-dgl extensions for DGL - cugraph-pyg - build the cugraph-dgl extensions for PyG - docs - build the docs - and is: - -v - verbose build mode - -g - build for debug - -n - do not install after a successful build - --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 - --without_cugraphops - do not build algos that require cugraph-ops - --cmake_default_generator - use the default cmake generator instead of ninja - --clean - clean an individual target (note: to do a complete rebuild, use the clean target described above) - -h - print this text - - default action (no args) is to build and install 'libcugraph' then 'libcugraph_etl' then 'pylibcugraph' then 'cugraph' then 'cugraph-service' targets - -examples: -$ ./build.sh clean # remove prior build artifacts (start over) -$ ./build.sh libcugraph -v # compile and install libcugraph with verbose output -$ ./build.sh libcugraph -g # compile and install libcugraph for debug -$ ./build.sh libcugraph -n # compile libcugraph but do not install - -# make parallelism options can also be defined: Example build jobs using 4 threads (make -j4) -$ PARALLEL_LEVEL=4 ./build.sh libcugraph - -Note that the libraries will be installed to the location set in `$PREFIX` if set (i.e. `export PREFIX=/install/path`), otherwise to `$CONDA_PREFIX`. +$ cd $CUGRAPH_HOME +$ ./build.sh libcugraph pylibcugraph cugraph --skip_cpp_tests ``` +There are several other options available on the build script for advanced +users. Refer to the output of `--help` for details. -## Building each section independently -#### Build and Install the C++/CUDA `libcugraph` Library -CMake depends on the `nvcc` executable being on your path or defined in `$CUDACXX`. - -This project uses cmake for building the C/C++ library. To configure cmake, run: - - ```bash - # Set the localtion to cuGraph in an environment variable CUGRAPH_HOME - export CUGRAPH_HOME=$(pwd)/cugraph - - cd $CUGRAPH_HOME - cd cpp # enter cpp directory - mkdir build # create build directory - cd build # enter the build directory - cmake .. -DCMAKE_INSTALL_PREFIX=$CONDA_PREFIX - - # now build the code - make -j # "-j" starts multiple threads - make install # install the libraries - ``` -The default installation locations are `$CMAKE_INSTALL_PREFIX/lib` and `$CMAKE_INSTALL_PREFIX/include/cugraph` respectively. +Note that libraries will be installed to the location set in `$PREFIX` if set +(i.e. `export PREFIX=/install/path`), otherwise to `$CONDA_PREFIX`. #### Updating the RAFT branch -`libcugraph` uses the [RAFT](https://github.com/rapidsai/raft) library and there are times when it might be desirable to build against a different RAFT branch, such as when working on new features that might span both RAFT and cuGraph. +`libcugraph` uses the [RAFT](https://github.com/rapidsai/raft) library and +there are times when it might be desirable to build against a different RAFT +branch, such as when working on new features that might span both RAFT and +cuGraph. -For local development, the `CPM_raft_SOURCE=` option can be passed to the `cmake` command to enable `libcugraph` to use the local RAFT branch. +For local development, the `CPM_raft_SOURCE=` option can +be passed to the `cmake` command to enable `libcugraph` to use the local RAFT +branch. The `build.sh` script calls `cmake` to build the C/C++ targets, but +developers can call `cmake` directly in order to pass it options like those +described here. Refer to the `build.sh` script to see how to call `cmake` and +other commands directly. -To have CI test a `cugraph` pull request against a different RAFT branch, modify the bottom of the `cpp/cmake/thirdparty/get_raft.cmake` file as follows: +To have CI test a `cugraph` pull request against a different RAFT branch, +modify the bottom of the `cpp/cmake/thirdparty/get_raft.cmake` file as follows: ```cmake # Change pinned tag and fork here to test a commit in CI @@ -167,24 +112,10 @@ find_and_configure_raft(VERSION ${CUGRAPH_MIN_VERSION_raft} ) ``` -When the above change is pushed to a pull request, the continuous integration servers will use the specified RAFT branch to run the cuGraph tests. After the changes in the RAFT branch are merged to the release branch, remember to revert the `get_raft.cmake` file back to the original cuGraph branch. - -### Building and installing the Python package - -2) Install the Python packages to your Python path: - -```bash -cd $CUGRAPH_HOME -cd python -cd pylibcugraph -python setup.py build_ext --inplace -python setup.py install # install pylibcugraph -cd ../cugraph -python setup.py build_ext --inplace -python setup.py install # install cugraph python bindings - -``` - +When the above change is pushed to a pull request, the continuous integration +servers will use the specified RAFT branch to run the cuGraph tests. After the +changes in the RAFT branch are merged to the release branch, remember to revert +the `get_raft.cmake` file back to the original cuGraph branch. ## Run tests @@ -240,7 +171,10 @@ Note: This conda installation only applies to Linux and Python versions 3.8/3.10 ### (OPTIONAL) Set environment variable on activation -It is possible to configure the conda environment to set environmental variables on activation. Providing instructions to set PATH to include the CUDA toolkit bin directory and LD_LIBRARY_PATH to include the CUDA lib64 directory will be helpful. +It is possible to configure the conda environment to set environment variables +on activation. Providing instructions to set PATH to include the CUDA toolkit +bin directory and LD_LIBRARY_PATH to include the CUDA lib64 directory will be +helpful. ```bash cd ~/anaconda3/envs/cugraph_dev @@ -271,7 +205,8 @@ unset LD_LIBRARY_PATH ## Creating documentation -Python API documentation can be generated from _./docs/cugraph directory_. Or through using "./build.sh docs" +Python API documentation can be generated from _./docs/cugraph directory_. Or +through using "./build.sh docs" ## Attribution Portions adopted from https://github.com/pytorch/pytorch/blob/master/CONTRIBUTING.md