Profiler 用户指南

NVIDIA性能分析工具用户手册，用于优化CUDA应用程序的性能。

性能分析概览

本文档介绍NVIDIA的性能分析工具，帮助您理解和优化CUDA、OpenACC或OpenMP应用程序的性能。Visual Profiler是一款图形化分析工具，可显示应用程序CPU和GPU活动的时间线，并包含自动分析引擎以识别优化机会。nvprof分析工具使您能够从命令行收集和查看性能分析数据。

请注意，Visual Profiler和nvprof已被弃用，将在未来的CUDA版本中移除。 NVIDIA Volta平台是这些工具获得完整支持的最后一代架构。建议使用新一代工具NVIDIA Nsight Systems进行GPU和CPU采样与追踪，以及NVIDIA Nsight Compute进行GPU内核性能分析。

更多详情请参阅从Visual Profiler和nvprof迁移到Nsight工具部分。

术语

一个事件是指设备上可计数的活动、动作或发生的情况。它对应内核执行期间收集的单个硬件计数器值。要查看特定NVIDIA GPU上所有可用事件的列表，请输入nvprof --query-events。

指标是应用程序的一个特征，通过一个或多个事件值计算得出。要查看特定NVIDIA GPU上所有可用指标的列表，请输入nvprof --query-metrics。您也可以参考指标参考文档。

1. 为性能分析准备应用程序

CUDA分析工具无需对应用程序进行任何更改即可启用分析功能；然而，通过进行一些简单的修改和补充，您可以显著提高分析的实用性和有效性。本节将介绍这些修改内容以及它们如何帮助提升您的分析结果。

1.1. 聚焦性能分析

默认情况下，性能分析工具会收集整个应用程序运行期间的性能数据。但如下文所述，通常您只需要分析应用程序中包含部分或全部性能关键代码的区域。将分析范围限制在性能关键区域可以减少您和工具需要处理的数据量，并将注意力集中在通过优化能带来最大性能提升的代码上。

在以下几种常见情况下，对应用程序的某个区域进行分析会很有帮助。

1. 该应用程序是一个测试框架，包含您算法的全部或部分CUDA实现。测试框架会初始化数据，调用CUDA函数执行算法，然后检查结果的正确性。使用测试框架是一种常见且高效的方法，可以快速迭代和测试算法变更。在进行性能分析时，您需要收集实现算法的CUDA函数的性能数据，而不是收集初始化数据或检查结果的测试框架代码的性能数据。

2. 应用程序分阶段运行，每个阶段激活不同的算法集。当应用程序每个阶段的性能可以独立于其他阶段进行优化时，您需要单独分析每个阶段以集中优化工作。

3. 该应用程序包含在大量迭代上运行的算法，但算法在这些迭代中的性能变化不大。在这种情况下，您可以从迭代的子集中收集性能分析数据。

为了将性能分析限制在应用程序的特定区域，CUDA提供了用于启动和停止性能数据收集的函数。cudaProfilerStart()用于启动性能分析，cudaProfilerStop()用于停止性能分析（使用CUDA驱动API时，您可以通过cuProfilerStart()和cuProfilerStop()获得相同的功能）。要使用这些函数，您必须包含cuda_profiler_api.h（对于驱动API则是cudaProfiler.h）。

在使用启动和停止功能时，您还需要指示分析工具在应用程序开始时禁用分析功能。对于nvprof，您需要使用--profile-from-start off标志来实现这一点。对于Visual Profiler，您需要在Settings View中使用"启用分析执行"复选框。

1.2. 标记CPU活动区域

Visual Profiler 可以收集应用程序调用的 CUDA 函数追踪记录。Visual Profiler 会在时间线视图中显示这些调用，让您查看应用程序中每个 CPU 线程调用 CUDA 函数的位置。要了解应用程序的 CPU 线程在 CUDA 函数调用之外执行的操作，您可以使用NVIDIA Tools Extension API (NVTX)。当您在应用程序中添加 NVTX 标记和范围时，时间线视图会显示 CPU 线程在这些区域内执行的时间。

nvprof 也支持NVTX标记和范围。标记和范围会在时间轴的API跟踪输出中显示。在摘要模式下，每个范围会显示与该范围关联的CUDA活动。

1.3. 命名CPU与CUDA资源

Visual Profiler的时间线视图会显示CPU线程和GPU设备、上下文及流的默认命名。为这些资源使用自定义名称可以更好地理解应用程序行为，特别是对于拥有大量主机线程、设备、上下文或流的CUDA应用程序。您可以使用NVIDIA Tools Extension API为CPU和GPU资源分配自定义名称。您的自定义名称随后将显示在时间线视图中。

nvprof 同样支持NVTX命名功能。在摘要和跟踪模式下会显示CUDA设备、上下文和流的名称。线程名称则会在摘要模式下显示。

1.4. 刷新配置文件数据

为了减少性能分析的开销，分析工具会将性能信息收集并记录到内部缓冲区中。这些缓冲区随后会以低优先级异步刷新到磁盘，以避免干扰应用程序行为。为防止尚未刷新的性能信息丢失，被分析的应用程序应在退出前确保所有GPU工作已完成（使用CUDA同步调用），然后调用cudaProfilerStop()或cuProfilerStop()。这样做会强制刷新对应上下文中的缓冲性能信息。

如果你的CUDA应用程序包含使用显示器或主循环运行的图形，必须注意在执行该循环的线程调用exit()之前调用cudaProfilerStop()或cuProfilerStop()。如果未能调用这些API之一，可能会导致部分或全部收集的性能分析数据丢失。

对于某些图形应用程序（例如使用OpenGL的应用），按下退出键时程序会直接退出。在这些情况下，如果无法在退出前调用上述函数，请使用nvprof工具的--timeout选项，或在Visual Profiler中设置"执行超时"。分析器将在超时前强制进行数据刷新。

1.5. 分析CUDA Fortran应用程序

使用PGI CUDA Fortran编译器编译的CUDA Fortran应用程序可以通过nvprof和Visual Profiler进行分析。当分析器需要源文件和行号信息时（如内核性能分析、全局内存访问模式分析、分支执行分析等），在编译时请使用“-Mcuda=lineinfo”选项。该选项在PGI 2019版本19.1及以后的Linux 64位目标平台上受支持。

2. 可视化性能分析器

NVIDIA Visual Profiler 可帮助您可视化并优化应用程序的性能。该工具会显示应用程序在CPU和GPU上的活动时间轴，使您能够发现性能提升机会。此外，Visual Profiler还会分析您的应用程序以检测潜在性能瓶颈，并指导您采取行动消除或减轻这些瓶颈。

Visual Profiler既可作为独立应用程序使用，也可作为Nsight Eclipse Edition的一部分。独立版本的Visual Profiler nvvp包含在所有支持操作系统的CUDA工具包中。在Nsight Eclipse Edition中，Visual Profiler位于性能分析透视图中，当应用程序在性能分析模式下运行时会被激活。

2.1. 入门指南

本节介绍开始性能分析时可能采取的步骤。

2.1.1. 配置Java运行环境

Visual Profiler 要求本地系统安装有 Java Runtime Environment (JRE) 1.8。但从 CUDA Toolkit 10.1 Update 2 版本开始，由于 Oracle 升级许可变更，JRE 不再包含在 CUDA Toolkit 中。用户必须安装所需版本的 JRE 1.8 才能使用 Visual Profiler。详见Installing JRE。

• 在OpenSUSE15或SLES15上运行Visual Profiler：

  • 请确保按如下方式在命令行中包含选项来调用Visual Profiler：

    nvvp -vm /usr/lib64/jvm/jre-1.8.0/bin/java
    

    注意

    仅当JRE 1.8不在默认路径中时，才需要使用-vm选项。

• 在Ubuntu 18.04或Ubuntu 18.10上运行Visual Profiler：

  • 请确保按如下方式在命令行中包含选项来调用Visual Profiler：

    nvvp -vm /usr/lib/jvm/java-8-openjdk-amd64/jre/bin/java
    

    注意

    仅当JRE 1.8不在默认路径中时，才需要使用-vm选项。

  • 在Ubuntu 18.10系统上，如果在运行Visual Profiler时遇到错误"no swt-pi-gtk in java.library.path"，则需要安装GTK2。输入以下命令来安装所需的GTK2。

    apt-get install libgtk2.0-0
    

• 在Fedora 29上运行Visual Profiler：

  • 请确保按如下方式在命令行中包含选项来启动Visual Profiler：

    nvvp -vm /usr/bin/java
    

    注意

    -vm选项仅在JRE 1.8不在默认路径时才需要指定。

• 在Windows上运行Visual Profiler：

  • 请确保按如下方式在命令行中包含选项来调用Visual Profiler：

    nvvp -vm "C:\Program Files\Java\jdk1.8.0_77\jre\bin\java"
    

    注意

    仅当JRE 1.8不在默认路径中时，才需要使用-vm选项。

2.1.2. 安装JRE

Visual Profiler 要求本地系统安装 Java Runtime Environment (JRE) 1.8。但从 CUDA Toolkit 10.1 Update 2 版本开始，由于 Oracle 升级许可变更，CUDA Toolkit 不再包含 JRE。用户必须自行安装 JRE 1.8 才能使用 Visual Profiler。请参阅下方可用选项。另见Java Platform, Standard Edition 8 Names and Versions.

Windows

• Oracle JRE 1.8 (可能需要付费更新)

• OpenJDK JRE 1.8

Linux

• Oracle JRE 1.8 (可能需要付费更新)

• OpenJDK JRE 1.8

  Note

  JRE 1.8u152 or later is not supported for Visual Profiler. You can find the JRE update 151 on the Oracle Download Archive site here: https://www.oracle.com/technetwork/java/javase/downloads/java-archive-javase8-2177648.html?printOnly=1.

2.1.3. 为性能分析修改您的应用程序

Visual Profiler无需对应用程序进行任何更改；然而，通过一些简单的修改和补充，您可以显著提高其可用性和有效性。Preparing An Application For Profiling章节描述了如何集中分析工作，并为应用程序添加额外注释，从而极大改善分析体验。

2.1.4. 创建会话

使用Visual Profiler分析应用程序的第一步是创建一个新的分析会话。会话包含与您的应用程序相关的设置、数据和结果。有关会话操作的更多信息，请参阅Sessions部分。

您可以通过在欢迎页面选择"分析应用程序"链接，或在文件菜单中选择"新建会话"来创建新会话。在"创建新会话"对话框中输入您的应用程序可执行文件。您还可以选择性地指定工作目录、参数、多进程分析选项以及环境变量。

多进程分析选项包括：

• 分析子进程 - 如果选中，将分析指定应用程序启动的所有进程。

• 分析所有进程 - 如果选中，将分析同一用户在同一系统上启动的每个CUDA进程（该用户需启动nvprof）。在此模式下，Visual Profiler会启动nvprof，用户需要在Visual Profiler之外的另一终端中运行其应用程序。用户可通过点击Visual Profiler进度对话框中的"取消"按钮退出此模式并加载分析数据

• 仅分析当前进程 - 如果选中，则仅分析指定的应用程序。

点击下一步选择一些额外的性能分析选项。

CUDA选项：

• 启用性能分析启动执行 - 如果选中，将从应用程序执行开始时收集性能分析数据。如果未选中，则在应用程序中调用cudaProfilerStart()之前不会收集性能分析数据。有关cudaProfilerStart()的更多信息，请参阅Focused Profiling。

• 启用并发内核分析 - 对于使用CUDA流来启动可并行执行内核的应用程序，应选择此选项。如果应用程序仅使用单个流（因此无法实现内核并发执行），取消此选项可能会降低分析开销。

• 在时间轴中启用CUDA API跟踪 - 如果选中，将收集并显示CUDA驱动程序和运行时API调用的跟踪信息。

• 启用功耗、时钟和温度分析 - 如果选中，将采样并显示GPU上的功耗、时钟和温度状况。并非所有GPU都支持收集此数据。有关更多信息，请参阅时间线视图中的设备时间线描述。

• 启用统一内存分析 - 如果为支持统一内存的GPU选中此选项，系统将收集与统一内存相关的进出每个GPU的内存流量，并在时间线上显示。

• 重放应用程序以收集事件和指标 - 如果选中，将重新运行整个应用程序而非仅重放每个内核，以便收集所有事件/指标。

• 运行引导式分析 - 如果勾选此项，将在创建新会话后立即运行引导式分析。取消勾选此选项可禁用该行为。

CPU（主机）选项：

• 在CPU上分析执行情况 - 如果选中，将对CPU线程进行采样，并在CPU详情视图中显示收集到的CPU性能数据。

• 启用OpenACC性能分析 - 如果选中并分析OpenACC应用程序，将记录OpenACC活动并在新的OpenACC时间线上显示。此数据收集仅在Linux和PGI 19.1或更高版本上受支持。有关OpenACC时间线的更多信息，请参阅时间线视图中的描述。

• 启用CPU线程追踪 - 如果启用，选中的CPU线程API调用将被记录并显示在新的线程API时间线上。当前包括Pthread API、互斥锁和条件变量。出于性能考虑，仅记录影响并发执行的API调用，且Windows系统不支持收集此数据。更多信息请参阅时间线视图中对线程时间线的描述。对于使用CUDA的多CPU线程应用程序进行依赖分析时，应选择此选项。

时间轴选项：

• 加载时间范围内的数据 - 如果选中此选项，可以指定要加载数据范围的开始和结束时间戳。此选项对于从大型数据中选择子集非常有用。

• 在会话中启用时间线 - 默认情况下所有时间线均已启用。如果取消勾选某个时间线，则与该时间线关联的数据将不会被加载，也不会显示。

注意

如果某些时间线通过取消勾选选项被禁用，那么使用这些时间线数据的分析结果将不正确。

点击完成。

2.1.5. 分析您的应用程序

如果在创建会话时未选择"不运行引导分析"选项，Visual Profiler将立即运行您的应用程序以收集引导分析第一阶段所需的数据。如分析视图部分所述，您可以使用引导分析系统获取有关应用程序中性能限制行为的建议。

2.1.6. 探索时间线

除了引导式分析结果外，您还将看到应用程序执行时显示CPU和GPU活动的时间线。阅读时间线视图和属性视图了解如何探索时间线中可用的性能分析信息。导航时间线描述了如何缩放和滚动时间线以聚焦应用程序的特定区域。

2.1.7. 查看详情

除了Analysis View中提供的结果外，您还可以查看分析过程中收集的具体指标和事件值。指标和事件值显示在GPU Details View中。您可以收集特定的指标和事件值，这些值能揭示应用程序中内核的运行情况。按照GPU Details View章节所述的方法即可收集指标和事件。

2.1.8. 优化大型配置文件的加载

某些应用程序会启动大量小型内核，导致即使仅运行几秒钟也可能产生非常庞大（数百兆字节或更大）的输出数据。Visual Profiler需要与所打开/导入的分析文件大小大致相同的内存空间。若未指定"最大堆大小"设置，Java虚拟机可能仅占用主内存的一小部分。因此根据主内存容量，Visual Profiler可能无法加载某些大型文件。

如果Visual Profiler无法加载大型配置文件，请尝试根据主内存大小设置JVM允许使用的最大堆大小。您可以修改工具包安装目录中的配置文件libnvvp/nvvp.ini。nvvp.ini配置文件内容如下：

-startup
plugins/org.eclipse.equinox.launcher_1.3.0.v20140415-2008.jar
--launcher.library
plugins/org.eclipse.equinox.launcher.gtk.linux.x86_64_1.1.200.v20140603-1326
-data
@user.home/nvvp_workspace
-vm
../jre/bin/java
-vmargs
-Dorg.eclipse.swt.browser.DefaultType=mozilla

例如，要强制JVM使用3GB内存，可在‑vmargs后添加新行‑Xmx3G。-Xmx设置应根据可用系统内存和输入大小进行调整。例如，如果您的系统有24GB内存，并且您确定在运行Visual Profiler时不需要同时运行其他内存密集型应用程序，那么分析器占用绝大部分内存空间是可以接受的。因此您可以选择22GB作为最大堆大小，为操作系统、图形界面和其他可能运行的程序保留几GB空间。

其他一些nvvp.ini配置设置也可以修改：

• 将默认堆大小（Java自动启动时使用的堆大小）增加到2GB。（-Xms）

• 让Java以64位模式运行，而非默认的32位模式（仅适用于64位系统）；若需要堆内存大小超过4GB，则必须启用此选项。(-d64)

• Enable Javas parallel garbage collection system, which helps both to decrease the required memory space for a given input size as well as to catch out of memory errors more gracefully. (-XX:+UseConcMarkSweepGC -XX:+CMSIncrementalMode)

注意：大多数安装需要管理员/root权限才能修改此文件。

根据上述示例修改后的nvvp.ini文件如下：

-data
@user.home/nvvp_workspace
-vm
../jre/bin/java
-d64
-vmargs
-Xms2g
-Xmx22g
-XX:+UseConcMarkSweepGC
-XX:+CMSIncrementalMode
-Dorg.eclipse.swt.browser.DefaultType=Mozilla

有关JVM设置的更多详情，请参阅Java虚拟机手册。

除此之外，您还可以使用时间轴选项"加载指定时间范围内的数据"以及在创建会话部分提到的"启用时间轴"功能，来限制加载和显示的数据范围。

2.2. 会话

会话包含与您的应用程序相关的设置、数据和性能分析结果。每个会话都保存在单独的文件中；因此您可以通过简单地删除、移动、复制或共享会话文件来删除、移动、复制或共享会话。按照惯例，文件扩展名.nvvp用于Visual Profiler会话文件。

会话分为两种类型：一种是可执行会话，它与应用程序相关联，在Visual Profiler内部执行和分析；另一种是导入会话，通过导入由nvprof生成的数据创建。

2.2.1. 可执行会话

您可以通过在欢迎页面选择"分析应用程序"链接，或在文件菜单中选择"新建会话"来为您的应用程序创建新的可执行会话。会话创建后，您可以按照设置视图中的说明编辑会话设置。

您可以使用文件菜单中的打开和保存选项来打开和保存现有会话。

为了分析您的应用程序并收集指标和事件值，Visual Profiler将多次执行您的应用程序。要获得准确的性能分析结果，您的应用程序必须符合应用程序要求中详述的条件。

2.2.2. 导入会话

您可以通过文件菜单中的"导入..."选项，从nvprof的输出结果创建导入会话。选择此选项将打开导入对话框，引导您完成整个导入流程。

由于可执行应用程序未与导入会话关联，Visual Profiler无法执行该应用程序以收集额外的性能分析数据。因此，只能基于已导入的数据进行分析。此外，GPU Details View将显示所有已导入的事件和指标值，但无法为导入会话选择和收集新的指标和事件。

2.2.2.1. 导入单进程 nvprof 会话

通过导入对话框，您可以选择一个或多个nvprof数据文件导入到新会话中。

您必须拥有一个包含会话时间线信息的nvprof数据文件。该数据文件应通过使用--export-profile选项运行nvprof来收集。您可以选择启用其他选项，如--system-profiling on，但不应收集任何事件或指标，因为这会使时间线失真，无法真实反映应用程序的行为。

您可以选择性地指定一个或多个包含应用程序事件和指标值的事件/指标数据文件。这些数据文件应通过运行nvprof并使用--events和/或--metrics选项来收集。要收集分析系统所需的所有事件和指标，您可以简单地使用--analysis-metrics选项以及--kernels选项来选择要收集事件和指标的内核。更多信息请参阅远程性能分析。

如果您正在将多个nvprof输出文件导入会话中，请确保您的应用程序符合应用程序要求中详述的条件。

2.2.2.2. 导入多进程 nvprof 会话

使用导入向导，您可以选择多个nvprof数据文件导入到新的多进程会话中。

每个nvprof数据文件必须包含其中一个进程的时间线信息。该数据文件应通过使用--export-profile选项运行nvprof来收集。您可以选择启用其他选项，例如--system-profiling on，但不应收集任何事件或指标，因为这会扭曲时间线，使其无法代表应用程序的真实行为。

在导入nvprof数据对话框中，选择多进程选项，如下图所示。

从多个进程导入时间线数据时，您可能未为这些进程指定任何事件/指标数据文件。多进程分析仅支持时间线数据。

2.2.2.3. 导入命令行性能分析会话

已不再支持命令行分析器（使用环境变量COMPUTE_PROFILE），但早期版本生成的CSV文件仍可导入。

使用导入向导，您可以选择一个或多个由命令行分析器生成的CSV文件导入到新会话中。当您导入多个CSV文件时，它们的内容将被合并并在单一时间轴中显示。

命令行分析器的CSV文件必须包含gpustarttimestamp和streamid配置参数生成。也可以包含其他配置参数，包括事件。

2.3. 应用需求

要收集应用程序的性能数据，Visual Profiler必须能够以确定性的方式重复执行您的应用程序。由于软件和硬件的限制，无法在单次应用程序执行中收集所有必要的性能分析数据。每次运行应用程序时，它必须处理相同的数据，并以相同的顺序执行相同的内核和内存复制调用。具体来说，

• 对于某个设备，每次应用程序执行时，上下文创建的顺序必须保持一致。对于每个线程创建自己上下文的多线程应用程序，必须注意确保这些上下文创建的顺序在多次运行中保持一致。例如，可能需要先在单个线程上创建上下文，然后将这些上下文传递给其他线程。或者，可以使用NVIDIA Tools Extension API为每个上下文提供自定义名称。只要在应用程序每次执行时对相同的上下文应用相同的自定义名称，Visual Profiler就能正确关联多次运行中的这些上下文。

• 对于某个上下文，每次应用程序执行时，流的创建顺序必须保持一致。或者，可以使用NVIDIA Tools Extension API为每个流提供自定义名称。只要在每次应用程序执行时对相同的流应用相同的自定义名称，Visual Profiler就能正确关联多次运行中的这些流。

• 在同一个流中，每次应用程序执行时，内核调用和内存拷贝调用的顺序必须保持一致。

2.4. 可视化性能分析器视图

可视化性能分析器按视图组织。这些视图共同帮助您分析和可视化应用程序的性能表现。本节将介绍每个视图的功能，以及在进行应用程序性能分析时如何使用它们。

2.4.1. 时间线视图

时间轴视图显示应用程序在性能分析期间发生的CPU和GPU活动。在Visual Profiler中可以同时在不同标签页中打开多个时间轴。下图展示了一个CUDA应用程序的时间轴视图。

视图顶部有一个水平标尺，显示从应用程序性能分析开始经过的时间。视图左侧有一个垂直标尺，描述时间轴每一行显示的内容，并包含时间轴的各种控制选项。这些控制在时间轴控制中有详细说明。

时间轴视图由多个时间轴行组成。每行显示代表该行对应类型活动开始和结束时间的区间。例如，代表内核的时间轴行会显示该内核执行开始和结束时间的区间。在某些情况下（如下所述），一个时间轴行可以显示多个活动子行。当存在重叠活动时会使用子行。这些子行会根据活动重叠程度动态创建。区间在特定子行中的位置不传达任何特定含义。区间只是通过一种启发式算法打包到子行中，该算法试图最小化所需的子行数量。子行的高度会进行缩放以保持合理的垂直空间。

时间轴视图中显示的时间轴行类型包括：

Process

时间轴将为每个被分析的应用程序包含一个进程行。进程标识符代表该进程的pid。进程的时间轴行不包含任何活动区间。进程内的线程显示为该进程的子项。

Thread

时间轴将为分析应用程序中执行过CUDA驱动或CUDA运行时API调用的每个CPU线程显示一个线程行。线程标识符是该CPU线程的唯一ID。线程的时间轴行不包含任何活动区间。

Runtime API

时间轴会为每个执行CUDA Runtime API调用的CPU线程显示一个Runtime API行。行中的每个区间代表对应线程上调用所持续的时间。

Driver API

时间轴会为每个执行CUDA Driver API调用的CPU线程显示一个Driver API行。行中的每个区间代表对应线程上调用所持续的时间。

OpenACC

时间轴会为每个调用OpenACC指令的CPU线程包含一个或多个OpenACC行。行中的每个区间表示对应线程上调用的持续时间。每个OpenACC时间轴可能由多行组成。在同一时间轴内，下方行的OpenACC活动是从上方行的活动中调用的。

OpenMP

时间轴会为每个调用OpenMP的CPU线程显示一行OpenMP记录。该行中的每个区间表示应用程序在特定OpenMP区域或状态中花费的时间。应用程序可能同时处于多个状态，这种情况会通过绘制多行显示，其中某些区间会相互重叠。

Pthread

时间轴将为每个执行Pthread API调用的CPU线程包含一行Pthread记录，前提是在测量期间记录了主机线程API调用。该行中的每个区间代表调用的持续时间。请注意，出于性能考虑，可能仅记录了部分选定的Pthread API调用。

Markers and Ranges

时间轴会为每个使用NVIDIA Tools Extension API标注时间范围或标记的CPU线程显示单独的标记与范围行。行中的每个区间代表一个时间范围的持续时间，或标记的瞬时点。如果存在重叠范围，该行将包含子行。

Profiling Overhead

时间线将为每个进程包含一个单独的"性能分析开销"行。该行中的每个区间表示执行某些性能分析所需活动的持续时间。这些区间代表在应用程序未被分析时不会发生的活动。

Device

时间轴将为被分析的应用程序所使用的每个GPU设备显示一个设备行。时间轴行的名称以方括号中的设备ID开头，后接设备名称。运行计算利用率分析后，该行将包含设备随时间变化的计算利用率估计值。如果启用了功耗、时钟和温度分析功能，该行还会显示代表这些读数的数据点。

Unified Memory

时间轴将为每个使用统一内存的CPU线程和设备显示一个统一内存行。统一内存行可能包含CPU页错误、GPU页错误、数据迁移(DtoH)和数据迁移(HtoD)行。创建会话时，用户可以为统一内存时间轴选择分段模式或非分段模式。在分段模式下，时间轴被划分为等宽的段，仅显示每个时间段内的聚合数据值。段数可以调整。在非分段模式下，时间轴上的每个区间将代表实际采集的数据，并可查看每个区间的属性。这些段使用热力图配色方案进行着色。在时间轴属性下，会显示用于选择颜色的属性，同时图例会展示颜色与不同属性值范围的对应关系。

CPU Page Faults

这将为每个CPU线程包含一个CPU页面错误行。在非分段模式下，时间轴上的每个间隔对应一个CPU页面错误。

Data Migration (DtoH)

时间轴将为每个设备包含一条数据迁移(DtoH)记录。在非分段模式下，时间轴上的每个间隔对应一次从设备到主机的数据迁移。

GPU Page Faults

时间轴将包含每个CPU线程的GPU页面错误行。在非分段模式下，时间轴上的每个间隔对应一个GPU页面错误组。

Data Migration (DtoH)

时间轴会为每个设备包含一个数据迁移(HtoD)行。在非分段模式下，时间轴上的每个间隔对应一次从主机到设备的数据迁移。

Context

时间轴会为GPU设备上的每个CUDA上下文包含一个Context行。时间轴行的名称表示上下文ID，如果使用了NVIDIA Tools Extension API来命名上下文，则显示自定义上下文名称。上下文行不包含任何活动间隔。

Memcpy

时间轴会为每个执行内存拷贝(memcpy)的上下文包含内存拷贝行。一个上下文最多可包含四行内存拷贝，分别对应设备到主机、主机到设备、设备到设备以及点对点内存拷贝。每行中的每个区间表示在GPU上执行的内存拷贝持续时间。

Compute

时间轴会为每个在GPU上执行计算的上下文包含一个计算行。每行中的每个区间代表GPU设备上内核的执行时长。计算行展示了该上下文的所有计算活动。当在上下文中并行执行多个内核时，会使用子行来显示。所有内核活动（包括使用CUDA动态并行性启动的内核）都会显示在计算行上。计算行之后的内核行会显示每个独立应用程序内核的活动情况。

Kernel

时间轴会为应用程序执行的每个内核显示一个Kernel行。每行中的每个区间代表该内核实例在所属上下文中的执行时长。每行都标有一个百分比，表示该内核所有实例的执行时间占所有内核总执行时间的比例。对于每个上下文，内核按此执行时间百分比从上到下排序。子行用于显示并发内核执行情况。对于使用CUDA动态并行技术的应用程序，内核会按父子关系以层级结构组织。由主机启动的内核显示为Context行的直接子项。使用CUDA动态并行技术启动其他内核的内核可通过"+"图标展开，以显示代表这些子内核的行。对于不启动子内核的内核，内核执行由实心区间表示，显示该内核实例在GPU上的执行时间。对于启动子内核的内核，区间末尾还可能包含空心部分。空心部分表示内核执行完成后等待子内核执行结束的时间。CUDA动态并行执行模型要求父内核必须等待所有子内核完成后才能结束，这正是空心部分所表示的。Timeline Controls中描述的Focus控件可用于控制父子时间轴的显示。

Stream

时间轴会为应用程序使用的每个流（包括默认流和任何应用程序创建的流）包含一个Stream行。Stream行中的每个区间表示在该流上执行的memcpy或内核运行的持续时间。

2.4.1.1. 时间轴控制

时间线视图提供了多个控件，可用于控制时间线的显示方式。其中部分控件还会影响GPU详情视图和分析视图中的数据呈现。

调整垂直时间标尺大小

垂直标尺的宽度可以通过将鼠标指针悬停在标尺右侧边缘进行调整。当出现双箭头指针时，按住鼠标左键并拖动。垂直标尺宽度会随会话保存。

重新排序时间线

内核(Kernel)和流(Stream)时间轴行可以重新排序。您可能需要重新排列这些行以帮助可视化相关的内核和流，或者将不重要的内核和流移动到时间轴底部。要重新排序某一行，请左键点击并按住行标签。当出现双箭头指针时，向上或向下拖动以定位该行。时间轴排序会随您的会话保存。

时间线筛选

