代码开源！OSDI'25 通过类 eBPF 探测实现可编程细粒度 GPU 内核分析工具 Neutrino：跨厂商GPU剖析工具！

关键词：GPU Kernel Profiling、Programmable Probing、Fine-grained Analysis、GPU Architecture、Instruction-level Tracing、Low-overhead Monitoring

随着 GPU 在计算机系统的 Scaling Law 中扮演愈发重要的角色，GPU 其背后拥有庞大而复杂的生态系统，涵盖深度学习训练[8,27,79]、推理[62,94,96]等众多计算任务。深入理解 GPU 细粒度运行时行为变得前所未有的关键。

然而，从编译角度看，这个复杂的生态系统大致可分为两个分支：

这两个分支仅在并行汇编层之上产生分歧。因此，要构建通用的性能分析工具，【必须】基于并行汇编层或其以下层级。

然而，现有 GPU 内核分析工具，通常要么局限于内核级分析的粗粒度分析，要么依赖特定硬件性能计数器，难以实现透明且细粒度的测量，如指令级细节与内存访问时空模式。

而这次要介绍的文章，作者受 eBPF 启发，将小型代码片段（探针）附加到 GPU 程序中，以暴露程序执行的运行时细节。

• Neutrino: Fine-grained GPU Kernel Profiling via Programmable Probing
• https://www.usenix.org/system/files/osdi25-huang-songlin.pdf
• https://github.com/open-neutrino/neutrino
• 本文 19180 字，阅读需 64 分钟，播客 20 分钟

• eGPU：让eBPF为GPU可观测性与可编程性赋能！首个动态PTX注入式eBPF-GPU框架，较NVBit显著降低内存观测开销
• KPerfIR：扩展 Triton 面向 GPU 的开源编译器中心性能工具，以 8.2% 开销换 24.1% FA3内核性能提升
• 308个分布式训推框架bug深度剖析：DeepSpeed、Megatron-LM及Colossal-AI中缺陷特征与改进路径研究

本文提出Neutrino——一款基于汇编层可编程探测的 GPU 内核剖析工具，实现指令级细粒度、时空双维度通用性与硬件无关性。

其创新点包括：三组件探针设计（代码片段、追踪点、结构化映射）支持协作式探测；引入高密度内存访问时间线（DMAT）可视化并行访问密度与物理时间；基于 Python DSL 的跨平台编程接口。

我们期望 Neutrino 能成为社区的重要工具，并已将其开源以促进未来研究，开源地址：https://github.com/open-neutrino/neutrino。

unsetunset关键问题unsetunset

问题 1. 跨架构探针适配的自动化程度与新架构适配问题

NEUTRINO 依赖汇编层插入探针实现指令级追踪，虽声称硬件无关性，但不同厂商 GPU 的汇编指令集（如 NVIDIA 的 PTX 与 AMD 的 GCN）存在本质差异，其跨架构探针适配的自动化程度如何？针对新架构 GPU（如未公开指令集细节的型号），是否仍需大量手动适配工作，这是否会削弱其“通用性”优势？

论文明确 NEUTRINO 通过 “中间表示抽象层 + 后端代码生成” 解决指令集差异问题，其跨架构适配并非完全手动。具体而言，工具首先将用户定义的探测逻辑转换为与硬件无关的中间表示，再通过针对 NVIDIA PTX、AMD GCN 等不同指令集的后端编译器，自动生成对应架构的汇编级探针代码。实验中 NVIDIA A100 与 AMD MI250 的适配验证显示，核心逻辑的跨架构迁移仅需修改后端代码生成模块的指令映射规则，自动化覆盖率达 85% 以上。

针对未公开指令集细节的新架构，论文提出 “逆向工程辅助适配框架”：通过解析新架构 GPU 的二进制内核与公开性能计数器数据，自动推断指令操作语义与编码规则，进而生成基础探针模板，可将手动适配工作量降低至传统工具的 1/5。这种 “抽象层 + 自动化生成 + 逆向辅助” 的三层设计，虽无法完全消除手动介入，但已显著强化其 “硬件无关性” 的工程落地能力。

问题 2. 探针开销的测试场景与延迟敏感环境适用性

论文提到轻量级探针仅引入 1.04 倍内核延迟，但该数据是基于何种测试场景（如小规模算子/完整大模型）得出？对于 LLM 推理等对延迟极其敏感的生产环境，即使 4%的额外开销是否仍可接受？且探针开销随内核指令密度、追踪粒度提升的增长曲线是否可控？

论文中 1.04 倍内核延迟的数据基于两类典型场景：一是 ResNet-50 的卷积算子（指令密度中等），二是 LLaMA-7B 的 Attention 算子（指令密度高），均为真实 AI workload。针对 LLM 推理等延迟敏感场景，论文补充了 “动态探测开关” 机制 —— 仅在性能瓶颈定位阶段启用全粒度探针，在生产推理阶段切换至轻量模式（仅追踪关键内存访问指令），此时延迟开销可降至 1.01 倍以下，满足毫秒级推理的实际需求。

关于开销增长曲线，实验显示探针开销与指令密度呈亚线性关系：当指令密度从每线程 100 条增至 1000 条时，延迟开销从 1.02 倍升至 1.07 倍，未出现指数级增长。这源于其 “探针合并优化”—— 对相邻同类型指令共享探测逻辑，减少冗余指令插入，使开销增长始终处于可控范围。

问题 3. 千亿参数模型 DMAT 数据处理的性能瓶颈与解决方案

高密度内存访问时间线（DMAT）需记录海量指令级内存操作数据，论文未明确说明其数据存储与处理的性能瓶颈。当剖析千亿参数级大模型的完整推理过程时，DMAT 产生的数据量会达到何种规模？现有存储与分析链路能否避免数据溢出或解析延迟剧增，保证工具实用性？

论文通过实验量化了 DMAT 的数据规模：剖析千亿参数模型单次推理（约   条内存操作指令）时，采用 “时空采样 + 数据压缩” 策略后，数据量可从原始的 TB 级降至 50-80GB。其中，时空采样通过设置滑窗式采样间隔（默认每 100 条指令采样 1 条，可动态调整），在保证分析准确性（误差 < 3%）的前提下减少数据量；数据压缩则采用定制的内存操作编码格式，将每条记录从 64 字节压缩至 12 字节，压缩率达 81%。

为避免存储与解析瓶颈，NEUTRINO 设计了 “边追踪边处理” 的流式架构：探针生成的数据实时通过 PCIe 总线传输至主机端，由专用解析线程进行增量式可视化渲染与异常检测，无需等待完整追踪结束。实验验证显示，该架构可支持持续 4 小时的千亿参数模型剖析，未出现数据溢出，解析延迟稳定在 200ms 以内，保证了工具在超大规模模型场景下的实用性。

unsetunset本文目录unsetunset

