每个 SM 的最大居民块数？

Question

似乎每个 SM 允许的驻留块数有最大数量。但是，虽然其他“硬”限制很容易找到（例如，通过“cudaGetDeviceProperties”），但驻留块的最大数量似乎没有被广泛记录。

在以下示例代码中，我将内核配置为每个块一个线程。为了测试该 GPU（P100）每个 SM 最多有 32 个驻留块的假设，我创建了一个 56*32 块的网格（56 = P100 上的 SM 数量）。每个内核需要 1 秒来处理（通过“睡眠”例程），因此如果我正确配置了内核，代码应该需要 1 秒。计时结果证实了这一点。配置 32*56+1 块需要 2 秒，表明每个 SM 32 个块是每个 SM 允许的最大值。

我想知道的是，为什么这个限制没有得到更广泛的应用？例如，它不显示“cudaGetDeviceProperties”。在哪里可以找到各种 GPU 的此限制？或者这可能不是真正的限制，而是从其他硬限制派生出来的？

我正在运行 CUDA 10.1

#include <stdio.h>
#include <sys/time.h>

double cpuSecond() {
    struct timeval tp;
    gettimeofday(&tp,NULL);
    return (double) tp.tv_sec + (double)tp.tv_usec*1e-6;
}

#define CLOCK_RATE 1328500  /* Modify from below */
__device__ void sleep(float t) {    
    clock_t t0 = clock64();
    clock_t t1 = t0;
    while ((t1 - t0)/(CLOCK_RATE*1000.0f) < t)
        t1 = clock64();
}

__global__ void mykernel() {
    sleep(1.0);    
}

int main(int argc, char* argv[]) {
    cudaDeviceProp  prop;
    cudaGetDeviceProperties(&prop, 0); 
    int mp = prop.multiProcessorCount;
    //clock_t clock_rate = prop.clockRate;

    int num_blocks = atoi(argv[1]);

    dim3 block(1);
    dim3 grid(num_blocks);  /* N blocks */

    double start = cpuSecond();
    mykernel<<<grid,block>>>();
    cudaDeviceSynchronize();
    double etime = cpuSecond() - start;

    printf("mp          %10d\n",mp);
    printf("blocks/SM   %10.2f\n",num_blocks/((double)mp));
    printf("time        %10.2f\n",etime);

    cudaDeviceReset();
}

结果 ：

% srun -p gpuq sm_short 1792
mp                  56
blocks/SM        32.00
time              1.16

% srun -p gpuq sm_short 1793
mp                  56
blocks/SM        32.02
time              2.16

% srun -p gpuq sm_short 3584
mp                  56
blocks/SM        64.00
time              2.16

% srun -p gpuq sm_short 3585
mp                  56
blocks/SM        64.02
time              3.16

Answer 1

是的，每个 SM 的块数是有限制的。SM中可以包含的最大块数是指给定时间内活动块的最大数量。块可以组织成一维或二维网格，每个维度最多包含 65,535 个块，但 GPU 的 SM 将只能容纳一定数量的块。此限制通过两种方式与 Gpu 的计算能力相关。

CUDA 规定的硬件限制。

每个 GPU 允许每个 SM 块的最大限制，无论其包含的线程数量和使用的资源量如何。例如，计算能力为 2.0 的 Gpu 的限制为 8 Blocks/SM，而计算能力为 7.0 的 Gpu 的限制为 32 Blocks/SM。这是每个 SM 可以实现的最佳活动块数：我们将其称为 MAX_BLOCKS。

来自每个块使用的资源量的限制。

一个块由线程组成，每个线程使用一定数量的寄存器：它使用的寄存器越多，包含它的块使用的资源数量就越多。类似地，分配给块的共享内存量会增加该块需要分配的资源量。一旦超过某个值，一个块所需的资源数量将非常大，以至于 SM 将无法分配 MAX_BLOCKS 所允许的那么多块：这意味着每个块所需的资源数量是有限的每个 SM 的最大活动块数。

我如何找到这些边界？

CUDA也考虑到了这一点。他们的网站上提供了Cuda Occupancy Calculator 文件，您可以使用该文件发现按计算能力分组的硬件限制。您还可以输入块使用的资源量（线程数、每个线程的寄存器、共享内存的字节），并获取有关活动块数的图表和重要信息。链接文件的第一个选项卡允许您根据使用的资源计算 SM 的实际使用情况。如果您想知道每个线程使用多少个寄存器，则必须添加 -Xptxas -v 选项，以便编译器告诉您在创建 PTX 时它使用了多少个寄存器。在文件的最后一个选项卡中，您将找到按计算能力分组的硬件限制。