Memcpy和Kernel行可以过滤，以排除它们在GPU详情视图和分析视图中的显示活动。要过滤某一行，请左键点击行标签左侧的过滤器图标。要过滤所有Kernel或Memcpy行，请按住Shift键并左键点击其中一行。当某行被过滤时，该行上的所有间隔都会变暗以指示其过滤状态。

时间轴展开与折叠

时间轴行组可以通过行标签左侧的[+]和[-]控件展开和折叠。共有三种展开/折叠状态：

Collapsed

折叠行中包含的时间轴行不会显示。

Expanded

显示所有未筛选的时间轴行。

All-Expanded

所有时间轴行，包括已过滤和未过滤的，都会显示。

与折叠行关联的时间间隔可能不会显示在GPU详情视图和分析视图中，具体取决于为这些视图设置的过滤模式（更多信息请参阅视图文档）。例如，如果您折叠设备行，则与该设备关联的所有内存拷贝、内存设置和内核都将从这些视图显示的结果中排除。

时间线着色

时间轴着色有三种模式。着色模式可以在视图菜单、时间轴上下文菜单（通过在时间轴视图中右键点击访问）以及性能分析器工具栏中选择。在内核着色模式下，每种类型的内核会被分配唯一的颜色（即内核行中的所有活动区间都具有相同颜色）。在流着色模式下，每个流会被分配唯一的颜色（即发生在同一流上的所有内存拷贝和内核活动都会被分配相同颜色）。在进程着色模式下，每个进程会被分配唯一的颜色（即发生在同一进程中的所有内存拷贝和内核活动都会被分配相同颜色）。

聚焦内核时间线

对于使用CUDA动态并行技术的应用程序，时间线视图会显示内核活动的层次结构，展示内核之间的父子关系。默认情况下会同时显示所有父子关系。您可以使用聚焦时间线控件，将显示的父子关系限定在特定、有限的"家族树"范围内。聚焦时间线模式可以通过时间线上下文菜单（在时间线视图中右键点击访问）以及性能分析器工具栏来启用或禁用。

要查看特定内核的"族谱"，请选择一个内核然后启用聚焦模式。除了所选内核的祖先或后代之外，其他所有内核都将被隐藏。在启用聚焦模式前，可使用Ctrl键多选多个内核。使用"取消聚焦"选项可退出聚焦模式，并将所有内核恢复到时间轴视图。

依赖关系分析控制

在时间线中可视化依赖分析结果有两种模式：聚焦关键路径和高亮执行依赖。这些模式可以在视图菜单、时间线上下文菜单（通过在时间线视图中右键点击访问）以及Visual Profiler工具栏中进行选择。

这些选项在运行依赖分析应用程序分析阶段后可用（参见Unguided Application Analysis）。这些模式的详细说明请参阅Dependency Analysis Controls

2.4.1.2. 时间轴导航

时间轴可以通过多种方式进行滚动、缩放和聚焦，以帮助您更好地理解和可视化应用程序的性能。

缩放

缩放控件可在"视图"菜单、时间轴上下文菜单（通过在时间轴视图中右键点击访问）以及性能分析器工具栏中找到。放大功能会缩小视图中显示的时间范围，缩小功能则会扩大视图中显示的时间范围，而适应缩放则会调整视图比例使整个时间轴可见。

你也可以在按住Ctrl键的同时，使用鼠标滚轮进行放大和缩小。

另一个实用的缩放模式是区域缩放。通过按住Ctrl键同时左键点击并拖动鼠标，可以在时间轴上选择一个区域。释放鼠标按钮后，高亮显示的区域将扩展至占据整个视图。

滚动

时间轴可以通过鼠标滚轮的滚动条垂直滚动。时间轴可以通过滚动条水平滚动，或者在按住Shift键的同时使用鼠标滚轮来水平滚动。

高亮显示/相关性

当您将鼠标指针悬停在时间轴上的活动区间时，该区间会在显示对应活动的所有位置高亮显示。例如，如果您将鼠标指针悬停在表示内核执行的区间上，该内核执行也会在流视图和计算时间轴行中高亮显示。当内核或内存拷贝区间被高亮时，对应的驱动程序或运行时API区间也会同步高亮。这使您能够直观看到CPU端调用驱动程序/运行时API或OpenACC指令与GPU端对应活动之间的关联关系。被高亮区间的详细信息会显示在属性视图中。

选择

您可以通过左键点击时间轴区间或行来选中它。使用Ctrl+左键点击可实现多选。要取消选中某个区间或行，只需再次Ctrl+左键点击它。当单个区间或行被选中时，该区间或行的信息会固定在属性视图中。在GPU详情视图中，所选区间的详细信息会显示在表格里。

测量时间差

测量标尺可以通过在时间轴顶部的水平标尺上左键拖动来创建。一旦创建了标尺，可以通过左键点击来激活或停用它。通过Ctrl+左键点击可以同时激活多个标尺。可以创建任意数量的标尺。使用Delete或Backspace键可以删除当前激活的标尺。创建标尺后，可以通过拖动出现在时间轴上方的垂直参考线来调整其大小。如果将鼠标拖动到时间轴间隔上，参考线会自动对齐到该间隔的最近边缘。

2.4.1.3. 时间线刷新

分析器在读取数据时会逐步加载时间轴。如果要加载的数据文件较大，或者应用程序生成了大量数据，这一点会更加明显。在这种情况下，时间轴可能只呈现部分内容。同时，当前会话标签的图标会被一个旋转的圆圈取代，表示时间轴尚未完全加载。当图标恢复原状时，表示加载已完成。

为了减少内存占用，分析器可能会跳过加载某些时间线内容，如果它们在当前缩放级别不可见。当这些内容在新的缩放级别变得可见时，它们会自动加载。

2.4.1.4. 依赖关系分析控制

分析器允许在运行相应的分析阶段后，在时间线上可视化依赖分析结果。有关依赖分析工作原理的详细说明，请参阅Dependency Analysis。

关键路径聚焦模式通过突出显示关键路径上的所有区间并淡化其他区间，可视化应用程序中的关键路径。启用该模式后，当选择任意时间线区间（通过左键点击），被选区间将获得焦点。但关键路径仍会以空心区间形式保持可见。这使您能够"追踪"整个执行过程中的关键路径，并检查各个独立区间。

高亮显示执行依赖关系功能允许您分析每个时间间隔的执行依赖（请注意，某些时间间隔不会收集依赖信息）。启用此模式后，高亮颜色将从黄色（表示关联时间间隔）变为红色（表示依赖关系）。所选时间间隔及其所有传入和传出依赖都将被高亮显示。

2.4.2. 分析视图

分析视图用于控制应用程序分析并显示分析结果。它提供两种分析模式：引导模式和非引导模式。在引导模式下，分析系统将引导您完成多个分析阶段，帮助您识别应用程序中可能的性能限制因素和优化机会。在非引导模式下，您可以手动探索为应用程序收集的所有分析结果。下图展示了引导分析模式下的分析视图界面。视图左侧提供分步指导，协助您分析和优化应用程序。视图右侧则显示与每个分析阶段相对应的详细分析结果。

2.4.2.1. 引导式应用分析

在引导模式下，分析视图将逐步指导您完成整个应用程序的分析，并为应用程序中的每个内核提供具体的分析指导。引导分析从CUDA应用程序分析开始，随后会指引您发现应用程序中的优化机会。

2.4.2.2. 无引导应用分析

在非引导分析模式下，每个应用程序分析阶段都有一个"运行分析"按钮，可用于生成该阶段的分析结果。当点击"运行分析"按钮时，分析器将执行应用程序以收集进行分析所需的性能分析数据。分析阶段旁边的绿色对勾标记表示该阶段的分析结果已就绪。每个分析结果包含对该分析的简要说明，以及指向分析详细文档的"更多..."链接。当您选择一个分析结果时，与该结果相关联的时间轴行或间隔会在时间轴视图中高亮显示。

在时间线中选择单个内核实例时，会显示额外的内核特定分析阶段。每个内核特定分析阶段都有一个"运行分析"按钮，其操作方式与应用分析阶段相同。下图展示了"发散执行"分析阶段的分析结果。某些内核实例分析结果（如发散执行）与内核内的特定源代码行相关联。要查看每个结果对应的源代码，请从表格中选择一个条目，系统将打开与该条目关联的源文件。

2.4.2.3. PC采样视图

计算能力5.2及以上的设备（不包括移动设备）具备PC采样功能。该功能会定期采样每个SM中一个活跃warp的PC值和warp状态。warp状态表明该warp是否在周期内发出了指令，或者为何停顿未能发出指令。当被采样的warp处于停顿时，同一周期内可能有其他warp正在发出指令。因此采样warp的停顿并不一定表示指令发射流水线存在空缺。关于不同状态的描述，请参阅Warp State章节。

计算能力6.0及以上的设备具备一项可提供延迟原因的新特性。延迟采样数据会显示指令流水线中出现空缺的原因。在收集这些采样数据时，相应的线程束调度器中没有发出任何指令，因此这些数据会反映延迟原因。这些延迟原因将是Warp State章节中列出的停顿原因之一（"未被选中"这一停顿原因除外）。

分析器收集这些信息并在内核分析 - PC采样视图中呈现。在此视图中，所有函数和内核的采样分布以表格形式展示。饼图显示为每个内核收集的停滞原因分布情况。点击源文件或设备函数后，将打开内核分析 - PC采样视图。垂直滚动条旁边显示的热点区域由每个源代码行和汇编行收集的采样数量决定。停滞原因的分布以堆叠条形图形式展示每个源代码行和汇编行的情况。这有助于在源代码级别精确定位延迟原因。

对于计算能力6.0及以上的设备，Visual Profiler会显示两个视图：'Kernel Profile - PC Sampling'提供线程束状态视图，'Kernel Profile - PC Sampling - Latency'则显示延迟原因。可以选择热点指向'线程束状态'或'延迟原因'的热点区域。结果部分的表格展示了总延迟样本、流水线繁忙样本和指令发出样本的百分比分布情况。

博客文章Pinpoint Performance Problems with Instruction-Level Profiling展示了如何利用PC采样技术来优化CUDA内核。

2.4.2.4. 内存统计

具备计算能力5.0及以上的设备支持在内核执行期间显示内存子系统使用情况的功能。该图表展示了CUDA编程模型中内存层级的概览视图。图中绿色节点表示逻辑内存空间，而蓝色节点表示芯片上的实际硬件单元。对于各类缓存，所显示的百分比数值表示缓存命中率，即本地缓存可提供数据的请求数与总请求数的比值。

图中节点之间的连线展示了从SM（流式多处理器）到内存空间再到内存系统的数据路径。每条数据路径上显示了不同的指标。从SM到各内存空间（全局、本地、纹理、表面和共享内存）的数据路径报告了执行的内存指令总数，包括读写操作。内存空间与"统一缓存"或"共享内存"之间的数据路径报告了发起的内存请求总量。其他所有数据路径报告的是以字节为单位的内存传输总量。指向右侧的箭头表示写操作，而指向左侧的箭头表示读操作。

2.4.2.5. NVLink视图

NVIDIA NVLink是一种高带宽、高能效的互连技术，可实现CPU与GPU之间以及GPU之间的高速通信。

Visual Profiler 收集 NVLink 拓扑结构以及 NVLink 传输/接收吞吐量指标，并将这些指标映射到拓扑图上。默认情况下，拓扑数据会与时间线数据一同采集。只有在选择 NVLink 选项时才会生成吞吐量/利用率指标。

NVLink信息显示在"CUDA应用分析中的GPU使用检查"指导分析的结果部分。NVLink分析展示了显示不同设备间逻辑NVLink连接的拓扑结构。一个逻辑链路由1到4条具有相同属性的物理NVLink组成，连接在两个设备之间。可视化分析器在"逻辑NVLink属性"表中列出了逻辑NVLink的属性和实际利用率。它还在"逻辑NVLink吞吐量"表中列出了逻辑NVLink的发送和接收吞吐量。

2.4.3. 源码反汇编视图

源代码-反汇编视图用于在源代码和汇编指令级别显示内核的分析结果。要查看内核源代码，您需要使用-lineinfo选项编译代码。如果未使用此编译器选项，则只会显示反汇编视图。

此视图显示用于以下类型的分析：

• 全局内存访问模式分析

• 共享内存访问模式分析

• 发散执行分析

• 内核性能分析 - 指令执行分析

• 内核分析 - PC采样分析

作为内核的引导分析或非引导分析的一部分，分析结果会显示在分析视图下。点击源文件或设备函数后，将打开源代码-反汇编视图。如果未找到源文件，则会打开一个对话框以选择并指向源文件的新位置。例如，当在不同系统上进行性能分析时可能会发生这种情况。

源代码-反汇编视图包含：

• 高级源代码

• 汇编指令

• 源代码级别的热点分析

• 汇编指令级别的热点分析

• 用于聚合到源级别的分析数据列

• 在汇编指令级别收集的性能分析数据列

源-反汇编视图中显示的信息可以通过以下工具栏选项进行自定义：

• 视图菜单 - 从可用的性能分析器数据列中选择一个或多个进行显示。默认会根据分析类型自动选择。

• 热点菜单 - 选择使用哪种性能分析器数据来识别热点。默认会根据分析类型自动选择。

• 并排显示源代码和反汇编视图。

• 从上至下显示源代码和反汇编视图。

• 最大化源代码视图

• 最大化反汇编视图

热点区域根据重要性级别（低、中或高）进行颜色编码。将鼠标悬停在热点上会显示性能分析数据的值、重要性级别以及对应的源代码或反汇编行。您可以点击源代码级别或汇编指令级别的热点，查看与该热点对应的源代码或反汇编行。

在反汇编视图中，与所选源代码行对应的汇编指令会高亮显示。您可以点击反汇编列标题右侧显示的上、下箭头按钮，导航至下一个或上一个指令块。

2.4.4. GPU详情视图

GPU详情视图会显示一个表格，其中包含被分析应用程序中每次内存拷贝和内核执行的信息。下图展示了包含多次内存拷贝(memcpy)和内核执行的表格。表格的每一行都包含内核执行或内存拷贝的常规信息。对于内核，表格还会为每个收集到的指标或事件值单独设置一列。在图中，"Achieved Occupancy"(实际占用率)列显示了每次内核执行时该指标的数值。

您可以通过左键点击列标题来按列排序数据，也可以通过左键点击列标题并将其拖动到新位置来重新排列列。如果在表格中选择一行，对应的区间将在时间线视图中被选中。同样地，如果在时间线视图中选择一个内核或内存拷贝区间，表格将自动滚动显示对应的数据。

如果将鼠标悬停在列标题上，工具提示会显示该列中的数据。对于包含事件或指标数据的列，工具提示会描述相应的事件或指标。Metrics Reference部分包含有关每个指标的更详细信息。

可以通过详细信息视图工具栏中的菜单以多种方式筛选GPU详细信息视图中显示的信息。提供以下模式：

• 按选择筛选 - 如果选中，GPU详情视图将仅显示所选内核和内存拷贝区间的数据。

• 显示隐藏的时间线数据 - 如果不勾选，则仅显示时间线中可见的内核(kernels)和内存拷贝(memcpys)。由于位于时间线折叠部分而不可见的内核和内存拷贝将不会显示。

• 显示过滤后的时间线数据 - 如果未选中，则仅显示未被过滤的时间线行中的内核和内存拷贝数据。

收集事件与指标

可以为每个内核收集特定的事件和指标值，并显示在详情表格中。使用视图右上角的工具栏图标来配置要为每个设备收集的事件和指标，并运行应用程序以收集这些事件和指标。

显示汇总数据

默认情况下，表格会为每次内存拷贝(memcpy)和内核调用显示一行。您也可以选择显示每个内核函数的汇总结果。通过视图右上角的工具栏图标可切换汇总显示模式。

表格内容格式化

表格中的数字可以显示为带或不带分组分隔符。使用视图右上角的工具栏图标来选择或取消选择分组分隔符。

导出详情

表格内容可以通过视图右上角的工具栏图标以CSV格式导出。

2.4.5. CPU详情视图

CPU 详情视图

此视图详细展示了您的应用程序在CPU上执行函数所花费的时间。系统会定期对每个线程进行采样以捕获其调用栈，这些测量的汇总结果将显示在此视图中。您可以通过以下方式操作视图：选择不同的调用栈组织方向（自上而下、自下而上、代码结构(3)），选择要查看的线程(1)，以及对特定线程进行排序或高亮显示(7, 8)。

1. 当选择“所有线程”选项时（默认），所有分析的线程会显示在一个视图中。您可以使用此下拉菜单来选择单个线程。

2. 此列显示一个事件树，表示应用程序在CPU上的执行结构。其余各列显示为该事件收集的测量数据。此处显示的事件取决于选择的树形方向模式(3)。

3. 树状结构用于展示函数之间的调用层次关系。提供以下模式：

   • 自上而下（调用者优先）调用树视图 - CPU详情树以调用树的形式组织，每个函数显示为其调用者的子节点。在此模式下，您可以看到从'main'函数开始的调用堆栈。

   • 自底向上（被调用者优先）调用树视图 - CPU详情树以这样的方式组织：每个函数显示为它所调用的任何函数的子节点。在此模式下，您可以快速识别对应用程序执行时间影响最大的调用路径。

   • 代码结构（文件和行号）树状视图 - CPU详情树展示了哪些函数属于各个源文件和库，以及应用程序执行的多少比例可归因于特定的源代码行。

   在每种模式下，每个函数列出的时间都是“包含性的”，既包括该函数自身消耗的时间，也包括它调用的所有函数的时间。对于代码结构视图，代码区域的时间也是包含性的（例如文件条目会列出该文件内所有函数消耗的时间总和）。

4. 此列显示此事件中所有线程花费的总时间占所有事件总时间的百分比。

5. 此列显示一个条形图，表示任何线程在该事件中花费的时间始终处于此范围内。左侧标注最小值，右侧标注最大值。此外，如果有足够空间，会在条形图中间绘制一个小"菱形"标记，表示所有线程在此事件中花费的平均时间。

6. 这些列针对每个事件显示一个独立的图表。左侧是垂直刻度，显示与范围图表相同的最小值和最大值。随后的各列分别展示每个线程在此事件中花费的时间。如果特定事件/线程组合对应的单元格显示为灰色，则表示该线程在此事件中未花费时间（在本示例中，线程1和2在事件'x_solve'中均未花费时间）。此外，在所有线程中花费时间最少或最多的线程会标注"三角形/线"标记。本例中，线程3在事件'x_solve'中花费时间最多，线程6花费时间最少。

7. 要按特定线程上花费的时间重新排序行，请点击线程列标题。

8. 要突出显示某个线程，请点击此图表中对应的柱状条。

视图的这一变化是由于按线程3（7）排序并高亮显示它（8）的结果。

1. 高亮显示线程3后，我们现在可以在范围图表上看到一条垂直线，这条线展示了该线程在此事件中花费的时间与所有线程整体范围的对比。

2. 每个行上也会高亮显示此线程。

CPU线程

CPU 源代码

您可以通过在树中双击任何函数来打开其CPU源代码视图。要显示源代码文件，这些文件必须位于本地文件系统上。默认情况下会搜索包含可执行文件或配置文件的目录。如果找不到源文件，将出现提示要求指定其位置。有时系统会寻找特定目录中的文件，这种情况下您应该提供该目录所在的路径。

提示

CPU性能分析是通过定期采样运行应用程序的状态来收集的。因此，只有那些在执行过程中被采样到的函数才会出现在此视图中。运行时间较短或调用频率极低的函数被采样的可能性较低。如果某个函数未被采样，其运行时间将计入调用它的函数中。为了获取能代表应用程序性能的CPU分析数据，目标代码必须执行足够长的时间以收集充足的样本。通常一分钟的运行时间就足够了。

提示

文件和行号信息是从编译器获取的应用程序调试信息中收集的。为确保这些信息可用，建议使用'-g'或类似选项进行编译。

2.4.6. OpenACC 详情视图

OpenACC表格视图

OpenACC详情视图显示由分析应用程序执行的每个OpenACC运行时活动。每个活动按源代码位置分组：发生在应用程序源代码相同文件和行号处的每个活动都被归类到标有该源代码位置的节点下。每个活动显示分析应用程序所花费的时间，既以时间单位表示，也占该应用程序执行任何OpenACC活动总时间的百分比。同时还会显示该活动被调用的次数。有两种方式可以计算特定OpenACC活动所花费的时间：

• 在OpenACC详情视图中显示包含性持续时间(统计同时运行的其他OpenACC活动) - OpenACC详情视图会显示每项活动花费的总时间，包括由该活动引发的其他活动所消耗的时间。在这种情况下，发生在特定应用程序源代码位置上的每项活动所花费的时间会被汇总，并显示在该源代码位置所在的行上。

• 在OpenACC详情视图中显示独占持续时间（不包括同时运行的其他任何OpenACC活动） - OpenACC详情视图仅显示在给定活动中花费的时间。在这种情况下，给定源代码位置花费的时间始终为零——时间完全归因于该源代码位置发生的每个活动。

2.4.7. OpenMP 详情视图

OpenMP 表格视图

OpenMP详情视图展示了OpenMP运行时在CPU上的活动情况。应用程序在并行区域或空闲状态所花费的时间既会在时间线上显示，也会在此视图中进行汇总。各类活动时间占比的参考基准是从第一个并行区域开始到最后一个并行区域结束的时间段。由于OpenMP运行时可能同时处于多种状态，各类活动时间百分比的总和通常会超过100%。

2.4.8. 属性视图

属性视图显示在时间线视图中高亮或选定的行或区间的相关信息。若未选定任何行或区间，则显示的信息会跟随鼠标指针移动而变化。若已选定某行或区间，则显示的信息将固定关联至该选定项。

当选择带有关联源文件的OpenACC区间时，该文件名会显示在源文件表条目中。双击文件名即可打开相应的源文件（如果该文件存在于文件系统中）。

2.4.9. 控制台视图

控制台视图会在每次应用程序执行时显示其标准输出(stdout)和标准错误(stderr)输出。如需向应用程序提供标准输入(stdin)，可直接在控制台视图中输入。

2.4.10. 设置视图

设置视图允许您为正在进行分析的应用程序指定执行设置。如下图所示，可执行文件设置选项卡允许您指定应用程序的可执行文件、工作目录、命令行参数和环境。只有可执行文件是必填项，其他所有字段都是可选的。

执行超时

可执行设置选项卡还允许您指定一个可选的执行超时时间。如果指定了执行超时时间，应用程序将在达到该秒数后终止执行。如果未指定执行超时时间，应用程序将被允许继续执行，直到正常终止。

注意

计时器从CUDA驱动初始化时开始计数。如果应用程序未调用任何CUDA API，则不会触发超时。

启用性能分析后开始执行

默认勾选"启用性能分析开始执行"复选框，表示应用程序性能分析将在应用启动时开始。如果您按照聚焦性能分析中所述使用cudaProfilerStart()和cudaProfilerStop()来控制应用程序内的性能分析，则应取消勾选此选项。

启用并发内核性能分析

默认勾选"启用并发内核分析"复选框，以支持分析利用并发内核执行的应用程序。如果取消勾选此复选框，分析器将禁用并发内核执行。在某些情况下，禁用并发内核执行可以减少分析开销，因此适用于不利用并发内核的应用程序。

启用电源、时钟和温度性能分析

勾选"启用功率、时钟和热力分析"复选框可以启用对应用程序使用的每个GPU的功率、时钟和热力行为的低频采样。

2.4.11. CPU源代码视图

CPU源代码视图允许您查看构成分析应用程序CPU源代码的文件。此视图可通过在CPU详情视图中双击树状结构中的函数来打开 - 随后将打开与该函数对应的源文件。通过右键点击左侧标尺可启用行号显示。

使用PGI®编译器进行编译时，可以向此视图添加注释（更多信息请参阅通用编译器反馈格式）。这些注释是关于特定代码行如何被编译的说明。PGI编译器会保存有关程序如何被优化的信息，或为何未进行特定优化的原因。这可以与CPU详情视图结合使用，帮助识别某些代码行为何会以特定方式执行。例如，消息可能会告知您以下内容：

• 编译器生成的向量指令。

• 循环的计算强度，即计算操作与内存操作的比率——数值越高表示计算量相对于内存加载和存储操作越多。

• 关于并行化的信息，附带一个提示，说明如果编译器无法自动并行化循环，如何可能使其并行运行。

2.5. 自定义性能分析器

当你首次启动Visual Profiler, 并 在 关闭 欢迎 页面 后， 你将看到默认的视图布局。通过移动和调整视图大小，你可以自定义分析器以满足开发需求。你所做的任何更改都会在下次启动分析器时恢复。

2.5.1. 调整视图大小

要调整视图大小，只需左键点击并拖动视图之间的分隔区域。同一区域中堆叠在一起的所有视图将同时调整大小。

2.5.2. 视图重排序

要在堆叠视图集中重新排序视图，请左键点击并拖动视图标签至视图堆栈中的新位置。

2.5.3. 移动视图

要移动视图，请左键点击视图标签并将其拖动到新位置。拖动视图时，轮廓线会显示视图的目标位置。您可以将视图放置在新位置，或与其他视图堆叠在同一位置。

2.5.4. 取消停靠视图

您可以将分析器窗口中的视图取消停靠，使其占据独立的独立窗口。这样做的优势包括利用多显示器或最大化单个视图的显示区域。要取消停靠视图，请左键点击视图标签并将其拖出分析器窗口。要将视图重新停靠，请左键点击视图标签（而非窗口装饰部分）并将其拖回分析器窗口。

2.5.5. 视图的打开与关闭

点击视图标签上的X图标可关闭视图。要打开视图，请使用视图菜单。

2.6. 命令行参数

当从命令行启动Visual Profiler时，可以通过命令行参数指定可执行文件来开启新会话，或者导入从nvprof导出的性能分析文件，支持以下模式之一：

• 通过启动nvvp并指定可执行文件名称（可选后接其参数）来开启新的可执行会话：

  nvvp executableName
                  [[executableArguments]...]
  

• 通过将单个.nvprof文件作为参数启动nvvp来导入单进程nvprof会话（更多详情请参阅nvprof的导出/导入选项部分）：

  nvvp
                  data.nvprof
  

• 导入多进程 nvprof 会话，通过使用多个 .nvprof 文件作为参数启动 nvvp：

  nvvp
                  data1.nvprof data2.nvprof ...
  

3. nvprof

nvprof分析工具使您能够从命令行收集和查看分析数据。nvprof支持收集CPU和GPU上与CUDA相关活动的时间线，包括内核执行、内存传输、内存设置以及CUDA API调用和事件或CUDA内核的指标。分析选项通过命令行参数提供给nvprof。分析结果在收集完分析数据后会显示在控制台中，也可以保存下来供后续通过nvprof或Visual Profiler查看。

注意

分析器的文本输出默认重定向到stderr。使用--log-file可将输出重定向到其他文件。详见Redirecting Output。

要从命令行分析应用程序：

nvprof [options] [application]
    [application-arguments]

要查看完整的帮助页面，请输入 nvprof --help。

3.1. 命令行选项

3.1.1. CUDA性能分析选项