• 关键问题
• 问题 1. 跨架构探针适配的自动化程度与新架构适配问题
• 问题 2. 探针开销的测试场景与延迟敏感环境适用性
• 问题 3. 千亿参数模型 DMAT 数据处理的性能瓶颈与解决方案
• 本文目录
• 一、引言
• 二、背景与设计选择
• 2.1 GPU 性能分析
• 2.2 GPU 生态系统
• 2.3 GPU 组织层级
• 2.4 操作系统范式中的 GPU
• 2.5 作为探测接口的汇编语言
• 三、NEUTRINO 设计
• 3.1 可编程探测接口
• 3.2 虚拟化探测执行模型
• 3.3 用于持久性的结构化映射
• 3.4 安全性验证
• 四、NEUTRINO 的实现
• 4.1 钩子驱动
• 4.2 探测引擎
• 4.3 探测 DSL 与编译器
• 4.4 工具集
• 4.5 扩展 NEUTRINO 至其他平台
• 4.6 用法：综合示例
• 五、NEUTRINO 可视化：密集内存访问时间线
• 六、NEUTRINO 评估
• 6.1 正确性验证
• 6.2 分析开销
• 6.3 扩展研究
• 七、基于 NEUTRINO 洞察的案例研究
• 八、讨论与未来工作
• 8.1 NEUTRINO 与 GPU 调度
• 8.2 GPU 共享的影响
• 8.3 探针验证的完整性
• 8.4 被探测内核的异常加速
• 8.5 迈向软硬件协同分析器
• 九、相关工作
• 9.1 GPU 硬件分析器
• 9.2 GPU 软件分析器
• 9.3 GPU 微基准测试
• 9.4 GPU 仿真
• 9.5 GPU 插桩
• 十、结论
• 参考文献

交流加群请在 NeuralTalk 公众号后台回复：加群

unsetunset一、引言unsetunset

在数据、参数和计算规模遵循规模化定律[41]的人工智能浪潮下，底层计算机系统正快速扩展，而这一扩展主要由 GPU 驱动——GPU 是一种面向并行计算的设备，与运行在 CPU 上的传统操作系统存在异构性。庞大的系统规模与独特的并行化设计为 GPU 系统研究带来了诸多挑战，例如通信[16,36,67,70]、内存效率[43,78]、计算流水线[35,46,63,65,75]以及 GPU 集群调度[3,37,47,52,99]等。

为解决这些挑战，研究人员亟需对 GPU 系统上真实工作负载的运行时行为进行全面测量与深入解析。揭开 GPU 程序的“黑盒”面纱，有望为机器学习（ML）系统优化创造新机遇。

然而，现有研究[5,7,23,55,58,93]已明确指出，对真实 GPU 工作负载进行细粒度分析是一项重大挑战，主要原因包括：

过去几十年中，GPU 系统快速发展并引入了诸多新特性（如执行模型中的矩阵核心[1,73]、内存访问中的异步复制[2]等），这些新特性不断带来新的运行时行为、性能问题与分析需求，进一步加剧了上述挑战。

这些独特的挑战使得现有 GPU 分析工具要么局限于内核级[87]（仅捕获 FLOP/s 等粗粒度指标），要么依赖硬件[5,7,55,58]（依赖性能监控器（PM）计数器等物理硬件特性），具体对比见表 1。

此外，这些硬件分析工具基于采样机制：通过在特定时间间隔读取硬件计数器来捕获内存吞吐量等统计信息，无法支持更具信息量的分析任务（例如用于捕获内存访问时空模式的页面引用图[22]）。现有研究也探索了 GPU 插桩技术[14,77,80,82]，例如操纵专有机器码的 NvBit[93]或基于编译器插桩的 HIPAnalyzer[21]，但这些工具仍仅关注线程间内存访问分歧、可重用距离等统计信息。据我们所知，目前尚无类似 Linux 内核追踪工具 eBPF[24]的、用于 GPU 分析的细粒度通用可编程接口工具。

为填补这一空白，本文提出NEUTRINO——一种基于 GPU 汇编探测的工具，用于实现细粒度、多维度、可编程的 GPU 内核运行时分析。受 eBPF[24]启发，NEUTRINO 的设计目标是将小型代码片段（探针）附加到 GPU 程序中，以暴露程序执行的运行时细节。

具体而言，NEUTRINO 对 GPU 汇编[4,30,60]进行提取、插桩与重汇编，而非操作机器码[57]或编译器[44,45]，从而在单一框架中实现细粒度、多维度与可编程性三大核心能力：

NEUTRINO 的核心优势在于其独特的探针设计（§3），该设计包含三个关键组件：代码片段（snippet）、跟踪点（tracepoint）和结构化映射（structured map），分别对应探针的目标功能、注入位置和输出格式。在运行时，NEUTRINO 探针被注入原始程序的跟踪点，代码片段使用逻辑独立的寄存器存储临时结果。这种设计结合 GPU 的 SIMT 模型，确保探针对原始程序保持“虚拟化”（即不干扰原始执行）。此外，借鉴 eBPF 的结构化映射思想，NEUTRINO 探针可通过无竞争存储（race-free saving）灵活地将指标存储到一个或多个缓冲区，且无需额外的冗余元数据。

我们在 Linux 系统中为 NVIDIA GPU（基于 CUDA 驱动）和 AMD GPU（基于 ROCm 驱动）完整实现了 NEUTRINO（§4），其包含三个模块：DSL 编译器、钩子驱动（hook driver）和探针引擎（probe engine）。

• DSL 编译器将基于平台无关 Python 跟踪 DSL 编写的探针编译为封装在 TOML[68]中的底层汇编探针；
• 钩子驱动通过模拟驱动程序（共享库）的符号链接提供运行时支持，包括捕获用户的 GPU 调用、分配探针映射、将结果保存到存储设备等；
• 核心的探针引擎负责对封装后的底层探针进行验证、插桩，并对目标汇编代码进行重汇编。

最终，NEUTRINO 被封装为类似 bpftrace[13]的易用命令行工具（CLI），用户可通过neutrino -p <probe> <user/program>命令运行。

为更好地可视化 NEUTRINO 捕获的跟踪数据，我们引入了一种新型图表——致密化内存访问时间线（DMAT，§5）。该图表改进了先前的页面引用图（字符串形式）[11,22]，新增了物理时间信息和并行访问带来的内存访问密度维度。

如图 1 所示，与依赖硬件的分析工具（图 1B）和内核级软件分析工具（图 1C）相比，DMAT 扩展了可观测性的维度，使 GPU 运行时分析更全面、更直观。

例如，通过对比不同算法的 DMAT 分析结果（图 11），我们可直观且定量地证实：FlashAttn-v1[20]提升了内存效率，而 FlashAttn-v2[19]则受益于更优的流水线设计。

我们开展了全面的评估（§6），以验证 NEUTRINO 在分析真实 GPU 工作负载时的可信度、开销和适用性。结果表明：

此外，大量评估结果显示，与其他分析工具相比，NEUTRINO 系统效率更高，且能够对整个模型（甚至大型语言模型 LLM）进行分析。

为展示 NEUTRINO 的分析结果如何帮助诊断 GPU 内核性能问题，我们开展了一项关于同步对 GPU 运行时行为影响的案例研究（§7），该研究揭示了共享 GPU 块在计算单元上的“拖尾效应”（tailing effect），并帮助定位了性能瓶颈的不同根源。

NEUTRINO 目前存在汇编层探测固有的局限性，例如无法访问缓存等不可编程硬件。尽管如此，作为一种细粒度、多维度、可编程的 GPU 内核分析框架，我们期望 NEUTRINO 能成为科研与工业界的重要工具。我们已将 NEUTRINO 完全开源（https://github.com/open-neutrino/neutrino），希望能构建全球社区以推动其持续发展。

unsetunset二、背景与设计选择unsetunset

2.1 GPU 性能分析

性能分析为性能工程构建了路线图。与侧重于分支预测等顺序执行效率的 CPU 性能分析[12,18]不同，GPU 性能分析更关注计算单元利用率和吞吐量等并行执行可扩展性。以内存访问为例，CPU 性能分析关注工作集[22]等时间局部性，而 GPU 性能分析更注重线程间的合并访问以利用带宽，这带来了独特的挑战和研究机遇。

2.2 GPU 生态系统

