Skip to content

Commit

Permalink
Fix race condition in raft::random::discrete (#1094)
Browse files Browse the repository at this point in the history
The stream needs to be passed to `cub::DeviceScan`.

Authors:
  - Louis Sugy (https://github.com/Nyrio)

Approvers:
  - Corey J. Nolet (https://github.com/cjnolet)

URL: #1094
  • Loading branch information
Nyrio authored Dec 13, 2022
1 parent c4132b8 commit 057fb37
Showing 1 changed file with 3 additions and 2 deletions.
5 changes: 3 additions & 2 deletions cpp/include/raft/random/detail/rng_impl.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -261,10 +261,11 @@ std::enable_if_t<std::is_integral_v<OutType>> discrete(RngState& rng_state,
// Compute the cumulative sums of the weights
size_t temp_storage_bytes = 0;
rmm::device_uvector<WeightType> weights_csum(len, stream);
cub::DeviceScan::InclusiveSum(nullptr, temp_storage_bytes, weights, weights_csum.data(), len);
cub::DeviceScan::InclusiveSum(
nullptr, temp_storage_bytes, weights, weights_csum.data(), len, stream);
rmm::device_uvector<uint8_t> temp_storage(temp_storage_bytes, stream);
cub::DeviceScan::InclusiveSum(
temp_storage.data(), temp_storage_bytes, weights, weights_csum.data(), len);
temp_storage.data(), temp_storage_bytes, weights, weights_csum.data(), len, stream);

// Sample indices with replacement
RAFT_CALL_RNG_FUNC(rng_state,
Expand Down

0 comments on commit 057fb37

Please sign in to comment.