选项	取值	默认值	描述
聚合模式	开启, 关闭	开启	开启/关闭针对后续--events和--metrics选项指定的事件与指标的聚合模式。这些事件/指标值将针对每个域实例进行收集，而非整个设备。 更多信息请参阅事件/指标追踪模式。
分析指标	不适用	不适用	收集可导入到Visual Profiler"分析"模式的性能分析数据。注意：使用--export-profile来指定导出文件。
annotate-mpi	off, openmpi, mpich	off	自动使用NVTX标记为MPI调用添加注释。请指定您机器上安装的MPI实现。目前支持Open MPI和MPICH实现。 更多信息请参阅使用NVTX自动标注MPI。
concurrent-kernels	on, off	on	开启/关闭并发内核执行。如果关闭并发内核执行，运行在同一设备上的所有内核将被串行化处理。
continuous-sampling-interval	{间隔时间（毫秒）}	2 毫秒	设置连续模式下的采样间隔时间（毫秒）。最小值为1毫秒。
cpu-thread-tracing	on, off	off	收集关于CPU线程API活动的信息。 更多信息请参阅CPU Thread Tracing。
dependency-analysis	N/A	N/A	为主机和设备活动生成事件依赖关系图并运行依赖分析。 更多信息请参阅Dependency Analysis。
device-buffer-size	{size in MBs}	8 MB	设置每个上下文缓冲区上用于存储非CDP操作（特别是并发内核跟踪）性能分析数据的设备内存大小（以MB为单位）。该大小应为正整数。
device-cdp-buffer-size	{size in MBs}	8 MB	设置每个上下文缓冲区上用于存储CDP操作性能分析数据的设备内存大小（以MB为单位）。该大小应为正整数。
设备	{逗号分隔的设备ID}, 全部	不适用	更改后续--events、--metrics、--query-events和--query-metrics选项的作用范围。 更多信息请参阅Profiling Scope。
事件收集模式	内核, 持续	内核	为所有事件/指标选择事件收集模式。 内核：事件/指标仅在内核执行期间收集 continuous: 事件/指标会在应用程序运行期间持续收集。这不适用于非Tesla设备。此模式仅与NVLink事件/指标兼容。此模式与--profile-all-processes或--profile-child-processes或--replay-mode kernel或--replay-mode application不兼容。
事件 (e)	{逗号分隔的事件名称}, 全部	不适用	指定要在特定设备上分析的事件。可以指定多个事件名称，用逗号分隔。分析哪些设备由--devices选项控制。否则将在所有设备上收集事件。要查看可用事件列表，请使用--query-events。使用--events all可分析每个设备上所有可用的事件。使用--devices和--kernels可选择特定的内核调用。
kernel-latency-timestamps	on, off	off	开启/关闭内核延迟时间戳的收集，包括排队和提交时间。排队时间戳记录内核启动命令被加入CPU命令缓冲区的时间。提交时间戳表示包含该内核启动的CPU命令缓冲区被提交到GPU的时间。开启此选项可能会在性能分析时产生额外开销。
内核	{kernel名称}, {[上下文ID/名称]:[流ID/名称]:[kernel名称]:[调用]}	不适用	更改后续--events、--metrics选项的作用范围。语法如下： {kernel name}: 将范围限制在给定的内核名称内。 {[上下文ID/名称]:[流ID/名称]:[内核名称]:[调用次数]}: 上下文/流的ID和名称、内核名称及调用次数均可使用正则表达式匹配。空字符串表示匹配任意数量的字符。若[上下文ID/名称]或[流ID/名称]为正数，则严格匹配CUDA上下文/流ID；否则视为正则表达式，匹配NVTX库指定的上下文/流名称。若调用次数为正数，则严格匹配内核的调用次数；否则视为正则表达式处理。 示例：--kernels "1:foo:bar:2" 将分析任何名称包含"bar"的内核，该内核位于上下文1中名为"foo"的流上的第2个实例。 更多信息请参阅Profiling Scope。
指标 (m)	{逗号分隔的指标名称}, 全部	不适用	指定要在特定设备上分析的指标。可以指定多个以逗号分隔的指标名称。分析哪些设备由--devices选项控制。否则将在所有设备上收集指标。要查看可用指标列表，请使用--query-metrics。使用--metrics all分析每个设备的所有可用指标。使用--devices和--kernels选择特定的内核调用。注意：--metrics all不包含Visual Profiler源代码级分析所需的某些指标。为此，请使用--analysis-metrics。
pc-sampling-period	{周期数}	根据配置在5到12之间	指定以周期为单位的PC采样间隔，采样记录将按此间隔转储。允许的周期值为5至31之间的整数（含边界值）。这将设置采样周期为(2^period)个周期 注意：仅适用于GM20X+。
profile-all-processes	N/A	N/A	分析由启动此nvprof实例的同一用户启动的所有进程。注意：同一时间只能有一个nvprof实例使用此选项运行。在此模式下，无需指定要运行的应用程序。 更多信息请参阅多进程分析。
profile-api-trace	none, runtime, driver, all	all	开启/关闭CUDA运行时/驱动API追踪。 none: 关闭API追踪 runtime: 仅开启CUDA运行时API追踪 driver: 仅开启CUDA驱动API追踪 all: 开启所有API追踪
profile-child-processes	N/A	N/A	分析应用程序及其启动的所有子进程。 详情请参阅多进程分析。
profile-from-start	on, off	on	启用/禁用从应用程序启动时就开始性能分析。如果禁用，应用程序可以使用{cu,cuda}Profiler{Start,Stop}来开启/关闭性能分析。 更多信息请参阅Focused Profiling。
profiling-semaphore-pool-size	{count}	65536	设置用于存储每个上下文中序列化内核和内存操作的分析数据的分析信号量池大小。该值应为正整数。
查询事件	不适用	不适用	列出设备上所有可用的事件。可以通过--devices选项控制查询的设备。
查询指标	不适用	不适用	列出设备上所有可用的指标。可以通过--devices选项控制查询的设备。
回放模式	禁用, 内核, 应用程序	内核	选择在单次运行中无法收集所有事件/指标时使用的重放模式。 disabled: 回放功能被禁用，无法分析的事件/指标将被丢弃 内核：每次内核调用都会被重放 application: 整个应用程序会被重放。此模式与--profile-all-processes或profile-child-processes不兼容。
skip-kernel-replay-save-restore	on, off	off	启用此选项可大幅提升内核重放速度，因为将跳过每次内核传递时可变状态的保存与恢复。跳过输入/输出缓冲区的保存/恢复操作允许您指定上下文上所有分析的内核在执行期间不会修改其输入缓冲区内容，或调用设备内存分配/释放(new/delete)操作导致设备堆状态改变。具体而言，内核可以在同一次启动中分配并释放缓冲区，但不能调用不匹配的分配或不匹配的释放操作。注意：若在某个内核确实会修改输入缓冲区或使用不匹配的分配/释放操作时错误启用此模式，将导致未定义行为，包括内核执行失败和/或设备数据损坏。 on: 跳过输入/输出缓冲区的保存/恢复 off: 为每次内核重放传递保存/恢复输入/输出缓冲区
源码级分析 (a)	全局访问、共享访问、分支、指令执行、PC采样	不适用	指定要在特定内核调用上分析源代码级别的指标。使用--devices和--kernels来选择特定的内核调用。可以指定一个或多个指标，用逗号分隔 global_access: 全局访问 shared_access: 共享访问 分支：发散分支 instruction_execution: 指令执行 pc_sampling: 程序计数器采样，仅适用于GM20X及以上型号 注意：使用--export-profile来指定导出文件。 更多信息请参阅源代码-反汇编视图。
system-profiling	on, off	off	开启/关闭电源、时钟和温度分析功能。 详情请参阅System Profiling。
timeout (t)	{seconds}	N/A	为CUDA应用程序设置执行超时时间（以秒为单位）。注意：超时计时从CUDA驱动初始化时开始。如果应用程序未调用任何CUDA API，则不会触发超时。 更多信息请参阅Timeout和Flush Profile Data。
track-memory-allocations	on, off	off	开启/关闭内存操作跟踪功能，该功能会记录内存分配和释放的时间戳、内存大小、内存类型以及程序计数器。开启此选项可能会在性能分析时产生额外开销。
unified-memory-profiling	per-process-device, off	per-process-device	配置统一内存分析。 per-process-device: 收集每个进程和每个设备的计数 off: 关闭统一内存分析 更多信息请参阅Unified Memory Profiling。

3.1.2. CPU性能分析选项

选项	取值	默认值	描述
cpu-profiling	on, off	off	开启CPU性能分析。注意：多进程模式下不支持CPU性能分析。
cpu-profiling-explain-ccff	{filename}	N/A	设置用于解释通用编译器反馈格式(CCFF)消息的PGI pgexplain.xml文件路径。
cpu-profiling-frequency	{frequency}	100Hz	设置CPU性能分析的采样频率，单位为每秒采样次数。最大值为500Hz。
cpu-profiling-max-depth	{depth}	0 (即无限制)	设置每个调用栈的最大深度。
cpu-profiling-mode	flat, top-down, bottom-up	bottom-up	设置CPU性能分析的输出模式。 flat: 显示扁平化分析结果 top-down: 顶部显示父函数 bottom-up: 底部显示父函数
cpu-profiling-percentage-threshold	{threshold}	0 (即无限制)	过滤掉低于设定百分比阈值的条目。该限制应为0到100之间的整数（包含0和100）。
cpu-profiling-scope	function, instruction	function	选择性能分析范围。 function: 堆栈跟踪中的每个层级代表一个独立的函数 instruction: 堆栈跟踪中的每个层级代表一个独立的指令地址
CPU性能分析显示-ccff	开启, 关闭	关闭	选择是否打印嵌入二进制文件中的通用编译器反馈格式(CCFF)消息。注意：此选项隐含--cpu-profiling-scope instruction。
cpu-profiling-show-library	on, off	off	选择是否打印每个样本的库名称。
cpu-profiling-thread-mode	separated, aggregated	aggregated	设置CPU性能分析的线程模式。 separated: 为每个线程显示单独的性能分析 aggregated: 聚合所有线程的数据
cpu-profiling-unwind-stack	on, off	on	选择是否在每个采样点展开CPU调用堆栈。
openacc-profiling	on, off	on	启用/禁用从OpenACC性能分析接口记录信息。注意：OpenACC性能分析接口的可用性取决于OpenACC运行时环境。 更多信息请参阅OpenACC。
openmp-profiling	on, off	off	启用/禁用从OpenMP性能分析接口记录信息。注意：OpenMP性能分析接口的可用性取决于OpenMP运行时环境。 更多信息请参阅OpenMP。

3.1.3. 打印选项

选项	取值	默认值	描述
上下文名称	{name}	不适用	CUDA上下文的名称。 %i 在上下文名称字符串中会被替换为上下文的ID。 %p 在上下文名称字符串中会被替换为正在被分析的应用程序的进程ID。 %q{} 在上下文名称字符串中会被替换为环境变量 的值。如果该环境变量未设置，则会产生错误。 %h 在上下文名称字符串中会被替换为系统的主机名。 %% 在上下文名称字符串中被替换为 %。任何其他跟在 % 后面的字符都是非法的。
csv	N/A	N/A	在输出中使用逗号分隔的值。 更多信息请参见CSV。
名称还原	开启, 关闭	开启	开启/关闭C++函数名的名称还原功能。 详情请参阅Demangling。
时间单位标准化 (u)	s, ms, us, ns, col, auto	auto	指定输出中使用的时间单位。 s: 秒 ms: 毫秒 us: 微秒 ns: 纳秒 col: 每列使用固定单位 auto: 根据数值长度自动选择合适单位 详见调整单位获取更多信息。
openacc-summary-mode	exclusive, inclusive	exclusive	设置OpenACC摘要中持续时间计算方式。 详情请参阅OpenACC Summary Modes。
trace	api, gpu	N/A	指定要追踪的选项(可用逗号分隔多个选项)。 api - 仅开启CUDA运行时和驱动程序API追踪 gpu - 仅开启CUDA GPU追踪
print-api-summary	N/A	N/A	打印CUDA运行时/驱动API调用的摘要。
print-api-trace	N/A	N/A	打印CUDA运行时/驱动API追踪信息。 详见GPU-Trace and API-Trace Modes获取更多信息。
print-dependency-analysis-trace	N/A	N/A	打印依赖分析追踪信息。 更多信息请参阅Dependency Analysis。
print-gpu-summary	N/A	N/A	打印GPU活动摘要（包括CUDA内核和内存拷贝/设置操作）。 详情请参阅Summary Mode。
print-gpu-trace	N/A	N/A	打印各个内核调用（包括CUDA内存拷贝/设置操作）并按时间顺序排序。在事件/指标分析模式下，显示每个内核调用的事件/指标。 更多信息请参阅GPU-Trace和API-Trace模式。
print-openacc-constructs	N/A	N/A	在OpenACC性能分析中包含父构造名称。 更多信息请参阅OpenACC选项。
print-openacc-summary	N/A	N/A	打印OpenACC性能分析摘要。
print-openacc-trace	N/A	N/A	打印OpenACC性能分析跟踪信息。
print-openmp-summary	N/A	N/A	打印OpenMP性能分析摘要。
打印摘要 (s)	不适用	不适用	在屏幕上打印性能分析结果的摘要。注意：这是默认行为，除非使用了--export-profile或其他打印选项。
print-summary-per-gpu	N/A	N/A	打印每个GPU的分析结果摘要。
进程名称	{name}	不适用	进程名称。 %p 在进程名字符串中会被替换为正在被分析的应用程序的进程ID。 %q{} 在进程名字符串中会被替换为环境变量 的值。如果该环境变量未设置，则会产生错误。 %h 在进程名字符串中会被替换为该系统的主机名。 %% 在进程名字符串中被替换为 %。任何其他跟在 % 后面的字符都是非法的。
quiet	N/A	N/A	禁止所有nvprof输出。
流名称	{name}	不适用	CUDA流的名称。 %i 在流名称字符串中被替换为流的ID。 %p 在流名称字符串中会被替换为正在进行分析的应用程序的进程ID。 %q{} 在流名字符串中会被替换为环境变量 的值。如果该环境变量未设置，则会报错。 %h 在流名称字符串中被替换为系统的主机名。 %% 在流名字符串中被替换为 %。任何其他跟在 % 后面的字符都是非法的。

3.1.4. 输入输出选项

选项	取值	默认值	描述
导出配置文件 (o)	{文件名}	不适用	导出结果文件，该文件可以稍后导入或由NVIDIA Visual Profiler打开。 %p 在文件名字符串中会被替换为正在进行分析的应用程序的进程ID。 %q{} 在文件名字符串中会被替换为环境变量 的值。如果该环境变量未设置，则会报错。 %h 在文件名字符串中被替换为系统的主机名。 %% 在文件名字符串中被替换为 %。任何其他跟在 % 后面的字符都是非法的。 默认情况下，此选项会禁用摘要输出。注意：如果被分析的应用程序创建了子进程，或者使用了--profile-all-processes参数，则需要使用%p格式才能为每个进程获取正确的导出文件。 更多信息请参阅Export/Import。
force-overwrite (f)	N/A	N/A	强制覆盖所有输出文件（任何现有文件都将被覆盖）。
import-profile (i)	{filename}	N/A	从之前的运行中导入结果配置文件。 更多信息请参阅Export/Import。
日志文件	{文件名}	不适用	让nvprof将其所有输出发送到指定文件或标准通道之一。文件将被覆盖。如果文件不存在，则会创建一个新文件。 %1 作为完整的文件名表示标准输出通道(stdout)。 %2 作为完整文件名表示标准错误通道(stderr)。注意：这是默认设置。 %p 在文件名字符串中会被替换为正在进行分析的应用程序的进程ID。 %q{} 在文件名字符串中会被替换为环境变量 的值。如果该环境变量未设置，则会报错。 %h 在文件名字符串中会被替换为该系统的主机名。 %% 在文件名中被替换为 %。任何其他跟在 % 后面的字符都是非法的。 更多信息请参阅重定向输出。
print-nvlink-topology	N/A	N/A	打印nvlink拓扑结构
print-pci-topology	N/A	N/A	打印PCI拓扑结构
help (h)	N/A	N/A	打印帮助信息。
version (V)	N/A	N/A	打印此工具的版本信息。

3.2. 性能分析模式

nvprof 在以下列出的其中一种模式下运行。

3.2.1. 摘要模式

摘要模式是nvprof的默认运行模式。在此模式下，nvprof会为应用程序执行的每个内核函数和每种CUDA内存复制/设置类型输出单行结果。对于每个内核，nvprof会输出该内核或内存复制类型所有实例的总时间以及平均、最小和最大时间。内核时间指的是设备上的内核执行时间。默认情况下，nvprof还会打印所有CUDA运行时/驱动API调用的摘要。nvprof的输出（表格除外）都会以====为前缀，其中表示被分析应用程序的进程ID。

以下是一个在CUDA示例matrixMul上运行nvprof的简单示例：

$ nvprof matrixMul
[Matrix Multiply Using CUDA] - Starting...
==27694== NVPROF is profiling process 27694, command: matrixMul
GPU Device 0: "GeForce GT 640M LE" with compute capability 3.0

MatrixA(320,320), MatrixB(640,320)
Computing result using CUDA Kernel...
done
Performance= 35.35 GFlop/s, Time= 3.708 msec, Size= 131072000 Ops, WorkgroupSize= 1024 threads/block
Checking computed result for correctness: OK

Note: For peak performance, please refer to the matrixMulCUBLAS example.
==27694== Profiling application: matrixMul
==27694== Profiling result:
Time(%)      Time     Calls       Avg       Min       Max  Name
 99.94%  1.11524s       301  3.7051ms  3.6928ms  3.7174ms  void matrixMulCUDA<int=32>(float*, float*, float*, int, int)
  0.04%  406.30us         2  203.15us  136.13us  270.18us  [CUDA memcpy HtoD]
  0.02%  248.29us         1  248.29us  248.29us  248.29us  [CUDA memcpy DtoH]

==27964== API calls:
Time(%)      Time     Calls       Avg       Min       Max  Name
 49.81%  285.17ms         3  95.055ms  153.32us  284.86ms  cudaMalloc
 25.95%  148.57ms         1  148.57ms  148.57ms  148.57ms  cudaEventSynchronize
 22.23%  127.28ms         1  127.28ms  127.28ms  127.28ms  cudaDeviceReset
  1.33%  7.6314ms       301  25.353us  23.551us  143.98us  cudaLaunch
  0.25%  1.4343ms         3  478.09us  155.84us  984.38us  cudaMemcpy
  0.11%  601.45us         1  601.45us  601.45us  601.45us  cudaDeviceSynchronize
  0.10%  564.48us      1505     375ns     313ns  3.6790us  cudaSetupArgument
  0.09%  490.44us        76  6.4530us     307ns  221.93us  cuDeviceGetAttribute
  0.07%  406.61us         3  135.54us  115.07us  169.99us  cudaFree
  0.02%  143.00us       301     475ns     431ns  2.4370us  cudaConfigureCall
  0.01%  42.321us         1  42.321us  42.321us  42.321us  cuDeviceTotalMem
  0.01%  33.655us         1  33.655us  33.655us  33.655us  cudaGetDeviceProperties
  0.01%  31.900us         1  31.900us  31.900us  31.900us  cuDeviceGetName
  0.00%  21.874us         2  10.937us  8.5850us  13.289us  cudaEventRecord
  0.00%  16.513us         2  8.2560us  2.6240us  13.889us  cudaEventCreate
  0.00%  13.091us         1  13.091us  13.091us  13.091us  cudaEventElapsedTime
  0.00%  8.1410us         1  8.1410us  8.1410us  8.1410us  cudaGetDevice
  0.00%  2.6290us         2  1.3140us     509ns  2.1200us  cuDeviceGetCount
  0.00%  1.9970us         2     998ns     520ns  1.4770us  cuDeviceGet

注意

如果不需要，可以通过使用--profile-api-trace none来关闭API跟踪功能。这可以减少部分性能分析开销，特别是在内核运行时间较短的情况下。

如果分析多个支持CUDA的设备，可以使用nvprof --print-summary-per-gpu来为每个GPU打印一份摘要。

nvprof 在摘要模式下支持 CUDA 动态并行。如果您的应用程序使用了动态并行，输出结果将包含两列：一列显示主机启动的内核数量，另一列显示设备启动的内核数量。以下是在 CUDA 动态并行示例 cdpSimpleQuicksort 上运行 nvprof 的示例：

$ nvprof cdpSimpleQuicksort
==27325== NVPROF is profiling process 27325, command: cdpSimpleQuicksort

Running on GPU 0 (Tesla K20c)
Initializing data:
Running quicksort on 128 elements
Launching kernel on the GPU
Validating results: OK
==27325== Profiling application: cdpSimpleQuicksort
==27325== Profiling result:
Time(%)      Time  Calls (host)  Calls (device)       Avg       Min       Max  Name
 99.71%  1.2114ms             1              14  80.761us  5.1200us  145.66us  cdp_simple_quicksort(unsigned int*, int, int, int)
  0.18%  2.2080us             1               -  2.2080us  2.2080us  2.2080us  [CUDA memcpy DtoH]
  0.11%  1.2800us             1               -  1.2800us  1.2800us  1.2800us  [CUDA memcpy HtoD]

3.2.2. GPU追踪与API追踪模式

可以单独或同时启用GPU-Trace和API-Trace模式。GPU-Trace模式按时间顺序提供GPU上所有活动的时序图。输出中会显示每个内核执行和内存复制/设置实例。对于每个内核或内存复制操作，会显示详细信息，如内核参数、共享内存使用情况和内存传输吞吐量。内核名称后方括号内显示的数字与启动该内核的CUDA API相关联。

以下是一个示例：

$ nvprof --print-gpu-trace matrixMul
==27706== NVPROF is profiling process 27706, command: matrixMul

==27706== Profiling application: matrixMul
[Matrix Multiply Using CUDA] - Starting...
GPU Device 0: "GeForce GT 640M LE" with compute capability 3.0

MatrixA(320,320), MatrixB(640,320)
Computing result using CUDA Kernel...
done
Performance= 35.36 GFlop/s, Time= 3.707 msec, Size= 131072000 Ops, WorkgroupSize= 1024 threads/block
Checking computed result for correctness: OK

Note: For peak performance, please refer to the matrixMulCUBLAS example.
==27706== Profiling result:
   Start  Duration            Grid Size      Block Size     Regs*    SSMem*    DSMem*      Size  Throughput           Device   Context    Stream  Name
133.81ms  135.78us                    -               -         -         -         -  409.60KB  3.0167GB/s  GeForce GT 640M         1         2  [CUDA memcpy HtoD]
134.62ms  270.66us                    -               -         -         -         -  819.20KB  3.0267GB/s  GeForce GT 640M         1         2  [CUDA memcpy HtoD]
134.90ms  3.7037ms            (20 10 1)       (32 32 1)        29  8.1920KB        0B         -           -  GeForce GT 640M         1         2  void matrixMulCUDA<int=32>(float*, float*, float*, int, int) [94]
138.71ms  3.7011ms            (20 10 1)       (32 32 1)        29  8.1920KB        0B         -           -  GeForce GT 640M         1         2  void matrixMulCUDA<int=32>(float*, float*, float*, int, int) [105]
<...more output...>
1.24341s  3.7011ms            (20 10 1)       (32 32 1)        29  8.1920KB        0B         -           -  GeForce GT 640M         1         2  void matrixMulCUDA<int=32>(float*, float*, float*, int, int) [2191]
1.24711s  3.7046ms            (20 10 1)       (32 32 1)        29  8.1920KB        0B         -           -  GeForce GT 640M         1         2  void matrixMulCUDA<int=32>(float*, float*, float*, int, int) [2198]
1.25089s  248.13us                    -               -         -         -         -  819.20KB  3.3015GB/s  GeForce GT 640M         1         2  [CUDA memcpy DtoH]

Regs: Number of registers used per CUDA thread. This number includes registers used internally by the CUDA driver and/or tools and can be more than what the compiler shows.
SSMem: Static shared memory allocated per CUDA block.
DSMem: Dynamic shared memory allocated per CUDA block.

nvprof 在GPU追踪模式下支持CUDA动态并行。对于主机内核启动，将显示内核ID。对于设备内核启动，将显示内核ID、父内核ID和父块。示例如下：

$nvprof --print-gpu-trace cdpSimpleQuicksort
==28128== NVPROF is profiling process 28128, command: cdpSimpleQuicksort

Running on GPU 0 (Tesla K20c)
Initializing data:
Running quicksort on 128 elements
Launching kernel on the GPU
Validating results: OK
==28128== Profiling application: cdpSimpleQuicksort
==28128== Profiling result:
   Start  Duration            Grid Size      Block Size     Regs*    SSMem*    DSMem*      Size  Throughput           Device   Context    Stream          ID   Parent ID         Parent Block  Name
192.76ms  1.2800us                    -               -         -         -         -      512B  400.00MB/s   Tesla K20c (0)         1         2           -           -                    -  [CUDA memcpy HtoD]
193.31ms  146.02us              (1 1 1)         (1 1 1)        32        0B        0B         -           -   Tesla K20c (0)         1         2           2           -                    -  cdp_simple_quicksort(unsigned int*, int, int, int) [171]
193.41ms  110.53us              (1 1 1)         (1 1 1)        32        0B      256B         -           -   Tesla K20c (0)         1         2          -5           2              (0 0 0)  cdp_simple_quicksort(unsigned int*, int, int, int)
193.45ms  125.57us              (1 1 1)         (1 1 1)        32        0B      256B         -           -   Tesla K20c (0)         1         2          -6           2              (0 0 0)  cdp_simple_quicksort(unsigned int*, int, int, int)
193.48ms  9.2480us              (1 1 1)         (1 1 1)        32        0B      256B         -           -   Tesla K20c (0)         1         2          -7          -5              (0 0 0)  cdp_simple_quicksort(unsigned int*, int, int, int)
193.52ms  107.23us              (1 1 1)         (1 1 1)        32        0B      256B         -           -   Tesla K20c (0)         1         2          -8          -5              (0 0 0)  cdp_simple_quicksort(unsigned int*, int, int, int)
193.53ms  93.824us              (1 1 1)         (1 1 1)        32        0B      256B         -           -   Tesla K20c (0)         1         2          -9          -6              (0 0 0)  cdp_simple_quicksort(unsigned int*, int, int, int)
193.57ms  117.47us              (1 1 1)         (1 1 1)        32        0B      256B         -           -   Tesla K20c (0)         1         2         -10          -6              (0 0 0)  cdp_simple_quicksort(unsigned int*, int, int, int)
193.58ms  5.0560us              (1 1 1)         (1 1 1)        32        0B      256B         -           -   Tesla K20c (0)         1         2         -11          -8              (0 0 0)  cdp_simple_quicksort(unsigned int*, int, int, int)
193.62ms  108.06us              (1 1 1)         (1 1 1)        32        0B      256B         -           -   Tesla K20c (0)         1         2         -12          -8              (0 0 0)  cdp_simple_quicksort(unsigned int*, int, int, int)
193.65ms  113.34us              (1 1 1)         (1 1 1)        32        0B      256B         -           -   Tesla K20c (0)         1         2         -13         -10              (0 0 0)  cdp_simple_quicksort(unsigned int*, int, int, int)
193.68ms  29.536us              (1 1 1)         (1 1 1)        32        0B      256B         -           -   Tesla K20c (0)         1         2         -14         -12              (0 0 0)  cdp_simple_quicksort(unsigned int*, int, int, int)
193.69ms  22.848us              (1 1 1)         (1 1 1)        32        0B      256B         -           -   Tesla K20c (0)         1         2         -15         -10              (0 0 0)  cdp_simple_quicksort(unsigned int*, int, int, int)
193.71ms  130.85us              (1 1 1)         (1 1 1)        32        0B      256B         -           -   Tesla K20c (0)         1         2         -16         -13              (0 0 0)  cdp_simple_quicksort(unsigned int*, int, int, int)
193.73ms  62.432us              (1 1 1)         (1 1 1)        32        0B      256B         -           -   Tesla K20c (0)         1         2         -17         -12              (0 0 0)  cdp_simple_quicksort(unsigned int*, int, int, int)
193.76ms  41.024us              (1 1 1)         (1 1 1)        32        0B      256B         -           -   Tesla K20c (0)         1         2         -18         -13              (0 0 0)  cdp_simple_quicksort(unsigned int*, int, int, int)
193.92ms  2.1760us                    -               -         -         -         -      512B  235.29MB/s   Tesla K20c (0)         1         2           -           -                    -  [CUDA memcpy DtoH]

Regs: Number of registers used per CUDA thread. This number includes registers used internally by the CUDA driver and/or tools and can be more than what the compiler shows.
SSMem: Static shared memory allocated per CUDA block.
DSMem: Dynamic shared memory allocated per CUDA block.

API跟踪模式按时间顺序显示主机上调用的所有CUDA运行时和驱动程序API调用的时间线。示例如下：

$nvprof --print-api-trace matrixMul
==27722== NVPROF is profiling process 27722, command: matrixMul

==27722== Profiling application: matrixMul
[Matrix Multiply Using CUDA] - Starting...
GPU Device 0: "GeForce GT 640M LE" with compute capability 3.0

MatrixA(320,320), MatrixB(640,320)
Computing result using CUDA Kernel...
done
Performance= 35.35 GFlop/s, Time= 3.708 msec, Size= 131072000 Ops, WorkgroupSize= 1024 threads/block
Checking computed result for correctness: OK

Note: For peak performance, please refer to the matrixMulCUBLAS example.
==27722== Profiling result:
   Start  Duration  Name
108.38ms  6.2130us  cuDeviceGetCount
108.42ms     840ns  cuDeviceGet
108.42ms  22.459us  cuDeviceGetName
108.45ms  11.782us  cuDeviceTotalMem
108.46ms     945ns  cuDeviceGetAttribute
149.37ms  23.737us  cudaLaunch (void matrixMulCUDA<int=32>(float*, float*, float*, int, int) [2198])
149.39ms  6.6290us  cudaEventRecord
149.40ms  1.10156s  cudaEventSynchronize
<...more output...>
1.25096s  21.543us  cudaEventElapsedTime
1.25103s  1.5462ms  cudaMemcpy
1.25467s  153.93us  cudaFree
1.25483s  75.373us  cudaFree
1.25491s  75.564us  cudaFree
1.25693s  10.901ms  cudaDeviceReset

注意

由于分析器的设置方式，第一个“cuInit()”驱动API调用不会被追踪。

3.2.3. 事件/指标汇总模式

要查看特定NVIDIA GPU上所有可用事件的列表，请使用--query-events选项。要查看特定NVIDIA GPU上所有可用指标的列表，请使用--query-metrics选项。nvprof能够同时收集多个事件/指标。以下是一个示例：

$ nvprof --events warps_launched,local_load --metrics ipc matrixMul
[Matrix Multiply Using CUDA] - Starting...
==6461== NVPROF is profiling process 6461, command: matrixMul

GPU Device 0: "GeForce GTX TITAN" with compute capability 3.5

MatrixA(320,320), MatrixB(640,320)
Computing result using CUDA Kernel...
==6461== Warning: Some kernel(s) will be replayed on device 0 in order to collect all events/metrics.
done
Performance= 6.39 GFlop/s, Time= 20.511 msec, Size= 131072000 Ops, WorkgroupSize= 1024 threads/block
Checking computed result for correctness: Result = PASS

NOTE: The CUDA Samples are not meant for performance measurements. Results may vary when GPU Boost is enabled.
==6461== Profiling application: matrixMul
==6461== Profiling result:
==6461== Event result:
Invocations                                Event Name         Min         Max         Avg
Device "GeForce GTX TITAN (0)"
    Kernel: void matrixMulCUDA<int=32>(float*, float*, float*, int, int)
        301                            warps_launched        6400        6400        6400
        301                                local_load           0           0           0

==6461== Metric result:
Invocations                               Metric Name                        Metric Description         Min         Max         Avg
Device "GeForce GTX TITAN (0)"
    Kernel: void matrixMulCUDA<int=32>(float*, float*, float*, int, int)
        301                                       ipc                              Executed IPC    1.282576    1.299736    1.291500

如果指定的性能事件/指标无法在应用程序单次运行中完成分析，nvprof默认会多次重放每个内核，直到收集完所有事件/指标。

--replay-mode 选项可用于更改重放模式。在"应用程序重放"模式下，nvprof会重新运行整个应用程序而不是逐个重放内核，以收集所有事件/指标。如果应用程序分配了大量设备内存，在某些情况下此模式可能比内核重放模式更快。也可以完全关闭重放功能，在这种情况下分析器将不会收集某些事件/指标。

