CUDA 流不是运行并行

Question

鉴于此代码：

void foo(cv::gpu::GpuMat const &src, cv::gpu::GpuMat *dst[], cv::Size const dst_size[], size_t numImages)
{
    cudaStream_t streams[numImages];
    for (size_t image = 0; image < numImages; ++image)
    {
        cudaStreamCreateWithFlags(&streams[image], cudaStreamNonBlocking);
        dim3 Threads(32, 16);
        dim3 Blocks((dst_size[image].width + Threads.x - 1)/Threads.x,
                    (dst_size[image].height + Threads.y - 1)/Threads.y);
        myKernel<<<Blocks, Threads, 0, streams[image]>>>(src, dst[image], dst_size[image]);
    }
    for (size_t image = 0; image < numImages; ++image)
    {
        cudaStreamSynchronize(streams[image]);
        cudaStreamDestroy(streams[image]);
    }
}

查看 nvvp 的输出，我看到几乎完美的串行执行，即使第一个流是一个冗长的过程，其他流应该能够与之重叠。

请注意，我的内核使用了 30 个寄存器，所有寄存器都报告了大约 0.87 的 "Achieved Occupancy"。对于最小的图像，网格大小为 [10,15,1]，块大小为 [32, 16,1]。

Answer 1

CUDA 编程指南 (link) 中给出了描述并发内核执行限制的条件，但其要点是您的 GPU 可能运行来自不同流的多个内核仅当您的 GPU 有足够的资源时。

在您的使用案例中，您说过您正在运行多次启动一个内核，其中包含 150 个块，每个块有 512 个线程。您的 GPU 有 12 个 SMM（我认为），并且每个 SMM 运行可以并发最多 4 个块（4 * 512 = 2048 个线程，这是 SMM 限制).所以你的 GPU 最多只能运行并发 4 * 12 = 48 个块。当命令管道中多次启动 150 个块时，似乎几乎没有（甚至没有）并发内核执行的机会。

如果您通过减小块大小来增加内核的调度粒度，您可能能够鼓励内核执行重叠。较小的块比较大的块更有可能找到可用资源和调度槽。同样，减少每个内核启动的总块数（可能通过增加每个线程的并行工作）可能也有助于增加多个内核重叠或并发执行的可能性。

CUDA 流不是运行并行

CUDA streams not running in parallel

c++

cuda

nvvp

CUDA 流不是 运行 并行

CUDA streams not running in parallel

c++

cuda

nvvp

CUDA 流不是运行并行