我熟悉使用nvprof访问基准的事件和指标,例如
nvprof --system-profiling on --print-gpu-trace -o (file name) --events inst_issued1 ./benchmarkname
system-profiling on --print-gpu-trace -o (filename)
命令给出了开始时间、内核结束时间、电源、临时时间的时间戳,并将信息保存在NVVP文件中,以便我们可以在可视化分析器中查看它。这使我们能够看到代码的任何部分发生了什么,特别是当一个特定的内核正在运行时。我的问题是——
是否有一种方法可以隔离仅在基准测试运行的一段时间内计数的事件,例如在内核执行期间?在上面的命令中,
--events inst_issued1
只给出整个可执行文件的指令计数。谢谢!
您可能需要阅读一下分析器文档。
您可以在可执行文件中打开和关闭分析。cuda运行时API为:
cudaProfilerStart()
cudaProfilerStop()
因此,如果您只想收集特定内核的概要信息,您可以这样做:
#include <cuda_profiler_api.h>
...
cudaProfilerStart();
myKernel<<<...>>>(...);
cudaProfilerStop();
(上面可以是调用内核的函数或代码,而不是内核调用)节选自文档:
当使用start和stop函数时,您还需要指示分析工具在应用程序开始时禁用分析。对于nvprof,您可以使用——profile-from-start - off标志来完成此操作。对于可视化分析器,您可以在设置视图中使用启用了分析的开始执行复选框。
同样从nvprof
的文档中,您可以使用命令行开关将事件/度量制表限制为单个内核:
--kernels <kernel name>
文档给出了更多的用法可能性。
同样的方法也适用于insight系统和insight compute。
CUDA分析器启动/停止功能的工作方式完全相同。insight系统文档解释了如何使用剖析器api管理的捕获控制来运行剖析器:
nsys [global-options] start -c cudaProfilerApi
或for insight compute:
ncu [options] --profile-from-start off
同样,可以通过命令行限制insight计算,仅配置特定的内核。此操作的主要开关是-k
,按名称选择内核。在重复的情况下,-c
开关可用于确定要进行分析的已命名内核的启动次数,-s
开关可用于在进行分析之前跳过一些启动。
这些方法不仅适用于事件和度量,而且适用于由各自的分析器执行的所有分析活动。
CUDA分析器API可以在任何可执行文件中使用,并且不需要使用nvcc
进行编译。
在对此进行了更深入的研究之后,发现通过使用
还可以为所有内核提供内核级信息(w/o使用--kernels
并具体指定它们)nvprof --events <event names> --metrics <metric names> ./<cuda benchmark>
实际上,它给出的输出形式是
"设备"、"内核"、"调用"、"活动名称","分","马克斯"、"Avg"
如果在基准测试中多次调用内核,这允许您查看这些内核运行所需事件的最小、最大和平均值。显然,Cuda 7.5 Profiler上的--kernels
选项允许指定每个内核的每次运行。