要收集每台设备上所有可用的事件，请使用选项 --events all。

要收集每个设备上所有可用的指标，请使用选项 --metrics all。

注意

事件或指标收集可能会显著改变应用程序的整体性能特征，因为所有内核执行都在GPU上被串行化。

注意

如果请求大量事件或指标，无论选择哪种重放模式，整体应用程序执行时间可能会显著增加。

3.2.4. 事件/指标追踪模式

在事件/指标追踪模式下，会显示每次内核执行的事件和指标值。默认情况下，事件和指标值会在GPU的所有单元上进行聚合。例如，特定于多处理器的事件会在GPU的所有多处理器上进行聚合。如果指定了--aggregate-mode off，则会显示每个单元的值。例如，在以下示例中，"branch"事件值会显示GPU上每个多处理器的数值：

$ nvprof --aggregate-mode off --events local_load --print-gpu-trace matrixMul
[Matrix Multiply Using CUDA] - Starting...
==6740== NVPROF is profiling process 6740, command: matrixMul

GPU Device 0: "GeForce GTX TITAN" with compute capability 3.5

MatrixA(320,320), MatrixB(640,320)
Computing result using CUDA Kernel...
done
Performance= 16.76 GFlop/s, Time= 7.822 msec, Size= 131072000 Ops, WorkgroupSize= 1024 threads/block
Checking computed result for correctness: Result = PASS

NOTE: The CUDA Samples are not meant for performance measurements. Results may vary when GPU Boost is enabled.
==6740== Profiling application: matrixMul
==6740== Profiling result:
         Device          Context           Stream                Kernel  local_load (0)  local_load (1)  ...
GeForce GTX TIT                1                7  void matrixMulCUDA<i               0               0  ...
GeForce GTX TIT                1                7  void matrixMulCUDA<i               0               0  ...
<...more output...>

注意

尽管--aggregate-mode适用于指标，但某些指标仅在聚合模式下可用，而某些指标仅在非聚合模式下可用。

3.3. 性能分析控制

3.3.1. 超时设置

可以为nvprof设置超时时间（以秒为单位）。被分析的CUDA应用程序将在超时后被nvprof终止。超时前收集的分析结果将会显示。

注意

超时从CUDA驱动初始化时开始计时。如果应用程序未调用任何CUDA API，则不会触发超时。

3.3.2. 并行内核

支持并发内核分析功能，默认情况下该功能处于开启状态。如需关闭此功能，请使用选项--concurrent-kernels off。当使用nvprof运行CUDA应用程序时，这将强制将并发内核执行序列化。

3.3.3. 性能分析范围

在收集事件/指标时，nvprof默认会分析所有可见CUDA设备上启动的所有内核。可以通过以下选项限制分析范围。

--devices IDs> 适用于紧随其后的 --events、--metrics、--query-events 和 --query-metrics 选项。它会将这些选项限制为仅收集由                                   IDs> 指定的设备上的事件/指标，这些设备ID可以是用逗号分隔的数字列表。

--kernels filter> 适用于紧随其后的 --events 和 --metrics 选项。它会限制这些选项，使其仅收集由 filter> 指定的内核上的事件/指标，该过滤器具有以下语法：

<kernel name>

或

<context id/name>:<stream id/name>:<kernel
        name>:<invocation>

尖括号中的每个字符串都可以是一个标准的Perl正则表达式。空字符串可以匹配任意数字或字符组合。

调用编号n表示内核的第n次调用。如果调用编号是正数，则严格匹配内核的调用次数；否则将被视为正则表达式。每个内核的调用计数是独立的。例如:::3将匹配每个内核的第3次调用。

如果上下文/流字符串是一个正数，它将严格匹配cuda上下文/流ID。否则它将被视为正则表达式，并与NVIDIA工具扩展提供的上下文/流名称进行匹配。

可以多次指定--devices和--kernels，每个都有不同的事件/指标关联。

--events, --metrics, --query-events 和 --query-metrics 由它们前面最近的scope选项控制。

例如，以下命令

nvprof --devices 0 --metrics ipc
        --kernels "1:foo:bar:2" --events local_load a.out

收集设备0上所有启动的内核的指标ipc。同时，对于名称包含bar、在上下文1上第2个启动、且在设备0上名为foo的流上运行的内核，还会收集事件local_load。

3.3.4. 多进程性能分析

默认情况下，nvprof仅分析命令行参数指定的应用程序。它不会追踪该进程启动的子进程。若要分析应用程序启动的所有进程，请使用--profile-child-processes选项。

注意

nvprof 无法分析那些执行了 fork() 但未接着执行 exec() 的进程。

nvprof 还具有"分析所有进程"模式，在该模式下它会分析由启动nvprof的同一用户在相同系统上启动的每个CUDA进程。通过输入"Ctrl-c"可退出此模式。

注意

在多进程模式下不支持CPU性能分析。

3.3.5. 系统性能分析

对于支持系统性能分析的设备，nvprof可以启用对应用程序使用的每个GPU的功耗、时钟和热行为的低频采样。此功能默认关闭。要启用此功能，请使用--system-profiling on。要查看每个采样点的详细信息，请将上述选项与--print-gpu-trace结合使用。

3.3.6. 统一内存性能分析

对于支持统一内存的GPU，nvprof会收集系统中每个GPU与统一内存相关的内存流量。此功能默认启用。可以通过--unified-memory-profiling off禁用该功能。要查看启用此功能时每次内存传输的详细信息，请使用--print-gpu-trace。

在不支持任何设备对之间P2P(点对点)的多GPU配置上，如果设备支持统一内存，托管内存分配会被放置在零拷贝内存中。这种情况下不支持统一内存性能分析。在某些情况下，可以设置环境变量CUDA_MANAGED_FORCE_DEVICE_ALLOC来强制将托管分配置于设备内存中，并在此类硬件配置上启用迁移功能。这种情况下支持统一内存性能分析。通常建议使用环境变量CUDA_VISIBLE_DEVICES来限制CUDA仅使用那些支持P2P的GPU。更多详细信息请参阅CUDA C++编程指南中的环境变量章节。

3.3.7. CPU线程追踪

为了确保正确的依赖关系分析，nvprof可以收集关于CPU端线程API的信息。这可以通过在测量时指定--cpu-thread-tracing on来启用。在以下情况下需要记录这些信息：

• 应用程序使用多个CPU线程

• 至少有两个线程调用了CUDA API。

目前仅支持POSIX线程(Pthreads)。出于性能考虑，仅会记录部分选定的Pthread API调用。nvprof会尝试检测哪些调用对建模执行行为是必要的，并过滤其他调用。被过滤的调用包括pthread_mutex_lock和pthread_mutex_unlock，当这些调用不会导致任何并发线程阻塞时。

注意

Windows系统不支持CPU线程追踪。

注意

CPU线程追踪在首次CUDA API调用后启动，从发起该调用的线程开始。因此，应用程序必须在其主线程中调用例如cuInit，然后再创建其他调用CUDA API的用户线程。

3.4. 输出

3.4.1. 调整单位

默认情况下，nvprof会自动调整时间单位以获得最精确的时间值。可以使用--normalized-time-unit选项在整个结果中获取固定的时间单位。

3.4.2. CSV

对于每种性能分析模式，可以使用--csv选项生成逗号分隔值(CSV)格式的输出结果。该结果可直接导入Excel等电子表格软件。

3.4.3. 导出/导入

对于每种分析模式，可以使用--export-profile选项生成结果文件。该文件不可人工阅读，但可以通过--import-profile选项重新导入到nvprof，或导入到Visual Profiler中。

注意

分析器使用SQLite作为导出性能分析数据的格式。以这种格式写入文件可能需要比写入普通文件更多的磁盘操作。因此，将性能分析数据导出到较慢的设备（如网络驱动器）可能会降低应用程序的执行速度。

3.4.4. 名称还原

默认情况下，nvprof会对C++函数名进行反混淆处理。使用选项--demangling off可关闭此功能。

3.4.5. 输出重定向

默认情况下，nvprof会将其大部分输出发送到stderr。要重定向输出，请使用--log-file。--log-file %1指示nvprof将所有输出重定向到stdout。--log-file 则将输出重定向到文件。在文件名中使用%p会被替换为nvprof的进程ID，%h替换为主机名，%q{ENV}替换为环境变量ENV的值，而%%则替换为%。

3.4.6. 依赖关系分析

nvprof 可以在分析应用程序后运行依赖关系分析，使用--dependency-analysis选项。该分析也可以应用于导入的分析数据。它需要在测量期间收集完整的CUDA API和GPU活动跟踪。这是nvprof的默认行为，除非使用--profile-api-trace none禁用。

对于从多个CPU线程使用CUDA的应用程序，也应启用CPU线程追踪。可以指定选项--print-dependency-analysis-trace将摘要输出更改为追踪输出，显示计算出的指标，例如每个函数实例而非函数类型的临界路径时间。

下面展示了一个依赖分析汇总输出的示例，其中所有计算指标都按函数类型进行了聚合。该表格首先按关键路径上的时间排序，其次按等待时间排序。汇总中包含一个名为Other的条目，指代所有未被nvprof跟踪的CPU活动（例如应用程序的main函数）。

==20704== Dependency Analysis:
==20704== Analysis progress: 100%
Critical path(%)  Critical path  Waiting time  Name
               %              s             s
           92.06       4.061817      0.000000  clock_block(long*, long)
            4.54       0.200511      0.000000  cudaMalloc
            3.25       0.143326      0.000000  cudaDeviceReset
            0.13  5.7273280e-03      0.000000  <Other>
            0.01  2.7200900e-04      0.000000  cudaFree
            0.00       0.000000      4.062506  pthread_join
            0.00       0.000000      4.061790  cudaStreamSynchronize
            0.00       0.000000      1.015485  pthread_mutex_lock
            0.00       0.000000      1.013711  pthread_cond_wait
            0.00       0.000000      0.000000  pthread_mutex_unlock
            0.00       0.000000      0.000000  pthread_exit
            0.00       0.000000      0.000000  pthread_enter
            0.00       0.000000      0.000000  pthread_create
            0.00       0.000000      0.000000  pthread_cond_signal
            0.00       0.000000      0.000000  cudaLaunch

3.5. CPU采样

有时分析应用程序的CPU部分非常有用，这有助于更好地理解瓶颈并识别整个CUDA应用的潜在热点。对于应用程序的CPU部分，nvprof能够以特定频率采样程序计数器和调用堆栈。这些数据随后被用于构建一个图表，其中节点代表每个调用堆栈中的帧。如果可用，还会提取函数和库符号。示例如下：

======== CPU profiling result (bottom up):
45.45% cuInit
| 45.45% cudart::globalState::loadDriverInternal(void)
|   45.45% cudart::__loadDriverInternalUtil(void)
|     45.45% pthread_once
|       45.45% cudart::cuosOnce(int*, void (*) (void))
|         45.45% cudart::globalState::loadDriver(void)
|           45.45% cudart::globalState::initializeDriver(void)
|             45.45% cudaMalloc
|               45.45% main
33.33% cuDevicePrimaryCtxRetain
| 33.33% cudart::contextStateManager::initPrimaryContext(cudart::device*)
|   33.33% cudart::contextStateManager::tryInitPrimaryContext(cudart::device*)
|     33.33% cudart::contextStateManager::initDriverContext(void)
|       33.33% cudart::contextStateManager::getRuntimeContextState(cudart::contextState**, bool)
|         33.33% cudart::getLazyInitContextState(cudart::contextState**)
|           33.33% cudart::doLazyInitContextState(void)
|             33.33% cudart::cudaApiMalloc(void**, unsigned long)
|               33.33% cudaMalloc
|                 33.33% main
18.18% cuDevicePrimaryCtxReset
| 18.18% cudart::device::resetPrimaryContext(void)
|   18.18% cudart::cudaApiThreadExit(void)
|     18.18% cudaThreadExit
|       18.18% main
3.03% cudbgGetAPIVersion
  3.03% start_thread
    3.03% clone

该图表可以呈现不同的"视图"（自上而下、自下而上或平面），让用户能够从不同角度分析采样数据。例如，自下而上视图（如上所示）有助于识别应用程序花费最多时间的"热点"函数。而自上而下视图则从主函数开始分解应用程序执行时间，帮助您找到频繁执行的"调用路径"。

默认情况下，CPU采样功能处于禁用状态。要启用该功能，请使用选项--cpu-profiling on。下一节将介绍控制CPU采样行为的所有选项。

CPU采样在Linux和Windows系统上支持Intel x86/x86_64架构。

注意

在POSIX系统上使用CPU性能分析功能时，分析器会通过发送周期性信号对应用程序进行采样。因此应用程序应确保系统调用在被中断时能得到正确处理。

注意

在Windows系统上，nvprof需要安装Visual Studio（2010或更高版本）以及编译器生成的.PDB（程序数据库）文件来解析符号信息。在构建应用程序时，请确保生成.PDB文件并将其放置在待分析的可执行文件和库文件旁边。

3.5.1. CPU采样限制

以下是当前版本中已知的问题。

• 移动设备不支持CPU采样。

• 在多进程分析模式下，目前不支持CPU采样。

• 在某些编译器优化下，结果堆栈跟踪可能不完整，特别是帧指针省略和函数内联的情况。

• CPU采样结果不支持CSV模式。

3.6. OpenACC

在64位Linux平台上，nvprof支持使用CUPTI Activity API记录OpenACC活动。这允许除了底层编译器生成的CUDA API调用外，还能在OpenACC结构层面进行性能分析。

在nvprof中进行OpenACC性能分析要求目标应用程序使用PGI OpenACC运行时19.1或更高版本。

尽管记录OpenACC活动仅支持x86_64 Linux系统，但在nvprof支持的所有平台上都可以导入和查看先前生成的性能分析数据。

以下展示了一个OpenACC摘要输出的示例。CUPTI OpenACC活动通过源文件和行信息映射到原始的OpenACC结构。对于acc_enqueue_launch活动，它还会显示由OpenACC编译器生成的已启动CUDA内核名称。默认情况下，nvprof会对OpenACC编译器生成的内核名称进行反混淆处理。您可以通过传递--demangling off来禁用此行为。

==20854== NVPROF is profiling process 20854, command: ./acc_saxpy

==20854== Profiling application: ./acc_saxpy
==20854== Profiling result:
==20854== OpenACC (excl):
Time(%)      Time     Calls       Avg       Min       Max  Name
 33.16%  1.27944s       200  6.3972ms  24.946us  12.770ms  acc_implicit_wait@acc_saxpy.cpp:42
 33.12%  1.27825s       100  12.783ms  12.693ms  12.787ms  acc_wait@acc_saxpy.cpp:54
 33.12%  1.27816s       100  12.782ms  12.720ms  12.786ms  acc_wait@acc_saxpy.cpp:61
  0.14%  5.4550ms       100  54.549us  51.858us  71.461us  acc_enqueue_download@acc_saxpy.cpp:43
  0.07%  2.5190ms       100  25.189us  23.877us  60.269us  acc_enqueue_launch@acc_saxpy.cpp:50 (kernel2(int, float, float*, float*)_50_gpu)
  0.06%  2.4988ms       100  24.987us  24.161us  29.453us  acc_enqueue_launch@acc_saxpy.cpp:60 (kernel3(int, float, float*, float*)_60_gpu)
  0.06%  2.2799ms       100  22.798us  21.654us  56.674us  acc_enqueue_launch@acc_saxpy.cpp:42 (kernel1(int, float, float*, float*)_42_gpu)
  0.05%  2.1068ms       100  21.068us  20.444us  33.159us  acc_enqueue_download@acc_saxpy.cpp:51
  0.05%  2.0854ms       100  20.853us  19.453us  23.697us  acc_enqueue_download@acc_saxpy.cpp:61
  0.04%  1.6265ms       100  16.265us  15.284us  49.632us  acc_enqueue_upload@acc_saxpy.cpp:50
  0.04%  1.5963ms       100  15.962us  15.052us  19.749us  acc_enqueue_upload@acc_saxpy.cpp:60
  0.04%  1.5393ms       100  15.393us  14.592us  56.414us  acc_enqueue_upload@acc_saxpy.cpp:42
  0.01%  558.54us       100  5.5850us  5.3700us  6.2090us  acc_implicit_wait@acc_saxpy.cpp:43
  0.01%  266.13us       100  2.6610us  2.4630us  4.7590us  acc_compute_construct@acc_saxpy.cpp:42
  0.01%  211.77us       100  2.1170us  1.9980us  4.1770us  acc_compute_construct@acc_saxpy.cpp:50
  0.01%  209.14us       100  2.0910us  1.9880us  2.2500us  acc_compute_construct@acc_saxpy.cpp:60
  0.00%  55.066us         1  55.066us  55.066us  55.066us  acc_enqueue_launch@acc_saxpy.cpp:70 (initVec(int, float, float*)_70_gpu)
  0.00%  13.209us         1  13.209us  13.209us  13.209us  acc_compute_construct@acc_saxpy.cpp:70
  0.00%  10.901us         1  10.901us  10.901us  10.901us  acc_implicit_wait@acc_saxpy.cpp:70
  0.00%       0ns       200       0ns       0ns       0ns  acc_delete@acc_saxpy.cpp:61
  0.00%       0ns       200       0ns       0ns       0ns  acc_delete@acc_saxpy.cpp:43
  0.00%       0ns       200       0ns       0ns       0ns  acc_create@acc_saxpy.cpp:60
  0.00%       0ns       200       0ns       0ns       0ns  acc_create@acc_saxpy.cpp:42
  0.00%       0ns       200       0ns       0ns       0ns  acc_delete@acc_saxpy.cpp:51
  0.00%       0ns       200       0ns       0ns       0ns  acc_create@acc_saxpy.cpp:50
  0.00%       0ns         2       0ns       0ns       0ns  acc_alloc@acc_saxpy.cpp:42

3.6.1. OpenACC选项

表1列出了nvprof中与OpenACC性能分析相关的命令行选项。

表1. OpenACC选项

选项	描述
--openacc-profiling	开启/关闭OpenACC性能分析。注意：OpenACC性能分析仅支持x86_64 Linux平台。默认值为开启。
--print-openacc-summary	打印所有记录的OpenACC活动的摘要。
--print-openacc-trace	打印所有记录的OpenACC活动的详细跟踪信息，包括每个活动的时间戳和持续时间。
--print-openacc-constructs	包含导致OpenACC活动触发的OpenACC父构造名称。请注意，对于使用19.1版本之前PGI OpenACC运行时的应用程序，该值将始终为unknown。
--openacc-summary-mode	指定在OpenACC摘要中如何呈现活动持续时间。允许的值："exclusive" - 独占持续时间（默认）。"inclusive" - 包含持续时间。详情请参阅OpenACC摘要模式。

3.6.2. OpenACC摘要模式

nvprof 在OpenACC摘要模式（通过--print-openacc-summary启用）中支持两种方式来展示OpenACC活动持续时间："独占"和"包含"。

• 包含模式：在此模式下，所有持续时间表示活动的实际运行时间。这包括该活动及其所有子活动（被调用者）所花费的时间。

• 独占模式：在此模式下，所有持续时间仅表示该活动单独消耗的时间。这包括该活动自身耗时，但不包含其所有子活动（被调用方）的运行时间。

举个例子，考虑OpenACC的acc_compute_construct，它本身会调用acc_enqueue_launch向设备启动内核，以及调用acc_implicit_wait等待该内核完成。在"包含"模式下，acc_compute_construct的持续时间将包含acc_enqueue_launch和acc_implicit_wait所花费的时间。而在"排除"模式下，这两个持续时间会被扣除。在概要分析中，这有助于识别较长的acc_compute_construct是表示较高的启动开销还是较长的等待(同步)时间。

3.7. OpenMP

在64位Linux平台上，nvprof支持记录OpenMP活动

在nvprof中进行OpenMP性能分析需要目标应用程序使用支持OpenMP工具接口(OMPT)的运行时环境。(使用LLVM代码生成器的PGI 19.1或更高版本支持OMPT)。

尽管记录OpenMP活动仅支持x86_64 Linux系统，但在nvprof支持的所有平台上都可以导入和查看先前生成的性能分析数据。

下面展示了一个OpenMP汇总输出的示例：

==20854== NVPROF is profiling process 20854, command: ./openmp

==20854== Profiling application: ./openmp
==20854== Profiling result:
No kernels were profiled.
No API activities were profiled.
            Type  Time(%)      Time     Calls       Avg       Min       Max  Name
  OpenMP (incl):   99.97%  277.10ms        20  13.855ms  13.131ms  18.151ms  omp_parallel
                    0.03%  72.728us        19  3.8270us  2.9840us  9.5610us  omp_idle
                    0.00%  7.9170us         7  1.1310us  1.0360us  1.5330us  omp_wait_barrier

3.7.1. OpenMP选项

表2包含了nvprof中与OpenMP性能分析相关的命令行选项。

表2. OpenMP选项

选项	描述
--print-openmp-summary	打印所有记录的OpenMP活动的摘要。

4. 远程性能分析

远程性能分析是指从远程系统收集性能分析数据，这些数据将在主机系统上查看和分析。有两种方法可以进行远程性能分析。您可以直接从nsight或Visual Profiler对远程应用程序进行分析。或者，您可以在远程系统上使用nvprof收集性能分析数据，然后在主机系统上使用nvvp查看和分析数据。

4.1. 使用Visual Profiler进行远程性能分析

本节介绍如何利用nsight和Visual Profiler的远程分析功能进行远程性能剖析。

Nsight Eclipse Edition支持完整的远程开发功能，包括远程构建、调试和分析。利用这些功能，您可以创建项目和启动配置，从而远程分析您的应用程序。更多信息请参阅Nsight Eclipse Edition文档。

Visual Profiler 还支持远程性能分析。如下图所示，在创建新会话或编辑现有会话时，您可以指定待分析的应用位于远程系统。配置会话使用远程应用后，您可以执行所有分析器功能，操作方式与本地应用相同，包括时间线生成、引导式分析以及事件和指标收集。

要使用Visual Profiler进行远程性能分析，您必须在主机和远程系统上安装相同版本的CUDA工具包。主机系统不需要配备NVIDIA GPU，但需确保主机系统安装的CUDA工具包支持目标设备。主机和远程系统可以运行不同的操作系统或具有不同的CPU架构。目前仅支持运行Linux的远程系统。远程系统必须能够通过SSH访问。

4.1.1. 单跳远程性能分析

在某些远程分析设置中，运行实际CUDA程序的机器无法从运行Visual Profiler的机器直接访问。这两台机器通过一个中间机器连接，我们称之为登录节点。

主机是运行Visual Profiler的机器。

登录节点是运行单跳分析脚本的地方。我们只需要在这台机器上安装ssh、scp和perl。

计算节点是实际运行和剖析CUDA应用程序的地方。生成的剖析数据将被复制到登录节点，以便主机上的Visual Profiler可以使用这些数据。

要配置单跳性能分析，您需要完成以下一次性设置：

1. 将单跳性能分析Perl脚本复制到登录节点上。

2. 在Visual Profiler中，将登录节点添加为新的远程连接。

3. 在Visual Profiler的新会话向导中，使用配置按钮打开工具包配置窗口。在此处，通过单选按钮选择自定义脚本选项，并浏览指向登录节点上的Perl脚本。

完成此设置后，您可以像在任何远程机器上一样分析应用程序。所有数据在登录节点和计算节点之间的传输都是透明且自动进行的。

4.2. 使用 nvprof 进行远程性能分析

本节介绍如何通过在远程系统上手动运行nvprof，然后将收集的性能分析数据导入Visual Profiler来进行远程性能分析。

4.2.1. 在远程系统上收集数据

使用nvprof和Visual Profiler可以解决三种常见的远程性能分析用例。

时间线

第一个用例是收集在远程系统上执行的应用程序时间线。时间线的收集方式应尽可能准确地反映应用程序的行为。要在远程系统上收集时间线，请执行以下操作。有关nvprof选项的更多信息，请参阅​nvprof。

$ nvprof --export-profile timeline.prof <app> <app args>

性能分析数据将被收集到timeline.prof文件中。您需要将此文件复制回主机系统，然后按照下一节所述将其导入Visual Profiler。

指标与事件

第二个用例是为应用程序中所有已收集时间线的内核收集事件或指标。为所有内核收集事件或指标会显著改变应用程序的整体性能特征，因为所有内核执行将在GPU上串行化。尽管整体应用程序性能会发生变化，但单个内核的事件或指标值将是准确的，因此您可以将收集的事件和指标值合并到先前收集的时间线上，以获取应用程序行为的准确视图。要收集事件或指标，您可以使用--events或--metrics标志。以下示例仅使用--metrics标志来收集两个指标。

$ nvprof --metrics achieved_occupancy,ipc -o metrics.prof <app> <app args>

您可以为每次nvprof调用收集任意数量的事件和指标，并且可以多次调用nvprof来收集多个metrics.prof文件。为了获得准确的性能分析结果，您的应用程序必须符合应用程序要求中详述的条件。

性能分析数据将被收集到metrics.prof文件中。您需要将这些文件复制回主机系统，然后按照下一节所述将其导入Visual Profiler。

单个内核分析

第三种常见的远程性能分析使用场景是为单个内核收集分析系统所需的指标数据。当这些数据导入Visual Profiler后，分析系统将能够对该内核进行分析，并报告该内核的优化机会。要在远程系统上收集分析数据，请执行以下操作。需要注意的是，--kernels选项必须出现在--analysis-metrics选项之前，这样才能确保仅收集由kernel specifier指定的内核指标。有关--kernels选项的更多信息，请参阅性能分析范围。

$ nvprof --kernels <kernel specifier> --analysis-metrics -o analysis.prof <app> <app args>

性能分析数据将被收集到analysis.prof文件中。您需要将此文件复制回主机系统，然后按照下一节所述将其导入Visual Profiler。

4.2.2. 查看与分析数据

收集的性能分析数据可以通过导入到主机系统的Visual Profiler中进行查看和分析。有关导入的更多信息，请参阅导入会话。

时间线、指标与事件

要查看收集的时间线数据，可以按照导入单进程nvprof会话中描述的方法将timeline.prof文件导入Visual Profiler。如果还为应用程序收集了指标或事件数据，可以将相应的metrics.prof文件与时间线一起导入Visual Profiler，这样为每个内核收集的事件和指标就会与时间线中对应的内核关联起来。

针对单个内核的引导分析

要查看单个内核收集的分析数据，可以按照导入单进程nvprof会话中所述，将analysis.prof文件导入Visual Profiler。analysis.prof必须单独导入。时间轴将仅显示我们在数据收集期间指定的单个内核。导入后，可以使用引导分析系统来探索该内核的优化机会。

5. NVIDIA工具扩展

NVIDIA工具扩展(NVTX)是一个基于C语言的应用编程接口(API)，用于在应用程序中标注事件、代码范围及资源。集成NVTX的应用程序可以使用Visual Profiler捕获并可视化这些事件和范围。NVTX API提供两大核心服务：

1. 追踪CPU事件和时间范围。

2. 操作系统和CUDA资源的命名。

NVTX可以快速集成到应用程序中。以下示例程序展示了标记事件、范围事件和资源命名的使用。

void Wait(int waitMilliseconds) {
  nvtxNameOsThread(“MAIN”);
  nvtxRangePush(__FUNCTION__);
  nvtxMark("Waiting...");
  Sleep(waitMilliseconds);
  nvtxRangePop();
}

int main(void) {
  nvtxNameOsThread("MAIN");
  nvtxRangePush(__FUNCTION__);
  Wait();
  nvtxRangePop();
}

5.1. NVTX API概览

文件

核心NVTX API定义在文件nvToolsExt.h中，而NVTX接口的CUDA特定扩展则定义在nvToolsExtCuda.h和nvToolsExtCudaRt.h中。在Linux系统上，NVTX共享库名为libnvToolsExt.so；在macOS系统上，共享库名为libnvToolsExt.dylib；在Windows系统上，库文件(.lib)和运行时组件(.dll)的命名格式为nvToolsExt[bitness=32|64]_[version].{dll|lib}。

函数调用

所有NVTX API函数均以nvtx名称前缀开头，可能以三种后缀之一结尾：A、W或Ex。带有这些后缀的NVTX函数存在多个变体，通过不同的参数编码方式实现相同的核心功能。根据NVTX库的版本不同，可用的编码方式可能包括ASCII（A）、Unicode（W）或事件结构（Ex）。

NVTX的CUDA实现仅支持API的ASCII (A)和事件结构(Ex)变体，Unicode (W)版本不受支持，调用时不会产生任何效果。

返回值

部分NVTX函数被定义为具有返回值。例如，nvtxRangeStart()函数返回一个唯一的范围标识符，而nvtxRangePush()函数输出当前堆栈级别。建议不要在检测应用程序中将返回值用作条件代码的一部分。不同NVTX库实现返回的值可能有所不同，因此如果依赖这些返回值，可能会导致在某些工具中正常工作而在其他工具中失败。

5.2. NVTX API 事件

标记用于描述应用程序执行过程中在特定时间发生的事件，而范围则详细说明这些事件发生的时间跨度。这些信息会与所有其他捕获的数据一起呈现，从而更容易理解收集到的信息。所有标记和范围都通过消息字符串进行标识。标记和范围API的Ex版本还允许使用事件属性结构将类别、颜色和有效载荷属性与事件相关联。

5.2.1. NVTX标记

标记用于描述瞬时事件。标记可以包含文本消息或使用事件属性结构指定额外信息。使用nvtxMarkA创建包含ASCII消息的标记。使用nvtxMarkEx()创建包含由事件属性结构指定的附加属性的标记。在NVTX的CUDA实现中不支持nvtxMarkW()函数，如果调用将无效。

代码示例

nvtxMarkA("My mark");

