Add R-mat generator

cpp/CMakeLists.txt

@@ -420,6 +420,7 @@ add_library(cugraph SHARED
     src/components/connectivity.cu
     src/centrality/katz_centrality.cu
     src/centrality/betweenness_centrality.cu
+    src/experimental/generate_rmat_edgelist.cu
     src/experimental/graph.cu
     src/experimental/graph_view.cu
     src/experimental/coarsen_graph.cu

cpp/include/experimental/graph_generator.hpp

@@ -0,0 +1,81 @@
+/*
+ * Copyright (c) 2021, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+#pragma once
+
+#include <raft/handle.hpp>
+#include <rmm/device_uvector.hpp>
+
+#include <cstdint>
+#include <tuple>
+
+namespace cugraph {
+namespace experimental {
+
+/**
+ * @brief generate an edge list for an R-mat graph.
+ *
+ * @p scramble_vertex_ids needs to be set to `true` to generate a graph conforming to the Graph 500
+ * specification (note that scrambling does not affect cuGraph's graph construction performance, so
+ * this is generally unnecessary). If `edge_factor` is given (e.g. Graph 500), set @p num_edges to
+ * (size_t{1} << @p scale) * `edge_factor`. To generate an undirected graph, set @p b == @p c and @p
+ * clip_and_flip = true. All the resulting edges will be placed in the lower triangular part
+ * (inculding the diagonal) of the graph adjacency matrix.
+ *
+ * For multi-GPU generation with `P` GPUs, @p seed should be set to different values in different
+ * GPUs to avoid every GPU generating the same set of edges. @p num_edges should be adjusted as
+ * well; e.g. assuming `edge_factor` is given, set @p num_edges = (size_t{1} << @p scale) *
+ * `edge_factor` / `P` + (rank < (((size_t{1} << @p scale) * `edge_factor`) % P) ? 1 : 0).
+ *
+ * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type.
+ * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
+ * handles to various CUDA libraries) to run graph algorithms.
+ * @param scale Scale factor to set the number of verties in the graph. Vertex IDs have values in
+ * [0, V), where V = 1 << @p scale.
+ * @param num_edges Number of edges to generate.
+ * @param a a, b, c, d (= 1.0 - (a + b + c)) in the R-mat graph generator (vist https://graph500.org
+ * for additional details). a, b, c, d should be non-negative and a + b + c should be no larger
+ * than 1.0.
+ * @param b a, b, c, d (= 1.0 - (a + b + c)) in the R-mat graph generator (vist https://graph500.org
+ * for additional details). a, b, c, d should be non-negative and a + b + c should be no larger
+ * than 1.0.
+ * @param c a, b, c, d (= 1.0 - (a + b + c)) in the R-mat graph generator (vist https://graph500.org
+ * for additional details). a, b, c, d should be non-negative and a + b + c should be no larger
+ * than 1.0.
+ * @param seed Seed value for the random number generator.
+ * @param clip_and_flip Flag controlling whether to generate edges only in the lower triangular part
+ * (including the diagonal) of the graph adjacency matrix (if set to `true`) or not (if set to
+ * `false`).
+ * @param scramble_vertex_ids Flag controlling whether to scramble vertex ID bits (if set to `true`)
+ * or not (if set to `false`); scrambling vertx ID bits breaks correlation between vertex ID values
+ * and vertex degrees. The scramble code here follows the algorithm in the Graph 500 reference
+ * implementation version 3.0.0.
+ * @return std::tuple<rmm::device_uvector<vertex_t>, rmm::device_uvector<vertex_t>> A tuple of
+ * rmm::device_uvector objects for edge source vertex IDs and edge destination vertex IDs.
+ */
+template <typename vertex_t>
+std::tuple<rmm::device_uvector<vertex_t>, rmm::device_uvector<vertex_t>> generate_rmat_edgelist(
+  raft::handle_t const& handle,
+  size_t scale,
+  size_t edge_factor       = 16,
+  double a                 = 0.57,
+  double b                 = 0.19,
+  double c                 = 0.19,
+  uint64_t seed            = 0,
+  bool clip_and_flip       = false,
+  bool scramble_vertex_ids = false);
+
+}  // namespace experimental
+}  // namespace cugraph

