我应该检查内核代码中的线程数吗？

Question

我是 CUDA 的初学者，我的同事总是使用以下包装来设计内核：

__global__ void myKernel(int nbThreads)
{
    int threadId = blockDim.x*blockIdx.y*gridDim.x  //rows preceeding current row in grid
            + blockDim.x*blockIdx.x             //blocks preceeding current block
            + threadIdx.x;

    if (threadId < nbThreads)
    {
        statement();
        statement();
        statement();
    }
}

他们认为在某些情况下，CUDA 可能会出于对齐/扭曲的原因启动比指定的线程更多的线程，因此我们每次都需要检查它。然而，到目前为止，我在互联网上还没有看到他们实际进行此验证的示例内核。

CUDA 实际上可以启动比指定的块/网格尺寸更多的线程吗？

Answer 1

CUDA 不会启动比块/网格尺寸指定的线程更多的线程。

然而，由于块尺寸的粒度（例如，希望块尺寸是32的倍数，并且其大小限制为1024或512），经常出现难以匹配线程网格的情况在数值上等于所需的问题大小。

在这些情况下，典型的行为是启动更多线程，根据块粒度有效舍入到下一个偶数大小，并使用内核中的“线程检查”代码来确保“额外线程”，即那些超出问题规模，不要做任何事情。

在您的示例中，可以通过编写以下内容来澄清这一点：

__global__ void myKernel(int problem_size)


    if (threadId < problem_size)

它传达了意图，即只有与问题大小相对应的线程（可能与启动的网格大小不匹配）才执行任何实际工作。

作为一个非常简单的示例，假设我想对长度为 10000 个元素的向量进行向量加法。10000 不是 32 的倍数，也不小于 1024，因此在典型的实现中，我会启动多个线程块来完成工作。

如果我希望每个线程块都是 32 的倍数，那么我无法选择能够提供 10000 个线程的线程块数量。因此，我可能会在一个线程块中选择 256 个线程，并启动 40 个线程块，总共有 10240 个线程。使用线程检查，我可以防止“额外”的 240 个线程执行任何操作。