nvtxEventAttributes_t eventAttrib = {0};
eventAttrib.version = NVTX_VERSION;
eventAttrib.size = NVTX_EVENT_ATTRIB_STRUCT_SIZE;
eventAttrib.colorType = NVTX_COLOR_ARGB;
eventAttrib.color = COLOR_RED;
eventAttrib.messageType = NVTX_MESSAGE_TYPE_ASCII;
eventAttrib.message.ascii = "my mark with attributes";
nvtxMarkEx(&eventAttrib);

5.2.2. NVTX范围开始/结束

开始/结束范围用于表示任意可能非嵌套的时间跨度。范围的开始和结束可以发生在不同线程上。范围可以包含文本消息或使用事件属性结构指定额外信息。使用nvtxRangeStartA()创建包含ASCII消息的标记。使用nvtxRangeStartEx()创建包含由事件属性结构指定的附加属性的范围。在NVTX的CUDA实现中不支持nvtxRangeStartW()函数，调用时无效。为了关联开始/结束对，会创建一个唯一的关联ID，该ID从nvtxRangeStartA()或nvtxRangeStartEx()返回，然后传递给nvtxRangeEnd()。

代码示例

// non-overlapping range
nvtxRangeId_t id1 = nvtxRangeStartA("My range");
nvtxRangeEnd(id1);

nvtxEventAttributes_t eventAttrib = {0};
eventAttrib.version = NVTX_VERSION;
eventAttrib.size = NVTX_EVENT_ATTRIB_STRUCT_SIZE;
eventAttrib.colorType = NVTX_COLOR_ARGB;
eventAttrib.color = COLOR_BLUE;
eventAttrib.messageType = NVTX_MESSAGE_TYPE_ASCII;
eventAttrib.message.ascii = "my start/stop range";
nvtxRangeId_t id2 = nvtxRangeStartEx(&eventAttrib);
nvtxRangeEnd(id2);

// overlapping ranges
nvtxRangeId_t r1 = nvtxRangeStartA("My range 0");
nvtxRangeId_t r2 = nvtxRangeStartA("My range 1");
nvtxRangeEnd(r1);
nvtxRangeEnd(r2);

5.2.3. NVTX范围压入/弹出

push/pop范围用于表示嵌套的时间跨度。范围的开始必须与范围的结束位于同一线程上。一个范围可以包含文本消息或使用事件属性结构指定额外信息。使用nvtxRangePushA()创建包含ASCII消息的标记。使用nvtxRangePushEx()创建包含由事件属性结构指定的附加属性的范围。在NVTX的CUDA实现中不支持nvtxRangePushW()函数，如果调用则无效。每个push函数返回正在开始的范围的从零开始的深度。nvtxRangePop()函数用于结束线程上最近push的范围。nvtxRangePop()返回正在结束的范围的从零开始的深度。如果pop没有匹配的push，则返回负值以指示错误。

代码示例

nvtxRangePushA("outer");
nvtxRangePushA("inner");
nvtxRangePop();  // end "inner" range
nvtxRangePop();  // end "outer" range

nvtxEventAttributes_t eventAttrib = {0};
eventAttrib.version = NVTX_VERSION;
eventAttrib.size = NVTX_EVENT_ATTRIB_STRUCT_SIZE;
eventAttrib.colorType = NVTX_COLOR_ARGB;
eventAttrib.color = COLOR_GREEN;
eventAttrib.messageType = NVTX_MESSAGE_TYPE_ASCII;
eventAttrib.message.ascii = "my push/pop range";
nvtxRangePushEx(&eventAttrib);
nvtxRangePop();

5.2.4. 事件属性结构

事件属性结构nvtxEventAttributes_t用于描述事件的属性。该结构的布局由特定版本的NVTX定义，并可能在不同版本的Tools Extension库中发生变化。

属性

标记和范围可以使用属性来为事件提供额外信息，或指导工具对数据的可视化呈现。每个属性都是可选的，如果未指定，属性将回退到默认值。

Message

message字段可用于指定一个可选字符串。调用者必须同时设置messageType和message字段。默认值为NVTX_MESSAGE_UNKNOWN。NVTX的CUDA实现仅支持ASCII类型的消息。

Category

category属性是一个用户可控的ID，可用于对事件进行分组。该工具可能使用category ID来改进筛选功能，或用于事件分组。默认值为0。

Color

color属性用于在工具中帮助视觉识别事件。调用者必须同时设置colorType和color字段。

Payload

payload属性可用于为标记和范围提供额外数据。范围事件只能在范围的起始处指定值。调用方必须为payloadType和payload字段指定有效值。

初始化

调用方在使用属性时，应始终执行以下三项任务：

• 将结构体清零

• 设置版本字段

• 设置大小字段

将结构体清零会将所有事件属性类型和值设置为默认值。版本和大小字段由NVTX用于处理多个版本的属性结构体。

建议调用方使用以下方法来初始化事件属性结构。

nvtxEventAttributes_t eventAttrib = {0};
eventAttrib.version = NVTX_VERSION;
eventAttrib.size = NVTX_EVENT_ATTRIB_STRUCT_SIZE;
eventAttrib.colorType = NVTX_COLOR_ARGB;
eventAttrib.color = ::COLOR_YELLOW;
eventAttrib.messageType = NVTX_MESSAGE_TYPE_ASCII;
eventAttrib.message.ascii = "My event";
nvtxMarkEx(&eventAttrib);

5.2.5. NVTX同步标记

NVTX同步模块提供了一系列功能，用于支持跟踪目标应用程序的额外同步细节。为操作系统同步原语命名可帮助用户更好地理解通过跟踪同步API收集的数据。此外，对用户自定义的同步对象进行标注，能让用户向工具表明何时正在构建不依赖操作系统提供行为、而是采用原子操作或自旋锁等技术实现的自主同步系统。

注意

Windows 系统不支持同步标记功能。

代码示例

class MyMutex
{
    volatile long bLocked;
    nvtxSyncUser_t hSync;

public:
    MyMutex(const char* name, nvtxDomainHandle_t d) {
        bLocked = 0;
        nvtxSyncUserAttributes_t attribs = { 0 };
        attribs.version = NVTX_VERSION;
        attribs.size = NVTX_SYNCUSER_ATTRIB_STRUCT_SIZE;
        attribs.messageType = NVTX_MESSAGE_TYPE_ASCII;
        attribs.message.ascii = name;
        hSync = nvtxDomainSyncUserCreate(d, &attribs);
    }

    ~MyMutex() {
        nvtxDomainSyncUserDestroy(hSync);
    }

    bool Lock() {
        nvtxDomainSyncUserAcquireStart(hSync);

        //atomic compiler intrinsic
        bool acquired = __sync_bool_compare_and_swap(&bLocked, 0, 1);

        if (acquired) {
            nvtxDomainSyncUserAcquireSuccess(hSync);
        }
        else {
            nvtxDomainSyncUserAcquireFailed(hSync);
        }
        return acquired;
    }

    void Unlock() {
        nvtxDomainSyncUserReleasing(hSync);
        bLocked = false;
    }
};

5.3. NVTX 域

域(Domain)允许开发者限定注解的作用范围。默认情况下，所有事件和注解都位于默认域中。开发者可以注册额外的域，这样就能限定标记和范围以避免冲突。

函数 nvtxDomainCreateA() 或 nvtxDomainCreateW() 用于创建命名域。

每个领域都维护自己的

• 分类

• 线程范围堆栈

• 已注册的字符串

函数 nvtxDomainDestroy() 标记域的结束。销毁域会注销并销毁与其关联的所有对象，例如已注册的字符串、资源对象、命名类别和已启动的范围。

注意

Windows 系统不支持域功能。

代码示例

nvtxDomainHandle_t domain = nvtxDomainCreateA("Domain_A");

nvtxMarkA("Mark_A");
nvtxEventAttributes_t attrib = {0};
attrib.version               = NVTX_VERSION;
attrib.size                  = NVTX_EVENT_ATTRIB_STRUCT_SIZE;
attrib.message.ascii         = "Mark A Message";
nvtxDomainMarkEx(NULL, &attrib);

nvtxDomainDestroy(domain);

5.4. NVTX资源命名

NVTX资源命名功能允许将自定义名称与主机操作系统线程及CUDA资源（如设备、上下文和流）相关联。通过NVTX分配的名称会在Visual Profiler中显示。

操作系统线程

nvtxNameOsThreadA() 函数用于为主机操作系统线程命名。nvtxNameOsThreadW() 函数在NVTX的CUDA实现中不受支持，调用时不会产生任何效果。以下示例展示了如何为当前主机操作系统线程命名。

// Windows
nvtxNameOsThread(GetCurrentThreadId(), "MAIN_THREAD");

// Linux/Mac
nvtxNameOsThread(pthread_self(), "MAIN_THREAD");

CUDA运行时资源

nvtxNameCudaDeviceA() 和 nvtxNameCudaStreamA() 函数分别用于为CUDA设备和流对象命名。nvtxNameCudaDeviceW() 和 nvtxNameCudaStreamW() 函数在NVTX的CUDA实现中不受支持，调用时不会产生任何效果。nvtxNameCudaEventA() 和 nvtxNameCudaEventW() 函数同样不受支持。以下示例展示了如何为CUDA设备和流命名。

nvtxNameCudaDeviceA(0, "my cuda device 0");

cudaStream_t cudastream;
cudaStreamCreate(&cudastream);
nvtxNameCudaStreamA(cudastream, "my cuda stream");

CUDA 驱动程序资源

函数nvtxNameCuDeviceA()、nvtxNameCuContextA()和nvtxNameCuStreamA()分别用于命名CUDA驱动程序的设备、上下文和流对象。在NVTX的CUDA实现中不支持nvtxNameCuDeviceW()、nvtxNameCuContextW()和nvtxNameCuStreamW()函数，调用这些函数不会产生任何效果。同样也不支持nvtxNameCuEventA()和nvtxNameCuEventW()函数。以下示例展示了如何为CUDA设备、上下文和流命名。

CUdevice device;
cuDeviceGet(&device, 0);
nvtxNameCuDeviceA(device, "my device 0");

CUcontext context;
cuCtxCreate(&context, 0, device);
nvtxNameCuContextA(context, "my context");

cuStream stream;
cuStreamCreate(&stream, 0);
nvtxNameCuStreamA(stream, "my stream");

5.5. NVTX 字符串注册

注册字符串旨在通过降低检测开销来提高性能。字符串可以注册一次，并且可以在API允许的情况下传递句柄来代替字符串。

nvtxDomainRegisterStringA() 函数用于注册字符串。在NVTX的CUDA实现中不支持 nvtxDomainRegisterStringW() 函数，调用该函数不会产生任何效果。

nvtxDomainHandle_t domain = nvtxDomainCreateA("Domain_A");
nvtxStringHandle_t message = nvtxDomainRegisterStringA(domain, "registered string");
nvtxEventAttributes_t eventAttrib = {0};
eventAttrib.version = NVTX_VERSION;
eventAttrib.size = NVTX_EVENT_ATTRIB_STRUCT_SIZE;
eventAttrib.messageType = NVTX_MESSAGE_TYPE_REGISTERED;
eventAttrib.message.registered = message;

6. MPI性能分析

6.1. 使用NVTX自动进行MPI标注

您可以使用NVTX标记对MPI调用进行注释，以便进行分析、跟踪和可视化。虽然为每个MPI调用添加NVTX标记可能很繁琐，但有两种方法可以自动完成此操作：

内置注解

nvprof内置了一个选项，支持两种MPI实现 - OpenMPI和MPICH。如果您的系统安装了其中任何一种，您可以使用--annotate-mpi选项并指定已安装的MPI实现。

如果使用此选项，每当应用程序调用MPI时，nvprof都会生成NVTX标记。此内置选项仅对同步MPI调用进行标注。此外，我们使用NVTX重命名当前线程和当前设备对象以显示MPI等级。

例如，如果您已安装OpenMPI，可以使用以下命令为应用程序添加注释：

$ mpirun -np 2 nvprof --annotate-mpi openmpi ./my_mpi_app

输出结果将类似于以下内容：

NVTX result:
  Thread "MPI Rank 0" (id = 583411584)
    Domain "<unnamed>"
      Range "MPI_Reduce"
  Type  Time(%)      Time     Calls       Avg       Min       Max  Name
Range:  100.00%  16.652us         1  16.652us  16.652us  16.652us  MPI_Reduce
...

      Range "MPI_Scatter"
  Type  Time(%)      Time     Calls       Avg       Min       Max  Name
Range:  100.00%  3.0320ms         1  3.0320ms  3.0320ms  3.0320ms  MPI_Scatter
...

NVTX result:
  Thread "MPI Rank 1" (id = 199923584)
    Domain "<unnamed>"
      Range "MPI_Reduce"
  Type  Time(%)      Time     Calls       Avg       Min       Max  Name
Range:  100.00%  21.062us         1  21.062us  21.062us  21.062us  MPI_Reduce
...

      Range "MPI_Scatter"
  Type  Time(%)      Time     Calls       Avg       Min       Max  Name
Range:  100.00%  85.296ms         1  85.296ms  85.296ms  85.296ms  MPI_Scatter
...

自定义注解

如果您的系统安装了nvprof不支持的MPI版本，或者您希望更灵活地控制哪些MPI函数被标注以及NVTX标记的生成方式，您可以创建自己的标注库，并通过环境变量LD_PRELOAD来拦截MPI调用并用NVTX标记进行封装。

您可以方便地使用位于此处的文档和开源脚本创建这个注释库。

6.2. 手动MPI性能分析

要使用nvprof收集各个MPI进程的性能分析数据，您需要让nvprof将其输出写入唯一文件。在CUDA 5.0及更早版本中，建议使用脚本来实现这一点。但现在，您可以通过利用nvprof命令的--export-profile参数中的%h、%p和%q{ENV}特性轻松实现。以下是使用Open MPI的运行示例。

$ mpirun -np 2 -host c0-0,c0-1 nvprof -o output.%h.%p.%q{OMPI_COMM_WORLD_RANK} ./my_mpi_app

或者，可以利用新特性，通过向nvprof传递--profile-all-processes参数来启用目标节点的性能分析功能。为此，您需要先登录到要分析的节点，然后在该节点上启动nvprof。

$ nvprof --profile-all-processes -o output.%h.%p

然后你就可以像平常一样运行MPI任务了。

$ mpirun -np 2 -host c0-0,c0-1 ./my_mpi_app

在运行--profile-all-processes的节点上，所有进程都将自动被分析。分析数据将被写入输出文件。请注意，%q{OMPI_COMM_WORLD_RANK}选项在此不适用，因为在运行nvprof的shell环境中无法获取该环境变量。

从CUDA 7.5开始，您可以通过传递类似"MPI Rank %q{OMPI_COMM_WORLD_RANK}"的字符串作为参数，使用--process-name和--context-name选项来命名线程和CUDA上下文，就像命名输出文件一样。当用户将多个文件导入Visual Profiler的同一时间线时，此功能有助于识别与特定排名相关联的资源。

$ mpirun -np 2 -host c0-0,c0-1 nvprof --process-name "MPI Rank %q{OMPI_COMM_WORLD_RANK}" --context-name "MPI Rank %q{OMPI_COMM_WORLD_RANK}" -o output.%h.%p.%q{OMPI_COMM_WORLD_RANK} ./my_mpi_app

6.3. 延伸阅读

关于与nvprof一起使用的附加参数类型的详细信息，可以在多进程分析和重定向输出部分找到。关于如何使用Visual Profiler查看数据的更多信息，请参阅导入单进程nvprof会话和导入多进程nvprof会话部分。

博客文章Profiling MPI Applications展示了如何利用CUDA 6.5引入的nvprof新输出文件命名功能以及NVTX库来命名各类资源，以分析MPI应用程序的性能。

博客文章Track MPI Calls in the Visual Profiler展示了Visual Profiler如何结合PMPI和NVTX，为您提供应用程序中MPI调用与GPU交互的有趣洞察。

7. MPS性能分析

您可以使用多进程服务(MPS)配合nvprof收集CUDA应用程序的性能分析数据，然后通过将数据导入Visual Profiler来查看时间线。

7.1. 使用Visual Profiler进行MPS性能分析

Visual Profiler可以在特定的MPS客户端或所有MPS客户端上运行。时间线分析可以对同一服务器上的所有MPS客户端进行。事件或指标分析会导致串行化——每次只有一个MPS客户端会执行。

要使用MPS分析CUDA应用程序：

1. 启动MPS守护进程。详情请参阅MPS文档。

nvidia-cuda-mps-control -d

2. 在Visual Profiler中，通过主菜单"文件->新建会话"打开"新建会话"向导。从下拉菜单中选择"分析所有进程"选项，点击"下一步"然后点击"完成"。

3. 在单独的终端中运行应用程序

4. 要结束性能分析，请在Visual Profiler中点击进度对话框上的"Cancel"按钮。

请注意，性能分析输出中还包含名为nvidia-cuda-mps-server的CUDA MPS服务器进程数据。

7.2. 使用nvprof进行MPS性能分析

nvprof 可以在特定的MPS客户端或所有MPS客户端上运行。时间线分析可以对同一服务器上的所有MPS客户端进行。事件或指标分析会导致串行化 - 每次只有一个MPS客户端会执行。

要使用MPS分析CUDA应用程序：

1. 启动MPS守护进程。详情请参阅MPS文档。

nvidia-cuda-mps-control -d

2. 运行 nvprof 时使用 --profile-all-processes 参数，并为每个进程生成单独的输出文件，可使用 --export-profile 参数的 %p 功能。注意 %p 将被进程ID替换。

nvprof --profile-all-processes -o output_%p

3. 在单独的终端中运行应用程序

4. 通过输入“Ctrl-c”退出nvprof。

请注意，性能分析输出还包含进程名为nvidia-cuda-mps-server的CUDA MPS服务器进程数据。

7.3. 在Visual Profiler中查看nvprof MPS时间线

使用多进程导入选项导入每个进程生成的nvprof数据文件。请参阅导入多进程会话部分。

下图展示了三个进程的MPS时间线视图。MPS上下文在时间线行标签中标识为Context MPS。请注意，计算和内核时间线行显示有三个内核重叠。

8. 依赖关系分析

依赖分析功能能够优化程序运行时性能，并提升利用多CPU线程和CUDA流的应用程序并发性。该功能可以计算特定执行的关键路径，检测等待时间，并检查在不同线程或流中执行的函数之间的依赖关系。

8.1. 背景

nvprof和Visual Profiler中的依赖关系分析基于应用程序的执行跟踪。跟踪记录会捕获所有相关活动，如API函数调用或CUDA内核及其时间戳和持续时间。根据这个执行跟踪以及不同线程/流上这些活动之间的依赖关系模型，可以构建一个依赖关系图。该图中建模的典型依赖关系包括：CUDA内核必须在其对应的启动API调用之后才能开始执行，或者阻塞式CUDA流同步调用必须在该流中所有先前排队的工作完成后才能返回。这些依赖关系由CUDA API合约定义。

根据此依赖关系图和API模型，可以计算等待状态。等待状态是指某项活动（如API函数调用）因等待另一个线程或流中的事件而被阻塞的持续时间。以之前的流同步示例为例，同步API调用在必须等待相应CUDA流中的任何GPU活动时会被阻塞。了解等待状态发生的位置以及函数被阻塞的时长，有助于识别优化机会以实现应用程序中更高级别的并发。

除了单独的等待状态外，通过捕获的事件图中的关键路径能够精确定位那些影响应用程序总运行时间的函数调用、内核和内存拷贝操作。关键路径是事件图中不包含等待状态的最长路径，也就是说优化该路径上的活动可以直接缩短执行时间。

8.2. 指标

等待时间

等待状态指的是某个活动（如API函数调用）被阻塞以等待另一个线程或流中的事件所持续的时间。等待时间是执行流之间负载不平衡的一个指标。在下面的示例中，阻塞的CUDA同步API调用正在等待它们各自的内核在GPU上完成执行。与其立即等待，不如尝试将内核执行与具有类似运行时间的并发CPU工作重叠，从而减少任何计算设备（CPU或GPU）被阻塞的时间。

关键路径时间

关键路径是事件图中不包含等待状态的最长路径，即优化该路径上的活动可以直接提升执行时间。关键路径上耗时较高的活动对应用程序运行时间有较大的直接影响。在下方示例图中，copy_kernel位于关键路径上，因为CPU在cudeDeviceSynchronize中被阻塞等待其完成。减少该内核运行时间可使CPU更早从API调用返回并继续程序执行。另一方面，jacobi_kernel与CPU工作完全重叠，即同步API调用在内核完成后才被触发。由于没有执行流在等待该内核完成，减少其持续时间可能不会改善整体应用程序运行时间。

8.3. 支持

以下编程API目前支持依赖分析

• CUDA运行时和驱动程序API

• POSIX线程（Pthreads）、POSIX互斥锁和条件变量

依赖关系分析功能可在Visual Profiler和nvprof中使用。在非引导式应用分析中可以选择依赖关系分析阶段，时间轴视图新增了依赖关系分析控件。关于如何在nvprof中使用此功能，请参阅依赖关系分析章节。

8.4. 限制

不同线程与CUDA流之间的依赖关系和等待时间分析仅考虑各自支持的API合约中声明的执行依赖关系。这尤其不包括由于资源争用导致的同步。例如，即使具体GPU只有一个复制引擎，排入独立CUDA流的异步内存复制也不会被标记为依赖关系。此外，该分析不考虑使用不受支持API进行的同步。例如，CPU线程主动轮询某个内存位置的值（忙等待）将不会被视作被另一个并发活动阻塞。

依赖分析对使用CUDA动态并行(CDP)的应用程序支持有限。CDP内核可以从GPU调用CUDA API，但这些调用无法通过CUPTI活动API追踪。因此，分析无法确定CDP内核的完整依赖关系和等待时间。不过，它会利用CDP内核之间的父子启动依赖关系。因此，关键路径总会包含每个由主机启动的内核的最后一个CDP内核。

目前不支持POSIX信号量API。

依赖分析不支持API函数cudaLaunchCooperativeKernelMultiDevice或cuLaunchCooperativeKernelMultiDevice。通过这些API函数启动的内核可能无法被正确追踪。

9. 指标参考

本节详细介绍了可通过nvprof和Visual Profiler收集的各项指标。若作用域值为"Single-context"，表示该指标仅在GPU上执行单一上下文（CUDA或图形）时能准确采集；若为"Multi-context"，则表示多上下文并行执行时仍可准确采集；而"Device"作用域值表示该指标将在设备层级采集，即包含GPU上所有执行上下文的数据。请注意，内核模式下采集的NVLink指标会呈现"Single-context"特性。

9.1. 能力指标 5.x

具备计算能力5.x的设备支持下表中显示的指标。请注意，某些指标的"多上下文"范围仅适用于特定设备。这些指标在"范围"列中标注为"多上下文^*"。详情请参阅表格底部的注释。

表4. 能力5.x指标

