无输入CUDA随机采样函数的最优线程/块数量探究

我发现现有关于优化CUDA函数线程与块数量的资料，大多围绕匹配固定问题规模展开，但我遇到的是无固定问题规模的任务：输入是随机数，最终要归约为单个标量。举个例子，用蒙特卡洛方法计算π——生成均匀随机点x∈(-1,1)和y∈(-1,1)，统计落入单位圆内的点数。我的思路是给每个线程分配独立的随机数生成器，用一个和总线程数一致的输出数组，让每个线程无需原子操作/竞争就能递增计数，最后对数组求和得到结果。

我原本觉得最优方案是用1个块、1024个线程，因为资料说NVIDIA GPU一个块最多跑1024个线程。我认为终止线程再开新块没优势，而且如果随机数生成器是有状态的，在线程内用for循环持续运行反而更高效。但我是不是忽略了什么？这类蒙特卡洛的用例很常见，为什么没人这么做？

编辑补充：实验结果与疑问 #

实际测试发现，更多块的配置性能反而更优（但需要为更多块准备随机种子）。我用Numba做了如下测试：

import cupy as cp
import numba as nb

import numba.cuda
from numba.cuda.random import create_xoroshiro128p_states
from numba.cuda.random import xoroshiro128p_uniform_float32

@nb.cuda.jit
def run_gpu(states, out, num_times):
    i = nb.cuda.grid(1)
    for _ in range(num_times):
        if xoroshiro128p_uniform_float32(states, i) > 0.5:
            out[i] += 1
global_seed = [0]
def run(num_times_total, num_blocks):
    global_seed[0] += 1
    states = create_xoroshiro128p_states(1024 * num_blocks, seed=global_seed[0])
    out = cp.zeros(1024 * num_blocks, dtype=np.float32)
    run_gpu[num_blocks, 1024](states, out, num_times_total // num_blocks)
    return out.sum() / 1024 / num_times_total

这个测试是对0.5到1的均匀分布积分（理论结果约为0.5），每次调用run时增加num_blocks，同时减少每个线程内循环的num_times，保证总计算量num_times_total恒定（测试中设为4*1024**2）。在GeForce RTX 3060上用%timeit得到的结果如下：

实验结果显示最优配置根本不是1个块，这是为什么？我到底忽略了什么？明明GPU一个块最多跑1024个线程，为什么把计算分散到更多块，比延长单个线程的for循环效率更高？

你之前的误解核心是混淆了单块线程上限和GPU的整体并行能力，具体原因如下：

1. GPU的SM与硬件并行逻辑
   NVIDIA GPU的核心是多个流式多处理器（SM），比如RTX 3060有35个SM。每个SM可以同时运行多个线程块（通常每个SM能容纳8-16个块，取决于资源占用）。你用1个块的话，只有1个SM在工作，剩下34个SM完全闲置——这才是性能低下的根本原因。

2. 线程块的调度特性
   GPU的调度单位是线程块，而非单个线程。当你启动多个块时，GPU会把这些块分配到不同的SM上并行执行。比如用512个块，GPU能把这些块填满所有SM，让所有计算核心同时工作，自然比单块单SM的效率高得多。

3. 线程内循环的串行本质
   单个线程的for循环是串行执行的，哪怕给一个线程分配10000次循环，它也只能一次处理一个计算任务。而多块多线程是真正的并行执行，多个线程同时做计算，总吞吐量自然远超单线程循环。

4. 随机数生成的状态问题
   你担心的有状态随机数生成器其实不是障碍——现代CUDA随机数库（比如你用的xoroshiro128p）支持给每个线程分配独立的状态，只需要初始化时生成足够多的种子即可，这也是你能顺利扩展块数量的原因。

5. 性能曲线的合理性
   从测试数据看，块数量增加到512-1024时，内核时间降到最低，这是因为此时所有SM都被充分利用；当块数量超过SM能容纳的上限后，GPU需要分批调度块，额外的调度开销导致性能下降，这就是2048块之后时间回升的原因。

总结来说，GPU的性能优势来自多SM的并行执行，单块配置完全浪费了大部分硬件资源，这就是多块配置性能远优于单块的核心原因。

内容的提问来源于stack exchange，提问作者Jim Pivarski