Skip to content

Commit fb14030

Browse files
GPUtesterGPUtester
authored
Merge pull request #1518 from rapidsai/branch-0.19
[gpuCI] Forward-merge branch-0.19 to branch-0.20 [skip ci]
2 parents 97a4b07 + 79c3ba0 commit fb14030

File tree

8 files changed

+411
-1
lines changed

8 files changed

+411
-1
lines changed

‎cpp/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -402,6 +402,7 @@ endif(BUILD_STATIC_FAISS)
402402
add_library(cugraph SHARED
403403
src/utilities/spmv_1D.cu
404404
src/utilities/cython.cu
405+
src/utilities/path_retrieval.cu
405406
src/structure/graph.cu
406407
src/linear_assignment/hungarian.cu
407408
src/link_analysis/gunrock_hits.cpp

‎cpp/include/utilities/path_retrieval.hpp‎

Lines changed: 45 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,45 @@
1+
/*
2+
* Copyright (c) 2021, NVIDIA CORPORATION.
3+
*
4+
* Licensed under the Apache License, Version 2.0 (the "License");
5+
* you may not use this file except in compliance with the License.
6+
* You may obtain a copy of the License at
7+
*
8+
* http://www.apache.org/licenses/LICENSE-2.0
9+
*
10+
* Unless required by applicable law or agreed to in writing, software
11+
* distributed under the License is distributed on an "AS IS" BASIS,
12+
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+
* See the License for the specific language governing permissions and
14+
* limitations under the License.
15+
*/
16+
17+
#pragma once
18+
19+
#include <raft/handle.hpp>
20+
21+
namespace cugraph {
22+
23+
/**
24+
* @brief Takes the results of BFS or SSSP function call and sums the given
25+
* weights along the path to the starting vertex.
26+
*
27+
* @tparam vertex_t Type of vertex identifiers. Needs to be an integral type.
28+
* @tparam weight_t Type of edge weights. Needs to be a floating point type.
29+
* @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
30+
* handles to various CUDA libraries) to run graph algorithms. Must have at least one worker stream.
31+
* @param vertices Pointer to vertex ids.
32+
* @param preds Pointer to predecessors.
33+
* @param info_weights Secondary weights along the edge from predecessor to vertex.
34+
* @param out Contains for each index the sum of weights along the path unfolding.
35+
* @param num_vertices Number of vertices.
36+
**/
37+
template <typename vertex_t, typename weight_t>
38+
void get_traversed_cost(raft::handle_t const &handle,
39+
vertex_t const *vertices,
40+
vertex_t const *preds,
41+
weight_t const *info_weights,
42+
weight_t *out,
43+
vertex_t stop_vertex,
44+
vertex_t num_vertices);
45+
} // namespace cugraph

‎cpp/src/utilities/path_retrieval.cu‎

Lines changed: 133 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,133 @@
1+
/*
2+
* Copyright (c) 2021, NVIDIA CORPORATION.
3+
*
4+
* Licensed under the Apache License, Version 2.0 (the "License");
5+
* you may not use this file except in compliance with the License.
6+
* You may obtain a copy of the License at
7+
*
8+
* http://www.apache.org/licenses/LICENSE-2.0
9+
*
10+
* Unless required by applicable law or agreed to in writing, software
11+
* distributed under the License is distributed on an "AS IS" BASIS,
12+
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+
* See the License for the specific language governing permissions and
14+
* limitations under the License.
15+
*/
16+
17+
#include <rmm/thrust_rmm_allocator.h>
18+
#include <rmm/device_uvector.hpp>
19+
20+
#include <raft/handle.hpp>
21+
22+
#include <utilities/error.hpp>
23+
#include <utilities/path_retrieval.hpp>
24+
25+
namespace cugraph {
26+
namespace detail {
27+
28+
template <typename vertex_t, typename weight_t>
29+
__global__ void get_traversed_cost_kernel(vertex_t const *vertices,
30+
vertex_t const *preds,
31+
vertex_t const *vtx_map,
32+
weight_t const *info_weights,
33+
weight_t *out,
34+
vertex_t stop_vertex,
35+
vertex_t num_vertices)
36+
{
37+
for (vertex_t i = threadIdx.x + blockIdx.x * blockDim.x; i < num_vertices;
38+
i += gridDim.x * blockDim.x) {
39+
weight_t sum = info_weights[i];
40+
vertex_t pred = preds[i];
41+
while (pred != stop_vertex) {
42+
vertex_t pos = vtx_map[pred];
43+
sum += info_weights[pos];
44+
pred = preds[pos];
45+
}
46+
out[i] = sum;
47+
}
48+
}
49+
50+
template <typename vertex_t, typename weight_t>
51+
void get_traversed_cost_impl(raft::handle_t const &handle,
52+
vertex_t const *vertices,
53+
vertex_t const *preds,
54+
weight_t const *info_weights,
55+
weight_t *out,
56+
vertex_t stop_vertex,
57+
vertex_t num_vertices)
58+
{
59+
auto stream = handle.get_stream();
60+
vertex_t max_blocks = handle.get_device_properties().maxGridSize[0];
61+
vertex_t max_threads = handle.get_device_properties().maxThreadsPerBlock;
62+
63+
dim3 nthreads, nblocks;
64+
nthreads.x = std::min<vertex_t>(num_vertices, max_threads);
65+
nthreads.y = 1;
66+
nthreads.z = 1;
67+
nblocks.x = std::min<vertex_t>((num_vertices + nthreads.x - 1) / nthreads.x, max_blocks);
68+
nblocks.y = 1;
69+
nblocks.z = 1;
70+
71+
rmm::device_uvector<vertex_t> vtx_map_v(num_vertices, stream);
72+
rmm::device_uvector<vertex_t> vtx_keys_v(num_vertices, stream);
73+
vertex_t *vtx_map = vtx_map_v.data();
74+
vertex_t *vtx_keys = vtx_keys_v.data();
75+
raft::copy(vtx_keys, vertices, num_vertices, stream);
76+
77+
thrust::sequence(rmm::exec_policy(stream)->on(stream), vtx_map, vtx_map + num_vertices);
78+
79+
thrust::stable_sort_by_key(
80+
rmm::exec_policy(stream)->on(stream), vtx_keys, vtx_keys + num_vertices, vtx_map);
81+
82+
get_traversed_cost_kernel<<<nblocks, nthreads>>>(
83+
vertices, preds, vtx_map, info_weights, out, stop_vertex, num_vertices);
84+
}
85+
} // namespace detail
86+
87+
template <typename vertex_t, typename weight_t>
88+
void get_traversed_cost(raft::handle_t const &handle,
89+
vertex_t const *vertices,
90+
vertex_t const *preds,
91+
weight_t const *info_weights,
92+
weight_t *out,
93+
vertex_t stop_vertex,
94+
vertex_t num_vertices)
95+
{
96+
CUGRAPH_EXPECTS(num_vertices > 0, "num_vertices should be strictly positive");
97+
CUGRAPH_EXPECTS(out != nullptr, "out should be of size num_vertices");
98+
cugraph::detail::get_traversed_cost_impl(
99+
handle, vertices, preds, info_weights, out, stop_vertex, num_vertices);
100+
}
101+
102+
template void get_traversed_cost<int32_t, float>(raft::handle_t const &handle,
103+
int32_t const *vertices,
104+
int32_t const *preds,
105+
float const *info_weights,
106+
float *out,
107+
int32_t stop_vertex,
108+
int32_t num_vertices);
109+
110+
template void get_traversed_cost<int32_t, double>(raft::handle_t const &handle,
111+
int32_t const *vertices,
112+
int32_t const *preds,
113+
double const *info_weights,
114+
double *out,
115+
int32_t stop_vertex,
116+
int32_t num_vertices);
117+
118+
template void get_traversed_cost<int64_t, float>(raft::handle_t const &handle,
119+
int64_t const *vertices,
120+
int64_t const *preds,
121+
float const *info_weights,
122+
float *out,
123+
int64_t stop_vertex,
124+
int64_t num_vertices);
125+
126+
template void get_traversed_cost<int64_t, double>(raft::handle_t const &handle,
127+
int64_t const *vertices,
128+
int64_t const *preds,
129+
double const *info_weights,
130+
double *out,
131+
int64_t stop_vertex,
132+
int64_t num_vertices);
133+
} // namespace cugraph

‎python/cugraph/tests/test_utils.py‎

Lines changed: 28 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,9 @@
1717
import pytest
1818

1919
import cugraph
20+
import cudf
2021
from cugraph.tests import utils
22+
import numpy as np
2123

2224

2325
def test_bfs_paths():
@@ -68,3 +70,29 @@ def test_bfs_paths_array():
6870
answer = cugraph.utils.get_traversed_path_list(df, 100)
6971

7072
assert "not in the result set" in str(ErrorMsg)
73+
74+
75+
@pytest.mark.parametrize("graph_file", utils.DATASETS)
76+
def test_get_traversed_cost(graph_file):
77+
cu_M = utils.read_csv_file(graph_file)
78+
79+
noise = cudf.Series(np.random.randint(10, size=(cu_M.shape[0])))
80+
cu_M['info'] = cu_M['2'] + noise
81+
82+
G = cugraph.Graph()
83+
G.from_cudf_edgelist(cu_M, source='0', destination='1', edge_attr='info')
84+
85+
# run SSSP starting at vertex 17
86+
df = cugraph.sssp(G, 16)
87+
88+
answer = cugraph.utilities.path_retrieval.get_traversed_cost(df, 16,
89+
cu_M['0'],
90+
cu_M['1'],
91+
cu_M['info']
92+
)
93+
94+
df = df.sort_values(by='vertex').reset_index()
95+
answer = answer.sort_values(by='vertex').reset_index()
96+
97+
assert df.shape[0] == answer.shape[0]
98+
assert np.allclose(df['distance'], answer['info'])

‎python/cugraph/utilities/__init__.py‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
# Copyright (c) 2019-2020, NVIDIA CORPORATION.
1+
# Copyright (c) 2019-2021, NVIDIA CORPORATION.
22
# Licensed under the Apache License, Version 2.0 (the "License");
33
# you may not use this file except in compliance with the License.
44
# You may obtain a copy of the License at
@@ -25,3 +25,4 @@
2525
is_cp_matrix_type,
2626
is_sp_matrix_type,
2727
)
28+
from cugraph.utilities.path_retrieval import get_traversed_cost

