c++ - CUDA block VS SM 和线程 VS SP

Question

我们都知道GPGPU有几个stream multiprocesssors(SM)，说到它的硬件架构，每个都有很多stream processors(SP)。但是它在NVDIA的block编程模型中引入了另一个概念thread和CUDA。

而且我们也知道block对应SM和thread对应SP，当我们启动一个CUDA kernel时，我们配置内核为kernel<<<blockNum, threadsNum>>>。这样的CUDA程序我已经写了将近两个月了。但是我还有很多问题。优秀的程序员永远不会满足于没有错误的程序，他们想要深入研究并了解程序的行为方式。

以下问题:

1. 假设一个 GPU 有 14 个 SM，每个 SM 有 48 个 SP，我们有这样一个内核:

   __global__ void double(int *data, int dataNum){
        unsigned int tid = threadIdx.x + blockIdx.x * blockDim.x;
        while(tid < dataNum){
            data[tid] *= 2;
            tid += blockDim.x + blockIdx.x;
        }
   }
   

   而data是一个1024 * 1024 int数的数组，内核配置为<<<128, 512>>>，表示网格有512 * 128个线程，每个内核会迭代(1024 * 1024)/(512 * 128) = 在其 while 循环中执行了 16 次。 但是只有14 * 48个SP，也就是说无论你配置多少 block 号或线程号，都只能同时运行14 * 48个线程，请问配置中的blockNum和threadNum有什么意义, 为什么不只是 <<<number of SMs, number of SPs>>> 。

2. <<<128, 512>>> 和 <<<64, 512>>> 有什么区别吗，也许前者会在 while 循环中迭代 16 次，而字母是 32 次，但前者有双 block 要调度。有没有办法知道什么是最好的配置，而不只是比较结果然后选择最好的，因为我们无法尝试所有组合，所以结果不是完全最好的，而是你尝试中最好的。

   <

3. 我们知道一次只有一个 block 可以运行一个 SM，但是 CUDA 将其他 block 的上下文存储在哪里，假设 512 个 block 和 14 个 SM，只有 14 个 block 在 SM 中有它们的上下文，另一个呢？ 498 block 的来龙去脉？

Answer 1

And we also know that block corresponds to SM and thread corresponds to SP

这是不正确的。一个SM可以同时处理多个 block ，一个SP可以同时处理多个线程。

1) 我认为您的问题可能是由于没有将应用程序需要完成的工作与可用于完成该工作的资源分开。当你启动一个内核时，你指定了你想要完成的工作。然后 GPU 使用其资源来执行工作。 GPU 拥有的资源越多，它可以并行执行的工作就越多。任何不能并行完成的工作都可以串行完成。

通过让程序员指定需要完成的工作，而无需将其绑定(bind)到给定 GPU 上的可用资源量，CUDA 提供了一种抽象，允许应用程序无缝扩展到任何 GPU。

But there are only 14 * 48 SPs, which says that only 14 * 48 threads can simultaneously run no matter how many block numbers or thread numbers in your configuration

SP 是流水线式的，因此它们同时处理许多线程。每个线程都处于不同的完成阶段。每个 SP 可以开始一个操作并在每个时钟产生一个操作的结果。不过，如您现在所见，即使您的陈述是正确的，也不会得出您的结论。

2) block 中的线程可以使用共享内存相互协作。如果您的应用程序未使用共享内存，则 block 大小的唯一含义就是性能。最初，您可以使用占用计算器获得 block 大小的良好值。之后，您可以通过测试不同大小来进一步微调 block 大小以提高性能。由于线程以 32 个为一组运行，称为扭曲，您希望 block 大小可以被 32 整除。因此没有那么多的 block 大小可供测试。

3) 一个SM可以同时运行多个区 block 。 block 的数量取决于每个 block 需要多少资源以及SM拥有多少资源。一个 block 使用许多不同的资源，其中一种资源成为同时运行多少个 block 的限制因素。占用率计算器会告诉您限制因素是什么。

只有同时运行的 block 才会消耗 SM 上的资源。我认为这些资源就是您所说的上下文。已完成的 block 和尚未开始的 block 不会消耗 SM 上的资源。