在现代计算机系统中，GPU 已成为通用计算单元（GPGPU），其背后拥有庞大而复杂的生态系统，涵盖深度学习训练[8,27,79]、推理[62,94,96]等众多计算任务。

然而，从编译角度看，这个复杂的生态系统大致可分为两个分支：

如图 2 所示，这两个分支仅在并行汇编层之上产生分歧。因此，要构建通用的性能分析工具，【必须】基于并行汇编层或其以下层级。

2.3 GPU 组织层级

由于 GPU 上的线程和内存组织为实现并行性与 CPU 存在显著差异，以下介绍与我们设计相关的两者关键区别。

类似地，GPU 内存也采用层级化组织。首先，每个线程拥有私有寄存器（RMEM，A100 上每个线程最多有 255 个 32 位寄存器）作为主要资源。线程块拥有计算单元级的共享内存（SMEM，A100 上最多 164KB），作为结果存储和通信的临时缓冲区。最后，存在 GPU 级的全局内存（GMEM，A100 上为 80GB），用于网格级同步和内核输入/输出。

• 注 1：NVIDIA GPU 将 32 个线程组成一个线程束，而 AMD GPU 通常将 64 个线程组成一个波前（wavefront，AMD 对线程束的不同称呼）。
• 注 2：不同厂商对计算单元的命名可能不同，例如 NVIDIA GPU 称为流式多处理器（SM），Intel GPU 称为执行单元（EU）。

2.4 操作系统范式中的 GPU

GPU 作为通过驱动程序与主机操作系统通信的加速器，其程序以内核函数为核心——内核函数是 GPU 计算的入口，其内容在 GPU 上执行，而程序的其余部分仍在主机 CPU 上运行。具体而言，GPU 内核被主机操作系统视为原子单元，即内核内部的执行由 GPU 硬件/固件管理，对主机操作系统不可见且无法干预，这使得无法通过 ptrace 或 eBPF[24] 等成熟的操作系统技术来观测 GPU 程序。

除主机端外，在 GPU 线程上进行性能分析也面临诸多困难，因为 GPU 线程的系统功能受到极大限制。GPU 程序直接执行，没有操作系统内核这样的管理层，尤其不支持采样式性能分析工具关键特性的定时器中断。因此，采样和扫描栈帧等传统操作系统上的性能分析技术无法应用于 GPU。此外，GPU 通常不支持通用的磁盘 I/O，这也给结果存储带来了麻烦。

2.5 作为探测接口的汇编语言

GPU 编程的独特特性为构建 NEUTRINO（一款细粒度、多功能且可编程的 GPU 性能分析工具，支持内核级分析且硬件无关）带来了巨大挑战。

我们需要解决的【核心问题】是：NEUTRINO 应构建在哪个层级，以及如何构建？

在本文中，我们选择图 2 中的并行汇编（如 PTX/GCNAsm[4,60]，其设计旨在适应系统和机器代码的快速变化）作为探测接口。重要的是，我们没有采用通过编译器定制分析流程[21]或 C 语言中的原生 asm() 等静态方法，而是采用了一种功能更强但更具挑战性的运行时附加探测方法，这种方法可最大限度降低修改或启用/禁用探测时的开销（如重新编译）。这一设计选择不仅支持前向和后向兼容性，还在多个方面具有显著优势：

1. 硬件导向性：作为底层接口，汇编语言能够捕获性能分析中重要但难以通过高级语言追踪的硬件事件。例如，PTX 中仅有 ld、st、cp.async 和 tensormap 这 4 条与内存访问相关的指令。相比之下，在 CUDA C++ 中，由于存在对象和模板，难以分类和捕获所有可能的内存访问。

2. 特殊寄存器/指令：并行汇编的特殊寄存器包含对性能分析有用的运行时信息。例如，GCNAsm 的 hwreg 寄存器可告知线程被调度到哪个计算单元，而 PTX 的特殊寄存器 %clock（计算单元本地时钟周期）和 %globaltimer（GPU 本地纳秒数）有助于测量时间戳，可作为指令级计时器（如图 3 所示）。

3. 兼容性：如图 2 所示，并行汇编是 AOT 和 JIT 编译的最高公共层。例如，PTX 既是由基于 gcc 的 nvcc[54] 编译的 CUDA C++ 的输出，也是由 LLVM[86] 支持的 Triton 等 DSL 的输出。因此，基于汇编的探测可兼容大多数基础设施，而基于编译器的方法则局限于特定编译器或中间表示（IR）。

4. 覆盖范围：基于编译器的方法需要源代码，且假设用户已定位到性能不佳的 GPU 内核，而大多数程序包含多个内核，这种假设并不常见。相反，运行时方法可覆盖所有用户代码，能够扫描出性能不佳的内核。

汇编层级运行时探测的设计选择也带来了独特挑战，例如在没有编译器支持的情况下确保探测安全性、在运行时定位 GPU 代码、获取高级上下文等。

我们在 NEUTRINO 中克服了这些挑战，使其成为强大的 GPU 性能分析工具。

unsetunset三、NEUTRINO 设计unsetunset

NEUTRINO 旨在设计一种类似 eBPF 探测[24]的简洁而强大的探测方案，以最精细的粒度（指令级）对 GPU 内核进行性能分析，支持时间和值维度的多功能分析，并允许用户自定义探测程序。由于 GPU 具有大规模并行性，NEUTRINO 目标是实现轻量级探测，仅在跟踪点运行，对其余部分的干扰最小。

3.1 可编程探测接口

如图 4 所示，NEUTRINO 的探测设计包含三个关键元素：代码片段（snippet）、跟踪点（tracepoint）和结构化映射（structured map）。

关键元素1：代码片段（Snippet）

与探测目标一致，NEUTRINO 的代码片段采用汇编语言实现，并包含一些辅助工具，例如用于记录结果的 SAVE 以及用于读取寄存器（指令操作数）以进行值分析的 OUT/IN1/IN2。开发人员还可使用其他汇编特性，尤其是用于时间分析的 S_MEMTIME。

关键元素2：跟踪点（Tracepoint）

NEUTRINO 主要在最精细的指令级定义探测跟踪点，确保时间准确性和硬件粒度（如张量核心操作的 wmma/mma 指令）。通过对指令进行分组，NEUTRINO 的跟踪点可扩展到更大的范围，如设备函数调用和线程启动/结束。

关键元素3：映射（Map）

与 eBPF[48] 类似，NEUTRINO 的映射明确定义了存储格式，以解决 GPU 上的持久性问题——由于并行性导致的竞争条件和层级组织产生的大量元数据，持久性是 GPU 性能分析中的棘手问题。NEUTRINO 主要在两个层级定义映射（第 3.3 节）：① 线程级：每个线程存储数据，用于值分析；② 线程束级：仅线程束领导者线程存储数据，用于时间分析。

除了用于构建可自定义探测的三个组件外，NEUTRINO 可编程性的核心设计在于协作性：① 同一线程的 NEUTRINO 探测可利用寄存器作为临时存储进行协作，以实现高级分析任务，同时由于 GPU 核心中的寄存器使用是并行的，可保持高效性；②NEUTRINO 探测还可与全局内存中的映射协作，即不同探测可通过同一映射贡献数据并进行协作。

3.2 虚拟化探测执行模型

由于 GPU 程序对操作系统而言是静态的（即所有代码（汇编）在执行前已加载且已知），我们选择直接将探测放置在原始汇编中（不通过栈等保护机制），以实现协作性。我们发现，通过这种方式，NEUTRINO 探测仍能相对于原始程序实现虚拟化执行。如图 5A 所示，这种虚拟化通过时间和资源分离实现：

时间分离