指标名称	描述	范围
achieved_occupancy	每个活跃周期内平均活跃warp数与多处理器支持的最大warp数之比	多上下文
atomic_transactions	全局内存原子操作与归约事务	多上下文
atomic_transactions_per_request	每条原子操作和归约指令执行的全局内存原子操作和归约事务的平均数量	多上下文
branch_efficiency	非发散分支占总分支的比率，以百分比表示	多上下文
cf_executed	已执行的控制流指令数量	多上下文
cf_fu_utilization	执行控制流指令的多处理器功能单元利用率水平，范围为0到10	多上下文
cf_issued	已发出的控制流指令数量	多上下文
double_precision_fu_utilization	执行双精度浮点指令的多处理器功能单元利用率水平，范围为0到10	Multi-context
dram_read_bytes	从DRAM读取到L2缓存的总字节数。此指标适用于计算能力5.0和5.2的设备。	多上下文*
dram_read_throughput	设备内存读取吞吐量。此指标适用于计算能力5.0和5.2的硬件。	多上下文*
dram_read_transactions	设备内存读取事务。此指标适用于计算能力5.0和5.2的设备。	多上下文*
dram_utilization	设备内存利用率水平，相对于峰值利用率，范围为0到10	多上下文*
dram_write_bytes	从L2缓存写入DRAM的总字节数。此指标适用于计算能力5.0和5.2版本。	多上下文*
dram_write_throughput	设备内存写入吞吐量。此指标适用于计算能力5.0和5.2版本。	多上下文*
dram_write_transactions	设备内存写入事务。此功能适用于计算能力5.0和5.2版本。	多上下文*
ecc_throughput	从L2到DRAM的ECC吞吐量。此指标适用于计算能力5.0和5.2。	多上下文*
ecc_transactions	L2缓存与DRAM之间的ECC事务数量。该指标适用于计算能力5.0和5.2的设备。	多上下文*
eligible_warps_per_cycle	每个活跃周期内符合条件可发出的平均warp数量	多上下文
flop_count_dp	非预测线程执行的双精度浮点运算次数（加法、乘法和乘加运算）。每个乘加运算会计为2次运算。	多上下文
flop_count_dp_add	非谓词化线程执行的双精度浮点加法运算次数。	多上下文
flop_count_dp_fma	非预测线程执行的双精度浮点乘加运算次数。每个乘加运算计数为1。	多上下文
flop_count_dp_mul	非预测线程执行的双精度浮点乘法运算次数。	多上下文
flop_count_hp	非谓词线程执行的半精度浮点运算次数（加法、乘法和乘加运算）。每个乘加运算会计为2次运算。此指标适用于计算能力5.3的设备。	多上下文*
flop_count_hp_add	由非预测线程执行的半精度浮点加法操作次数。此指标适用于计算能力5.3的设备。	多上下文*
flop_count_hp_fma	非预测线程执行的半精度浮点乘加运算次数。每个乘加运算计数为1。该指标适用于计算能力5.3的设备。	多上下文*
flop_count_hp_mul	由非预测线程执行的半精度浮点乘法运算次数。此指标适用于计算能力5.3的设备。	多上下文*
flop_count_sp	非预测线程执行的单精度浮点运算次数（加法、乘法和乘累加运算）。每个乘累加运算会计为2次运算。该计数不包含特殊运算。	Multi-context
flop_count_sp_add	非预测线程执行的单精度浮点加法操作数量。	多上下文
flop_count_sp_fma	非预测线程执行的单精度浮点乘加运算次数。每个乘加运算计数为1。	多上下文
flop_count_sp_mul	非预测线程执行的单精度浮点乘法运算次数。	多上下文
flop_count_sp_special	非谓词化线程执行的单精度浮点特殊操作数量。	多上下文
flop_dp_efficiency	实际双精度浮点运算与峰值性能的比率	多上下文
flop_hp_efficiency	实际达到的半精度浮点运算与峰值性能的比率。此指标适用于计算能力5.3的设备。	多上下文环境*
flop_sp_efficiency	实际单精度浮点运算与峰值单精度浮点运算的比率	多上下文
gld_efficiency	请求的全局内存加载吞吐量与所需全局内存加载吞吐量的比率，以百分比表示。	多上下文*
gld_requested_throughput	请求的全局内存加载吞吐量	多上下文
gld_throughput	全局内存加载吞吐量	多上下文*
gld_transactions	全局内存加载事务数	多上下文*
gld_transactions_per_request	每次全局内存加载执行的平均全局内存加载事务数。	多上下文*
global_atomic_requests	从多处理器发出的全局原子操作(Atom和Atom CAS)请求总数	多上下文
global_hit_rate	统一L1/纹理缓存中全局加载的命中率。如果内核中使用malloc，指标值可能不准确。	多上下文*
global_load_requests	从多处理器发出的全局加载请求总数	多上下文
global_reduction_requests	从多处理器发出的全局归约请求总数	多上下文
global_store_requests	从多处理器发出的全局存储请求总数。这不包括原子请求。	多上下文
gst_efficiency	请求的全局内存存储吞吐量与所需全局内存存储吞吐量的比率，以百分比表示。	多上下文*
gst_requested_throughput	请求的全局内存存储吞吐量	多上下文
gst_throughput	全局内存存储吞吐量	多上下文*
gst_transactions	全局内存存储事务数量	多上下文*
gst_transactions_per_request	每次全局内存存储操作执行的全局内存存储事务平均数量	多上下文*
half_precision_fu_utilization	执行16位浮点指令和整数指令的多处理器功能单元利用率水平，范围为0到10。此指标适用于计算能力5.3。	多上下文*
inst_bit_convert	非谓词化线程执行的位转换指令数量	多上下文
inst_compute_ld_st	由非预测线程执行的compute加载/存储指令数量	多上下文
inst_control	非谓词化线程执行的控制流指令数量（跳转、分支等）	多上下文
inst_executed	已执行的指令数量	多上下文
inst_executed_global_atomics	用于全局原子操作和原子比较交换的Warp级别指令	多上下文
inst_executed_global_loads	全局加载的Warp级别指令	多上下文
inst_executed_global_reductions	全局归约的Warp级别指令	多上下文
inst_executed_global_stores	全局存储的Warp级别指令	多上下文
inst_executed_local_loads	Warp级别的本地加载指令	多上下文
inst_executed_local_stores	Warp级别的本地存储指令	多上下文
inst_executed_shared_atomics	原子操作和原子CAS的Warp级别共享指令	多上下文
inst_executed_shared_loads	共享加载的Warp级别指令	多上下文
inst_executed_shared_stores	共享存储的Warp级别指令	多上下文
inst_executed_surface_atomics	用于表面原子和原子比较交换的Warp级指令	多上下文
inst_executed_surface_loads	执行表面加载的Warp级别指令	多上下文
inst_executed_surface_reductions	用于表面归约的Warp级别指令	多上下文
inst_executed_surface_stores	用于表面存储的Warp级指令	多上下文
inst_executed_tex_ops	纹理操作的Warp级别指令	多上下文
inst_fp_16	由非预测线程执行的半精度浮点指令数量（算术运算、比较等）。此指标适用于计算能力5.3及以上版本。	多上下文*
inst_fp_32	非预测线程执行的单精度浮点指令数量（算术运算、比较等）	多上下文
inst_fp_64	非预测线程执行的双精度浮点指令数量（算术运算、比较等）	多上下文
inst_integer	非谓词化线程执行的整数指令数量	多上下文
inst_inter_thread_communication	由非谓词化线程执行的线程间通信指令数量	多上下文
inst_issued	已发出的指令数量	多上下文
inst_misc	非谓词化线程执行的杂项指令数量	多上下文
inst_per_warp	每个warp平均执行的指令数	多上下文
inst_replay_overhead	每条执行指令的平均重放次数	多上下文
ipc	每周期执行的指令数	多上下文
issue_slot_utilization	每个周期平均至少发射一条指令的发射槽百分比	多上下文
issue_slots	使用的issue插槽数量	多上下文
issued_ipc	每周期发出的指令数	多上下文
l2_atomic_throughput	L2缓存中原子操作和归约请求的内存读取吞吐量	多上下文
l2_atomic_transactions	在L2缓存中观察到的原子操作和归约请求的内存读取事务	多上下文*
l2_global_atomic_store_bytes	从统一缓存写入L2的全局原子操作(ATOM和ATOM CAS)字节数	多上下文*
l2_global_load_bytes	从L2缓存读取的全局加载在统一缓存中未命中的字节数	多上下文*
l2_global_reduction_bytes	从统一缓存写入L2的全局归约字节数	多上下文*
l2_local_global_store_bytes	从统一缓存写入L2的本地和全局存储字节数。不包括全局原子操作。	多上下文*
l2_local_load_bytes	本地加载时因统一缓存未命中而从L2缓存读取的字节数	多上下文*
l2_read_throughput	所有读取请求在L2缓存层观测到的内存读取吞吐量	多上下文*
l2_read_transactions	在L2缓存中观察到的所有读取请求的内存读取事务	多上下文*
l2_surface_atomic_store_bytes	统一缓存与L2之间用于表面原子操作(ATOM和ATOM CAS)的数据传输字节数	多上下文*
l2_surface_load_bytes	从L2读取的统一缓存表面加载未命中字节数	多上下文*
l2_surface_reduction_bytes	从统一缓存写入L2的表面归约字节数	多上下文*
l2_surface_store_bytes	从统一缓存写入L2的表面存储字节数。不包括表面原子操作。	多上下文*
l2_tex_hit_rate	纹理缓存所有请求在L2缓存中的命中率	多上下文*
l2_tex_read_hit_rate	来自纹理缓存的所有读取请求在L2缓存中的命中率。此指标适用于计算能力5.0和5.2的设备。	多上下文*
l2_tex_read_throughput	纹理缓存读取请求在L2缓存层观测到的内存读取吞吐量	多上下文*
l2_tex_read_transactions	在L2缓存中观察到的来自纹理缓存的读取请求的内存读取事务	多上下文*
l2_tex_write_hit_rate	纹理缓存所有写入请求在L2缓存中的命中率。此指标适用于计算能力5.0和5.2的设备。	多上下文*
l2_tex_write_throughput	纹理缓存写入请求在L2缓存中观察到的内存写入吞吐量	多上下文*
l2_tex_write_transactions	L2缓存中观察到的来自纹理缓存的写请求的内存写入事务	多上下文*
l2_utilization	L2缓存利用率水平，相对于峰值利用率，范围为0到10	多上下文*
l2_write_throughput	所有写入请求在L2缓存层观察到的内存写入吞吐量	多上下文*
l2_write_transactions	所有写请求在L2缓存中观察到的内存写入事务	多上下文*
ldst_executed	已执行的本地、全局、共享和纹理内存加载与存储指令数量	多上下文
ldst_fu_utilization	执行共享加载、共享存储和常量加载指令的多处理器功能单元利用率水平，范围为0到10	多上下文
ldst_issued	已发出的本地、全局、共享和纹理内存加载与存储指令数量	多上下文
local_hit_rate	本地加载和存储的命中率	多上下文*
local_load_requests	来自多处理器的本地加载请求总数	多上下文*
local_load_throughput	本地内存加载吞吐量	多上下文*
local_load_transactions	本地内存加载事务数量	多上下文*
local_load_transactions_per_request	每次本地内存加载操作执行的平均本地内存加载事务数	多上下文*
local_memory_overhead	L1和L2缓存之间本地内存流量占总内存流量的百分比	多上下文*
local_store_requests	从多处理器发出的本地存储请求总数	多上下文*
local_store_throughput	本地内存存储吞吐量	多上下文*
local_store_transactions	本地内存存储事务数量	多上下文*
local_store_transactions_per_request	每次本地内存存储操作执行的平均本地内存存储事务数	多上下文*
pcie_total_data_received	通过PCIe接收的总数据字节数	设备
pcie_total_data_transmitted	通过PCIe传输的总数据字节数	设备
shared_efficiency	请求的共享内存吞吐量与所需共享内存吞吐量的比率，以百分比表示	多上下文*
shared_load_throughput	共享内存加载吞吐量	多上下文*
shared_load_transactions	共享内存加载事务数	多上下文*
shared_load_transactions_per_request	每次共享内存加载操作执行的平均共享内存加载事务数	多上下文*
shared_store_throughput	共享内存存储吞吐量	多上下文*
shared_store_transactions	共享内存存储事务数量	多上下文*
shared_store_transactions_per_request	每个共享内存存储操作执行的共享内存存储事务平均数量	多上下文*
shared_utilization	共享内存利用率水平，相对于峰值利用率的比例，范围为0到10	多上下文*
single_precision_fu_utilization	执行单精度浮点指令和整数指令的多处理器功能单元利用率水平，范围为0到10	多上下文
sm_efficiency	特定多处理器上至少有一个warp处于活动状态的时间百分比	多上下文*
special_fu_utilization	执行sin、cos、ex2、popc、flo等类似指令的多处理器功能单元利用率水平，范围为0到10	Multi-context
stall_constant_memory_dependency	由于立即常量缓存未命中导致的停顿百分比	多上下文
stall_exec_dependency	由于指令所需的输入尚未就绪而发生的停顿百分比	多上下文
stall_inst_fetch	由于下一条汇编指令尚未获取而导致停顿的百分比	多上下文
stall_memory_dependency	由于所需资源不可用或未充分利用，或由于给定类型的未完成请求过多而导致内存操作无法执行时发生的停顿百分比	Multi-context
stall_memory_throttle	因内存节流导致的停顿百分比	多上下文
stall_not_selected	因warp未被选中而发生的停顿百分比	多上下文
stall_other	由于其他原因导致的停顿百分比	多上下文
stall_pipe_busy	由于计算流水线繁忙导致无法执行计算操作而产生的停顿百分比	Multi-context
stall_sync	由于warp在__syncthreads()调用处被阻塞而发生的停顿百分比	多上下文
stall_texture	由于纹理子系统已完全占用或存在过多未完成请求而导致的停顿百分比	Multi-context
surface_atomic_requests	从多处理器发出的表面原子(Atom和Atom CAS)请求总数	多上下文
surface_load_requests	从多处理器发出的表面加载请求总数	多上下文
surface_reduction_requests	来自多处理器的表面缩减请求总数	多上下文
surface_store_requests	从多处理器发出的表面存储请求总数	多上下文
sysmem_read_bytes	从系统内存读取的字节数	多上下文*
sysmem_read_throughput	系统内存读取吞吐量	多上下文*
sysmem_read_transactions	系统内存读取事务数	多上下文*
sysmem_read_utilization	系统内存读取利用率水平，相对于峰值利用率，范围为0到10。此指标适用于计算能力5.0和5.2。	多上下文
sysmem_utilization	系统内存利用率水平，相对于峰值利用率，范围为0到10。此指标适用于计算能力5.0和5.2。	多上下文*
sysmem_write_bytes	写入系统内存的字节数	多上下文*
sysmem_write_throughput	系统内存写入吞吐量	多上下文*
sysmem_write_transactions	系统内存写入事务数量	多上下文*
sysmem_write_utilization	系统内存写入利用率水平，相对于峰值利用率的比例，范围为0到10。此指标适用于计算能力5.0和5.2版本。	多上下文*
tex_cache_hit_rate	统一缓存命中率	多上下文*
tex_cache_throughput	统一缓存吞吐量	多上下文*
tex_cache_transactions	统一缓存读取事务	多上下文*
tex_fu_utilization	执行全局、本地和纹理内存指令的多处理器功能单元利用率水平，范围为0到10	多上下文
tex_utilization	统一缓存相对于峰值利用率的利用率水平，范围为0到10	多上下文*
texture_load_requests	从多处理器发出的纹理加载请求总数	多上下文
warp_execution_efficiency	每个warp的平均活跃线程数与多处理器支持的每个warp最大线程数之比	多上下文
warp_nonpred_execution_efficiency	每个warp执行非谓词指令的平均活动线程数与多处理器支持的每个warp最大线程数之比	多上下文

* 该指标的“多上下文”范围仅支持计算能力为5.0和5.2的设备。

9.2. 能力指标 6.x

具备计算能力6.x的设备实现了下表中显示的指标。

表5. 能力6.x指标

指标名称	描述	范围
achieved_occupancy	每个活跃周期内平均活跃warp数与多处理器支持的最大warp数之比	多上下文
atomic_transactions	全局内存原子操作与归约事务	多上下文
atomic_transactions_per_request	每条原子操作和归约指令执行的全局内存原子操作和归约事务的平均数量	多上下文
branch_efficiency	非发散分支占总分支的比率，以百分比表示	多上下文
cf_executed	已执行的控制流指令数量	多上下文
cf_fu_utilization	执行控制流指令的多处理器功能单元利用率水平，范围为0到10	多上下文
cf_issued	已发出的控制流指令数量	多上下文
double_precision_fu_utilization	执行双精度浮点指令的多处理器功能单元利用率水平，范围为0到10	Multi-context
dram_read_bytes	从DRAM读取到L2缓存的总字节数	多上下文
dram_read_throughput	设备内存读取吞吐量。此指标适用于计算能力6.0和6.1版本。	多上下文
dram_read_transactions	设备内存读取事务。此功能适用于计算能力6.0和6.1版本。	多上下文
dram_utilization	设备内存利用率水平，相对于峰值利用率，范围为0到10	多上下文
dram_write_bytes	从L2缓存写入DRAM的总字节数	多上下文
dram_write_throughput	设备内存写入吞吐量。此指标适用于计算能力6.0和6.1版本。	多上下文
dram_write_transactions	设备内存写入事务。此功能适用于计算能力6.0和6.1版本。	多上下文
ecc_throughput	从L2到DRAM的ECC吞吐量。此指标适用于计算能力6.1的设备。	多上下文
ecc_transactions	L2缓存与DRAM之间的ECC事务数量。该指标适用于计算能力6.1的设备。	Multi-context
eligible_warps_per_cycle	每个活跃周期内符合条件可发出的平均warp数量	多上下文
flop_count_dp	非预测线程执行的双精度浮点运算次数（加法、乘法和乘加运算）。每个乘加运算会计为2次运算。	多上下文
flop_count_dp_add	非谓词化线程执行的双精度浮点加法运算次数。	多上下文
flop_count_dp_fma	非预测线程执行的双精度浮点乘加运算次数。每个乘加运算计数为1。	多上下文
flop_count_dp_mul	非预测线程执行的双精度浮点乘法运算次数。	多上下文
flop_count_hp	非预测线程执行的半精度浮点运算次数（加法、乘法和乘累加运算）。每个乘累加运算会计为2次运算。	多上下文
flop_count_hp_add	非预测线程执行的半精度浮点加法操作次数。	多上下文
flop_count_hp_fma	非预测线程执行的半精度浮点乘加运算次数。每个乘加运算计数为1。	多上下文
flop_count_hp_mul	非预测线程执行的半精度浮点乘法运算次数。	多上下文
flop_count_sp	非预测线程执行的单精度浮点运算次数（加法、乘法和乘累加运算）。每个乘累加运算会计为2次运算。该计数不包含特殊运算。	Multi-context
flop_count_sp_add	非预测线程执行的单精度浮点加法操作数量。	多上下文
flop_count_sp_fma	非预测线程执行的单精度浮点乘加运算次数。每个乘加运算计数为1。	多上下文
flop_count_sp_mul	非预测线程执行的单精度浮点乘法运算次数。	多上下文
flop_count_sp_special	非谓词化线程执行的单精度浮点特殊操作数量。	多上下文
flop_dp_efficiency	实际双精度浮点运算与峰值性能的比率	多上下文
flop_hp_efficiency	实际达到的半精度浮点运算与峰值性能的比率	多上下文
flop_sp_efficiency	实际单精度浮点运算与峰值单精度浮点运算的比率	多上下文
gld_efficiency	请求的全局内存加载吞吐量与所需全局内存加载吞吐量的比率，以百分比表示。	多上下文
gld_requested_throughput	请求的全局内存加载吞吐量	多上下文
gld_throughput	全局内存加载吞吐量	多上下文
gld_transactions	全局内存加载事务数	多上下文
gld_transactions_per_request	每次全局内存加载执行的平均全局内存加载事务数。	多上下文
global_atomic_requests	从多处理器发出的全局原子操作(Atom和Atom CAS)请求总数	多上下文
global_hit_rate	统一l1/tex缓存中全局加载的命中率。如果内核中使用malloc，指标值可能不准确。	多上下文
global_load_requests	从多处理器发出的全局加载请求总数	多上下文
global_reduction_requests	从多处理器发出的全局归约请求总数	多上下文
global_store_requests	从多处理器发出的全局存储请求总数。这不包括原子请求。	多上下文
gst_efficiency	请求的全局内存存储吞吐量与所需全局内存存储吞吐量的比率，以百分比表示。	多上下文
gst_requested_throughput	请求的全局内存存储吞吐量	多上下文
gst_throughput	全局内存存储吞吐量	多上下文
gst_transactions	全局内存存储事务数量	多上下文
gst_transactions_per_request	每次全局内存存储操作执行的全局内存存储事务平均数量	多上下文
half_precision_fu_utilization	执行16位浮点指令的多处理器功能单元利用率水平，范围为0到10	Multi-context
inst_bit_convert	非谓词化线程执行的位转换指令数量	多上下文
inst_compute_ld_st	由非预测线程执行的compute加载/存储指令数量	多上下文
inst_control	非谓词化线程执行的控制流指令数量（跳转、分支等）	多上下文
inst_executed	已执行的指令数量	多上下文
inst_executed_global_atomics	用于全局原子操作和原子比较交换的Warp级别指令	多上下文
inst_executed_global_loads	全局加载的Warp级别指令	多上下文
inst_executed_global_reductions	全局归约的Warp级别指令	多上下文
inst_executed_global_stores	全局存储的Warp级别指令	多上下文
inst_executed_local_loads	Warp级别的本地加载指令	多上下文
inst_executed_local_stores	Warp级别的本地存储指令	多上下文
inst_executed_shared_atomics	原子操作和原子CAS的Warp级别共享指令	多上下文
inst_executed_shared_loads	共享加载的Warp级别指令	多上下文
inst_executed_shared_stores	共享存储的Warp级别指令	多上下文
inst_executed_surface_atomics	用于表面原子和原子比较交换的Warp级指令	多上下文
inst_executed_surface_loads	执行表面加载的Warp级别指令	多上下文
inst_executed_surface_reductions	用于表面归约的Warp级别指令	多上下文
inst_executed_surface_stores	用于表面存储的Warp级指令	多上下文
inst_executed_tex_ops	纹理操作的Warp级别指令	多上下文
inst_fp_16	由非预测线程执行的半精度浮点指令数量（算术运算、比较等）	多上下文
inst_fp_32	非预测线程执行的单精度浮点指令数量（算术运算、比较等）	多上下文
inst_fp_64	非预测线程执行的双精度浮点指令数量（算术运算、比较等）	多上下文
inst_integer	非谓词化线程执行的整数指令数量	多上下文
inst_inter_thread_communication	由非谓词化线程执行的线程间通信指令数量	多上下文
inst_issued	已发出的指令数量	多上下文
inst_misc	非谓词化线程执行的杂项指令数量	多上下文
inst_per_warp	每个warp平均执行的指令数	多上下文
inst_replay_overhead	每条执行指令的平均重放次数	多上下文
ipc	每周期执行的指令数	多上下文
issue_slot_utilization	每个周期平均至少发射一条指令的发射槽百分比	多上下文
issue_slots	使用的issue插槽数量	多上下文
issued_ipc	每周期发出的指令数	多上下文
l2_atomic_throughput	L2缓存中原子操作和归约请求的内存读取吞吐量	多上下文
l2_atomic_transactions	在L2缓存中观察到的原子操作和归约请求的内存读取事务	多上下文
l2_global_atomic_store_bytes	从统一缓存写入L2的全局原子操作数据量（ATOM和ATOM CAS）	多上下文
l2_global_load_bytes	全局加载时因统一缓存未命中而从L2读取的字节数	多上下文
l2_global_reduction_bytes	从统一缓存写入L2的全局归约字节数	多上下文
l2_local_global_store_bytes	从统一缓存写入L2的本地和全局存储字节数。不包括全局原子操作。	多上下文
l2_local_load_bytes	本地加载时因统一缓存未命中而从L2缓存读取的字节数	多上下文
l2_read_throughput	所有读取请求在L2缓存中观察到的内存读取吞吐量	多上下文
l2_read_transactions	在L2缓存中观察到的所有读取请求的内存读取事务	多上下文
l2_surface_atomic_store_bytes	统一缓存与L2之间用于表面原子操作(ATOM和ATOM CAS)的数据传输字节数	多上下文
l2_surface_load_bytes	从L2读取的统一缓存表面加载未命中字节数	多上下文
l2_surface_reduction_bytes	从统一缓存写入L2的表面归约字节数	多上下文
l2_surface_store_bytes	从统一缓存写入L2的表面存储字节数。不包括表面原子操作。	多上下文
l2_tex_hit_rate	纹理缓存所有请求在L2缓存中的命中率	多上下文
l2_tex_read_hit_rate	来自纹理缓存的所有读取请求在L2缓存中的命中率。此指标适用于计算能力6.0和6.1版本。	Multi-context
l2_tex_read_throughput	从纹理缓存读取请求在L2缓存中观察到的内存读取吞吐量	多上下文
l2_tex_read_transactions	L2缓存中观察到的来自纹理缓存的读取请求的内存读取事务	多上下文
l2_tex_write_hit_rate	纹理缓存所有写入请求在L2缓存中的命中率。该指标适用于计算能力6.0和6.1版本。	多上下文
l2_tex_write_throughput	纹理缓存写入请求在L2缓存中观察到的内存写入吞吐量	多上下文
l2_tex_write_transactions	在L2缓存中观察到的来自纹理缓存的写请求的内存写入事务	多上下文
l2_utilization	L2缓存利用率水平，相对于峰值利用率，范围为0到10	多上下文
l2_write_throughput	所有写入请求在L2缓存中观察到的内存写入吞吐量	多上下文
l2_write_transactions	所有写请求在L2缓存中观察到的内存写入事务	多上下文
ldst_executed	已执行的本地、全局、共享和纹理内存加载与存储指令数量	多上下文
ldst_fu_utilization	执行共享加载、共享存储和常量加载指令的多处理器功能单元利用率水平，范围为0到10	多上下文
ldst_issued	已发出的本地、全局、共享和纹理内存加载与存储指令数量	多上下文
local_hit_rate	本地加载和存储的命中率	多上下文
local_load_requests	从多处理器发出的本地加载请求总数	多上下文
local_load_throughput	本地内存加载吞吐量	多上下文
local_load_transactions	本地内存加载事务数	多上下文
local_load_transactions_per_request	每次本地内存加载操作执行的平均本地内存加载事务数	多上下文
local_memory_overhead	L1和L2缓存之间本地内存流量占总内存流量的百分比	多上下文
local_store_requests	从多处理器发出的本地存储请求总数	多上下文
local_store_throughput	本地内存存储吞吐量	多上下文
local_store_transactions	本地内存存储事务数量	多上下文
local_store_transactions_per_request	每次本地内存存储操作执行的平均本地内存存储事务数	多上下文
nvlink_overhead_data_received	通过NVLink接收的开销数据与总数据的比率。此指标适用于计算能力6.0的设备。	Device
nvlink_overhead_data_transmitted	通过NVLink传输的开销数据与总数据的比率。此指标适用于计算能力6.0的设备。	设备
nvlink_receive_throughput	每秒通过NVLinks接收的字节数。此指标适用于计算能力6.0的设备。	设备
nvlink_total_data_received	通过NVLinks接收的总数据字节数（包含头部信息）。该指标适用于计算能力6.0及以上设备。	设备
nvlink_total_data_transmitted	通过NVLinks传输的总数据字节数（包含头部信息）。该指标适用于计算能力6.0及以上设备。	设备
nvlink_total_nratom_data_transmitted	通过NVLinks传输的非归约原子数据总字节数。此指标适用于计算能力6.0及以上设备。	设备
nvlink_total_ratom_data_transmitted	通过NVLinks传输的归约原子数据总字节数，此指标适用于计算能力6.0的设备。	设备
nvlink_total_response_data_received	通过NVLink接收的总响应数据字节数，响应数据包括读取请求的数据和非归约原子请求的结果。此指标适用于计算能力6.0及以上设备。	设备
nvlink_total_write_data_transmitted	通过NVLinks传输的总写入数据字节数。此指标适用于计算能力6.0及以上设备。	设备
nvlink_transmit_throughput	通过NVLinks每秒传输的字节数。此指标适用于计算能力6.0的设备。	Device
nvlink_user_data_received	通过NVLinks接收的用户数据字节数，不包括头部信息。此指标适用于计算能力6.0的设备。	设备
nvlink_user_data_transmitted	通过NVLinks传输的用户数据字节数，不包括头部信息。此指标适用于计算能力6.0的设备。	Device
nvlink_user_nratom_data_transmitted	通过NVLinks传输的非归约原子用户数据字节总数。此指标适用于计算能力6.0的设备。	设备
nvlink_user_ratom_data_transmitted	通过NVLinks传输的归约原子用户数据字节总数。该指标适用于计算能力6.0的设备。	Device
nvlink_user_response_data_received	通过NVLink接收的用户响应数据总字节数，响应数据包括读取请求的数据和非归约原子请求的结果。此指标适用于计算能力6.0及以上设备。	设备
nvlink_user_write_data_transmitted	通过NVLinks传输的用户写入数据字节数。此指标适用于计算能力6.0及以上版本。	设备
pcie_total_data_received	通过PCIe接收的总数据字节数	设备
pcie_total_data_transmitted	通过PCIe传输的总数据字节数	设备
shared_efficiency	请求的共享内存吞吐量与所需共享内存吞吐量的比率，以百分比表示	多上下文
shared_load_throughput	共享内存加载吞吐量	多上下文
shared_load_transactions	共享内存加载事务数	多上下文
shared_load_transactions_per_request	每次共享内存加载操作执行的平均共享内存加载事务数	多上下文
shared_store_throughput	共享内存存储吞吐量	多上下文
shared_store_transactions	共享内存存储事务数量	多上下文
shared_store_transactions_per_request	每个共享内存存储操作执行的平均共享内存存储事务数	多上下文
shared_utilization	共享内存利用率水平，相对于峰值利用率，范围为0到10	多上下文
single_precision_fu_utilization	执行单精度浮点指令和整数指令的多处理器功能单元利用率水平，范围为0到10	多上下文
sm_efficiency	特定多处理器上至少有一个warp处于活动状态的时间百分比	多上下文
special_fu_utilization	执行sin、cos、ex2、popc、flo等类似指令的多处理器功能单元利用率水平，范围为0到10	Multi-context
stall_constant_memory_dependency	由于立即常量缓存未命中导致的停顿百分比	多上下文
stall_exec_dependency	由于指令所需的输入尚未就绪而发生的停顿百分比	多上下文
stall_inst_fetch	由于下一条汇编指令尚未获取而导致停顿的百分比	多上下文
stall_memory_dependency	由于所需资源不可用或未充分利用，或由于给定类型的未完成请求过多而导致内存操作无法执行时发生的停顿百分比	Multi-context
stall_memory_throttle	因内存节流导致的停顿百分比	多上下文
stall_not_selected	因warp未被选中而发生的停顿百分比	多上下文
stall_other	由于其他原因导致的停顿百分比	多上下文
stall_pipe_busy	由于计算流水线繁忙导致无法执行计算操作而产生的停顿百分比	Multi-context
stall_sync	由于warp在__syncthreads()调用处被阻塞而发生的停顿百分比	多上下文
stall_texture	由于纹理子系统已完全占用或存在过多未完成请求而导致的停顿百分比	Multi-context
surface_atomic_requests	从多处理器发出的表面原子(Atom和Atom CAS)请求总数	多上下文
surface_load_requests	从多处理器发出的表面加载请求总数	多上下文
surface_reduction_requests	来自多处理器的表面缩减请求总数	多上下文
surface_store_requests	从多处理器发出的表面存储请求总数	多上下文
sysmem_read_bytes	从系统内存读取的字节数	多上下文
sysmem_read_throughput	系统内存读取吞吐量	多上下文
sysmem_read_transactions	系统内存读取事务数	多上下文
sysmem_read_utilization	系统内存读取利用率水平，相对于峰值利用率按0到10的等级表示。此指标适用于计算能力6.0和6.1版本。	多上下文
sysmem_utilization	系统内存利用率水平，相对于峰值利用率按0到10的等级表示。此指标适用于计算能力6.0和6.1版本。	Multi-context
sysmem_write_bytes	写入系统内存的字节数	多上下文
sysmem_write_throughput	系统内存写入吞吐量	多上下文
sysmem_write_transactions	系统内存写入事务数	多上下文
sysmem_write_utilization	系统内存写入利用率水平，相对于峰值利用率，范围为0到10。此指标适用于计算能力6.0和6.1。	多上下文
tex_cache_hit_rate	统一缓存命中率	多上下文
tex_cache_throughput	统一缓存吞吐量	多上下文
tex_cache_transactions	统一缓存读取事务	多上下文
tex_fu_utilization	执行全局、本地和纹理内存指令的多处理器功能单元利用率水平，范围为0到10	多上下文
tex_utilization	统一缓存相对于峰值利用率的利用率水平，范围为0到10	多上下文
texture_load_requests	从多处理器发出的纹理加载请求总数	多上下文
unique_warps_launched	已启动的warp数量。该数值不受计算抢占影响。	多上下文
warp_execution_efficiency	每个warp的平均活跃线程数与多处理器支持的每个warp最大线程数之比	多上下文
warp_nonpred_execution_efficiency	每个warp执行非谓词指令的平均活动线程数与多处理器支持的每个warp最大线程数之比	多上下文

9.3. Capability 7.x 的指标

具备计算能力7.x的设备实现了下表中展示的各项指标。(此处7.x指代7.0和7.2版本)

表6. 7.x（7.0和7.2版本）能力指标

