c++ - CUDA流未并行运行

给出以下代码:

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。对于最小的图像，网格大小为[10,15,1]，块大小为[32，16,1]。

最佳答案

CUDA编程指南(link)中给出了描述并发内核执行限制的条件，但要点是，只有在GPU具有足够资源来运行GPU的情况下，GPU才可能潜在地运行来自不同流的多个内核。

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

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

关于c++ - CUDA流未并行运行，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/34847798/