主页 / coding / 问题 / 78827628
Accepted
Iceberglet
Asked: 2024-08-03 08:59:39 +0800 CST

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

  • 772

我是 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. 根据块数而变化的性能。我读到网格/块只是开发人员的观点，没有影响（特别是对于没有共享/减少的完全隔离的线程）。那么为什么这很重要，以及如何选择最佳块大小？
cuda

1 个回答

  1. Best Answer

    您没有显示相关代码或提供 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 上执行。因此，256 个线程和 1024 个块大小意味着所有线程都在同一处理器上执行，而其他处理器上的计算资源则处于闲置状态。

    总的来说，我认为这个指标毫无意义。100-1000 个线程太少了，您需要查看所有线程的吞吐量，即工作项数除以内核总执行时间。

    根据块数而变化的性能。我读到网格/块只是开发人员的观点，没有影响（特别是对于没有共享/减少的完全隔离的线程）。那么为什么这很重要，以及如何选择最佳块大小？

    错了。块大小确实很重要。同样，如果您查看计算能力，每个 SM 的块数以及每个 SM 的线程数都有限制。如果您的块大小低于 64，则不会达到 100% 的占用率。块大小不是 warp 大小的倍数也会因停用线程而浪费资源。当然，在 CUDA 8.6-8.9 设备上，每个 SM 有 1536 个线程，1024 的块大小永远不会达到超过 2/3 的占用率。

    块也有启动开销，因此较少的块可能是有益的，但块太大也会产生负面影响。只有在旧块的所有线程完成后，新块才能启动。大块意味着在内核的末尾（或在__syncthreads()屏障处），许多线程可以等待很少的落后者，占用资源的时间比必要的时间更长。

    根据经验，每个块使用大约 128-256 个线程。如果需要，可以对不同大小进行基准测试。但您需要使其成为有意义的基准。查看整体吞吐量并充分利用 GPU。

    • 2

相关问题