中文内容
如果多次使用相同的输入数据运行某项计算均能生成逐位完全一致的结果，则该计算被视为具有确定性。尽管这看似易于保证的特性，但在实践中往往难以实现，尤其是在并行编程与浮点运算中。这是因为浮点加法与乘法并不严格遵循结合律——即 (a + b) + c 可能不等于 a + (b + c)——其原因在于中间结果以有限精度存储时会发生舍入。
伴随 NVIDIA CUDA Core Compute Libraries (CCCL) 3.1 的发布，CUB——一个面向极限性能并行设备算法的底层 CUDA 库——新增了一个接受执行环境的单阶段 API，使用户能够自定义算法行为。我们可以利用该环境来配置 reduce 算法的确定性属性。此操作只能通过新的单阶段 API 实现，因为两阶段 API 不接受执行环境。
以下代码展示了如何在 CUB 中指定确定性级别（可使用 compiler explorer 在线查找完整示例）。
auto input = thrust::device_vector<float>{0.0f, 1.0f, 2.0f, 3.0f};
auto output = thrust::device_vector<float>(1);
auto env = cuda::execution::require(cuda::execution::determinism::not_guaranteed); // can be not_guaranteed, run_to_run (default), or gpu_to_gpu
auto error = cub::DeviceReduce::Sum(input.begin(), output.begin(), input.size(), env);
if (error != cudaSuccess)
{
std::cerr << "cub::DeviceReduce::Sum failed with status: " << error << std::endl;
}
assert(output[0] == 6.0f);
我们首先指定输入和输出向量。随后，我们使用 cuda::execution::require() 构造一个 cuda::std::execution::env 对象，并将确定性级别设置为 not_guaranteed。
可用于归约的确定性级别共有三种，分别为：
- 不保证
- 运行间
- GPU间
不保证确定性
在浮点归约运算中，结果可能取决于元素结合的顺序。若两次运行以不同顺序应用归约算子，最终值可能会略有差异。在许多应用场景中，此类微小差异是可以接受的。通过放宽对严格确定性的要求，归约实现可按任意顺序重排运算，从而提升运行时性能。
在 CUB 中，not_guaranteed 放宽了确定性级别。这使得原子操作——其在线程间的无序执行会导致各次运行间的操作顺序不同——能够同时计算块级部分聚合结果与最终归约值。整个归约过程也可在单次内核启动中完成，因为原子操作会直接将块级部分聚合结果合并至最终结果中。
非确定性归约变体通常比运行间确定性版本更快——尤其是对于较小的输入数组，因为在单次内核中执行归约能够减少多次内核启动带来的延迟、最大限度地减少额外的数据移动，并避免额外的同步。其代价在于，由于缺乏确定性行为，重复运行可能会产生略微不同的结果。
运行间确定性
尽管非确定性归约能带来潜在的性能提升，但 CUB 也提供了一种保证各次运行结果一致的模式。默认情况下，cub::DeviceReduce 具备运行间确定性，这对应于在单阶段 API 中将确定性级别设置为 run_to_run。在此模式下，使用相同的输入、内核启动配置及 GPU 进行多次调用，将产生完全相同的输出。
这种确定性是通过将归约构建为固定的层次树来实现的，而非依赖原子操作（其更新顺序在不同运行间可能发生变化）。在归约的每个阶段，元素首先在单个线程内进行合并。随后，利用 shuffle 指令在 warp 内的线程间对中间结果进行归约，再通过共享内存执行块级归约。最后，第二个内核会聚合各块的结果以生成最终输出。由于该流程是预先确定且独立于线程执行的相对时序的，因此在相同的输入、内核配置及 GPU 下，将产生逐位完全相同的结果。
GPU 到 GPU 确定性
对于需要最高级别可重复性的应用，CUB 还提供了 GPU 到 GPU 确定性，可保证在相同输入条件下，不同 GPU 上的多次运行结果完全一致。此模式对应将确定性级别设置为 gpu_to_gpu 。
为实现该级别的确定性，CUB 采用了可重复浮点累加器（Reproducible Floating-point Accumulator, RFA），该方案基于 NVIDIA GTC 2024 会议演讲《Restoring the Scientific Method to HPC: High Performance Reproducible Parallel Reductions》。RFA 通过将所有输入值划分至固定数量的指数范围（默认分为三个 bin）来应对浮点运算的非结合性问题（该现象源于对不同指数的数值进行相加）。这种固定的结构化累加顺序确保了最终结果独立于 GPU 架构。
最终结果的精度取决于 bin 的数量：bin 越多，精度越高，但也会增加中间求和的次数，从而可能降低性能。当前实现默认将 bin 数量设置为三个，这是在性能与精度之间取得平衡的最佳默认值。值得注意的是，该配置不仅具备严格的确定性，还能保证数值计算的正确性，其提供的误差界限比并行归约中传统使用的标准成对求和法更为严格。
结果随确定性级别的变化情况
这三种确定性级别在多次运行中产生的差异程度有所不同：
- 不保证确定性在每次调用时会产生略有不同的求和值。
- 运行间确定性确保在单张 GPU 上每次调用均得到相同的值，但更换不同 GPU 时结果可能会有所不同。
- GPU 间确定性保证每次调用的求和值完全一致，无论具体由哪张 GPU 执行归约操作。
如图 1 所示，各确定性级别的数组求和结果（分别以绿色、蓝色和红色圆圈表示）随运行次数绘制。呈现为一条水平直线表示归约操作产生了相同的结果。
确定性性能对比
所选的确定性级别会影响 cub::DeviceReduce 的性能。不保证确定性因其要求较为宽松，可提供最高性能。默认的跨运行确定性表现良好，但速度略低于不保证确定性。跨GPU确定性强制要求在不同GPU间实现最严格的可重现性，会显著降低性能，在处理大规模问题时会使执行时间增加20%至30%。
图2对比了在 NVIDIA H200 GPU 上针对不同确定性要求处理 float32 和 float64 输入时的性能（数值越低越好）。图表清晰地展示了确定性级别的选择对不同数据类型执行时间的影响。
结论
随着单阶段 API 和显式确定性级别的引入，CUB 提供了一个功能增强的工具箱，用于控制归约算法的行为与性能。用户可根据自身需求选择最适合的确定性级别：从高性能且灵活但不作保证的模式，到可靠的跨次运行一致性默认模式，直至最严格的 GPU 间可复现性。
CUB 中的确定性不仅限于归约操作。我们计划将这些能力扩展至更多算法，使开发者能够控制更广泛的并行 CUDA 原语的可复现性。如需获取更新和参与讨论，请参阅 GitHub 上关于扩展确定性支持的现有议题，以跟踪我们的开发路线图，并就您希望推出确定性版本的算法提供反馈。
标签
