每个 SM 的最大驻留块数?

Posted

技术标签:

【中文标题】每个 SM 的最大驻留块数?【英文标题】:Maximum number of resident blocks per SM? 【发布时间】:2020-08-06 06:45:54 【问题描述】:

似乎每个 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

【问题讨论】:

docs.nvidia.com/cuda/cuda-c-programming-guide/… @talonmies 谢谢 - 这是我第一次看到硬限制。我想知道为什么 Nvidia 不通过cudaGetDeviceProperties 提供这个号码。 因为它大多是无关紧要的。这就是为什么 仅供参考,从 CUDA 11 开始,此限制公开为 cudaDevAttrMaxBlocksPerMultiprocessorcudaDeviceProp 中的 maxBlocksPerMultiProcessor。在此处查看最新文档:docs.nvidia.com/cuda/cuda-runtime-api/…。 【参考方案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 file,您可以通过它发现按计算能力分组的硬件限制。您还可以输入块使用的资源量(线程数、每个线程的寄存器、共享内存的字节数)并获取有关活动块数的图表和重要信息。 链接文件的第一个选项卡允许您根据使用的资源计算 SM 的实际使用情况。如果你想知道你使用的每个线程有多少个寄存器,你必须添加 -Xptxas -v 选项让编译器告诉你它在创建 PTX 时使用了多少个寄存器。 在文件的最后一个选项卡中,您将找到按计算能力分组的硬件限制。

【讨论】:

我了解限制块数的因素有很多。我在询问有关硬限制的更多信息,这似乎没有被广泛记录(尽管上面的链接清楚地表明存在限制)。

以上是关于每个 SM 的最大驻留块数?的主要内容,如果未能解决你的问题,请参考以下文章

以编程方式检索每个多处理器的最大块数

你能以编程方式知道 GPU 中每个块的最大块数和线程数吗?

CUDA:网格中的最大块数!= CU_DEVICE_ATTRIBUTE_MAX_GRID_DIM_X?

2020杭电多校第二场题解

表空间支持的最大数据文件大小的算法数据库限制数据文件文件头保留数据块数

51nod1053&&1052 最大M子段和