在 CUDA 中管理占用率

Question

早安。

我开始学习 cuda 编程，我正在学习性能。我在 CUDA 网站上读到，为了获得良好的性能，我们应该考虑四件事：

http://docs.nvidia.com/gameworks/content/developertools/desktop/analysis/report/cudaexperiments/kernellevel/achievedoccupancy.htm

-warps per SM（系统多处理器） - 每个 SM 块 - 注册每个 SM - 每个 SM 共享内存

所以我首先要回顾一下，根据 GPU，我根据每个 SM 的最大扭曲和每个 SM 的块来定义内核的尺寸。我的任务是执行一亿次求和来衡量哪种方法更好。

我所做的是一个 for 循环，在该循环中，我在每次迭代时启动一个内核以最大化占用率。例如，对于 NVidia 1080 GPU，我读到：

int max_blocks = 32; //maximum number of active blocks per SM int max_threads_per_Block = 64; //maximum number of active threads per SM int max_threads = 2048;

这为每个 SM 提供了总共 2048 个线程，并保证了最大占用率。这个 GPU 可以有 64 个活动 warp，每个有 32 个线程。在这个 GPU 中，一个活动块有 2 个扭曲，这意味着每个块可以同时拥有 64 个活动线程。有了这个我启动内核如下：

dim3 threadsPerBlock(max_threads_per_Block); dim3 numBlocks(max_blocks); VecAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C,max_threads);

令我惊讶的是，如果我像这样直接启动这个内核：

int N = total_ops; //in this case one thousand millions dim3 threadsPerBlock(256); dim3 numBlocks(2*N / threadsPerBlock.x); VecAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C,);

性能更好（消耗时间）。我在相同的执行中启动了 5 次相同的实验以避免异常值。我的问题是：有没有办法管理占用以获得比编译器和运行时 API 更好的结果？我知道我尝试做的优化已经以某种方式由 GPU 管理。我知道，如果有一个文档解释了我们应该如何启动软件（上面的链接）以实现良好的性能，它应该是一种控制它的方法。

谢谢

Answer 1

在你的第一个例子中，

int max_blocks = 32;            //maximum number of active blocks per SM
int max_threads_per_Block = 64; //maximum number of active threads per SM
int max_threads = 2048;

dim3 threadsPerBlock(max_threads_per_Block);
dim3 numBlocks(max_blocks);
VecAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C,max_threads);

您将根据需要为每个块启动尽可能多的块和线程以完全加载 一个 SM。但是您的 GTX 1080 有 20 个 SM，因此您的占用率只有 1/20 = 5%。

在第二个例子中，

int N = total_ops;              //in this case one thousand millions
dim3 threadsPerBlock(256);
dim3 numBlocks(2*N / threadsPerBlock.x);
VecAdd<<<numBlocks, threadsPerBlock>>>(d_A, d_B, d_C,);

您正在启动大量块，这使 GPU 可以根据需要并行执行尽可能多的块，以达到 100% 的占用率（资源允许，这在简单的向量添加的情况下应该不是问题）。因此性能更好。

虽然您可以在第一个示例中将块数乘以 20 以获得与第二个示例相同的性能，但首选第二个示例中的模式，因为它不涉及所使用 GPU 的特定配置。因此，代码将完全加载大量 GPU 中的任何一个。

附带说明，作为内存绑定算法的向量加法并不是特别适合演示占用的影响。但是，您仍然会看到差异，因为完全加载内存子系统（由内存带宽乘以内存访问的延迟确定）需要一定的最小内存事务数量，并且 5% 的占用率示例低于此最小值.