关于简单非共享任务的 CUDA 性能问题

Question

我是 GPU 加速方面的新手。刚刚尝试了一个简单的内核在 CUDA 上的基本 LWJGL 绑定，没有共享内存，函数签名如下

__global__ void compute(
unsigned int n,
unsigned long long int* timeMs,
double* a, double* b, double* c, 
double *g_odata)

内核函数基本上是从上面的数组（timeMs、a、b、c 等）中检索线程 id 的数据并进行一些数学运算，然后将结果放入适当线程 id 上的 g_odata 数组中。 n 是要计算的线程数（当然它会检查线程 id 是否超过 n）。没有共享内存或减少。

现在，当我测量内核完成所需的总时间时，这是关于 n（总线程大小/并行度）和块大小的奇怪情况 （我有一个带有 80 个多处理器的 GPU）

通过

clock64()

我在内核函数前后添加了时间戳，并收集了每个线程的总时间，很明显，线程越多，相同任务的速度就越慢

现在问题：

1. 为什么总时间在大约 100 个线程之后开始减少？考虑到 80 个多处理器和 10K+ cuda 核心，我预计这个数字会更大，所以可能存在一些配置问题？
2. 为什么内核函数对于更多线程需要花费更多时间？执行是交错的（即调度程序可以在其中一个完成之前暂停其中一个并执行另一个）
3. 为什么在线程数达到 100 后会出现平稳行为？以及为什么它再次起飞
4. 基于区块数量的不同性能。我读到网格/块只是开发人员的视角，没有任何影响（特别是对于我完全隔离的线程，没有共享/减少）。那么为什么它很重要，以及如何选择最佳的块大小？

Answer 1

您没有显示相关代码或给我们 GPU 的类型。这使得回答具体问题变得困难。首先，如果这是消费级 GPU，则不要使用“cuda 核心”，因为它们仅适用于单精度 (

float

)。只是调出随机 GPU 的规格：根据 TechPowerUp 优秀的 GPU 数据库，RTX-4090 单精度为 82.58 TFLOPS，双精度仅为 1.29 TFLOPS。

为什么内核函数对于更多线程需要花费更多时间？执行是交错的（即调度程序可以在其中一个完成之前暂停其中一个并执行另一个）

是的，GPU 就是这样运作的。如果您查看 CUDA 编程指南中的“计算能力”表，您会发现 SM（流式多处理器）通常具有 1024-2048 个线程，但是当您将其与“算术指令表”进行比较时，吞吐量仅为大约每个时钟周期 128 个单精度指令。它的工作原理就是通过过度使用 GPU 资源来隐藏延迟。 为什么在线程数达到 100 后会出现平稳行为？以及为什么它再次起飞

对数刻度很难解释，但看起来可能是 256 个线程的凸起（？）。可能是调度程序在每个时钟周期都找不到空闲的双精度执行单元。视觉配置文件

Nsight Compute

应该能够告诉你。

请注意，单个块始终在单个 SM 上执行。因此，具有 1024 块大小的 256 个线程意味着所有线程都在同一处理器上执行，从而使其他处理器上的计算资源未被占用。 总的来说，我认为这个指标无论如何都是没有意义的。 100-1000 个线程太少了，您需要查看所有线程的吞吐量，这意味着工作项数除以总内核执行时间。

基于区块数量的不同性能。我读到网格/块只是开发人员的视角，没有任何影响（特别是对于我完全隔离的线程，没有共享/减少）。那么为什么它很重要，以及如何选择最佳的块大小？

这是错误的。块大小确实很重要。同样，如果您查看计算能力，就会发现每个 SM 的块数量以及每个 SM 的线程数量都有限制。如果您的块大小低于 64，您将无法达到 100% 的占用率。不是扭曲大小倍数的块大小也会因停用的线程而浪费资源。

区块也有启动开销，因此较少的区块可能是有益的，但区块太大也会产生负面影响。只有当旧块的所有线程完成后，新块才能启动。大块意味着在内核末尾（或在

__syncthreads()

屏障处），许多线程可以等待很少的线程，从而占用资源的时间超过必要的时间。

根据经验，每个块使用大约 128-256 个线程。如果需要，可以对不同的尺寸进行基准测试。但你需要让它成为一个有意义的基准。查看总体吞吐量并完全占用 GPU。