CUDA - 没有块，只是未定义尺寸的线程-Java 学习之路

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

我设计了一个带有（x; y）块和（x; y）线程的CUDA内核 .

由于矩阵的宽度/高度不是我的维度的倍数，因此让事情起作用是一种可怕的痛苦，并且代码变得越来越复杂以获得合并内存读取 .

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

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

我可以消除块并使用大量线程还是必须进行块细分？

2 回答

2
你实际上不能只创建一个“线程网格”，因为你必须将线程组织成块，每个块最多可以有512个线程 . 但是，您可以通过每个块使用1个线程来有效地执行此操作，这将导致1 x 1块的X x Y网格 . 但是，由于以下几个因素，这将导致非常糟糕的性能：
- 根据CUDA编程指南，SM可以随时处理最多8个块 . 这将限制每个SM 8个线程，这甚至不足以填充单个warp . 如果你有48个CUDA核心，你将只能在任何给定的时间处理384个线程 .
- 在SM上只有8个线程可用，将会有太少的warp来隐藏内存延迟 . GPU将花费大部分时间等待内存访问完成，而不是进行任何计算 .
- 您将无法合并内存读取和写入，从而导致内存带宽使用不足 .
- 您实际上无法利用共享内存，因为这是块中线程之间的共享资源 .
虽然必须确保块中线程的正确性是令人讨厌的，但是您的性能将比您的“线程网格”理念要好得多 .
回复于 2024-04-29T10:36:37+08:00

这是我用来将需要 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
    ...
}

我不认为这需要很多努力/注册/代码行来检查是否正确 .

回复于 2024-04-29T10:36:37+08:00

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

2 回答

相关问题