Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Clean up the use of cuco hash functions #4707

Merged
merged 2 commits into from
Oct 14, 2024
Merged

Clean up the use of cuco hash functions #4707

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
974bb1a
Use public cuco hash functions
PointKernel Oct 9, 2024
b4e2b3f
Use cuda::std::byte in device functions
PointKernel Oct 9, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
14 changes: 7 additions & 7 deletions cpp/src/detail/graph_partition_utils.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,7 +43,7 @@ struct compute_gpu_id_from_ext_vertex_t {

__host__ __device__ int operator()(vertex_t v) const
{
cuco::detail::MurmurHash3_32<vertex_t> hash_func{};
cuco::murmurhash3_32<vertex_t> hash_func{};
auto vertex_partition_id = static_cast<int>(hash_func(v) % comm_size);
return partition_manager::compute_global_comm_rank_from_vertex_partition_id(
major_comm_size, minor_comm_size, vertex_partition_id);
Expand All @@ -58,7 +58,7 @@ struct compute_gpu_id_from_ext_edge_id_t {

__host__ __device__ int operator()(edge_t e) const
{
cuco::detail::MurmurHash3_32<edge_t> hash_func{};
cuco::murmurhash3_32<edge_t> hash_func{};
auto vertex_partition_id = static_cast<int>(hash_func(e) % comm_size);
return partition_manager::compute_global_comm_rank_from_vertex_partition_id(
major_comm_size, minor_comm_size, vertex_partition_id);
Expand Down Expand Up @@ -88,7 +88,7 @@ struct compute_vertex_partition_id_from_ext_vertex_t {

__host__ __device__ int operator()(vertex_t v) const
{
cuco::detail::MurmurHash3_32<vertex_t> hash_func{};
cuco::murmurhash3_32<vertex_t> hash_func{};
return hash_func(v) % comm_size;
}
};
Expand All @@ -114,7 +114,7 @@ struct compute_gpu_id_from_ext_edge_endpoints_t {

__host__ __device__ int operator()(vertex_t major, vertex_t minor) const
{
cuco::detail::MurmurHash3_32<vertex_t> hash_func{};
cuco::murmurhash3_32<vertex_t> hash_func{};
auto major_vertex_partition_id = static_cast<int>(hash_func(major) % comm_size);
auto minor_vertex_partition_id = static_cast<int>(hash_func(minor) % comm_size);
auto major_comm_rank = major_vertex_partition_id % major_comm_size;
Expand All @@ -126,7 +126,7 @@ struct compute_gpu_id_from_ext_edge_endpoints_t {
__host__ __device__ int operator()(
thrust::tuple<vertex_t, vertex_t> pair /* major, minor */) const
{
cuco::detail::MurmurHash3_32<vertex_t> hash_func{};
cuco::murmurhash3_32<vertex_t> hash_func{};
auto major_vertex_partition_id = static_cast<int>(hash_func(thrust::get<0>(pair)) % comm_size);
auto minor_vertex_partition_id = static_cast<int>(hash_func(thrust::get<1>(pair)) % comm_size);
auto major_comm_rank = major_vertex_partition_id % major_comm_size;
Expand Down Expand Up @@ -192,15 +192,15 @@ struct compute_edge_partition_id_from_ext_edge_endpoints_t {

__host__ __device__ int operator()(vertex_t major, vertex_t minor) const
{
cuco::detail::MurmurHash3_32<vertex_t> hash_func{};
cuco::murmurhash3_32<vertex_t> hash_func{};
return (hash_func(major) % comm_size) * minor_comm_size +
(hash_func(minor) % comm_size) / major_comm_size;
}

__host__ __device__ int operator()(
thrust::tuple<vertex_t, vertex_t> pair /* major, minor */) const
{
cuco::detail::MurmurHash3_32<vertex_t> hash_func{};
cuco::murmurhash3_32<vertex_t> hash_func{};
return (hash_func(thrust::get<0>(pair)) % comm_size) * minor_comm_size +
(hash_func(thrust::get<1>(pair)) % comm_size) / major_comm_size;
}
Expand Down
5 changes: 3 additions & 2 deletions cpp/src/structure/remove_multi_edges_impl.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,6 +28,7 @@

#include <rmm/device_uvector.hpp>

#include <cuda/std/cstddef>
#include <thrust/binary_search.h>
#include <thrust/distance.h>
#include <thrust/iterator/zip_iterator.h>
Expand All @@ -53,8 +54,8 @@ struct hash_src_dst_pair {
vertex_t pair[2];
pair[0] = thrust::get<0>(t);
pair[1] = thrust::get<1>(t);
cuco::detail::MurmurHash3_32<vertex_t*> hash_func{};
return hash_func.compute_hash(reinterpret_cast<std::byte*>(pair), 2 * sizeof(vertex_t)) %
cuco::murmurhash3_32<vertex_t*> hash_func{};
return hash_func.compute_hash(reinterpret_cast<cuda::std::byte*>(pair), 2 * sizeof(vertex_t)) %
num_groups;
}
};
Expand Down
2 changes: 1 addition & 1 deletion cpp/tests/prims/mg_count_if_v.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,7 +48,7 @@ struct test_predicate {
test_predicate(int mod_count) : mod(mod_count) {}
__device__ bool operator()(vertex_t, const vertex_t& val)
{
cuco::detail::MurmurHash3_32<vertex_t> hash_func{};
cuco::murmurhash3_32<vertex_t> hash_func{};
return (0 == (hash_func(val) % mod));
}
};
Expand Down
2 changes: 1 addition & 1 deletion cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_intersection.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -137,7 +137,7 @@ class Tests_MGPerVPairTransformDstNbrIntersection
cugraph::get_dataframe_buffer_begin(mg_vertex_pair_buffer),
cugraph::get_dataframe_buffer_end(mg_vertex_pair_buffer),
[comm_rank, num_vertices = mg_graph_view.number_of_vertices()] __device__(size_t i) {
cuco::detail::MurmurHash3_32<size_t>
cuco::murmurhash3_32<size_t>
hash_func{}; // use hash_func to generate arbitrary vertex pairs
auto v0 = static_cast<vertex_t>(hash_func(i + comm_rank) % num_vertices);
auto v1 = static_cast<vertex_t>(hash_func(i + num_vertices + comm_rank) % num_vertices);
Expand Down
2 changes: 1 addition & 1 deletion cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_weighted_intersection.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -163,7 +163,7 @@ class Tests_MGPerVPairTransformDstNbrIntersection
cugraph::get_dataframe_buffer_begin(mg_vertex_pair_buffer),
cugraph::get_dataframe_buffer_end(mg_vertex_pair_buffer),
[comm_rank, num_vertices = mg_graph_view.number_of_vertices()] __device__(size_t i) {
cuco::detail::MurmurHash3_32<size_t>
cuco::murmurhash3_32<size_t>
hash_func{}; // use hash_func to generate arbitrary vertex pairs
auto v0 = static_cast<vertex_t>(hash_func(i + comm_rank) % num_vertices);
auto v1 = static_cast<vertex_t>(hash_func(i + num_vertices + comm_rank) % num_vertices);
Expand Down
2 changes: 1 addition & 1 deletion cpp/tests/prims/mg_transform_reduce_v.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -53,7 +53,7 @@ struct v_op_t {

__device__ auto operator()(vertex_t, vertex_t val) const
{
cuco::detail::MurmurHash3_32<vertex_t> hash_func{};
cuco::murmurhash3_32<vertex_t> hash_func{};
return cugraph::test::detail::make_property_value<property_t>(hash_func(val) % mod);
}
};
Expand Down
4 changes: 2 additions & 2 deletions cpp/tests/utilities/property_generator_kernels.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -60,7 +60,7 @@ struct vertex_property_transform {
{
static_assert(cugraph::is_thrust_tuple_of_arithmetic<property_t>::value ||
std::is_arithmetic_v<property_t>);
cuco::detail::MurmurHash3_32<vertex_t> hash_func{};
cuco::murmurhash3_32<vertex_t> hash_func{};
return make_property_value<property_t>(hash_func(v) % mod);
}
};
Expand All @@ -74,7 +74,7 @@ struct edge_property_transform {
{
static_assert(cugraph::is_thrust_tuple_of_arithmetic<property_t>::value ||
std::is_arithmetic_v<property_t>);
cuco::detail::MurmurHash3_32<vertex_t> hash_func{};
cuco::murmurhash3_32<vertex_t> hash_func{};
return make_property_value<property_t>(hash_func(src + dst) % mod);
}
};
Expand Down
Loading