NPP支持重叠流吗？

Question

我试图使用NVIDIA的NPP库方法nppiFilterBorder_32f_C1R_Ctx，在一个具有多个过滤器的图像上执行多个异步2D卷积。然而，即使在创建多个流并将它们分配给NPPI的方法之后，重叠也不会发生；NVIDIA的nvvp也通知了相同的消息：

​

​

也就是说，我很困惑NPP是否支持重叠的上下文操作。

下面是我的代码的简化，只显示异步方法调用和相关变量：

std::vector<NppStreamContext> streams(n_filters);

for(size_t stream_idx=0; stream_idx<n_filters; stream_idx++)
{
  cudaStreamCreateWithFlags(&(streams[stream_idx].hStream), cudaStreamNonBlocking);
  streams[stream_idx].nStreamFlags = cudaStreamNonBlocking;
  // fill up NppStreamContext remaining fields 
  // malloc image and filter pointers
}

for(size_t stream_idx=0; stream_idx<n_filters; stream_idx++)
{
  cudaMemcpyAsync(..., streams[stream_idx].hStream);
  nppiFilterBorder_32f_C1R_Ctx(..., streams[stream_idx]);
  cudaMemcpy2DAsync(..., streams[stream_idx].hStream);
}

for(size_t stream_idx=0; stream_idx<n_filters; stream_idx++)
{
  cudaStreamSynchronize(streams[stream_idx].hStream);
  cudaStreamDestroy(streams[stream_idx].hStream);
}

注意：输出图像和输入过滤器的所有设备指针都存储在std::vector中，在std::vector中，我通过当前流索引(例如，float *ptr_filter_d = filters[stream_idx])访问它们。

Answer 1

总结并添加注释:配置文件确实显示了少量的重叠，因此标题问题的答案显然是肯定的。

重叠如此小的原因仅仅是每个NPP内核在其大部分运行时已经需要使用的GPU的所有资源。在每个内核的末尾，人们可能会看到尾效应(即块的数量不是每个时刻可以驻留在SMs中的块数的倍数)，因此来自下一个内核的块正在被调度，并且有一些重叠。

有时，强制大内核(首先启动并使用完整的设备)和随后只需要少量资源的小内核之间的重叠(即优化)是很有用的。在这种情况下，可以通过cudaStreamCreateWithPriority使用流优先级来提示调度程序在第一个内核的块之前对第二个内核的块进行调度。这方面的一个例子可以在这个多GPU示例(permalink)中找到。

然而，在这种情况下，由于内核的大小是相同的，并且没有理由将它们中的任何一个优先于其他内核，因此强制这样的重叠不会减少整个运行时，因为计算资源是有限的。在分析器视图中，内核可能会显示更多的重叠，但每一个都需要更多的时间。这就是为什么调度程序不重叠内核的原因，即使您允许它使用多个流(参见异步流还是并行流)。

为了提高性能，可以编写一个定制的CUDA内核，在一个内核启动时执行所有过滤器。在这种情况下，这可能比使用NPP更好的主要原因是所有NPP内核都采用相同的输入映像。因此，一个内核可以通过只读取一次输入映像的每个块(对共享内存，尽管L1缓存可能足够)，显着地减少对全局内存的访问次数，然后依次或并行地应用所有过滤器(通过将线程块拆分成更小的单元)并写出结果。