cpp/src/experimental/generate_rmat_edgelist.cu

@@ -0,0 +1,149 @@
+/*
+ * Copyright (c) 2021, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#include <experimental/scramble.cuh>
+
+#include <experimental/graph_generator.hpp>
+#include <utilities/error.hpp>
+
+#include <rmm/thrust_rmm_allocator.h>
+#include <raft/handle.hpp>
+#include <raft/random/rng.cuh>
+#include <rmm/device_uvector.hpp>
+
+#include <thrust/iterator/zip_iterator.h>
+#include <thrust/tuple.h>
+
+#include <tuple>
+
+namespace cugraph {
+namespace experimental {
+
+template <typename vertex_t>
+std::tuple<rmm::device_uvector<vertex_t>, rmm::device_uvector<vertex_t>> generate_rmat_edgelist(
+  raft::handle_t const& handle,
+  size_t scale,
+  size_t num_edges,
+  double a,
+  double b,
+  double c,
+  uint64_t seed,
+  bool clip_and_flip,
+  bool scramble_vertex_ids)
+{
+  CUGRAPH_EXPECTS(size_t{1} << scale <= std::numeric_limits<vertex_t>::max(),
+                  "Invalid input argument: scale to large for vertex_t.");
+  CUGRAPH_EXPECTS((a >= 0.0) && (b >= 0.0) && (c >= 0.0) && (a + b + c <= 1.0),
+                  "Invalid input argument: a, b, c should be non-negative and a + b + c should not "
+                  "be larger than 1.0.");
+
+  raft::random::Rng rng(seed + 10);
+  // to limit memory footprint (1024 is a tuning parameter)
+  auto max_edges_to_generate_per_iteration =
+    static_cast<size_t>(handle.get_device_properties().multiProcessorCount) * 1024;
+  rmm::device_uvector<float> rands(
+    std::min(num_edges, max_edges_to_generate_per_iteration) * 2 * scale, handle.get_stream());
+
+  rmm::device_uvector<vertex_t> srcs(num_edges, handle.get_stream());
+  rmm::device_uvector<vertex_t> dsts(num_edges, handle.get_stream());
+
+  size_t num_edges_generated{0};
+  while (num_edges_generated < num_edges) {
+    auto num_edges_to_generate =
+      std::min(num_edges - num_edges_generated, max_edges_to_generate_per_iteration);
+    auto pair_first = thrust::make_zip_iterator(thrust::make_tuple(srcs.begin(), dsts.begin())) +
+                      num_edges_generated;
+    rng.uniform<float, size_t>(
+      rands.data(), num_edges_to_generate * 2 * scale, 0.0f, 1.0f, handle.get_stream());
+    thrust::transform(
+      rmm::exec_policy(handle.get_stream())->on(handle.get_stream()),
+      thrust::make_counting_iterator(size_t{0}),
+      thrust::make_counting_iterator(num_edges_to_generate),
+      pair_first,
+      // if a + b == 0.0, a_norm is irrelevant, if (1.0 - (a+b)) == 0.0, c_norm is irrelevant
+      [scale,
+       clip_and_flip,
+       rands    = rands.data(),
+       a_plus_b = a + b,
+       a_norm   = (a + b) > 0.0 ? a / (a + b) : 0.0,
+       c_norm   = (1.0 - (a + b)) > 0.0 ? c / (1.0 - (a + b)) : 0.0] __device__(auto i) {
+        vertex_t src{0};
+        vertex_t dst{0};
+        for (size_t bit = scale - 1; bit != 0; --bit) {
+          auto r0          = rands[i * 2 * scale + 2 * bit];
+          auto r1          = rands[i * 2 * scale + 2 * bit + 1];
+          auto src_bit_set = r0 > a_plus_b;
+          auto dst_bit_set = r1 > (src_bit_set ? c_norm : a_norm);
+          if (clip_and_flip) {
+            if (src == dst) {
+              if (!src_bit_set && dst_bit_set) {
+                src_bit_set = !src_bit_set;
+                dst_bit_set = !dst_bit_set;
+              }
+            }
+          }
+          src += src_bit_set ? static_cast<vertex_t>(1 << bit) : 0;
+          dst += dst_bit_set ? static_cast<vertex_t>(1 << bit) : 0;
+        }
+        return thrust::make_tuple(src, dst);
+      });
+    num_edges_generated += num_edges_to_generate;
+  }
+
+  if (scramble_vertex_ids) {
+    rands.resize(0, handle.get_stream());
+    rands.shrink_to_fit(handle.get_stream());
+
+    auto pair_first = thrust::make_zip_iterator(thrust::make_tuple(srcs.begin(), dsts.begin()));
+    thrust::transform(rmm::exec_policy(handle.get_stream())->on(handle.get_stream()),
+                      pair_first,
+                      pair_first + srcs.size(),
+                      pair_first,
+                      [scale] __device__(auto pair) {
+                        return thrust::make_tuple(scramble(thrust::get<0>(pair), scale),
+                                                  scramble(thrust::get<1>(pair), scale));
+                      });
+  }
+
+  return std::make_tuple(std::move(srcs), std::move(dsts));
+}
+
+// explicit instantiation
+
+template std::tuple<rmm::device_uvector<int32_t>, rmm::device_uvector<int32_t>>
+generate_rmat_edgelist<int32_t>(raft::handle_t const& handle,
+                                size_t scale,
+                                size_t num_edges,
+                                double a,
+                                double b,
+                                double c,
+                                uint64_t seed,
+                                bool clip_and_flip,
+                                bool scramble_vertex_ids);
+
+template std::tuple<rmm::device_uvector<int64_t>, rmm::device_uvector<int64_t>>
+generate_rmat_edgelist<int64_t>(raft::handle_t const& handle,
+                                size_t scale,
+                                size_t num_edges,
+                                double a,
+                                double b,
+                                double c,
+                                uint64_t seed,
+                                bool clip_and_flip,
+                                bool scramble_vertex_ids);
+
+}  // namespace experimental
+}  // namespace cugraph

cpp/src/experimental/scramble.cuh

@@ -0,0 +1,88 @@
+/* Copyright (C) 2009-2010 The Trustees of Indiana University.             */
+/*                                                                         */
+/* Use, modification and distribution is subject to the Boost Software     */
+/* License, Version 1.0. (See accompanying file LICENSE_1_0.txt or copy at */
+/* http://www.boost.org/LICENSE_1_0.txt)                                   */
+/*                                                                         */
+/*  Authors: Jeremiah Willcock                                             */
+/*           Andrew Lumsdaine                                              */
+
+/*
+ * Copyright (c) 2021, NVIDIA CORPORATION.
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+#pragma once
+
+#include <cassert>
+#include <cstdint>
+
+/* Apply a permutation to scramble vertex numbers; a randomly generated
+ * permutation is not used because applying it at scale is too expensive. */
+template <typename vertex_t>
+__device__ std::enable_if_t<sizeof(vertex_t) == 8, vertex_t> scramble(vertex_t value, size_t lgN)
+{
+  assert(std::is_unsigned<vertex_t>::value || lgN < 64);
+  assert(value >= 0);
+
+  constexpr uint64_t scramble_value0{606610977102444280};    // randomly generated
+  constexpr uint64_t scramble_value1{11680327234415193037};  // randomly generated
+
+  auto v = static_cast<uint64_t>(value);
+  v += scramble_value0 + scramble_value1;
+  v *= (scramble_value0 | uint64_t{0x4519840211493211});
+  v = __brevll(v) >> (64 - lgN);
+  v *= (scramble_value1 | uint64_t{0x3050852102C843A5});
+  v = __brevll(v) >> (64 - lgN);
+  return static_cast<vertex_t>(v);
+}
+
+/* Apply a permutation to scramble vertex numbers; a randomly generated
+ * permutation is not used because applying it at scale is too expensive. */
+template <typename vertex_t>
+__device__ std::enable_if_t<sizeof(vertex_t) == 4, vertex_t> scramble(vertex_t value, size_t lgN)
+{
+  assert(std::is_unsigned<vertex_t>::value || lgN < 32);
+  assert(value >= 0);
+
+  constexpr uint32_t scramble_value0{282475248};   // randomly generated
+  constexpr uint32_t scramble_value1{2617694917};  // randomly generated
+
+  auto v = static_cast<uint32_t>(value);
+  v += scramble_value0 + scramble_value1;
+  v *= (scramble_value0 | uint32_t{0x11493211});
+  v = __brev(v) >> (32 - lgN);
+  v *= (scramble_value1 | uint32_t{0x02C843A5});
+  v = __brev(v) >> (32 - lgN);
+  return static_cast<vertex_t>(v);
+}
+
+/* Apply a permutation to scramble vertex numbers; a randomly generated
+ * permutation is not used because applying it at scale is too expensive. */
+template <typename vertex_t>
+__device__ std::enable_if_t<sizeof(vertex_t) == 2, vertex_t> scramble(vertex_t value, size_t lgN)
+{
+  assert(std::is_unsigned<vertex_t>::value || lgN < 16);
+  assert(value >= 0);
+
+  constexpr uint32_t scramble_value0{0};     // randomly generated
+  constexpr uint32_t scramble_value1{8620};  // randomly generated
+
+  auto v = static_cast<uint16_t>(value);
+  v += scramble_value0 + scramble_value1;
+  v *= (scramble_value0 | uint16_t{0x3211});
+  v = static_cast<uint16_t>(__brev(v) >> 16) >> (16 - lgN);
+  v *= (scramble_value1 | uint16_t{0x43A5});
+  v = static_cast<uint16_t>(__brev(v) >> 16) >> (16 - lgN);
+  return static_cast<vertex_t>(v);
+}

