关于 CUDA 中每个 SM 分配的寄存器数量

Question

第一个问题。 CUDA C 编程指南的编写如下。

相同的片上存储器用于 L1 和共享存储器：它可以 配置为 48 KB 的共享内存和 16 KB 的 L1 高速缓存或 16 KB 共享内存和 48 KB 一级缓存

但是，设备查询显示“每个块可用的寄存器总数：32768”。 我使用 GTX580。（CC 是 2.0） 该指南说默认缓存大小为 16KB，但 32768 表示 32768*4(byte) = 131072 Bytes = 128 KBytes。其实，我不知道哪个是正确的。

第二个问题。 我设置如下，

dim3    grid(32, 32);            //blocks in a grid
dim3    block(16, 16);           //threads in a block
kernel<<<grid,block>>>(...);

那么，每个块的线程数是 256。=> 每个块需要 256*N 个寄存器。 N 表示每个线程所需的寄存器数。 (256*N)*blocks 是每个 SM 的寄存器数。（不是字节） 因此，如果默认大小为 16KB 并且线程/SM 为 MAX(1536)，则 N 不能超过 2。因为“每个多处理器的最大线程数：1536”。 16KB/4Bytes = 4096 个寄存器，4096/1536 = 2.66666...

在更大的缓存 48KB 的情况下，N 不能超过 8。 48KB/4Bytes = 12288 个寄存器，12288/1536 = 8

这是真的吗？其实我很困惑。

---

实际上，我几乎完整的代码在这里。 我认为，当块尺寸为 16x16 时，内核得到了优化。 但是，在 8x8 的情况下，比 16x16 或类似的要快。 我不知道为什么。

每个线程的寄存器数为16，共享内存为80+16字节。

我问过同样的问题，但我无法得到确切的解决方案。： The result of an experiment different from CUDA Occupancy Calculator

#define WIDTH 512
#define HEIGHT 512
#define TILE_WIDTH 8
#define TILE_HEIGHT 8
#define CHANNELS 3
#define DEVICENUM 1 
#define HEIGHTs HEIGHT/DEVICENUM

__global__ void PRINT_POLYGON( unsigned char *IMAGEin, int *MEMin, char a, char b, char c){
        int Col = blockIdx.y*blockDim.y+ threadIdx.y;           //Col is y coordinate
        int Row = blockIdx.x*blockDim.x+ threadIdx.x;           //Row is x coordinate
        int tid_in_block = threadIdx.x + threadIdx.y*blockDim.x;
        int bid_in_grid = blockIdx.x + blockIdx.y*gridDim.x;
        int threads_per_block = blockDim.x * blockDim.y;
        int tid_in_grid = tid_in_block + threads_per_block * bid_in_grid;

        float result_a, result_b;
        __shared__ int M[15];
        for(int k = 0; k < 5; k++){
                M[k] = MEMin[a*5+k];
                M[k+5] = MEMin[b*5+k];
                M[k+10] = MEMin[c*5+k];
        }

        int result_a_up = (M[11]-M[1])*(Row-M[0]) - (M[10]-M[0])*(Col-M[1]);
        int result_b_up = (M[6] -M[1])*(M[0]-Row) - (M[5] -M[0])*(M[1]-Col);

        int result_down = (M[11]-M[1])*(M[5]-M[0]) - (M[6]-M[1])*(M[10]-M[0]);

        result_a = (float)result_a_up / (float)result_down;
        result_b = (float)result_b_up / (float)result_down;

        if((0 <= result_a && result_a <=1) && ((0 <= result_b && result_b <= 1)) && ((0 <= (result_a+result_b) && (result_a+result_b) <= 1))){
                IMAGEin[tid_in_grid*CHANNELS] += M[2] + (M[7]-M[2])*result_a + (M[12]-M[2])*result_b;      //Red Channel
                IMAGEin[tid_in_grid*CHANNELS+1] += M[3] + (M[8]-M[3])*result_a + (M[13]-M[3])*result_b;    //Green Channel
                IMAGEin[tid_in_grid*CHANNELS+2] += M[4] + (M[9]-M[4])*result_a + (M[14]-M[4])*result_b;    //Blue Channel
        }
}

struct DataStruct {
    int                 deviceID;
    unsigned char       IMAGE_SEG[WIDTH*HEIGHTs*CHANNELS];
};

