CUDA 流未并行运行答案

【问题标题】：CUDA streams not running in parallelCUDA 流未并行运行
【发布时间】：2016-01-18 05:25:38
【问题描述】：

鉴于此代码：

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”。对于最小的图像，Grid Size 为 [10,15,1]，Block Size 为 [32, 16,1]。

【问题讨论】：

每个内核启动 150 个块。您的 GPU 必须有足够的可用资源才能安排来自另一个流的块。可以？这里有什么问题？
我有一台 GTX 980M，每个多处理器有 2048 个线程，每个块有 1024 个线程。我是否创建了如此多的线程或块，以至于一次只能运行一个内核 - 这就是 87% 的占用率告诉我的吗？
据我估计，您的 GPU 最多可以在每个 SMM 上运行 4 个块，而您的 GPU 中大约有 12 个 SMM。因此，您的并发块限制为 48。如果您的内核启动是 150 个块，则并非单个内核启动中的所有块都可以同时启动。多个内核启动同时运行的范围相当有限

标签： c++ cuda nvvp

【解决方案1】：

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

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

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

【讨论】：

感谢您的解释；即使是最大的图像，我的实际代码也只需要 400 μs 即可运行，并且我确实看到使用 nvvp 时有轻微的重叠（大约 2 或 3 μs），这可能是由于某些 SMM 的完成速度略快于其他 SMM。我现在正在开发另一个更复杂的内核，稍后将移植到更强大的卡上，所以如果我在那里遇到性能问题，这个背景资料会有所帮助。
@keny-n：如果这个答案对你有用，你能接受吗？除其他外，它使它们在搜索中更加可见