cpp/tests/CMakeLists.txt

@@ -321,6 +321,14 @@ set(MST_TEST_SRC
 ConfigureTest(MST_TEST "${MST_TEST_SRC}")
 
 
+###################################################################################################
+# - Experimental R-mat graph generation tests -----------------------------------------------------
+
+set(EXPERIMENTAL_GENERATE_RMAT_TEST_SRCS
+    "${CMAKE_CURRENT_SOURCE_DIR}/experimental/generate_rmat_test.cpp")
+
+ConfigureTest(EXPERIMENTAL_GENERATE_RMAT_TEST "${EXPERIMENTAL_GENERATE_RMAT_TEST_SRCS}" "")
+
 ###################################################################################################
 # - Experimental Graph tests ----------------------------------------------------------------------
 

cpp/tests/experimental/coarsen_graph_test.cpp

@@ -36,20 +36,6 @@
 #include <type_traits>
 #include <vector>
 
-template <typename vertex_t>
-std::enable_if_t<std::is_signed<vertex_t>::value, bool> is_valid_vertex(vertex_t num_vertices,
-                                                                        vertex_t v)
-{
-  return (v >= 0) && (v < num_vertices);
-}
-
-template <typename vertex_t>
-std::enable_if_t<std::is_unsigned<vertex_t>::value, bool> is_valid_vertex(vertex_t num_vertices,
-                                                                          vertex_t v)
-{
-  return v < num_vertices;
-}
-
 template <typename vertex_t, typename edge_t, typename weight_t>
 void check_coarsened_graph_results(edge_t* org_offsets,
                                    vertex_t* org_indices,
@@ -68,13 +54,13 @@ void check_coarsened_graph_results(edge_t* org_offsets,
   ASSERT_TRUE(std::count_if(org_indices,
                             org_indices + org_offsets[num_org_vertices],
                             [num_org_vertices](auto nbr) {
-                              return !is_valid_vertex(num_org_vertices, nbr);
+                              return !cugraph::test::is_valid_vertex(num_org_vertices, nbr);
                             }) == 0);
   ASSERT_TRUE(std::is_sorted(coarse_offsets, coarse_offsets + num_coarse_vertices));
   ASSERT_TRUE(std::count_if(coarse_indices,
                             coarse_indices + coarse_offsets[num_coarse_vertices],
                             [num_coarse_vertices](auto nbr) {
-                              return !is_valid_vertex(num_coarse_vertices, nbr);
+                              return !cugraph::test::is_valid_vertex(num_coarse_vertices, nbr);
                             }) == 0);
   ASSERT_TRUE(num_coarse_vertices <= num_org_vertices);