NVidia Visual Profiler 的分析器开销过多答案

【问题标题】：Excessive profiler overhead with NVidia Visual ProfilerNVidia Visual Profiler 的分析器开销过多
【发布时间】：2017-10-23 23:53:20
【问题描述】：

尝试使用nvvp（或nvprof）分析我的代码时，我得到了很多分析开销：

总时间为 98 毫秒，我在第一次内核启动时获得了 85 毫秒的“仪器”。

如何减少这种分析开销或仅放大我感兴趣的部分？

背景

我在未选中“启用性能分析的情况下开始执行”的情况下运行此程序，并且我使用 cudaProfilerStart/cudaProfilerStop 限制了性能分析，如下所示：

/* --- generate data etc --- */
// Call the function once to warm up the FFT plan cache
applyConvolution( T, N, stride, plans, yData, phiW, fData, y_dwt );
gpuErrchk( cudaDeviceSynchronize() );
// Call it once for profiling
cudaProfilerStart();
applyConvolution( T, N, stride, plans, yData, phiW, fData, y_dwt );
gpuErrchk( cudaDeviceSynchronize() );
cudaProfilerStop();

applyConvolution() 是我正在分析的函数。

我在带有 GTX 1080 的 Ubuntu 16.04 上使用 CUDA Toolkit 8.0。

【问题讨论】：

标签： nvvp

【解决方案1】：

在我写这个问题的时候，我想我会尝试弄乱分析器设置来尝试抢占一些潜在的评论回答材料。

令我惊讶的是，禁用“启用并发内核分析”完全消除了分析器开销：

但也许这并不令人感到意外：

启用并发内核分析 - 应选择此选项对于使用 CUDA 流启动内核的应用程序，可以同时执行。如果应用程序仅使用单个流（因此不能有并发内核执行），取消选择此选项可能会减少分析开销。

（取自http://docs.nvidia.com/cuda/profiler-users-guide/）

CUDA Profiler 用户指南的早期版本也在“Profiling Limitations”部分中指出：

如果在以下内核上使用并发内核模式会增加大量开销执行大量块且执行时间短。

哦，好吧。无论如何发布这个问题/答案，以防它帮助其他人避免这种烦恼。

【讨论】：

【解决方案2】：

我看到类似的东西，但可能只是模糊相关。但由于上述答案有帮助，我将添加我的观察。

在对 Quadro GV100 进行性能分析时，与 pascal-gen 卡（例如 1080）相比，相当简单的内核的明显性能发生了巨大变化。我也在运行 nvvp 并禁用分析并在我感兴趣的部分代码中激活它。然后我不小心忽略了打开它，我得到的只是我们的手动事件标记（使用 nvtxRangePush 和 nvtxRangePop）。你知道吗，十倍加速。也就是说; 在 Quadro GV100 上存在大量的分析开销，这是早期 GPU 所没有的。

像您一样禁用并发分析没有有帮助，但禁用 API 跟踪可以。

虽然与手动 nvtx 相比仍有很大的开销，但至少它允许对 GV100 上的内核性能有所了解。较大的内核似乎受到的影响较小，如果它与固定成本开销或 API 跟踪有关，这是很自然的。剩下的未知数就是为什么 API 跟踪在 GV100 上的成本如此之高，但我无法推测，至少现在还没有。

我为上述测试使用 gcc/5.4 和 cuda/9.0 编译了特定于 sm 的二进制文件，并为一个简单的测试用例运行了 RELION 单线程。

【讨论】：