Something very minor, but should we place input vectors (e.g. biases) before output vectors (e.g. output)? AFAIK, that is a convention in the C++ API.

I may delete this function. I implemented a different way. After I've refactored the implementation I'll revisit whether we need this function or not. If we keep it I'll make that change.

Is false negative here a right terminology? AFAIK, false negative is something that should be reported is missing. Here, isn't it the opposite? (edges that shouldn't appear actually appear).

Changed to remove_existing_edges

Did you commit the change? I still see false negatives.

Missed the change in the documentation... I'll search the documentation and push a fix once I resolve the testing issue on one of the build configurations.

Are these false negatives? Should we better say something like @param remove_positive_samples If true, remove positive samples (edges that exist in the input graph).

False negatives can be mis-interpreted as something that should be reported but missing. I guess here false negatives mean something that is reported as negative samples but should not be reported (as it is positive samples, this actually means false positive negative samples).

Changed to remove_existing_edges

Sorry for nitpicking but better be src_biases and dst_biases to be consistent with the reset of C++ API? (use plural forms for vectors with multiple elements?)

Changed to be plural.

What are the ranges for multi-GPU?

Added comment on that.

Disable expensive_check once validated.

You can use stencil.

https://nvidia.github.io/cccl/thrust/api/function_group__stream__compaction_1ga557e8dd3130229b1a6193b3acc82ee5e.html

auto edge_first = thrust::make_zip_iterator(batch_src.begin(), batch_dst.begin());
auto new_end = thrust::remove_if(handle.get_thrust_policy(),  edge_first, edge_first + batch_src.size(), has_edge_flags.begin(), []__device__(auto flag) { return !flag; });

Change made

We can use thrust::merge if remove_duplicates is true. If I am not mistaken, both (src, dst) and (batch_src, batch_dst) are sorted.

If remove_duplicates is false, we don't need to sort, right?

Switched to merge.

Why should we sort here?

Deleted

This won't work.

edge_src|dst_property_t stores edge src|dst property values in either linear array or (key, value) pairs.

This works only if property values are store in an linear array, but with this approach, memory footprint scales as a function of O(V/sqrt(P)), so this won't scale in large scale for graphs with a relatively low average vertex degree. Using key, value pairs ensure that memory footprint stays as min(O(V/sqrt(P)), O(E/P)).

And I am uncomfortable in exploiting this implementation details outside the primitives. And if we are only visiting the edges that exist in the input graph, we can store in (key, value) pairs, but if we need to cover the entire src/dst range, edge_property_t is not a right data structure.

One robust (but wastes some communication) approach is this.

We locally sum src biases and dst biases in each rank, then call host_scalar_allgather.

We first generate edge sources in two steps.

1. In each rank, find how many sources to generate in each V/P segment (host_scalar_allgather output of local src bias sum). Then call host_scalar_all_to_all (host_scalar_multicast). Based on this, each rank generates sources for its local vertex partition range. Then, send the source values back.
2. Do the same for destinations.
3. Shuffle edges.

A key draw back is that we need to send sources & destinations back and shuffle edges as well. More communication than necessary but no memory footprint issue and minimally exploits any cugraph implementation details.

A more efficient approach is to create a P by sqrt(P) (or minor_comm_size to be exact) array (this requires understanding how the adjacency matrix is partitoined). By just collecting local bias sum values, we can compute how many edges should be generated in each box. We can send these nuumbers to the rank owning the boxes.

When we generate sources, we can generate sources in minor_com_size iterations; in each iteration we hold source biases only for the V/P vertices. When we generate destination vertices. We first just compute how many to generate per each V/P range in each box. Sum it, send to the owning GPUs (or in loops, we can bring V/P dst bias values based on which incur less communication), and get the destination vertices back. This will incur less communication but more involved.

I am inclined to implement the first approach, and if that turns out to be a performance bottleneck, we may consider implementing the second approach may be in the primitive space.... Let me know what you think.

I'll revisit this.

