CUDA 中的 GPU 使用率低答案

【问题标题】：Low GPU usage in CUDACUDA 中的 GPU 使用率低
【发布时间】：2017-02-02 21:43:56
【问题描述】：

我实现了一个程序，它使用来自不同 CPU 线程的不同 CUDA 流。内存复制是通过cudaMemcpyAsync 使用这些流实现的。内核启动也使用这些流。该程序正在执行双精度计算（我怀疑这是罪魁祸首，但是，cuBlas 达到了 75-85% 的 CPU 使用率，用于 doubles 的矩阵相乘）。还有减少操作，但是它们是通过if(threadIdx.x < s) 实现的，s 在每次迭代中减少 2 次，因此其他块应该可以使用停滞的扭曲。该应用程序是 GPU 和 CPU 密集型的，一旦前一个工作完成，它就会开始另一个工作。所以我希望它能够达到 100% 的 CPU 或 GPU。

问题在于，如果信任 GPU-Z 1.9.0，我的程序会产生 30-40% 的 GPU 负载（以及大约 50% 的 CPU 负载）。内存控制器负载为 9-10%，总线接口负载为 6%。这是因为 CPU 线程数等于 CPU 内核数。如果我将 CPU 线程数加倍，负载将保持不变（包括 CPU 负载）。

那是为什么呢？瓶颈在哪里？

我正在使用 GeForce GTX 560 Ti、CUDA 8RC、MSVC++2013、Windows 10。

我的一个猜测是，Windows 10 采用了一些积极的节能措施，即使 GPU 和 CPU 温度较低，电源计划设置为“高性能”并且电源为 700W，而功耗与最大 CPU 和 GPU TDP大约是550W。

另一个猜测是双精度速度是单精度速度的1/12，因为我的卡上每12个单精度CUDA核心就有1个双精度CUDA核心，GPU-Z占100%使用所有单精度和双精度内核时的情况。但是，数字并不完全匹配。

【问题讨论】：

您真的不希望有人会根据您提供的信息告诉您可能限制应用程序性能的因素吗？ NSight 专门为此目的提供了指导性的性能分析和指标。使用它们。
@talonmies，我认为这是一个常见问题，特别是如果我对双精度的猜测是正确的。 NSight 最近放弃了对 5xx 系列的支持，所以我不能使用它。
Linux 中的 NSight 仍然支持 Fermi 卡，您可以尝试使用它。或者只是购买一张新卡。低于预期的性能是一个很常见的问题，好吧，但从来没有一个单一的原因。如果没有代码或性能指标，您如何想象有人可能会告诉您在您的案例中可能发生的情况？
@talonmies，我认为低占用率是一个常见的问题，而过度使用共享内存或寄存器是常见的原因。我想到了共享内存限制，但忘记了寄存器。

标签： performance cuda load gpu gpgpu

【解决方案1】：

显然原因是由于 CUDA 线程默认使用太多寄存器而导致占用率低。要告诉编译器每个线程的寄存器数量限制，可以使用__launch_bounds__，如here 所述。因此，为了能够在 560 Ti 中启动所有 1536 个线程，对于块大小 256，可以指定以下内容：

_global__ void __launch_bounds__(256, 6) MyKernel(...) { ... }

在限制每个 CUDA 线程的寄存器数量后，我的 GPU 使用率提高到了 60%。

顺便说一句，NSight v5.1 for Visual Studio 仍然支持 5xx 系列卡。 It can be downloaded 来自archive。

编辑：在使用来自多个 CPU 线程的多个 GPU 流的应用程序中，以下标志进一步将 GPU 使用率增加到 70%：

cudaSetDeviceFlags(cudaDeviceScheduleYield | cudaDeviceMapHost | cudaDeviceLmemResizeToMax);

cudaDeviceScheduleYield 让其他线程在 CPU 运行时执行线程正在等待 GPU 操作，而不是为结果。
cudaDeviceLmemResizeToMax，据我所知，是内核异步启动自身并避免过多的本地内存分配和解除分配。

【讨论】：