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

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 实际上可以启动比指定的 block /网格维度更多的线程吗？

Answer 1

CUDA 不会启动比 block /网格维度指定的线程更多的线程。

但是，由于 block 维度的粒度(例如，希望 block 维度是 32 的倍数，并且大小限制为 1024 或 512)，通常情况下很难匹配线程网格在数值上等于所需的问题大小。

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

在你的例子中，这可以通过写作来澄清:

__global__ void myKernel(int problem_size)


    if (threadId < problem_size)

它传达了预期的内容，即只有与问题大小(可能与启动的网格大小不匹配)相对应的线程才能执行任何实际工作。

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

如果我希望每个线程 block 都是 32 的倍数，则没有可供我选择的线程 block 数量，这将恰好为我提供 10000 个线程。因此，我可能会在一个线程 block 中选择 256 个线程，并启动 40 个线程 block ，总共有 10240 个线程。使用线程检查，我阻止了“额外的”240 个线程执行任何操作。