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

Question

我是 CUDA 新手，需要帮助理解一些事情。我需要帮助并行化这两个 for 循环。具体如何设置 dimBlock 和 dimGrid 以使其运行更快。我知道这看起来像 sdk 中的 vector 添加示例，但该示例仅适用于方阵，当我尝试为我的 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。我已阅读指南，但仍然不明白如何获得正确的编号。网格/ block /线程的数量并有效地使用它们。

Answer 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 个 block ，每个包含 512 个线程将被启动到 GPU 硬件上，以并行执行数组加法操作。请注意，如果输入数据大小不能表示为 block 大小的整数倍，则需要将 block 数向上取整以覆盖整个输入数据集。

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