Update frontier v push if out nbr prim test (#1985)

- Added testing for update_frontier_v_push_if_out_nbr primitive.
- Fixed extract_if_e testing. Added missing assert.
- Fix clang tools version for 11.5 conda recipe.

Authors:
  - Kumar Aatish (https://github.com/kaatish)

Approvers:
  - Seunghwa Kang (https://github.com/seunghwak)
  - AJ Schmidt (https://github.com/ajschmidt8)

URL: #1985

conda/environments/cugraph_dev_cuda11.5.yml

@@ -19,8 +19,8 @@ dependencies:
 - ucx-proc=*=gpu
 - scipy
 - networkx>=2.5.1
-- clang=11.0.0
-- clang-tools=11.0.0
+- clang=11.1.0
+- clang-tools=11.1.0
 - cmake>=3.20.1
 - python>=3.6,<3.9
 - notebook>=0.5.0

cpp/include/cugraph/prims/extract_if_e.cuh

@@ -27,6 +27,7 @@
 #include <raft/handle.hpp>
 
 #include <cstdint>
+#include <numeric>
 #include <optional>
 #include <tuple>
 #include <type_traits>

cpp/include/cugraph/prims/property_op_utils.cuh

@@ -15,6 +15,7 @@
  */
 #pragma once
 
+#include <cugraph/utilities/error.hpp>
 #include <cugraph/utilities/thrust_tuple_utils.cuh>
 
 #include <raft/comms/comms.hpp>
@@ -148,7 +149,7 @@ struct property_op<thrust::tuple<Args...>, Op>
 };
 
 template <typename T, typename F>
-constexpr auto op_dispatch(raft::comms::op_t op, F&& f)
+auto op_dispatch(raft::comms::op_t op, F&& f)
 {
   switch (op) {
     case raft::comms::op_t::SUM: {

cpp/include/cugraph/prims/transform_reduce_e.cuh

@@ -16,6 +16,7 @@
 #pragma once
 
 #include <cugraph/graph_view.hpp>
+#include <cugraph/matrix_partition_device_view.cuh>
 #include <cugraph/matrix_partition_view.hpp>
 #include <cugraph/prims/property_op_utils.cuh>
 #include <cugraph/utilities/error.hpp>

cpp/tests/CMakeLists.txt

@@ -526,6 +526,11 @@ if(BUILD_CUGRAPH_MG_TESTS)
         # - MG PRIMS COUNT_IF_V tests -------------------------------------------------------------
         ConfigureTestMG(MG_COUNT_IF_V_TEST prims/mg_count_if_v.cu)
 
+        ###########################################################################################
+        # - MG PRIMS UPDATE_FRONTIER_V_PUSH_IF_OUT_NBR tests --------------------------------------
+        ConfigureTestMG(MG_UPDATE_FRONTIER_V_PUSH_IF_OUT_NBR_TEST
+                        prims/mg_update_frontier_v_push_if_out_nbr.cu)
+
         ###########################################################################################
         # - MG PRIMS REDUCE_V tests ---------------------------------------------------------------
         ConfigureTestMG(MG_REDUCE_V_TEST prims/mg_reduce_v.cu)

cpp/tests/prims/mg_extract_if_e.cu

@@ -49,6 +49,19 @@ __device__ auto make_type_casted_tuple_from_scalar(T val, std::index_sequence<Is
     static_cast<typename thrust::tuple_element<Is, TupleType>::type>(val)...);
 }
 
+template <typename property_t, typename T>
+__device__ __host__ auto make_property_value(T val)
+{
+  property_t ret{};
+  if constexpr (cugraph::is_thrust_tuple_of_arithmetic<property_t>::value) {
+    ret = make_type_casted_tuple_from_scalar<property_t>(
+      val, std::make_index_sequence<thrust::tuple_size<property_t>::value>{});
+  } else {
+    ret = static_cast<property_t>(val);
+  }
+  return ret;
+}
+
 template <typename vertex_t, typename property_t>
 struct property_transform_t {
   int mod{};
@@ -58,14 +71,7 @@ struct property_transform_t {
     static_assert(cugraph::is_thrust_tuple_of_arithmetic<property_t>::value ||
                   std::is_arithmetic_v<property_t>);
     cuco::detail::MurmurHash3_32<vertex_t> hash_func{};
-    property_t ret{};
-    if constexpr (cugraph::is_thrust_tuple_of_arithmetic<property_t>::value) {
-      ret = make_type_casted_tuple_from_scalar<property_t>(
-        hash_func(v) % mod, std::make_index_sequence<thrust::tuple_size<property_t>::value>{});
-    } else {
-      ret = static_cast<property_t>(hash_func(v) % mod);
-    }
-    return ret;
+    return make_property_value<property_t>(hash_func(v) % mod);
   }
 };
 
@@ -311,12 +317,13 @@ class Tests_MG_ExtractIfE
                        mg_edge_first + mg_aggregate_edgelist_srcs.size());
           thrust::sort(
             handle.get_thrust_policy(), sg_edge_first, sg_edge_first + sg_edgelist_srcs.size());
-          thrust::equal(handle.get_thrust_policy(),
-                        mg_edge_first,
-                        mg_edge_first + mg_aggregate_edgelist_srcs.size(),
-                        sg_edge_first,
-                        compare_equal_t<
-                          typename thrust::iterator_traits<decltype(mg_edge_first)>::value_type>{});
+          ASSERT_TRUE(thrust::equal(
+            handle.get_thrust_policy(),
+            mg_edge_first,
+            mg_edge_first + mg_aggregate_edgelist_srcs.size(),
+            sg_edge_first,
+            compare_equal_t<
+              typename thrust::iterator_traits<decltype(mg_edge_first)>::value_type>{}));
         } else {
           auto mg_edge_first = thrust::make_zip_iterator(thrust::make_tuple(
             mg_aggregate_edgelist_srcs.begin(), mg_aggregate_edgelist_dsts.begin()));
@@ -327,12 +334,13 @@ class Tests_MG_ExtractIfE
                        mg_edge_first + mg_aggregate_edgelist_srcs.size());
           thrust::sort(
             handle.get_thrust_policy(), sg_edge_first, sg_edge_first + sg_edgelist_srcs.size());
-          thrust::equal(handle.get_thrust_policy(),
-                        mg_edge_first,
-                        mg_edge_first + mg_aggregate_edgelist_srcs.size(),
-                        sg_edge_first,
-                        compare_equal_t<
-                          typename thrust::iterator_traits<decltype(mg_edge_first)>::value_type>{});
+          ASSERT_TRUE(thrust::equal(
+            handle.get_thrust_policy(),
+            mg_edge_first,
+            mg_edge_first + mg_aggregate_edgelist_srcs.size(),
+            sg_edge_first,
+            compare_equal_t<
+              typename thrust::iterator_traits<decltype(mg_edge_first)>::value_type>{}));
         }
       }
     }