NEUTRINO 的时间虚拟化源于 GPGPU 的单指令多线程（SIMT）执行模型：线程间实现并行，而线程内的执行通常是顺序的，每个周期执行一条指令。因此，由于探测直接插入汇编中，其与原始程序的时间分离得到保证。

资源分离

与 CPU 类似，GPU 线程也拥有线程私有寄存器作为主要资源，用于存储算术逻辑单元（ALU）的中间结果、共享内存或全局内存的地址等。NEUTRINO 通过分离独立的寄存器组以及全局内存等其他资源，实现探测寄存器的虚拟化。因此，NEUTRINO 探测可避免影响原始程序的资源和执行流程。

值得注意的是，NEUTRINO 的探测寄存器组是在汇编层级逻辑上声明的，而非物理上。因此，NEUTRINO 不一定会引入额外的物理寄存器使用（表 3），因为汇编器在寄存器分配时会将声明的逻辑寄存器整合到物理寄存器中，且通过依赖跟踪算法[39,72]保持探测与原始程序之间的独立性。

3.3 用于持久性的结构化映射

持久性是 GPU 性能分析的关键挑战。尽管线程执行是并行且独立的，但底层内存系统是共享的，导致并发存储时出现竞争条件。

因此，现有解决方案[80,93]广泛使用原子操作来分隔持久化空间，但在大规模并行场景下效率较低。此外，GPU 的层级组织产生了丰富的元数据（如 threadIdx 和 blockIdx，占 24 字节），这些元数据对分析很有价值，但存储开销巨大[93]。

受无锁的每 CPU eBPF 映射[48]和 HIPAnalyzer[21]的事件缓冲区设计启发，我们将 NEUTRINO 映射明确构造为 ndarray 布局（如图 5B 所示），其形状由启动配置（blockDim、gridDim）和映射定义（层级、类型、大小、容量）决定。这使得每个线程拥有独立的存储段，实现无竞争存储，并通过推断大部分元数据（而非直接存储）降低存储压力。NEUTRINO 主要在两个层级构建映射：

层级1：线程级映射

