diff --git a/cpp/include/cugraph/detail/decompress_edge_partition.cuh b/cpp/include/cugraph/detail/decompress_edge_partition.cuh index cc4445af392..aa682b01dfd 100644 --- a/cpp/include/cugraph/detail/decompress_edge_partition.cuh +++ b/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 edge_partition, vertex_t const* input_majors, edge_t const* input_major_start_offsets, + vertex_t input_majors_size, std::vector 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(major - edge_partition.major_range_first()); vertex_t const* indices{nullptr}; thrust::optional 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, diff --git a/cpp/src/c_api/uniform_neighbor_sampling.cpp b/cpp/src/c_api/uniform_neighbor_sampling.cpp index ed458eaf1cd..6d35dfe2dbc 100644 --- a/cpp/src/c_api/uniform_neighbor_sampling.cpp +++ b/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(graph)->vertex_type_ == + reinterpret_cast(start) + ->type_, + CUGRAPH_INVALID_INPUT, + "vertex type of graph and start must match", + *error); + + CAPI_EXPECTS( + reinterpret_cast(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); diff --git a/cpp/src/sampling/detail/sampling_utils_impl.cuh b/cpp/src/sampling/detail/sampling_utils_impl.cuh index 8a88b274c94..47c02054466 100644 --- a/cpp/src/sampling/detail/sampling_utils_impl.cuh +++ b/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(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(active_majors.size()), *majors_segments, majors.data(), minors.data(),