深入GPU性能分析工具：现代加速器追踪工具的实现调研

对异构加速器（GPU、DPU和APU）进行性能分析和追踪对于优化现代系统的性能至关重要。本调研提供了一个深入的实现导向的综述，介绍了当前最先进的工具如何捕获低级执行细节。我们面向两类受众：(1) 工具开发者，他们希望了解现有性能分析工具的构建方式——包括它们钩住的接口以及如何追踪CPU和加速器；(2) 系统工程师，他们需要根据功能和开销决定集成哪些工具。我们探讨了关键性能分析工具的内部架构、它们拦截的运行时库和驱动API、它们是否追踪CPU、GPU或两者上的事件，以及它们如何处理跨设备关联。然后我们检查实现策略（插桩vs采样、性能计数器的使用、动态二进制插桩、内核钩子如eBPF）以及GPU特定技术（如warp级指令采样）。我们比较了数据处理和可视化方法，包括追踪数据格式、存储和导出方法，以及GUI/CLI/Web界面（与TensorBoard或Grafana等仪表板集成）。我们列出了这些工具收集的性能指标——从内核延迟和吞吐量到SM占用率、缓存未命中和互连利用率——解释了原始硬件计数器如何转换为派生指标以及涉及的粒度/精度权衡。接下来，我们讨论用户可以从可视化（时间线、火焰图、依赖图等）中获得的见解，例如检测流水线停顿、内存瓶颈或CPU-GPU不同步。我们回顾了工具的可扩展性（对插件或用户定义插桩的支持），以及不同测量技术的相对开销和侵入性。我们还讨论了安全和沙盒问题——例如，在容器化或多租户环境中的性能分析——并概述了已知限制和未来方向（包括对新兴DPU和集成CPU-GPU APU的支持）。本调研的目标是提供一个全面的系统级概述，超越功能列表，比较这些性能分析工具在底层的工作方式，包括开源工具使用的任何逆向工程或未记录的方法。

关键性能分析工具的内部架构

了解每个工具的内部设计揭示了它能够观察软件/硬件栈的哪些层次。下面我们回顾具有代表性的GPU性能分析和追踪工具——包括厂商提供的和开源的——重点关注它们的架构、数据收集机制和目标平台。

NVIDIA Nsight Systems

NVIDIA Nsight Systems是一个系统级性能分析工具，为CUDA应用程序提供CPU和GPU活动的时间线视图。在内部，Nsight Systems结合了轻量级的CPU采样和GPU活动追踪。它利用Linux性能事件接口来捕获CPU线程调度并对CPU指令指针进行采样。实际上，这意味着它依赖于perf_event_open进行最小开销的采样，在Linux上需要适当的perf_event_paranoid设置。对于GPU追踪，Nsight Systems通过NVIDIA的CUDA分析工具接口(CUPTI)与之交互，以在GPU操作发生时（内核启动、内存复制等）接收回调"活动"记录。通过钩入CUDA驱动程序/运行时（通过CUPTI），Nsight Systems记录每个GPU内核执行和内存传输的时间戳。它通过一个公共时间戳基准同步CPU和GPU时间线（CUPTI事件被时间戳标记并与CPU事件合并），使两个处理器上的事件可以在统一的时间线中关联起来。该工具还追踪CUDA之外的各种运行时：它可以捕获操作系统事件、存储I/O和GPU加速库（如cuDNN、cuBLAS）（当通过NVTX注释进行插桩时）。Nsight Systems主要对GPU使用基于事件的插桩（通过CUPTI回调）和对CPU使用统计采样，使其成为一种混合方法。这种设计实现了低开销——通常只有几个百分点——足以对完整应用程序进行分析。数据被记录到专有的跟踪文件（如.nsys-rep）中，可以在Nsight Systems GUI中查看或导出（CLI可以生成JSON或CSV摘要统计信息）。它支持多线程和多GPU追踪，在时间线上显示上下文切换和并发性。然而，作为厂商特定的工具，Nsight Systems专注于NVIDIA GPU以及CUDA或相关NVIDIA API；它不直接支持AMD或Intel GPU运行时。

NVIDIA Nsight Compute

Nsight Compute是NVIDIA的内核级性能分析器，用于对单个GPU内核进行深入分析。与提供时间线不同，它为每个内核启动提供详细的指标报告。在内部，Nsight Compute广泛地对GPU硬件性能计数器进行插桩和测量。它使用NVIDIA的分析接口（如CUPTI的Event和Metric API或更新的CUDA Profiling API）来编程GPU性能监视器，以计算执行的指令、内存事务、warp占用率等事件。由于GPU可用计数器有限，Nsight Compute通常会多次重放每个内核以收集所有请求的指标。例如，一个内核可能会运行数十次（例如46次）以收集不同的计数器集，这些计数器稍后会被聚合。这种事后插桩方法产生了全面的指标（包括派生指标，如已实现的占用率或内存吞吐量），但会产生显著的运行时开销——因此Nsight Compute通常用于小型内核或选择性分析会话。该工具的架构包括一个内核注入机制：它在每个内核启动时暂停目标程序，设置计数器（甚至在内核周围插入计时工具），然后恢复执行并收集结果。它仅支持NVIDIA GPU（从Volta架构开始）和主要是CUDA内核（尽管它也可以通过钩住这些调用来分析来自GPU加速库或CUDA图的内核）。结果存储在报告文件中或显示在交互式GUI/CLI中，指标按"节"分组。Nsight Compute的指标计算系统非常先进——它使用工具中嵌入的公式从低级计数器计算高级统计数据（例如，已实现的占用率、停滞原因细分）。输出可以导出为文本或CSV以供进一步分析。虽然非常详细，但Nsight Compute与NVIDIA驱动程序的紧密集成意味着它不能扩展到其他供应商，而且必须离线使用（对于连续监控来说太具侵入性）。

AMD Omnitrace

Omnitrace是由AMD研究院开发的开源性能分析和追踪工具。其设计强调通过动态二进制插桩和采样对CPU和GPU执行进行全面追踪。Omnitrace可以附加到未修改的应用程序上，在运行时对函数调用（包括HIP和OpenCL API调用）进行插桩。在底层，它使用动态二进制插桩框架在目标程序的函数中插入探针（文档暗示它可能使用Dyninst或类似工具）。用户可以在执行前配置要插桩的函数或模块，使他们能够通过只选择特定区域来限制开销。在运行时，插桩会记录进入/退出事件和时间戳，而并行的调用栈采样组件定期对所有线程进行采样，以捕获CPU执行热点。对于GPU工作，Omnitrace钩入AMD的ROCm栈：它使用ROCm的roctracer回调来追踪AMD GPU上的HIP运行时事件和GPU内核执行。（目前，Omnitrace仅通过HIP和HSA API支持AMD GPU——尚不支持NVIDIA GPU，反映了其AMD专注点）。收集的追踪数据以开放的Perfetto格式（基于protobuf）写入，可以在Chrome/Perfetto网页查看器中可视化。这种格式选择使得可以在任何浏览器中查看丰富的时间线视图（Perfetto UI）。此外，Omnitrace生成聚合的性能分析（例如，每个函数的总时间、调用图）以JSON和文本文件形式。JSON输出与Hatchet等分析工具集成（它将数据加载到pandas中进行自定义分析）。由于其开源性质，Omnitrace的架构具有高度可扩展性：用户可以修改它以支持新事件或使用其用户API标记自定义区间。它有效地合并了统计性能分析器和追踪器的功能，提供了高级摘要和完整执行追踪。由于它在二进制级别进行插桩，存在一些开销，但这通过让用户选择采样频率和要插桩的函数来缓解。总体而言，Omnitrace提供了跨平台性能分析器的蓝图（它设计为也可在仅CPU的应用程序上工作，支持任何CPU供应商，而GPU支持目前是AMD特定的）。

Rice HPCToolkit

HPCToolkit是一个长期存在的开源性能分析工具包，以其能够以低开销分析大型HPC代码而著称。它最近的扩展通过结合CPU采样和GPU活动追踪来支持GPU加速程序。HPCToolkit的内部架构使用供应商中立的GPU监控基础。这个基础连接到特定供应商的驱动程序：在NVIDIA上，它使用CUPTI；在AMD上，它使用ROC-tracer；对于Intel GPU，它包装了Level Zero API。本质上，HPCToolkit注册回调函数，以便每当GPU运行时（CUDA、HIP、Level0、OpenCL）启动内核或执行数据传输时，HPCToolkit会收到通知。此外，HPCToolkit为OpenCL和Level Zero安装自己的拦截包装器，以捕获在没有内置跟踪器的平台上的事件。所有GPU事件（内核启动开始/结束等）都被时间戳标记并记录在GPU操作跟踪中。同时，应用程序的CPU端通常通过采样进行分析——HPCToolkit的hpcrun使用定时器或硬件计数器中断对CPU线程的调用栈进行采样，将成本归因到代码区域。这两个数据源在分析过程中合并，以关联CPU调用路径与GPU活动。HPCToolkit更进一步，提供了细粒度GPU分析：在NVIDIA GPU上，它可以对GPU本身执行PC采样。自Maxwell以来的NVIDIA硬件可以在执行期间周期性地对warp的程序计数器进行采样。HPCToolkit通过CUPTI利用这一点来收集GPU指令地址的统计分析，这些地址稍后映射回源代码行和函数（HPCToolkit对CUDA CUBIN文件进行二进制分析以恢复函数名和行映射）。这产生了GPU代码的层次分析，类似于CPU采样如何识别源级别的热点。（在AMD GPU上，由于缺乏硬件支持，目前无法进行这种细粒度采样，因此HPCToolkit在那里只追踪内核级计时。）该工具还可以通过包装这些API来追踪GPU内存分配调用和其他API使用情况。HPCToolkit收集的数据存储在一个结构化格式中（包含二进制性能数据文件的实验目录）。可视化通过HPCToolkit的专用查看器完成：HPC‐Viewer用于聚合分析，HPC‐TraceViewer用于时间线。该设计强调事后分析——数据以低干扰方式收集，然后分析以产生调用路径分析、追踪和分析的混合。HPCToolkit的开放设计允许添加对新平台的支持（最近版本添加了Intel oneAPI Level0甚至OpenMP卸载追踪）。它的优势在于将GPU工作关联回CPU上下文（重建GPU调用上下文树）并以供应商无关的方式进行（例如，将warp和wavefront统一呈现为抽象线程）。一个限制是配置和处理跟踪可能很复杂，细粒度GPU采样（特别是在NVIDIA上）可能会扰乱执行（早期CUDA版本将启用PC采样的内核序列化，导致开销）。尽管如此，HPCToolkit展示了混合方法：采样+追踪，大量使用供应商接口但也使用静态分析将性能指标归因到代码。

