CUDA 优化网格步幅循环的块数

Question

我已经开始使用 CUDA 实现一个简单的一维数组计算。按照文档，我首先尝试定义最佳块数和块大小

...
int N_array = 1000000
...
int n_threads = 256;

int n_blocks = ceil(float(N_array / n_threads));
dim3 grid(n_blocks, 1, 1);
dim3 block(n_threads, 1, 1);
...

对于内核，我使用了 nvidia 博客中建议的网格步长方法

...
int global_idx = blockIdx.x * blockDim.x + threadIdx.x;
int stride = gridDim.x * blockDim.x;
int threadsInBlock = blockDim.x;

for (unsigned long long n = global_idx; n < N_array; n += stride) {
    ...

我的问题是：

1. 可以像以前一样定义块数吗？还是应该将它们定义为请求的线程总数小于可用的 CUDA 内核数？ （认为​​这种方式的块将通过进行更多计算来利用网格步长循环）。

2. 由于对于这个大型数组，请求的线程数大于 CUDA 内核数，如果有许多块处于非活动状态，是否会受到惩罚？与请求更少的块并保持大部分活跃相比？ （这与1有关。）

Answer 1

传统观点认为，网格步长循环的网格中线程数的大小应大致匹配相关 GPU 的线程承载能力。这样做的原因是最大化暴露的并行性，这是任何 CUDA 程序员的两个最重要的目标之一。这使机器有最大的机会进行延迟隐藏。这与 CUDA 内核的数量不同。对于这些类型的设计问题，请不要考虑 GPU 中的 CUDA 内核数量。 CUDA 内核的数量与此查询不相关。

GPU的线程承载能力，忽略占用限制，是GPU中的SM数量乘以每个SM的最大线程数。

这两个数量都可以通过编程方式检索，deviceQuery sample code 演示了如何。

如果你想更精确，你可以对你的内核进行占用分析，以确定实际可以驻留在 SM 上的最大线程数，然后将其乘以 SM 的数量。占用分析可以使用作为 CUDA 工具包的一部分提供的the occupancy calculator spreadsheet 静态完成，也可以使用the occupancy API 动态完成。 （您也可以在事后使用 nsight 计算分析器检查/测量占用率。）cuda SO 标签上已经有很多问题在讨论这些事情，the programming guide 对此进行了介绍，所以我不会提供入住教程在这里。您通过占用分析得出的数字是上限，计算方法是 SM 数乘以每个 SM 的最大线程数。

您将希望根据允许达到最大值的值来选择每个块的线程和块值的数量。例如，在每个 SM 最多 1536 个线程的 cc8.6 GPU 上，您可能希望选择每个块 512 个线程，然后块数等于 GPU 中 SM 数量的 3 倍。您还可以选择每个块 256 个线程和 6 倍的 SM 数量。在此特定示例中，选择每个块 1024 个线程的值并忽略占用考虑，可能不是一个好的选择。