启动多个内核（>1000）时 CUDA 流被阻塞

Question

我发现当我启动大量内核（超过 1000 个）时，CUDA 流会阻塞。我想知道有什么配置可以更改吗？

在我的实验中，我启动了一个小内核 10000 次。该内核运行时间很短（大约 190us）。启动前 1000 个内核时，内核启动速度非常快。启动内核需要 4~5us。但此后，启动过程变得缓慢。启动新内核大约需要 190us。 CUDA 流似乎等待前一个内核完成，缓冲区大小约为 1000 个内核。 当我创建 3 个流时，每个流可以启动 1000 个内核异步。

我想让这个缓冲区更大。我尝试设置

cudaLimitDevRuntimePendingLaunchCount

，但不起作用。有什么办法吗

#include <stdio.h>
#include "cuda_runtime.h"

#define CUDACHECK(cmd) do {                                   \
            cudaError_t e = cmd;                              \
            if (e != cudaSuccess) {                           \
                printf("Failed: Cuda error %s:%d '%s'\n",     \
                    __FILE__,__LINE__,cudaGetErrorString(e)); \
                exit(EXIT_FAILURE);                           \
            }                                                 \
        } while (0)

// a dummy kernel for test
__global__ void add(float *a, int n) {
    int id = threadIdx.x + blockIdx.x * blockDim.x;
    for (int i = 0; i < n; i++) {
        a[id] = sqrt(a[id] + 1);
    }
}

int main(int argc, char* argv[])
{
    // managing 1 devices
    int nDev = 1;
    int nStream = 1;
    int size = 32*1024*1024;

    // allocating and initializing device buffers
    float** buffer = (float**)malloc(nDev * sizeof(float*));
    cudaStream_t* s = (cudaStream_t*)malloc(sizeof(cudaStream_t)*nDev*nStream);

    for (int i = 0; i < nDev; ++i) {
        CUDACHECK(cudaSetDevice(i));
        // CUDACHECK(cudaDeviceSetLimit(cudaLimitDevRuntimePendingLaunchCount, 10000));
        CUDACHECK(cudaMalloc(buffer + i, size * sizeof(float)));
        CUDACHECK(cudaMemset(buffer[i], 1, size * sizeof(float)));
        for (int j = 0; j < nStream; j++) {
            CUDACHECK(cudaStreamCreate(s+i*nStream+j));
        }
    }

    for (int i = 0; i < nDev; ++i) {
        CUDACHECK(cudaSetDevice(i));
        for (int j=0; j < 10000; j++) {
            for (int k=0; k < nStream; k++) {
                add<<<32, 1024, 0, s[i*nStream+k]>>>(buffer[i], 1000);
            }
        }
    }

    for (int i = 0; i < nDev; ++i) {
        CUDACHECK(cudaSetDevice(i));
        cudaDeviceSynchronize();
    }

    // free device buffers
    for (int i = 0; i < nDev; ++i) {
        CUDACHECK(cudaSetDevice(i));
        CUDACHECK(cudaFree(buffer[i]));
    }

    printf("Success \n");
    return 0;
}

这是 nvprof 结果：

当我创建 3 个流时，前 3000 个内核启动很快，然后变慢

当我创建 1 个流时，前 1000 个内核启动很快，然后变慢

Answer 1

您所目睹的行为是预期行为。如果您在

cuda

标签上搜索“队列”或“启动队列”，您会发现许多其他涉及它的问题。 CUDA 有一个内核启动进入的队列（显然是每个流）。只要未完成的启动计数小于队列深度，启动过程就会是异步的。

然而，当未完成的（即未完成的）启动超过队列深度时，启动过程将变为一种同步行为（尽管不是通常意义上的同步）。具体来说，当未完成的内核启动次数超过队列深度时，启动进程将阻塞正在执行下一次启动的 CPU 线程，直到队列中打开启动槽（实际上意味着内核已在队列的另一端退出）队列）。

您无法了解这一点（无法查询队列中打开的插槽数量），也无法查看或控制队列深度。我在这里叙述的大部分信息都是通过检查获得的；据我所知，它尚未在 CUDA 文档中正式发布。

正如评论中已经讨论的那样，减轻您对多设备场景中启动的担忧的一种可能方法是启动广度优先而不是深度优先。我的意思是，您应该修改启动循环，以便在设备 0 上启动下一个内核之前，将内核启动到设备 0，然后是设备 1，然后是设备 2，依此类推。这将为您提供最佳性能。所有 GPU 将在启动序列中尽早参与处理。

如果您想查看 CUDA 行为或文档的更改，一般建议是成为developer.nvidia.com 上的注册开发人员，然后登录您的帐户并提交错误，使用可通过单击访问的错误归档流程来提交错误在右上角您的帐户名上。

Answer 2

如果我们仍然使用深度优先策略，但采用CPU多线程来为不同设备启动内核呢？这样能解决问题吗？