GPUprobe (基于eBPF的CUDA追踪)

GPUprobe是一个最近的开源工具，展示了使用Linux eBPF进行零插桩GPU监控。与修改目标应用程序或需要供应商库不同，GPUprobe使用操作系统级的动态追踪。在内部，它将eBPF uprobes（用户级探针）附加到NVIDIA CUDA运行时库（libcudart.so）中的函数上。例如，它在cudaMalloc、cudaFree和cudaLaunchKernel上放置探针，以拦截应用程序何时调用这些函数。当调用这样的函数时，uprobe会触发内核中的一个小型BPF程序，该程序可以记录参数（例如分配的大小）和时间戳。这些数据存储在内核缓冲区或队列中。然后GPUprobe的用户空间组件定期从这个缓冲区检索事件。因此，该工具记录GPU内存分配模式、内核启动时间和其他CUDA API使用情况，无需对目标进程进行任何代码更改或重新编译。这种方法利用了Linux内核能够以最小开销"窥探"用户级函数调用的能力——本质上是通过BPF插入"微断点"。相对而言，开销很低，因为像cudaMalloc()或启动内核这样的调用已经有相当大的成本；额外的上下文切换和日志记录（BPF高效完成）只增加了很小的延迟。在基准测试中，GPUprobe报告的典型使用开销低于4%。收集的指标可以实时显示或导出。GPUprobe包括一种定期输出到stdout的模式和一个HTTP Prometheus指标端点（:9000/metrics）以与监控系统集成。本质上，GPUprobe填补了全功能性能分析器和最小监控之间的空白：它不捕获完整的GPU时间线细节或硬件计数器，但它让开发者可以快速查看内存泄漏、分配大小、内核调用次数和持续时间，零侵入性。作为开源（用Rust编写，带有BPF C片段），它可以扩展以钩住额外的CUDA API甚至其他GPU库。值得注意的是，GPUprobe专用于CUDA（NVIDIA）并需要带有BPF的Linux；如果类似的用户级钩子和符号可用，其概念理论上可以扩展到AMD或其他加速器。它是使用现代内核追踪（eBPF）进行GPU可观察性而不需要供应商特定SDK的典型例子。

NVIDIA DCGM (数据中心GPU管理器)

DCGM不是针对单个应用程序的性能分析器，而是一个在集群/节点级别的GPU遥测和监控套件。它的架构是面向服务的：DCGM作为一个后台主机引擎服务（nv-hostengine）运行，或者可以通过其libdcgm.so API嵌入到管理应用程序中。核心DCGM库与NVIDIA驱动程序接口交互，收集关于GPU的广泛指标——包括利用率（SM忙碌%，内存使用情况）、温度、功耗、时钟速度、内存错误，甚至吞吐量统计如PCIe传输字节数。这些指标中的许多来自与nvidia-smi相同的底层计数器（通过NVML，NVIDIA管理库）。DCGM聚合服务器上所有GPU的这些数据。它可以被配置为定期记录指标并基于策略触发警报或应用功率上限。开放核心设计意味着核心库和API是开源的，而一些诊断模块仍然是专有的。DCGM的主机引擎可以通过CLI (dcgmi)接受命令或通过DCGM-Exporter提供数据，后者与Kubernetes集成。DCGM-Exporter提供一个HTTP端点，持续以Prometheus格式导出GPU指标，允许轻松抓取并在Grafana仪表板中可视化。这使DCGM非常适合在容器化环境中进行长期运行监控。在内部追踪方面，DCGM主要使用驱动程序计数器轮询和某些健康事件的事件回调（如双位ECC错误）。它不是关于追踪单个内核启动，而是关于随时间聚合的性能和健康。例如，DCGM可以报告作业的平均SM占用率或内存带宽，但不能报告哪个函数导致了缓存未命中。它的操作粒度比Nsight或HPCToolkit粗糙——通常以秒为时间尺度——并且设计为具有非常低的开销（"作业运行时低开销、非侵入性健康监控"）。在安全方面，DCGM可以以有限的权限运行：默认情况下，可能泄露详细信息的GPU性能计数器被限制为管理员，但配置（如设置NVreg_RestrictProfilingToAdminUsers=0）可以允许非root DCGM使用。总之，DCGM的架构服务于系统管理员：它为跨节点的GPU指标收集提供了统一的API和服务，与集群管理器和云监控的集成点。它通过专注于全局GPU行为和健康而不是特定于代码的性能事件来补充每个应用程序的性能分析器。

TAU (调优和分析工具)

TAU是一个可移植的性能分析/追踪工具包，在HPC中广泛使用，它已经发展到支持GPU卸载，以及CPU和网络追踪。TAU的架构是模块化的：它可以使用源代码级插桩、二进制重写或拦截包装器来对应用程序进行插桩，并支持许多编程模型（MPI、OpenMP、Pthreads等）。对于GPU支持，TAU利用现有的性能工具接口：它通过CUPTI支持NVIDIA GPU，通过ROCm的ROCprofiler或roctracer支持AMD GPU，通过Level Zero工具支持Intel GPU。具体来说，TAU的运行时库可以订阅CUPTI的回调API来追踪CUDA内核启动、内存复制等，类似于HPCToolkit的方法。它还支持MPI工具接口（MPI_T）和Kokkos分析钩子，将高级并行运行时事件与GPU活动关联起来。TAU通常通过启动程序的包装器（tau_exec）使用，该包装器可以自动注入插桩。例如，tau_exec -cuda将启用基于CUPTI的所有CUDA API调用和内核的插桩。TAU为每个函数记录计时器，因此如果拦截了GPU内核启动，TAU将记录在CPU上花费在该启动调用中的时间以及通过CUPTI的异步活动记录在GPU上的内核执行时间。该工具还可以追踪数据传输大小和计时（例如cudaMemcpy）。除了显式插桩外，TAU还支持基于事件的采样。它可以配置硬件计数器或定时器定期对程序计数器进行采样，然后使用二进制调试信息（通过libdwarf/BFD的DWARF）将这些样本映射到源代码行。这对于不使用函数插桩的代码分析或捕获系统级异常很有用。TAU将性能结果存储在紧凑的二进制格式或数据库中（它有一个用于数据管理的PerfDMF框架）。可视化可以使用TAU的ParaProf GUI完成（显示每个函数的指标表、每个线程的时间线、调用图等），或者通过将追踪转换为OTF2等格式以使用Vampir等工具。TAU是可扩展的——用户可以使用TAU API标记自定义事件或代码中的阶段，这些事件或阶段将显示在性能分析中。类似插件的功能存在于TAU可以相对容易地与新的后端测量源接口（TAU团队添加了对新GPU甚至OMPT的支持，而不改变面向用户的API）。TAU的设计理念是通过标准工具接口（CUDA/CUPTI、ROCm工具、MPI PMPI、OpenMP OMPT）尽可能多地拦截，以便它不需要修改目标代码。这使它成为一个多平台"胶水"性能分析器，但也意味着TAU在给定平台上的能力受到该平台API暴露的限制。例如，TAU可以追踪在GPU内核和数据移动中花费的时间，但对于NVIDIA内核内部细节，它将依赖于CUPTI指标（TAU可以类似于Nsight Compute收集这些指标，尽管TAU通常专注于时间线和聚合时间而不是每个硬件指标）。总之，TAU的内部架构结合了插桩包装器（自动插入）和采样，并在性能数据库中统一数据。它提供了跨CPU、GPU和MPI的整体视图，这在应用程序跨越所有这些层的HPC设置中是有价值的。

(其他工具：除了上述工具外，还有特定供应商的GPU性能分析器，如AMD的Radeon GPU Profiler (RGP)用于低级AMD GPU追踪，Intel的VTune Profiler和Graphics Performance Analyzers用于Intel GPU，以及研究工具如rocProf（AMD上的ROCm命令行性能分析器）。这些通常遵循类似的模式：例如，RGP使用驱动程序级插桩和GPU着色器性能计数器来生成图形/计算波的时间线，VTune使用Intel的驱动程序在Intel GPU上追踪Level0或OpenCL内核。由于篇幅限制，我们重点介绍上述主要工具，它们涵盖了一系列实现策略。)

这些工具追踪的API和运行时

性能分析工具通过钩入软件栈的各个层来实现覆盖。表1总结了每个工具能够追踪的运行时、库或驱动程序API：

