fix non-deterministic bug in uniform neighborhood sampling (rapidsai#…

…2477)

Found a bug in uniform neighborhood sampling.  Called a function that returns an optional.  In SG mode the optional is always null.  It was being dereferenced and passed to a function where random memory contents were being used.

Fixed the SG code to not use those values.

Closes rapidsai#2446

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

Approvers:
  - Alex Barghi (https://github.com/alexbarghi-nv)
  - Seunghwa Kang (https://github.com/seunghwak)

URL: rapidsai#2477

cpp/include/cugraph/detail/decompress_edge_partition.cuh

@@ -285,6 +285,7 @@ void partially_decompress_edge_partition_to_fill_edgelist(
   edge_partition_device_view_t<vertex_t, edge_t, weight_t, multi_gpu> edge_partition,
   vertex_t const* input_majors,
   edge_t const* input_major_start_offsets,
+  vertex_t input_majors_size,
   std::vector<vertex_t> const& segment_offsets,
   vertex_t* majors,
   vertex_t* minors,
@@ -458,7 +459,7 @@ void partially_decompress_edge_partition_to_fill_edgelist(
     thrust::for_each(
       handle.get_thrust_policy(),
       thrust::make_counting_iterator(vertex_t{0}),
-      thrust::make_counting_iterator(edge_partition.major_range_size()),
+      thrust::make_counting_iterator(input_majors_size),
       [edge_partition,
        input_majors,
        input_major_start_offsets,
@@ -471,13 +472,10 @@ void partially_decompress_edge_partition_to_fill_edgelist(
        global_edge_index] __device__(auto idx) {
         auto major        = input_majors[idx];
         auto major_offset = input_major_start_offsets[idx];
-        auto major_partition_offset =
-          static_cast<size_t>(major - edge_partition.major_range_first());
         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_partition_offset);
+        thrust::tie(indices, weights, local_degree) = edge_partition.local_edges(major);
 
         // FIXME: This can lead to thread divergence if local_degree varies significantly
         //        within threads in this warp
@@ -497,7 +495,7 @@ void partially_decompress_edge_partition_to_fill_edgelist(
                        major_input_property);
         }
         if (global_edge_index) {
-          auto adjacency_list_offset = thrust::get<0>(*global_edge_index)[major_partition_offset];
+          auto adjacency_list_offset = thrust::get<0>(*global_edge_index)[major];
           auto minor_map             = thrust::get<1>(*global_edge_index);
           thrust::sequence(thrust::seq,
                            minor_map + major_offset,

cpp/src/c_api/uniform_neighbor_sampling.cpp

@@ -165,6 +165,21 @@ extern "C" cugraph_error_code_t cugraph_uniform_neighbor_sample(
   cugraph_sample_result_t** result,
   cugraph_error_t** error)
 {
+  CAPI_EXPECTS(
+    reinterpret_cast<cugraph::c_api::cugraph_graph_t*>(graph)->vertex_type_ ==
+      reinterpret_cast<cugraph::c_api::cugraph_type_erased_device_array_view_t const*>(start)
+        ->type_,
+    CUGRAPH_INVALID_INPUT,
+    "vertex type of graph and start must match",
+    *error);
+
+  CAPI_EXPECTS(
+    reinterpret_cast<cugraph::c_api::cugraph_type_erased_host_array_view_t const*>(fan_out)
+        ->type_ == INT32,
+    CUGRAPH_INVALID_INPUT,
+    "fan_out should be of type int",
+    *error);
+
   uniform_neighbor_sampling_functor functor{
     handle, graph, start, fan_out, with_replacement, do_expensive_check};
   return cugraph::c_api::run_algorithm(graph, functor, result, error);

cpp/src/sampling/detail/sampling_utils_impl.cuh

@@ -775,6 +775,7 @@ gather_one_hop_edgelist(
         partition,
         active_majors.cbegin(),
         active_majors_out_offsets.cbegin(),
+        static_cast<vertex_t>(active_majors.size()),
         majors_segments,
         output_offset + majors.data(),
         output_offset + minors.data(),
@@ -834,6 +835,7 @@ gather_one_hop_edgelist(
       partition,
       active_majors.cbegin(),
       active_majors_out_offsets.cbegin(),
+      static_cast<vertex_t>(active_majors.size()),
       *majors_segments,
       majors.data(),
       minors.data(),