CUDA Kernel Profiling Guide LLM翻译
2. 内核分析指南
Nsight Compute 分析指南
本指南涵盖了与 NVIDIA Nsight Compute 及其命令行界面（CLI）相关的各种分析主题。这些内容大多适用于工具的图形用户界面（UI）和命令行版本。
2.1. 引言
为了有效使用这些工具，建议您阅读本指南以及《CUDA 编程指南》中的至少以下章节：
- [相关章节链接]
之后，只需查阅《NVIDIA Nsight Compute》或《NVIDIA Nsight Compute CLI》文档的快速入门章节，即可开始使用这些工具。
2.1.1. 应用程序分析
通常情况下，CUDA 应用程序由用户启动，并直接与 CUDA 用户模式驱动程序及可能的 CUDA 运行时库通信。
然而，当使用 NVIDIA Nsight Compute 分析应用程序时，流程有所不同。用户在主机系统上启动 Nsight Compute 的前端（无论是 UI 还是 CLI），前端随后在目标系统上作为新进程启动实际的应用程序。尽管主机和目标系统常为同一台机器，但目标也可以是具有不同操作系统的远程系统。
该工具将其测量库插入到应用程序进程中，使得分析器能够拦截与 CUDA 用户模式驱动程序的通信。此外，一旦检测到内核启动，这些库还可以从 GPU 收集请求的性能指标。结果随后被传回前端。
2.2. 性能指标收集
性能指标的收集是 NVIDIA Nsight Compute 的核心功能。由于可用指标数量庞大，使用工具预定义的集合或部分来收集常用子集通常更为简便。用户可根据需要自由调整为哪些内核收集哪些指标，但重要的是要考虑到数据收集带来的开销。
2.2.1. 集合与部分
NVIDIA Nsight Compute 使用“部分集”（简称为“集合”）在非常高的层次上决定要收集的指标数量。
每个集合包含一个或多个“部分”，每个部分指定几个逻辑上相关的指标。
例如，一个部分可能只包括高级的 SM 和内存利用率指标，而另一个部分可能包括与内存单元或硬件调度器相关的指标。
部分指定的指标数量和类型对分析期间的开销有重大影响。为了使您能快速选择快速但细节较少的配置文件或较慢但更全面的分析，您可以选择相应的部分集。更多关于分析开销的信息，请参见“开销”。
默认情况下，收集的指标数量相对较少，主要包括高级利用率信息以及静态启动和占用数据。后两者通常无需重放内核启动即可获得。当命令行中未传递 --set、--section 和 --metrics 选项时，将收集基本集。使用 --set full 可收集完整部分集。使用 --list-sets 查看当前可用的集合列表，使用 --list-sections 查看当前可用的部分列表。默认搜索目录和预定义部分文件的位置也称为 sections/。所有相关的命令行选项可在《NVIDIA Nsight Compute CLI 文档》中找到。
在任何目录下放置名为 .ncu-ignore 的文件，可让工具在查找部分（和规则）文件时忽略该目录的内容。即使存在此文件，在递归添加部分目录时，仍会搜索子目录。
2.2.2. 部分与规则
2.2.3. 重放
根据需要收集的指标，内核可能需要重放一次或多次，因为并非所有指标都能在一次遍历中收集。例如，GPU 同时可以收集的来自硬件（HW）性能计数器的指标数量有限。此外，基于补丁的软件（SW）性能计数器可能会严重影响内核运行时间，并导致 HW 计数器的结果偏差。
内核重放
在内核重放中，NVIDIA Nsight Compute 为特定内核实例请求的所有指标被分组为一个或多个遍。对于第一次遍历，内核可以访问的所有 GPU 内存都将被保存。第一次遍历后，确定内核所写入的内存子集。在每次遍历（除第一次外）之前，这个子集会在原始位置恢复，以便内核在每次重放遍历中访问相同的内容。
NVIDIA Nsight Compute 尝试使用最快可用的存储位置进行这种保存和恢复策略。例如，如果数据分配在设备内存中，且仍有足够的设备内存可用，则直接存储在那里。如果设备内存耗尽，数据将被传输到 CPU 主机内存。类似地，如果分配源自 CPU 主机内存，工具首先尝试将其保存在同一内存位置（如果可能的话）。
如“开销”中解释，这所需的时间随着内核访问（尤其是写入）的内存越多，开销增加越大。如果 Nsight Compute 确定只需一次重放遍历即可收集请求的指标，则根本不执行保存和恢复，以减少开销。
应用程序重放
在应用程序重放中，NVIDIA Nsight Compute 为特定内核启动请求的所有指标被分组为一个或多个遍。与内核重放不同，整个应用程序将运行多次，以便在每次运行中为每个内核收集那些遍的指标。
为了正确识别并结合从单个内核启动的多个应用程序重放遍中收集的性能计数器，应用程序需要在内核活动及其分配给 GPU、上下文、流以及潜在的 NVTX 范围方面具有确定性。通常，这也意味着应用程序的整体执行也需要是确定性的。
应用程序重放的好处是，内核访问的内存不需要通过工具保存和恢复，因为每个内核启动在应用程序进程的生命周期中仅执行一次。除了避免内存保存和恢复的开销外，应用程序重放还允许禁用缓存控制。这对于应用程序利用其他 GPU 活动来设置缓存到某种预期状态的情况特别有用。
此外，应用程序重放还可以支持分析在执行期间与主机有相互依赖关系的内核。使用内核重放时，这类内核通常在分析时会挂起，因为在除第一次遍历外的所有遍历中，缺少来自主机的必要响应。相比之下，应用程序重放确保了程序执行在每次遍历中的正确行为。
与内核重放不同，通过应用程序重放收集的多个遍意味着应用程序的所有主机端活动也被重复。如果应用程序需要大量时间进行设置或文件系统访问等操作，开销将相应增加。
跨应用程序重放遍，NVIDIA Nsight Compute 会匹配选定内核启动的指标数据。匹配策略可以通过 --app-replay-match 选项选择。匹配时，仅考虑同一进程内且在相同设备上运行的内核。默认情况下，使用网格策略，根据内核名称和网格大小匹配启动。当多个启动具有相同的属性（例如名称和网格大小）时，它们按执行顺序匹配。#######范围重放（Range Replay）
在NVIDIA Nsight Compute中，范围重放将所有请求的度量指标分组到一个或多个传递中。与内核重放和应用程序重放不同，范围重放捕获并重放被分析应用程序中完整的CUDA API调用和内核启动序列。这样，度量指标就不再关联于单个内核，而是与整个范围相关联。这使得工具可以在不进行序列化的情况下执行内核，从而支持对那些为了正确性或性能原因应并发运行的内核进行剖析。
###定义范围
范围重放需要你在应用程序中指定要剖析的范围。一个范围由开始和结束标记定义，并包括这两个标记之间的所有CPU线程中的CUDA API调用和启动的内核。应用程序负责在这些线程间插入适当的同步，以确保预期的API调用集被捕获。范围标记可以通过以下选项之一设置：
范围必须满足几个要求：
此外，为了保证正确的捕获和重放，范围还应遵循几项建议：
###支持的API
范围重放支持CUDA API的一子集用于捕获和重放。此页面列出了受支持的功能以及可能适用的任何进一步的、特定于API的限制。如果在捕获的范围内检测到不受支持的API调用，将报告错误，且该范围无法被剖析。下面列出的组别与CUDA驱动程序API文档中的组别相匹配。
通常，范围重放仅捕获和重放CUDA驱动程序API调用。当CUDA运行时API调用内部仅生成受支持的CUDA驱动程序API调用时，可以捕获它们。已弃用的API不受支持。
###错误处理、初始化、版本管理、设备管理等API的详细支持情况在此省略，但总体上涵盖了大部分常用功能，同时明确指出某些特定功能如虚拟内存管理、流排序内存分配器、统一寻址、图管理、图形互操作性等是不被支持的。
###应用程序范围重放
在应用程序范围重放中，NVIDIA Nsight Compute中所有请求的度量指标同样被分组到一个或多个传递中。类似于范围重放，度量指标与整个选定的范围相关联，而不是单个内核，这允许工具在无序列化的情况下执行工作负载（内核、CUDA图等），从而支持对那些为了正确性或性能必须并发运行的工作负载进行剖析。
与范围重放不同，范围不是直接为每个传递显式捕获和执行，而是整个应用程序多次重新运行，每次应用程序执行中为每个范围收集一个传递。这样做的好处是不需要为每个范围观察和捕获应用程序状态，且范围内的API调用不需要明确支持，因为范围的正确执行由应用程序自身处理。
定义要剖析的范围与范围重放相同。所要剖析范围的CUDA上下文必须在定义范围起始的线程中当前有效，并且在整个范围内保持活动状态。
###图剖析
在多种重放模式下，NVIDIA Nsight Compute可以将CUDA图作为单一工作负载实体进行剖析，而不是单独剖析各个内核节点。此行为可以在相应的命令行或UI选项中切换。
启用此模式的主要用例包括：
注意，当启用图剖析时，某些度量指标如指令级源代码度量不可用，这同样适用于在图外剖析的内核。
###兼容性、配置文件系列、开销以及度量指标指南的部分在此省略，但主要涉及根据GPU工作负载类型确定可用的重放模式和度量指标、通过自动改变参数配置来快速便捷地分析单个内核性能、剖析过程中的运行时开销因素，以及关于如何理解和使用Nsight Compute提供的各种度量指标的指南。####### 2.3.1 硬件模型
计算模型
所有NVIDIA GPU设计支持一种通用的异构并行编程模型，通常称为计算模型。该模型将GPU从传统的图形管线中分离出来，使其成为一种通用的并行多处理器。异构计算模型意味着存在主机（Host）和设备（Device），在此情况下分别对应CPU和GPU。从高层次来看，主机（CPU）负责管理自身与设备之间的资源，并将工作发送到设备上并行执行。
计算模型的核心是网格(Grid)、块(Block)、线程(Thread)这一层次结构，它定义了GPU上计算任务如何组织。这一层次结构自上而下为：
- 网格：包含多个块。
- 块：包含多个线程，也称为线程块(Thread Block)或协同线程阵列(CTA)。
- 线程：执行的基本单元。
这一层次结构旨在暴露线程间的局部性概念，即协同线程阵列(CTA)。在CUDA中，CTA被称为线程块。架构可以利用这种局部性，为单个CTA内的线程提供快速的共享内存和屏障。当一个网格被启动时，架构保证同一个CTA中的所有线程将在同一流式多处理器(SM)上并发运行。关于网格和块的信息可以在“启动统计”部分找到。
流式多处理器(SM)
流式多处理器是GPU中的核心处理单元，优化以应对广泛的工作负载，包括通用计算、深度学习、光线追踪以及光照和着色。SM设计为同时执行多个CTA，这些CTA可以来自不同的网格启动。
SM实现了单指令多线程(SIMT)执行模型，允许个别线程拥有独特的控制流程，同时作为线程束的一部分执行。图灵架构的SM继承了伏打架构的独立线程调度模型，每个线程维护执行状态，包括程序计数器(PC)和调用栈。独立线程调度允许GPU暂停任何线程的执行，以更高效地使用执行资源或让线程等待其他线程（可能在同一线程束内）产生的数据。通过“源计数器”收集的信息，可以在源页面上检查指令执行和预测细节，以及采样信息。
每个SM被划分为四个处理单元，称为SM子分区。子分区是SM上的主要处理元素，每个子分区包含以下单元：
- 算术逻辑单元(ALU)
- 专用功能单元
- 分支执行单元
- 共享内存和寄存器文件
SM内部共享的资源有：
- L1缓存
- 载入/存储单元(LSU)
一个线程束分配给一个子分区，并从启动到完成都驻留在该子分区上。线程束在映射到子分区时被称为活动或驻留。每个子分区管理一个固定大小的线程束池。在伏打架构中，池的大小为16个线程束；在图灵架构中，池的大小为8个线程束。活动线程束在准备好发出指令时处于可调度状态，这需要线程束有解码指令、所有输入依赖解决，且功能单元可用。“调度器统计”部分可以收集活动、可调度和发出指令的线程束统计信息。
线程束在等待以下情况时会停滞：
- 数据依赖解决
- 指令调度
- 内存访问完成等
“线程束调度状态”列出了可分析的停滞原因，“线程状态统计”部分提供了内核执行中发现的线程状态摘要。
编译器控制的重要资源是内核使用的寄存器数量。每个子分区有一组32位寄存器，由硬件以固定大小的块进行分配。“启动统计”部分显示了内核的寄存器使用情况。
计算抢占
计算抢占避免长时间运行的内核独占GPU，但可能会带来上下文切换开销。执行上下文（寄存器、共享内存等）在抢占时保存并在之后恢复。上下文切换以指令级粒度发生。在支持的系统上，可以使用独占进程计算模式来避免上下文切换。
内存
- 全局内存是一个49位虚拟地址空间，映射到设备上的物理内存、固定的系统内存或对等内存。全局内存对GPU中的所有线程可见，通过SM L1缓存和GPU L2缓存访问。
- 局部内存是执行线程的私有存储，对线程外部不可见，用于存储线程栈和寄存器溢出等线程本地数据。局部内存访问延迟与全局内存相同，但连续的32位字按线程ID顺序排列，只要线程束中所有线程访问相同的相对地址，访问就能完全合并。
- 共享内存位于芯片上，带宽高且延迟低。它可以跨计算CTA共享。使用共享内存共享数据的CTA必须在存储和加载之间使用同步操作（如
__syncthreads()）确保一个线程写入的数据对CTA中的其他线程可见。
缓存
所有GPU单元通过二级缓存（L2）与主内存通信。L2缓存位于片上内存客户端和帧缓冲区之间，工作在物理地址空间，除了提供缓存功能外，还包含硬件以执行压缩和全局原子操作。
- 一级数据缓存（L1） 在处理全局、局部、共享、纹理和表面内存的读写，以及归约和原子操作方面起关键作用。在伏打和图灵架构中，每个TPC有两个L1缓存，每个SM一个。
纹理/表面
TEX单元执行纹理获取和过滤。除了简单的纹理内存访问，TEX还负责执行地址计算、LOD、环绕、滤波和格式转换等操作，将纹理读取请求转换为结果。
纹理和表面内存空间位于设备内存中，并被L1缓存。通过纹理或表面内存读取设备内存具有一些优势，使其成为从全局或常量内存读取内存的有利替代方案。
有关纹理和表面内存的详细信息可以在“内存工作负载分析”部分找到。####### 2.3.2. 度量结构概览
度量概述
NVIDIA Nsight Compute 使用了一套先进的度量计算系统，旨在帮助您了解程序运行的情况（计数器和度量值），以及程序达到GPU峰值性能的程度（通过吞吐量百分比表示）。每个计数器在数据库中都有相应的峰值速率，以便计算其作为百分比的吞吐量。
吞吐量度量返回其构成计数器的最大百分比值。这些构成部分经过精心挑选，代表了决定峰值性能的GPU管道部分。虽然所有计数器都可以转换为峰值的百分比，但并非所有计数器都适合于峰值性能分析；不适合的例子包括活动的限定子集和工作负载驻留时间计数器。使用吞吐量度量确保了分析的有意义和可操作性。
每个计数器有两种类型的峰值速率：突发和持续。突发速率是在单个时钟周期内可报告的最大速率。持续速率是对于“典型”操作，在无限长的测量周期内可达到的最大速率。对于许多计数器，突发速率等于持续速率。由于不能超过突发速率，因此其百分比总是小于100%。持续速率的百分比在边缘情况下偶尔可以超过100%。
度量实体
在NVIDIA Nsight Compute中，所有性能计数器都被命名为度量值，它们还可以进一步细分为具有特定属性的组。通过PerfWorks测量库收集的度量值，存在以下实体：
- 计数器可以是来自GPU的原始计数器，也可以是计算出的计数器值。每个计数器下有四个子度量，也称为汇总：
- 平均：计数器的平均值。
- 最大：计数器的最大值。
- 最小：计数器的最小值。
- 总和：计数器的总和。
- 汇总还包括以下内置的子度量：
- 峰值突发吞吐量（新加入）。
- 峰值持续吞吐量（新加入）。
比例有三个子度量：
吞吐量表明GPU的某部分达到了接近峰值速率的程度。每个吞吐量具有以下子度量：
此外，还说明了已废弃的计数器和吞吐量子度量，因为它们对性能优化没有帮助。
除了PerfWorks度量之外，NVIDIA Nsight Compute还使用其他几个生成自己度量值的测量提供者，这些在度量参考中解释。
度量示例与命名约定
周期度量
实例化度量
Nsight Compute收集的度量值可以有一个聚合值、多个实例值或两者都有。实例允许度量值有多个子值，例如代表作为工作负载一部分执行的函数地址范围内的源度量值的每个指令偏移处的值。如果一个度量值有实例值，它通常也有每个实例的关联ID。关联ID和值形成映射，使工具能够在上下文中关联这些值。对于源度量值，该上下文通常是作为工作负载一部分执行的函数的地址范围。
2.3.3. 度量解码器
接下来解释在NVIDIA Nsight Compute度量名称中发现的术语，这些术语在“度量结构”中引入。
简化的Volta及更新架构中L1TEX处理模型如下：当SM执行全局或局部内存指令为一个warp时，会向L1TEX发送一个请求。这个请求传达了这个warp（最多32个）所有参与线程的信息。对于局部和全局内存，根据访问模式和参与线程，请求需要访问一定数量的缓存行和这些缓存行内的扇区。L1TEX单元内部有多个在管道中操作的处理阶段。
波前（wavefront）是每周期能通过那个管道阶段的最大单位。如果所有的缓存行或扇区不能在一个波前中被访问，那么就会创建多个波前并逐一发送进行处理，即以序列化的方式。波前内部工作的限制可能包括需要一致的内存空间、可访问的最大缓存行数，以及其他各种原因。然后每个波前通过L1TEX管道，并获取在该波前中处理的扇区。在这个模型中，三个关键值之间的关系是：请求：扇区为1:N，波前：扇区为1:N，请求：波前为1:N。
波前被描述为一次可以处理的工作包，即在L1TEX中有一种每周期处理一个波前的概念。波前因此代表了处理请求所需的周期数，而每个请求的扇区数是内存指令的所有参与线程的访问模式的属性。例如，有可能一个内存指令每请求需要4个扇区且在1个波前中。但同样也可能有内存指令每请求需要4个扇区，却需要2个或更多波前来处理。####### 2.3.4. 测量范围与精度概览
一般来说，如果度量值超出了指标的预期逻辑范围，这可能归因于以下一个或多个根本原因。如果值超过了这样的范围，工具故意不将其限制在预期值，以确保剖析报告的其余部分保持自洽。
异步GPU活动
除了被指标测量的GPU引擎（如显示、复制引擎、视频编码器、视频解码器等）之外，在剖析过程中可能会访问共享资源。这些芯片全局共享资源包括L2缓存、DRAM、PCIe和NVLINK。如果内核启动较小，其他引擎可能会对DRAM结果造成显著混淆，因为无法隔离SM的DRAM流量。为了减少这类异步单元的影响，建议在没有活动显示且没有其他进程同时访问GPU的情况下进行剖析。
多遍数据收集
当剖析器重放内核启动以收集指标时，如果工作分配在多次重放中显著不同，就常会出现超出范围的指标。像命中率（命中数/查询数）这样的指标，如果在不同的重放遍次中收集了命中和查询，并且内核未能使GPU达到稳定状态（通常大于20微秒），就会出现较大误差。同样，当工作负载本质上是可变的，比如在自旋循环的情况下，也可能显示意外的值。
为缓解此问题，若适用，请尝试增加测量的工作负载，让GPU在每次启动时都能达到稳定状态。同时减少同时收集的指标数量也能通过提高同一指标的计数器在同一遍中被收集的可能性来提升精度。
工具问题
如果遵循上述指南后仍然观察到指标问题，请联系我们并描述您的问题。
2.4. 指标参考概览
大多数NVIDIA Nsight Compute中的指标可以通过ncu命令行界面的–query-metrics选项查询。
以下指标可以明确收集，但不被–query-metrics列出，也不遵循在指标结构中解释的命名方案。它们应按原样使用。
launch__*指标是按内核启动收集的，不需要额外的重放遍次。它们作为内核启动参数的一部分可用（如网格大小、块大小等），或者使用CUDA占用率计算器计算得出。
设备属性
device__attribute_*指标代表CUDA设备属性。收集它们不需要额外的内核重放遍次，因为其值可以从CUDA驱动程序为每个CUDA设备获得。
请参见下文自定义device__attribute_*指标。
通过线程调度状态采样收集。无论调度器是否在同周期发出指令，它们都会递增。这些指标的实例值从函数地址（uint64）映射到样本数（uint64）。
通过线程调度状态采样收集。它们仅在调度器未发出指令的周期递增。这些指标的实例值从函数地址（uint64）映射到样本数（uint64）。
大多数是通过SASS修补收集的。这些指标的实例值从函数地址（uint64）映射到关联值（uint64）。memory_[access_]type指标映射到字符串值。
通过SASS修补收集。这些指标的实例值从SASS操作码（字符串）映射到执行次数（uint64）。
由工具本身生成的指标，用于提供剖析期间的统计信息或问题。
2.5. 采样
NVIDIA Nsight Compute可以通过固定间隔的采样来收集某些性能数据。
2.5.1. 性能监视器（PM）采样
NVIDIA Nsight Compute支持定期以固定间隔采样GPU的性能监视器（PM），以收集许多指标。得到的指标是实例化的，每个样本由其值和采集时的（GPU）时间戳组成。这使得工具可以在时间线上可视化数据，帮助您了解剖析工作负载在运行时的行为如何变化。
2.5.2. 线程采样
NVIDIA Nsight Compute支持周期性地采样线程程序计数器和线程调度状态。在固定的周期数上，每个流式多处理器中的采样器选择一个活跃的线程并输出程序计数器和线程调度状态。工具为设备选择最小的间隔。在小型设备上，这可能是每32个周期一次。在具有更多多处理器的更大芯片上，这可能是2048个周期。采样器随机选择一个活跃的线程。在同一流水线周期，调度器可能会选择另一个线程来发射指令。
2.6. 可重复性
为了在应用程序运行之间提供可操作且确定的结果，NVIDIA Nsight Compute采用多种方法调整指标的收集方式。这包括序列化内核启动、在每次内核重放前清空GPU缓存或调整GPU时钟。#######2.6.1. 序列化
NVIDIA Nsight Compute 在被分析的应用程序中对内核启动进行序列化，这可能涉及到同时由一个或多个工具实例分析的多个进程。
跨进程的序列化是必要的，因为为了收集硬件性能指标，某些GPU和驱动程序对象一次只能被一个进程获取。为此，使用了锁定文件 TMPDIR/nsight-compute-lock。在Windows上，TMPDIR是Windows GetTempPath API函数返回的路径。在其他平台上，它是从列表TMPDIR、TMP、TEMP、TEMPDIR中的第一个环境变量提供的路径。如果这些都找不到，在QNX上是/var/nvidia，在其他情况下是/tmp。
进程内的序列化对于大多数指标正确映射到相应的内核是必需的。此外，如果不进行序列化，当内核在同一设备上并发执行时，性能指标值可能会有较大波动。
目前还无法禁用此工具行为。请参考常见问题解答中的解决方法。
###2.6.2. 时钟控制
对于许多指标，它们的值直接受到当前GPU SM（流式多处理器）和内存时钟频率的影响。例如，如果分析的是应用程序中之前已经执行过其他内核的内核实例，GPU可能已经处于更高的时钟状态，测量到的内核持续时间以及其他指标都会受到影响。同样，如果一个内核实例是应用程序中第一个启动的内核，GPU时钟通常会较低。另外，由于内核重放的存在，指标值可能依赖于在哪个重放阶段收集，因为后续的重放阶段会导致更高的时钟状态。
为了减轻这种非确定性，NVIDIA Nsight Compute尝试将GPU时钟频率限制在其基本值。因此，指标值较少受到内核在应用程序中的位置或特定重放次数的影响。
但是，对于内核分析，这种行为可能不理想，例如当使用外部工具固定时钟频率，或者分析内核在应用程序中的行为时。为了解决这个问题，用户可以通过调整--clock-control选项来指定是否应由工具固定任何时钟频率。
###2.6.3. 缓存控制
如内核重放部分所述，为了收集所有请求的指标，可能需要多次重放内核。虽然NVIDIA Nsight Compute可以为每次重放保存和恢复内核访问的GPU设备内存内容，但它无法对硬件缓存（如L1和L2缓存）的内容做同样的操作。
这可能会导致后续重放通行证比首次通行证有更好的或更差的性能，因为缓存可能已经被内核最后访问的数据预加载。类似地，首次通行证收集的硬件性能计数器的值可能取决于在测量的内核启动之前执行了哪些内核（如果有）。
为了让硬件性能计数器的值更加确定，NVIDIA Nsight Compute默认在每次重放前刷新所有GPU缓存。因此，在每次重放中，内核都将访问干净的缓存，行为就如同内核完全独立执行一样。
对于性能分析，这种行为可能不理想，特别是如果测量集中在大型应用程序执行中的某个内核上，并且收集的数据针对的是以缓存为中心的指标。在这种情况下，您可以使用--cache-control none来禁用工具对任何硬件缓存的刷新。
###2.6.4. 持续模式
NVIDIA内核模式驱动程序必须在任何用户与目标GPU设备交互之前运行并连接到该设备。驱动程序的行为根据操作系统而异。通常，在Linux上，如果内核模式驱动程序尚未运行或连接到目标GPU，任何尝试与该GPU交互的程序的调用都会透明地导致驱动程序加载和/或初始化GPU。当所有GPU客户端终止时，驱动程序随后会解除初始化GPU。
如果未启用持久模式（作为操作系统的一部分或由用户），触发GPU初始化的应用程序可能会产生短暂的启动成本。此外，在某些配置中，当应用程序结束时GPU被解除初始化时，也可能存在关闭成本。
在使用NVIDIA Nsight Compute进行分析之前，建议在适用的操作系统上启用持久模式，以获得更一致的应用程序行为。
###2.7. 特殊配置
###2.7.1. 多实例GPU
多实例GPU（MIG）是一个功能，允许GPU被划分为多个CUDA设备。划分在两个层面上进行：首先，GPU可以被分割成一个或多个GPU实例。每个GPU实例声明拥有一个或多个流式多处理器（SM）、GPU总内存的一部分，以及可能的其他GPU资源，如视频编码器/解码器。其次，每个GPU实例可以进一步划分为一个或多个计算实例。每个计算实例对其分配的GPU实例的SM具有独占所有权。然而，同一GPU实例中的所有计算实例共享该GPU实例的内存和内存带宽。每个计算实例都像一个具有唯一设备ID的CUDA设备那样行动和操作。更多关于如何配置MIG实例的信息，请参见驱动程序发行说明以及nvidia-smi命令行工具的文档。
在分析方面，计算实例可以是两种类型之一：隔离的或共享的。
一个隔离的计算实例拥有其分配的所有资源，并不与其他计算实例共享任何GPU单元。换句话说，计算实例与其父GPU实例大小相同，因此没有其他兄弟计算实例。对隔离的计算实例的分析工作照常进行。
一个共享的计算实例使用的GPU资源可能也会被同一GPU实例中的其他计算实例访问。由于这种资源共享，从这些共享单元收集分析数据是不允许的。尝试从共享单元收集指标会失败，并显示错误信息“ERROR 无法访问以下指标”。在MIG实例上进行分析时，无法从与其他MIG实例共享的GPU单元收集指标，之后会列出失败的指标。仅从共享计算实例独占拥有的GPU单元收集指标仍然是可行的。
###锁定时钟
NVIDIA Nsight Compute无法为任何计算实例设置分析时的时钟频率。您可以继续在没有固定时钟频率的情况下分析内核（使用--clock-control none; 更多细节请见此处）。如果您有足够的权限，可以使用nvidia-smi通过调用nvidia-smi --lock-gpu-clocks=tdp,tdp为整个GPU配置固定的频率。这会将GPU时钟设置为基础TDP频率，直到您通过调用nvidia-smi --reset-gpu-clocks重置时钟为止。
###MIG在裸金属（非vGPU）上
GPU上的所有计算实例共享相同的时钟频率。
###MIG在NVIDIA vGPU上
为VM启用分析将使VM能够访问GPU的全局性能计数器，这可能包括来自在同一天GPU上执行的其他VM的活动。为VM启用分析也允许VM锁定GPU上的时钟，这会影响在同一天GPU上执行的所有其他VM，包括MIG计算实例。
###2.8. 屋顶线图
屋顶线图提供了一种非常有用的方式来可视化复杂处理单元（如GPU）上的实际性能。本节介绍了在分析报告中呈现的屋顶线图。####### 2.8.1. 概览
内核性能不仅仅依赖于GPU的运算速度。因为内核需要数据来处理，所以性能也取决于GPU向内核提供数据的速度。典型的屋顶线图结合了GPU的峰值性能和内存带宽，以及一个称为算术强度（工作与内存流量之间的比率）的指标，形成一个单一图表，以更现实地表示所分析内核的实现性能。一个简单的屋顶线图可能如下所示：
此图实际上展示了两条不同的屋顶线。然而，每个屋顶线都可识别出以下组件：
2.8.2. 分析
屋顶线图对于指导特定内核的性能优化工作非常有帮助。
如图所示，脊点将屋顶线图分为两个区域。斜率下的蓝色阴影部分为“内存受限”区域，而峰值性能边界下的绿色阴影部分为“计算受限”区域。所达成值所在的区域决定了内核性能当前的限制因素。
所达成值到各自屋顶线边界的距离（本图中以白色虚线表示），代表了性能改进的机会。所达成值越接近屋顶线边界，其性能就越优化。如果所达成值位于内存带宽边界上但尚未达到脊点的高度，则表明进一步提高总体FLOP/s仅在同时增加算术强度的情况下才有可能。
结合基线功能使用屋顶线图是跟踪多个内核执行过程中优化进度的好方法。
2.9. 内存图表
内存图表以图形逻辑方式显示GPU内外存子单元的性能数据。性能数据包括传输大小、命中率、指令或请求的数量等。
2.9.1. 概览
逻辑单元（绿色）
逻辑单元以绿色显示。
物理单元（蓝色）
物理单元以蓝色显示。
根据具体的GPU架构，显示的单位集可能会有所不同，因为并非所有GPU都具备所有单位。
链接
内核与其他逻辑单元之间的链接代表针对相应单元执行的指令数（Inst）。例如，内核与全局之间的链接代表加载自或存储至全局内存空间的指令。使用NVIDIA A100的Load Global Store Shared范式的指令会单独显示，因为它们的寄存器或缓存访问行为可能与常规全局加载或共享存储不同。
逻辑单元与蓝色物理单元之间的链接代表因各自指令产生的请求（Req）数量。例如，从L1/TEX缓存到全局的链接显示由于全局加载指令生成的请求数量。
每条链接的颜色代表相应通信路径的峰值利用率百分比。图表右侧的颜色图例显示了从未使用（0%）到峰值性能运行（100%）应用的颜色渐变。图例左侧的三角标记对应图表中的链接。这些标记相对于颜色渐变本身提供了更准确的峰值性能实现值估计。
一个单元通常共享一个公共数据端口用于输入和输出流量。虽然共享端口的链接可能远低于各自的峰值性能，但该单元的数据端口可能已经达到了其峰值。端口利用率通过图表中位于输入和输出链接处的单元内部的彩色矩形显示。端口使用与数据链接相同颜色渐变，并且在图例左侧也有相应的标记。
内存表中报告的峰值值与内存图表中的端口之间的关联示例如下所示。
度量
此图表的度量可通过命令行使用--set full, --section MemoryWorkloadAnalysis_Chart或--metrics group:memory__chart收集。
（后续章节省略以保持回答简洁性）
以上内容概述了GPU性能优化中的关键概念，包括屋顶线图分析、内存图表、内存表格等，这些工具和技术有助于开发者理解和优化其应用程序在GPU上的执行效率。