CUDA驱动程序/运行时(NVIDIA)： 几乎所有工具（Nsight Systems/Compute、HPCToolkit、TAU、GPUprobe）都支持CUDA。Nsight和TAU使用CUPTI拦截CUDA调用，HPCToolkit使用CUPTI进行活动，GPUprobe通过uprobes钩入libcudart。开源的Omnitrace目前不支持CUDA（专注于AMD），而DCGM通过NVML在更高级别监控（它看到整体CUDA使用情况但不是单独的调用）。
HIP/ROCm (AMD)： Omnitrace和TAU通过ROCtracer/ROCprofiler支持AMD的HIP运行时。HPCToolkit也使用ROC-tracer进行AMD GPU内核。这些允许追踪HIP内核启动、内存复制等。像Nsight（NVIDIA）这样的工具不支持AMD API。AMD的RGP和rocProf当然本地针对ROCm/HIP。
OpenCL： HPCToolkit有函数包装器来拦截OpenCL API调用，用于缺乏更好接口的GPU。TAU也支持OpenCL追踪。Nsight Systems可以在NVIDIA平台上追踪OpenCL（如果OpenCL实现通过CUDA驱动程序路由，则通过CUPTI）。Omnitrace可能可以通过其二进制插桩处理AMD上的OpenCL（如果配置为插桩这些调用）。
Vulkan/OpenGL（图形API）： Nsight Systems和NVIDIA工具可以追踪源自图形API的GPU工作（在时间线上显示GPU工作，如Vulkan启动），但我们的重点是计算。像RGP这样的工具也针对图形。一般来说，以HPC为中心的工具（HPCToolkit、TAU）默认不追踪图形API。
Intel oneAPI Level Zero / DPC++： HPCToolkit和TAU已经添加了对Intel GPU卸载的支持。HPCToolkit包装了Level Zero（oneAPI）API以捕获内核提交。TAU通过OMPT（OpenMP卸载）接口和Level Zero拦截器支持oneAPI。Intel的VTune也使用驱动程序来追踪这些，但第三方支持正在兴起。
驱动程序级钩子： 一些工具在驱动程序级别集成：DCGM通过NVML与GPU驱动程序接口以获取指标，可能直接与内核驱动程序交互以处理节流或健康事件。大多数单个操作的追踪是在用户API级别完成的（第三方的驱动程序级插桩很少见，因为GPU驱动程序是专有的闭源，除了AMD的开放驱动程序，ROCtracer可以插入其中）。

总结来说，NVIDIA的CUPTI是Nsight、HPCToolkit、TAU（以前还有nvprof、Visual Profiler）用来在驱动程序API级别追踪CUDA并收集计数器的常见接口。AMD的ROCm工具（rocTracer/rocProfiler）对HIP/HSA起着类似的作用。Intel的Level Zero更新；HPCToolkit/TAU中已有支持。这些工具也覆盖了CPU线程库（pthreads、通过OMPT的OpenMP、通过PMPI/MPI_T的MPI），以将CPU端阶段与GPU活动相关联。

CPU与GPU的追踪及事件关联

这些性能分析器的一个重要方面是它们是捕获CPU端的事件，GPU端的事件，还是两者都有——以及它们如何合并这些时间线。CPU端追踪通常涉及记录应用程序线程何时调用某些API或某些函数何时执行。例如，基于插桩的工具如TAU或Omnitrace会在CPU上调用CUDA内核启动函数时记录时间戳。基于采样的工具如Nsight Systems或HPCToolkit也记录CPU端样本（可以指示CPU何时忙碌或空闲）。GPU端追踪意味着捕获在GPU设备上发生的事件，如GPU上内核的实际执行间隔，或其他GPU引擎活动（GPU引擎上的内存复制等）。

大多数工具依赖于供应商支持来获取GPU端事件计时。例如，NVIDIA上的CUPTI提供内核在GPU上的开始和停止时间戳，Nsight、HPCToolkit和TAU都使用这些信息。类似地，在AMD上，ROCm的活动回调提供GPU内核持续时间。通过收集这些，工具可以绘制实际GPU利用率时间线。GPU硬件通常有一个单独的时钟域，但性能分析接口确保时间戳可以转换为主机时钟（通常通过在性能分析开始时采样共同时钟）。

关联然后通过按时间合并CPU和GPU事件流来实现。Nsight Systems在其时间线中做到这一点：可以看到CPU线程启动一个内核，然后稍后（在同一时间线上）看到GPU活动发生，通过时间戳对齐。HPCToolkit采用不同的方法，通过调用上下文进行关联：它记录特定的CPU调用路径启动了内核X，并在分析中将内核的GPU时间归因于该调用路径，有效地将GPU时间链接到CPU调用上下文树中。这回答了"哪个CPU函数导致了这个GPU工作，它花了多长时间？"这样的问题。TAU同样将GPU内核时间归因于启动它的CPU代码区域（具有像TAU_CALLSITE这样的功能来精确定位源代码行）。

处理上下文切换和并发是另一个挑战。Nsight Systems明确地通过perf或内核跟踪点追踪CPU上的上下文切换事件，因此它可以显示哪个线程正在运行与被抢占。在与GPU事件关联时，工具必须考虑GPU内核异步运行的事实：例如，CPU可能将内核排队并继续执行，而GPU稍后执行它。性能分析器通常包括流同步信息，以便它们可以显示在CPU调度之后仍在进行的GPU工作。

如果多个CPU线程向多个GPU流启动工作，关联确保正确的匹配。唯一ID（如CUDA流ID或HIP队列ID）在内部用于标记事件。例如，CUPTI将报告带有ID的内核启动，该ID对应于特定的CPU线程和流，使工具能够将其与CPU端启动API调用记录匹配。

总之，高级性能分析器追踪CPU事件（函数调用、线程调度等）和GPU事件（内核执行间隔、设备内存传输）。它们通过时间和元数据对齐这些事件，实现一个连贯的视图。这对于识别流水线延迟至关重要（例如，如果CPU正在等待GPU或反之亦然）。没有GPU端数据，人们可能只知道CPU何时启动了内核，而不知道它何时在GPU上完成。相反，仅GPU的追踪将缺乏CPU在做什么的上下文。因此，所有被审查的工具使用某种形式的组合追踪：例如，HPCToolkit的分析显示组合的CPU+GPU调用路径，而Nsight Systems使用覆盖CPU核心和GPU时间线的时间线并行显示。确保时钟同步（CUPTI处理CUDA的这一点）和处理异步性是它们实现的关键部分。

插桩与采样：数据收集策略

性能分析工具通常采用两种策略中的一种（或混合）来收集性能数据：插桩（事件追踪）或统计采样。每种策略对开销、准确性和侵入性都有影响。

基于插桩的性能分析： 这涉及插入钩子或探针，以记录感兴趣的事件每次发生时。例如，用开始/停止计时器包装函数（如TAU所做的），或使用在每次内核启动时触发的CUPTI回调。插桩可以在编译时（在源代码或二进制文件中插入代码）、运行时注入（使用LD_PRELOAD拦截库调用，或像Omnitrace那样进行动态二进制插桩）或在操作系统级别（使用GPUprobe所做的kprobes/uprobes）完成。优势是完整的信息：记录每个事件，提供确切的计数和计时。缺点是如果事件频繁，开销可能会很高。例如，对细粒度应用程序中的每个函数调用进行插桩可能会显著减慢它。像TAU和Omnitrace这样的工具使用插桩来获取详细的追踪（Omnitrace通过二进制插桩，TAU通过包装器库和基于编译器的插桩）。GPU驱动程序API通常通过官方回调接口（CUPTI、ROCtracer）进行插桩，以追踪GPU事件而不修改供应商代码。插桩还可以使用动态运行时方法：例如，LD_PRELOAD用于通过在加载时覆盖符号来拦截malloc或CUDA调用；二进制重写可以永久地在代码中添加探针；而uprobes/kprobes允许在函数入口处插入断点（GPUprobe使用uprobes在运行时附加到cudaLaunchKernel等）。每种方法在灵活性与性能之间有权衡。
基于采样的性能分析： 这种技术在间隔（基于时间或事件）而不是记录每个事件时收集信息。对于CPU，这通常意味着使用硬件计时器或性能计数器溢出中断定期采样程序计数器（PC）和调用栈。对于GPU，采样可能意味着定期查询利用率，或者在NVIDIA的情况下，使用GPU上的PC采样，如前所述。Nsight Systems使用定时采样来采样CPU线程（例如，以一定间隔捕获堆栈跟踪，以统计推断热点代码）。HPCToolkit严重依赖采样来分析CPU（并可选择使用GPU PC采样进行细粒度GPU分析）。采样的优势是开销低得多——例如，每秒捕获1000个样本几乎不会导致减速，但提供了统计覆盖，显示时间花在哪里。它还自动过滤掉非常快的事件（如果一个函数执行100ns，它可能永远不会被采样，从而有效地专注于主要成本）。缺点是精度降低：你可能会错过不频繁的事件或引入采样误差。此外，对于事件计时（如精确的内核持续时间），采样不合适——它更适合聚合指标，如"哪些代码平均消耗了多少周期"。
混合方法： 许多工具结合了两者。例如，Nsight Systems对GPU事件进行插桩（以精确计时内核），但对CPU活动进行采样（以降低开销）。HPCToolkit在粗粒度级别（GPU内核）进行插桩，并对其他一切使用采样。Omnitrace对函数进入/退出进行插桩，并对调用栈进行采样。这种混合策略允许捕获关键事件，同时保持开销可控。

在现代GPU性能分析中，事件追踪通常用于GPU操作，因为GPU工作通常是稀疏的（内核、内存复制）并由运行时API明确定义——记录每个内核启动是可行的。同时，采样通常用于CPU计算或长时间GPU内核内部，以了解其内部行为（如内核内的PC采样，本质上是在GPU执行中采样事件）。一个特殊情况是硬件计数器采样：一些工具设置硬件PMU来采样特定事件，例如计算缓存未命中并在溢出时采样指令地址（这合并了两个世界：设置计数器的插桩，采样来捕获事件）。

