CUDA如何获取网格、块、线程大小和并行化非方阵计算

Posted 2023-04-15

【中文标题】CUDA如何获取网格、块、线程大小和并行化非方阵计算

【英文标题】：CUDA how to get grid, block, thread size and parallalize non square matrix calculation

【发布时间】：2011-08-04 07:56:56

【问题描述】：

我是 CUDA 的新手，需要帮助理解一些事情。我需要帮助并行化这两个 for 循环。具体如何设置 dimBlock 和 dimGrid 以使其运行更快。我知道这看起来像 sdk 中的向量添加示例，但该示例仅适用于方阵，当我尝试为我的 128 x 1024 矩阵修改该代码时，它无法正常工作。

__global__ void mAdd(float* A, float* B, float* C)

    for(int i = 0; i < 128; i++)
    
        for(int j = 0; j < 1024; j++)
        
            C[i * 1024 + j] = A[i * 1024 + j] + B[i * 1024 + j];
        
    

这段代码是更大循环的一部分，也是代码中最简单的部分，所以我决定尝试并行化 thia 并同时学习 CUDA。我已阅读指南，但仍然不明白如何获得正确的编号。网格/块/线程的数量并有效地使用它们。

【问题讨论】：

在pycuda 中只是C[i] = A[i] + B[i] demo.py

【参考方案1】：

正如您所写，该内核是完全串行的。每个启动执行它的线程都将执行相同的工作。

CUDA（以及 OpenCL 和其他类似的“单程序、多数据”类型编程模型）背后的主要思想是您采用“数据并行”操作 - 因此必须执行许多相同的、很大程度上独立的操作次 - 并编写一个执行该操作的内核。然后启动大量（半）自治线程以跨输入数据集执行该操作。

在你的数组加法示例中，数据并行操作是

C[k] = A[k] + B[k];

对于0到128 * 1024之间的所有k。每个加法操作是完全独立的，没有顺序要求，因此可以由不同的线程执行。为了在 CUDA 中表达这一点，可以这样编写内核：

__global__ void mAdd(float* A, float* B, float* C, int n)

    int k = threadIdx.x + blockIdx.x * blockDim.x;

    if (k < n)
        C[k] = A[k] + B[k];

[免责声明：代码在浏览器中编写，未经测试，使用风险自负]

这里，串行代码中的内循环和外循环每次操作都被一个 CUDA 线程替换，并且我在代码中添加了限制检查，以便在启动的线程多于所需操作的情况下，不会出现缓冲区溢出发生。如果内核是这样启动的：

const int n = 128 * 1024;
int blocksize = 512; // value usually chosen by tuning and hardware constraints
int nblocks = n / blocksize; // value determine by block size and total work

madd<<<nblocks,blocksize>>>mAdd(A,B,C,n);

然后 256 个块，每个包含 512 个线程将被启动到 GPU 硬件上，以并行执行数组加法操作。请注意，如果输入数据大小不能表示为块大小的整数倍，则需要将块数向上取整以覆盖整个输入数据集。

以上所有内容都是对 CUDA 范式的一个非常简单的操作的非常简化的概述，但也许它为您提供了足够的洞察力来继续自己。如今，CUDA 相当成熟，网上有很多免费的优质教育材料，您可能可以使用它来进一步阐明我在此答案中所掩盖的编程模型的许多方面。

【讨论】：

int k = threadIdx.x + gridDim.x * blockDim.x;这肯定是不正确的？在您的示例中，gridDim.x * blockDim.x 将始终为 256*512。应该是 int k = threadIdx.x + blockIdx.x * blockDim.x;我试图编辑它但被拒绝了。

对浏览器的警告：nblocks = ceil(n / nthreads); // 如果你的数据没有完美分割。

@ofer.sheffer：我确实写过“请注意，如果输入数据大小不能表示为块大小的整数倍，则需要将块数四舍五入以覆盖完整的输入数据集。”。这还不够清楚吗？

@talonmies，您的回答非常好，我赞成。另一方面，当我阅读它时，我在想“他错过了+1”，以防数据分布不均……然后我继续阅读了其他一些内容，然后回到这里完成阅读我注意到你写了它。作为一个通常只是先看代码并考虑稍后阅读每个单词的略读读者——我认为我的警告会对我未来的自己有所帮助。

我怎么知道nthreads？ blocksize不是线程数吗？