‎python/cugraph/utilities/path_retrieval.pxd‎

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,30 @@
1+
# Copyright (c) 2021, NVIDIA CORPORATION.
2+
# Licensed under the Apache License, Version 2.0 (the "License");
3+
# you may not use this file except in compliance with the License.
4+
# You may obtain a copy of the License at
5+
#
6+
# http://www.apache.org/licenses/LICENSE-2.0
7+
#
8+
# Unless required by applicable law or agreed to in writing, software
9+
# distributed under the License is distributed on an "AS IS" BASIS,
10+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
11+
# See the License for the specific language governing permissions and
12+
# limitations under the License.
13+
14+
# cython: profile=False
15+
# distutils: language = c++
16+
# cython: embedsignature = True
17+
# cython: language_level = 3
18+
19+
from cugraph.structure.graph_primtypes cimport *
20+
21+
cdef extern from "utilities/path_retrieval.hpp" namespace "cugraph":
22+
23+
cdef void get_traversed_cost[vertex_t, weight_t](const handle_t &handle,
24+
const vertex_t *vertices,
25+
const vertex_t *preds,
26+
const weight_t *info_weights,
27+
weight_t *out,
28+
vertex_t stop_vertex,
29+
vertex_t num_vertices) except +
30+

‎python/cugraph/utilities/path_retrieval.py‎

