CUDA - 没有块，只有未定义尺寸的线程答案

【问题标题】：CUDA - no blocks, just threads for undefined dimensionsCUDA - 没有块，只有未定义尺寸的线程
【发布时间】：2011-07-25 15:48:55
【问题描述】：

我有一些未知大小的矩阵在两个方向上从 10 到 20.000 不等。

我设计了一个带有 (x;y) 块和 (x;y) 线程的 CUDA 内核。

由于矩阵的宽度/高度不是我的维度的倍数，所以让事情工作起来非常痛苦，而且代码变得越来越复杂以获取合并内存读取。

除此之外，内核的大小还在增长，使用越来越多的寄存器来检查正确性......所以我认为这不是我应该采用的方式。

我的问题是：如果我完全消除块并创建一个 x;y 线程网格怎么办？ SM单元没有很多块会不会有问题？

我可以消除块并使用大量线程还是需要块细分？

【问题讨论】：

标签： c++ matrix cuda

【解决方案1】：

您实际上不能只创建“线程网格”，因为您必须将线程组织成块，并且每个块最多可以有 512 个线程。但是，您可以通过每个块使用 1 个线程来有效地做到这一点，这将导致 1x1 块的 X x Y 网格。但是，由于以下几个因素，这将导致非常糟糕的性能：

根据 CUDA 编程指南，一个 SM 在任何时候最多可以处理 8 个块。这会将每个 SM 限制为 8 个线程，这甚至不足以填充单个经线。如果您有 48 个 CUDA 内核，那么您在任何给定时间只能处理 384 个线程。
由于 SM 上只有 8 个线程可用，因此扭曲太少，无法隐藏内存延迟。 GPU 将花费大部分时间等待内存访问完成，而不是进行任何计算。
您将无法合并内存读取和写入，从而导致内存带宽使用率降低。
您将无法有效利用共享内存，因为这是块中线程之间的共享资源。

虽然必须确保块中线程的正确性很烦人，但您的性能将大大优于“线程网格”的想法。

【讨论】：

@Paul：检查边界只有一行设备代码。
这不是我的情况，我正在研究 nvidia SDK 的合并可分离卷积滤波器，顺便说一句，我会看看我能做什么

【解决方案2】：

这是我用来将需要num_threads 的给定任务划分为块和网格的代码。是的，您最终可能会启动许多块（但只有极少数），并且您最终可能会拥有比所需更多的实际线程，但这种方式既简单又高效。请参阅下面的第二个代码示例，了解我的简单内核边界检查。

PS：我一直有block_size == 128，因为它在多核占用、寄存器使用、共享内存要求和我所有内核的合并访问之间取得了很好的折衷。

计算良好网格大小（主机）的代码：

#define GRID_SIZE 65535

//calculate grid size (store result in grid/block)
void kernelUtilCalcGridSize(unsigned int num_threads, unsigned int block_size, dim3* grid, dim3* block) {


    //block
    block->x = block_size;
    block->y = 1;
    block->z = 1;


    //number of blocks
    unsigned int num_blocks = kernelUtilCeilDiv(num_threads, block_size);
    unsigned int total_threads = num_blocks * block_size;
    assert(total_threads >= num_threads);

    //calculate grid size
    unsigned int gy = kernelUtilCeilDiv(num_blocks, GRID_SIZE);
    unsigned int gx = kernelUtilCeilDiv(num_blocks, gy);
    unsigned int total_blocks = gx * gy;
    assert(total_blocks >= num_blocks);

    //grid
    grid->x = gx;
    grid->y = gy;
    grid->z = 1;
}

//ceil division (rounding up)
unsigned int kernelUtilCeilDiv(unsigned int numerator, unsigned int denominator) {
    return (numerator + denominator - 1) / denominator;
}

计算唯一线程ID并检查边界（设备）的代码：

//some kernel
__global__ void kernelFoo(unsigned int num_threads, ...) {


    //calculate unique id
    const unsigned int thread_id = threadIdx.x;
    const unsigned int block_id = blockIdx.x + blockIdx.y * gridDim.x;
    const unsigned int unique_id = thread_id + block_id * blockDim.x;


    //check range
    if (unique_id >= num_threads) return;

    //do the actual work
    ...
}

我不认为检查正确性需要花费很多精力/寄存器/代码行。

【讨论】：