策略总结： 如果开发者需要精确的事件时间线和计数，首选插桩（用于追踪工具、时间线可视化等）。如果需要对大型应用程序进行低开销性能分析以查找热点，首选采样（用于HPC性能分析器、统计性能分析器）。许多工具提供两种模式。实现必须仔细管理插桩的启用/禁用（例如，支持选择性插桩——Omnitrace允许配置要插桩的函数，TAU允许通过标志切换MPI、I/O、GPU的插桩）。这有助于针对分析并减少数据量。

硬件性能计数器（PMU）的使用

硬件性能监控单元（PMU）对于在CPU和GPU上收集无法通过高级计时获得的低级指标（如指令计数、缓存未命中、内存吞吐量等）至关重要。

CPU PMU： 大多数CPU通过Linux perf等接口暴露计数器。像HPCToolkit和TAU这样的工具可以使用PAPI库来配置CPU计数器（周期、缓存未命中、浮点运算等）并在溢出时进行采样。例如，HPCToolkit的命令hpcrun -e PAPI_TOT_CYC@500000 -e PAPI_L2_TCM@200000将在总周期和L2缓存未命中达到特定计数后进行采样。通过使用这些计数器作为采样触发器，工具将硬件事件归因到代码位置。TAU也与PAPI集成，如果进行了插桩，它可以记录每个函数调用的计数器（在每个区域内累加计数）。在内核级分析（如操作系统调度）中，也可能通过ftrace或BPF使用CPU PMU，但这在面向用户的性能分析器中不太常见。

GPU PMU： 现代GPU有广泛的性能计数器（NVIDIA上每个流式多处理器，AMD上每个计算单元）。对这些的访问是通过供应商库实现的：

NVIDIA提供CUPTI/Event和Metric API，允许工具列出可用的计数器（例如，"dram_read_transactions"、"active_warps"）并读取它们。像Nsight Compute这样的工具编程这些计数器并在内核执行后读取它们。一些计数器可以在内核之后读取（给出该内核的聚合计数），使用CUDA的分析模式。其他可能会持续采样（NVIDIA支持"流式性能监视器"模式进行持续收集，但这通常是内部的）。
AMD通过ROCprofiler或GPUPerfAPI（GPAPI）提供类似的接口来访问性能计数器。Radeon GPU Profiler（RGP）和rocprof可以收集AMD GPU的计数器，如wavefront占用率、缓存命中等。Omnitrace可能在底层使用ROCprofiler（因为它是为AMD构建的）来获取计数器，除了跟踪回调。
Intel GPU（Xe）通过VTune或Metrics Discovery API暴露计数器，HPC工具才刚开始支持这些。HPCToolkit提到通过Level Zero以有限方式支持Intel GPU指标（可能链接到Intel的GT-Pin或使用oneAPI VTune后端）。

内核内计数器和PMU： 一些性能分析方法也使用操作系统级计数器。例如，Linux有GPU调度器事件的跟踪点（在一些开源驱动程序上），可用于跟踪GPU引擎上的上下文切换。此外，NVIDIA的驱动程序跟踪诸如内存页迁移、NVLink使用等，DCGM可以检索这些信息。DCGM使用一部分计数器来监控健康状况（如内存ECC计数）和吞吐量（PCIe字节），这些通常不向应用性能分析器暴露。

PMU数据使用： 原始计数器值通常需要后处理才有意义：

派生指标： 许多工具提供计算比率或百分比。例如，SM占用率派生自"活动warp"与"理论最大warp"计数器的比较。缓存命中率派生自缓存命中和未命中计数器。Nsight Compute明确计算这些指标，甚至提供指导性分析规则（例如，如果某个停滞原因计数器很高，就建议优化）。HPCToolkit以中立方式呈现原始指标（计数或百分比），但不专注于指导性提示。
采样与聚合： 当计数器被采样（如设置计数器触发中断）时，你得到基于样本的归因（如"20个L2未命中样本发生在函数X中"）。当计数器按内核聚合（如Nsight Compute所做）时，你得到该内核的确切计数，但必须归因于整个内核调用。一些工具可能使用PC采样结合计数器在内核内按源代码行进行细分（例如，将PC样本与"已发出指令"计数指标相关联，以估计代码每个部分执行了多少指令）。

开销： 访问GPU计数器通常需要对GPU性能监控硬件进行独占控制，这就是为什么一次只能运行一个性能分析会话。在NVIDIA上，如果一个计数器分析器（Nsight Compute或CUPTI指标）正在运行，另一个就不能运行，默认情况下只有特权用户可以使用计数器，因为可能存在旁路泄漏。AMD的ROCm可能允许更多并发使用，但仍有开销。工具必须错开计数器读取，以避免扰乱GPU流水线（因此Nsight Compute的内核重放确保原始执行不会因一次读取太多计数器而减慢）。

总之，PMU是低级指标的基础。工具实现者通过供应商API（CUPTI、ROCprofiler等）与这些交互，这些API抽象了直接寄存器操作。在某些情况下，开源工具已经逆向工程了部分内容：例如，通过未记录的perf事件读取Nvidia GPU性能计数器——但通常由于复杂性，使用官方API是常态。PMU的使用使得可以深入了解内存带宽（通过计算传输的字节数）、执行效率（已发出与已完成指令）、停滞（硬件可以在Nvidia上自Volta以来按类别计数停滞原因）等。这些指标大大丰富了性能分析器的输出，超越简单的计时，但代价是更复杂的数据收集。

插桩的钩子和技术

为了实现插桩，工具可能钩入软件栈的不同层。以下是常见技术：

用户模式API包装： 许多工具创建包装库，拦截对标准API的调用。例子包括TAU对MPI的拦截（通过PMPI），或包装CUDA运行时调用。这可以通过提供定义相同符号的库（例如，一个在记录后调用真实函数的cudaLaunchKernel）并使用LD_PRELOAD或链接顺序确保使用包装器来完成。HPCToolkit的基础设施有OpenCL和Level Zero的函数包装器，可能以这种方式操作。这种方法直接且不需要特殊权限，但你必须为每个感兴趣的函数实现一个包装器。它用于高级事件（例如，记录内核已启动，然后调用实际启动）。
编译器/源代码插桩： TAU可以插桩源代码或使用编译器标志（如使用-finstrument-functions或使用Clang/OMPT进行OpenMP）。这将探针编译到程序中。优势在于丰富的上下文（你可以在循环或函数入口处插桩，计时器开销最小），但需要重新编译或至少添加额外的编译步骤。
动态二进制插桩（DBI）： 像Omnitrace这样的工具执行DBI，意味着它们在运行时或预先将插桩注入二进制文件，无需源代码。DBI框架包括Dyninst、Intel PIN或DynamoRIO。Omnitrace被描述为对函数进行动态插桩。这很强大，因为它可以插桩即使是第三方库代码（源代码不可用）并可以动态配置。代价是复杂性和一些开销，因为代码修补和在DBI环境下运行。
内核探针（kprobes）和eBPF： 对于系统范围或内核级事件，性能分析器可以使用kprobes（插桩内核函数）或跟踪点。例如，要测量GPU驱动程序事件，可以在GPU驱动程序中插入kprobes，如果你知道符号（在调试之外不常见）。更实际地，eBPF被GPUprobe用来附加到用户级函数（uprobes），也可以附加到内核跟踪点（如调度或GPU中断）。eBPF的优势是它是动态的（无需应用程序重建）并且在生产中运行是安全的（在内核中沙盒化）。我们看到GPUprobe使用uprobes；类似地，可以将kprobes附加到，例如，GPU调度器的作业提交函数，以知道内核何时在硬件上调度。除了GPUprobe外，很少有性能分析工具目前使用eBPF，但这种技术正在引起人们对内核和用户事件低开销插桩的兴趣。例如，可以想象一个基于eBPF的性能分析器，通过每隔几毫秒轮询设备文件来采样GPU利用率，或跟踪影响GPU的操作系统事件（如GPU中断、DMA完成事件）。
驱动程序/内核插桩钩子： 一些供应商允许在驱动程序级别插入钩子。NVIDIA的CUPTI本质上在CUDA驱动程序内部注册钩子——NVIDIA为性能分析器构建了这个。AMD的ROCtracer通过与ROCm驱动程序栈接口类似地做到这一点。如果源代码可用（AMD的运行时是开源的），甚至可以修改或扩展它以添加自定义探针，但通常提供的钩子就足够了。在CPU上，操作系统提供跟踪点（如上下文切换事件、系统调用事件），Nsight Systems等工具用它们来跟踪调度。这些本质上是内核中的插桩点（例如，sched:sched_switch跟踪点）。工具可能通过perf子系统或ftrace接口启用这些。
GPU指令注入： 在某些情况下，插桩可以插入到GPU代码本身。这在商业工具中不常见（由于改变行为的风险），但研究工具或某些调试模式会这样做。例如，可以修改PTX（CUDA汇编）在某些点调用计数器增量。NVIDIA较旧的性能分析（CUPTI PC采样）是比显式插桩更好的方法，但NVIDIA可以在warp-trap指令处进行插桩。AMD的着色器性能分析可能在代码中插入"wave时间戳"。然而，通常供应商避免修改用户内核进行性能分析，而是在特殊模式下运行它们（如单warp重放或序列化执行以进行分析）。
Uprobe/Uretprobe用于参数/返回： 如GPUprobe所示，有时需要同时插桩函数的入口和出口，以获取参数和返回值。Uretprobes捕获函数返回，但由于寄存器可能被破坏，使用入口uprobe（存储参数）和返回探针的组合。这种技术需要测量诸如分配大小（入口参数）和分配指针（退出时返回）等内容，而不改变函数。
事件缓冲区和异步日志记录： 插桩通常将数据记录到缓冲区以最小化干扰。例如，CUPTI将活动记录写入缓冲区，性能分析器稍后读取。eBPF将事件写入BPF环形缓冲区或映射，由用户空间异步消费。这种解耦对性能至关重要：插桩函数触发快速日志写入（例如，只在无锁队列中放置事件）然后继续，而另一个线程或进程处理诸如写入磁盘或网络等重任务。

