Copy v transform reduce out test

cpp/include/cugraph/prims/reduce_v.cuh

@@ -55,9 +55,9 @@ T reduce_v(raft::handle_t const& handle,
     handle.get_thrust_policy(),
     vertex_value_input_first,
     vertex_value_input_first + graph_view.get_number_of_local_vertices(),
-    ((GraphViewType::is_multi_gpu) && (handle.get_comms().get_rank() == 0)) ? init : T{},
+    ((GraphViewType::is_multi_gpu) && (handle.get_comms().get_rank() != 0)) ? T{} : init,
     property_add<T>());
-  if (GraphViewType::is_multi_gpu) {
+  if constexpr (GraphViewType::is_multi_gpu) {
     ret =
       host_scalar_allreduce(handle.get_comms(), ret, raft::comms::op_t::SUM, handle.get_stream());
   }
@@ -93,9 +93,9 @@ T reduce_v(raft::handle_t const& handle,
     handle.get_thrust_policy(),
     input_first,
     input_last,
-    ((GraphViewType::is_multi_gpu) && (handle.get_comms().get_rank() == 0)) ? init : T{},
+    ((GraphViewType::is_multi_gpu) && (handle.get_comms().get_rank() != 0)) ? T{} : init,
     property_add<T>());
-  if (GraphViewType::is_multi_gpu) {
+  if constexpr (GraphViewType::is_multi_gpu) {
     ret =
       host_scalar_allreduce(handle.get_comms(), ret, raft::comms::op_t::SUM, handle.get_stream());
   }

cpp/include/cugraph/prims/transform_reduce_e.cuh

@@ -410,11 +410,10 @@ T transform_reduce_e(raft::handle_t const& handle,
   property_add<T> edge_property_add{};
 
   auto result_buffer = allocate_dataframe_buffer<T>(1, handle.get_stream());
-  thrust::fill(
-    handle.get_thrust_policy(),
-    get_dataframe_buffer_begin(result_buffer),
-    get_dataframe_buffer_begin(result_buffer) + 1,
-    ((GraphViewType::is_multi_gpu) && (handle.get_comms().get_rank() == 0)) ? init : T{});
+  thrust::fill(handle.get_thrust_policy(),
+               get_dataframe_buffer_begin(result_buffer),
+               get_dataframe_buffer_begin(result_buffer) + 1,
+               T{});
 
   for (size_t i = 0; i < graph_view.get_number_of_local_adj_matrix_partitions(); ++i) {
     auto matrix_partition =
@@ -510,13 +509,14 @@ T transform_reduce_e(raft::handle_t const& handle,
     }
   }
 
-  auto result = thrust::reduce(handle.get_thrust_policy(),
-                               get_dataframe_buffer_begin(result_buffer),
-                               get_dataframe_buffer_begin(result_buffer) + 1,
-                               T{},
-                               edge_property_add);
+  auto result = thrust::reduce(
+    handle.get_thrust_policy(),
+    get_dataframe_buffer_begin(result_buffer),
+    get_dataframe_buffer_begin(result_buffer) + 1,
+    ((GraphViewType::is_multi_gpu) && (handle.get_comms().get_rank() != 0)) ? T{} : init,
+    edge_property_add);
 
-  if (GraphViewType::is_multi_gpu) {
+  if constexpr (GraphViewType::is_multi_gpu) {
     result = host_scalar_allreduce(
       handle.get_comms(), result, raft::comms::op_t::SUM, handle.get_stream());
   }

cpp/include/cugraph/prims/transform_reduce_v.cuh

@@ -60,7 +60,7 @@ T transform_reduce_v(raft::handle_t const& handle,
     vertex_value_input_first,
     vertex_value_input_first + graph_view.get_number_of_local_vertices(),
     v_op,
-    ((GraphViewType::is_multi_gpu) && (handle.get_comms().get_rank() == 0)) ? init : T{},
+    ((GraphViewType::is_multi_gpu) && (handle.get_comms().get_rank() != 0)) ? T{} : init,
     property_add<T>());
   if (GraphViewType::is_multi_gpu) {
     ret =
@@ -104,7 +104,7 @@ T transform_reduce_v(raft::handle_t const& handle,
     input_first,
     input_last,
     v_op,
-    ((GraphViewType::is_multi_gpu) && (handle.get_comms().get_rank() == 0)) ? init : T{},
+    ((GraphViewType::is_multi_gpu) && (handle.get_comms().get_rank() != 0)) ? T{} : init,
     property_add<T>());
   if (GraphViewType::is_multi_gpu) {
     ret =

cpp/tests/CMakeLists.txt

@@ -468,7 +468,7 @@ if(BUILD_CUGRAPH_MG_TESTS)
 
         ###########################################################################################
         # - MG PRIMS COPY_V_TRANSFORM_REDUCE_OUT tests --------------------------------------------
-        ConfigureTestMG(MG_COPY_V_TRANSFORM_REDUCE_OUT_TEST prims/mg_copy_v_transform_reduce_out_nbr.cu)
+        ConfigureTestMG(MG_COPY_V_TRANSFORM_REDUCE_INOUT_NBR_TEST prims/mg_copy_v_transform_reduce_inout_nbr.cu)
     else()
        message(FATAL_ERROR "OpenMPI NOT found, cannot build MG tests.")
     endif()

cpp/tests/prims/mg_copy_v_transform_reduce_out_nbr.cu → cpp/tests/prims/mg_copy_v_transform_reduce_inout_nbr.cu

@@ -222,10 +222,10 @@ struct Prims_Usecase {
 };
 
 template <typename input_usecase_t>
-class Tests_MG_CopyVTransformReduceOutNbr
+class Tests_MG_CopyVTransformReduceInOutNbr
   : public ::testing::TestWithParam<std::tuple<Prims_Usecase, input_usecase_t>> {
  public:
-  Tests_MG_CopyVTransformReduceOutNbr() {}
+  Tests_MG_CopyVTransformReduceInOutNbr() {}
   static void SetupTestCase() {}
   static void TearDownTestCase() {}
 
@@ -281,16 +281,18 @@ class Tests_MG_CopyVTransformReduceOutNbr
     // 3. run MG transform reduce
 
     const int hash_bin_count = 5;
-    const int initial_value  = 0;
+    const int initial_value  = 4;
 
     auto property_initial_value = generate<result_t>::initial_value(initial_value);
     using property_t            = decltype(property_initial_value);
     auto vertex_property_data =
       generate<result_t>::vertex_property((*d_mg_renumber_map_labels), hash_bin_count, handle);
     auto col_prop =
       generate<result_t>::column_property(handle, mg_graph_view, vertex_property_data);
-    auto row_prop = generate<result_t>::row_property(handle, mg_graph_view, vertex_property_data);
-    auto result   = cugraph::allocate_dataframe_buffer<property_t>(
+    auto row_prop   = generate<result_t>::row_property(handle, mg_graph_view, vertex_property_data);
+    auto out_result = cugraph::allocate_dataframe_buffer<property_t>(
+      mg_graph_view.get_number_of_local_vertices(), handle.get_stream());
+    auto in_result = cugraph::allocate_dataframe_buffer<property_t>(
       mg_graph_view.get_number_of_local_vertices(), handle.get_stream());
 
     if (cugraph::test::g_perf) {
@@ -299,6 +301,35 @@ class Tests_MG_CopyVTransformReduceOutNbr
       hr_clock.start();
     }
 
+    copy_v_transform_reduce_in_nbr(
+      handle,
+      mg_graph_view,
+      row_prop.device_view(),
+      col_prop.device_view(),
+      [] __device__(auto row, auto col, weight_t wt, auto row_property, auto col_property) {
+        if (row_property < col_property) {
+          return row_property;
+        } else {
+          return col_property;
+        }
+      },
+      property_initial_value,
+      cugraph::get_dataframe_buffer_begin(in_result));
+
+    if (cugraph::test::g_perf) {
+      CUDA_TRY(cudaDeviceSynchronize());  // for consistent performance measurement
+      handle.get_comms().barrier();
+      double elapsed_time{0.0};
+      hr_clock.stop(&elapsed_time);
+      std::cout << "MG copy v transform reduce in took " << elapsed_time * 1e-6 << " s.\n";
+    }
+
+    if (cugraph::test::g_perf) {
+      CUDA_TRY(cudaDeviceSynchronize());  // for consistent performance measurement
+      handle.get_comms().barrier();
+      hr_clock.start();
+    }
+
     copy_v_transform_reduce_out_nbr(
       handle,
       mg_graph_view,
@@ -312,7 +343,7 @@ class Tests_MG_CopyVTransformReduceOutNbr
         }
       },
       property_initial_value,
-      cugraph::get_dataframe_buffer_begin(result));
+      cugraph::get_dataframe_buffer_begin(out_result));
 
     if (cugraph::test::g_perf) {
       CUDA_TRY(cudaDeviceSynchronize());  // for consistent performance measurement
@@ -340,8 +371,9 @@ class Tests_MG_CopyVTransformReduceOutNbr
         generate<result_t>::column_property(handle, sg_graph_view, sg_vertex_property_data);
       auto sg_row_prop =
         generate<result_t>::row_property(handle, sg_graph_view, sg_vertex_property_data);
+      result_compare comp{handle};
 
-      auto global_result = cugraph::allocate_dataframe_buffer<property_t>(
+      auto global_out_result = cugraph::allocate_dataframe_buffer<property_t>(
         sg_graph_view.get_number_of_local_vertices(), handle.get_stream());
       copy_v_transform_reduce_out_nbr(
         handle,
@@ -356,70 +388,90 @@ class Tests_MG_CopyVTransformReduceOutNbr
           }
         },
         property_initial_value,
-        cugraph::get_dataframe_buffer_begin(global_result));
-      auto expected_result = permute_result(handle, global_result, (*d_mg_renumber_map_labels));
-      result_compare comp{handle};
-      ASSERT_TRUE(comp(result, expected_result));
+        cugraph::get_dataframe_buffer_begin(global_out_result));
+      auto expected_out_result =
+        permute_result(handle, global_out_result, (*d_mg_renumber_map_labels));
+      ASSERT_TRUE(comp(out_result, expected_out_result));
+
+      auto global_in_result = cugraph::allocate_dataframe_buffer<property_t>(
+        sg_graph_view.get_number_of_local_vertices(), handle.get_stream());
+      copy_v_transform_reduce_in_nbr(
+        handle,
+        sg_graph_view,
+        sg_row_prop.device_view(),
+        sg_col_prop.device_view(),
+        [] __device__(auto row, auto col, weight_t wt, auto row_property, auto col_property) {
+          if (row_property < col_property) {
+            return row_property;
+          } else {
+            return col_property;
+          }
+        },
+        property_initial_value,
+        cugraph::get_dataframe_buffer_begin(global_in_result));
+      auto expected_in_result =
+        permute_result(handle, global_in_result, (*d_mg_renumber_map_labels));
+      ASSERT_TRUE(comp(in_result, expected_in_result));
     }
   }
 };
 
-using Tests_MG_CopyVTransformReduceOutNbr_File =
-  Tests_MG_CopyVTransformReduceOutNbr<cugraph::test::File_Usecase>;
-using Tests_MG_CopyVTransformReduceOutNbr_Rmat =
-  Tests_MG_CopyVTransformReduceOutNbr<cugraph::test::Rmat_Usecase>;
+using Tests_MG_CopyVTransformReduceInOutNbr_File =
+  Tests_MG_CopyVTransformReduceInOutNbr<cugraph::test::File_Usecase>;
+using Tests_MG_CopyVTransformReduceInOutNbr_Rmat =
+  Tests_MG_CopyVTransformReduceInOutNbr<cugraph::test::Rmat_Usecase>;
 
-TEST_P(Tests_MG_CopyVTransformReduceOutNbr_File, CheckInt32Int32FloatTupleIntFloatTransposeFalse)
+TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_File, CheckInt32Int32FloatTupleIntFloatTransposeFalse)
 {
   auto param = GetParam();
   run_current_test<int32_t, int32_t, float, std::tuple<int, float>, false>(std::get<0>(param),
                                                                            std::get<1>(param));
 }
 
