c++ - 在 CUDA 中管理占用

Question

早上好。

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

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

-warps per SM(系统多处理器) - 每个 SM block -每个SM注册 -每个SM的共享内存

因此，我首先要处理的事情是，根据 GPU，我根据每个 SM 的最大扭曲和每个 SM 的 block 定义了内核的维度。我的任务是执行一亿次求和来衡量哪种方法更好。

我所做的是一个 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，每个 warp 有 32 个线程。在这个 GPU 中，一个事件 block 有 2 个扭曲，这意味着每个 block 可以同时有 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);

您将根据需要启动尽可能多的 block 和每个 block 的线程以完全加载一个 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,);

您正在启动大量 block ，这允许 GPU 并行执行所需数量的 block 以达到 100% 的占用率(资源允许，这在简单 vector 加法的情况下应该不是问题)。因此性能更好。

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

附带说明一下， vector 加法作为内存限制算法并不是特别适合演示占用的影响。然而，您仍然看到差异，因为需要一定的最小内存事务数量才能完全加载内存子系统(由内存带宽乘以内存访问延迟决定)，5% 的占用率示例未达到此最小值.