总体而言，实施插桩是关于仔细选择在哪里拦截（用户API、库或内核）、如何注入探针（源代码vs二进制vs动态）以及如何处理数据（立即打印vs缓冲）。我们讨论的工具说明了这些选择：例如，TAU主要使用用户API包装，Omnitrace使用DBI，GPUprobe使用内核uprobes，Nsight依赖驱动程序提供的钩子，HPCToolkit使用驱动程序钩子和自己的包装器的混合。

GPU指令级追踪技术

在GPU上进行指令级或warp级别的性能分析极具挑战性，因为操作量庞大且可观察性有限。然而，存在一些方法：

GPU PC采样： 如前所述，NVIDIA GPU支持在每个SM上定期采样程序计数器。这类似于CPU PC采样，提供关于哪些指令（或源代码行）在GPU上频繁执行的统计视图。HPCToolkit使用这一点将时间或周期归因到GPU内核中的代码行。CUPTI PC采样机制还可以归因采样warp的停滞原因（例如，warp是否因内存、执行依赖等而停滞）——提供关于为何性能损失的洞察。这有效地提供了指令级洞察（代码的哪些部分是热点，它们引起什么停滞）。开销可能不小，早期实现会序列化执行，但较新的GPU通过将样本收集卸载到硬件缓冲区更有效地处理PC采样。
硬件流水线统计： 与其追踪每条指令，GPU有计数器来统计"每周期发出的warp指令数"或"因X原因停滞的周期数"等。Nsight Compute读取这些来生成指令流水线利用率分析（例如，显示IPC、发出停滞细分）。虽然不是逐条指令的追踪，但它提供了指令级别类别的聚合视图。其中一些计数器实质上将warp的周期分类（内存、计算等），这几乎就像是随时间追踪流水线气泡的原因。
Wavefront/warp追踪调试模式： 在调试上下文中，GPU可以被置于模式中，让单个warp运行到完成或记录特定事件。AMD的GPU PerfStudio有低级模式，NVIDIA有如Nsight Graphics这样的工具，可以捕获每个绘制调用的计时直至着色器指令（用于图形）。但对于通用计算，由于开销，很少使用这种指令追踪。
逆向工程和未记录的方法： 开源开发者有时会逆向工程着色器ISA以获取指令信息。例如，HPCToolkit的开发者编写了自己的CUBIN分析器，将指令地址映射到行号，因为NVIDIA不提供SASS的源代码行映射。类似地，在AMD上，可能解析代码对象（其ISA）以与HSAIL或源代码关联。这些是离线分析，但对于指令级性能分析至关重要，因为工具必须将性能指标归因到单个指令或行。如果没有供应商文档，这涉及大量自定义工具开发。

总之，GPU上的指令级追踪是通过智能采样和计数器而非逐事件追踪（这会使任何系统不堪重负）。PC采样和专用计数器的组合提供了指令追踪的近似，足以突出低效指令或代码中的停滞部分。这种详细程度通常只在优化特定内核时需要（因此使用Nsight Compute或学术项目如GPU性能模拟器等工具）。大多数通用性能分析器不会输出执行的每条指令的日志（那更属于模拟器或非常特殊的调试模式的领域）。

数据收集、传输和存储

性能分析器收集大量数据——它们如何管理这些数据与收集本身一样重要。关键考虑因素是：数据如何从GPU移动到主机，如何缓冲，使用什么文件格式，以及用户最终如何访问它。

从GPU到CPU的数据传输： 收集GPU性能数据（计数器或跟踪）时，数据必须传输到主机。例如，CUPTI的Activity API在缓冲区中累积记录，然后通过驱动程序回调将它们刷新到主机内存。这通常在内核完成后或缓冲区填满时发生。对于连续指标，DCGM或其他工具通过驱动程序调用定期查询GPU（在底层通过PCIe读取寄存器）。数据传输的开销通过批处理减轻：性能分析器使用大缓冲区，只偶尔中断GPU以排出数据。一些高级用法可能会固定缓冲区，让GPU通过DMA输出数据（NVIDIA的NVPW流式计数器做类似的事情）。但细节通常由供应商库抽象。
缓冲和流式传输： 大多数工具实现双缓冲或队列，使数据收集和写入磁盘不会阻塞应用程序。例如，HPCToolkit将GPU事件记录到内存缓冲区，在运行后写入文件。Nsight Systems增量写入内存映射文件（如果连接，甚至可以流式传输到UI）。GPUprobe使用BPF队列，用户空间每隔几秒读取一次。实时流式传输跟踪数据很棘手，因为数据量大，但一些工具（如某些调试模式或Intel GPA）通过将事件通过套接字发送到GUI允许实时查看。在HPC中，通常数据被写入本地存储以进行事后分析，因为数据量很大（多GB跟踪）。
文件格式： 有多种格式：
专有二进制：Nsight Systems和Compute使用自己的二进制格式（Nsight Systems中的.nsys-rep实际上是SQLite数据库；Nsight Compute的.ncu-rep是带有部分的结构化二进制）。这些针对它们的GUI进行优化，不打算手动编辑。
开放跟踪格式：一些开源工具使用标准格式如OTF2（开放跟踪格式）或CTF。例如，TAU可以导出为OTF2。这些格式设计用于合并HPC中多个等级的跟踪。
Perfetto/Chrome JSON：Perfetto格式（protobuf跟踪）或较旧的Chrome Trace Event格式（JSON）被多个工具使用，因为Chrome的跟踪基础设施是事实上的标准。例如，PyTorch的性能分析器输出Chrome JSON跟踪，TensorBoard或chrome://tracing可以显示。Omnitrace写入Perfetto .proto文件，可在Perfetto UI中查看。GPUprobe不产生完整时间线，但可以扩展以输出事件的Chrome跟踪JSON。
纯文本/CSV：许多工具也生成人类可读的摘要报告。Nsight Compute CLI可以输出每个内核的CSV指标。HPCToolkit生成顶级热点的文本报告。DCGM的dcgmi可以将当前指标打印到控制台，DCGM导出器为Prometheus输出文本。
数据库：一些工具将数据存储在SQL或自定义数据库中。如前所述，Nsight Systems在底层使用SQLite；TAU的PerfDMF历史上使用数据库存储性能分析数据以供查询。使用数据库可以便于事后查询特定指标。
导出路径和遥测：工具在如何检索数据方面有所不同。
基于文件的工作流：例如，运行nsys profile ...生成稍后打开的文件；运行hpcrun然后hpcstruct/hpcprof生成分析数据库。
实时UI：例如，Nsight Systems也可以与附加到远程应用程序的GUI一起运行，流式传输数据。Intel VTune可以附加到正在运行的进程，定期在UI中更新统计信息。
Web仪表板：一些现代性能分析器有Web界面。例如，NVIDIA在其云中推出了Web版Nsight Systems，Omnitrace结果可以轻松加载到Perfetto网页查看器中。DCGM的Kubernetes集成专门用于仪表板：DCGM导出器提供Prometheus数据，在Grafana或其他工具中可视化。
与TensorBoard集成：针对ML的工具（如PyTorch的Kineto或TensorFlow的性能分析器）导出TensorBoard性能分析插件可以显示的数据。这些通常使用Chrome跟踪格式。例如，使用Kineto的PyTorch可以捕获CUDA内核时间线并在TensorBoard中显示，有效地提供类似Nsight的视图，但在ML工具环境中。底层，Kineto使用CUPTI（与Nsight Systems类似的数据），但转换为Chrome格式。
自定义遥测API：少数工具允许程序化访问。DCGM有C和Python API，所以你可以从应用程序或脚本中查询指标。CUPTI提供了一种程序化方式来开始/停止追踪并获取记录（所以可以将CUPTI嵌入到应用程序中以收集自己的跟踪）。一些HPC监控系统通过钩入这些API集成性能分析。
数据量和减少：单次运行可能产生巨大的数据（每个内核启动的跟踪等）。工具实现过滤和详细级别。Nsight Systems有选项限制收集（例如，只跟踪CUDA，如果不需要则忽略操作系统调度）。HPCToolkit可以配置为只跟踪GPU而不收集CPU样本等。它们还有选项包括/排除特定内核或API调用。这很重要，因为在长时间运行中写出每个事件可能不可行。一些工具，如Score-P（在大规模场景中与TAU一起使用），允许在性能分析（聚合计数）和跟踪（详细事件日志）模式之间切换，因为完整跟踪超过一定点就不可扩展。

总结，这些工具中的数据处理涉及事件的仔细缓冲，使用高效的文件格式（通常是二进制或压缩的），并提供直接可视化或导出到其他分析框架的方式。趋势是标准化使用Perfetto/Chromium trace格式用于时间线数据，JSON/CSV用于摘要指标，以便于集成。像Omnitrace这样的工具明确拥抱网页可视化，生成Perfetto跟踪，而HPC工具保持更加自定义但提供转换工具。在集群监控中，流式导出（通过DCGM导出器使用Prometheus）对实时可观察性至关重要。因此，我们看到事后分析工作流（HPC，桌面性能分析）和实时监控（数据中心，云）根据用例的不同而共存。