-TEST_P(Tests_MG_CopyVTransformReduceOutNbr_Rmat, CheckInt32Int32FloatTupleIntFloatTransposeFalse)
+TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_Rmat, CheckInt32Int32FloatTupleIntFloatTransposeFalse)
 {
   auto param = GetParam();
   run_current_test<int32_t, int32_t, float, std::tuple<int, float>, false>(
     std::get<0>(param),
     cugraph::test::override_Rmat_Usecase_with_cmd_line_arguments(std::get<1>(param)));
 }
 
-TEST_P(Tests_MG_CopyVTransformReduceOutNbr_File, CheckInt32Int32FloatTupleIntFloatTransposeTrue)
+TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_File, CheckInt32Int32FloatTupleIntFloatTransposeTrue)
 {
   auto param = GetParam();
   run_current_test<int32_t, int32_t, float, std::tuple<int, float>, true>(std::get<0>(param),
                                                                           std::get<1>(param));
 }
 
-TEST_P(Tests_MG_CopyVTransformReduceOutNbr_Rmat, CheckInt32Int32FloatTupleIntFloatTransposeTrue)
+TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_Rmat, CheckInt32Int32FloatTupleIntFloatTransposeTrue)
 {
   auto param = GetParam();
   run_current_test<int32_t, int32_t, float, std::tuple<int, float>, true>(
     std::get<0>(param),
     cugraph::test::override_Rmat_Usecase_with_cmd_line_arguments(std::get<1>(param)));
 }
 