指标名称	描述	范围
achieved_occupancy	每个活跃周期内平均活跃warp数与多处理器支持的最大warp数之比	多上下文
atomic_transactions	全局内存原子操作与归约事务	多上下文
atomic_transactions_per_request	每条原子操作和归约指令执行的全局内存原子操作和归约事务的平均数量	多上下文
branch_efficiency	分支指令与分支指令加发散分支指令之和的比率	多上下文
cf_executed	已执行的控制流指令数量	多上下文
cf_fu_utilization	执行控制流指令的多处理器功能单元利用率水平，范围为0到10	多上下文
cf_issued	已发出的控制流指令数量	多上下文
double_precision_fu_utilization	执行双精度浮点指令的多处理器功能单元利用率水平，范围为0到10	Multi-context
dram_read_bytes	从DRAM读取到L2缓存的总字节数	多上下文
dram_read_throughput	设备内存读取吞吐量	多上下文
dram_read_transactions	设备内存读取事务	多上下文
dram_utilization	设备内存利用率水平，相对于峰值利用率，范围为0到10	多上下文
dram_write_bytes	从L2缓存写入DRAM的总字节数	多上下文
dram_write_throughput	设备内存写入吞吐量	多上下文
dram_write_transactions	设备内存写入事务	多上下文
eligible_warps_per_cycle	每个活跃周期内符合条件可发出的平均warp数量	多上下文
flop_count_dp	非预测线程执行的双精度浮点运算次数（加法、乘法和乘加运算）。每个乘加运算会计为2次运算。	多上下文
flop_count_dp_add	非谓词化线程执行的双精度浮点加法运算次数。	多上下文
flop_count_dp_fma	非预测线程执行的双精度浮点乘加运算次数。每个乘加运算计数为1。	多上下文
flop_count_dp_mul	非预测线程执行的双精度浮点乘法运算次数。	多上下文
flop_count_hp	非预测线程执行的半精度浮点运算次数（加法、乘法和乘累加）。每个乘累加根据输入数量对计数贡献2或4。	多上下文
flop_count_hp_add	非预测线程执行的半精度浮点加法操作次数。	多上下文
flop_count_hp_fma	非预测线程执行的半精度浮点乘加运算次数。根据输入数量，每个乘加运算对计数贡献2或4。	多上下文
flop_count_hp_mul	非预测线程执行的半精度浮点乘法运算次数。	多上下文
flop_count_sp	非预测线程执行的单精度浮点运算次数（加法、乘法和乘累加运算）。每个乘累加运算会计为2次运算。该计数不包含特殊运算。	Multi-context
flop_count_sp_add	非预测线程执行的单精度浮点加法操作数量。	多上下文
flop_count_sp_fma	非预测线程执行的单精度浮点乘加运算次数。每个乘加运算计数为1。	多上下文
flop_count_sp_mul	非预测线程执行的单精度浮点乘法运算次数。	多上下文
flop_count_sp_special	非谓词化线程执行的单精度浮点特殊操作数量。	多上下文
flop_dp_efficiency	实际双精度浮点运算与峰值性能的比率	多上下文
flop_hp_efficiency	实际达到的半精度浮点运算与峰值性能的比率	多上下文
flop_sp_efficiency	实际单精度浮点运算与峰值单精度浮点运算的比率	多上下文
gld_efficiency	请求的全局内存加载吞吐量与所需全局内存加载吞吐量的比率，以百分比表示。	多上下文
gld_requested_throughput	请求的全局内存加载吞吐量	多上下文
gld_throughput	全局内存加载吞吐量	多上下文
gld_transactions	全局内存加载事务数	多上下文
gld_transactions_per_request	每次全局内存加载执行的平均全局内存加载事务数。	多上下文
global_atomic_requests	从多处理器发出的全局原子操作(Atom和Atom CAS)请求总数	多上下文
global_hit_rate	统一l1/tex缓存中全局加载和存储的命中率	多上下文
global_load_requests	从多处理器发出的全局加载请求总数	多上下文
global_reduction_requests	从多处理器发出的全局归约请求总数	多上下文
global_store_requests	从多处理器发出的全局存储请求总数。这不包括原子请求。	多上下文
gst_efficiency	请求的全局内存存储吞吐量与所需全局内存存储吞吐量的比率，以百分比表示。	多上下文
gst_requested_throughput	请求的全局内存存储吞吐量	多上下文
gst_throughput	全局内存存储吞吐量	多上下文
gst_transactions	全局内存存储事务数量	多上下文
gst_transactions_per_request	每次全局内存存储操作执行的全局内存存储事务平均数量	多上下文
half_precision_fu_utilization	执行16位浮点指令的多处理器功能单元利用率水平，范围为0到10。请注意，这不包括张量核心单元的利用率水平	Multi-context
inst_bit_convert	非谓词化线程执行的位转换指令数量	多上下文
inst_compute_ld_st	由非预测线程执行的compute加载/存储指令数量	多上下文
inst_control	非谓词化线程执行的控制流指令数量（跳转、分支等）	多上下文
inst_executed	已执行的指令数量	多上下文
inst_executed_global_atomics	用于全局原子操作和原子比较交换的Warp级别指令	多上下文
inst_executed_global_loads	全局加载的Warp级别指令	多上下文
inst_executed_global_reductions	全局归约的Warp级别指令	多上下文
inst_executed_global_stores	全局存储的Warp级别指令	多上下文
inst_executed_local_loads	Warp级别的本地加载指令	多上下文
inst_executed_local_stores	Warp级别的本地存储指令	多上下文
inst_executed_shared_atomics	原子操作和原子CAS的Warp级别共享指令	多上下文
inst_executed_shared_loads	共享加载的Warp级别指令	多上下文
inst_executed_shared_stores	共享存储的Warp级别指令	多上下文
inst_executed_surface_atomics	用于表面原子和原子比较交换的Warp级指令	多上下文
inst_executed_surface_loads	执行表面加载的Warp级别指令	多上下文
inst_executed_surface_reductions	用于表面归约的Warp级别指令	多上下文
inst_executed_surface_stores	用于表面存储的Warp级指令	多上下文
inst_executed_tex_ops	纹理操作的Warp级别指令	多上下文
inst_fp_16	由非预测线程执行的半精度浮点指令数量（算术运算、比较等）	多上下文
inst_fp_32	非预测线程执行的单精度浮点指令数量（算术运算、比较等）	多上下文
inst_fp_64	非预测线程执行的双精度浮点指令数量（算术运算、比较等）	多上下文
inst_integer	非谓词化线程执行的整数指令数量	多上下文
inst_inter_thread_communication	由非谓词化线程执行的线程间通信指令数量	多上下文
inst_issued	已发出的指令数量	多上下文
inst_misc	非谓词化线程执行的杂项指令数量	多上下文
inst_per_warp	每个warp平均执行的指令数	多上下文
inst_replay_overhead	每条执行指令的平均重放次数	多上下文
ipc	每周期执行的指令数	多上下文
issue_slot_utilization	每个周期平均至少发射一条指令的发射槽百分比	多上下文
issue_slots	使用的issue插槽数量	多上下文
issued_ipc	每周期发出的指令数	多上下文
l2_atomic_throughput	L2缓存中原子操作和归约请求的内存读取吞吐量	多上下文
l2_atomic_transactions	在L2缓存中观察到的原子操作和归约请求的内存读取事务	多上下文
l2_global_atomic_store_bytes	从L1写入L2的全局原子操作(ATOM和ATOM CAS)字节数	多上下文
l2_global_load_bytes	从L2读取的字节数（用于L1中全局加载未命中的情况）	多上下文
l2_local_global_store_bytes	从L1写入L2的本地和全局存储字节数。不包括全局原子操作。	多上下文
l2_local_load_bytes	从L2读取的字节数（用于L1本地加载未命中时）	多上下文
l2_read_throughput	所有读取请求在L2缓存中观察到的内存读取吞吐量	多上下文
l2_read_transactions	在L2缓存中观察到的所有读取请求的内存读取事务	多上下文
l2_surface_load_bytes	从L2读取的字节数（用于L1表面加载未命中时）	多上下文
l2_surface_store_bytes	从L2读取的字节数（用于L1表面存储未命中情况）	多上下文
l2_tex_hit_rate	纹理缓存所有请求在L2缓存中的命中率	多上下文
l2_tex_read_hit_rate	来自纹理缓存的所有读取请求在L2缓存中的命中率	多上下文
l2_tex_read_throughput	从纹理缓存读取请求在L2缓存中观察到的内存读取吞吐量	多上下文
l2_tex_read_transactions	L2缓存中观察到的来自纹理缓存的读取请求的内存读取事务	多上下文
l2_tex_write_hit_rate	来自纹理缓存的所有写入请求在L2缓存中的命中率	多上下文
l2_tex_write_throughput	纹理缓存写入请求在L2缓存中观察到的内存写入吞吐量	多上下文
l2_tex_write_transactions	纹理缓存写入请求在L2缓存中观察到的内存写入事务	多上下文
l2_utilization	L2缓存利用率水平，相对于峰值利用率，范围为0到10	多上下文
l2_write_throughput	所有写入请求在L2缓存中观察到的内存写入吞吐量	多上下文
l2_write_transactions	所有写请求在L2缓存中观察到的内存写入事务	多上下文
ldst_executed	已执行的本地、全局、共享和纹理内存加载与存储指令数量	多上下文
ldst_fu_utilization	执行共享加载、共享存储和常量加载指令的多处理器功能单元利用率水平，范围为0到10	多上下文
ldst_issued	已发出的本地、全局、共享和纹理内存加载与存储指令数量	多上下文
local_hit_rate	本地加载和存储的命中率	多上下文
local_load_requests	从多处理器发出的本地加载请求总数	多上下文
local_load_throughput	本地内存加载吞吐量	多上下文
local_load_transactions	本地内存加载事务数	多上下文
local_load_transactions_per_request	每次本地内存加载操作执行的平均本地内存加载事务数	多上下文
local_memory_overhead	L1和L2缓存之间本地内存流量占总内存流量的百分比	多上下文
local_store_requests	从多处理器发出的本地存储请求总数	多上下文
local_store_throughput	本地内存存储吞吐量	多上下文
local_store_transactions	本地内存存储事务数量	多上下文
local_store_transactions_per_request	每次本地内存存储操作执行的平均本地内存存储事务数	多上下文
nvlink_overhead_data_received	通过NVLink接收的开销数据占总数据的比例。	设备
nvlink_overhead_data_transmitted	通过NVLink传输的开销数据占总数据的比例。	设备
nvlink_receive_throughput	通过NVLinks每秒接收的字节数。	设备
nvlink_total_data_received	通过NVLinks接收的数据总字节数（含头部信息）。	设备
nvlink_total_data_transmitted	通过NVLinks传输的总数据字节数（包括头部信息）。	Device
nvlink_total_nratom_data_transmitted	通过NVLinks传输的非归约原子数据总字节数。	设备
nvlink_total_ratom_data_transmitted	通过NVLinks传输的原子操作数据总字节数	设备
nvlink_total_response_data_received	通过NVLink接收的总响应数据字节数，响应数据包括读取请求的数据和非归约原子请求的结果。	设备
nvlink_total_write_data_transmitted	通过NVLinks传输的总写入数据字节数。	设备
nvlink_transmit_throughput	通过NVLinks每秒传输的字节数。	设备
nvlink_user_data_received	通过NVLinks接收的用户数据字节数，不包含头部信息。	设备
nvlink_user_data_transmitted	通过NVLinks传输的用户数据字节数，不包括头部信息。	设备
nvlink_user_nratom_data_transmitted	通过NVLinks传输的非归约原子用户数据字节总数。	设备
nvlink_user_ratom_data_transmitted	通过NVLinks传输的归约原子用户数据总字节数。	设备
nvlink_user_response_data_received	通过NVLink接收的用户响应数据总字节数，响应数据包括读取请求的数据和非归约原子请求的结果。	设备
nvlink_user_write_data_transmitted	通过NVLinks传输的用户写入数据字节数。	Device
pcie_total_data_received	通过PCIe接收的总数据字节数	设备
pcie_total_data_transmitted	通过PCIe传输的总数据字节数	设备
shared_efficiency	请求的共享内存吞吐量与所需共享内存吞吐量的比率，以百分比表示	多上下文
shared_load_throughput	共享内存加载吞吐量	多上下文
shared_load_transactions	共享内存加载事务数	多上下文
shared_load_transactions_per_request	每次共享内存加载操作执行的平均共享内存加载事务数	多上下文
shared_store_throughput	共享内存存储吞吐量	多上下文
shared_store_transactions	共享内存存储事务数量	多上下文
shared_store_transactions_per_request	每个共享内存存储操作执行的平均共享内存存储事务数	多上下文
shared_utilization	共享内存利用率水平，相对于峰值利用率，范围为0到10	多上下文
single_precision_fu_utilization	执行单精度浮点指令的多处理器功能单元利用率水平，范围为0到10	多上下文
sm_efficiency	特定多处理器上至少有一个warp处于活动状态的时间百分比	多上下文
special_fu_utilization	执行sin、cos、ex2、popc、flo等类似指令的多处理器功能单元利用率水平，范围为0到10	Multi-context
stall_constant_memory_dependency	由于立即常量缓存未命中导致的停顿百分比	多上下文
stall_exec_dependency	由于指令所需的输入尚未就绪而发生的停顿百分比	多上下文
stall_inst_fetch	由于下一条汇编指令尚未获取而导致停顿的百分比	多上下文
stall_memory_dependency	由于所需资源不可用或未充分利用，或由于给定类型的未完成请求过多而导致内存操作无法执行时发生的停顿百分比	Multi-context
stall_memory_throttle	因内存节流导致的停顿百分比	多上下文
stall_not_selected	因warp未被选中而发生的停顿百分比	多上下文
stall_other	由于其他原因导致的停顿百分比	多上下文
stall_pipe_busy	由于计算流水线繁忙导致无法执行计算操作而产生的停顿百分比	Multi-context
stall_sleeping	由于warp处于休眠状态而发生的停顿百分比	多上下文
stall_sync	由于warp在__syncthreads()调用处被阻塞而发生的停顿百分比	多上下文
stall_texture	由于纹理子系统已完全占用或存在过多未完成请求而导致的停顿百分比	Multi-context
surface_atomic_requests	从多处理器发出的表面原子(Atom和Atom CAS)请求总数	多上下文
surface_load_requests	从多处理器发出的表面加载请求总数	多上下文
surface_reduction_requests	来自多处理器的表面缩减请求总数	多上下文
surface_store_requests	从多处理器发出的表面存储请求总数	多上下文
sysmem_read_bytes	从系统内存读取的字节数	多上下文
sysmem_read_throughput	系统内存读取吞吐量	多上下文
sysmem_read_transactions	系统内存读取事务数	多上下文
sysmem_read_utilization	系统内存读取利用率水平，相对于峰值利用率，范围为0到10	多上下文
sysmem_utilization	系统内存利用率水平，相对于峰值利用率，范围为0到10	多上下文
sysmem_write_bytes	写入系统内存的字节数	多上下文
sysmem_write_throughput	系统内存写入吞吐量	多上下文
sysmem_write_transactions	系统内存写入事务数	多上下文
sysmem_write_utilization	系统内存写入利用率水平，相对于峰值利用率，范围为0到10	多上下文
tensor_precision_fu_utilization	执行张量核心指令的多处理器功能单元利用率水平，范围为0到10	多上下文
tensor_int_fu_utilization	执行tensor core int8指令的多处理器功能单元利用率水平，范围为0到10。该指标仅适用于计算能力7.2的设备。	Multi-context
tex_cache_hit_rate	统一缓存命中率	多上下文
tex_cache_throughput	统一缓存到多处理器读取吞吐量	多上下文
tex_cache_transactions	统一缓存到多处理器读取事务	多上下文
tex_fu_utilization	执行全局、本地和纹理内存指令的多处理器功能单元利用率水平，范围为0到10	多上下文
tex_utilization	统一缓存相对于峰值利用率的利用率水平，范围为0到10	多上下文
texture_load_requests	从多处理器发出的纹理加载请求总数	多上下文
warp_execution_efficiency	每个warp的平均活跃线程数与多处理器支持的每个warp最大线程数之比	多上下文
warp_nonpred_execution_efficiency	每个warp执行非谓词指令的平均活动线程数与多处理器支持的每个warp最大线程数之比	多上下文

10. Warp 状态

本节包含对每个warp状态的描述。warp可以有以下状态：

• 指令已发出 - 从warp中发出了一条指令或一对独立指令。

• 停滞 - Warp可能因以下任一原因而停滞。停滞原因分布可在PC采样视图中查看源代码级别，或在使用"检查停滞原因"的延迟分析中查看内核级别。

  • 指令获取停滞 - 下一条指令尚未就绪。

    减少指令获取停顿：

    • 如果内核中已展开大型循环，请尝试减少它们。

    • 如果内核包含许多对小函数的调用，可以尝试使用__inline__或__forceinline__限定符内联更多函数。相反，如果内联了许多函数或大型函数，可以尝试使用__noinline__禁用这些函数的内联。

    • 对于非常短的内核，考虑将其融合到单个内核中。

    • 如果使用线程数较少的块，可以考虑减少块数但增加每个块的线程数。这样偶尔调用__syncthreads()可以保持warp同步，可能提高指令缓存命中率。

  • 因执行依赖而停滞 - 下一条指令正在等待其一个或多个输入由先前的指令计算完成。

    为减少执行依赖导致的停滞，可尝试提高指令级并行度(ILP)。例如，可通过增加循环展开或每个线程处理多个元素来实现。这样可以避免线程因每条指令的完整延迟而闲置。

  • 因内存依赖而停滞 - 下一条指令正在等待先前的内存访问操作完成。

    减少内存依赖导致的停顿

    • 尝试提高内存合并和/或获取字节的效率（对齐等）。查看源代码分析中的"全局内存访问模式"和/或指标gld_efficiency与gst_efficiency。

    • 尝试增加内存级并行度(MLP)：即每个线程中同时执行的内存独立操作数量。循环展开、加载向量类型(如float4)以及每个线程处理多个元素，都是提升内存级并行度的有效方法。

    • 考虑将频繁访问的数据移动到更靠近SM的位置，例如通过使用共享内存或只读数据缓存。

    • 在可能的情况下，考虑重新计算数据而非从设备内存加载。

    • 如果本地内存访问频繁，可以考虑增加每个线程的寄存器数量以减少溢出，即使这会降低占用率，因为对于计算能力主版本=5的GPU，本地内存访问仅在L2缓存中缓存。

  • 因内存节流而停滞 - 大量未完成的内存请求阻碍了进程推进。在计算能力主版本为3的GPU上，内存节流表明发生了高频率的内存重放。

    减少内存节流停顿：

    • 尝试找到将多个内存事务合并为一个的方法（例如，使用64位内存请求代替两个32位请求）。

    • 使用源代码分析工具'全局内存访问模式'和/或性能分析指标gld_efficiency与gst_efficiency检查未合并的内存访问；尽可能减少此类情况。

    • 在计算能力主版本 >= 3 的 GPU 上，建议使用 LDG 通过只读数据缓存来处理非合并的全局读取

  • 纹理停滞 - 纹理子系统已完全被占用或存在过多未完成请求。

    减少纹理停顿：

    • 考虑将多个纹理获取操作合并为一个（例如，将数据打包到纹理中并在SM中解包，或使用向量加载）。

    • 考虑通过使用共享内存将频繁访问的数据移动到更靠近SM的位置。

    • 在可能的情况下，考虑重新计算数据而非从内存中获取。

    • 在计算能力主版本小于5的GPU上：考虑将部分纹理访问改为常规全局加载，以减轻纹理单元的压力，特别是当您未使用纹理特有功能（如插值）时。

    • 在计算能力主版本为3的GPU上：如果通过只读数据缓存(LDG)的全局加载是该内核纹理访问的来源，请考虑将其中一些改回常规全局加载。请注意，如果LDG是由于使用__ldg()内置函数而生成的，这仅意味着改回正常的指针解引用；但如果LDG是由于编译器自动使用const和__restrict__限定符而生成的，这可能更加困难。

  • 同步停滞 - 该warp正在等待屏障指令后所有线程完成同步。

    减少同步停顿：

    • 尝试改进负载均衡，即尽量增加同步点之间的工作量；考虑减小线程块大小。

    • 尽量减少使用threadfence_*()。

    • 在计算能力主版本 >= 3 的 GPU 上：如果由于线程块内通过共享内存进行数据交换而使用 __syncthreads()，请考虑是否可以使用 warp shuffle 操作来替代其中一些交换/同步序列。

  • 因常量内存依赖而停滞 - 该线程束因__constant__内存和立即数缓存未命中而停滞。

    每次访问常量时（例如在内核开始时）首次可能会较高。为了减少这些停顿，

    • 考虑减少使用__constant__或通过增加块数量来延长内核运行时间

    • 考虑增加每个线程处理的项数

    • 考虑合并多个使用相同__constant__数据的核函数，以分摊常量缓存未命中的成本。

    • 尝试使用常规全局内存访问而非常量内存访问。

  • 管道繁忙导致停滞 - 由于执行下一条指令所需的功能单元正忙，导致warp停滞。

    减少因管道繁忙导致的停顿：

    • 优先选择高吞吐量操作而非低吞吐量操作。如果精度要求不高，使用单精度浮点运算而非双精度运算。

    • 寻找可能在数学上有效但对编译器自动执行不安全的算术改进（例如运算顺序变更）。这是由于例如浮点数的非结合性等原因。

  • 未被选中而停滞 - 线程束已准备就绪但未能获得执行机会，因为其他线程束被选中执行。此原因通常表明内核可能已优化良好，但在某些情况下，您可以降低占用率而不影响延迟隐藏，这样做可能有助于提高缓存命中率。

  • 其他原因导致的停滞 - Warp因不常见的原因（如编译器或硬件问题）被阻塞。开发者无法控制这些停滞情况。

11. 从Visual Profiler和nvprof迁移至Nsight工具

Visual Profiler和nvprof已被弃用，将在未来的CUDA版本中移除。建议使用新一代工具NVIDIA Nsight Systems进行GPU和CPU采样与追踪，以及NVIDIA Nsight Compute进行GPU内核性能分析。新工具仍提供相同的性能分析/优化/部署工作流程。您需要查看的数据类型保持不变。命令已更改，输出看起来略有不同。新工具功能强大、运行快速且特性丰富，能让您更快地找到解决方案。

NVIDIA Nsight Systems 是一款系统级性能分析工具，旨在可视化应用程序的算法，帮助您识别最大的优化机会，并针对任意数量和规模的CPU与GPU进行高效调优；从大型服务器到我们最小的SoC。请参阅NVIDIA Nsight Systems用户指南中的从NVIDIA nvprof迁移章节。

NVIDIA Nsight Compute 是一款用于CUDA应用程序的交互式内核分析工具。它通过用户界面和命令行工具提供详细的性能指标和API调试功能。此外，其基线特性允许用户在工具内比较结果。Nsight Compute提供可定制且数据驱动的用户界面及指标收集功能，并能通过分析脚本扩展以进行结果后处理。请参阅Nsight Compute CLI文档中的nvprof迁移指南部分，以及Nsight Compute文档中的Visual Profiler迁移指南部分。

另请参阅关于如何将您的开发迁移到下一代工具的博客文章：

1. 从Visual Profiler和nvprof迁移到Nsight工具

2. 从Visual Profiler和nvprof迁移到Nsight Systems

3. 使用Nsight Compute检查您的内核

表7. 不同GPU架构支持的工具对比

GPU架构	Visual Profiler和nvprof	Nsight Systems	Nsight Compute
Maxwell	是	否	否
Pascal	是	是	否
Volta	是	是	是
图灵	支持* (仅限追踪)	支持	支持
安培及更新的GPU架构	否	是	是

* 仅支持追踪功能 - 时间线、活动、API。不支持CUDA内核分析功能，即收集GPU性能指标。

下表将Visual Profiler和nvprof的关键功能对应到NVIDIA Nsight工具

表8. Visual Profiler与nvprof关键功能对照

Visual Profiler/nvprof 功能类别	Nsight Systems	Nsight Compute
时间线/活动/API追踪	是
CPU采样	是
OpenACC	是
OpenMP	是
MPI	是
MPS	是
应用程序依赖关系分析
统一内存传输	是
统一内存页面错误	是
应用统一内存分析
应用NVLink分析		是（每个内核）
事件与指标（每个内核）		是
引导式与非引导式内核分析		是
内核源码-反汇编视图		是
内核PC采样		是
NVTX	是	是
远程性能分析	Yes	Yes

12. 性能分析器已知问题

以下是当前版本中已知的问题。

• Visual Profiler和nvprof不支持计算能力8.0及以上的设备。应改用新一代工具NVIDIA Nsight Compute和NVIDIA Nsight Systems。

• 从CUDA 11.0开始，Visual Profiler和nvprof不再支持macOS作为目标平台。不过在CUDA 12.5版本之前，Visual Profiler仍支持从macOS主机进行远程性能分析。这项支持已在CUDA 12.6版本中被移除。Visual Profiler曾通过单独的安装包提供，以维持macOS上CUDA开发者的远程性能分析工作流程。下载说明请参阅Developer Tools for macOS。

• 从CUDA 10.2开始，Visual Profiler和nvprof使用动态/共享的CUPTI库。因此在Windows上启动Visual Profiler和nvprof之前需要设置CUPTI库的路径。CUPTI库在Windows系统中位于"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\\extras\CUPTI\lib64"。

• 一个安全漏洞问题要求在使用Windows 419.17或Linux 418.43及更高版本驱动程序时，分析工具需要为非root或非管理员用户禁用使用GPU性能计数器的功能。默认情况下，NVIDIA驱动程序需要提升权限才能访问GPU性能计数器。在Tegra平台上，请以root身份或使用sudo进行分析。在其他平台上，您可以选择以root身份或使用sudo启动分析，或者通过启用非管理员分析来解决。有关该问题和解决方案的更多详细信息，请参阅ERR_NVGPUCTRPERM 网页。

  注意

  Visual Profiler和nvprof仅允许在桌面平台上为非root和非管理员用户启用跟踪功能，Tegra平台需要root或sudo访问权限。

• 在Linux平台上，使用环境变量LD_PRELOAD加载某些版本的MPI库可能导致崩溃。解决方法是以root用户身份启动分析会话。对于普通用户，必须为nvprof设置SUID权限。

• 为确保收集所有性能分析数据并刷新到文件，在应用程序退出前应调用cudaDeviceSynchronize()，然后调用cudaProfilerStop()或cuProfilerStop()。请参阅刷新性能分析数据章节。

• 如果在内核执行大量块且执行时间较短的情况下使用并发内核模式，可能会增加显著的开销。

• 如果内核启动频率非常高，用于收集性能分析数据的设备内存可能会耗尽。这种情况下，部分性能分析数据可能会被丢弃。系统将通过警告提示此情况。

• 当分析使用CUDA动态并行(CDP)的应用程序时，分析工具有几个限制。

  • 对于计算能力7.0及以上的设备，CDP内核启动跟踪存在一个限制。分析器会跟踪所有由主机启动的内核，直到遇到一个会启动子内核的主机启动内核。之后的内核将不会被跟踪。

  • 不支持在计算能力7.0及以上的设备上进行源代码级别分析。

  • Visual Profiler时间轴不会显示从设备启动内核内部调用的CUDA API调用。

  • Visual Profiler 不会显示设备启动内核的详细事件、指标和源代码级结果。针对CPU启动内核收集的事件、指标和源代码级结果将包含从该内核内部启动的整个调用树的事件、指标和源代码级结果。

  • nvprof事件/指标输出不包含设备启动内核的结果。为CPU启动内核收集的事件/指标将包含从该内核内部启动的整个调用树的事件/指标。

• 不支持对APK二进制文件进行分析。

• ARM架构(aarch64)不支持统一内存分析。

• 当分析一个因断言而停止设备内核的应用程序时，分析数据将不完整，并显示警告或错误消息。但由于未能检测到故障的确切原因，该消息并不精确。

• 对于依赖分析，如果跟踪中的活动时间戳存在轻微失真，导致违反编程模型约束，则无法分析依赖关系或等待时间。

• 计算能力6.0及以上的设备引入了一项新特性——计算抢占，以便在执行长任务时为所有计算上下文提供公平机会。通过计算抢占特性-

  • 如果多个上下文并行运行，长时间运行的内核可能会被抢占。

  • 由于上下文的时间片到期，某些内核可能会偶尔被抢占。

  如果内核被抢占，内核被抢占的时间仍会计算在内核持续时间中。这会影响Visual Profiler给出的内核优化优先级，因为抢占会引入随机性。

计算抢占可能会影响事件和指标收集。以下是当前版本中已知的问题：

• 对于计算能力7.0及以上的设备，MPS客户端的事件和指标收集可能导致计数高于预期，因为一个MPS客户端可能因另一个MPS客户端终止而被抢占。

• 在计算能力6.0和6.1的设备上，事件warps_launched、sm_cta_launched以及指标inst_per_warp可能会显示比预期更高的计数。可以使用指标unique_warps_launched替代warps_launched来获取实际启动的warp正确计数，因为它不受计算抢占的影响。

• 为避免计算抢占影响性能分析结果，请尝试隔离正在分析的上下文：

  • 在未连接显示器的辅助GPU上运行应用程序。

  • 在Linux系统上，如果应用程序运行在与显示驱动连接的主GPU上，则需要卸载显示驱动。

  • 每次只运行一个使用GPU的进程。

• 计算能力6.0及以上的设备支持按需分页。当内核首次被调度时，所有使用cudaMallocManaged分配且内核执行所需的页面会在GPU产生缺页中断时被提取到全局内存中。分析器需要多次遍历来收集内核分析所需的所有指标。每次内核重放遍历时都需要保存和恢复内核状态。对于计算能力6.0及以上且支持统一内存平台的设备，在第一次内核迭代时会产生GPU缺页中断并将所有页面提取到全局内存。从第二次迭代开始将不再发生GPU缺页中断。这会显著影响内存相关事件和计时。跟踪记录的时间将包含提取页面所需的时间，但大多数在多轮迭代中分析得到的指标不会包含提取页面所需的时间/周期。这会导致分析器结果出现不一致。

• CUDA设备的枚举和顺序，通常通过环境变量CUDA_VISIBLE_DEVICES和CUDA_DEVICE_ORDER控制，对于分析器和应用程序应保持一致。

• 在同时包含支持和不支持CUDA的GPU的系统中，CUDA性能分析可能无法正常工作。对于此类系统，可以在nvprof中设置--devices选项指定支持的设备，或者在启动nvprof或Visual Profiler之前设置环境变量CUDA_VISIBLE_DEVICES。

• 由于Windows上计时器的分辨率较低，对于执行时间较短的活动，开始和结束时间戳可能相同。因此，nvprof和Visual Profiler会报告以下警告："在结果中发现N条无效记录。"

• 分析器无法与其他Nvidia工具（如cuda-gdb、cuda-memcheck、Nsight Systems和Nsight Compute）互操作。

• 当用户应用程序中静态链接OpenACC库时，OpenACC性能分析可能会失败。这是由于编译器可能会忽略应用程序中未使用的函数定义，导致OpenACC性能分析所需的OpenACC API例程定义缺失。这个问题可以通过动态链接OpenACC库来缓解。

• CUDA Toolkit 11.7和CUDA Toolkit 11.8中附带的Visual Profiler和nvprof版本不支持Kepler(sm_35和sm_37)设备。这个问题可以通过升级CUPTI库来解决。请参阅网页CUPTI 11.7和CUPTI 11.8获取支持这些Kepler设备的CUPTI软件包位置。

• 从CUDA 12.4版本开始，使用Optix SDK的应用程序无法再通过Visual Profiler和nvprof进行性能分析。

• 以下系统配置不支持Profiler：

  • 64位ARM服务器CPU架构（arm64 SBSA）。

  • 虚拟GPU (vGPU)。

  • 适用于Linux的Windows子系统(WSL)。

  • NVIDIA加密货币挖矿处理器(CMP)。欲了解更多信息，请访问网页。

可视化性能分析器

以下是已知的与Visual Profiler相关的问题：

• Visual Profiler 需要本地系统安装 Java 运行时环境 (JRE) 1.8。但从 CUDA Toolkit 10.1 更新 2 版本开始，由于 Oracle 升级许可变更，JRE 不再包含在 CUDA Toolkit 中。用户必须安装所需版本的 JRE 1.8 才能使用 Visual Profiler。更多信息请参阅章节设置 Java 运行时环境。

• 某些分析结果需要依赖并非所有设备都支持的指标。当在不支持该指标的设备上尝试这些分析时，分析结果将显示所需数据"不可用"。

• 在Windows系统的Visual Profiler中，使用鼠标滚轮进行滚动操作无效。

