Add do_expensive_check to graph primitives (#2274)

Partially address #2003

Add the `do_expensive_check` input parameter to graph primitives.

Authors:
  - Seunghwa Kang (https://github.com/seunghwak)

Approvers:
  - Chuck Hastings (https://github.com/ChuckHastings)

URL: #2274

cpp/include/cugraph/prims/count_if_e.cuh

@@ -52,6 +52,7 @@ namespace cugraph {
  * @param e_op Quaternary (or quinary) operator takes edge source, edge destination, (optional edge
  * weight), property values for the source, and property values for the destination and returns if
  * this edge should be included in the returned count.
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
  * @return GraphViewType::edge_type Number of times @p e_op returned true.
  */
 template <typename GraphViewType,
@@ -63,11 +64,16 @@ typename GraphViewType::edge_type count_if_e(
   GraphViewType const& graph_view,
   EdgePartitionSrcValueInputWrapper edge_partition_src_value_input,
   EdgePartitionDstValueInputWrapper edge_partition_dst_value_input,
-  EdgeOp e_op)
+  EdgeOp e_op,
+  bool do_expensive_check = false)
 {
   using vertex_t = typename GraphViewType::vertex_type;
   using edge_t   = typename GraphViewType::edge_type;
 
+  if (do_expensive_check) {
+    // currently, nothing to do
+  }
+
   return transform_reduce_e(handle,
                             graph_view,
                             edge_partition_src_value_input,

cpp/include/cugraph/prims/count_if_v.cuh

@@ -63,16 +63,22 @@ struct count_if_call_v_op_t {
  * @param v_op Binary operator takes vertex ID and *(@p vertex_value_input_first + i) (where i is
  * [0, @p graph_view.local_vertex_partition_range_size())) and returns true if this vertex should be
  * included in the returned count.
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
  * @return GraphViewType::vertex_type Number of times @p v_op returned true.
  */
 template <typename GraphViewType, typename VertexValueInputIterator, typename VertexOp>
 typename GraphViewType::vertex_type count_if_v(raft::handle_t const& handle,
                                                GraphViewType const& graph_view,
                                                VertexValueInputIterator vertex_value_input_first,
-                                               VertexOp v_op)
+                                               VertexOp v_op,
+                                               bool do_expensive_check = false)
 {
   using vertex_t = typename GraphViewType::vertex_type;
 
+  if (do_expensive_check) {
+    // currently, nothing to do
+  }
+
   auto it = thrust::make_transform_iterator(
     thrust::make_counting_iterator(vertex_t{0}),
     detail::count_if_call_v_op_t<vertex_t, VertexValueInputIterator, VertexOp>{

cpp/include/cugraph/prims/extract_if_e.cuh

@@ -113,6 +113,7 @@ struct call_e_op_t {
  * weight), property values for the source, and property values for the destination and returns a
  * boolean value to designate whether to include this edge in the returned edge list (if true is
  * returned) or not (if false is returned).
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
  * @return std::tuple<rmm::device_uvector<typename GraphViewType::vertex_type>,
  * rmm::device_uvector<typename GraphViewType::vertex_type>,
  * std::optional<rmm::device_uvector<typename GraphViewType::weight_type>>> Tuple storing an
@@ -129,12 +130,17 @@ extract_if_e(raft::handle_t const& handle,
              GraphViewType const& graph_view,
              EdgePartitionSrcValueInputWrapper edge_partition_src_value_input,
              EdgePartitionDstValueInputWrapper edge_partition_dst_value_input,
-             EdgeOp e_op)
+             EdgeOp e_op,
+             bool do_expensive_check = false)
 {
   using vertex_t = typename GraphViewType::vertex_type;
   using edge_t   = typename GraphViewType::edge_type;
   using weight_t = typename GraphViewType::weight_type;
 
+  if (do_expensive_check) {
+    // currently, nothing to do
+  }
+
   std::vector<size_t> edgelist_edge_counts(graph_view.number_of_local_edge_partitions(), size_t{0});
   for (size_t i = 0; i < edgelist_edge_counts.size(); ++i) {
     edgelist_edge_counts[i] =

cpp/include/cugraph/prims/per_src_dst_key_transform_reduce_e.cuh

@@ -591,6 +591,7 @@ per_src_dst_key_transform_reduce_e(
  * transformed value to be reduced to (source key, value) pairs.
  * @param init Initial value to be added to the value in each transform-reduced (source key, value)
  * pair.
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
  * @return std::tuple Tuple of rmm::device_uvector<typename GraphView::vertex_type> and
  * rmm::device_uvector<T> (if T is arithmetic scalar) or a tuple of rmm::device_uvector objects (if
  * T is a thrust::tuple type of arithmetic scalar types, one rmm::device_uvector object per scalar
@@ -609,12 +610,17 @@ auto per_src_key_transform_reduce_e(
   EdgePartitionDstValueInputWrapper edge_partition_dst_value_input,
   EdgePartitionSrcKeyInputWrapper edge_partition_src_key_input,
   EdgeOp e_op,
-  T init)
+  T init,
+  bool do_expensive_check = false)
 {
   static_assert(is_arithmetic_or_thrust_tuple_of_arithmetic<T>::value);
   static_assert(std::is_same<typename EdgePartitionSrcKeyInputWrapper::value_type,
                              typename GraphViewType::vertex_type>::value);
 
+  if (do_expensive_check) {
+    // currently, nothing to do
+  }
+
   return detail::per_src_dst_key_transform_reduce_e<true>(handle,
                                                           graph_view,
                                                           edge_partition_src_value_input,
@@ -663,6 +669,7 @@ auto per_src_key_transform_reduce_e(
  * transformed value to be reduced to (destination key, value) pairs.
  * @param init Initial value to be added to the value in each transform-reduced (destination key,
  * value) pair.
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
  * @return std::tuple Tuple of rmm::device_uvector<typename GraphView::vertex_type> and
  * rmm::device_uvector<T> (if T is arithmetic scalar) or a tuple of rmm::device_uvector objects (if
  * T is a thrust::tuple type of arithmetic scalar types, one rmm::device_uvector object per scalar
@@ -681,12 +688,17 @@ auto per_dst_key_transform_reduce_e(
   EdgePartitionDstValueInputWrapper edge_partition_dst_value_input,
   EdgePartitionDstKeyInputWrapper edge_partition_dst_key_input,
   EdgeOp e_op,
-  T init)
+  T init,
+  bool do_expensive_check = false)
 {
   static_assert(is_arithmetic_or_thrust_tuple_of_arithmetic<T>::value);
   static_assert(std::is_same<typename EdgePartitionDstKeyInputWrapper::value_type,
                              typename GraphViewType::vertex_type>::value);
 
+  if (do_expensive_check) {
+    // currently, nothing to do
+  }
+
   return detail::per_src_dst_key_transform_reduce_e<false>(handle,
                                                            graph_view,
                                                            edge_partition_src_value_input,

cpp/include/cugraph/prims/per_v_transform_reduce_dst_key_aggregated_outgoing_e.cuh

@@ -198,6 +198,7 @@ struct reduce_with_init_t {
  * first (inclusive) vertex (assigned to this process in multi-GPU). `vertex_value_output_last`
  * (exclusive) is deduced as @p vertex_value_output_first + @p
  * graph_view.local_vertex_partition_range_size().
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
  */
 template <typename GraphViewType,
           typename EdgePartitionSrcValueInputWrapper,
@@ -219,7 +220,8 @@ void per_v_transform_reduce_dst_key_aggregated_outgoing_e(
   KeyAggregatedEdgeOp key_aggregated_e_op,
   T init,
   ReduceOp reduce_op,
-  VertexValueOutputIterator vertex_value_output_first)
+  VertexValueOutputIterator vertex_value_output_first,
+  bool do_expensive_check = false)
 {
   static_assert(!GraphViewType::is_storage_transposed,
                 "GraphViewType should support the push model.");
@@ -234,6 +236,46 @@ void per_v_transform_reduce_dst_key_aggregated_outgoing_e(
 
   double constexpr load_factor = 0.7;
 
+  if (do_expensive_check) {
+    rmm::device_uvector<vertex_t> keys(thrust::distance(map_unique_key_first, map_unique_key_last),
+                                       handle.get_stream());
+    thrust::copy(
+      handle.get_thrust_policy(), map_unique_key_first, map_unique_key_last, keys.begin());
+    thrust::sort(handle.get_thrust_policy(), keys.begin(), keys.end());
+    auto has_duplicates =
+      (thrust::unique(handle.get_thrust_policy(), keys.begin(), keys.end()) != keys.end());
+
+    if constexpr (GraphViewType::is_multi_gpu) {
+      auto& comm           = handle.get_comms();
+      auto const comm_size = comm.get_size();
+      auto const comm_rank = comm.get_rank();
+
+      auto num_invalid_keys = thrust::count_if(
+        handle.get_thrust_policy(),
+        map_unique_key_first,
+        map_unique_key_last,
+        [comm_rank,
+         key_func = detail::compute_gpu_id_from_ext_vertex_t<vertex_t>{
+           comm_size}] __device__(auto key) { return key_func(key) != comm_rank; });
+      num_invalid_keys =
+        host_scalar_allreduce(comm, num_invalid_keys, raft::comms::op_t::SUM, handle.get_stream());
+      CUGRAPH_EXPECTS(
+        num_invalid_keys == 0,
+        "Invalid input argument: map (unique key, value) pairs should be pre-shuffled.");
+
+      has_duplicates =
+        host_scalar_allreduce(
+          comm, has_duplicates ? int{1} : int{0}, raft::comms::op_t::MAX, handle.get_stream()) ==
+            int{1}
+          ? true
+          : false;
+    }
+
+    CUGRAPH_EXPECTS(has_duplicates == false,
+                    "Invalid input argument: there are duplicates in [map_unique_key_first, "
+                    "map_unique_key_last).");
+  }
+
   auto total_global_mem = handle.get_device_properties().totalGlobalMem;
   auto element_size     = sizeof(vertex_t) * 2 + sizeof(weight_t);
   auto constexpr mem_frugal_ratio =

cpp/include/cugraph/prims/per_v_transform_reduce_incoming_outgoing_e.cuh

@@ -916,6 +916,7 @@ void per_v_transform_reduce_e(raft::handle_t const& handle,
  * (inclusive) vertex (assigned to this process in multi-GPU). `vertex_value_output_last`
  * (exclusive) is deduced as @p vertex_value_output_first + @p
  * graph_view.local_vertex_partition_range_size().
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
  */
 template <typename GraphViewType,
           typename EdgePartitionSrcValueInputWrapper,
@@ -930,8 +931,13 @@ void per_v_transform_reduce_incoming_e(
   EdgePartitionDstValueInputWrapper edge_partition_dst_value_input,
   EdgeOp e_op,
   T init,
-  VertexValueOutputIterator vertex_value_output_first)
+  VertexValueOutputIterator vertex_value_output_first,
+  bool do_expensive_check = false)
 {
+  if (do_expensive_check) {
+    // currently, nothing to do
+  }
+
   detail::per_v_transform_reduce_e<true>(handle,
                                          graph_view,
                                          edge_partition_src_value_input,
@@ -975,6 +981,7 @@ void per_v_transform_reduce_incoming_e(
  * first (inclusive) vertex (assigned to this process in multi-GPU). `vertex_value_output_last`
  * (exclusive) is deduced as @p vertex_value_output_first + @p
  * graph_view.local_vertex_partition_range_size().
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
  */
 template <typename GraphViewType,
           typename EdgePartitionSrcValueInputWrapper,
@@ -989,8 +996,13 @@ void per_v_transform_reduce_outgoing_e(
   EdgePartitionDstValueInputWrapper edge_partition_dst_value_input,
   EdgeOp e_op,
   T init,
-  VertexValueOutputIterator vertex_value_output_first)
+  VertexValueOutputIterator vertex_value_output_first,
+  bool do_expensive_check = false)
 {
+  if (do_expensive_check) {
+    // currently, nothing to do
+  }
+
   detail::per_v_transform_reduce_e<false>(handle,
                                           graph_view,
                                           edge_partition_src_value_input,

cpp/include/cugraph/prims/reduce_v.cuh

@@ -52,14 +52,16 @@ namespace cugraph {
  * future) implementations of graph primitives may check whether @p ReduceOp is a known type (or has
  * known member variables) to take a more optimized code path. See the documentation in the
  * reduce_op.cuh file for instructions on writing custom reduction operators.
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
  * @return T Reduced input vertex property values.
  */
 template <typename GraphViewType, typename ReduceOp, typename VertexValueInputIterator, typename T>
 T reduce_v(raft::handle_t const& handle,
            GraphViewType const& graph_view,
            VertexValueInputIterator vertex_value_input_first,
            T init,
-           ReduceOp reduce_op)
+           ReduceOp reduce_op,
+           bool do_expensive_check = false)
 {
   using vertex_t = typename GraphViewType::vertex_type;
 
@@ -68,6 +70,10 @@ T reduce_v(raft::handle_t const& handle,
       std::remove_cv_t<typename thrust::iterator_traits<VertexValueInputIterator>::value_type>,
       std::remove_cv_t<T>>);
 
+  if (do_expensive_check) {
+    // currently, nothing to do
+  }
+
   if (graph_view.number_of_vertices() == vertex_t{0}) { return init; }
 
   T ret{};
@@ -183,14 +189,20 @@ T reduce_v(raft::handle_t const& handle,
  * (inclusive) vertex (assigned to this process in multi-GPU). `vertex_value_input_last` (exclusive)
  * is deduced as @p vertex_value_input_first + @p graph_view.local_vertex_partition_range_size().
  * @param init Initial value to be added to the reduced input vertex property values.
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
  * @return T Reduced input vertex property values.
  */
 template <typename GraphViewType, typename VertexValueInputIterator, typename T>
 T reduce_v(raft::handle_t const& handle,
            GraphViewType const& graph_view,
            VertexValueInputIterator vertex_value_input_first,
-           T init)
+           T init,
+           bool do_expensive_check = false)
 {
+  if (do_expensive_check) {
+    // currently, nothing to do
+  }
+
   return reduce_v(handle, graph_view, vertex_value_input_first, init, reduce_op::plus<T>{});
 }
 
@@ -209,16 +221,22 @@ T reduce_v(raft::handle_t const& handle,
  * @param vertex_value_input_first Iterator pointing to the vertex property values for the first
  * (inclusive) vertex (assigned to this process in multi-GPU). `vertex_value_input_last` (exclusive)
  * is deduced as @p vertex_value_input_first + @p graph_view.local_vertex_partition_range_size().
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
  * @return Reduced input vertex property values.
  */
 template <typename GraphViewType, typename VertexValueInputIterator>
 auto reduce_v(raft::handle_t const& handle,
               GraphViewType const& graph_view,
-              VertexValueInputIterator vertex_value_input_first)
+              VertexValueInputIterator vertex_value_input_first,
+              bool do_expensive_check = false)
 {
   using T =
     std::remove_cv_t<typename thrust::iterator_traits<VertexValueInputIterator>::value_type>;
 
+  if (do_expensive_check) {
+    // currently, nothing to do
+  }
+
   return reduce_v(handle, graph_view, vertex_value_input_first, T{}, reduce_op::plus<T>{});
 }
 

cpp/include/cugraph/prims/transform_reduce_e.cuh

@@ -384,6 +384,7 @@ __global__ void trasnform_reduce_e_high_degree(
  * weight), property values for the source, and property values for the destination and returns a
  * value to be reduced.
  * @param init Initial value to be added to the reduced @p edge_op outputs.
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
  * @return T Transform-reduced @p edge_op outputs.
  */
 template <typename GraphViewType,
@@ -396,14 +397,19 @@ T transform_reduce_e(raft::handle_t const& handle,
                      EdgePartitionSrcValueInputWrapper edge_partition_src_value_input,
                      EdgePartitionDstValueInputWrapper edge_partition_dst_value_input,
                      EdgeOp e_op,
-                     T init)
+                     T init,
+                     bool do_expensive_check = false)
 {
   static_assert(is_arithmetic_or_thrust_tuple_of_arithmetic<T>::value);
 
   using vertex_t = typename GraphViewType::vertex_type;
   using edge_t   = typename GraphViewType::edge_type;
   using weight_t = typename GraphViewType::weight_type;
 
+  if (do_expensive_check) {
+    // currently, nothing to do
+  }
+
   property_op<T, thrust::plus> edge_property_add{};
 
   auto result_buffer = allocate_dataframe_buffer<T>(1, handle.get_stream());
@@ -547,6 +553,7 @@ T transform_reduce_e(raft::handle_t const& handle,
  * @param e_op Quaternary (or quinary) operator takes edge source, edge destination, (optional edge
  * weight), property values for the source, and property values for the destination and returns a
  * value to be reduced.
+ * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
  * @return Transform-reduced @p edge_op outputs.
  */
 template <typename GraphViewType,
@@ -557,7 +564,8 @@ auto transform_reduce_e(raft::handle_t const& handle,
                         GraphViewType const& graph_view,
                         EdgePartitionSrcValueInputWrapper edge_partition_src_value_input,
                         EdgePartitionDstValueInputWrapper edge_partition_dst_value_input,
-                        EdgeOp e_op)
+                        EdgeOp e_op,
+                        bool do_expensive_check = false)
 {
   using vertex_t    = typename GraphViewType::vertex_type;
   using weight_t    = typename GraphViewType::weight_type;
@@ -566,6 +574,10 @@ auto transform_reduce_e(raft::handle_t const& handle,
   using T           = typename detail::
     edge_op_result_type<vertex_t, vertex_t, weight_t, src_value_t, dst_value_t, EdgeOp>::type;
 
+  if (do_expensive_check) {
+    // currently, nothing to do
+  }
+
   return transform_reduce_e(
     handle, graph_view, edge_partition_src_value_input, edge_partition_dst_value_input, e_op, T{});
 }