在 CUDA 中管理占用

Question

早上好。

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

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

-每个 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 个活动扭曲，每个扭曲有 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% 占用率示例低于此最小值.