-TEST_P(Tests_MG_CopyVTransformReduceOutNbr_File, CheckInt32Int32FloatTransposeFalse)
+TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_File, CheckInt32Int32FloatTransposeFalse)
 {
   auto param = GetParam();
   run_current_test<int32_t, int32_t, float, int, false>(std::get<0>(param), std::get<1>(param));
 }
 
-TEST_P(Tests_MG_CopyVTransformReduceOutNbr_Rmat, CheckInt32Int32FloatTransposeFalse)
+TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_Rmat, CheckInt32Int32FloatTransposeFalse)
 {
   auto param = GetParam();
   run_current_test<int32_t, int32_t, float, int, false>(
     std::get<0>(param),
     cugraph::test::override_Rmat_Usecase_with_cmd_line_arguments(std::get<1>(param)));
 }
 
-TEST_P(Tests_MG_CopyVTransformReduceOutNbr_File, CheckInt32Int32FloatTransposeTrue)
+TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_File, CheckInt32Int32FloatTransposeTrue)
 {
   auto param = GetParam();
   run_current_test<int32_t, int32_t, float, int, true>(std::get<0>(param), std::get<1>(param));
 }
 
-TEST_P(Tests_MG_CopyVTransformReduceOutNbr_Rmat, CheckInt32Int32FloatTransposeTrue)
+TEST_P(Tests_MG_CopyVTransformReduceInOutNbr_Rmat, CheckInt32Int32FloatTransposeTrue)
 {
   auto param = GetParam();
   run_current_test<int32_t, int32_t, float, int, true>(
@@ -429,7 +481,7 @@ TEST_P(Tests_MG_CopyVTransformReduceOutNbr_Rmat, CheckInt32Int32FloatTransposeTr
 
 INSTANTIATE_TEST_SUITE_P(
   file_test,
-  Tests_MG_CopyVTransformReduceOutNbr_File,
+  Tests_MG_CopyVTransformReduceInOutNbr_File,
   ::testing::Combine(
     ::testing::Values(Prims_Usecase{true}),
     ::testing::Values(cugraph::test::File_Usecase("test/datasets/karate.mtx"),
@@ -439,14 +491,14 @@ INSTANTIATE_TEST_SUITE_P(
 
 INSTANTIATE_TEST_SUITE_P(
   rmat_small_test,
-  Tests_MG_CopyVTransformReduceOutNbr_Rmat,
+  Tests_MG_CopyVTransformReduceInOutNbr_Rmat,
   ::testing::Combine(::testing::Values(Prims_Usecase{true}),
                      ::testing::Values(cugraph::test::Rmat_Usecase(
                        10, 16, 0.57, 0.19, 0.19, 0, false, false, 0, true))));
 
 INSTANTIATE_TEST_SUITE_P(
   rmat_large_test,
-  Tests_MG_CopyVTransformReduceOutNbr_Rmat,
+  Tests_MG_CopyVTransformReduceInOutNbr_Rmat,
   ::testing::Combine(::testing::Values(Prims_Usecase{false}),
                      ::testing::Values(cugraph::test::Rmat_Usecase(
                        20, 32, 0.57, 0.19, 0.19, 0, false, false, 0, true))));

cpp/tests/prims/mg_count_if_e.cu

@@ -221,10 +221,7 @@ class Tests_MG_TransformCountIfE
     // 3. run MG count_if_e
 
     const int hash_bin_count = 5;
-    const int initial_value  = 0;
 
-    auto property_initial_value = generate<result_t>::initial_value(initial_value);
-    using property_t            = decltype(property_initial_value);
     auto vertex_property_data =
       generate<result_t>::vertex_property((*d_mg_renumber_map_labels), hash_bin_count, handle);
     auto col_prop =

cpp/tests/prims/mg_transform_reduce_e.cu

@@ -259,7 +259,7 @@ class Tests_MG_TransformReduceE
     // 3. run MG transform reduce
 
     const int hash_bin_count = 5;
-    const int initial_value  = 0;
+    const int initial_value  = 4;
 
     auto property_initial_value = generate<result_t>::initial_value(initial_value);
     using property_t            = decltype(property_initial_value);