Clean up the use of cuco hash functions (#4707)

This PR improves the use of cuco hashers by replacing detail APIs with public ones and updating device code to use `cuda::std::byte` instead of `std::byte`.

Authors:
  - Yunsong Wang (https://github.com/PointKernel)

Approvers:
  - Chuck Hastings (https://github.com/ChuckHastings)
  - Seunghwa Kang (https://github.com/seunghwak)

URL: #4707

cpp/src/detail/graph_partition_utils.cuh

@@ -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);
@@ -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);
@@ -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;
   }
 };
@@ -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;
@@ -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;
@@ -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;
   }

cpp/src/structure/remove_multi_edges_impl.cuh

@@ -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>
@@ -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;
   }
 };

cpp/tests/prims/mg_count_if_v.cu

@@ -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));
   }
 };

cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_intersection.cu

@@ -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);

cpp/tests/prims/mg_per_v_pair_transform_dst_nbr_weighted_intersection.cu

@@ -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);

cpp/tests/prims/mg_transform_reduce_v.cu

@@ -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);
   }
 };

cpp/tests/utilities/property_generator_kernels.cuh

@@ -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);
   }
 };
@@ -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);
   }
 };