My first attempt at refactoring did something akin to your first suggestion. I ran into some problems in debugging and then the current approach occurred to me. My second attempt was something akin to your second approach, but then I ran into the issue that the source edge partitions are separate. This was my third attempt and works [as long as the internal implementation is an array and not a k-v implementation... I guess I didn't have a test that reached that point].

I agree that doing this properly and efficiently might need some work in primitive space.

Refactored in latest push to use option 1. Note that's only required for biased... if either src or dst is uniform that can be done directly.

I suspect the original problem I had with that approach was actually the raft bug that I later worked around. Code is much simpler this way, and I think the performance should be fine - at least for now.

Something minor, but should we better call this function only when biases is valid? (instead of taking an optional biases variable?).

I will adjust. This is residual from when we were doing the 2D partitioning... in that case we needed to do some of the setup (what became this function) even if it was uniform biases. That's no longer required... so I will fix.

I assume we don't need both the above host_scalar_allreduce and device_allgatherv here.

If we just call allgatherv on every GPU's local sum, we can compute aggregate_sum from there. We have every GPU's local sum and the aggregate sum, we can compute everything without calling an additional host_scalar_allreduce.

Cleaned up in next push.

1.1 => 1?

gpu_biases->set_element_async(gpu_biases->size() - 1, force_to_one, handle.get_stream());
handle.sync_stream(); <= this is necessary as force_to_one can become out-of-scope before this update finishes.

Or if we really want to ensure that no value is larger than 1, we may call thrust::transform on the entire gpu_biases and set value to cuda::std::min(org_value, 1.0).

I am not sure that floating point arithmetic guarantees that running inclusive scans on local_sum / aggregate_sum values guarantee that no value exceeds 1 (and if the local_sum on the last GPU is zero, non-last element may have a value larger than 1 as well).

The problem is the bug in raft's float random number generator which allows for the generation of a random number value 1.0. When I do thrust::upper_bound below, if the random number is 1.0 then the upper_bound call goes out of bounds on the gpu_counts array. Forcing this to a value larger than 1.0 lets us handle this condition.

But you are correct, what I want to do is set all values at the end of the array (where counts are 0) to this value. I will ponder a bit.

Refactored to support the case where some GPUs might have no edges.

size_t => auto const (comm_size & comm_rank are actually int values following MPI convention, I guess we won''t need to consider more than 2^31-1 GPUs in our life time.

Fixed

Unused?

Removed

You need to just collect sample counts for comm_size GPUs (no need to collect for the entire comm_size * comm_size matrix). Use device_multicast_sendrecv to send to comm_size GPUs and receive from comm_size GPUs.

Then, the code below will also become simpler.

Refactored to use shuffle_values, which is essentially this.

I will adjust. This is residual from when we were doing the 2D partitioning... in that case we needed to do some of the setup (what became this function) even if it was uniform biases. That's no longer required... so I will fix.

Cleaned up in next push.

The problem is the bug in raft's float random number generator which allows for the generation of a random number value 1.0. When I do thrust::upper_bound below, if the random number is 1.0 then the upper_bound call goes out of bounds on the gpu_counts array. Forcing this to a value larger than 1.0 lets us handle this condition.

But you are correct, what I want to do is set all values at the end of the array (where counts are 0) to this value. I will ponder a bit.

Removed

Fixed

Refactored to use shuffle_values, which is essentially this.

Refactored to support the case where some GPUs might have no edges.

Sort of our convention is to list scalar parameters at the end. We may better move num_samples after dst_biases.

Moved in next push

Delete commented out code.l

Deleted in next push

We should better test with edge masked graphs as well.

What would behavior be in an edge masked graph? That if an edge is masked out it would be allowed to be in the negative sampled edges if remove_existing_edges is set to true? I can add that tests for that.

Yes. If you call has_edge() on masked out edges, you will get false. So, those edges can be included in the negative sampling output

Done in next push

We should better test with edge masked graphs as well.

Done in next push