在NVIDIA CCCL中控制浮点确定性
2026年3月5日
作者：Nader Al Awar 和 Srinivas Yadav Singanaboina
概述
若多次运行相同输入数据产生相同的按位结果，则该计算被视为确定性的。虽然这看起来是一个简单的属性，但在实践中很难实现，尤其是在并行编程和浮点算术中。这是因为浮点加法和乘法并非严格满足结合律——即由于中间结果以有限精度存储时会发生舍入，(a + b) + c 可能不等于 a + (b + c)。
在NVIDIA CUDA核心计算库(CCCL) 3.1中，CUB——一个用于实现极致并行设备算法速度的底层CUDA库——新增了一个接受执行环境的单阶段API，使用户能够自定义算法行为。我们可以使用该环境配置reduce算法的确定性属性。这只能通过新的单阶段API完成，因为两阶段API不接受执行环境。
以下代码展示了如何在CUB中指定确定性级别（完整示例可通过编译器浏览器在线查看）：
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); // 可选 not_guaranteed, run_to_run (默认), 或 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。
规约操作可用的确定性级别有三种：
not_guaranteed（不保证确定性）run_to_run（运行间确定性）gpu_to_gpu（GPU间确定性）
不保证确定性
在浮点规约中，结果可能取决于元素组合的顺序。如果两次运行以不同顺序应用规约操作符，最终值可能略有差异。在许多应用中，这些微小差异是可以接受的。通过放宽对严格确定性的要求，规约实现可以任意重新排列操作顺序，从而提高运行时性能。
在CUB中，not_guaranteed放松了确定性级别。这允许使用原子操作——由于线程间无序执行导致不同运行间操作顺序不同——来计算块级部分聚合和最终规约值。整个规约也可以在一次内核启动中完成，因为原子操作将块级部分聚合合并到结果中。
非确定性规约变体通常比运行间确定性版本更快——特别是对于较小的输入数组，单次内核执行规约减少了多次内核启动带来的延迟，最小化了额外数据移动，并避免了额外同步。代价是，由于缺乏确定性行为，重复运行可能产生略有不同的结果。
运行间确定性
虽然非确定性规约提供了潜在的性能提升，CUB也提供了一种保证跨运行结果一致的模式。默认情况下，cub::DeviceReduce是运行间确定性的，这对应于在单阶段API中将确定性级别设置为run_to_run。在此模式下，使用相同输入、内核启动配置和GPU的多次调用将产生相同的输出。
这种确定性是通过将规约结构化为固定的层次树而非依赖原子操作来实现的，因为原子操作的更新顺序可能随运行而变化。在规约的每个阶段，元素首先在单个线程内组合。然后使用shuffle指令在warp内跨线程归约中间结果，接着使用共享内存进行块级归约。最后，第二个内核聚合每个块的结果以产生最终输出。由于该序列是预先确定的且与线程执行的相对时序无关，相同的输入、内核配置和GPU会产生相同的按位结果。
GPU间确定性
对于需要最高级别可重现性的应用，CUB还提供了GPU间确定性，保证在不同GPU上使用相同输入进行多次运行产生相同结果。此模式对应于将确定性级别设置为gpu_to_gpu。
为实现此确定性级别，CUB使用了可重现浮点累加器(RFA)，该解决方案基于NVIDIA GTC 2024的会议《将科学方法恢复至HPC：高性能可重现并行规约》。RFA通过将所有输入值分组到固定数量的指数范围（默认为三个桶）中，来对抗因指数不同的数字相加而产生的浮点非结合性。这种固定的结构化累加顺序确保了最终结果与GPU架构无关。
最终结果的精度取决于桶的数量：更多的桶提供更高的精度，但也增加了中间求和的次数，可能降低性能。当前实现将桶数默认为3，这是一个在性能和精度之间提供平衡的最优默认值。值得注意的是，此配置不仅是严格确定性的，而且保证了数值正确的结果，提供比传统并行规约中使用的标准成对求和更严格的误差界。
不同确定性级别下的结果差异
三种确定性级别在多次运行产生的差异量上有所不同：
- 不保证确定性：每次调用产生略有不同的求和值。
- 运行间确定性：确保在单个GPU上的每次调用产生相同值，但如果使用不同的GPU，结果可能不同。
- GPU间确定性：保证无论哪个GPU执行规约，每次调用的求和值都相同。
如下图所示，每个确定性级别（分别用绿色、蓝色和红色圆圈表示）的数组求和值随运行次数变化。水平的直线表示规约产生相同结果。
确定性性能比较
所选确定性级别影响cub::DeviceReduce的性能。不保证确定性由于其放宽的要求，提供了最高的性能。默认的运行间确定性提供良好的性能，但比不保证确定性稍慢。GPU间确定性强制了跨不同GPU最严格的可重现性，可能显著降低性能，对于大规模问题，执行时间增加20%到30%。
下图比较了在NVIDIA H200 GPU上，float32和float64输入在不同确定性要求下的性能（越低越好）。这些图清楚地显示了确定性级别的选择如何影响不同数据类型下的执行时间。
结论
随着单阶段API和显式确定性级别的引入，CUB提供了一个增强的工具箱，用于控制规约算法的行为和性能。用户可以选择最适合其需求的确定性级别：从高性能灵活的不保证模式，到可靠的运行间确定性默认模式，直至最严格的GPU间可重现性。
CUB中的确定性不仅限于规约操作。计划将这些能力扩展到更多算法，让开发者能够控制更广泛的并行CUDA原语的可重现性。有关更新和讨论，请参阅GitHub上关于扩展确定性支持的持续议题，以跟踪路线图并提供关于您希望看到确定性版本的算法的反馈。FINISHED
