Delete old nbr sampling software (#2371)

Time ran out during 22.06 development to replace the old code.

This PR renames the experimental version of uniform neighborhood sampling to the regular name and eliminates the original implementation.

Tagged as `breaking` since we eliminate a function from the C++ API.

Authors:
  - Chuck Hastings (https://github.com/ChuckHastings)

Approvers:
  - Seunghwa Kang (https://github.com/seunghwak)
  - Brad Rees (https://github.com/BradReesWork)

URL: #2371

cpp/CMakeLists.txt

@@ -186,10 +186,8 @@ add_library(cugraph
     src/community/legacy/egonet.cu
     src/sampling/neighborhood.cu
     src/sampling/random_walks.cu
-    src/sampling/detail/gather_utils_impl.cu
     src/sampling/detail/sampling_utils_mg.cu
     src/sampling/detail/sampling_utils_sg.cu
-    src/sampling/nbr_sampling_mg.cu
     src/sampling/uniform_neighbor_sampling_mg.cpp
     src/sampling/uniform_neighbor_sampling_sg.cpp
     src/cores/legacy/core_number.cu

cpp/include/cugraph/algorithms.hpp

@@ -1468,42 +1468,6 @@ void core_number(raft::handle_t const& handle,
                  size_t k_last           = std::numeric_limits<size_t>::max(),
                  bool do_expensive_check = false);
 
-/**
- * @brief Multi-GPU Uniform Neighborhood Sampling.
- * @deprecated will be removed later in this release (22.06)
- *
- * @tparam graph_view_t Type of graph view.
- * @tparam gpu_t Type of rank (GPU) indices;
- * @tparam index_t Type used for indexing; typically edge_t
- * @param handle RAFT handle object to encapsulate resources (e.g. CUDA stream, communicator, and
- * handles to various CUDA libraries) to run graph algorithms.
- * @param graph_view Graph View object to generate NBR Sampling on.
- * @param ptr_d_starting_vertices Device array of starting vertex IDs for the NBR Sampling.
- * @param ptr_d_ranks Device array of: rank IDs (GPU IDs) for the NBR Sampling.
- * @param num_starting_vertices size of starting vertex set
- * @param h_fan_out vector of branching out (fan-out) degree per source vertex for each level
- * parameter used for obtaining local out-degree information
- * @param with_replacement boolean flag specifying if random sampling is done with replacement
- * (true); or, without replacement (false); default = true;
- * @return tuple of tuple of device vectors and counts:
- * ((vertex_t source_vertex, vertex_t destination_vertex, int rank, edge_t index), rx_counts)
- */
-template <typename graph_view_t,
-          typename gpu_t,
-          typename index_t = typename graph_view_t::edge_type>
-std::tuple<std::tuple<rmm::device_uvector<typename graph_view_t::vertex_type>,
-                      rmm::device_uvector<typename graph_view_t::vertex_type>,
-                      rmm::device_uvector<gpu_t>,
-                      rmm::device_uvector<index_t>>,
-           std::vector<size_t>>
-uniform_nbr_sample(raft::handle_t const& handle,
-                   graph_view_t const& graph_view,
-                   typename graph_view_t::vertex_type const* ptr_d_starting_vertices,
-                   gpu_t const* ptr_d_ranks,
-                   size_t num_starting_vertices,
-                   std::vector<int> const& h_fan_out,
-                   bool with_replacement = true);
-
 /**
  * @brief Uniform Neighborhood Sampling.
  *

cpp/include/cugraph/detail/decompress_edge_partition.cuh

@@ -255,7 +255,7 @@ __global__ void partially_decompress_to_edgelist_mid_degree(
       edge_partition.local_edges(major_partition_offset);
 
     auto major_offset = input_major_start_offsets[idx];
-    for (edge_t i = threadIdx.x; i < local_degree; i += blockDim.x) {
+    for (edge_t i = lane_id; i < local_degree; i += raft::warp_size()) {
       output_majors[major_offset + i] = major;
       output_minors[major_offset + i] = indices[i];
 
@@ -290,7 +290,10 @@ void partially_decompress_edge_partition_to_fill_edgelist(
   vertex_t* minors,
   thrust::optional<weight_t*> weights,
   thrust::optional<thrust::tuple<prop_t const*, prop_t*>> property,
-  thrust::optional<thrust::tuple<edge_t const*, edge_t*>> global_edge_index)
+  thrust::optional<thrust::tuple<edge_t const*, edge_t*>> global_edge_index,
+  // FIXME:  Once PR 2356 is merged, this parameter could go away because
+  // major_hypersparse_first will be part of edge_partition
+  std::optional<std::vector<vertex_t>> local_edge_partition_segment_offsets)
 {
   auto execution_policy = handle.get_thrust_policy();
   static_assert(detail::num_sparse_segments_per_vertex_partition == 3);
@@ -408,6 +411,10 @@ void partially_decompress_edge_partition_to_fill_edgelist(
                           ? thrust::make_optional(thrust::make_tuple(
                         thrust::get<0>(*property) + segment_offsets[3], thrust::get<1>(*property)))
                           : thrust::nullopt,
+       // FIXME:  Once PR 2356 is merged, this parameter could go away because
+       // major_hypersparse_first will be part of edge_partition
+       segment_offsets_last =
+         (*local_edge_partition_segment_offsets)[detail::num_sparse_segments_per_vertex_partition],
        global_edge_index] __device__(auto idx) {
         auto major        = input_majors[idx];
         auto major_offset = input_major_start_offsets[idx];
@@ -416,7 +423,10 @@ void partially_decompress_edge_partition_to_fill_edgelist(
           vertex_t const* indices{nullptr};
           thrust::optional<weight_t const*> weights{thrust::nullopt};
           edge_t local_degree{};
-          thrust::tie(indices, weights, local_degree) = edge_partition.local_edges(*major_idx);
+          // FIXME:  Once PR 2356 is merged, this computation should be changed to use
+          // major_hypersparse_first which will be part of edge_partition
+          thrust::tie(indices, weights, local_degree) =
+            edge_partition.local_edges(segment_offsets_last + *major_idx);
           thrust::fill(
             thrust::seq, majors + major_offset, majors + major_offset + local_degree, major);
           thrust::copy(thrust::seq, indices, indices + local_degree, minors + major_offset);

cpp/include/cugraph_c/sampling_algorithms.h

@@ -122,9 +122,6 @@ typedef struct {
  * @param [in]  graph        Pointer to graph.  NOTE: Graph might be modified if the storage
  *                           needs to be transposed
  * @param [in]  start        Device array of start vertices for the sampling
- * @param [in]  start_label  Device array of start labels.  These labels will propagate to the
- * results so that the result can be properly organized when the input needs to be sent back to
- * different callers (different processes or different gpus).
  * @param [in]  fanout       Host array defining the fan out at each step in the sampling algorithm
  * @param [in]  with_replacement
  *                           Boolean value.  If true selection of edges is done with
@@ -136,37 +133,7 @@ typedef struct {
  *                           be populated if error code is not CUGRAPH_SUCCESS
  * @return error code
  */
-// FIXME: This older API will be phased out this release in favor of the experimental one below
 cugraph_error_code_t cugraph_uniform_neighbor_sample(
-  const cugraph_resource_handle_t* handle,
-  cugraph_graph_t* graph,
-  const cugraph_type_erased_device_array_view_t* start,
-  const cugraph_type_erased_device_array_view_t* start_label,
-  const cugraph_type_erased_host_array_view_t* fan_out,
-  bool_t with_replacement,
-  bool_t do_expensive_check,
-  cugraph_sample_result_t** result,
-  cugraph_error_t** error);
-
-/**
- * @brief     Uniform Neighborhood Sampling
- *
- * @param [in]  handle       Handle for accessing resources
- * @param [in]  graph        Pointer to graph.  NOTE: Graph might be modified if the storage
- *                           needs to be transposed
- * @param [in]  start        Device array of start vertices for the sampling
- * @param [in]  fanout       Host array defining the fan out at each step in the sampling algorithm
- * @param [in]  with_replacement
- *                           Boolean value.  If true selection of edges is done with
- *                           replacement.  If false selection is done without replacement.
- * @param [in]  do_expensive_check
- *                           A flag to run expensive checks for input arguments (if set to true)
- * @param [in]  result       Output from the uniform_neighbor_sample call
- * @param [out] error        Pointer to an error object storing details of any error.  Will
- *                           be populated if error code is not CUGRAPH_SUCCESS
- * @return error code
- */
-cugraph_error_code_t cugraph_experimental_uniform_neighbor_sample(
   const cugraph_resource_handle_t* handle,
   cugraph_graph_t* graph,
   const cugraph_type_erased_device_array_view_t* start,