可视化和分析界面

性能数据的呈现方式对于获取洞察至关重要。不同的可视化帮助回答不同的问题：

时间线图： 这些由Nsight Systems、Chrome/Perfetto跟踪查看器、HPCToolkit的TraceViewer等提供。时间线在水平轴上显示时间，在垂直轴上显示不同的活动（线程、GPU流、网络传输）。这对于识别并发问题和空闲间隙非常出色。例如，可以看到GPU在CPU执行某些操作时是否空闲（表明潜在的CPU瓶颈或排队工作不足），或者CPU线程是否空闲等待GPU（可能是同步延迟）。在时间线中，你可能会看到重叠的矩形，其中一个内核在另一个CPU线程运行时执行——这是良好的利用率——相对于大的间隙，这表示流水线停滞。依赖关系或事件链接可以被绘制（Nsight显示从CPU启动事件到它触发的GPU内核执行事件的箭头），所以你可以看到关系和任何排队延迟。时间线还可以显示CPU上的上下文切换（上下文切换前后使用不同的线程颜色），帮助识别是否有太多上下文切换影响性能。总体而言，时间线可视化提供了流水线瓶颈（例如，GPU等待来自磁盘/CPU的数据）、重叠计算/通信和跨设备的负载平衡等洞察。
火焰图/调用栈图： 火焰图通常由采样数据产生——它可视化随时间或聚合的调用栈分析。HPCToolkit的HPCViewer和TAU的ParaProf可以显示类似火焰图的内容：基本上，哪些函数调用了哪些函数，以及每个函数花费了多少时间（通常绘制为堆叠条或"火焰"，宽度对应于时间）。这有助于根据调用上下文识别热点：例如，你可以看到函数foo()占用了30%的时间，主要是在从main()调用时，而不是从init()调用时。它还可以暴露递归或昂贵的调用链。在GPU上下文中，HPCToolkit可以显示GPU内核执行的火焰图，归因于启动它们的CPU调用路径——这是独特的，因为它将设备时间连接到火焰图中。火焰图（如Brendan Gregg的flamegraph工具生成的）通常是样本数据的静态可视化；许多性能工程师使用它们来快速了解时间去向。
有向无环图（依赖图）： 一些工具（尤其是在并行运行时上下文中）可以可视化任务图或依赖关系。例如，在GPU操作的流水线中，图视图可能显示内核A和B并行运行，然后流入C。NVIDIA的Nsight Compute不这样做，但Nsight Systems通过时间线隐式显示依赖关系。其他分析工具可能重建计算和数据移动的图（这更为小众，通常手动完成或通过自定义分析，但不是列出工具的主要功能，除了可能作为可视化异步依赖关系的箭头部分）。
散点图： 当你有许多类似事件并想查看分布时，这些很有用。例如，绘制内核执行时间与内核启动索引的关系可以显示后面的内核是否变慢（可能由于热节流或争用）。散点图还可能显示两个指标之间的相关性——例如，可以绘制内核的已实现占用率与执行时间的关系，以查看低占用率是否与较长的运行时间相关。工具通常不在其UI中包含散点图，但用户导出数据进行此类分析。一个例外是一些研究/专业工具，它们允许在内核之间绘制一个指标与另一个指标的关系。
热图： 热图可能用于显示，例如，GPU单元利用率的时间线热图（一些图形工具显示每个SM的利用率随时间变化的热图）。另一个用途是可视化内存访问模式——例如，GPU内存地址与时间。这些是专门的，通常不在通用性能分析器中找到，但可以生成，例如，缓存未命中率随时间片段的热图，以发现计算阶段。在性能分析文献中，"GPU利用率热图"可能指的是显示每个引擎随时间的忙碌程度。
指标仪表板： 一些可视化仅仅是指标的表格或图表。Nsight Compute的GUI显示指标表，并可能以红色突出显示有问题的指标。DCGM与Grafana将显示指标随时间变化的线图（如GPU温度或利用率%）。这些有助于监控和趋势分析（作业是否饱和内存带宽？GPU利用率是否在某个点下降？）。

每种可视化提供不同的洞察：

从时间线中，可以观察到CPU-GPU不同步：例如，CPU启动内核然后什么都不做（空闲）直到结果返回——表明CPU可以做其他工作或内核是瓶颈。或者相反：GPU空闲是因为CPU没有提供足够的工作（可能是单线程CPU代码导致GPU饥饿）。这些场景在时间线上明显表现为大的间隙。
流水线瓶颈：例如，在深度学习中，时间线可能显示CPU上的数据加载很慢，使GPU等待（常见问题）。或者如果有多个GPU，可能会看到一个GPU提前完成然后在同步点等待——表明负载不平衡。
内存带宽争用：如果性能分析器捕获内存吞吐量，可能会看到它在某些内核运行时达到最大（接近100%利用率），解释为什么这些内核不会进一步加速（内存受限）。DCGM和Nsight可以报告PCIe或NVLink使用情况；高使用率可能解释延迟，如果作业移动大量数据。如果运行多个进程，时间线甚至可能显示重叠的数据传输，这些传输使总线饱和。
缓存未命中和停滞：像Nsight Compute这样的工具列出缓存命中率和停滞原因。它们可能不会可视化它们，但报告有效地突出了内存争用（例如，高L2未命中率和许多内存停滞周期）。用户可以从这些指标推断内存延迟正在影响性能。
CPU线程问题：火焰图或时间线可能显示过多的上下文切换或锁争用（如果插桩）。Nsight Systems可以显示CPU线程的互斥等待时间，帮助找到CPU上的并发问题，这些问题间接减慢了GPU供给。
依赖图（如果手动分析）可以突出显示异步工作中的关键路径——例如，内核可能可以并发运行，但由于可优化的依赖关系而按顺序调度。
在GPU代码上的火焰图（通过PC采样）可以揭示内核代码的哪一部分是热点。例如，HPCToolkit的GPU PC采样可能显示内核中40%的时间花在特定循环（源代码行X）上，暗示该循环是内存受限或计算密集型，适合优化。

实际上，许多工具集成了多种视图。Nsight Systems有时间线，加上放大查看特定时刻每个线程的调用栈的能力。TAU的ParaProf可以显示分析表和时间线（如果启用跟踪）。HPCToolkit有单独的查看器用于聚合分析（如自上而下的调用树与指标）和事件时间线（如果需要精细细节）。这种组合很强大：人们可能使用分析找到主要罪魁祸首函数，然后检查时间线看看它为什么慢（可能是并发问题），然后使用硬件指标视图看看它是CPU受限还是内存受限等。

可扩展性和自定义插桩

不同工具提供不同程度的可扩展性——即能够整合新的事件源或定义自定义指标和插件：

用户标记和注释： 许多性能分析器允许用户标记代码的部分。NVIDIA提供NVTX（NVIDIA工具扩展）——一个开发者可以注释代码区域或事件的库。Nsight Systems将在时间线上显示这些NVTX范围（用于标记应用程序的阶段很有用）。类似地，对于CPU，有像Intel的ITT（插桩和跟踪技术）这样的注释API，一些工具（如VTune或oneAPI工具）使用；TAU有自己的API（例如，TAU_PROFILE_TIMER_START(name)）。这些不是工具的扩展本身，但它们允许用户注入工具将记录和可视化的自定义事件。这对于理解高级阶段（如在时间线上的"数据预处理"、"训练步骤"等）至关重要。
插件架构： 一些工具有模块化设计，可以添加新模块。例如，DCGM的诊断是模块化的——可以添加新的诊断插件（尽管编写一个可能需要NVIDIA的合作，因为不是完全开源）。HPCToolkit不是真正基于插件的，但由于是开源的，用户已经添加了功能（比如可以通过为其API编写新连接器来添加对新GPU的支持）。TAU相当可扩展；它是学术性能工具社区的一部分（与MPI_T、OMPT等标准一致）。如果出现新的运行时（比如新的任务库），TAU可以为其添加接口。
定义新事件： 在像TAU和HPCToolkit这样的工具中，你可以通过使用工具的API（对于应用开发者）或由工具维护者为新库函数添加包装器来插桩新事件。例如，当OpenACC和OMPT接口出现时，TAU迅速集成了这些，这意味着它们有效地"教会"了TAU处理新事件类型（OpenACC卸载事件、OpenMP运行时事件）。类似地，HPCToolkit添加Level Zero支持是实现新回调处理程序和包装器的问题。因为HPCToolkit的设计是供应商中立的，添加一个API意味着用另一个连接器扩展其基础设施。Omnitrace作为AMD研究项目，通过添加CUPTI集成（如果有人实现）可能扩展到支持CUDA——框架已经存在（二进制插桩和Perfetto输出），它只需要来自NVIDIA的事件源。
数据收集的定制： 一些性能分析器允许脚本或配置来决定收集什么。例如，Nsight Compute有"部分"和规则——高级用户可以在Nsight Compute的部分文件中定义自定义指标集甚至自定义公式来派生指标（有一种方式在Nsight中编写自定义分析规则）。这是一种针对指标的有限形式的可扩展性。类似地，可以选择要收集哪些计数器（减少开销）。像rocprof这样的AMD工具允许用户通过配置文件指定启用哪些计数器和哪些API回调。
输出集成： 可扩展性还意味着数据可以多么容易地用于其他地方。提供标准格式（JSON、CSV、Prometheus等）的工具实际上是可扩展的，因为它们允许将数据插入到自定义管道中（如自定义ML性能仪表板或自动回归测试器）。Omnitrace与Hatchet（这是一个分析库）的兼容性意味着用户可以通过编写Python来过滤或比较分析来扩展分析。TAU有一个接口来合并多次运行的分析，可用于自动化测试和分析。
可扩展性的限制： 值得注意的是，像Nsight Systems/Compute这样的闭源工具不支持插件——你不能向Nsight添加超出NVIDIA提供的新事件类型。你限于使用NVTX标记进行自定义注释。相比之下，开源工具（HPCToolkit、TAU、Omnitrace、GPUprobe）允许你修改或扩展它们，但需要编程。例如，GPUprobe可以通过在其源代码中添加更多BPF探针来扩展追踪新的CUDA函数甚至其他库。
脚本和自动化： 一些性能分析框架（尤其是在HPC中）支持自动化性能分析实验。例如，使用TAU的Python接口或脚本TAU commander在不同运行中收集不同的指标，然后综合结果。这不是插件，但它表明该工具可以集成到更大的工作流程中（如自动调优循环或持续集成以捕获性能回归）。

