From 996c7a46526c24c2f7d046cb587f72c5f5da4ff1 Mon Sep 17 00:00:00 2001
From: Ralph Liu <137829296+nv-rliu@users.noreply.github.com>
Date: Fri, 30 Aug 2024 20:45:08 -0400
Subject: [PATCH 1/2] Add `nx-cugraph` Benchmarking Scripts (#4616)
Closes https://github.com/rapidsai/graph_dl/issues/596
This PR adds scripts written by @rlratzel to generate benchmarking numbers for `nx-cugraph`
Authors:
- Ralph Liu (https://github.com/nv-rliu)
- Rick Ratzel (https://github.com/rlratzel)
Approvers:
- Don Acosta (https://github.com/acostadon)
- Rick Ratzel (https://github.com/rlratzel)
URL: https://github.com/rapidsai/cugraph/pull/4616
---
benchmarks/nx-cugraph/pytest-based/README.md | 45 +++
.../nx-cugraph/pytest-based/bench_algos.py | 31 +-
.../create_results_summary_page.py | 291 ++++++++++++++++++
.../pytest-based/get_graph_bench_dataset.py | 35 +++
.../pytest-based/run-main-benchmarks.sh | 58 ++++
5 files changed, 458 insertions(+), 2 deletions(-)
create mode 100644 benchmarks/nx-cugraph/pytest-based/README.md
create mode 100644 benchmarks/nx-cugraph/pytest-based/create_results_summary_page.py
create mode 100644 benchmarks/nx-cugraph/pytest-based/get_graph_bench_dataset.py
create mode 100755 benchmarks/nx-cugraph/pytest-based/run-main-benchmarks.sh
diff --git a/benchmarks/nx-cugraph/pytest-based/README.md b/benchmarks/nx-cugraph/pytest-based/README.md
new file mode 100644
index 00000000000..4ea0f127a51
--- /dev/null
+++ b/benchmarks/nx-cugraph/pytest-based/README.md
@@ -0,0 +1,45 @@
+## `nx-cugraph` Benchmarks
+
+### Overview
+
+This directory contains a set of scripts designed to benchmark NetworkX with the `nx-cugraph` backend and deliver a report that summarizes the speed-up and runtime deltas over default NetworkX.
+
+Our current benchmarks provide the following datasets:
+
+| Dataset | Nodes | Edges | Directed |
+| -------- | ------- | ------- | ------- |
+| netscience | 1,461 | 5,484 | Yes |
+| email-Eu-core | 1,005 | 25,571 | Yes |
+| cit-Patents | 3,774,768 | 16,518,948 | Yes |
+| hollywood | 1,139,905 | 57,515,616 | No |
+| soc-LiveJournal1 | 4,847,571 | 68,993,773 | Yes |
+
+
+
+### Scripts
+
+#### 1. `run-main-benchmarks.sh`
+This script allows users to run selected algorithms across multiple datasets and backends. All results are stored inside a sub-directory (`logs/`) and output files are named based on the combination of parameters for that benchmark.
+
+NOTE: If running with all algorithms, datasets, and backends, this script may take a few hours to finish running.
+
+**Usage:**
+ ```bash
+ bash run-main-benchmarks.sh # edit this script directly
+ ```
+
+#### 2. `get_graph_bench_dataset.py`
+This script downloads the specified dataset using `cugraph.datasets`.
+
+**Usage:**
+ ```bash
+ python get_graph_bench_dataset.py [dataset]
+ ```
+
+#### 3. `create_results_summary_page.py`
+This script is designed to be run after `run-gap-benchmarks.sh` in order to generate an HTML page displaying a results table comparing default NetworkX to nx-cugraph. The script also provides information about the current system.
+
+**Usage:**
+ ```bash
+ python create_results_summary_page.py > report.html
+ ```
diff --git a/benchmarks/nx-cugraph/pytest-based/bench_algos.py b/benchmarks/nx-cugraph/pytest-based/bench_algos.py
index d40b5130827..f88d93c3f17 100644
--- a/benchmarks/nx-cugraph/pytest-based/bench_algos.py
+++ b/benchmarks/nx-cugraph/pytest-based/bench_algos.py
@@ -271,9 +271,8 @@ def bench_from_networkx(benchmark, graph_obj):
# normalized_param_values = [True, False]
-# k_param_values = [10, 100]
normalized_param_values = [True]
-k_param_values = [10]
+k_param_values = [10, 100, 1000]
@pytest.mark.parametrize(
@@ -282,6 +281,10 @@ def bench_from_networkx(benchmark, graph_obj):
@pytest.mark.parametrize("k", k_param_values, ids=lambda k: f"{k=}")
def bench_betweenness_centrality(benchmark, graph_obj, backend_wrapper, normalized, k):
G = get_graph_obj_for_benchmark(graph_obj, backend_wrapper)
+
+ if k > G.number_of_nodes():
+ pytest.skip(reason=f"{k=} > {G.number_of_nodes()=}")
+
result = benchmark.pedantic(
target=backend_wrapper(nx.betweenness_centrality),
args=(G,),
@@ -305,6 +308,10 @@ def bench_edge_betweenness_centrality(
benchmark, graph_obj, backend_wrapper, normalized, k
):
G = get_graph_obj_for_benchmark(graph_obj, backend_wrapper)
+
+ if k > G.number_of_nodes():
+ pytest.skip(reason=f"{k=} > {G.number_of_nodes()=}")
+
result = benchmark.pedantic(
target=backend_wrapper(nx.edge_betweenness_centrality),
args=(G,),
@@ -473,6 +480,26 @@ def bench_pagerank_personalized(benchmark, graph_obj, backend_wrapper):
assert type(result) is dict
+def bench_shortest_path(benchmark, graph_obj, backend_wrapper):
+ """
+ This passes in the source node with the highest degree, but no target.
+ """
+ G = get_graph_obj_for_benchmark(graph_obj, backend_wrapper)
+ node = get_highest_degree_node(graph_obj)
+
+ result = benchmark.pedantic(
+ target=backend_wrapper(nx.shortest_path),
+ args=(G,),
+ kwargs=dict(
+ source=node,
+ ),
+ rounds=rounds,
+ iterations=iterations,
+ warmup_rounds=warmup_rounds,
+ )
+ assert type(result) is dict
+
+
def bench_single_source_shortest_path_length(benchmark, graph_obj, backend_wrapper):
G = get_graph_obj_for_benchmark(graph_obj, backend_wrapper)
node = get_highest_degree_node(graph_obj)
diff --git a/benchmarks/nx-cugraph/pytest-based/create_results_summary_page.py b/benchmarks/nx-cugraph/pytest-based/create_results_summary_page.py
new file mode 100644
index 00000000000..f1cc4b06ccc
--- /dev/null
+++ b/benchmarks/nx-cugraph/pytest-based/create_results_summary_page.py
@@ -0,0 +1,291 @@
+# Copyright (c) 2024, 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.
+
+
+import re
+import pathlib
+import json
+import platform
+import psutil
+import socket
+import subprocess
+
+
+def get_formatted_time_value(time):
+ res = ""
+ if time < 1:
+ if time < 0.001:
+ units = "us"
+ time *= 1e6
+ else:
+ units = "ms"
+ time *= 1e3
+ else:
+ units = "s"
+ return f"{time:.3f}{units}"
+
+
+def get_all_benchmark_info():
+ benchmarks = {}
+ # Populate benchmarks dir from .json files
+ for json_file in logs_dir.glob("*.json"):
+ try:
+ data = json.loads(open(json_file).read())
+ except json.decoder.JSONDecodeError:
+ continue
+
+ for benchmark_run in data["benchmarks"]:
+ # example name: "bench_triangles[ds=netscience-backend=cugraph-preconverted]"
+ name = benchmark_run["name"]
+
+ algo_name = name.split("[")[0]
+ if algo_name.startswith("bench_"):
+ algo_name = algo_name[6:]
+ # special case for betweenness_centrality
+ match = k_patt.match(name)
+ if match is not None:
+ algo_name += f", k={match.group(1)}"
+
+ match = dataset_patt.match(name)
+ if match is None:
+ raise RuntimeError(
+ f"benchmark name {name} in file {json_file} has an unexpected format"
+ )
+ dataset = match.group(1)
+ if dataset.endswith("-backend"):
+ dataset = dataset[:-8]
+
+ match = backend_patt.match(name)
+ if match is None:
+ raise RuntimeError(
+ f"benchmark name {name} in file {json_file} has an unexpected format"
+ )
+ backend = match.group(1)
+ if backend == "None":
+ backend = "networkx"
+
+ runtime = benchmark_run["stats"]["mean"]
+ benchmarks.setdefault(algo_name, {}).setdefault(backend, {})[
+ dataset
+ ] = runtime
+ return benchmarks
+
+
+def compute_perf_vals(cugraph_runtime, networkx_runtime):
+ speedup_string = f"{networkx_runtime / cugraph_runtime:.3f}X"
+ delta = networkx_runtime - cugraph_runtime
+ if abs(delta) < 1:
+ if abs(delta) < 0.001:
+ units = "us"
+ delta *= 1e6
+ else:
+ units = "ms"
+ delta *= 1e3
+ else:
+ units = "s"
+ delta_string = f"{delta:.3f}{units}"
+
+ return (speedup_string, delta_string)
+
+
+def get_mem_info():
+ return round(psutil.virtual_memory().total / (1024**3), 2)
+
+
+def get_cuda_version():
+ output = subprocess.check_output("nvidia-smi", shell=True).decode()
+ try:
+ return next(
+ line.split("CUDA Version: ")[1].split()[0]
+ for line in output.splitlines()
+ if "CUDA Version" in line
+ )
+ except subprocess.CalledProcessError:
+ return "Failed to get CUDA version."
+
+
+def get_first_gpu_info():
+ try:
+ gpu_info = (
+ subprocess.check_output(
+ "nvidia-smi --query-gpu=name,memory.total,memory.free,memory.used --format=csv,noheader",
+ shell=True,
+ )
+ .decode()
+ .strip()
+ )
+ if gpu_info:
+ gpus = gpu_info.split("\n")
+ num_gpus = len(gpus)
+ first_gpu = gpus[0] # Get the information for the first GPU
+ gpu_name, mem_total, _, _ = first_gpu.split(",")
+ return f"{num_gpus} x {gpu_name.strip()} ({round(int(mem_total.strip().split()[0]) / (1024), 2)} GB)"
+ else:
+ print("No GPU found or unable to query GPU details.")
+ except subprocess.CalledProcessError:
+ print("Failed to execute nvidia-smi. No GPU information available.")
+
+
+def get_system_info():
+ print('
')
+ print(f"
Hostname: {socket.gethostname()}
")
+ print(
+ f'
Operating System: {platform.system()} {platform.release()}
'
+ )
+ print(f'
Kernel Version : {platform.version()}
')
+ with open("/proc/cpuinfo") as f:
+ print(
+ f'
CPU: {next(line.strip().split(": ")[1] for line in f if "model name" in line)} ({psutil.cpu_count(logical=False)} cores)
'
+ )
+ print(f'
Memory: {get_mem_info()} GB
')
+ print(f"
GPU: {get_first_gpu_info()}
")
+ print(f"
CUDA Version: {get_cuda_version()}
")
+
+
+if __name__ == "__main__":
+ logs_dir = pathlib.Path("logs")
+
+ dataset_patt = re.compile(".*ds=([\w-]+).*")
+ backend_patt = re.compile(".*backend=(\w+).*")
+ k_patt = re.compile(".*k=(10*).*")
+
+ # Organize all benchmark runs by the following hierarchy: algo -> backend -> dataset
+ benchmarks = get_all_benchmark_info()
+
+ # dump HTML table
+ ordered_datasets = [
+ "netscience",
+ "email_Eu_core",
+ "cit-patents",
+ "hollywood",
+ "soc-livejournal1",
+ ]
+ # dataset, # Node, # Edge, Directed info
+ dataset_meta = {
+ "netscience": ["1,461", "5,484", "Yes"],
+ "email_Eu_core": ["1,005", "25,571", "Yes"],
+ "cit-patents": ["3,774,768", "16,518,948", "Yes"],
+ "hollywood": ["1,139,905", "57,515,616", "No"],
+ "soc-livejournal1": ["4,847,571", "68,993,773", "Yes"],
+ }
+
+ print(
+ """
+
+
+
+
+
+
+
+ Dataset Nodes Edges Directed | """
+ )
+ for ds in ordered_datasets:
+ print(
+ f" {ds} {dataset_meta[ds][0]} {dataset_meta[ds][1]} {dataset_meta[ds][2]}
| "
+ )
+ print(
+ """
+
+
+ """
+ )
+ for algo_name in sorted(benchmarks):
+ algo_runs = benchmarks[algo_name]
+ print(" ")
+ print(f" {algo_name} | ")
+ # Proceed only if any results are present for both cugraph and NX
+ if "cugraph" in algo_runs and "networkx" in algo_runs:
+ cugraph_algo_runs = algo_runs["cugraph"]
+ networkx_algo_runs = algo_runs["networkx"]
+ datasets_in_both = set(cugraph_algo_runs).intersection(networkx_algo_runs)
+
+ # populate the table with speedup results for each dataset in the order
+ # specified in ordered_datasets. If results for a run using a dataset
+ # are not present for both cugraph and NX, output an empty cell.
+ for dataset in ordered_datasets:
+ if dataset in datasets_in_both:
+ cugraph_runtime = cugraph_algo_runs[dataset]
+ networkx_runtime = networkx_algo_runs[dataset]
+ (speedup, runtime_delta) = compute_perf_vals(
+ cugraph_runtime=cugraph_runtime,
+ networkx_runtime=networkx_runtime,
+ )
+ nx_formatted = get_formatted_time_value(networkx_runtime)
+ cg_formatted = get_formatted_time_value(cugraph_runtime)
+ print(
+ f" {nx_formatted} / {cg_formatted} {speedup} {runtime_delta} | "
+ )
+ else:
+ print(f" | ")
+
+ # If a comparison between cugraph and NX cannot be made, output empty cells
+ # for each dataset
+ else:
+ for _ in range(len(ordered_datasets)):
+ print(" | ")
+ print("
")
+ print(
+ """
+ \n
+
+
+
Table Format:
+
+ - NetworkX time / nx-cugraph time
+ - Speed-up of using nx-cugraph
+ - Time-delta
+
+
"""
+ )
+ get_system_info()
+ print("""
\n
\n""")
diff --git a/benchmarks/nx-cugraph/pytest-based/get_graph_bench_dataset.py b/benchmarks/nx-cugraph/pytest-based/get_graph_bench_dataset.py
new file mode 100644
index 00000000000..5a0a15da8ee
--- /dev/null
+++ b/benchmarks/nx-cugraph/pytest-based/get_graph_bench_dataset.py
@@ -0,0 +1,35 @@
+# Copyright (c) 2024, 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.
+
+"""
+Checks if a particular dataset has been downloaded inside the datasets dir
+(RAPIDS_DATAEST_ROOT_DIR). If not, the file will be downloaded using the
+datasets API.
+
+Positional Arguments:
+ 1) dataset name (e.g. 'email_Eu_core', 'cit-patents')
+ available datasets can be found here: `python/cugraph/cugraph/datasets/__init__.py`
+"""
+
+import sys
+
+import cugraph.datasets as cgds
+
+
+if __name__ == "__main__":
+ # download and store dataset (csv) by using the Datasets API
+ dataset = sys.argv[1].replace("-", "_")
+ dataset_obj = getattr(cgds, dataset)
+
+ if not dataset_obj.get_path().exists():
+ dataset_obj.get_edgelist(download=True)
diff --git a/benchmarks/nx-cugraph/pytest-based/run-main-benchmarks.sh b/benchmarks/nx-cugraph/pytest-based/run-main-benchmarks.sh
new file mode 100755
index 00000000000..1a81fe4b80a
--- /dev/null
+++ b/benchmarks/nx-cugraph/pytest-based/run-main-benchmarks.sh
@@ -0,0 +1,58 @@
+#!/bin/bash
+# Copyright (c) 2024, 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.
+
+
+# location to store datasets used for benchmarking
+export RAPIDS_DATASET_ROOT_DIR=/datasets/cugraph
+mkdir -p logs
+
+# list of algos, datasets, and back-ends to use in combinations
+algos="
+ pagerank
+ betweenness_centrality
+ louvain
+ shortest_path
+ weakly_connected_components
+ triangles
+ bfs_predecessors
+"
+datasets="
+ netscience
+ email_Eu_core
+ cit_patents
+ hollywood
+ soc-livejournal
+"
+# None backend is default networkx
+# cugraph-preconvert backend is nx-cugraph
+backends="
+ None
+ cugraph-preconverted
+"
+
+for algo in $algos; do
+ for dataset in $datasets; do
+ python get_graph_bench_dataset.py $dataset
+ for backend in $backends; do
+ name="${backend}__${algo}__${dataset}"
+ echo "Running: $backend, $dataset, bench_$algo"
+ # command to preproduce test
+ # echo "RUNNING: \"pytest -sv -k \"$backend and $dataset and bench_$algo and not 1000\" --benchmark-json=\"logs/${name}.json\" bench_algos.py"
+ pytest -sv \
+ -k "$backend and $dataset and bench_$algo and not 1000" \
+ --benchmark-json="logs/${name}.json" \
+ bench_algos.py 2>&1 | tee "logs/${name}.out"
+ done
+ done
+done
From 338e5e0944eeb7533288a274624720b730a21a81 Mon Sep 17 00:00:00 2001
From: Seunghwa Kang <45857425+seunghwak@users.noreply.github.com>
Date: Tue, 3 Sep 2024 09:07:14 -0700
Subject: [PATCH 2/2] Heterogeneous renumbering implementation (#4602)
This PR implements heterogeneous renumbering for GNN.
In addition,
* Update the existing (homogeneous) sampling post processing function test file extension from .cu to .cpp.
* Remove the unused `renumber_sampled_edgelist` function (breaking because this function is removed)
* Add a `stride_fill` utility function (thrust wrapper)
* Add test utility functions to generate edge types & IDs.
Closes #4412
Authors:
- Seunghwa Kang (https://github.com/seunghwak)
Approvers:
- Chuck Hastings (https://github.com/ChuckHastings)
- Alex Barghi (https://github.com/alexbarghi-nv)
URL: https://github.com/rapidsai/cugraph/pull/4602
---
cpp/CMakeLists.txt | 2 -
.../cugraph/detail/utility_wrappers.hpp | 22 +
cpp/include/cugraph/graph_functions.hpp | 57 -
cpp/include/cugraph/sampling_functions.hpp | 84 +-
cpp/src/detail/utility_wrappers_32.cu | 17 +
cpp/src/detail/utility_wrappers_64.cu | 12 +
cpp/src/detail/utility_wrappers_impl.cuh | 16 +
.../renumber_sampled_edgelist_impl.cuh | 719 -------
.../renumber_sampled_edgelist_sg_v32_e32.cu | 37 -
.../renumber_sampled_edgelist_sg_v64_e64.cu | 37 -
.../sampling_post_processing_impl.cuh | 1638 ++++++++++++++--
.../sampling_post_processing_sg_v32_e32.cu | 56 +
.../sampling_post_processing_sg_v32_e64.cu | 56 +
.../sampling_post_processing_sg_v64_e64.cu | 56 +
cpp/tests/CMakeLists.txt | 11 +-
.../sampling_post_processing_validate.cu | 1738 +++++++++++++++++
.../sampling_post_processing_validate.hpp | 101 +
...ing_heterogeneous_post_processing_test.cpp | 828 ++++++++
...t.cu => sampling_post_processing_test.cpp} | 1151 ++++-------
.../property_generator_utilities.hpp | 23 +
.../property_generator_utilities_impl.cuh | 98 +
cpp/tests/utilities/thrust_wrapper.cu | 69 +
cpp/tests/utilities/thrust_wrapper.hpp | 14 +
23 files changed, 4971 insertions(+), 1871 deletions(-)
delete mode 100644 cpp/src/sampling/renumber_sampled_edgelist_impl.cuh
delete mode 100644 cpp/src/sampling/renumber_sampled_edgelist_sg_v32_e32.cu
delete mode 100644 cpp/src/sampling/renumber_sampled_edgelist_sg_v64_e64.cu
create mode 100644 cpp/tests/sampling/detail/sampling_post_processing_validate.cu
create mode 100644 cpp/tests/sampling/detail/sampling_post_processing_validate.hpp
create mode 100644 cpp/tests/sampling/sampling_heterogeneous_post_processing_test.cpp
rename cpp/tests/sampling/{sampling_post_processing_test.cu => sampling_post_processing_test.cpp} (52%)
diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt
index 26b710247f6..b8eaba9d575 100644
--- a/cpp/CMakeLists.txt
+++ b/cpp/CMakeLists.txt
@@ -338,8 +338,6 @@ set(CUGRAPH_SOURCES
src/sampling/negative_sampling_mg_v32_e64.cu
src/sampling/negative_sampling_mg_v32_e32.cu
src/sampling/negative_sampling_mg_v64_e64.cu
- src/sampling/renumber_sampled_edgelist_sg_v64_e64.cu
- src/sampling/renumber_sampled_edgelist_sg_v32_e32.cu
src/sampling/sampling_post_processing_sg_v64_e64.cu
src/sampling/sampling_post_processing_sg_v32_e32.cu
src/sampling/sampling_post_processing_sg_v32_e64.cu
diff --git a/cpp/include/cugraph/detail/utility_wrappers.hpp b/cpp/include/cugraph/detail/utility_wrappers.hpp
index 61ac1bd2804..3d99b85556b 100644
--- a/cpp/include/cugraph/detail/utility_wrappers.hpp
+++ b/cpp/include/cugraph/detail/utility_wrappers.hpp
@@ -87,6 +87,28 @@ void sequence_fill(rmm::cuda_stream_view const& stream_view,
size_t size,
value_t start_value);
+/**
+ * @brief Fill a buffer with a sequence of values with the input stride
+ *
+ * Fills the buffer with the sequence with the input stride:
+ * {start_value, start_value+stride, start_value+stride*2, ..., start_value+stride*(size-1)}
+ *
+ * @tparam value_t type of the value to operate on
+ *
+ * @param[in] stream_view stream view
+ * @param[out] d_value device array to fill
+ * @param[in] size number of elements in array
+ * @param[in] start_value starting value for sequence
+ * @param[in] stride input stride
+ *
+ */
+template
+void stride_fill(rmm::cuda_stream_view const& stream_view,
+ value_t* d_value,
+ size_t size,
+ value_t start_value,
+ value_t stride);
+
/**
* @brief Compute the maximum vertex id of an edge list
*
diff --git a/cpp/include/cugraph/graph_functions.hpp b/cpp/include/cugraph/graph_functions.hpp
index 7f6543ccab8..866ab16ee97 100644
--- a/cpp/include/cugraph/graph_functions.hpp
+++ b/cpp/include/cugraph/graph_functions.hpp
@@ -988,63 +988,6 @@ rmm::device_uvector select_random_vertices(
bool sort_vertices,
bool do_expensive_check = false);
-/**
- * @brief renumber sampling output
- *
- * @deprecated This API will be deprecated and will be replaced by the
- * renumber_and_compress_sampled_edgelist and renumber_and_sort_sampled_edgelist functions in
- * sampling_functions.hpp.
- *
- * This function renumbers sampling function (e.g. uniform_neighbor_sample) outputs satisfying the
- * following requirements.
- *
- * 1. If @p edgelist_hops is valid, we can consider (vertex ID, flag=src, hop) triplets for each
- * vertex ID in @p edgelist_srcs and (vertex ID, flag=dst, hop) triplets for each vertex ID in @p
- * edgelist_dsts. From these triplets, we can find the minimum (hop, flag) pairs for every unique
- * vertex ID (hop is the primary key and flag is the secondary key, flag=src is considered smaller
- * than flag=dst if hop numbers are same). Vertex IDs with smaller (hop, flag) pairs precede vertex
- * IDs with larger (hop, flag) pairs in renumbering. Ordering can be arbitrary among the vertices
- * with the same (hop, flag) pairs.
- * 2. If @p edgelist_hops is invalid, unique vertex IDs in @p edgelist_srcs precede vertex IDs that
- * appear only in @p edgelist_dsts.
- * 3. If label_offsets.has_value() is ture, edge lists for different labels will be renumbered
- * separately.
- *
- * This function is single-GPU only (we are not aware of any practical multi-GPU use cases).
- *
- * @tparam vertex_t Type of vertex identifiers. Needs to be an integral type.
- * @tparam label_t Type of labels. 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 edgelist_srcs A vector storing original edgelist source vertices.
- * @param edgelist_dsts A vector storing original edgelist destination vertices (size = @p
- * edgelist_srcs.size()).
- * @param edgelist_hops An optional pointer to the array storing hops for each edge list (source,
- * destination) pairs (size = @p edgelist_srcs.size() if valid).
- * @param label_offsets An optional tuple of unique labels and the input edge list (@p
- * edgelist_srcs, @p edgelist_hops, and @p edgelist_dsts) offsets for the labels (siez = # unique
- * labels + 1).
- * @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
- * @return Tuple of vectors storing renumbered edge sources (size = @p edgelist_srcs.size()) ,
- * renumbered edge destinations (size = @p edgelist_dsts.size()), renumber_map to query original
- * verties (size = # unique vertices or aggregate # unique vertices for every label), and
- * renumber_map offsets (size = std::get<0>(*label_offsets).size() + 1, valid only if @p
- * label_offsets.has_value() is true).
- */
-template
-std::tuple,
- rmm::device_uvector,
- rmm::device_uvector,
- std::optional>>
-renumber_sampled_edgelist(
- raft::handle_t const& handle,
- rmm::device_uvector&& edgelist_srcs,
- rmm::device_uvector&& edgelist_dsts,
- std::optional> edgelist_hops,
- std::optional, raft::device_span>>
- label_offsets,
- bool do_expensive_check = false);
-
/**
* @brief Remove self loops from an edge list
*
diff --git a/cpp/include/cugraph/sampling_functions.hpp b/cpp/include/cugraph/sampling_functions.hpp
index 4e5596d06e0..783cd3a7e2b 100644
--- a/cpp/include/cugraph/sampling_functions.hpp
+++ b/cpp/include/cugraph/sampling_functions.hpp
@@ -476,12 +476,12 @@ renumber_and_sort_sampled_edgelist(
* 1. If @p edgelist_hops is valid, we can consider (vertex ID, hop, flag=major) triplets for each
* vertex ID in edge majors (@p edgelist_srcs if @p src_is_major is true, @p edgelist_dsts if false)
* and (vertex ID, hop, flag=minor) triplets for each vertex ID in edge minors. From these triplets,
- * we can find the minimum (hop, flag) pairs for every unique vertex ID (hop is the primary key and
+ * we can find the minimum (hop, flag) pair for every unique vertex ID (hop is the primary key and
* flag is the secondary key, flag=major is considered smaller than flag=minor if hop numbers are
* same). Vertex IDs with smaller (hop, flag) pairs precede vertex IDs with larger (hop, flag) pairs
* in renumbering (if their vertex types are same, vertices with different types are renumbered
* separately). Ordering can be arbitrary among the vertices with the same (vertex type, hop, flag)
- * triplets. If @p seed_vertices.has-value() is true, we assume (hop=0, flag=major) for every vertex
+ * triplets. If @p seed_vertices.has_value() is true, we assume (hop=0, flag=major) for every vertex
* in @p *seed_vertices in renumbering (this is relevant when there are seed vertices with no
* neighbors).
* 2. If @p edgelist_hops is invalid, unique vertex IDs in edge majors precede vertex IDs that
@@ -495,11 +495,15 @@ renumber_and_sort_sampled_edgelist(
* Edge IDs are renumbered fulfilling the following requirements (This is relevant only when @p
* edgelist_edge_ids.has_value() is true).
*
- * 1. If @p edgelist_edge_types.has_value() is true, unique (edge type, edge ID) pairs are
- * renumbered to consecutive integers starting from 0 for each edge type. If @p
- * edgelist_edge_types.has_value() is true, unique edge IDs are renumbered to consecutive inetgers
- * starting from 0.
- * 2. If edgelist_label_offsets.has_value() is true, edge lists for different labels will be
+ * 1. If @p edgelist_hops is valid, we can consider (edge ID, hop) pairs. From these pairs, we can
+ * find the minimum hop value for every unique edge ID. Edge IDs with smaller hop values precede
+ * edge IDs with larger hop values in renumbering (if their edge types are same, edges with
+ * different edge types are renumbered separately). Ordering can be arbitrary among the edge IDs
+ * with the same (edge type, hop) pairs.
+ * 2. If @p edgelist_edge_hops.has_value() is false, unique edge IDs (for each edge type is @p
+ * edgelist_edge_types.has_value() is true) are mapped to consecutive integers starting from 0. The
+ * ordering can be arbitrary.
+ * 3. If edgelist_label_offsets.has_value() is true, edge lists for different labels will be
* renumbered separately.
*
* The renumbered edges are sorted based on the following rules.
@@ -510,6 +514,11 @@ renumber_and_sort_sampled_edgelist(
* true.
* 2. Edges in each label are sorted independently if @p edgelist_label_offsets.has_value() is true.
*
+ * This function assumes that there is a single edge source vertex type and a single edge
+ * destination vertex type for each edge. If @p edgelist_edge_types.has_value() is false (i.e. there
+ * is only one edge type), there should be only one edge source vertex type and only one edge
+ * destination vertex type; the source & destination vertex types may or may not coincide.
+ *
* This function is single-GPU only (we are not aware of any practical multi-GPU use cases).
*
* @tparam vertex_t Type of vertex identifiers. Needs to be an integral type.
@@ -530,19 +539,16 @@ renumber_and_sort_sampled_edgelist(
* edgelist_srcs.size() if valid).
* @param edgelist_hops An optional vector storing edge list hop numbers (size = @p
* edgelist_srcs.size() if valid). @p edgelist_hops should be valid if @p num_hops >= 2.
- * @param edgelist_label_offsets An optional pointer to the array storing label offsets to the input
- * edges (size = @p num_labels + 1). @p edgelist_label_offsets should be valid if @p num_labels
- * >= 2.
* @param seed_vertices An optional pointer to the array storing seed vertices in hop 0.
* @param seed_vertex_label_offsets An optional pointer to the array storing label offsets to the
* seed vertices (size = @p num_labels + 1). @p seed_vertex_label_offsets should be valid if @p
* num_labels >= 2 and @p seed_vertices is valid and invalid otherwise.
- * ext_vertices A pointer to the array storing external vertex IDs for the local internal vertices.
- * The local internal vertex range can be obatined bgy invoking a graph_view_t object's
- * local_vertex_partition_range() function. ext_vertex_type offsets A pointer to the array storing
- * vertex type offsets for the entire external vertex ID range (array size = @p num_vertex_types +
- * 1). For example, if the array stores [0, 100, 200], external vertex IDs [0, 100) has vertex type
- * 0 and external vertex IDs [100, 200) has vertex type 1.
+ * @param edgelist_label_offsets An optional pointer to the array storing label offsets to the input
+ * edges (size = @p num_labels + 1). @p edgelist_label_offsets should be valid if @p num_labels
+ * >= 2.
+ * @param vertex_type offsets A pointer to the array storing vertex type offsets for the entire
+ * vertex ID range (array size = @p num_vertex_types + 1). For example, if the array stores [0, 100,
+ * 200], vertex IDs [0, 100) has vertex type 0 and vertex IDs [100, 200) has vertex type 1.
* @param num_labels Number of labels. Labels are considered if @p num_labels >=2 and ignored if @p
* num_labels = 1.
* @param num_hops Number of hops. Hop numbers are considered if @p num_hops >=2 and ignored if @p
@@ -552,31 +558,36 @@ renumber_and_sort_sampled_edgelist(
* @param src_is_major A flag to determine whether to use the source or destination as the
* major key in renumbering and sorting.
* @param do_expensive_check A flag to run expensive checks for input arguments (if set to `true`).
- * @return Tuple of vectors storing edge sources, edge destinations, optional edge weights (valid
- * only if @p edgelist_weights.has_value() is true), optional edge IDs (valid only if @p
- * edgelist_edge_ids.has_value() is true), optional edge types (valid only if @p
- * edgelist_edge_types.has_value() is true), optional (label, hop) offset values to the renumbered
- * and sorted edges (size = @p num_labels * @p num_hops + 1, valid only when @p
- * edgelist_hops.has_value() or @p edgelist_label_offsetes.has_value() is true), renumber_map to
- * query original vertices (size = # unique or aggregate # unique vertices for each label), and
- * label offsets to the renumber map (size = @p num_labels + 1, valid only if @p
- * edgelist_label_offsets.has_value() is true).
+ * @return Tuple of vectors storing renumbered edge sources, renumbered edge destinations, optional
+ * edge weights (valid only if @p edgelist_weights.has_value() is true), optional renumbered edge
+ * IDs (valid only if @p edgelist_edge_ids.has_value() is true), optional (label, edge type, hop)
+ * offset values to the renumbered and sorted edges (size = @p num_labels * @p num_edge_types * @p
+ * num_hops + 1, valid only when @p edgelist_edge_types.has_value(), @p edgelist_hops.has_value(),
+ * or @p edgelist_label_offsetes.has_value() is true), renumber_map to query original vertices (size
+ * = # unique or aggregate # unique vertices for each label), (label, vertex type) offsets to the
+ * vertex renumber map (size = @p num_labels * @p num_vertex_types + 1), optional renumber_map to
+ * query original edge IDs (size = # unique (edge_type, edge ID) pairs, valid only if @p
+ * edgelist_edge_ids.has_value() is true), and optional (label, edge type) offsets to the edge ID
+ * renumber map (size = @p num_labels + @p num_edge_types + 1, valid only if @p
+ * edgelist_edge_ids.has_value() is true). We do not explicitly return edge source & destination
+ * vertex types as we assume that source & destination vertex type are implicilty determined for a
+ * given edge type.
*/
template
std::tuple<
- rmm::device_uvector, // srcs
- rmm::device_uvector, // dsts
- std::optional>, // weights
- std::optional>, // edge IDs
- std::optional>, // edge types
- std::optional>, // (label, edge type, hop) offsets to the edges
- rmm::device_uvector, // vertex renumber map
- std::optional>, // (label, type) offsets to the vertex renumber map
+ rmm::device_uvector, // srcs
+ rmm::device_uvector, // dsts
+ std::optional>, // weights
+ std::optional>, // edge IDs
+ std::optional>, // (label, edge type, hop) offsets to the edges
+ rmm::device_uvector, // vertex renumber map
+ rmm::device_uvector, // (label, vertex type) offsets to the vertex renumber map
std::optional>, // edge ID renumber map
- std::optional>> // (label, type) offsets to the edge ID renumber map
+ std::optional<
+ rmm::device_uvector>> // (label, edge type) offsets to the edge ID renumber map
heterogeneous_renumber_and_sort_sampled_edgelist(
raft::handle_t const& handle,
rmm::device_uvector&& edgelist_srcs,
@@ -585,11 +596,10 @@ heterogeneous_renumber_and_sort_sampled_edgelist(
std::optional>&& edgelist_edge_ids,
std::optional>&& edgelist_edge_types,
std::optional>&& edgelist_hops,
- std::optional> edgelist_label_offsets,
std::optional> seed_vertices,
std::optional> seed_vertex_label_offsets,
- raft::device_span ext_vertices,
- raft::device_span ext_vertex_type_offsets,
+ std::optional> edgelist_label_offsets,
+ raft::device_span vertex_type_offsets,
size_t num_labels,
size_t num_hops,
size_t num_vertex_types,
diff --git a/cpp/src/detail/utility_wrappers_32.cu b/cpp/src/detail/utility_wrappers_32.cu
index 72dee4a19a5..de407f12493 100644
--- a/cpp/src/detail/utility_wrappers_32.cu
+++ b/cpp/src/detail/utility_wrappers_32.cu
@@ -68,6 +68,23 @@ template void sequence_fill(rmm::cuda_stream_view const& stream_view,
size_t size,
int32_t start_value);
+template void sequence_fill(rmm::cuda_stream_view const& stream_view,
+ uint32_t* d_value,
+ size_t size,
+ uint32_t start_value);
+
+template void stride_fill(rmm::cuda_stream_view const& stream_view,
+ int32_t* d_value,
+ size_t size,
+ int32_t start_value,
+ int32_t stride);
+
+template void stride_fill(rmm::cuda_stream_view const& stream_view,
+ uint32_t* d_value,
+ size_t size,
+ uint32_t start_value,
+ uint32_t stride);
+
template int32_t compute_maximum_vertex_id(rmm::cuda_stream_view const& stream_view,
int32_t const* d_edgelist_srcs,
int32_t const* d_edgelist_dsts,
diff --git a/cpp/src/detail/utility_wrappers_64.cu b/cpp/src/detail/utility_wrappers_64.cu
index e7254d97c4d..2c136d5902b 100644
--- a/cpp/src/detail/utility_wrappers_64.cu
+++ b/cpp/src/detail/utility_wrappers_64.cu
@@ -71,6 +71,18 @@ template void sequence_fill(rmm::cuda_stream_view const& stream_view,
size_t size,
uint64_t start_value);
+template void stride_fill(rmm::cuda_stream_view const& stream_view,
+ int64_t* d_value,
+ size_t size,
+ int64_t start_value,
+ int64_t stride);
+
+template void stride_fill(rmm::cuda_stream_view const& stream_view,
+ uint64_t* d_value,
+ size_t size,
+ uint64_t start_value,
+ uint64_t stride);
+
template int64_t compute_maximum_vertex_id(rmm::cuda_stream_view const& stream_view,
int64_t const* d_edgelist_srcs,
int64_t const* d_edgelist_dsts,
diff --git a/cpp/src/detail/utility_wrappers_impl.cuh b/cpp/src/detail/utility_wrappers_impl.cuh
index ce8549db9f8..074d7044261 100644
--- a/cpp/src/detail/utility_wrappers_impl.cuh
+++ b/cpp/src/detail/utility_wrappers_impl.cuh
@@ -72,6 +72,22 @@ void sequence_fill(rmm::cuda_stream_view const& stream_view,
thrust::sequence(rmm::exec_policy(stream_view), d_value, d_value + size, start_value);
}
+template
+void stride_fill(rmm::cuda_stream_view const& stream_view,
+ value_t* d_value,
+ size_t size,
+ value_t start_value,
+ value_t stride)
+{
+ thrust::transform(rmm::exec_policy(stream_view),
+ thrust::make_counting_iterator(size_t{0}),
+ thrust::make_counting_iterator(size),
+ d_value,
+ cuda::proclaim_return_type([start_value, stride] __device__(size_t i) {
+ return static_cast(start_value + stride * i);
+ }));
+}
+
template
vertex_t compute_maximum_vertex_id(rmm::cuda_stream_view const& stream_view,
vertex_t const* d_edgelist_srcs,
diff --git a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh b/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh
deleted file mode 100644
index f5bc3ef6d2e..00000000000
--- a/cpp/src/sampling/renumber_sampled_edgelist_impl.cuh
+++ /dev/null
@@ -1,719 +0,0 @@
-/*
- * Copyright (c) 2023-2024, 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 "prims/kv_store.cuh"
-
-#include
-#include
-#include
-
-#include
-
-#include
-
-#include
-#include
-#include
-#include
-#include
-#include
-#include
-#include
-#include
-#include
-#include
-#include
-#include
-#include
-
-#include
-
-// FIXME: deprecated, to be deleted
-namespace cugraph {
-
-namespace {
-
-// output sorted by (primary key:label_index, secondary key:vertex)
-template
-std::tuple> /* label indices */,
- rmm::device_uvector /* vertices */,
- std::optional> /* minimum hops for the vertices */,
- std::optional> /* label offsets for the output */>
-compute_min_hop_for_unique_label_vertex_pairs(
- raft::handle_t const& handle,
- raft::device_span vertices,
- std::optional> hops,
- std::optional> label_indices,
- std::optional> label_offsets)
-{
- auto approx_edges_to_sort_per_iteration =
- static_cast(handle.get_device_properties().multiProcessorCount) *
- (1 << 20) /* tuning parameter */; // for segmented sort
-
- if (label_indices) {
- auto num_labels = (*label_offsets).size() - 1;
-
- rmm::device_uvector tmp_label_indices((*label_indices).size(),
- handle.get_stream());
- thrust::copy(handle.get_thrust_policy(),
- (*label_indices).begin(),
- (*label_indices).end(),
- tmp_label_indices.begin());
-
- rmm::device_uvector tmp_vertices(0, handle.get_stream());
- std::optional> tmp_hops{std::nullopt};
-
- if (hops) {
- tmp_vertices.resize(vertices.size(), handle.get_stream());
- thrust::copy(
- handle.get_thrust_policy(), vertices.begin(), vertices.end(), tmp_vertices.begin());
- tmp_hops = rmm::device_uvector((*hops).size(), handle.get_stream());
- thrust::copy(handle.get_thrust_policy(), (*hops).begin(), (*hops).end(), (*tmp_hops).begin());
-
- auto triplet_first = thrust::make_zip_iterator(
- tmp_label_indices.begin(), tmp_vertices.begin(), (*tmp_hops).begin());
- thrust::sort(
- handle.get_thrust_policy(), triplet_first, triplet_first + tmp_label_indices.size());
- auto key_first = thrust::make_zip_iterator(tmp_label_indices.begin(), tmp_vertices.begin());
- auto num_uniques = static_cast(
- thrust::distance(key_first,
- thrust::get<0>(thrust::unique_by_key(handle.get_thrust_policy(),
- key_first,
- key_first + tmp_label_indices.size(),
- (*tmp_hops).begin()))));
- tmp_label_indices.resize(num_uniques, handle.get_stream());
- tmp_vertices.resize(num_uniques, handle.get_stream());
- (*tmp_hops).resize(num_uniques, handle.get_stream());
- tmp_label_indices.shrink_to_fit(handle.get_stream());
- tmp_vertices.shrink_to_fit(handle.get_stream());
- (*tmp_hops).shrink_to_fit(handle.get_stream());
- } else {
- rmm::device_uvector segment_sorted_vertices(vertices.size(), handle.get_stream());
-
- rmm::device_uvector d_tmp_storage(0, handle.get_stream());
-
- auto [h_label_offsets, h_edge_offsets] = detail::compute_offset_aligned_element_chunks(
- handle, *label_offsets, vertices.size(), approx_edges_to_sort_per_iteration);
- auto num_chunks = h_label_offsets.size() - 1;
-
- for (size_t i = 0; i < num_chunks; ++i) {
- size_t tmp_storage_bytes{0};
-
- auto offset_first =
- thrust::make_transform_iterator((*label_offsets).data() + h_label_offsets[i],
- detail::shift_left_t{h_edge_offsets[i]});
- cub::DeviceSegmentedSort::SortKeys(static_cast(nullptr),
- tmp_storage_bytes,
- vertices.begin() + h_edge_offsets[i],
- segment_sorted_vertices.begin() + h_edge_offsets[i],
- h_edge_offsets[i + 1] - h_edge_offsets[i],
- h_label_offsets[i + 1] - h_label_offsets[i],
- offset_first,
- offset_first + 1,
- handle.get_stream());
-
- if (tmp_storage_bytes > d_tmp_storage.size()) {
- d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream());
- }
-
- cub::DeviceSegmentedSort::SortKeys(d_tmp_storage.data(),
- tmp_storage_bytes,
- vertices.begin() + h_edge_offsets[i],
- segment_sorted_vertices.begin() + h_edge_offsets[i],
- h_edge_offsets[i + 1] - h_edge_offsets[i],
- h_label_offsets[i + 1] - h_label_offsets[i],
- offset_first,
- offset_first + 1,
- handle.get_stream());
- }
- d_tmp_storage.resize(0, handle.get_stream());
- d_tmp_storage.shrink_to_fit(handle.get_stream());
-
- auto pair_first =
- thrust::make_zip_iterator(tmp_label_indices.begin(), segment_sorted_vertices.begin());
- auto num_uniques = static_cast(thrust::distance(
- pair_first,
- thrust::unique(
- handle.get_thrust_policy(), pair_first, pair_first + tmp_label_indices.size())));
- tmp_label_indices.resize(num_uniques, handle.get_stream());
- segment_sorted_vertices.resize(num_uniques, handle.get_stream());
- tmp_label_indices.shrink_to_fit(handle.get_stream());
- segment_sorted_vertices.shrink_to_fit(handle.get_stream());
-
- tmp_vertices = std::move(segment_sorted_vertices);
- }
-
- rmm::device_uvector tmp_label_offsets(num_labels + 1, handle.get_stream());
- tmp_label_offsets.set_element_to_zero_async(0, handle.get_stream());
- thrust::upper_bound(handle.get_thrust_policy(),
- tmp_label_indices.begin(),
- tmp_label_indices.end(),
- thrust::make_counting_iterator(size_t{0}),
- thrust::make_counting_iterator(num_labels),
- tmp_label_offsets.begin() + 1);
-
- return std::make_tuple(std::move(tmp_label_indices),
- std::move(tmp_vertices),
- std::move(tmp_hops),
- std::move(tmp_label_offsets));
- } else {
- rmm::device_uvector tmp_vertices(vertices.size(), handle.get_stream());
- thrust::copy(
- handle.get_thrust_policy(), vertices.begin(), vertices.end(), tmp_vertices.begin());
-
- if (hops) {
- rmm::device_uvector tmp_hops((*hops).size(), handle.get_stream());
- thrust::copy(handle.get_thrust_policy(), (*hops).begin(), (*hops).end(), tmp_hops.begin());
-
- auto pair_first = thrust::make_zip_iterator(
- tmp_vertices.begin(), tmp_hops.begin()); // vertex is a primary key, hop is a secondary key
- thrust::sort(handle.get_thrust_policy(), pair_first, pair_first + tmp_vertices.size());
- tmp_vertices.resize(
- thrust::distance(tmp_vertices.begin(),
- thrust::get<0>(thrust::unique_by_key(handle.get_thrust_policy(),
- tmp_vertices.begin(),
- tmp_vertices.end(),
- tmp_hops.begin()))),
- handle.get_stream());
- tmp_hops.resize(tmp_vertices.size(), handle.get_stream());
-
- return std::make_tuple(
- std::nullopt, std::move(tmp_vertices), std::move(tmp_hops), std::nullopt);
- } else {
- thrust::sort(handle.get_thrust_policy(), tmp_vertices.begin(), tmp_vertices.end());
- tmp_vertices.resize(
- thrust::distance(
- tmp_vertices.begin(),
- thrust::unique(handle.get_thrust_policy(), tmp_vertices.begin(), tmp_vertices.end())),
- handle.get_stream());
- tmp_vertices.shrink_to_fit(handle.get_stream());
-
- return std::make_tuple(std::nullopt, std::move(tmp_vertices), std::nullopt, std::nullopt);
- }
- }
-}
-
-template
-std::tuple, std::optional>>
-compute_renumber_map(raft::handle_t const& handle,
- raft::device_span edgelist_srcs,
- raft::device_span edgelist_dsts,
- std::optional> edgelist_hops,
- std::optional> label_offsets)
-{
- auto approx_edges_to_sort_per_iteration =
- static_cast(handle.get_device_properties().multiProcessorCount) *
- (1 << 20) /* tuning parameter */; // for segmented sort
-
- std::optional> edgelist_label_indices{std::nullopt};
- if (label_offsets) {
- edgelist_label_indices =
- detail::expand_sparse_offsets(*label_offsets, label_index_t{0}, handle.get_stream());
- }
-
- auto [unique_label_src_pair_label_indices,
- unique_label_src_pair_vertices,
- unique_label_src_pair_hops,
- unique_label_src_pair_label_offsets] =
- compute_min_hop_for_unique_label_vertex_pairs(
- handle,
- edgelist_srcs,
- edgelist_hops,
- edgelist_label_indices ? std::make_optional>(
- (*edgelist_label_indices).data(), (*edgelist_label_indices).size())
- : std::nullopt,
- label_offsets);
-
- auto [unique_label_dst_pair_label_indices,
- unique_label_dst_pair_vertices,
- unique_label_dst_pair_hops,
- unique_label_dst_pair_label_offsets] =
- compute_min_hop_for_unique_label_vertex_pairs(
- handle,
- edgelist_dsts,
- edgelist_hops,
- edgelist_label_indices ? std::make_optional>(
- (*edgelist_label_indices).data(), (*edgelist_label_indices).size())
- : std::nullopt,
- label_offsets);
-
- edgelist_label_indices = std::nullopt;
-
- if (label_offsets) {
- auto num_labels = (*label_offsets).size() - 1;
-
- rmm::device_uvector renumber_map(0, handle.get_stream());
- rmm::device_uvector renumber_map_label_indices(0, handle.get_stream());
-
- renumber_map.reserve(
- (*unique_label_src_pair_label_indices).size() + (*unique_label_dst_pair_label_indices).size(),
- handle.get_stream());
- renumber_map_label_indices.reserve(renumber_map.capacity(), handle.get_stream());
-
- auto num_chunks = (edgelist_srcs.size() + (approx_edges_to_sort_per_iteration - 1)) /
- approx_edges_to_sort_per_iteration;
- auto chunk_size = (num_chunks > 0) ? ((num_labels + (num_chunks - 1)) / num_chunks) : 0;
-
- size_t copy_offset{0};
- for (size_t i = 0; i < num_chunks; ++i) {
- auto src_start_offset =
- (*unique_label_src_pair_label_offsets).element(chunk_size * i, handle.get_stream());
- auto src_end_offset =
- (*unique_label_src_pair_label_offsets)
- .element(std::min(chunk_size * (i + 1), num_labels), handle.get_stream());
- auto dst_start_offset =
- (*unique_label_dst_pair_label_offsets).element(chunk_size * i, handle.get_stream());
- auto dst_end_offset =
- (*unique_label_dst_pair_label_offsets)
- .element(std::min(chunk_size * (i + 1), num_labels), handle.get_stream());
-
- rmm::device_uvector merged_label_indices(
- (src_end_offset - src_start_offset) + (dst_end_offset - dst_start_offset),
- handle.get_stream());
- rmm::device_uvector merged_vertices(merged_label_indices.size(),
- handle.get_stream());
- rmm::device_uvector merged_flags(merged_label_indices.size(), handle.get_stream());
-
- if (edgelist_hops) {
- rmm::device_uvector merged_hops(merged_label_indices.size(), handle.get_stream());
- auto src_quad_first =
- thrust::make_zip_iterator((*unique_label_src_pair_label_indices).begin(),
- unique_label_src_pair_vertices.begin(),
- (*unique_label_src_pair_hops).begin(),
- thrust::make_constant_iterator(int8_t{0}));
- auto dst_quad_first =
- thrust::make_zip_iterator((*unique_label_dst_pair_label_indices).begin(),
- unique_label_dst_pair_vertices.begin(),
- (*unique_label_dst_pair_hops).begin(),
- thrust::make_constant_iterator(int8_t{1}));
- thrust::merge(handle.get_thrust_policy(),
- src_quad_first + src_start_offset,
- src_quad_first + src_end_offset,
- dst_quad_first + dst_start_offset,
- dst_quad_first + dst_end_offset,
- thrust::make_zip_iterator(merged_label_indices.begin(),
- merged_vertices.begin(),
- merged_hops.begin(),
- merged_flags.begin()));
-
- auto unique_key_first =
- thrust::make_zip_iterator(merged_label_indices.begin(), merged_vertices.begin());
- merged_label_indices.resize(
- thrust::distance(
- unique_key_first,
- thrust::get<0>(thrust::unique_by_key(
- handle.get_thrust_policy(),
- unique_key_first,
- unique_key_first + merged_label_indices.size(),
- thrust::make_zip_iterator(merged_hops.begin(), merged_flags.begin())))),
- handle.get_stream());
- merged_vertices.resize(merged_label_indices.size(), handle.get_stream());
- merged_hops.resize(merged_label_indices.size(), handle.get_stream());
- merged_flags.resize(merged_label_indices.size(), handle.get_stream());
- auto sort_key_first = thrust::make_zip_iterator(
- merged_label_indices.begin(), merged_hops.begin(), merged_flags.begin());
- thrust::sort_by_key(handle.get_thrust_policy(),
- sort_key_first,
- sort_key_first + merged_label_indices.size(),
- merged_vertices.begin());
- } else {
- auto src_triplet_first =
- thrust::make_zip_iterator((*unique_label_src_pair_label_indices).begin(),
- unique_label_src_pair_vertices.begin(),
- thrust::make_constant_iterator(int8_t{0}));
- auto dst_triplet_first =
- thrust::make_zip_iterator((*unique_label_dst_pair_label_indices).begin(),
- unique_label_dst_pair_vertices.begin(),
- thrust::make_constant_iterator(int8_t{1}));
- thrust::merge(
- handle.get_thrust_policy(),
- src_triplet_first + src_start_offset,
- src_triplet_first + src_end_offset,
- dst_triplet_first + dst_start_offset,
- dst_triplet_first + dst_end_offset,
- thrust::make_zip_iterator(
- merged_label_indices.begin(), merged_vertices.begin(), merged_flags.begin()));
-
- auto unique_key_first =
- thrust::make_zip_iterator(merged_label_indices.begin(), merged_vertices.begin());
- merged_label_indices.resize(
- thrust::distance(
- unique_key_first,
- thrust::get<0>(thrust::unique_by_key(handle.get_thrust_policy(),
- unique_key_first,
- unique_key_first + merged_label_indices.size(),
- merged_flags.begin()))),
- handle.get_stream());
- merged_vertices.resize(merged_label_indices.size(), handle.get_stream());
- merged_flags.resize(merged_label_indices.size(), handle.get_stream());
- auto sort_key_first =
- thrust::make_zip_iterator(merged_label_indices.begin(), merged_flags.begin());
- thrust::sort_by_key(handle.get_thrust_policy(),
- sort_key_first,
- sort_key_first + merged_label_indices.size(),
- merged_vertices.begin());
- }
-
- renumber_map.resize(copy_offset + merged_vertices.size(), handle.get_stream());
- thrust::copy(handle.get_thrust_policy(),
- merged_vertices.begin(),
- merged_vertices.end(),
- renumber_map.begin() + copy_offset);
- renumber_map_label_indices.resize(copy_offset + merged_label_indices.size(),
- handle.get_stream());
- thrust::copy(handle.get_thrust_policy(),
- merged_label_indices.begin(),
- merged_label_indices.end(),
- renumber_map_label_indices.begin() + copy_offset);
-
- copy_offset += merged_vertices.size();
- }
-
- renumber_map.shrink_to_fit(handle.get_stream());
- renumber_map_label_indices.shrink_to_fit(handle.get_stream());
-
- return std::make_tuple(std::move(renumber_map), std::move(renumber_map_label_indices));
- } else {
- if (edgelist_hops) {
- rmm::device_uvector merged_vertices(
- unique_label_src_pair_vertices.size() + unique_label_dst_pair_vertices.size(),
- handle.get_stream());
- rmm::device_uvector merged_hops(merged_vertices.size(), handle.get_stream());
- rmm::device_uvector merged_flags(merged_vertices.size(), handle.get_stream());
- auto src_triplet_first = thrust::make_zip_iterator(unique_label_src_pair_vertices.begin(),
- (*unique_label_src_pair_hops).begin(),
- thrust::make_constant_iterator(int8_t{0}));
- auto dst_triplet_first = thrust::make_zip_iterator(unique_label_dst_pair_vertices.begin(),
- (*unique_label_dst_pair_hops).begin(),
- thrust::make_constant_iterator(int8_t{1}));
- thrust::merge(handle.get_thrust_policy(),
- src_triplet_first,
- src_triplet_first + unique_label_src_pair_vertices.size(),
- dst_triplet_first,
- dst_triplet_first + unique_label_dst_pair_vertices.size(),
- thrust::make_zip_iterator(
- merged_vertices.begin(), merged_hops.begin(), merged_flags.begin()));
-
- unique_label_src_pair_vertices.resize(0, handle.get_stream());
- unique_label_src_pair_vertices.shrink_to_fit(handle.get_stream());
- unique_label_src_pair_hops = std::nullopt;
- unique_label_dst_pair_vertices.resize(0, handle.get_stream());
- unique_label_dst_pair_vertices.shrink_to_fit(handle.get_stream());
- unique_label_dst_pair_hops = std::nullopt;
-
- merged_vertices.resize(
- thrust::distance(merged_vertices.begin(),
- thrust::get<0>(thrust::unique_by_key(
- handle.get_thrust_policy(),
- merged_vertices.begin(),
- merged_vertices.end(),
- thrust::make_zip_iterator(merged_hops.begin(), merged_flags.begin())))),
- handle.get_stream());
- merged_hops.resize(merged_vertices.size(), handle.get_stream());
- merged_flags.resize(merged_vertices.size(), handle.get_stream());
-
- auto sort_key_first = thrust::make_zip_iterator(merged_hops.begin(), merged_flags.begin());
- thrust::sort_by_key(handle.get_thrust_policy(),
- sort_key_first,
- sort_key_first + merged_hops.size(),
- merged_vertices.begin());
-
- return std::make_tuple(std::move(merged_vertices), std::nullopt);
- } else {
- rmm::device_uvector output_vertices(unique_label_dst_pair_vertices.size(),
- handle.get_stream());
- auto output_last = thrust::set_difference(handle.get_thrust_policy(),
- unique_label_dst_pair_vertices.begin(),
- unique_label_dst_pair_vertices.end(),
- unique_label_src_pair_vertices.begin(),
- unique_label_src_pair_vertices.end(),
- output_vertices.begin());
-
- auto num_unique_srcs = unique_label_src_pair_vertices.size();
- auto renumber_map = std::move(unique_label_src_pair_vertices);
- renumber_map.resize(
- renumber_map.size() + thrust::distance(output_vertices.begin(), output_last),
- handle.get_stream());
- thrust::copy(handle.get_thrust_policy(),
- output_vertices.begin(),
- output_last,
- renumber_map.begin() + num_unique_srcs);
-
- return std::make_tuple(std::move(renumber_map), std::nullopt);
- }
- }
-}
-
-} // namespace
-
-template
-std::tuple,
- rmm::device_uvector,
- rmm::device_uvector,
- std::optional>>
-renumber_sampled_edgelist(
- raft::handle_t const& handle,
- rmm::device_uvector&& edgelist_srcs,
- rmm::device_uvector&& edgelist_dsts,
- std::optional> edgelist_hops,
- std::optional, raft::device_span>>
- label_offsets,
- bool do_expensive_check)
-{
- using label_index_t = uint32_t;
-
- // 1. check input arguments
-
- CUGRAPH_EXPECTS(!label_offsets || (std::get<0>(*label_offsets).size() <=
- std::numeric_limits::max()),
- "Invalid input arguments: current implementation assumes that the number of "
- "unique labels is no larger than std::numeric_limits::max().");
-
- CUGRAPH_EXPECTS(
- edgelist_srcs.size() == edgelist_dsts.size(),
- "Invalid input arguments: edgelist_srcs.size() and edgelist_dsts.size() should coincide.");
- CUGRAPH_EXPECTS(!edgelist_hops.has_value() || (edgelist_srcs.size() == (*edgelist_hops).size()),
- "Invalid input arguments: if edgelist_hops is valid, (*edgelist_hops).size() and "
- "edgelist_srcs.size() should coincide.");
- CUGRAPH_EXPECTS(!label_offsets.has_value() ||
- (std::get<1>(*label_offsets).size() == std::get<0>(*label_offsets).size() + 1),
- "Invalid input arguments: if label_offsets is valid, "
- "std::get<1>(label_offsets).size() (size of the offset array) should be "
- "std::get<0>(label_offsets).size() (number of unique labels) + 1.");
-
- if (do_expensive_check) {
- if (label_offsets) {
- CUGRAPH_EXPECTS(thrust::is_sorted(handle.get_thrust_policy(),
- std::get<1>(*label_offsets).begin(),
- std::get<1>(*label_offsets).end()),
- "Invalid input arguments: if label_offsets is valid, "
- "std::get<1>(*label_offsets) should be sorted.");
- size_t back_element{};
- raft::update_host(
- &back_element,
- std::get<1>(*label_offsets).data() + (std::get<1>(*label_offsets).size() - 1),
- size_t{1},
- handle.get_stream());
- handle.get_stream();
- CUGRAPH_EXPECTS(back_element == edgelist_srcs.size(),
- "Invalid input arguments: if label_offsets is valid, the last element of "
- "std::get<1>(*label_offsets) and edgelist_srcs.size() should coincide.");
- }
- }
-
- // 2. compute renumber_map
-
- auto [renumber_map, renumber_map_label_indices] = compute_renumber_map(
- handle,
- raft::device_span(edgelist_srcs.data(), edgelist_srcs.size()),
- raft::device_span(edgelist_dsts.data(), edgelist_dsts.size()),
- edgelist_hops,
- label_offsets ? std::make_optional>(std::get<1>(*label_offsets))
- : std::nullopt);
-
- // 3. compute renumber map offsets for each label
-
- std::optional> renumber_map_label_offsets{};
- if (label_offsets) {
- auto num_unique_labels = thrust::count_if(
- handle.get_thrust_policy(),
- thrust::make_counting_iterator(size_t{0}),
- thrust::make_counting_iterator((*renumber_map_label_indices).size()),
- detail::is_first_in_run_t{(*renumber_map_label_indices).data()});
- rmm::device_uvector unique_label_indices(num_unique_labels, handle.get_stream());
- rmm::device_uvector vertex_counts(num_unique_labels, handle.get_stream());
- thrust::reduce_by_key(handle.get_thrust_policy(),
- (*renumber_map_label_indices).begin(),
- (*renumber_map_label_indices).end(),
- thrust::make_constant_iterator(size_t{1}),
- unique_label_indices.begin(),
- vertex_counts.begin());
-
- renumber_map_label_offsets =
- rmm::device_uvector(std::get<0>(*label_offsets).size() + 1, handle.get_stream());
- thrust::fill(handle.get_thrust_policy(),
- (*renumber_map_label_offsets).begin(),
- (*renumber_map_label_offsets).end(),
- size_t{0});
- thrust::scatter(handle.get_thrust_policy(),
- vertex_counts.begin(),
- vertex_counts.end(),
- unique_label_indices.begin(),
- (*renumber_map_label_offsets).begin() + 1);
-
- thrust::inclusive_scan(handle.get_thrust_policy(),
- (*renumber_map_label_offsets).begin(),
- (*renumber_map_label_offsets).end(),
- (*renumber_map_label_offsets).begin());
- }
-
- // 4. renumber input edges
-
- if (label_offsets) {
- rmm::device_uvector new_vertices(renumber_map.size(), handle.get_stream());
- thrust::tabulate(handle.get_thrust_policy(),
- new_vertices.begin(),
- new_vertices.end(),
- [label_indices = raft::device_span(
- (*renumber_map_label_indices).data(), (*renumber_map_label_indices).size()),
- renumber_map_label_offsets = raft::device_span(
- (*renumber_map_label_offsets).data(),
- (*renumber_map_label_offsets).size())] __device__(size_t i) {
- auto label_index = label_indices[i];
- auto label_start_offset = renumber_map_label_offsets[label_index];
- return static_cast(i - label_start_offset);
- });
-
- (*renumber_map_label_indices).resize(0, handle.get_stream());
- (*renumber_map_label_indices).shrink_to_fit(handle.get_stream());
-
- auto num_labels = std::get<0>(*label_offsets).size();
-
- rmm::device_uvector segment_sorted_renumber_map(renumber_map.size(),
- handle.get_stream());
- rmm::device_uvector segment_sorted_new_vertices(new_vertices.size(),
- handle.get_stream());
-
- rmm::device_uvector d_tmp_storage(0, handle.get_stream());
-
- auto approx_edges_to_sort_per_iteration =
- static_cast(handle.get_device_properties().multiProcessorCount) *
- (1 << 20) /* tuning parameter */; // for segmented sort
-
- auto [h_label_offsets, h_edge_offsets] = detail::compute_offset_aligned_element_chunks(
- handle,
- raft::device_span{(*renumber_map_label_offsets).data(),
- (*renumber_map_label_offsets).size()},
- renumber_map.size(),
- approx_edges_to_sort_per_iteration);
- auto num_chunks = h_label_offsets.size() - 1;
-
- for (size_t i = 0; i < num_chunks; ++i) {
- size_t tmp_storage_bytes{0};
-
- auto offset_first =
- thrust::make_transform_iterator((*renumber_map_label_offsets).data() + h_label_offsets[i],
- detail::shift_left_t{h_edge_offsets[i]});
- cub::DeviceSegmentedSort::SortPairs(static_cast(nullptr),
- tmp_storage_bytes,
- renumber_map.begin() + h_edge_offsets[i],
- segment_sorted_renumber_map.begin() + h_edge_offsets[i],
- new_vertices.begin() + h_edge_offsets[i],
- segment_sorted_new_vertices.begin() + h_edge_offsets[i],
- h_edge_offsets[i + 1] - h_edge_offsets[i],
- h_label_offsets[i + 1] - h_label_offsets[i],
- offset_first,
- offset_first + 1,
- handle.get_stream());
-
- if (tmp_storage_bytes > d_tmp_storage.size()) {
- d_tmp_storage = rmm::device_uvector(tmp_storage_bytes, handle.get_stream());
- }
-
- cub::DeviceSegmentedSort::SortPairs(d_tmp_storage.data(),
- tmp_storage_bytes,
- renumber_map.begin() + h_edge_offsets[i],
- segment_sorted_renumber_map.begin() + h_edge_offsets[i],
- new_vertices.begin() + h_edge_offsets[i],
- segment_sorted_new_vertices.begin() + h_edge_offsets[i],
- h_edge_offsets[i + 1] - h_edge_offsets[i],
- h_label_offsets[i + 1] - h_label_offsets[i],
- offset_first,
- offset_first + 1,
- handle.get_stream());
- }
- new_vertices.resize(0, handle.get_stream());
- d_tmp_storage.resize(0, handle.get_stream());
- new_vertices.shrink_to_fit(handle.get_stream());
- d_tmp_storage.shrink_to_fit(handle.get_stream());
-
- auto edgelist_label_indices = detail::expand_sparse_offsets(
- std::get<1>(*label_offsets), label_index_t{0}, handle.get_stream());
-
- auto pair_first =
- thrust::make_zip_iterator(edgelist_srcs.begin(), edgelist_label_indices.begin());
- thrust::transform(
- handle.get_thrust_policy(),
- pair_first,
- pair_first + edgelist_srcs.size(),
- edgelist_srcs.begin(),
- [renumber_map_label_offsets = raft::device_span(
- (*renumber_map_label_offsets).data(), (*renumber_map_label_offsets).size()),
- old_vertices = raft::device_span(segment_sorted_renumber_map.data(),
- segment_sorted_renumber_map.size()),
- new_vertices = raft::device_span(
- segment_sorted_new_vertices.data(),
- segment_sorted_new_vertices.size())] __device__(auto pair) {
- auto old_vertex = thrust::get<0>(pair);
- auto label_index = thrust::get<1>(pair);
- auto label_start_offset = renumber_map_label_offsets[label_index];
- auto label_end_offset = renumber_map_label_offsets[label_index + 1];
- auto it = thrust::lower_bound(thrust::seq,
- old_vertices.begin() + label_start_offset,
- old_vertices.begin() + label_end_offset,
- old_vertex);
- assert(*it == old_vertex);
- return *(new_vertices.begin() + thrust::distance(old_vertices.begin(), it));
- });
-
- pair_first = thrust::make_zip_iterator(edgelist_dsts.begin(), edgelist_label_indices.begin());
- thrust::transform(
- handle.get_thrust_policy(),
- pair_first,
- pair_first + edgelist_dsts.size(),
- edgelist_dsts.begin(),
- [renumber_map_label_offsets = raft::device_span(
- (*renumber_map_label_offsets).data(), (*renumber_map_label_offsets).size()),
- old_vertices = raft::device_span(segment_sorted_renumber_map.data(),
- segment_sorted_renumber_map.size()),
- new_vertices = raft::device_span(
- segment_sorted_new_vertices.data(),
- segment_sorted_new_vertices.size())] __device__(auto pair) {
- auto old_vertex = thrust::get<0>(pair);
- auto label_index = thrust::get<1>(pair);
- auto label_start_offset = renumber_map_label_offsets[label_index];
- auto label_end_offset = renumber_map_label_offsets[label_index + 1];
- auto it = thrust::lower_bound(thrust::seq,
- old_vertices.begin() + label_start_offset,
- old_vertices.begin() + label_end_offset,
- old_vertex);
- assert(*it == old_vertex);
- return new_vertices[thrust::distance(old_vertices.begin(), it)];
- });
- } else {
- kv_store_t kv_store(renumber_map.begin(),
- renumber_map.end(),
- thrust::make_counting_iterator(vertex_t{0}),
- std::numeric_limits::max(),
- std::numeric_limits::max(),
- handle.get_stream());
- auto kv_store_view = kv_store.view();
-
- kv_store_view.find(
- edgelist_srcs.begin(), edgelist_srcs.end(), edgelist_srcs.begin(), handle.get_stream());
- kv_store_view.find(
- edgelist_dsts.begin(), edgelist_dsts.end(), edgelist_dsts.begin(), handle.get_stream());
- }
-
- return std::make_tuple(std::move(edgelist_srcs),
- std::move(edgelist_dsts),
- std::move(renumber_map),
- std::move(renumber_map_label_offsets));
-}
-
-} // namespace cugraph
diff --git a/cpp/src/sampling/renumber_sampled_edgelist_sg_v32_e32.cu b/cpp/src/sampling/renumber_sampled_edgelist_sg_v32_e32.cu
deleted file mode 100644
index dee28c593ad..00000000000
--- a/cpp/src/sampling/renumber_sampled_edgelist_sg_v32_e32.cu
+++ /dev/null
@@ -1,37 +0,0 @@
-/*
- * Copyright (c) 2023-2024, 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 "renumber_sampled_edgelist_impl.cuh"
-
-#include
-
-// FIXME: deprecated, to be deleted
-namespace cugraph {
-
-template std::tuple,
- rmm::device_uvector,
- rmm::device_uvector,
- std::optional>>
-renumber_sampled_edgelist(
- raft::handle_t const& handle,
- rmm::device_uvector&& edgelist_srcs,
- rmm::device_uvector&& edgelist_dsts,
- std::optional> edgelist_hops,
- std::optional, raft::device_span>>
- label_offsets,
- bool do_expensive_check);
-
-} // namespace cugraph
diff --git a/cpp/src/sampling/renumber_sampled_edgelist_sg_v64_e64.cu b/cpp/src/sampling/renumber_sampled_edgelist_sg_v64_e64.cu
deleted file mode 100644
index 99293c68f0c..00000000000
--- a/cpp/src/sampling/renumber_sampled_edgelist_sg_v64_e64.cu
+++ /dev/null
@@ -1,37 +0,0 @@
-/*
- * Copyright (c) 2023-2024, 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 "renumber_sampled_edgelist_impl.cuh"
-
-#include
-
-// FIXME: deprecated, to be deleted
-namespace cugraph {
-
-template std::tuple,
- rmm::device_uvector,
- rmm::device_uvector,
- std::optional>>
-renumber_sampled_edgelist(
- raft::handle_t const& handle,
- rmm::device_uvector&& edgelist_srcs,
- rmm::device_uvector&& edgelist_dsts,
- std::optional> edgelist_hops,
- std::optional, raft::device_span>>
- label_offsets,
- bool do_expensive_check);
-
-} // namespace cugraph
diff --git a/cpp/src/sampling/sampling_post_processing_impl.cuh b/cpp/src/sampling/sampling_post_processing_impl.cuh
index b0b3bb5f4f2..4624e6d4a5e 100644
--- a/cpp/src/sampling/sampling_post_processing_impl.cuh
+++ b/cpp/src/sampling/sampling_post_processing_impl.cuh
@@ -49,9 +49,10 @@ namespace cugraph {
namespace {
-template
+template
struct edge_order_t {
thrust::optional> edgelist_label_offsets{thrust::nullopt};
+ thrust::optional> edgelist_edge_types{thrust::nullopt};
thrust::optional> edgelist_hops{thrust::nullopt};
raft::device_span edgelist_majors{};
raft::device_span edgelist_minors{};
@@ -72,6 +73,12 @@ struct edge_order_t {
if (l_label != r_label) { return l_label < r_label; }
}
+ if (edgelist_edge_types) {
+ auto l_type = (*edgelist_edge_types)[l_idx];
+ auto r_type = (*edgelist_edge_types)[r_idx];
+ if (l_type != r_type) { return l_type < r_type; }
+ }
+
if (edgelist_hops) {
auto l_hop = (*edgelist_hops)[l_idx];
auto r_hop = (*edgelist_hops)[r_idx];
@@ -151,6 +158,7 @@ struct optionally_compute_label_index_t {
template
@@ -164,8 +172,11 @@ void check_input_edges(raft::handle_t const& handle,
std::optional> seed_vertices,
std::optional> seed_vertex_label_offsets,
std::optional> edgelist_label_offsets,
+ std::optional> vertex_type_offsets,
size_t num_labels,
size_t num_hops,
+ size_t num_vertex_types,
+ std::optional num_edge_types,
bool do_expensive_check)
{
CUGRAPH_EXPECTS(
@@ -193,6 +204,7 @@ void check_input_edges(raft::handle_t const& handle,
"(size of the offset array) should be num_labels + 1.");
if (edgelist_majors.size() > 0) {
+ static_assert(std::is_same_v);
CUGRAPH_EXPECTS((num_labels >= 1) && (num_labels <= std::numeric_limits::max()),
"Invalid input arguments: num_labels should be a positive integer and the "
"current implementation assumes that the number of unique labels is no larger "
@@ -209,13 +221,16 @@ void check_input_edges(raft::handle_t const& handle,
CUGRAPH_EXPECTS(
(num_hops == 1) || edgelist_hops.has_value(),
"Invalid input arguments: edgelist_hops.has_value() should be true if num_hops >= 2.");
- } else {
- CUGRAPH_EXPECTS(
- "num_labels == 0",
- "Invalid input arguments: num_labels should be 0 if the input edge list is empty.");
+
+ static_assert(std::is_same_v);
CUGRAPH_EXPECTS(
- "num_hops == 0",
- "Invalid input arguments: num_hops should be 0 if the input edge list is empty.");
+ (num_vertex_types >= 1) && (num_vertex_types <= std::numeric_limits::max()),
+ "Invalid input arguments: num_vertex_types should be a positive integer and the "
+ "current implementation assumes that the number of vertex types is no larger "
+ "than std::numeric_limits::max().");
+ CUGRAPH_EXPECTS((num_vertex_types == 1) || vertex_type_offsets.has_value(),
+ "Invalid input arguments: vertex_type_offsets.has_value() should be true if "
+ "num_vertex_types >= 2.");
}
CUGRAPH_EXPECTS((!seed_vertices.has_value() && !seed_vertex_label_offsets.has_value()) ||
@@ -257,6 +272,174 @@ void check_input_edges(raft::handle_t const& handle,
"*edgelist_label_offsets and edgelist_(srcs|dsts).size() should coincide.");
}
+ if (edgelist_edge_types && num_edge_types) {
+ CUGRAPH_EXPECTS(
+ thrust::count_if(handle.get_thrust_policy(),
+ (*edgelist_edge_types).begin(),
+ (*edgelist_edge_types).end(),
+ [num_edge_types = static_cast(*num_edge_types)] __device__(
+ edge_type_t edge_type) { return edge_type >= num_edge_types; }) == 0,
+ "Invalid input arguments: edgelist_edge_type is valid but contains out-of-range edge type "
+ "values.");
+ if constexpr (std::is_signed_v) {
+ CUGRAPH_EXPECTS(thrust::count_if(handle.get_thrust_policy(),
+ (*edgelist_edge_types).begin(),
+ (*edgelist_edge_types).end(),
+ [] __device__(edge_type_t edge_type) {
+ return edge_type < edge_type_t{0};
+ }) == 0,
+ "Invalid input arguments: edgelist_edge_type is valid but contains "
+ "negative edge type values.");
+ }
+ }
+
+ if (vertex_type_offsets) {
+ CUGRAPH_EXPECTS(
+ thrust::is_sorted(
+ handle.get_thrust_policy(), (*vertex_type_offsets).begin(), (*vertex_type_offsets).end()),
+ "Invalid input arguments: if vertex_type_offsets is valid, "
+ "*vertex_type_offsets should be sorted.");
+ vertex_t front_element{};
+ raft::update_host(
+ &front_element, (*vertex_type_offsets).data(), size_t{1}, handle.get_stream());
+ vertex_t back_element{};
+ raft::update_host(&back_element,
+ (*vertex_type_offsets).data() + num_vertex_types,
+ size_t{1},
+ handle.get_stream());
+ handle.sync_stream();
+ CUGRAPH_EXPECTS(
+ front_element == vertex_t{0},
+ "Invalid input arguments: if vertex_type_offsets is valid, the first element of "
+ "*vertex_type_offsets should be 0.");
+ vertex_t max_v = std::max(thrust::reduce(handle.get_thrust_policy(),
+ edgelist_majors.begin(),
+ edgelist_majors.end(),
+ vertex_t{0},
+ thrust::maximum{}),
+ thrust::reduce(handle.get_thrust_policy(),
+ edgelist_minors.begin(),
+ edgelist_minors.end(),
+ vertex_t{0},
+ thrust::maximum{}));
+ CUGRAPH_EXPECTS(
+ back_element > max_v,
+ "Invalid input arguments: if vertex_type_offsets is valid, the last element of "
+ "*vertex_type_offsets should be larger than the maximum vertex ID in edgelist_majors & "
+ "edgelist_minors.");
+
+ rmm::device_uvector tmp_majors(edgelist_majors.size(), handle.get_stream());
+ rmm::device_uvector tmp_minors(edgelist_minors.size(), handle.get_stream());
+ thrust::copy(handle.get_thrust_policy(),
+ edgelist_majors.begin(),
+ edgelist_majors.end(),
+ tmp_majors.begin());
+ thrust::copy(handle.get_thrust_policy(),
+ edgelist_minors.begin(),
+ edgelist_minors.end(),
+ tmp_minors.begin());
+ if (edgelist_edge_types) {
+ rmm::device_uvector tmp_edge_types((*edgelist_edge_types).size(),
+ handle.get_stream());
+ thrust::copy(handle.get_thrust_policy(),
+ (*edgelist_edge_types).begin(),
+ (*edgelist_edge_types).end(),
+ tmp_edge_types.begin());
+ auto triplet_first =
+ thrust::make_zip_iterator(tmp_edge_types.begin(), tmp_majors.begin(), tmp_minors.begin());
+ thrust::sort(handle.get_thrust_policy(), triplet_first, triplet_first + tmp_majors.size());
+ CUGRAPH_EXPECTS(
+ thrust::count_if(
+ handle.get_thrust_policy(),
+ thrust::make_counting_iterator(size_t{0}),
+ thrust::make_counting_iterator(tmp_majors.size()),
+ [vertex_type_offsets = *vertex_type_offsets, triplet_first] __device__(size_t i) {
+ if (i > 0) {
+ auto prev = *(triplet_first + i - 1);
+ auto cur = *(triplet_first + i);
+ if (thrust::get<0>(prev) == thrust::get<0>(cur)) { // same edge type
+ auto prev_major_v_type =
+ thrust::distance(vertex_type_offsets.begin() + 1,
+ thrust::upper_bound(thrust::seq,
+ vertex_type_offsets.begin() + 1,
+ vertex_type_offsets.end(),
+ thrust::get<1>(prev)));
+ auto cur_major_v_type =
+ thrust::distance(vertex_type_offsets.begin() + 1,
+ thrust::upper_bound(thrust::seq,
+ vertex_type_offsets.begin() + 1,
+ vertex_type_offsets.end(),
+ thrust::get<1>(cur)));
+ if (prev_major_v_type != cur_major_v_type) { return true; }
+ auto prev_minor_v_type =
+ thrust::distance(vertex_type_offsets.begin() + 1,
+ thrust::upper_bound(thrust::seq,
+ vertex_type_offsets.begin() + 1,
+ vertex_type_offsets.end(),
+ thrust::get<2>(prev)));
+ auto cur_minor_v_type =
+ thrust::distance(vertex_type_offsets.begin() + 1,
+ thrust::upper_bound(thrust::seq,
+ vertex_type_offsets.begin() + 1,
+ vertex_type_offsets.end(),
+ thrust::get<2>(cur)));
+ if (prev_minor_v_type != cur_minor_v_type) { return true; }
+ }
+ }
+ return false;
+ }) == 0,
+ "Invalid input arguments: if vertex_type_offsets and edgelist_edge_types are valid, the "
+ "entire set of input edge source vertices for each edge type should have an identical "
+ "vertex type, and the entire set of input edge destination vertices for each type should "
+ "have an identical vertex type.");
+ } else {
+ auto pair_first = thrust::make_zip_iterator(tmp_majors.begin(), tmp_minors.begin());
+ thrust::sort(handle.get_thrust_policy(), pair_first, pair_first + tmp_majors.size());
+ CUGRAPH_EXPECTS(
+ thrust::count_if(
+ handle.get_thrust_policy(),
+ thrust::make_counting_iterator(size_t{0}),
+ thrust::make_counting_iterator(tmp_majors.size()),
+ [vertex_type_offsets = *vertex_type_offsets, pair_first] __device__(size_t i) {
+ if (i > 0) {
+ auto prev = *(pair_first + i - 1);
+ auto cur = *(pair_first + i);
+ auto prev_src_v_type =
+ thrust::distance(vertex_type_offsets.begin() + 1,
+ thrust::upper_bound(thrust::seq,
+ vertex_type_offsets.begin() + 1,
+ vertex_type_offsets.end(),
+ thrust::get<0>(prev)));
+ auto cur_src_v_type =
+ thrust::distance(vertex_type_offsets.begin() + 1,
+ thrust::upper_bound(thrust::seq,
+ vertex_type_offsets.begin() + 1,
+ vertex_type_offsets.end(),
+ thrust::get<0>(cur)));
+ if (prev_src_v_type != cur_src_v_type) { return true; }
+ auto prev_dst_v_type =
+ thrust::distance(vertex_type_offsets.begin() + 1,
+ thrust::upper_bound(thrust::seq,
+ vertex_type_offsets.begin() + 1,
+ vertex_type_offsets.end(),
+ thrust::get<1>(prev)));
+ auto cur_dst_v_type =
+ thrust::distance(vertex_type_offsets.begin() + 1,
+ thrust::upper_bound(thrust::seq,
+ vertex_type_offsets.begin() + 1,
+ vertex_type_offsets.end(),
+ thrust::get<1>(cur)));
+ if (prev_dst_v_type != cur_dst_v_type) { return true; }
+ }
+ return false;
+ }) == 0,
+ "Invalid input arguments: if vertex_type_offsets is valid (but "
+ "edgelist_edge_types is invalid), the entire set of input edge source "
+ "vertices should have an identical vertex type, and the entire set of "
+ "input edge destination vertices should have an identical vertex type.");
+ }
+ }
+
if (seed_vertices) {
for (size_t i = 0; i < num_labels; ++i) {
rmm::device_uvector this_label_seed_vertices(0, handle.get_stream());
@@ -356,7 +539,7 @@ compute_min_hop_for_unique_label_vertex_pairs(
std::optional> seed_vertex_label_offsets,
std::optional> edgelist_label_offsets)
{
- auto approx_edges_to_sort_per_iteration =
+ auto approx_items_to_sort_per_iteration =
static_cast(handle.get_device_properties().multiProcessorCount) *
(1 << 18) /* tuning parameter */; // for segmented sort
@@ -369,7 +552,7 @@ compute_min_hop_for_unique_label_vertex_pairs(
detail::compute_offset_aligned_element_chunks(handle,
*edgelist_label_offsets,
edgelist_vertices.size(),
- approx_edges_to_sort_per_iteration);
+ approx_items_to_sort_per_iteration);
auto num_chunks = h_label_offsets.size() - 1;
if (edgelist_hops) {
@@ -406,28 +589,28 @@ compute_min_hop_for_unique_label_vertex_pairs(
}
tmp_indices.resize(
- thrust::distance(
- tmp_indices.begin(),
- thrust::unique(handle.get_thrust_policy(),
- tmp_indices.begin(),
- tmp_indices.end(),
- [edgelist_label_offsets = *edgelist_label_offsets,
- edgelist_vertices,
- edgelist_hops = *edgelist_hops] __device__(size_t l_idx, size_t r_idx) {
- auto l_it = thrust::upper_bound(thrust::seq,
- edgelist_label_offsets.begin() + 1,
- edgelist_label_offsets.end(),
- l_idx);
- auto r_it = thrust::upper_bound(thrust::seq,
- edgelist_label_offsets.begin() + 1,
- edgelist_label_offsets.end(),
- r_idx);
- if (l_it != r_it) { return false; }
-
- auto l_vertex = edgelist_vertices[l_idx];
- auto r_vertex = edgelist_vertices[r_idx];
- return l_vertex == r_vertex;
- })),
+ thrust::distance(tmp_indices.begin(),
+ thrust::unique(handle.get_thrust_policy(),
+ tmp_indices.begin(),
+ tmp_indices.end(),
+ [edgelist_label_offsets = *edgelist_label_offsets,
+ edgelist_vertices] __device__(size_t l_idx, size_t r_idx) {
+ auto l_it =
+ thrust::upper_bound(thrust::seq,
+ edgelist_label_offsets.begin() + 1,
+ edgelist_label_offsets.end(),
+ l_idx);
+ auto r_it =
+ thrust::upper_bound(thrust::seq,
+ edgelist_label_offsets.begin() + 1,
+ edgelist_label_offsets.end(),
+ r_idx);
+ if (l_it != r_it) { return false; }
+
+ auto l_vertex = edgelist_vertices[l_idx];
+ auto r_vertex = edgelist_vertices[r_idx];
+ return l_vertex == r_vertex;
+ })),
handle.get_stream());
tmp_label_indices.resize(tmp_indices.size(), handle.get_stream());
@@ -859,17 +1042,23 @@ compute_min_hop_for_unique_label_vertex_pairs(
}
}
-template
-std::tuple, std::optional>>
-compute_renumber_map(raft::handle_t const& handle,
- raft::device_span edgelist_majors,
- raft::device_span edgelist_minors,
- std::optional> edgelist_hops,
- std::optional> seed_vertices,
- std::optional> seed_vertex_label_offsets,
- std::optional> edgelist_label_offsets)
+// returns renumber map & optional (label, type) offsets
+// indices are non-descedning)
+template
+std::tuple, std::optional>>
+compute_vertex_renumber_map(
+ raft::handle_t const& handle,
+ raft::device_span edgelist_majors,
+ raft::device_span edgelist_minors,
+ std::optional> edgelist_hops,
+ std::optional> seed_vertices,
+ std::optional> seed_vertex_label_offsets,
+ std::optional> edgelist_label_offsets,
+ std::optional> vertex_type_offsets,
+ size_t num_labels,
+ size_t num_vertex_types)
{
- auto approx_edges_to_sort_per_iteration =
+ auto approx_items_to_sort_per_iteration =
static_cast(handle.get_device_properties().multiProcessorCount) *
(1 << 20) /* tuning parameter */; // for segmented sort
@@ -892,10 +1081,9 @@ compute_renumber_map(raft::handle_t const& handle,
compute_min_hop_for_unique_label_vertex_pairs(
handle, edgelist_minors, edgelist_hops, std::nullopt, std::nullopt, edgelist_label_offsets);
+ rmm::device_uvector renumber_map(0, handle.get_stream());
+ std::optional> renumber_map_label_type_offsets{std::nullopt};
if (edgelist_label_offsets) {
- auto num_labels = (*edgelist_label_offsets).size() - 1;
-
- rmm::device_uvector renumber_map(0, handle.get_stream());
rmm::device_uvector renumber_map_label_indices(0, handle.get_stream());
renumber_map.reserve((*unique_label_major_pair_label_indices).size() +
@@ -903,8 +1091,8 @@ compute_renumber_map(raft::handle_t const& handle,
handle.get_stream());
renumber_map_label_indices.reserve(renumber_map.capacity(), handle.get_stream());
- auto num_chunks = (edgelist_majors.size() + (approx_edges_to_sort_per_iteration - 1)) /
- approx_edges_to_sort_per_iteration;
+ auto num_chunks = (edgelist_majors.size() + (approx_items_to_sort_per_iteration - 1)) /
+ approx_items_to_sort_per_iteration;
auto chunk_size = (num_chunks > 0) ? ((num_labels + (num_chunks - 1)) / num_chunks) : 0;
size_t copy_offset{0};
@@ -963,12 +1151,37 @@ compute_renumber_map(raft::handle_t const& handle,
merged_vertices.resize(merged_label_indices.size(), handle.get_stream());
merged_hops.resize(merged_label_indices.size(), handle.get_stream());
merged_flags.resize(merged_label_indices.size(), handle.get_stream());
- auto sort_key_first = thrust::make_zip_iterator(
- merged_label_indices.begin(), merged_hops.begin(), merged_flags.begin());
- thrust::sort_by_key(handle.get_thrust_policy(),
- sort_key_first,
- sort_key_first + merged_label_indices.size(),
- merged_vertices.begin());
+ if (vertex_type_offsets) {
+ auto quadraplet_first = thrust::make_zip_iterator(merged_label_indices.begin(),
+ merged_vertices.begin(),
+ merged_hops.begin(),
+ merged_flags.begin());
+ thrust::sort(
+ handle.get_thrust_policy(),
+ quadraplet_first,
+ quadraplet_first + merged_vertices.size(),
+ [offsets = *vertex_type_offsets] __device__(auto lhs, auto rhs) {
+ auto lhs_v_type = thrust::distance(
+ offsets.begin() + 1,
+ thrust::upper_bound(
+ thrust::seq, offsets.begin() + 1, offsets.end(), thrust::get<1>(lhs)));
+ auto rhs_v_type = thrust::distance(
+ offsets.begin() + 1,
+ thrust::upper_bound(
+ thrust::seq, offsets.begin() + 1, offsets.end(), thrust::get<1>(rhs)));
+ return thrust::make_tuple(
+ thrust::get<0>(lhs), lhs_v_type, thrust::get<2>(lhs), thrust::get<3>(lhs)) <
+ thrust::make_tuple(
+ thrust::get<0>(rhs), rhs_v_type, thrust::get<2>(rhs), thrust::get<3>(rhs));
+ });
+ } else {
+ auto sort_key_first = thrust::make_zip_iterator(
+ merged_label_indices.begin(), merged_hops.begin(), merged_flags.begin());
+ thrust::sort_by_key(handle.get_thrust_policy(),
+ sort_key_first,
+ sort_key_first + merged_label_indices.size(),
+ merged_vertices.begin());
+ }
} else {
auto major_triplet_first =
thrust::make_zip_iterator((*unique_label_major_pair_label_indices).begin(),
@@ -999,12 +1212,33 @@ compute_renumber_map(raft::handle_t const& handle,
handle.get_stream());
merged_vertices.resize(merged_label_indices.size(), handle.get_stream());
merged_flags.resize(merged_label_indices.size(), handle.get_stream());
- auto sort_key_first =
- thrust::make_zip_iterator(merged_label_indices.begin(), merged_flags.begin());
- thrust::sort_by_key(handle.get_thrust_policy(),
- sort_key_first,
- sort_key_first + merged_label_indices.size(),
- merged_vertices.begin());
+ if (vertex_type_offsets) {
+ auto triplet_first = thrust::make_zip_iterator(
+ merged_label_indices.begin(), merged_vertices.begin(), merged_flags.begin());
+ thrust::sort(
+ handle.get_thrust_policy(),
+ triplet_first,
+ triplet_first + merged_vertices.size(),
+ [offsets = *vertex_type_offsets] __device__(auto lhs, auto rhs) {
+ auto lhs_v_type = thrust::distance(
+ offsets.begin() + 1,
+ thrust::upper_bound(
+ thrust::seq, offsets.begin() + 1, offsets.end(), thrust::get<1>(lhs)));
+ auto rhs_v_type = thrust::distance(
+ offsets.begin() + 1,
+ thrust::upper_bound(
+ thrust::seq, offsets.begin() + 1, offsets.end(), thrust::get<1>(rhs)));
+ return thrust::make_tuple(thrust::get<0>(lhs), lhs_v_type, thrust::get<2>(lhs)) <
+ thrust::make_tuple(thrust::get<0>(rhs), rhs_v_type, thrust::get<2>(rhs));
+ });
+ } else {
+ auto sort_key_first =
+ thrust::make_zip_iterator(merged_label_indices.begin(), merged_flags.begin());
+ thrust::sort_by_key(handle.get_thrust_policy(),
+ sort_key_first,
+ sort_key_first + merged_label_indices.size(),
+ merged_vertices.begin());
+ }
}
renumber_map.resize(copy_offset + merged_vertices.size(), handle.get_stream());
@@ -1025,7 +1259,41 @@ compute_renumber_map(raft::handle_t const& handle,
renumber_map.shrink_to_fit(handle.get_stream());
renumber_map_label_indices.shrink_to_fit(handle.get_stream());
- return std::make_tuple(std::move(renumber_map), std::move(renumber_map_label_indices));
+ renumber_map_label_type_offsets =
+ rmm::device_uvector(num_labels * num_vertex_types + 1, handle.get_stream());
+ (*renumber_map_label_type_offsets).set_element_to_zero_async(0, handle.get_stream());
+ if (vertex_type_offsets) {
+ auto label_type_pair_first = thrust::make_zip_iterator(
+ renumber_map_label_indices.begin(),
+ thrust::make_transform_iterator(
+ renumber_map.begin(),
+ cuda::proclaim_return_type(
+ [offsets = *vertex_type_offsets] __device__(auto v) {
+ return static_cast(thrust::distance(
+ offsets.begin() + 1,
+ thrust::upper_bound(thrust::seq, offsets.begin() + 1, offsets.end(), v)));
+ })));
+ auto value_first = thrust::make_transform_iterator(
+ thrust::make_counting_iterator(size_t{0}),
+ cuda::proclaim_return_type>(
+ [num_vertex_types] __device__(size_t i) {
+ return thrust::make_tuple(static_cast(i / num_vertex_types),
+ static_cast(i % num_vertex_types));
+ }));
+ thrust::upper_bound(handle.get_thrust_policy(),
+ label_type_pair_first,
+ label_type_pair_first + renumber_map.size(),
+ value_first,
+ value_first + (num_labels * num_vertex_types),
+ (*renumber_map_label_type_offsets).begin() + 1);
+ } else {
+ thrust::upper_bound(handle.get_thrust_policy(),
+ renumber_map_label_indices.begin(),
+ renumber_map_label_indices.end(),
+ thrust::make_counting_iterator(label_index_t{0}),
+ thrust::make_counting_iterator(static_cast(num_labels)),
+ (*renumber_map_label_type_offsets).begin() + 1);
+ }
} else {
if (edgelist_hops) {
rmm::device_uvector merged_vertices(
@@ -1067,13 +1335,34 @@ compute_renumber_map(raft::handle_t const& handle,
merged_hops.resize(merged_vertices.size(), handle.get_stream());
merged_flags.resize(merged_vertices.size(), handle.get_stream());
- auto sort_key_first = thrust::make_zip_iterator(merged_hops.begin(), merged_flags.begin());
- thrust::sort_by_key(handle.get_thrust_policy(),
- sort_key_first,
- sort_key_first + merged_hops.size(),
- merged_vertices.begin());
+ if (vertex_type_offsets) {
+ auto triplet_first = thrust::make_zip_iterator(
+ merged_vertices.begin(), merged_hops.begin(), merged_flags.begin());
+ thrust::sort(
+ handle.get_thrust_policy(),
+ triplet_first,
+ triplet_first + merged_vertices.size(),
+ [offsets = *vertex_type_offsets] __device__(auto lhs, auto rhs) {
+ auto lhs_v_type = thrust::distance(
+ offsets.begin() + 1,
+ thrust::upper_bound(
+ thrust::seq, offsets.begin() + 1, offsets.end(), thrust::get<0>(lhs)));
+ auto rhs_v_type = thrust::distance(
+ offsets.begin() + 1,
+ thrust::upper_bound(
+ thrust::seq, offsets.begin() + 1, offsets.end(), thrust::get<0>(rhs)));
+ return thrust::make_tuple(lhs_v_type, thrust::get<1>(lhs), thrust::get<2>(lhs)) <
+ thrust::make_tuple(rhs_v_type, thrust::get<1>(rhs), thrust::get<2>(rhs));
+ });
+ } else {
+ auto sort_key_first = thrust::make_zip_iterator(merged_hops.begin(), merged_flags.begin());
+ thrust::sort_by_key(handle.get_thrust_policy(),
+ sort_key_first,
+ sort_key_first + merged_hops.size(),
+ merged_vertices.begin());
+ }
- return std::make_tuple(std::move(merged_vertices), std::nullopt);
+ renumber_map = std::move(merged_vertices);
} else {
rmm::device_uvector output_vertices(unique_label_minor_pair_vertices.size(),
handle.get_stream());
@@ -1085,7 +1374,7 @@ compute_renumber_map(raft::handle_t const& handle,
output_vertices.begin());
auto num_unique_majors = unique_label_major_pair_vertices.size();
- auto renumber_map = std::move(unique_label_major_pair_vertices);
+ renumber_map = std::move(unique_label_major_pair_vertices);
renumber_map.resize(
renumber_map.size() + thrust::distance(output_vertices.begin(), output_last),
handle.get_stream());
@@ -1094,9 +1383,370 @@ compute_renumber_map(raft::handle_t const& handle,
output_last,
renumber_map.begin() + num_unique_majors);
- return std::make_tuple(std::move(renumber_map), std::nullopt);
+ if (vertex_type_offsets) {
+ thrust::stable_sort(
+ handle.get_thrust_policy(),
+ renumber_map.begin(),
+ renumber_map.end(),
+ [offsets = *vertex_type_offsets] __device__(auto lhs, auto rhs) {
+ auto lhs_v_type = thrust::distance(
+ offsets.begin() + 1,
+ thrust::upper_bound(
+ thrust::seq, offsets.begin() + 1, offsets.end(), thrust::get<0>(lhs)));
+ auto rhs_v_type = thrust::distance(
+ offsets.begin() + 1,
+ thrust::upper_bound(
+ thrust::seq, offsets.begin() + 1, offsets.end(), thrust::get<0>(rhs)));
+ return lhs_v_type < rhs_v_type;
+ });
+ }
+ }
+
+ if (vertex_type_offsets) {
+ renumber_map_label_type_offsets =
+ rmm::device_uvector(num_vertex_types + 1, handle.get_stream());
+ (*renumber_map_label_type_offsets).set_element_to_zero_async(0, handle.get_stream());
+ auto type_first = thrust::make_transform_iterator(
+ renumber_map.begin(),
+ cuda::proclaim_return_type(
+ [offsets = *vertex_type_offsets] __device__(auto v) {
+ return static_cast(thrust::distance(
+ offsets.begin() + 1,
+ thrust::upper_bound(thrust::seq, offsets.begin() + 1, offsets.end(), v)));
+ }));
+ thrust::upper_bound(
+ handle.get_thrust_policy(),
+ type_first,
+ type_first + renumber_map.size(),
+ thrust::make_counting_iterator(vertex_type_t{0}),
+ thrust::make_counting_iterator(static_cast(num_vertex_types)),
+ (*renumber_map_label_type_offsets).begin() + 1);
+ }
+ }
+
+ return std::make_tuple(std::move(renumber_map), std::move(renumber_map_label_type_offsets));
+}
+
+// returns renumber map & optional (label, type) offsets
+template
+std::tuple, std::optional>>
+compute_edge_id_renumber_map(
+ raft::handle_t const& handle,
+ raft::device_span edgelist_edge_ids,
+ std::optional> edgelist_edge_types,
+ std::optional> edgelist_hops,
+ std::optional> edgelist_label_offsets,
+ size_t num_labels,
+ size_t num_edge_types)
+{
+ rmm::device_uvector renumber_map(0, handle.get_stream());
+ std::optional> renumber_map_label_type_offsets{std::nullopt};
+ if (edgelist_label_offsets) {
+ auto approx_items_to_sort_per_iteration =
+ static_cast(handle.get_device_properties().multiProcessorCount) *
+ (1 << 20) /* tuning parameter */; // for segmented sort
+
+ auto [h_label_offsets, h_edge_offsets] =
+ detail::compute_offset_aligned_element_chunks(handle,
+ *edgelist_label_offsets,
+ edgelist_edge_ids.size(),
+ approx_items_to_sort_per_iteration);
+ auto num_chunks = h_label_offsets.size() - 1;
+
+ rmm::device_uvector tmp_indices(edgelist_edge_ids.size(), handle.get_stream());
+ thrust::sequence(handle.get_thrust_policy(), tmp_indices.begin(), tmp_indices.end(), size_t{0});
+
+ // cub::DeviceSegmentedSort currently does not suuport thrust::tuple type keys, sorting in
+ // chunks still helps in limiting the binary search range and improving memory locality
+ for (size_t i = 0; i < num_chunks; ++i) {
+ // sort by (label, (type), id, (hop))
+
+ thrust::sort(
+ handle.get_thrust_policy(),
+ tmp_indices.begin() + h_edge_offsets[i],
+ tmp_indices.begin() + h_edge_offsets[i + 1],
+ [edgelist_label_offsets =
+ raft::device_span((*edgelist_label_offsets).data() + h_label_offsets[i],
+ (h_label_offsets[i + 1] - h_label_offsets[i]) + 1),
+ edgelist_edge_types = detail::to_thrust_optional(edgelist_edge_types),
+ edgelist_edge_ids,
+ edgelist_hops = detail::to_thrust_optional(edgelist_hops)] __device__(size_t l_idx,
+ size_t r_idx) {
+ auto l_it = thrust::upper_bound(
+ thrust::seq, edgelist_label_offsets.begin() + 1, edgelist_label_offsets.end(), l_idx);
+ auto r_it = thrust::upper_bound(
+ thrust::seq, edgelist_label_offsets.begin() + 1, edgelist_label_offsets.end(), r_idx);
+ if (l_it != r_it) { return l_it < r_it; }
+
+ if (edgelist_edge_types) {
+ auto l_type = (*edgelist_edge_types)[l_idx];
+ auto r_type = (*edgelist_edge_types)[r_idx];
+ if (l_type != r_type) { return l_type < r_type; }
+ }
+
+ auto l_id = edgelist_edge_ids[l_idx];
+ auto r_id = edgelist_edge_ids[r_idx];
+ if (l_id != r_id) { return l_id < r_id; }
+
+ if (edgelist_hops) {
+ auto l_hop = (*edgelist_hops)[l_idx];
+ auto r_hop = (*edgelist_hops)[r_idx];
+ return l_hop < r_hop;
+ }
+
+ return false;
+ });
+
+ // find unique (label, (type), id, (min_hop)) tuples
+
+ auto last = thrust::unique(
+ handle.get_thrust_policy(),
+ tmp_indices.begin() + h_edge_offsets[i],
+ tmp_indices.begin() + h_edge_offsets[i + 1],
+ [edgelist_label_offsets = *edgelist_label_offsets,
+ edgelist_edge_types = detail::to_thrust_optional(edgelist_edge_types),
+ edgelist_edge_ids] __device__(size_t l_idx, size_t r_idx) {
+ auto l_it = thrust::upper_bound(
+ thrust::seq, edgelist_label_offsets.begin() + 1, edgelist_label_offsets.end(), l_idx);
+ auto r_it = thrust::upper_bound(
+ thrust::seq, edgelist_label_offsets.begin() + 1, edgelist_label_offsets.end(), r_idx);
+ if (l_it != r_it) { return false; }
+
+ if (edgelist_edge_types) {
+ auto l_type = (*edgelist_edge_types)[l_idx];
+ auto r_type = (*edgelist_edge_types)[r_idx];
+ if (l_type != r_type) { return false; }
+ }
+
+ auto l_id = edgelist_edge_ids[l_idx];
+ auto r_id = edgelist_edge_ids[r_idx];
+ return l_id == r_id;
+ });
+
+ // sort by (label, (type), (min_hop), id)
+
+ if (edgelist_hops) {
+ thrust::sort(
+ handle.get_thrust_policy(),
+ tmp_indices.begin() + h_edge_offsets[i],
+ last,
+ [edgelist_label_offsets =
+ raft::device_span((*edgelist_label_offsets).data() + h_label_offsets[i],
+ (h_label_offsets[i + 1] - h_label_offsets[i]) + 1),
+ edgelist_edge_types = detail::to_thrust_optional(edgelist_edge_types),
+ edgelist_edge_ids,
+ edgelist_hops = detail::to_thrust_optional(edgelist_hops)] __device__(size_t l_idx,
+ size_t r_idx) {
+ auto l_it = thrust::upper_bound(
+ thrust::seq, edgelist_label_offsets.begin() + 1, edgelist_label_offsets.end(), l_idx);
+ auto r_it = thrust::upper_bound(
+ thrust::seq, edgelist_label_offsets.begin() + 1, edgelist_label_offsets.end(), r_idx);
+ if (l_it != r_it) { return l_it < r_it; }
+
+ if (edgelist_edge_types) {
+ auto l_type = (*edgelist_edge_types)[l_idx];
+ auto r_type = (*edgelist_edge_types)[r_idx];
+ if (l_type != r_type) { return l_type < r_type; }
+ }
+
+ if (edgelist_hops) {
+ auto l_hop = (*edgelist_hops)[l_idx];
+ auto r_hop = (*edgelist_hops)[r_idx];
+ return l_hop < r_hop;
+ }
+
+ auto l_id = edgelist_edge_ids[l_idx];
+ auto r_id = edgelist_edge_ids[r_idx];
+ if (l_id != r_id) { return l_id < r_id; }
+
+ return false;
+ });
+ }
+
+ // mark invalid indices
+
+ thrust::fill(handle.get_thrust_policy(),
+ last,
+ tmp_indices.begin() + h_edge_offsets[i + 1],
+ std::numeric_limits