void* routine( void *pvoidData ) { 
        DataStruct  *data = (DataStruct*)pvoidData;
        unsigned char *dev_IMAGE;
        int *dev_MEM;
        unsigned char *IMAGE_SEG = data->IMAGE_SEG;

        HANDLE_ERROR(cudaSetDevice(5));

        //initialize array
        memset(IMAGE_SEG, 0, WIDTH*HEIGHTs*CHANNELS);
        cudaDeviceSetCacheConfig(cudaFuncCachePreferL1);
        printf("Device %d Starting..\n", data->deviceID);

        //Evaluate Time
        cudaEvent_t start, stop;
        cudaEventCreate( &start );
        cudaEventCreate( &stop );

        cudaEventRecord(start, 0); 

        HANDLE_ERROR( cudaMalloc( (void **)&dev_MEM, sizeof(int)*35) );
        HANDLE_ERROR( cudaMalloc( (void **)&dev_IMAGE, sizeof(unsigned char)*WIDTH*HEIGHTs*CHANNELS) );

        cudaMemcpy(dev_MEM, MEM, sizeof(int)*35, cudaMemcpyHostToDevice);
        cudaMemset(dev_IMAGE, 0, sizeof(unsigned char)*WIDTH*HEIGHTs*CHANNELS);

        dim3    grid(WIDTH/TILE_WIDTH, HEIGHTs/TILE_HEIGHT);            //blocks in a grid
        dim3    block(TILE_WIDTH, TILE_HEIGHT);                         //threads in a block

        PRINT_POLYGON<<<grid,block>>>( dev_IMAGE, dev_MEM, 0, 1, 2);
        PRINT_POLYGON<<<grid,block>>>( dev_IMAGE, dev_MEM, 0, 2, 3);
        PRINT_POLYGON<<<grid,block>>>( dev_IMAGE, dev_MEM, 0, 3, 4);
        PRINT_POLYGON<<<grid,block>>>( dev_IMAGE, dev_MEM, 0, 4, 5);
        PRINT_POLYGON<<<grid,block>>>( dev_IMAGE, dev_MEM, 3, 2, 4);
        PRINT_POLYGON<<<grid,block>>>( dev_IMAGE, dev_MEM, 2, 6, 4);

        HANDLE_ERROR( cudaMemcpy( IMAGE_SEG, dev_IMAGE, sizeof(unsigned char)*WIDTH*HEIGHTs*CHANNELS, cudaMemcpyDeviceToHost ) );
        HANDLE_ERROR( cudaFree( dev_MEM ) );
        HANDLE_ERROR( cudaFree( dev_IMAGE ) );

        cudaEventRecord(stop, 0); 
        cudaEventSynchronize(stop);

        cudaEventElapsedTime( &elapsed_time_ms[data->deviceID], start, stop );
        cudaEventDestroy(start);
        cudaEventDestroy(stop);


        elapsed_time_ms[DEVICENUM] += elapsed_time_ms[data->deviceID];
        printf("Device %d Complete!\n", data->deviceID);

        return 0;
}

Answer 1

blockDim 8x8 比 16x16 快，因为当您增加块大小时，内存访问中的地址分歧会增加。

在具有 15 个 SM 的 GTX480 上收集的指标。

metric                         8x8         16x16
duration                        161µs       114µs
issued_ipc                     1.24        1.31
executed_ipc                    .88         .59
serialization                 54.61%      28.74%

指令重播的数量提示我们可能存在错误的内存访问模式。

achieved occupancy            88.32%      30.76%
0 warp schedulers issues       8.81%       7.98%
1 warp schedulers issues       2.36%      29.54%
2 warp schedulers issues      88.83%      52.44%

16x16 似乎使 warp 调度程序保持忙碌。但是，它使调度程序忙于重新发出指令。

l1 global load trans          524,407     332,007
l1 global store trans         401,224     209,139
l1 global load trans/request    3.56        2.25
l1 global store trans/request  16.33        8.51

首要任务是减少每个请求的事务。 Nsight VSE 源视图可以显示每条指令的内存统计信息。内核中的主要问题是 IMAGEin[] += 值的交错 U8 加载和存储。在 16x16 时，每个请求产生 16.3 个事务，但对于 8x8 配置只有 8.3 个。

改变 IMAGEin[(i*HEIGHTs+j)*CHANNELS] += ...

连续将 16x16 的性能提高 3 倍。我想将通道增加到 4 个并在内核中处理打包将提高缓存性能和内存吞吐量。

如果您固定每个请求的内存事务数量，您可能需要查看执行依赖关系并尝试增加 ILP。

Answer 2

块大小为 8x8 时速度更快，因为它是 32 的较小倍数，如下图所示，有 32 个 CUDA 内核绑定在一起，两个不同的 warp 调度程序实际上调度相同的事情。所以每个执行周期在这 32 个内核上执行相同的指令。

为了更好地阐明这一点，在第一种情况下（8x8），每个块由两个扭曲（64 个线程）组成，因此它仅在两个执行周期内完成，但是，当您使用 (16x16) 作为块大小时，每个都需要 8 个 warp（256 个线程），因此执行周期增加了 4 倍，从而导致复合速度变慢。

然而，在某些情况下，用更多的 warp 填充一个 SM 会更好，当内存访问量很高并且每个 warp 都可能进入内存停顿（即从内存中获取其操作数）时，它将被另一个 warp 替换直到内存操作完成。因此导致 SM 的占用率更高。

您当然应该在计算中加入每个 SM 的块数和 SM 总数，例如，将超过 8 个块分配给单个 SM 可能会降低其占用率，但在您的情况下，您可能不是面对这些问题，因为 256 通常比 64 更好，因为它会平衡 SM 之间的块，而使用 64 线程会导致在同一个 SM 中执行更多块。

编辑：此答案基于我的推测，有关更科学的方法，请参阅 Greg Smiths 的答案。

寄存器池不同于共享内存/缓存，在它们的架构的最底层！

寄存器由Flip-flops组成，L1缓存可能是SRAM。

只是想了解一下，看看下面代表 FERMI 架构的图片，然后更新您的问题以进一步说明您面临的问题。

作为注释，您可以通过将选项 --ptxas-options = -v 传递给 nvcc 来查看函数占用了多少寄存器和共享内存 (smem)。