总之，可扩展性范围从简单的应用内注释（让终端用户标记事件）到能够集成新后端（让工具开发者支持新硬件）。调查的工具显示，开放框架和HPC工具优先考虑灵活性（因为它们需要针对不断发展的平台），而供应商GUI优先考虑在固定范围内的完善体验。对于构建新性能分析器的人来说，利用标准（如CUPTI、ROCtracer、OMPT）是一个很好的起点，因为它立即授予对各种事件的支持，并设计工具接受插件模块或配置文件来获取新指标可以使其面向未来。

开销和侵入性权衡

性能分析本质上会扰乱目标程序——关键是最小化并理解这种开销。不同工具有不同的影响：

轻量级采样 – 低开销： 低频率采样（例如，每秒100-1000个样本）通常产生低于5%的开销，往往不引人注意。HPCToolkit的CPU采样属于这一类别，使其适用于大型HPC作业，在这些作业中即使增加10%的开销也太多。HPCToolkit明确目标是每个线程每秒几百个样本，以保持合理的开销。类似地，Nsight Systems的CPU采样设计为非常低的开销，使用为此优化的操作系统级perf。权衡是你可能不会捕获非常短暂的函数或获得精确的计数。
函数插桩 – 中到高开销： 在每个函数进入/退出或每个CUDA调用上包装或插入代码可能会增加显著的时间，特别是如果这些函数很频繁。TAU的开销取决于插桩的粒度：只插桩外部循环或MPI调用成本低，但插桩每个C++内联函数可能开销巨大。TAU允许选择级别（例如，如果试图减少开销，你可能不会插桩每个GPU内核启动，只插桩整体区域）。Omnitrace的动态插桩在启动时会有一些开销，以及每个插桩函数调用的运行时开销（他们试图最小化它，但它确实存在）。Omnitrace文档强调将其用于整体执行分析，而不是每个内核的超精细细节（实际上建议如果想深入了解单个内核，使用另一个工具——暗示Omnitrace专注于广泛分析，而非微优化单个内核）。
GPU追踪开销： 仅通过CUPTI追踪内核事件开销很低——每个内核记录事件几微秒。如果内核数量众多（每秒数千个），这可能开始变得重要。CUPTI的活动缓冲通常可以高效处理许多事件。但当你开启更详细的GPU分析（如计数器或PC采样）时，开销增加。CUPTI的PC采样可能降低应用程序性能，尤其是在旧GPU中，它阻止了内核的重叠。工具通常对此发出警告；例如，Nsight Compute重放内核，因为试图在一次运行中收集所有指标会极大减慢内核或由于计数器限制而不可能。所以Nsight Compute是有意侵入的——它牺牲实时性能换取数据。
内存和存储开销： 性能分析还可能增加内存使用（缓冲区）并产生大型输出文件。如果应用程序对内存敏感，大型追踪缓冲区可能会扰乱缓存使用或内存分配模式。例如，启用详细追踪可能会为缓冲区分配数百MB，可能导致应用程序更多缓存未命中。大多数工具让你配置缓冲区大小，在丢失事件与内存开销之间权衡。
并发性扰动： 侵入性插桩可能改变调度。如果性能分析导致函数花费2倍时间，那可能改变线程如何交错或GPU流如何获取工作。例如，CPU上的大量插桩可能导致向GPU提供工作变慢，使GPU看起来利用不足，而实际上是性能分析器的开销导致的。好的性能分析器尝试量化或限制其开销，以避免误导用户。有些提供"开销补偿"，或至少测量在性能分析代码本身中花费的时间（TAU有时包括"性能分析器开销"指标）。
准确性与侵入性： 你想要的准确性/细节越多，通常产生的开销越大。例如，记录每个内存分配（如GPUprobe所做）可能减慢频繁分配的应用程序，但GPUprobe发现在典型CUDA代码中，相对开销很小，因为这些调用不是极其频繁且已经需要时间。如果有人运行每秒调用cudaMalloc数千次的微基准测试，GPUprobe的开销可能变得更明显（由于每次调用的上下文切换）。类似地，采样有可调频率：如果你在每个CPU核心上以10 kHz采样，可能开始看到几个百分点或更多的开销，更令人不安的是，你可能稍微干扰CPU涡轮频率或功耗。通常，选择一个安全范围作为默认值（如1 kHz）。
并行可扩展性的开销： 在分布式运行（带GPU的MPI）中，还需要考虑性能分析工具可能增加通信或I/O开销。例如，如果每个等级尝试写入大型追踪文件，I/O负载可能减慢速度。一些工具合并数据以避免太多输出（Score-P合并追踪等），但这可能在程序结束时造成同步开销。如果不小心，开销可能随规模成倍增加，这就是为什么在大规模HPC中通常首选采样（每个进程固定开销）。

工具比较：

Nsight Systems: 设计用于低开销系统性能分析，默认设置下通常<5%开销（如果增加CPU采样频率或追踪极多事件则更高）。它通常可以在生产中短时间运行。
Nsight Compute: 对目标内核有非常高的开销（可能是数量级），不适合生产运行——它是开发者的诊断工具。
HPCToolkit: 采样的开销非常低（适合长时间HPC运行）。如果使用GPU PC采样，GPU上的开销可能中等（根据采样率，该内核可能减速2倍）。但HPCToolkit默认只追踪内核，几乎不增加开销。
Omnitrace: 中等开销——潜在的许多函数动态插桩和采样结合。它针对大型应用优化，但仍在做很多事。对于中等插桩，他们可能实现低几十个百分点的开销，但用户应该选择更窄的焦点以保持开销可控。
GPUprobe: 对于它所做的事情，开销非常低（<4%），因为它只钩入几个慢调用。它明确宣传为轻量级。
DCGM: 极低开销。它在后台运行并只轮询指标（这在GPU上通常可忽略——读取计数器寄存器不是免费的，但微小，而且DCGM针对最小影响优化）。除非你设置极高的采样频率，否则它不会对GPU或CPU造成太大困扰。它适合始终开启的监控。
TAU: 差异很大。如果只使用采样或粗粒度插桩（MPI调用、高级循环），开销很低。如果使用自动源代码插桩以细粒度，开销可能很高（可能是2倍或更多运行时间）。TAU提供许多选项来调整这一点，如从插桩中过滤出小函数。

权衡总结： 对于日常性能工程，通常从低开销广泛工具（如Nsight Systems或HPCToolkit采样）开始找主要问题，然后使用更高开销的针对性工具（如Nsight Compute或在TAU中添加更多插桩）深入特定内核或代码区域。新工具的实现者必须决定在这个光谱上的位置——新的始终开启监控工具应该是低开销的（可能基于采样，如流式GPU性能计数器），而新的深度分析工具可以是侵入性的但谨慎使用。用户通常需要意识到开销，不要将性能分析器导致的减速与真实行为混淆（这就是为什么使用多个工具交叉检查或以不同采样率运行是良好实践）。

安全和沙盒考虑

性能分析工具通常需要低级访问，这在共享系统或云中引发安全和隔离问题：

