Remove a bunch of legacy code that's no longer used (#4609)

Supersedes #4553.  That PR identified some issues with legacy code that's not really necessary anymore.

This PR removes the `graph_utils.cuh` file and gets rid of a few straggling functions that still used it.

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

Approvers:
  - Seunghwa Kang (https://github.com/seunghwak)
  - Robert Maynard (https://github.com/robertmaynard)

URL: #4609

Author: ChuckHastings

cpp/src/community/egonet_impl.cuh

@@ -17,8 +17,6 @@
 
 // #define TIMING
 
-#include "utilities/graph_utils.cuh"
-
 #include <cugraph/algorithms.hpp>
 #include <cugraph/graph.hpp>
 #include <cugraph/graph_functions.hpp>

cpp/src/components/legacy/connectivity.cu

@@ -15,7 +15,6 @@
  */
 
 #include "scc_matrix.cuh"
-#include "utilities/graph_utils.cuh"
 #include "weak_cc.cuh"
 
 #include <cugraph/algorithms.hpp>

cpp/src/layout/legacy/barnes_hut.cuh

@@ -19,7 +19,6 @@
 #include "bh_kernels.cuh"
 #include "converters/legacy/COOtoCSR.cuh"
 #include "fa2_kernels.cuh"
-#include "utilities/graph_utils.cuh"
 #include "utils.hpp"
 
 #include <cugraph/detail/utility_wrappers.hpp>

cpp/src/layout/legacy/fa2_kernels.cuh

@@ -17,7 +17,9 @@
 #pragma once
 #define restrict __restrict__
 
-#include "utilities/graph_utils.cuh"
+// From old graph_utils.cuh
+#define CUDA_MAX_BLOCKS         65535
+#define CUDA_MAX_KERNEL_THREADS 256  // kernel will launch at most 256 threads per block
 
 namespace cugraph {
 namespace detail {

cpp/src/sampling/random_walks.cuh

@@ -18,8 +18,6 @@
 //
 #pragma once
 
-#include "utilities/graph_utils.cuh"
-
 #include <cugraph/algorithms.hpp>
 #include <cugraph/detail/utility_wrappers.hpp>
 #include <cugraph/graph.hpp>

cpp/src/sampling/rw_traversals.hpp

@@ -18,8 +18,6 @@
 //
 #pragma once
 
-#include "utilities/graph_utils.cuh"
-
 #include <cugraph/api_helpers.hpp>
 #include <cugraph/graph.hpp>
 

cpp/src/structure/legacy/graph.cu

@@ -14,17 +14,18 @@
  * limitations under the License.
  */
 
-#include "utilities/graph_utils.cuh"
-
 #include <cugraph/legacy/graph.hpp>
 #include <cugraph/utilities/error.hpp>
 
+#include <raft/core/device_span.hpp>
 #include <raft/util/device_atomics.cuh>
 
 #include <rmm/exec_policy.hpp>
 
+#include <thrust/execution_policy.h>
 #include <thrust/for_each.h>
 #include <thrust/iterator/counting_iterator.h>
+#include <thrust/sequence.h>
 
 namespace {
 
@@ -69,15 +70,40 @@ namespace legacy {
 template <typename VT, typename ET, typename WT>
 void GraphViewBase<VT, ET, WT>::get_vertex_identifiers(VT* identifiers) const
 {
-  cugraph::detail::sequence<VT>(number_of_vertices, identifiers);
+  thrust::sequence(thrust::device,
+                   thrust::device_pointer_cast(identifiers),
+                   thrust::device_pointer_cast(identifiers + number_of_vertices),
+                   VT{0});
+  RAFT_CHECK_CUDA(nullptr);
 }
 
+// FIXME: Need to get rid of this function... still used in python
 template <typename VT, typename ET, typename WT>
 void GraphCompressedSparseBaseView<VT, ET, WT>::get_source_indices(VT* src_indices) const
 {
   CUGRAPH_EXPECTS(offsets != nullptr, "No graph specified");
-  cugraph::detail::offsets_to_indices<ET, VT>(
-    offsets, GraphViewBase<VT, ET, WT>::number_of_vertices, src_indices);
+  rmm::cuda_stream_view stream_view;
+
+  raft::device_span<VT> indices_span(src_indices, GraphViewBase<VT, ET, WT>::number_of_edges);
+
+  if (indices_span.size() > 0) {
+    thrust::fill(rmm::exec_policy(stream_view), indices_span.begin(), indices_span.end(), VT{0});
+
+    thrust::for_each(rmm::exec_policy(stream_view),
+                     offsets + 1,
+                     offsets + GraphViewBase<VT, ET, WT>::number_of_vertices,
+                     [indices_span] __device__(ET offset) {
+                       if (offset < static_cast<ET>(indices_span.size())) {
+                         cuda::atomic_ref<VT, cuda::thread_scope_device> atomic_counter(
+                           indices_span.data()[offset]);
+                         atomic_counter.fetch_add(VT{1}, cuda::std::memory_order_relaxed);
+                       }
+                     });
+    thrust::inclusive_scan(rmm::exec_policy(stream_view),
+                           indices_span.begin(),
+                           indices_span.end(),
+                           indices_span.begin());
+  }
 }
 
 template <typename VT, typename ET, typename WT>
@@ -152,6 +178,4 @@ void GraphCompressedSparseBaseView<VT, ET, WT>::degree(ET* degree, DegreeDirecti
 }  // namespace legacy
 }  // namespace cugraph
 
-#include "utilities/eidir_graph_utils.hpp"
-
 #include <cugraph/legacy/eidir_graph.hpp>

cpp/src/traversal/extract_bfs_paths_impl.cuh

@@ -17,7 +17,6 @@
 
 #include "detail/graph_partition_utils.cuh"
 #include "utilities/collect_comm.cuh"
-#include "utilities/graph_utils.cuh"
 
 #include <cugraph/algorithms.hpp>
 #include <cugraph/detail/utility_wrappers.hpp>