Lines changed: 100 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,100 @@
1+
# Copyright (c) 2021, NVIDIA CORPORATION.
2+
# Licensed under the Apache License, Version 2.0 (the "License");
3+
# you may not use this file except in compliance with the License.
4+
# You may obtain a copy of the License at
5+
#
6+
# http://www.apache.org/licenses/LICENSE-2.0
7+
#
8+
# Unless required by applicable law or agreed to in writing, software
9+
# distributed under the License is distributed on an "AS IS" BASIS,
10+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
11+
# See the License for the specific language governing permissions and
12+
# limitations under the License.
13+
14+
import numpy as np
15+
import cudf
16+
17+
from cugraph.structure.symmetrize import symmetrize
18+
from cugraph.structure.number_map import NumberMap
19+
from cugraph.utilities import path_retrieval_wrapper
20+
21+
22+
def get_traversed_cost(df, source, source_col, dest_col, value_col):
23+
"""
24+
Take the DataFrame result from a BFS or SSSP function call and sums
25+
the given weights along the path to the starting vertex.
26+
The source_col, dest_col identifiers need to match with the vertex and
27+
predecessor columns of df.
28+
29+
Input Parameters
30+
----------
31+
df : cudf.DataFrame
32+
The dataframe containing the results of a BFS or SSSP call
33+
source: int
34+
Index of the source vertex.
35+
source_col : cudf.DataFrame
36+
This cudf.Series wraps a gdf_column of size E (E: number of edges).
37+
The gdf column contains the source index for each edge.
38+
Source indices must be an integer type.
39+
dest_col : cudf.Series
40+
This cudf.Series wraps a gdf_column of size E (E: number of edges).
41+
The gdf column contains the destination index for each edge.
42+
Destination indices must be an integer type.
43+
value_col : cudf.Series
44+
This cudf.Series wraps a gdf_column of size E (E: number of edges).
45+
The gdf column contains values associated with this edge.
46+
Weight should be a floating type.
47+
48+
Returns
49+
---------
50+
df : cudf.DataFrame
51+
DataFrame containing two columns 'vertex' and 'info'.
52+
Unreachable vertices will have value the max value of the weight type.
53+
"""
54+
55+
if 'vertex' not in df.columns:
56+
raise ValueError("DataFrame does not appear to be a BFS or "
57+
"SSP result - 'vertex' column missing")
58+
if 'distance' not in df.columns:
59+
raise ValueError("DataFrame does not appear to be a BFS or "
60+
"SSP result - 'distance' column missing")
61+
if 'predecessor' not in df.columns:
62+
raise ValueError("DataFrame does not appear to be a BFS or "
63+
"SSP result - 'predecessor' column missing")
64+
65+
src, dst, val = symmetrize(source_col,
66+
dest_col,
67+
value_col)
68+
69+
symmetrized_df = cudf.DataFrame()
70+
symmetrized_df['source'] = src
71+
symmetrized_df['destination'] = dst
72+
symmetrized_df['weights'] = val
73+
74+
input_df = df.merge(symmetrized_df,
75+
left_on=['vertex', 'predecessor'],
76+
right_on=['source', 'destination'],
77+
how="left"
78+
)
79+
80+
# Set unreachable vertex weights to max float and source vertex weight to 0
81+
max_val = np.finfo(val.dtype).max
82+
input_df[['weights']] = input_df[['weights']].fillna(max_val)
83+
input_df.loc[input_df['vertex'] == source, 'weights'] = 0
84+
85+
# Renumber
86+
renumbered_gdf, renumber_map = NumberMap.renumber(input_df,
87+
["vertex"],
88+
["predecessor"],
89+
preserve_order=True)
90+
renumbered_gdf = renumbered_gdf.rename(columns={'src': 'vertex',
91+
'dst': 'predecessor'})
92+
stop_vertex = renumber_map.to_internal_vertex_id(cudf.Series(-1)).values[0]
93+
94+
out_df = path_retrieval_wrapper.get_traversed_cost(renumbered_gdf,
95+
stop_vertex)
96+
97+
# Unrenumber
98+
out_df['vertex'] = renumber_map.unrenumber(renumbered_gdf, 'vertex',
99+
preserve_order=True)["vertex"]
100+
return out_df

0 commit comments

Comments
 (0)