diff --git a/cub/agent/single_pass_scan_operators.cuh b/cub/agent/single_pass_scan_operators.cuh index 06a0fec4c..f461dd3f2 100644 --- a/cub/agent/single_pass_scan_operators.cuh +++ b/cub/agent/single_pass_scan_operators.cuh @@ -154,6 +154,20 @@ __device__ __forceinline__ void delay_or_prevent_hoisting(int ns) ((void)ns; __threadfence_block();)); } +struct rng_t +{ + static constexpr unsigned int a = 16807; + static constexpr unsigned int c = 0; + static constexpr unsigned int m = 1u << 31; + + unsigned int seed; + + __device__ unsigned int operator()(unsigned int min, unsigned int max) + { + return (seed = (a * seed + c) % m) % (max + 1 - min) + min; + } +}; + } /** @@ -314,11 +328,12 @@ struct ScanTileState tile_descriptor = reinterpret_cast(alias); } + detail::rng_t rng{static_cast(tile_idx)}; int delay = CUB_DETAIL_INITIAL_L2_BACKOFF_NS; while (WARP_ANY((tile_descriptor.status == SCAN_TILE_INVALID), 0xffffffff)) { - detail::delay_or_prevent_hoisting(delay <<= 1); + detail::delay_or_prevent_hoisting(rng(0, delay <<= 1)); TxnWord alias = detail::load_relaxed(d_tile_descriptors + TILE_STATUS_PADDING + tile_idx); tile_descriptor = reinterpret_cast(alias); }