权限要求： 如前所述，在GPU上访问硬件计数器通常默认限制为管理员（NVIDIA驱动程序默认NVreg_RestrictProfilingToAdminUsers=1防止非root用户使用CUPTI指标或Nsight）。许多工具（Nsight、HPCToolkit）如果不禁用该限制或以root运行，将无法捕获某些指标。原因是性能计数器可能被用作旁路通道（例如，计时另一个进程的GPU使用情况）。DCGM同样可能需要提升权限进行某些操作，但其开源核心在适当配置下可以在用户空间运行。
容器化： 在容器（Docker/Kubernetes）内运行性能分析器很棘手。按设计，容器限制了对性能监控的访问。例如，perf_event_paranoid sysctl在整个系统范围内适用——在容器内部，如果主机设置不够宽松，可能不允许性能分析。像Nsight这样的工具如果获得正确的权限可以在容器中运行（NVIDIA容器工具包可以传递GPU性能分析能力）。Kubernetes中有一个官方的DCGM容器用于监控GPU，它以特权方式运行以收集数据，然后通过网络端点以安全方式向非特权消费者提供指标。如果在没有权限的集群上尝试在容器中运行HPCToolkit或TAU，它们可能无法附加到GPU事件。一些云提供商出于安全考虑明确禁用其GPU上的性能分析。
多租户GPU： 如果两个用户共享一个GPU（通过MPS或虚拟化），性能分析可能失败或潜在地暴露另一个用户的活动。通常，供应商驱动程序通过仅限管理员的限制防止这种情况。AMD的ROCm可能允许多个用户进程，但rocm性能分析可能只能看到你自己的队列（尽管通常仍然需要root访问性能计数器）。在多租户场景中，运营商通常会允许监控（如DCGM聚合统计信息），但不允许任意用户进行详细追踪。这是新性能分析工具必须考虑的事项——你可能必须与权限模型集成。例如，站点可能以特权服务运行性能分析器，用户可以请求数据，而不是将二进制文件提供给用户运行。
插桩方法的安全性： eBPF和ptrace等技术（一些工具用于附加到进程）受Linux功能限制。默认情况下，一个进程不能ptrace另一个不同用户的进程，非特权BPF可能受限（最近的内核允许一些带限制的非特权BPF，但uprobes可能需要CAP_SYS_ADMIN，取决于配置）。因此，GPUprobe可能需要root或至少BPF功能才能将uprobes附加到另一个进程的libcudart。这限制了在生产中使用这样的工具，除非安排了适当的权限。
数据净化： 性能分析数据可能无意中捕获敏感信息（例如，地址、暗示数据大小的分配大小，甚至可能是专有算法的内核名称）。因此，共享性能分析结果必须谨慎。一些工具在需要时匿名化某些数据（例如，可以过滤掉内核名称或内存地址）。
DPU和隔离环境： 在DPU（如Nvidia BlueField，它有Arm CPU并运行自己的操作系统）中，你可以像在任何Linux服务器上一样对DPU进行性能分析（因为BlueField运行Ubuntu）。但如果DPU在主机卸载任务到它的模式下运行，分析两者之间的交互很复杂。可能需要分别分析主机和DPU，然后关联。出于安全考虑，DPU通常隔离主机视图——主机可能无法轻易地在没有权限的情况下窥探DPU的进程。因此，分析DPU可能需要登录到DPU的操作系统。这是一个新兴领域；像DCGM这样的工具目前不涵盖DPU（DCGM是GPU特定的）。未来的"DPU分析器"可能需要考虑网络安全（确保追踪网络数据包不会暴露数据内容）和多租户（DPU可能为多个VM提供服务）。
副作用： 一些插桩（如写入/proc或加载内核模块）在加固环境中可能被禁止。例如，要分析GPU，可能需要加载内核模块（Linux nvidia 驱动程序有一些性能分析钩子）。在HPC中心，管理员通常预先安装并启用需要的配置。在云中，可能需要容器权限或特殊编排。

简而言之，在沙盒环境中进行性能分析通常需要管理员有意启用。内部可观察性平台可能使用特权代理（如DCGM导出器）收集指标，并以受控方式向用户暴露，而不是让用户在共享硬件上运行任意性能分析器。许多供应商正致力于使容器中的性能分析更易访问（随着NVIDIA Docker集成，NVIDIA对NVTX和某些CUPTI功能在容器中的支持有所改善）。

对于构建新工具的人来说，明智的做法是了解这些限制，并可能设计工具使其能够：

在没有特殊权限的情况下以减少功能模式运行（例如，只使用用户空间计时，而非硬件计数器），或者
部署为多个用户可以安全使用的特权服务（可能通过经过身份验证的API）。

限制、注意事项和未来方向

尽管功能强大，当前的GPU/DPU性能分析工具仍有限制。了解这些有助于识别未来工作的领域：

一个工具中有限的供应商支持： 没有单一工具能很好地覆盖所有加速器。Nsight仅限NVIDIA；Omnitrace目前仅限AMD；Intel的工具仅限Intel。像HPCToolkit和TAU这样的工具尝试多平台，但即使它们也需要每个供应商不同的后端，可能不会暴露每个供应商特定的指标。真正统一的性能分析标准仍然缺乏。OneAPI朝着CPU/GPU方向迈出了一步（通过Level Zero和插桩接口），但采用仍在进行中。未来的工作可以开发通用追踪格式或API，使一个工具（或其插件）能够无缝处理来自多个供应商的GPU。
深度与开销权衡： 如前所述，捕获所有内容（长时间运行的完整追踪）可能不可行。因此，工具要么限制持续时间、频率，要么限制详细程度。用户需要注意，例如，Nsight Systems在缓冲区溢出时可能会丢弃事件，或者HPCToolkit可能不会记录每个微小的内核（如果它们发生得太频繁，通常会过滤掉非常短的内核以避免数据淹没）。这意味着性能分析可能不是100%准确。需要注意的是，启用性能分析可能会改变性能；例如，分析行为可能会使一些GPU操作序列化或改变时间，因此用户必须将结果视为近似值。未来的性能分析器可能利用更多硬件支持来减少开销（例如，硬件追踪缓冲区，如CPU中的Intel PT；也许有一天GPU会有一个性能分析器可以收获的硬件追踪）。
DPU和异构工作负载性能分析： 分析DPU（数据处理单元）和紧密集成的系统（CPU和GPU共享内存的APU）是一个不断发展的领域。对于DPU，由于它们结合了网络、存储和计算卸载，性能分析器需要捕获网络事件（数据包或协议处理）以及DPU的Arm核心上的CPU事件。今天，可能在DPU上使用标准Linux工具（如用于CPU的perf和可能通过ethtool的专用NIC计数器）。目前还没有广泛使用的专用DPU性能分析器，可以显示，例如，"NIC中的这个流导致了DPU上的这个CPU处理"。未来的工作可以整合NIC遥测与CPU性能分析——例如，使用类似DPDK或DOCA库在DPU上检测数据包处理。
对于APU（一个芯片上的CPU+GPU，共享内存），一个限制是许多性能分析工具仍然分别对待CPU和GPU。但APU可以启用新型性能分析：由于它们共享内存，可以系统范围地追踪内存访问。AMD的HSA旨在统一CPU/GPU内存空间，实际上，HSA的工具可以测量CPU对缓冲区的访问和GPU访问是否干扰等。我们还没有看到主流工具显示APU上CPU+GPU的组合缓存配置。未来研究可能允许追踪APU中CPU和GPU之间的缓存一致性事件或内存流量，以查看争用（例如，如果CPU和GPU争夺内存带宽，如何在性能分析中检测？今天可能会看到CPU和GPU都显示高内存使用率，但手动关联它们）。
另一个APU特定的挑战是时间线对齐——在APU上，CPU和GPU时钟可能更紧密耦合，甚至可能是相同的振荡器。这可能简化相关性（不考虑PCIe延迟），但工具尚未利用任何特殊的APU功能。
可扩展性： 随着HPC迈向百亿亿次计算，同时分析数百或数千个GPU变得困难。HPCToolkit和TAU通过专注于采样和分析（减少数据）而非巨大追踪来解决这个问题。但即使每秒从100,000个节点收集一个样本也是大量数据。未来方向包括在线分析（性能分析器即时分析和丢弃数据，只保留摘要或异常）和流式方法，只将有趣的事件发送到中央收集器。也许AI技术将应用于性能数据，智能决定记录什么。
自动化和洞察生成： 目前，性能分析器呈现数据，由用户推断洞察（工具提供一些提示）。Nsight Compute有一个规则引擎来提示常见问题（"内存受限，尝试增加并发性"等），但大多数工具不会自动确定根本原因。未来方向是更多自动性能诊断。例如，关联模式："每当GPU空闲时，CPU线程X在互斥锁中——可能是同步瓶颈。"一些研究工具或原型功能致力于这种专家系统。TAU与调优框架集成，可能可以自动对可疑热点应用插桩。我们可能会看到更多性能异常检测内置于性能分析器中（特别是对于始终开启的监控，标记某次迭代比平常慢并给出可能原因）。
能源和效率指标： 随着加速器的增多，功率和能源性能分析变得重要。DCGM提供功率使用情况，但很少有工具将其与性能集成。许多性能分析器的一个限制是它们不容易将能源使用归因到每个内核。未来，工具可以使用功率遥测来识别能源浪费的阶段（GPU功率高但利用率低）。AMD和Intel通过RAPL或ROCm-smi支持一些功率指标；将这些集成到性能分析中正在进行。
新型加速器： 除了DPU之外，还有TPU（谷歌的）、FPGA、定制ASIC。每种类型往往有自己的工具（谷歌TPU有集成在TensorBoard中的性能分析器）。社区面临的挑战是通过标准接口将性能分析推广到任何加速器。朝这个方向发展的一个项目是开源Kokkos Tools，用于节点级性能分析，它抽象了设备。它允许通过通用API分析任何后端（CUDA、HIP等）。像HPCToolkit和TAU这样的工具参与这些计划，为新设备做好准备（如果加速器供应商提供类似OMPT或CUPTI的API，这些工具可以接入）。
需要逆向工程： 开源工具经常需要追赶未记录的功能。例如，在NVIDIA提供官方支持之前，HPCToolkit开发者必须弄清楚如何通过解析CUBIN中的ELF符号将GPU PC映射到源代码，以及如何通过CUPTI的不断发展的API收集GPU指标。这意味着当新硬件到来时，开源工具落后于供应商工具（供应商可能不会立即更新CUPTI等）。未来理想的情况是供应商开放更多驱动程序的插桩钩子或标准化它们。同时，可以预期，对于前沿功能（如新GPU内存类型或新DPU卸载引擎），早期性能分析可能涉及创造性黑客技术或有限信息，直到官方支持赶上。

总之，这一领域的未来工作可能会专注于统一和简化异构性能分析的工具链，进一步减少开销（可能通过硬件支持），将性能分析与监控集成（在生产中持续捕获问题），以及更智能的分析来帮助开发者更快地找出问题。景观正在向开放标准（如用于追踪的Perfetto，也许有一天会出现类似PAPI的GPU标准）移动，这将使构建插入的自定义性能分析器更加容易。当前工具实现的洞察——它们的成功和痛点——将指导这些发展。随着硬件继续发展（芯片GPU、统一内存层次结构等），性能分析工具需要适应，但基本技术（插桩与采样、硬件计数器使用、多流关联）将仍然是任何性能分析工具包的基石。