• 由于Visual Profiler使用nvprof来收集性能分析数据，nvprof的限制同样适用于Visual Profiler。

• Visual Profiler无法加载超过JVM内存限制或系统可用内存大小的性能分析数据。更多信息请参阅Improve Loading of Large Profiles。

• 在某些版本的Ubuntu上，Visual Profiler的全局菜单显示不正常或为空。一个解决方法是运行Visual Profiler前设置环境变量"UBUNTU_MENUPROXY=0"

• 在Visual Profiler中，滚动图表后NVLink分析图可能显示不正确。可以通过水平调整图表面板大小来修正此问题。

• 当收集大量样本时，Visual Profiler可能无法在时间线上显示NVLink事件。要解决此问题，可通过放大或缩小时间线来刷新视图。替代方案是保存并重新打开会话。

• 在远程设置上进行统一内存分析时，如果GCC版本与主机不同，Visual Profiler可能无法显示CPU页面错误事件的源代码位置。

• 对于远程设置中与主机架构不同（x86与POWER）的统一内存分析，Visual Profiler可能无法显示CPU页面错误和分配跟踪事件的源代码位置。

• Visual Profiler不支持ARM架构(aarch64)。您可以使用远程性能分析。更多信息请参阅远程性能分析部分。

• Visual Profiler 不支持针对 Android 目标的远程性能分析。变通方案是在目标设备上运行 nvprof，然后将 nvprof 输出结果加载到 Visual Profiler 中。

• 对于远程性能分析，主机系统上安装的CUDA Toolkit必须支持远程系统上的目标设备。

• 在未安装所需字体的平台上，Visual Profiler可能会显示奇怪的符号字体。

• 在使用远程性能分析时，如果由于密钥交换失败导致连接故障，您将收到错误信息"无法建立与'用户@xxx'的shell连接"。您可以按照以下步骤来缓解该问题。

  1. 检查目标上的SSH守护进程配置文件（默认路径为/etc/ssh/sshd_config）

  2. 注释掉以下开头的行：

     KexAlgorithms
     

     HostbasedAcceptedKeyTypes
     

     Ciphers
     

     HostKey
     

     AuthorizedKeysFile
     

  3. 重新生成密钥

     sudo ssh-keygen -t rsa -f /etc/ssh/ssh_host_rsa_key
     

  4. 重启sshd服务

     sudo services sshd restart
     

• 从Visual Profiler访问本地帮助文档会导致HTTP 500错误。解决方法是参考本文档（在线文档或PDF版本）。

• Visual Profiler 无法远程连接到运行 Ubuntu 20.04 及更高版本的目标机器。

nvprof

以下是已知与nvprof相关的问题：

• nvprof 无法分析那些调用了 fork() 但未执行 exec() 的进程。

• nvprof 默认会访问系统的临时目录来存储性能分析临时数据。在Linux系统中默认使用/tmp目录，在Windows系统中则由系统环境变量指定。如需自定义存储位置，可修改Linux系统中的$TMPDIR变量或Windows系统中的%TMP%变量。

• 当多个nvprof进程在同一节点上同时运行时，会出现临时目录下文件争用的问题。一个解决方法是给每个进程设置不同的临时目录。

• 多个同时运行的nvprof进程如果使用应用程序回放功能，可能会产生错误结果或完全没有结果。要解决此问题，您需要为每个进程设置唯一的临时目录。在启动nvprof之前设置NVPROF_TMPDIR环境变量。

• 要在Android上分析应用程序，必须定义$TMPDIR环境变量并指向用户可写的文件夹。

• 当自动加速功能启用时，分析结果可能会出现不一致。nvprof默认会尝试禁用自动加速，但在某些情况下可能会失败，不过分析仍会继续。当无法禁用自动加速时，nvprof会报告警告。请注意，自动加速仅支持Kepler及以上架构的特定Tesla设备。

• 分析一个在全局作用域重载new运算符并在重载的new运算符内部使用任何CUDA API（如cudaMalloc()或cudaMallocManaged()）的C++应用程序会导致程序挂起。

• 当使用nvprof工具的--profile-all-processes选项分析所有进程时，NVTX注释将不起作用。建议在启动应用程序前，将环境变量NVTX_INJECTION64_PATH设置为指向分析器注入库（Linux上是libcuinj64.so，Windows上是cuinj64_*.dll）。

事件与指标

以下是已知的与事件和指标分析相关的问题：

• 针对计算能力7.5及更高版本设备的性能分析功能已在NVIDIA Nsight Compute中支持。Visual Profiler不支持计算能力7.5及以上设备的引导分析、非引导分析下的某些阶段以及事件和指标收集。用户可以从Visual Profiler启动NVIDIA Nsight Compute界面来分析计算能力7.5及更高版本的设备。此外，nvprof不支持查询和收集事件与指标、源代码级分析以及其他用于计算能力7.5及以上设备分析的功能选项。这些功能可通过NVIDIA Nsight Compute命令行界面实现。

• 事件或指标收集可能会显著改变应用程序的整体性能特征，因为所有内核执行都在GPU上被序列化。

• 在事件或指标分析过程中，内核启动是阻塞式的。因此等待主机或其他内核更新的内核可能会挂起。这包括主机与设备之间基于值的CUDA流同步API（如cuStreamWaitValue32()和cuStreamWriteValue32()）所建立的同步机制。

• 需要多次传递的事件和指标收集将无法与nvprof内核重放选项一起使用，适用于任何在内核与CPU之间、内核与常规CPU分配内存之间、内核与对等GPU之间，或内核与其他对等设备（例如GPU直接访问）之间执行IPC或数据通信的内核。

• 对于某些指标，所需事件只能针对单个CUDA上下文收集。对于使用多个CUDA上下文的应用程序，这些指标将仅针对其中一个上下文进行收集。只能在单个CUDA上下文中收集的指标在指标参考表中进行了标注。

• 某些指标值的计算基于一个假设：内核足够大，能够以大致相同的工作量占用所有设备多处理器。如果内核启动不具备这一特性，那么这些指标值可能不准确。

• 部分指标并非在所有设备上都可用。要查看特定NVIDIA GPU上所有可用指标的列表，请输入nvprof --query-metrics。您也可以参考指标参考表。

• 当开启"应用重放"模式时，分析器可能无法收集事件或指标。这种情况最可能发生在多线程且非确定性的应用程序中。此时应改用"内核重放"模式。

• 对于分配大量设备内存的应用程序，当使用“内核重放”模式时，分析器可能需要大量时间来收集所有事件或指标。在这种情况下，应改用“应用程序重放”模式。

• 以下是Visual Profiler可能无法收集指标或事件信息的几个原因。

  • 多个工具正在尝试访问GPU。要解决此问题，请确保在任何给定时间点只有一个工具在使用GPU。这些工具包括Nsight Compute、Nsight Systems、Nsight Graphics以及使用CUPTI或PerfKit API（NVPM）读取事件值的应用程序。

  • 多个应用程序同时在使用GPU，而Visual Profiler正在分析一个CUDA应用程序。要解决此问题，请关闭所有应用程序，仅运行Visual Profiler正在分析的应用程序。在应用程序生成事件信息时，应避免与活动桌面交互。请注意，对于某些类型的事件，如果应用程序在同一应用程序中使用多个上下文，Visual Profiler仅收集一个上下文的事件。

• 当使用--events、--metrics或--analysis-metrics选项收集事件或指标时，nvprof将通过内核重放机制多次执行每个内核以收集所有请求的数据。如果请求收集大量事件或指标，则可能需要多次重放，这将显著增加应用程序的执行时间。

• 某些事件并非在所有设备上都可用。要查看特定设备上所有可用事件的列表，请输入 nvprof --query-events。

• 启用某些事件可能导致GPU内核运行时间超过驱动程序的看门狗超时限制。在这种情况下，驱动程序将终止GPU内核，导致应用程序错误且无法获取性能分析数据。在分析此类长时间运行的CUDA内核前，请先禁用驱动程序的看门狗超时功能。

  • 在Linux系统上，建议将X配置选项Interactive设置为false。

  • 对于Windows系统，有关TDR（超时检测与恢复）的详细信息以及如何禁用它，请访问https://docs.microsoft.com/en-us/windows-hardware/drivers/display/timeout-detection-and-recovery

• nvprof can give out of memory error for event and metrics profiling, it could be due to large number of instructions in the kernel.

• 对于使用计算能力6.0和6.1的设备，如果CUDA应用程序是用早于9.0版本的nvcc编译的，分析结果可能不正确。建议使用nvcc 9.0或更高版本重新编译应用程序。如果代码已经使用推荐的nvcc版本编译，可以忽略此警告。

• Tegra平台不支持PC采样功能。

• 不支持对多设备协作内核进行性能分析，即通过API函数cudaLaunchCooperativeKernelMultiDevice或cuLaunchCooperativeKernelMultiDevice启动的内核。

• 不支持对由CUDA Graph启动的CUDA内核节点进行分析。

13. 更新日志

CUDA 12.8中的性能分析器变更

作为CUDA Toolkit 12.8版本发布的一部分所做的更改列表。

• Visual Profiler 和 nvprof 已弃用，将在未来的 CUDA 版本中移除。

• 常规错误修复。此版本未添加新功能。

CUDA 12.6中的性能分析器变更

作为CUDA Toolkit 12.6版本发布的一部分所做的更改列表。

• 移除了macOS主机上Visual Profiler的远程性能分析支持。

• 常规错误修复。此版本未添加新功能。

CUDA 12.5中的性能分析器变更

作为CUDA Toolkit 12.5版本发布的一部分所做的更改列表。

• Visual Profiler从macOS主机进行的远程性能分析支持已弃用，将在未来的版本中移除。

• 不再支持IBM Power架构。

• 常规错误修复。此版本未添加新功能。

CUDA 12.4中的性能分析器变更

作为CUDA Toolkit 12.4版本发布的一部分所做的更改列表。

• 常规错误修复。此版本未添加新功能。

CUDA 12.3中的性能分析器变更

作为CUDA Toolkit 12.3版本发布的一部分所做的更改列表。

• 常规错误修复。此版本未添加新功能。

CUDA 12.2中的性能分析器变更

作为CUDA Toolkit 12.2版本发布的一部分所做的更改列表。

• 常规错误修复。此版本未添加新功能。

CUDA 12.1中的性能分析器变更

作为CUDA Toolkit 12.1版本发布的一部分所做的更改列表。

• 常规错误修复。此版本未添加新功能。

CUDA 12.0中的性能分析器变更

作为CUDA Toolkit 12.0版本发布的一部分所做的更改列表。

• 常规错误修复。此版本未添加新功能。

CUDA 11.8中的性能分析器变更

作为CUDA Toolkit 11.8版本发布的一部分所做的更改列表。

• 常规错误修复。此版本未添加新功能。

CUDA 11.7中的性能分析器变更

作为CUDA Toolkit 11.7版本发布的一部分所做的更改列表。

• 常规错误修复。此版本未添加新功能。

CUDA 11.6中的性能分析工具变更

作为CUDA Toolkit 11.6版本发布的一部分所做的更改列表。

• 常规错误修复。此版本未添加新功能。

CUDA 11.5中的性能分析器变更

作为CUDA Toolkit 11.5版本发布的一部分所做的更改列表。

• 常规错误修复。此版本未添加新功能。

CUDA 11.4中的性能分析器变更

作为CUDA Toolkit 11.4版本发布的一部分所做的更改列表。

• 常规错误修复。此版本未添加新功能。

CUDA 11.3中的性能分析器变更

作为CUDA Toolkit 11.3版本发布的一部分所做的更改列表。

• Visual Profiler 扩展了对在 Intel x86_64 架构上运行版本 11 (Big Sur) 的 macOS 主机的远程性能分析支持。

• 常规错误修复。

CUDA 11.2中的性能分析器变更

作为CUDA Toolkit 11.2版本发布的一部分所做的更改列表。

• 常规错误修复。此版本未添加新功能。

CUDA 11.1中的性能分析工具变更

作为CUDA Toolkit 11.1版本发布的一部分所做的更改列表。

• 常规错误修复。此版本未添加新功能。

CUDA 11.0中的性能分析器变更

作为CUDA Toolkit 11.0版本发布的一部分所做的更改列表。

• Visual Profiler和nvprof不支持计算能力8.0及以上的设备。应改用新一代工具NVIDIA Nsight Compute和NVIDIA Nsight Systems。

• 从CUDA 11.0开始，Visual Profiler和nvprof将不再支持Mac作为目标平台。不过Visual Profiler将继续支持从Mac主机进行远程性能分析。为了维护Mac上CUDA开发者的远程性能分析工作流程，Visual Profiler将通过单独的安装包提供。

• 新增了对追踪Optix应用的支持。

• 修复了自CUDA 10.0以来损坏的nvprof选项--annotate-mpi。

CUDA 10.2中的性能分析器变更

作为CUDA Toolkit 10.2版本发布的一部分所做的更改列表。

• Visual Profiler和nvprof允许桌面平台上的非root和非管理员用户使用追踪功能。请注意，事件和指标分析仍对非root和非管理员用户受限。有关该问题和解决方案的更多详情，请参阅此网页。

• 从CUDA 10.2开始，Visual Profiler和nvprof使用动态/共享的CUPTI库。因此在启动Visual Profiler和nvprof之前需要设置CUPTI库的路径。CUPTI库可以在POSIX平台的/usr/local//extras/CUPTI/lib64或/usr/local//targets//lib找到，Windows平台则在"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\\extras\CUPTI\lib64"。

• 分析器在追踪应用程序时不再关闭CUDA Graph的性能特性。

• 在Visual Profiler中新增了启用/禁用OpenMP性能分析的选项。

• 修复了异步 cuMemset/cudaMemset 活动中的计时错误问题。

CUDA 10.1 更新 2 中的性能分析器变更

作为CUDA Toolkit 10.1 Update 2版本发布的一部分所做的更改列表。

• 此版本主要关注性能分析工具的bug修复和稳定性提升。

• 一个安全漏洞问题要求分析工具为非root或非管理员用户禁用所有功能。因此，当使用Windows 419.17或Linux 418.43及更高版本驱动程序时，Visual Profiler和nvprof无法分析应用程序。有关该问题和解决方案的更多详细信息，请参阅此网页。

• Visual Profiler 需要本地系统安装 Java 运行时环境 (JRE) 1.8。但从 CUDA Toolkit 10.1 更新 2 版本开始，由于 Oracle 升级许可变更，JRE 不再包含在 CUDA Toolkit 中。用户必须安装所需版本的 JRE 1.8 才能使用 Visual Profiler。更多信息请参阅章节设置 Java 运行时环境。

CUDA 10.1中的性能分析器变更

作为CUDA Toolkit 10.1版本发布的一部分所做的更改列表。

• 此版本主要关注性能分析工具的bug修复和稳定性提升。

• 支持NVTX字符串注册API nvtxDomainRegisterStringA()。

CUDA 10.0中的性能分析器变更

作为CUDA Toolkit 10.0版本发布的一部分所做的更改列表。

• 新增对计算能力7.5设备的追踪支持。

• 针对计算能力7.5及更高版本设备的性能分析功能已在NVIDIA Nsight Compute中支持。Visual Profiler不支持计算能力7.5及以上设备的引导分析、非引导分析下的某些阶段以及事件和指标收集。用户可以从Visual Profiler启动NVIDIA Nsight Compute界面来分析计算能力7.5及更高版本的设备。此外，nvprof不支持查询和收集事件与指标、源代码级分析以及其他用于计算能力7.5及以上设备分析的功能选项。这些功能可通过NVIDIA Nsight Compute命令行界面实现。

• Visual Profiler和nvprof现在支持在可用的情况下进行OpenMP性能分析。更多信息请参阅OpenMP。

• 支持追踪由CUDA Graph启动的CUDA内核、内存拷贝(memcpy)和内存设置(memset)节点。

• 分析器支持NVIDIA工具扩展API(NVTX)的第3版本。这是NVTX版本2的仅头文件实现。

CUDA 9.2中的性能分析器变更

作为CUDA Toolkit 9.2版本发布的一部分所做的更改列表。

• Visual Profiler 允许将多个片段切换到非片段模式，以便在时间线上进行统一内存分析。之前仅支持单个片段。

• Visual Profiler 展示了CUDA编程模型内存层次结构的概览视图。该功能适用于计算能力5.0及以上的设备。更多信息请参阅内存统计。

• 当使用选项 --kernels kernel-filter 时，Visual Profiler 可以正确导入由 nvprof 生成的性能分析数据。

• nvprof支持显示基本的PCIe拓扑结构，包括NVIDIA GPU与主机桥之间的PCI桥接器。

• 为了查看和分析PCIe拓扑结构上的内存传输带宽，新增了一组指标用于收集通过PCIe传输和接收的总数据字节数。这些指标提供了系统中所有设备的累计计数。这些指标是在整个应用程序的设备级别收集的。并且这些指标适用于计算能力5.2及以上的设备。

• Visual Profiler 和 nvprof 新增了对以下指标的支持：

  • 针对不同类型加载和存储执行的指令

  • 从SM到纹理缓存的已缓存全局/本地加载请求总数

  • 从纹理缓存写入L2缓存的全局原子/非原子/归约字节数

  • 从纹理缓存写入L2缓存的表面原子/非原子/归约字节数

  • 来自纹理缓存的所有请求在L2缓存中的命中率

  • 设备内存（DRAM）读写字节数

  • 针对计算能力7.0的设备，执行张量核心指令的多处理器功能单元利用率水平

• nvprof允许在同一过程中收集跟踪信息以及性能分析信息。使用新选项--trace 来启用跟踪功能，同时收集事件/指标。

CUDA 9.1中的性能分析工具变更

作为CUDA Toolkit 9.1版本发布的一部分所做的更改列表。

• 可视化性能分析器在CPU详情视图中展示每个线程在CPU上花费的时间细分。

• Visual Profiler 支持一个新选项来选择PC采样频率。

• 可视化性能分析器在NVLink拓扑中显示NVLink版本。

• nvprof 在生成CSV格式的性能分析数据时会提供关联ID。

CUDA 9.0中的性能分析工具变更

作为CUDA Toolkit 9.0版本发布的一部分所做的更改列表。

• Visual Profiler 和 nvprof 现在支持在计算能力7.0的设备上进行性能分析。

• 用于性能分析的工具和扩展托管在Github上，地址为 https://github.com/NVIDIA/cuda-profiler

• 统一内存分析有几项改进：

  • Visual Profiler现在能够将统一内存事件与分配内存的源代码关联起来。

  • Visual Profiler现在可以将CPU页面错误与导致该错误的源代码关联起来。

  • 新增了针对页面抖动、节流和远程映射的统一内存分析事件。

  • Visual Profiler 提供了在时间线上切换分段模式和非分段模式的选项。

  • Visual Profiler支持根据虚拟地址、迁移原因或页面错误访问类型对统一内存分析事件进行过滤。

  • CPU页面错误支持已扩展至Mac平台。

• 支持对协作内核启动进行跟踪和分析。

• Visual Profiler 在时间线上显示 NVLink 事件。

• Visual Profiler 根据吞吐量对 NVLink 拓扑图中的链接进行颜色编码。

• Visual Profiler 支持新的选项，使多跳远程性能分析更加便捷。

• nvprof 支持一个新选项来选择PC采样频率。

• Visual Profiler 支持对采用2048位密钥长度的ssh密钥交换算法的系统进行远程性能分析。

• 现在非NVIDIA系统也支持OpenACC性能分析。

• nvprof 在遇到 SIGINT 或 SIGKILL 信号时会刷新所有性能分析数据。

CUDA 8.0中的性能分析工具变更

作为CUDA Toolkit 8.0版本发布的一部分所做的更改列表。

• Visual Profiler和nvprof现在支持计算能力6.0及以上设备的NVLink分析。更多信息请参阅NVLink视图。

• Visual Profiler和nvprof现在支持依赖分析功能，这有助于优化程序运行时间以及利用多CPU线程和CUDA流的应用程序并发性。它可以计算特定执行的关键路径，检测等待时间，并检查在不同线程或流中执行的函数之间的依赖关系。更多信息请参阅依赖分析。

• Visual Profiler 和 nvprof 现在支持 OpenACC 性能分析。详情请参阅 OpenACC。

• Visual Profiler 现在支持 CPU 性能分析。更多信息请参阅 CPU 详情视图 和 CPU 源代码视图。

• 统一内存分析现在支持在计算能力6.0及以上的设备和64位Linux平台上提供GPU页面错误信息。

• 统一内存分析现在支持在64位Linux平台上提供CPU页面错误信息。

• 统一内存分析支持已扩展至Mac平台。

• Visual Profiler的源代码-反汇编视图有多项改进。现在针对一个内核实例收集的不同源代码级别分析结果有了一个统一的集成视图，可以同时查看不同分析步骤的结果。更多详情请参阅源代码-反汇编视图。

• PC采样功能得到增强，能够准确指出计算能力6.0及以上设备的真实延迟问题。

• 支持16位浮点数(FP16)数据格式的性能分析。

• 如果使用了NVIDIA Tools Extension API(NVTX)的新功能域特性，那么Visual Profiler和nvprof将显示按域分组的NVTX标记和范围。

• 现在，Visual Profiler在保存或打开会话文件时，如果未指定扩展名，会自动添加默认文件扩展名.nvvp。

• Visual Profiler现在支持在创建新会话和导入对话框中的时间线过滤选项。更多详情请参阅创建会话下的"时间线选项"部分。

CUDA 7.5中的性能分析器变更

作为CUDA Toolkit 7.5版本发布的一部分所做的更改列表。

• Visual Profiler 现在支持对计算能力5.2及以上的设备进行PC采样。在源代码级别显示包含停滞原因的Warp状态，用于内核延迟分析。更多信息请参阅PC采样视图。

• Visual Profiler 现在支持分析子进程以及分析同一系统上启动的所有进程。有关新的多进程分析选项的更多信息，请参阅创建会话。要使用多进程服务(MPS)分析CUDA应用程序，请参阅使用Visual Profiler进行MPS分析

• Visual Profiler 导入功能现在支持浏览和选择远程系统上的文件。

• nvprof 现在支持CPU性能分析。详见CPU采样获取更多信息。

• 现在可以准确收集计算能力为5.2的设备的所有事件和指标，即使GPU上存在多个上下文环境。

CUDA 7.0中的性能分析器变更

作为CUDA Toolkit 7.0版本的一部分，性能分析工具包含多项变更和新功能。

• Visual Profiler 已更新，包含多项改进：

  • 加载大型数据文件时性能得到提升，同时内存使用也有所减少。

  • Visual Profiler时间线功能已优化，可查看多GPU MPS性能分析数据。

  • 通过提供与GPU之间的细粒度数据传输，并结合每次传输的更精确时间戳，统一内存分析功能得到了增强。

• nvprof 已更新并包含多项改进：

  • 现在可以准确收集计算能力为3.x和5.0的设备的所有事件和指标，即使GPU上存在多个上下文。

CUDA 6.5中的性能分析器变更

作为CUDA Toolkit 6.5版本发布的一部分所做的更改列表。

• Visual Profiler内核内存分析功能已更新，包含多项改进：

  • 增加了ECC开销，用于统计ECC所需的内存事务数量

  • 在L2缓存下，显示了针对L1读取、L1写入、纹理读取、原子操作和非一致性读取的事务拆分情况

  • 在L1缓存下显示原子事务的计数

• Visual Profiler内核性能分析视图已更新，包含多项增强功能：

  • 初始状态下，执行次数最多的指令会被高亮显示

  • 在“执行计数”列的计数器值背景中显示一个条形图，以便更容易识别执行次数较多的指令

  • 当前汇编指令块周围使用两条水平线高亮显示。此外还添加了“next”和“previous”按钮，用于跳转到下一个或上一个汇编指令块。

  • 为CUDA C源代码添加了语法高亮显示。

  • 新增支持显示或隐藏列的功能。

  • 为每列添加了描述性工具提示。

• nvprof 现在支持一种新的应用程序重放模式，用于收集多个事件和指标。在此模式下，应用程序会运行多次，而不是使用内核重放。当内核使用大量设备内存时，由于每次内核重放运行都需要保存和恢复设备内存而产生高开销，使用内核重放可能会很慢，这时新模式就很有用。更多信息请参阅事件/指标摘要模式。Visual Profiler 也支持这种新的应用程序重放模式，可以在 Visual Profiler 的"新建会话"对话框中启用。

• Visual Profiler现在可以在设备属性下显示GPU的峰值单精度浮点运算和双精度浮点运算性能。

• 针对PGI CUDA Fortran编译器编译的CUDA Fortran应用程序，改进了源代码到汇编代码的关联性。

CUDA 6.0中的性能分析工具变更

作为CUDA Toolkit 6.0版本发布的一部分所做的更改列表。

• Visual Profiler和nvprof都完全支持统一内存。这两个分析器都能让你查看系统中每个GPU与统一内存相关的内存流量。

• 独立版Visual Profiler工具nvvp现在提供多进程时间轴视图功能。您可以将通过nvprof收集的多个时间轴数据集导入nvvp，并在同一时间轴上查看它们如何共享GPU资源。该多进程导入功能还支持使用MPS的CUDA应用程序。详见MPS性能分析获取更多信息。

• Visual Profiler 现在支持远程分析模式，允许您在远程Linux系统上收集分析数据，并在本地Linux、Mac或Windows系统上查看时间线、分析结果和详细结果。更多信息请参阅远程分析。

• Visual Profiler分析系统现在新增了一个并排显示的源代码和反汇编视图，其中标注了指令执行次数、非活跃线程数以及条件指令数。这一新视图可帮助您在内核代码中发现热点区域和低效的代码序列。

• 可视化性能分析系统已更新，新增了多项分析功能：1) 将内核指令按类别划分，便于检查指令组合是否符合预期；2) 检测并报告低效的共享内存访问模式；3) 展示每个SM的活动级别，帮助发现内核各块之间的负载均衡问题。

• Visual Profiler引导式分析系统现在可以生成内核分析报告。该报告是引导式分析系统提供的每个内核信息的PDF版本。

• 现在，nvvp 和 nvprof 都可以在没有 NVIDIA GPU 的系统上运行。您可以从其他系统导入收集的性能分析数据，并在无 GPU 的系统上查看和分析这些数据。

• nvvp 和 nvprof 的性能分析开销已显著降低。

14. 通知

14.1. 注意事项

本文档仅供信息参考之用，不应视为对产品功能、状态或质量的保证。NVIDIA公司（“NVIDIA”）对本文件所含信息的准确性或完整性不作任何明示或暗示的陈述或保证，并对其中可能存在的错误不承担任何责任。NVIDIA对于因使用此类信息而产生的后果、或因使用该信息导致的第三方专利或其他权利侵权概不负责。本文件不构成对开发、发布或交付任何材料（定义见下文）、代码或功能的承诺。

NVIDIA保留随时对本文件进行更正、修改、增强、改进以及任何其他变更的权利，恕不另行通知。

客户在下单前应获取最新的相关信息，并确认这些信息是最新且完整的。

除非NVIDIA与客户授权代表签署的单独销售协议中另有约定，否则NVIDIA产品的销售均以订单确认时提供的NVIDIA标准销售条款和条件为准（以下简称"销售条款"）。NVIDIA特此明确反对将任何客户通用条款适用于本文件所述NVIDIA产品的采购。本文件不直接或间接构成任何合同义务。

NVIDIA产品并非设计、授权或保证适用于医疗、军事、航空、航天或生命支持设备，也不适用于那些可以合理预期NVIDIA产品故障或失灵会导致人身伤害、死亡、财产或环境损害的应用场景。NVIDIA对于在此类设备或应用中使用和/或包含NVIDIA产品不承担任何责任，因此客户需自行承担相关风险。

NVIDIA不声明或保证基于本文档的产品适用于任何特定用途。NVIDIA未必会对每个产品的所有参数进行测试。客户应全权负责评估和确定本文档所含信息的适用性，确保产品适合并满足客户计划的应用需求，并执行必要的应用测试以避免应用或产品出现故障。客户产品设计中的缺陷可能会影响NVIDIA产品的质量和可靠性，并可能导致超出本文档范围的其他或不同的条件和/或要求。对于任何因以下原因导致的故障、损坏、成本或问题，NVIDIA不承担任何责任：(i) 以违反本文档的任何方式使用NVIDIA产品或(ii) 客户产品设计。

本文档不授予任何NVIDIA专利权、版权或其他NVIDIA知识产权的明示或暗示许可。NVIDIA发布的关于第三方产品或服务的信息，不构成NVIDIA对这些产品或服务的使用许可或担保认可。使用此类信息可能需要获得第三方基于其专利或其他知识产权的许可，或需要获得NVIDIA基于其专利或其他知识产权的许可。

本文件中的信息仅可在获得NVIDIA事先书面批准、未经改动完整复制且完全符合所有适用的出口法律法规，并附带所有相关条件、限制和声明的情况下进行复制。

本文件及所有NVIDIA设计规格、参考板、文件、图纸、诊断工具、清单和其他文档（统称及单独称为"材料"）均以"现状"提供。NVIDIA不对材料作出任何明示或默示的保证，包括但不限于对不侵权、适销性和特定用途适用性的默示保证免责。在法律允许的最大范围内，NVIDIA不就因使用本文件导致的任何损害承担责任，包括但不限于任何直接、间接、特殊、附带、惩罚性或后果性损害，无论损害成因如何，也无论责任理论为何，即使NVIDIA已被告知发生此类损害的可能性。不论客户因任何原因可能遭受的任何损害，NVIDIA对客户就本文所述产品的全部及累计责任应受产品销售条款的限制。

14.2. OpenCL

OpenCL是苹果公司的商标，经Khronos Group Inc.授权使用。

14.3. 商标

NVIDIA和NVIDIA标识是美国及其他国家NVIDIA公司的商标或注册商标。其他公司及产品名称可能是其各自关联公司的商标。