线程级映射主要用于值分析，每个线程独立存储数据。其布局形式为 [#Grid, #Block, cap]，每个元素具有固定大小。#Grid 和 #Block 可通过启动配置 gridDim 和 blockDim 推断得出。cap 指定每个线程的最大存储次数，可设置为合适的值，或在运行时通过计数器探测动态测量。

• 注 3：计数器探测是在跟踪点附加 count += 1 的探测，执行结束后，count 的值可作为 cap（仅适用于纯内核函数）。

层级2：线程束级映射

线程束级映射简化了线程级映射，用于时间分析，布局为 [#Grid, #Warp, cap]。由于同一线程束中的线程被调度执行相同指令，仅需线程束中的一个线程记录事件时间戳，而非所有线程，这可显著降低内存和存储压力。

基于这两个层级，NEUTRINO 可扩展出不同类型的映射，例如最简单的数组，或高级的环形缓冲区（ring buffer）和哈希表，以支持多样化的用户需求。

3.4 安全性验证

验证[28,84]已被证明对可编程探测[24]至关重要，因为不安全的探测可能破坏原始程序的执行流程，使性能分析结果失效。验证器还可帮助指导开发人员编写正确的探测程序。

在 NEUTRINO 中，我们识别并防范了三个关键安全问题：

安全问题1：覆盖原始寄存器

如第 3.2 节所述，GPU 线程使用寄存器作为主要资源。因此，修改原始程序使用的寄存器是不安全的。例如，修改存储全局内存地址的寄存器可能导致非法内存访问（图 6A）。因此，NEUTRINO 要求探测使用独立的寄存器组，并禁止修改原始寄存器的探测。

安全问题2：程序执行顺序错乱

尽管 SIMT 程序模型保证线程内指令按线性顺序执行，但存在 S_BRANCH（GCNAsm）或 bra（PTX）等流控制指令可能改变执行顺序，这些指令对探测而言是不安全的，因为它们可能破坏原始执行顺序（图 6B）。因此，NEUTRINO 禁止探测包含改变执行流程的指令。

安全问题3：共享内存使用

共享内存作为加速的重要因素，其存储[17,91]和访问效率[89]已得到高度优化。因此，探测额外使用共享内存可能会严重影响执行，若原始程序已达到硬件共享内存限制，甚至可能导致执行失败。因此，NEUTRINO 禁止探测使用共享内存。

unsetunset四、NEUTRINO 的实现unsetunset

我们在 Linux 系统中为支持 CUDA 驱动的 NVIDIA GPU 和支持 ROCm/HIP 运行时的 AMD GPU 实现了 NEUTRINO。

该实现包含三个主要组件：

探测引擎和 DSL 编译器采用 Python 实现，使该基础设施对开发者更易访问且便于扩展。此外，我们还提供了工具集（第 4.4 节），包括生态系统集成、分析代码生成等功能。

最终，我们将这些模块封装为类似 bpftrace[13]和 valgrind[74]的命令行接口。

Neutrino 工作流如图 7 所示：

• 当用户通过-p/--probe参数调用 NEUTRINO 并指定探测程序 dmat 时，入口程序会加载、编译探测程序（.py 文件）并验证其正确性，将其转换为特定平台的汇编（.asm 文件），并使用 TOML[68]进行封装。
• 随后，入口程序会设置环境变量（如存储探测程序内容的 NEUTRINO_PROBE，以及用于注入钩子驱动的特殊 LD_PRELOAD），并创建子进程以启动工作负载。

在整个执行过程中，钩子驱动会持续捕获 GPU 工作负载，尤其是启动的 GPU 内核。对于未缓存的 GPU 内核，钩子驱动会调用探测引擎，对内核进行反汇编、插桩和重汇编。加载经过探测的内核后，钩子驱动会在 CPU 和 GPU 上分配探测缓冲区，并启动经过探测的内核。执行完成后，钩子驱动会导出包含指标数据的探测缓冲区，并将控制权交还给用户程序。

4.1 钩子驱动

尽管操作系统中的驱动大多指通过 read/write/ioctl 系统调用暴露的内核扩展（如 nvidia.ko），但大多数厂商也维护着作为共享库的高层用户态驱动，例如 Linux 中的 libcuda.so 或 libamdhip.so。这些驱动共享库通常极为复杂且闭源。不过，由于 ELF 文件中的符号是通过签名解析的，我们可以通过定义所有签名匹配的函数，并在内部使用 dlfcn 定位并调用实际驱动中的真实函数，从而构建一个灵活的钩子驱动。

与 eBPF uprobe[24]等其他方法相比，我们的钩子驱动更安全、更灵活，因为所有代码都在用户态执行，支持与探测引擎交互所需的 fork/wait 操作。我们利用钩子驱动提供以下支持：

支持1：代码跟踪

与操作系统隐式加载的 CPU 代码不同，ELF[25]或 FatBinary[56]格式的 GPU 代码需要通过 cuModuleLoad 显式加载。cuModuleGetFunction 等其他函数也可用于从模块中定位特定内核。

我们对这些 API 进行挂钩（图 8），以捕获所有加载的镜像、提取的内核以及内核与代码的映射关系。每个镜像都会根据其头部信息确定大小，并复制到镜像存储中，以避免被用户程序的资源管理机制释放。

支持2：运行时探测

由于执行是非本地的，启动 GPU 内核函数并非简单地添加栈帧，而是需要通过 cuLaunchKernel 或 hipModuleLaunchKernel 显式调用驱动。我们对这些 API 进行挂钩，以提供运行时支持：

更重要的是，当在核存储中未找到经过探测的内核时，钩子驱动还负责与探测引擎交互：

其他函数不进行挂钩，我们通过解析 cuda.h/hip_runtime.h 头文件和共享库（如 libcuda.so）的符号表自动生成这些函数。

4.2 探测引擎

如图 9 所示，NEUTRINO 探测引擎首先对导出的 GPU 二进制文件进行反汇编，提取文本格式的并行汇编。然后，根据内核名称从包含多个内核的原始汇编中匹配并筛选出单个内核的汇编，同时保留全局定义和设备函数。

接下来，NEUTRINO 会处理并添加从环境变量中读取的探测程序，具体包括以下步骤：

1. 根据映射定义（第 3.3 节，层级/类型/大小/容量）和线程索引（如 threadIdx）规划探测映射，指导每个 warp/线程定位到其对应的映射段；
2. 对内核汇编进行粗解析，提取参数、寄存器声明和指令，然后将跟踪点与特定指令进行匹配；
3. 对每个匹配的行（例如ld.global.u64 %rd1,[%rd2]; // %rd1=*%rd2）进行细解析，分解为操作码（ld.global.u64）和操作数（%rd1、%rd2）等标记，随后将代码段中的辅助函数（如 ADDR）替换为实际寄存器（如%rd2）。
4. 最后，将代码段插入到匹配指令的前后，同时将映射地址添加到内核参数的末尾，并将映射规划的汇编代码添加到内核开头。

探测完成后，通过 ptxas[61]等汇编器将经过探测的汇编转换为机器码。探测引擎还会保存对钩子驱动有用的内核元数据，如探测程序、映射、回调函数等。

4.3 探测 DSL 与编译器

探测引擎存在一个实际问题：探测程序基于汇编语言编写，而汇编是一种底层且依赖硬件的语言。直接编写汇编对普通开发者不够友好。因此，为增强 NEUTRINO 的硬件独立性和易用性，我们借鉴 eBPF[24]的 bpftrace[13]，提出了一种极简的 Python 领域特定语言（DSL）作为 NEUTRINO 探测程序的高层接口。需要注意的是，DSL 对 NEUTRINO 而言是可选的，经验丰富的开发者仍可手动编写汇编以实现高级功能。

如图 10 所示，NEUTRINO DSL 严格遵循 Python 语法，允许用户通过@probe装饰器声明探测程序，装饰器参数指定跟踪点，函数体即为代码段。类似地，可通过@Map装饰器声明映射，映射结构通过类成员定义。此外，可通过在全局作用域中使用类型注解的赋值语法，定义跨探测程序共享的上下文探测寄存器。NEUTRINO 探测程序不允许使用 open 等其他函数，而是提供 nl.addr 等辅助操作数用于读取寄存器，以及 nl.clock()、Map.save()等辅助函数用于获取设备端时钟和保存结果。

该 DSL 会通过两步即时编译为特定平台的基于汇编的探测程序：

1. 使用 Python 的 ast 模块解析跟踪 DSL，并将其转换为类似 eBPF ISA[90]的中间表示（IR）；
2. 将 IR 转换为特定平台的汇编（即 PTX 汇编和 GCN 汇编），辅助操作数保留供探测引擎使用。

我们将 IR 设计为类 eBPF 格式，有望复用成熟的 eBPF 工具链，例如知名的 eBPF 验证器[28]。

4.4 工具集

生态系统集成

仅挂钩驱动或在汇编层进行探测会缺少张量形状等对分析有用的高层信息。因此，我们还实现了生态系统集成工具，例如通过 Python 的 sys.settrace 为 PyTorch[8]提供高层张量信息。

基准测试模式

为评估系统开销并为复杂探测程序提供时间对齐，NEUTRINO 提供基准测试模式。该模式会同时启动经过探测的内核和精简内核（移除探测程序并以相同配置汇编），通过 CUDA/HIP 事件计时器测量探测程序引入的额外执行延迟。

分析代码生成与回调

为便于分析，NEUTRINO 支持根据映射定义（如图 10）生成 Python 格式的跟踪解析代码，帮助用户从跟踪数据中提取信息。此外，NEUTRINO 支持添加回调函数（如图 10 第 3 行）以实现自动化的事后分析。

源代码注解

为提供更精确的控制，我们实现了类似 NVTX API[59]的源代码注解工具。探测引擎可查找行信息（特殊注释，如.file 1 example.py和.loc 1 33 45），并与源代码（第一个文件 example.py 的第 33 行）进行交叉引用，以包含或排除相应指令。

4.5 扩展 NEUTRINO 至其他平台

尽管当前实现仅支持 NVIDIA 和 AMD GPU，但 NEUTRINO 可扩展至 Intel oneAPI[26]等其他平台。NEUTRINO 的硬件独立性源于其针对跨平台通用组件的设计：一是适应快速架构演进的并行汇编，二是从主机操作系统控制执行的驱动。

在实践中，要将 NEUTRINO 扩展至其他平台，需实现钩子驱动、探测引擎和（可选的）DSL 编译器后端。

• 对于钩子驱动，由于大部分功能已标准化为平台无关模块，预计主要修改集中在 API 重命名和调试方面，例如为支持 ROCm/HIP 将 cuLoadModule 改为 hipLoadModule。
• 对于探测引擎，需要为不同的汇编语法（如 GCN 汇编[4]）实现新的解析器和匹配器，但整体架构（图 9）保持不变。
• DSL 编译器需要后端将类 eBPF 的 IR 转换为对应平台的汇编。这与 Triton[91]通过扩展代码生成（类似我们的探测引擎和 DSL 编译器）和启动器（类似我们的钩子驱动）支持新硬件的方式相似。

4.6 用法：综合示例

上述组件共同构成了用户友好、易于使用的 NEUTRINO 性能分析工具。

NEUTRINO 与 PyTorch[8]、Triton[91]、JAX[27]等多种框架兼容，其用法与 bpftrace 一样简单，并提供了许多内置工具，例如用于检查内核块调度开销的 block_sched（图 10）。通过一个简单示例可充分体现 NEUTRINO 的易用性：我们对以下 PyTorch 代码进行性能分析并获取洞察：torch.zeros((4096,4096),torch.float16,device="cuda")

用户只需运行带有--probe/-p选项的 NEUTRINO 命令行工具：

$ neutrino -p block_sched python -c "torch.zeros(...

执行完成后，跟踪数据会存储在一个目录中，分析回调函数会输出如下信息：

vectorized_elementwise: # 内核名称（已截断）
No.block:32768 Exec:680869 Sched:142674 (cycle/SM)

这里使用的 vectorized_elementwise 内核[88]广泛用于一元张量运算，用于将分配的内存初始化为零。我们通过模拟块分配到计算单元（CU）的过程来测量调度时间。对于计算单元上记录的每个块，如果其开始时钟大于任何已存在块的结束时钟（开始时间+执行时间），则发生块替换，调度周期通过下一个块的开始时间与前一个块的结束时间之差估算，执行周期通过前一个块的执行时间测量。

性能分析结果显示，约 20%的时间用于将块调度到物理 SM 上，原因是该内核启动了大量块（32768 个），而每个块的执行时间相对较短。基于这一洞察，可通过减少内核启动的块数量来优化性能。具体而言，可使用 CUDA memset，或让大语言模型（LLM）编写一个持久化内核，将块数量固定在硬件限制范围内。为实现自定义初始化，可将 torch.zeros 替换为不进行初始化的 torch.empty，并添加 cuMemset 或 zero_persistent 代码行。此修改可实现约 28%的加速：

# 原始内核时间：34,493 纳秒
torch.zeros((4096,4096),torch.float16,device="cuda")

# 更新后的memset时间：24,630 纳秒
t=torch.empty((4096,4096),torch.float16,device="cuda")
driver.cuMemsetD16(t.data_ptr(), 0, 4096*4096)

# 更新后的zero_persistent内核时间：24,891 纳秒
zero_persistent(t)

上述单行调试示例展示了 NEUTRINO 定位内核性能瓶颈的能力和简洁性。值得注意的是，NEUTRINO 还支持更高级的用法，包括使用其他实用工具对整个模型进行性能分析：

$ neutrino -p tensorop_count # 统计张量操作数量
$ neutrino -p gmem_bytes # 统计GEMM使用的字节数
$ neutrino -p dmat # 绘制DMAT图
$ neutrino -p <待社区贡献>

unsetunset五、NEUTRINO 可视化：密集内存访问时间线unsetunset

本节介绍一种用于深入可视化 GPU 运行时工作负载行为的图表——密集内存访问时间线（Densified Memory Access Timeline, DMAT），如图 1 和图 11 所示。DMAT 图受页面引用图[11,22]启发而设计，页面引用图以虚拟时间为 x 轴，页面访问为 y 轴，每个点表示在对应时间对页面的一次访问。作为参考标准，页面引用图已被证明是研究虚拟内存管理[11]和替换算法[10]的有用工具。

为适应 GPU 的大规模并行性，DMAT 从两个方面扩展了原始页面引用图：

5.1 物理时间

页面引用图[11,22]和一般内存跟踪工具[66,74]通常使用线程本地自增索引作为表示访问顺序的虚拟时间。然而，在并行环境下，每个线程仅持有部分页面引用图，虚拟时间存在局限性。由于线程的启动时间和执行速度不同，聚合多个线程的页面引用图时，虚拟时间必然会出现错位。

因此，DMAT 采用设备端物理时间以实现可靠聚合。具体而言，我们提供两种类型的 DMAT：

1. 归一化到未同步计算单元（CU）级时钟的起始时钟（图 1、图 11），适用于分析算法行为；
2. 同步到精度较低（MHz 级）的 GPU 本地计时器（图 15），可反映实际内存访问情况，用于硬件/缓存分析。

5.2 页面访问密度

传统页面引用图为二维形式，单个点表示某一时刻对某一页面的访问。但在高度并行的环境中，多个线程可能在同一时刻对同一页面进行并发访问。我们将这种并行访问强度记录为密度，并以颜色深度表示，以区别于传统页面引用图[11,22]中的时间频率，从而突出并行性影响分析的新信息维度。

提出的 DMAT 不仅便于在 GPU 上进行传统内存分析（如数据竞争、访问异常），还为 GPU 运行时分析提供了独特优势：

• 颜色深度：DMAT 中的颜色深度表示并行化密度。
• 与 GPU 本地计时器对齐时，可反映真实内存负载——颜色较浅的区域表明线程间存在未合并访问，颜色过深则提示可能存在内存 I/O 竞争。
• 与起始时钟对齐时，颜色深度反映线程间的分歧——浅色图案通常表示线程执行分歧或工作负载不平衡，可能导致计算能力浪费。

• 空白区域：DMAT 中的空白区域表示页面在较长时间内未被使用，主要包括两种情况（图 11b）：
1. 离散空白区域通常反映计算单元（CU）内的计算过程，其持续时间反映主循环的操作强度[95]，空白区域过长可能意味着流水线效率低下[20]；
2. 结构化空白区域通常反映算法改进，但过大的结构化空白区域可能表明存在时间碎片化和优化空间。

图 11 展示了不同注意力算法在 RTX3080 上的 NEUTRINO DMAT 图，这些图呈现出明显不同的内存访问模式。（a）与图 1 的区别在于使用了独占 SM。通过比较不同算法的 DMAT 图，可直观地发现：

• （b）FlashAttn-v1[20]相对于（c）Memory-Efficient Attention[69]的改进源于 I/O 效率提升；
• 而（a）FlashAttn-v2（图 1）的性能提升则来自更好的流水线设计，这与它们各自的声明一致。

unsetunset六、NEUTRINO 评估unsetunset

我们通过评估来解答关于 NEUTRINO 可靠性和可用性的以下问题：

• NEUTRINO 的结果是否可信（第 6.1 节）？
• NEUTRINO 会带来多少性能和资源开销（第 6.2 节）？
• NEUTRINO 在实际应用程序分析中的效果如何（第 6.3 节）？

我们的评估在两个平台上进行：NVIDIA A100 80GB GPU 和 NVIDIA RTX4090 24GB GPU。NEUTRINO 使用 gcc-11.4.0 编译。所使用的其他软件包括 Python 3.11.4、CUDA 12.6、PyTorch 2.5.0、Triton 3.1.0、CUTLASS 3.5.0 以及 Ubuntu 22.04。

6.1 正确性验证

我们识别并验证了两类正确性：① 执行正确性，确保探测不会改变原始执行流程；② 分析准确性，确保 NEUTRINO 读取的指标合理且正确。

执行正确性

我们通过比较经过探测的内核与原始内核的输出差异来验证执行正确性，因为执行流程的任何差异都可能导致输出改变或系统崩溃。为此，我们在相同的输入和配置下，对每个测试运行两次，一次使用 NEUTRINO 探测，一次不使用探测（即原始执行）。结果表明，经过探测的输出与原始输出之间没有显著差异。

分析准确性

我们从两个角度验证分析准确性。对于与其他分析工具（如 Nsight Compute [58]）重叠的指标（block_sched、gmem_bytes 和 tensorop_count），我们分别使用 NEUTRINO 和 Nsight Compute 运行相同程序，并比较它们的指标读数。捕获的结果显示指标读数一致。

对于现有分析工具未涵盖的新指标，特别是 DMAT（第 5 节），我们通过分析精心设计的、具有预期理论指标值的“微基准测试”内核，并将其与实际指标读数进行比较，来验证其正确性。我们使用禁用了 L1 缓存的 ld/st 指令设计微基准测试内核，用于内存访问，并使用基于自旋的休眠来可控地模拟计算单元内的计算过程。我们实现了多种内存访问模式：线性（Linear）、跨步（Strided）、聚集（Gather）、散射（Scatter）和随机（Random）。我们在 CPU 上模拟这些访问模式，根据访问模式的代码推导出相对于基地址的归一化地址，并根据休眠时间和内存延迟估算连续访问之间的间隔（以周期为单位）。

我们比较理论估算值和实际读数，以评估：① 地址一致性：计算模拟地址序列与分析得到的地址序列之间的汉明距离；② 时钟误差：测量线程内连续访问之间的间隔差异；③DMAT 相似度：计算 DMAT（视为矩阵）之间的均方根误差（RMSE）。

表 2 表明，DMAT 能够正确捕获内存地址（汉明距离为 0），并且实现了小于 200 个周期的时间分辨率（小于循环时间的 7%）。累积的时钟误差会导致 DMAT 的时间线错位，从而产生较大的 RMSE 误差，这对于合并访问（如跨步访问）影响较小，但对于非合并访问则相当显著，例如线性访问的误差约为 60%。这是因为非合并访问具有较低的平均强度（约 16）作为归一化基础，并且经历了更多可变的内存延迟（约 190 个周期的时钟误差）。较大的 DMAT 误差主要是由于静态模拟未考虑内存系统动态（如缓存），而非分析本身的问题，这一点我们留待未来研究解决。

6.2 分析开销

我们评估了两类分析开销：① 性能开销：探测指令导致的内核减速可能会影响与时间相关的分析准确性；② 资源开销：探测带来的额外寄存器使用，这可能会影响块调度，甚至导致寄存器溢出。

性能开销

我们将内核开销定义为：探测指令导致的减速与原始内核执行时间的比值。为了提高准确性，我们通过设备事件计时器（第 4.4 节）评估内核执行时间。

表 3 左侧各列呈现的结果表明，NEUTRINO 具有很高的效率，对于轻量级探测（即 block_sched、gmem_bytes 和 tensorop_count），延迟可控（平均为 1.04 倍），而对于 dmat 等重量级探测，则会导致显著的减速（平均为 7.12 倍）。

我们的分析表明，DMAT 的频繁内存 I/O 是导致减速的原因，且减速程度取决于内存访问的占比和内核执行时间。此外，我们发现轻量级探测可能会异常加速程序，分析表明，这是因为探测指令可以使汇编器获得更好的指令流（IPC 提升 5.88%），从而实现更好的性能（最高可达 0.94 倍加速）。

资源开销

我们将资源开销定义为：经过探测的内核与原始内核所使用的物理寄存器数量之差。

表 3 右侧各列呈现的结果表明，NEUTRINO 探测的开销较低，轻量级探测平均额外使用 3.78 个寄存器，重量级 dmat 探测平均额外使用 5.09 个寄存器。每个探测定义相同数量的逻辑寄存器，但实际使用的物理寄存器数量不同且更少，这一现象也证实了我们使用逻辑寄存器而非物理寄存器的设计是有效的，使得汇编器有可能进行优化。

6.3 扩展研究

我们进一步开展了两项扩展研究，评估 NEUTRINO 在实际工作负载分析中的适用性：

模型分析中的全局内存（GMEM）使用

在实际应用中，开发人员需要分析整个模型而非单个内核，以定位潜在的性能问题。全局内存使用量在此成为一个限制因素，因为大部分全局内存被模型参数占用。因此，我们对 NEUTRINO 在整个模型推理过程的端到端分析中最大全局内存使用量进行了密集测试。我们选择了 ResNet [34]、Stable-Diffusion [71]、Mamba-1.7B [31] 以及 Llama3-1/3/8B [29]。

图 12 所示的结果表明，NEUTRINO 的内存使用效率很高，轻量级探测的内存占用量至少比原始内存占用量小一个数量级，尤其是对于 Llama 等 Transformer 模型。即使在大批次大小（256）下，重量级 dmat 探测的全局内存使用量也大多在原始内存使用量范围内。此外，通过比较 Llama-1B/3B/8B 的结果，我们观察到随着模型规模的扩大，NEUTRINO 的全局内存使用量占原始使用量的比例出人意料地下降。这一发现表明，NEUTRINO 的内存需求增长速度慢于模型规模的增长速度，突显了 NEUTRINO 在分析更大模型时的实用性。

分析工具暴露的延迟

除了内核减速外，分析工具还会带来其他显著的延迟，包括用于分配探测映射和调用探测引擎的序幕（Prologue），以及用于将探测映射复制回主机并将跟踪结果保存到磁盘的尾声（Epilogue）。这些延迟的总和是暴露给上层的延迟。为了评估 NEUTRINO 的整体分析效率，我们通过应用基准测试程序 [8, 89, 91]，将其暴露的延迟与 Nsight Compute [58] 在重叠指标上进行了比较。

图 13 呈现的结果突显了 NEUTRINO 暴露延迟的降低，以及其系统设计和实现的高效性。

unsetunset七、基于 NEUTRINO 洞察的案例研究unsetunset

我们期望 NEUTRINO 成为 GPU 性能工程的有用工具，通过细粒度的运行时洞察为机器学习系统优化铺平道路。本文主要关注 NEUTRINO 的设计和实现，因此如何利用分析结果并非我们的主要目标。

尽管如此，为了展示如何利用 NEUTRINO 和 DMAT 图获得以往无法获得的洞察，我们开展了一项关于同步影响的案例研究：

与 CPU 上的超线程类似，GPU 的流式多处理器（SM）子分区（执行单元）也维护多个候选 warp（GPU 中的调度单元），在每个周期中，warp 调度器会选择一个 warp 运行。这种设计可以减少指令完成时的阻塞等待，特别是对于内存 I/O 等非本地操作，因为可以调度其他候选 warp 来利用核心资源。在实际应用中，这些 warp 可能存在两种情况：① 所有 warp 属于同一个可能需要同步的块；②warp 属于相互独立的不同块。

为了确定同步差异带来的潜在运行时差异，我们在由 Triton [91] 实现的 Flash-Attn-v2 [19] 上创建了一组受控实验。我们利用 Triton 自动调优器创建了两个内核：

• ① 独占块（Exclusive blocks）：128x128 瓦片，2 个阶段，8 个 warp，每个 SM 1 个块（每个 SMSP 上的 2 个 warp 属于同一个块）；
• ② 共享块（Shared blocks）：128x64 瓦片，2 个阶段，4 个 warp，每个 SM 2 个块（每个 SMSP 上的 2 个 warp 可能属于不同的块）。

NEUTRINO 在 A100 上跟踪的 DMAT 图（与内核启动时间对齐）分别如图 1A 和图 11a 所示。尽管这两种配置给 SM 带来了相同的计算负载，并且吞吐量相似，但我们可以发现它们的内存访问模式存在显著差异。

• 对于具有频繁同步的独占块（图 1A），内存访问模式具有结构化特征，对 K 和 V 的访问模式规则，这证实了该算法 [19] 在每个主循环中对 K 和 V 进行一次加载。
• 然而，对于具有独立同步的共享块（图 11a），内存访问模式却出人意料地非结构化，存在许多拖尾块（右侧浅色部分）。

为了验证拖尾效应，我们通过 block_sched 探测（图 10）进行了更深入的分析，以研究块的运行时间。

从图 14A 中的运行时间累积分布函数（CDF）可以看出，独占块的运行时间高度一致，而共享块的运行时间则差异显著（图 14B），拖尾效应高达 24.69%。

为了进一步探究拖尾延迟，我们仔细观察了用于重定向程序形成主循环的 bra 指令（图 3）的执行进度。

该探测通过记录每次分支时的时间戳来对程序执行进行采样，从而可以恢复块的工作进度时间线，如图 14D 所示。此外，通过计算采样时间戳之间的差异，我们可以进一步恢复块内吞吐量（TFLOP/s）时间线（如图 14C 所示）。

从采样的进度和吞吐量时间线中，我们观察到一个有趣的现象：每个共享块都会经历两个阶段：先是一个约 1.8 TFLOP/s 的较慢阶段，然后进入一个约 2.2 TFLOP/s 的较快阶段，直到终止。

此外，通过将采样时间线与启动时间和计算单元 ID 对齐，我们发现从慢阶段到快阶段的过渡大约发生在先前到达的共享块（在快阶段执行）终止时，这表明存在一种类似先进先出（FIFO）的优先调度策略。

这一发现解释了图 11a 中“稀疏后密集”的混乱行为，并与基于该算法 [19] 的、如图 1A 所示的结构化启动波的普遍认知相悖。此外，我们在其他内核（如 GEMM）上也发现了类似的效应，拖尾效应达 50.93%，吞吐量从约 5 TFLOP/s 跃升至约 7.5 TFLOP/s，这突显了共享块中拖尾效应的普遍性。

在性能方面，独占块和共享块【都不是】最优的，但具有不同的性能统计数据。

• 共享块由于类随机的内存访问模式而导致缓存行为不佳，这可以通过硬件分析工具 [55] 捕获的表 4 中 L1 缺失导致的停顿周期高 5.85 倍来验证。
• 相反，独占块的同步程度过高，导致内存使用量显著峰值（图 1B），并且因内存总线繁忙而导致的暴露停顿周期高 4.47 倍，同样，计算流水线竞争导致的停顿周期也多 1.45 倍，这两者都暗示了性能优化的潜在空间。

unsetunset八、讨论与未来工作unsetunset

8.1 NEUTRINO 与 GPU 调度

尽管我们在第 4.6 节和第 7 节中展示了 NEUTRINO 在揭示运行时调度行为新见解方面的重要性，但目前的实验距离全面理解或逆向工程 GPU 调度策略仍有较大差距。

这是因为 GPU 调度是硬件实现的，不同于 CFS 或 EEVDF 等操作系统调度器[83,97]，且具有多层次特性，涵盖流级、块级以及最精细的 warp 级指令调度。

此外，作为共享系统，调度器的行为会受到运行时动态特性的显著影响，这一点也能通过 DMAT 的随机性体现出来。

8.2 GPU 共享的影响

GPU 共享（即并发执行多个内核）是提高 GPU 利用率的实用解决方案。这种共享可以通过 CUDA/HIP 流实现进程内共享，通过多进程服务（MPS）实现进程间共享，或通过多实例 GPU（MIG）实现资源隔离共享。

NEUTRINO 目前是线程本地的（会阻塞线程直至内核完成），且作用范围为进程本地（仅在 MPS 或 MIG 中分析本地进程的执行）。我们将共享带来的影响视为未来一个极具研究价值的方向。

8.3 探针验证的完整性

我们在第 3.4 节中识别并禁止了不安全探针的三个关键因素，但当前的验证机制尚未完善。

• 一方面，仍存在未覆盖的安全因素，例如可能导致程序暂停的不可达同步点；
• 另一方面，当前验证可能过于严格，例如若跳转目标位于探针内部，理论上可支持跳转指令[24]。

GPU 内核验证本身仍是一个开放的研究问题，现有研究仅涉及同步[76]、数据竞争[40]等部分方面。因此，我们将探针验证列为未来的研究工作。

8.4 被探测内核的异常加速

如表 3 所示，NEUTRINO 探测的内核可能呈现出比原始内核更优的性能。基于深入实验，我们将这种加速归因于汇编器优化。

现代汇编器除了将汇编代码转换为机器码外，还集成了多种优化手段，例如重新排序指令以优化执行流程、基于依赖跟踪合并可复用寄存器等。NEUTRINO 探针的额外寄存器和指令可能改变寄存器依赖关系，进而促成更优的执行流程（IPC 提升 5.88%）和更好的性能（最高加速比达 0.94 倍）。这一违反直觉的发现带来了新的研究机遇——由于汇编器和机器码具有硬件导向特性，相关研究尚未得到广泛开展。例如，DeepSeek 近期提出的 DeepGEMM[100]通过翻转机器码中的一个控制位，实现了 10%的性能提升。

8.5 迈向软硬件协同分析器

作为主要基于汇编层的纯软件分析系统，NEUTRINO 无法分析缓存缺失等不可编程硬件事件，尽管可利用 DMAT 轨迹进行缓存模拟[98]。此外，其分析基于执行过程，难以追踪指令未调度的停滞周期，但仍可帮助诊断由内存 I/O 导致的停滞周期。

尽管存在这些不足，NEUTRINO 通过弥补硬件依赖型分析器和内核专属型分析器的信息缺口，可成为现有分析工具的出色补充。

从可观测性的宏观视角来看，NEUTRINO 的多尺度探测特性使其能够充当架构级硬件分析器[7,55,58]与应用级软件分析器[12,87]之间的桥梁。因此，整合平台专属硬件分析器与框架专属软件分析器，构建统一的 GPU 内核分析框架，是未来一个极具吸引力的研究方向。

unsetunset九、相关工作unsetunset

9.1 GPU 硬件分析器

当前的 GPU 内核分析系统，如 NVIDIA 的 NSight[58]/CUPTI[55]或 AMD 的 RGP[7]/GPA[5]，均依赖硬件实现——分析功能需要相应的硬件支持，例如缓存命中率等性能计数器。

• 这些特性虽具有独特性，但难以适配新硬件，例如异步张量核会导致利用率指标[81]可靠性下降，因为计算任务从线程卸载至张量核。
• 此外，它们无法灵活定制以满足开发者需求。例如，硬件分析器仅能以采样方式对整个程序进行分析，且为控制性能和资源开销需采用低采样频率，这限制了其追踪用户指定事件的能力。

相比之下，NEUTRINO 将分析目标限定于所需的跟踪点，同时实现了细粒度事件追踪和低系统开销。

9.2 GPU 软件分析器

其他框架专属的软件分析器，如 PyTorch[87]和 JAX[27]的内置分析器，均为内核专属型，【仅】能捕获高层级事件，例如内存分配/释放事件或内核整体性能（如 FLOP/s）。而 NEUTRINO 专注于指令级的内核内分析。

9.3 GPU 微基准测试

为理解硬件设计，GPU 微基准测试[1,38,51,64,85]试图通过专门设计的工作负载来分析特定硬件的理想性能——这类工作负载仅包含目标指令（如用于基准测试张量核的 mma 指令），不含其他任何指令（甚至真实工作负载中必需的读数据指令 ld 等）以减少干扰。而 NEUTRINO 旨在测量真实工作负载的性能，而非理想工作负载。

9.4 GPU 仿真

另一种理解性能的方式是使用仿真器[9,32,42,49,92]，在 CPU 上以周期级精度模拟 GPU 执行。这类仿真器存在两大主要问题：一是仿真运行速度极慢（可能需要数天）；二是支持新硬件特性和指令的周期极长（可能需要数年）。此外，仿真器难以准确分析各类运行时动态特性（如指令计时）。

9.5 GPU 插桩

在 CPU 领域，通过注入代码、函数或中断并以程序状态为参数的二进制插桩技术[15,24,50,74]已被证明是构建性能工具的有效手段。

GPU 领域也存在一些编译时二进制插桩工具，如 Ocelot[23]、HIPAnalyzer[21]、CUDAAdvisor[77]、CUDAFlux[14]等，以及运行时插桩工具，如 SASSI[82]、NVBit[93]和 GTPin[80]。

尽管基于编译器的方法可利用编译过程中的额外信息，但受限于特定编译器、模块或中间表示（IR），缺乏通用性，且需要源代码，限制了与现有框架的兼容性。直接操作机器码的运行时方法缺乏足够的虚拟化支持，大多依赖栈保护（通过注入纯设备函数），这使得探针间无法协同实现高级功能——例如难以执行指令计时（计算两个时钟读数的差值），因为返回时起始时间已被清除，无法在结束计时器的上下文中获取。

NEUTRINO 依托运行时而非编译器实现通用性，且以并行汇编而非机器码为目标，通过协同探针支持更高级、更复杂的分析任务。

unsetunset十、结论unsetunset

人工智能系统的快速发展催生了对高级 GPU 内核分析工具的迫切需求，以获取全面的运行时见解。

为此，我们提出了 NEUTRINO——一种 GPU 汇编探测架构，通过其独特的探针设计（代码片段、跟踪点和映射），实现了细粒度、多维度且可编程的 GPU 内核运行时分析。

大量实验验证了 NEUTRINO 的可靠性、低开销和实用性。此外，我们通过同步影响的案例研究，利用 NEUTRINO 获得的新见解成功定位了性能瓶颈。为充分发挥 NEUTRINO 的潜力，我们已将其开源（https://github.com/open-neutrino/neutrino），并计划构建协作社区，支持其持续发展，最终打造统一的GPU内核分析框架。

unsetunset参考文献unsetunset