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

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

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

以下问题:

假设一个 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 整数数组，内核配置为 <<>>，这意味着网格有 512 * 128 个线程，每个内核将迭代 (1024 * 1024)/(512 * 128) = 16 次在它的 while 循环中。 但是只有 14 * 48 个 SP，也就是说不管你的配置中有多少块号或线程号，只有 14 * 48 个线程可以同时运行 ，配置中的 blockNum 和 threadNum 有什么意义，为什么不只是 <<<number of SMs, number of SPs>>> .

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

我们知道只有一个块可以运行一次 SM，但是 CUDA 在哪里存储其他块的上下文，假设 512 个块和 14 个 SM，只有 14 个块在 SM 中有它们的上下文，那么其他 498 个块的上下文呢？

最佳答案

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

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

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

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

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

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

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

关于c++ - CUDA 块 VS SM 和线程 VS SP，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/20768545/

sm

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