Skip to content

Commit db58f54

Browse files
committed
clarifies determinism requirements
1 parent 2db9c5e commit db58f54

File tree

1 file changed

+5
-2
lines changed

1 file changed

+5
-2
lines changed

cub/cub/device/device_topk.cuh

Lines changed: 5 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -118,10 +118,13 @@ CUB_RUNTIME_FUNCTION static cudaError_t dispatch_topk_hub(
118118
//! DeviceTopK can process all of the built-in C++ numeric primitive types (`unsigned char`, `int`, `double`, etc.) as
119119
//! well as CUDA's `__half` and `__nv_bfloat16` 16-bit floating-point types.
120120
//!
121-
//! Stability
121+
//! Determinism
122122
//! ++++++++++++++++++++++++++
123123
//!
124-
//! DeviceTopK currently only provides an unstable version.
124+
//! DeviceTopK currently only supports unordered output, which may be non-deterministic for certain inputs.
125+
//! That is, if there are multiple items across the k-th position that compare equal, the subset of tied elements that
126+
//! ends up in the returned top‑k is not uniquely defined and may vary between runs. This behavior has to be explicitly
127+
//! acknowledged by the user by passing `cuda::execution::determinism::not_guaranteed`.
125128
//!
126129
//! Usage Considerations
127130
//! ++++++++++++++++++++++++++

0 commit comments

Comments
 (0)