Neutrino: Fine-grained GPU Kernel Profiling via Programmable Probing

Songlin Huang and Chenshu Wu, The University of Hong Kong

A1 主要贡献

本论文介绍了NEUTRINO，这是一个可编程的GPU内核剖析接口，通过汇编层探测实现指令级细粒度剖析、在时间和值域上的剖析多功能性，以及硬件独立性。为了更好地可视化NEUTRINO捕获的丰富细节，我们引入了稠密内存访问时间线（DMAT），这是一种新型表示形式，为GPU运行时行为提供了新洞见。我们在Linux上为NVIDIA和AMD GPU实现了NEUTRINO，并进行了广泛的评估和分析。结果展示了NEUTRINO在GPU内核剖析方面的优越能力，且开销较低。我们设想NEUTRINO作为社区的有价值工具，并已开源以促进未来研究。

随着GPU在缩放定律时代中在计算机系统中扮演越来越重要的角色，理解细粒度的GPU运行时行为比以往任何时候都更关键。然而，现有的GPU内核剖析器通常是内核排他性的或硬件依赖的，往往无法捕获细粒度测量。本文提出了NEUTRINO，这是一个GPU汇编探测工具，用于细粒度、多功能和可编程的GPU内核运行时剖析。受eBPF【24，eBPF, 2024】启发，NEUTRINO的设计旨在将小代码片段（探测）附加到GPU程序中，以暴露程序执行的运行时细节。具体而言，NEUTRINO提取、插桩和重新组装GPU汇编【4，GCN Assembly, 2024】【30，SPIR: The Standard IR for Parallel Compute and Graphics, 2024】【60，PTX ISA 8.5, 2024】，而非机器码【57，Instruction Set Reference, 2024】或编译器【44，LLVM: A compilation framework for lifelong program analysis & transformation, 2004】，从而在一个框架中实现细粒度、多功能性和可编程性。

细粒度：NEUTRINO直接在汇编（最低软件层）上工作，提供指令级的最细粒度，可以有效映射到特定硬件单元，如张量核心和内存I/O。

多功能性：NEUTRINO支持从值域（捕获运行时值，如内存地址）和时间域（记录事件时间戳或甚至通过差分时间戳进行内核内微基准测试）两个视角的GPU内核剖析。通过覆盖这两个维度，NEUTRINO支持从warp/block调度到内存访问模式的多功能剖析任务。

可编程性：NEUTRINO通过利用寄存器作为探测之间的临时存储，将先前GPU插桩框架【14，Cuda flux: A lightweight instruction profiler for cuda applications, 2019】【21，Low-overhead trace collection and profiling on gpu compute kernels, 2024】【77，Cudaadvisor: Llvm-based runtime profiling for modern gpus, 2018】【80，Flexible binary instrumentation framework to profile code running on intel gpus, 2022】【82，Flexible software profiling of gpu architectures, 2015】【93，Nvbit: A dynamic binary instrumentation framework for nvidia gpus, 2019】的可编程性扩展到协作探测。通过这样做，NEUTRINO通过在不同跟踪点和时间的协作探测启用更复杂和灵活的剖析任务。

NEUTRINO以其独特的探测设计脱颖而出，包括三个关键组件：片段、跟踪点和结构化映射，分别对应探测的目标功能、注入点和输出格式。在运行时，NEUTRINO探测将被注入到原始程序的跟踪点中，片段使用逻辑独立的寄存器放置临时结果。这种设计与GPU SIMT模型一起，确保探测对原始程序是虚拟的。此外，通过类似于eBPF的结构化映射，NEUTRINO探测可以灵活地将指标存储到一个或多个缓冲区中，通过无竞争保存而无需昂贵的元数据。

我们为NVIDIA GPU使用CUDA驱动程序和AMD GPU使用ROCm驱动程序在Linux上完全实现了NEUTRINO，包括三个模块：DSL编译器、钩子驱动程序和探测引擎。DSL编译器将以平台独立的Python跟踪DSL编写的探测编译成用TOML【68，TOML, 2024】包装的原始低级汇编探测。钩子驱动程序模拟到驱动程序（共享库）的符号链接，以提供运行时支持，包括从用户捕获GPU调用、分配探测映射并将结果保存到存储。核心探测引擎验证、插桩并从包装的低级探测重新组装探测后的汇编代码。最后，NEUTRINO被封装成类似于bpftrace【13，bpftrace: High-level tracing language for Linux, 2025】的易用CLI，可以通过neutrino -p 运行。

为了更好地可视化NEUTRINO捕获的跟踪，我们引入了一种名为稠密内存访问时间线（DMAT）的新型图表，它改进了先前的页面引用映射（字符串）【11，An anomaly in space-time characteristics of certain programs running in a paging machine, 1969】【22，Working set analytics, 2021】，添加了物理时间信息和来自并行性的内存访问密度。如图1所示，DMAT相对于硬件依赖剖析器（图1B）和内核排他性软件剖析器（图1C）扩展了新的可观测维度，实现了更全面和直观的GPU运行时分析。例如，通过比较它们的DMAT剖面（图11），我们可以视觉和定量地确认FlashAttn-v1【20，Flashattention: Fast and memory-efficient exact attention with io-awareness, 2022】改进了内存效率，而FlashAttn-v2【19，FlashAttention-2: Faster attention with better parallelism and work partitioning, 2024】受益于更好的流水线。

我们进行了全面评估，以验证NEUTRINO在剖析真实GPU工作负载时的可信度、开销和适用性。结果表明NEUTRINO确保了执行正确性（即探测不会改变原始执行流）和剖析准确性（即剖面是可信的）。它还在内核减速（大多数探测仅1.04x）和额外寄存器使用（平均4.11个更多寄存器）方面产生了低开销。此外，我们的广泛评估突出了相对于其他剖析器的高系统效率以及剖析整个模型的能力，甚至对于LLM。为了展示NEUTRINO剖面洞见如何帮助诊断GPU内核中的性能问题，我们对同步对GPU运行时行为的影响进行了案例研究，这揭示了共享GPU块在计算单元上的未注意到的尾随效应，并帮助 pinpoint 性能瓶颈的不同根因。

NEUTRINO目前具有汇编级探测的固有限制，例如无法访问不可编程硬件如缓存。尽管如此，作为一个细粒度、多功能和可编程的GPU内核剖析框架，我们设想NEUTRINO作为研究和工业社区的有价值工具。我们已在https://github.com/open-neutrino/neutrino完全开源NEUTRINO，并希望培养一个全球社区以支持其持续发展。

表1: NEUTRINO与其他GPU剖析器的比较。NEUTRINO是唯一的平台独立运行时GPU内核剖析器，并提供许多独特功能，例如DMAT、通过Python DSL或TOML的可编程性、协作探测、类似于eBPF的映射等。PM代表性能监视器，PC代表程序计数器，PRM代表线程局部页面引用映射。

A3 背景知识/关键Observation/设计原则

GPU剖析：剖析为性能工程构建路线图。与强调分支预测等顺序执行效率的CPU剖析【12，Triangulating python performance issues with Scalene, 2023】【18，LDB: An efficient latency profiling tool for multithreaded applications, 2024】不同，GPU剖析优先考虑计算单元利用率和吞吐量等并行执行可扩展性。以内存访问为例，CPU剖析关注工作集【22，Working set analytics, 2021】等时间局部性，而GPU剖析更关注线程间的合并访问以利用带宽，这开辟了独特的挑战和研究机会。

GPU生态系统：在现代计算机系统中，GPU已成为通用计算单元（GPGPU），由庞大而复杂的生态系统支持，涵盖众多计算任务，如深度学习训练【8，Pytorch 2: Faster machine learning through dynamic python bytecode transformation and graph compilation, 2024】【27，Compiling machine learning programs via high-level tracing, 2018】【79，Megatron-lm: Training multi-billion parameter language models using model parallelism, 2020】和推理【62，TensorRT, 2024】【94，Diffusers: State-of-the-art diffusion models, 2022】【96，Transformers: State-of-the-art natural language processing, 2020】。然而，从编译角度来看，这个复杂的生态系统大致可分为两个分支：❶提前编译（AOT）的操作符库，由手写代码【53，CUDA C++, 2024】通过C++编译器【6，HCC Compiler for ROCm, 2024】【54，CUDA Compiler Driver NVCC, 2024】编译，例如ATen【8，Pytorch 2: Faster machine learning through dynamic python bytecode transformation and graph compilation, 2024】；❷即时编译（JIT）的领域特定语言，如triton【91，Triton: an intermediate language and compiler for tiled neural network computations, 2019】由LLVM【44，LLVM: A compilation framework for lifelong program analysis & transformation, 2004】/MLIR【45，MLIR: Scaling compiler infrastructure for domain specific computation, 2021】编译。如图2所示，这两个分支仅在并行汇编之上分歧。因此，要构建通用剖析器，必须在并行汇编层或以下层构建。

GPU组织层次结构：由于GPU上的线程和内存组织与CPU为并行性而显著不同，我们以下呈现CPU和GPU之间的一些关键差异，这些差异与我们的设计相关。首先，GPU上的并行性是层次化的：32或64个线程分组为一个warp，这是GPU的调度单元，即这些线程共享单个PC，必须在同一周期执行相同的指令。Warp进一步分组为块，这是并发执行单元，即同一块中的线程在一个物理计算单元（CU）中执行，用于通信和同步。最后，块分组为网格，映射到同一GPU，这是主机侧的管理单元，以及内核排他性剖析器的测量单元【87，PyTorch Profiler, 2024】。

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

操作系统范式中的GPU：作为与主机OS通过驱动程序通信的加速器，GPU程序以内核函数为中心，这是GPU计算的入口，其内容将在GPU上执行，而程序的其余部分保留在主机CPU上。具体而言，GPU内核对主机OS被视为原子，即内核内的执行由GPU硬件/固件管理，对主机OS是不可见和不可触及的，这禁止通过成熟的OS技术如ptrace或eBPF【24，eBPF, 2024】观察GPU程序。

除了主机，在GPU线程上剖析也困难，因为GPU线程的系统功能高度受限。GPU程序是直接执行，没有像OS内核那样的管理层。特别是，没有定时中断的支持，这是基于采样的剖析器的关键特征。因此，传统OS上的剖析技术，例如采样和扫描栈帧，不适用于GPU。而且，没有常用支持的磁盘I/O，使得保存结果变得麻烦。

作为探测接口的汇编：GPU编程的独特特性为构建NEUTRINO（一个细粒度、多功能和可编程的GPU剖析工具，既内核包容又硬件独立）带来了巨大挑战。我们需要回答的关键问题是：NEUTRINO应该在哪个层构建以及如何构建？在本文中，在图2的所有选择中，我们选择并行汇编，如PTX/GCNAsm【4，GCN Assembly, 2024】【60，PTX ISA 8.5, 2024】，设计用于适应系统和机器码的快速变化，作为探测接口。重要的是，我们采用更强大但具有挑战性的方法，即在运行时附加探测，而不是静态方法，如通过编译器自定义剖析通道【21，Low-overhead trace collection and profiling on gpu compute kernels, 2024】或C中的朴素asm()。这种设计选择不仅允许向前和向后兼容性，还在各种方面承诺了明显的优势。

❶面向硬件：作为低级接口，汇编可以捕获对性能分析重要但用高级语言难以跟踪的硬件事件。例如，PTX中只有四个与内存访问相关的指令，即ld、st、cp.async和tensormap。相比之下，在CUDA C++中使用对象和模板，难以分类和捕获所有可能的内存访问。

❷特殊寄存器/指令：并行汇编的特殊寄存器包含有用的运行时信息用于剖析。例如，GCNAsm的hwreg告诉线程调度在哪个计算单元上，而PTX的特殊寄存器%clock（以CU本地时钟周期）和%globaltimer（以GPU本地纳秒）有助于测量时间戳，并可作为指令级定时器，如图3所示。

❸兼容性：如图2所示，并行汇编是AOT和JIT编译的最高共同层。例如，PTX是由基于gcc的nvcc【54，CUDA Compiler Driver NVCC, 2024】编译的CUDA C++和DSL（如由LLVM【86，LLVM PTX Backend, 2024】支持的Triton）的共同输出。因此，在汇编上探测可以与大多数基础设施兼容，而基于编译器的方法限于特定编译器或IR(s)。

❹覆盖范围：基于编译器的方法需要源代码，假设用户已定位性能差的GPU内核，这是不常见的，因为大多数程序有许多内核。相反，运行时方法可以覆盖所有用户代码，能够扫描性能差的内核。

汇编层运行时探测的设计选择也带来了独特的挑战，例如在没有编译器支持的情况下保护探测、在运行时定位GPU代码、获取高级上下文等。我们在NEUTRINO中克服了这些挑战，并使其成为强大的GPU剖析器。

A2 方法细节

3.1 可编程探测接口

可编程探测接口：如图4所示，NEUTRINO在其探测设计中具有三个关键元素：片段、跟踪点和结构化映射。

片段：与探测目标相同，NEUTRINO的片段是汇编，具有一些助手，如SAVE用于记录结果，以及OUT/IN1/IN2用于读取寄存器（指令操作数）用于值剖析。开发者还可以使用其他汇编功能，特别是S_MEMTIME用于时间剖析。

跟踪点：NEUTRINO主要在最细指令级制定探测跟踪点，这确保了时间准确性和硬件粒度，如用于张量核心操作的wmma/mma。通过分组指令，NEUTRINO的跟踪点可以扩展到更大规模，如设备函数调用和线程开始/结束。

映射：类似于eBPF【48，Understanding performance of ebpf maps, 2024】，NEUTRINO的映射显式结构化保存格式，以解决持久性问题，这是GPU上由于并行性和层次化组织的巨大元数据而麻烦的问题。NEUTRINO主要在两个级别定义映射（§3.3）：❶线程级：每个线程保存，用于值剖析；❷warp级：仅warp领导线程保存，用于时间剖析。

除了制定自定义探测的三个组件之外，NEUTRINO可编程性的关键设计是协作性：❶同一线程的NEUTRINO探测可以通过利用寄存器作为临时存储来协作，用于高级剖析任务，同时保持效率，因为寄存器使用在GPU核心中是并行的。❷NEUTRINO探测还可以通过GMEM中的映射协作，即不同探测可以贡献并通过同一映射协作。

3.2 虚拟化探测执行模型

虚拟化探测执行模型：由于GPU程序对OS是静态的，即所有代码（汇编）在执行前加载并已知，我们选择直接将探测放置在原始汇编中，而不进行保护，例如栈，以实现协作性。我们识别出，通过这样做，NEUTRINO探测仍然从原始程序虚拟执行。如图5A所示，这种虚拟化通过时间和资源分离实现。

时间分离：NEUTRINO的时间虚拟化源于GPGPU的SIMT执行模型，其中并行发生在线程之间，而线程内的执行通常是顺序的，每个周期一条指令。因此，由于探测直接插入到汇编中，它们与原始程序的时间分离将被保证。

资源分离：类似于CPU，GPU线程也有线程私有寄存器作为其主要资源，其中包含来自ALU的中间结果、共享或全局内存的地址等。NEUTRINO通过分离独立的寄存器组以及其他资源如GMEM来虚拟化探测寄存器。因此，NEUTRINO探测可以避免影响原始程序的资源和执行流。值得注意的是，NEUTRINO的探测寄存器组是在汇编级逻辑声明的，而不是物理的。因此，NEUTRINO可能不一定引入额外的物理寄存器使用（表3），因为声明的逻辑寄存器将被汇编器在寄存器分配中集成到物理寄存器中，通过依赖跟踪算法【39，A novel renaming scheme to exploit value temporal locality through physical register reuse and unification, 1998】【72，Register integration: a simple and efficient implementation of squash reuse, 2000】保留探测和原始程序之间的独立性。

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

用于持久性的结构化映射：持久性是GPU剖析的关键挑战。虽然线程执行是并行和独立的，但底层内存系统是共享的，导致并发保存之间的竞争条件。因此，先前解决方案【80，Flexible binary instrumentation framework to profile code running on intel gpus, 2022】【93，Nvbit: A dynamic binary instrumentation framework for nvidia gpus, 2019】广泛使用原子操作来分离持久空间，这在大量并行性下可能变得低效。而且，GPU层次化组织创建了丰富的元数据，如threadIdx和blockIdx（24字节），这对分析有用但可能消耗存储【93，Nvbit: A dynamic binary instrumentation framework for nvidia gpus, 2019】。

受无锁per-cpu eBPF映射【48，Understanding performance of ebpf maps, 2024】和HIPAnalyzer的事件缓冲区【21，Low-overhead trace collection and profiling on gpu compute kernels, 2024】的启发，我们显式将NEUTRINO映射结构化为ndarray布局，如图5B所示，其形状由启动配置（blockDim, gridDim）和映射定义（level, type, size, cap）确定。这实现了无竞争保存，因为每个线程有独立的段，并减少了存储压力，因为大多数元数据可以推断而不是直接保存。NEUTRINO主要在两个级别制定映射。

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

Warp级：warp级映射简化了线程级映射，用于时间剖析，布局为[#Grid, #Warp, cap]。因为同一warp内的线程一起调度执行同一指令，记录事件时间戳只需要warp内的一个线程，而不是所有，这可以显著减少内存和存储压力。

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

3.4 用于安全的验证

用于安全的验证：验证【28，Simple and precise static analysis of untrusted linux kernel extensions, 2019】【84，Validating the eBPF verifier via state embedding, 2024】已被证明对可编程探测至关重要【24，eBPF, 2024】，因为不安全的探测可能破坏原始程序的执行流并使剖析无效。验证器还可以帮助开发者编写正确的探测。在NEUTRINO中，我们识别并防止三个关键安全问题：覆盖原始寄存器：如§3.2所述，GPU线程使用寄存器作为主要资源。因此，修改原始程序使用的寄存器是不安全的。例如，修改持有全局内存地址的寄存器可能导致非法内存访问（图6A）。因此，NEUTRINO要求探测使用独立的寄存器组，并禁止修改原始寄存器的探测。

程序乱序：虽然SIMT程序模型保证线程内的指令线性执行，但有流控制指令，如GCNAsm的S_BRANCH或PTX的bra，可能改变执行顺序，这对探测是不安全的，因为它们可能破坏原始执行顺序（图6B）。因此，NEUTRINO禁止探测使用改变执行流的指令。

共享内存：作为加速的重要因素，共享内存已被高度优化用于存储【17，Tvm: an automated end-to-end optimizing compiler for deep learning, 2018】【91，Triton: an intermediate language and compiler for tiled neural network computations, 2019】和访问效率【89，CUTLASS, 2023】。因此，来自探测的额外共享内存使用可能大大影响或失败执行，如果原始使用已达到硬件限制。因此，NEUTRINO禁止探测使用共享内存。

4.1 钩子驱动程序

钩子驱动程序：虽然OS中的驱动程序大多指通过read/write/ioctl系统调用暴露的内核扩展，如nvidia.ko，但大多数供应商还维护更高层的用户空间驱动程序作为共享库，如Linux中的libcuda.so或libamdhip.so。这些驱动程序共享库通常高度复杂且闭源。然而，给定ELF中的符号通过其签名解析，我们可以通过定义具有匹配签名的所有函数并在内部使用dlfcn定位和调用实际驱动程序的真实函数来构建巧妙的钩子驱动程序（附录B中有更多细节）。与eBPF uprobe【24，eBPF, 2024】等其他方法相比，我们的钩子驱动程序更安全且更灵活，因为所有代码在用户模式执行，支持fork/wait，这对与探测引擎的交互很重要。我们利用钩子驱动程序提供以下支持。

代码跟踪：与OS隐式加载的CPU代码相比，GPU代码以ELF【25，Executable and Linkable Format, 2024】或FatBinary【56，Fat Binaries, 2024】格式，需要通过cuModuleLoad显式加载。其他函数，如cuModuleGetFunction，也可用于从模块定位特定内核。我们钩住这些API（图8）来捕获所有加载的图像、提取的内核以及从内核到代码的映射。每个图像从其头部memcpy到图像存储中，以避免被用户程序的资源管理释放。

运行时探测：由于执行是非本地的，启动GPU内核函数不是简单添加栈帧，而是需要显式驱动程序调用cuLaunchKernel或hipModuleLaunchKernel。我们钩住这些API以提供运行时支持：❶它通过原始内核（指针）在内核存储中搜索探测后的内核；❷它根据元数据分配探测缓冲区；❸它启动探测后的内核并同步完成，即GPU上的探测缓冲区是指标读数；❹它memcpy探测缓冲区回CPU，然后fwrite并释放探测缓冲区。

更重要的是，当在内核存储中未找到探测后的内核时，钩子驱动程序还负责与探测引擎交互：❶它在图像存储中搜索包含内核的二进制文件。❷它在目录中fwrite二进制文件并fork子进程调用探测引擎。❸等待后，它查找目录并加载回探测后的内核和元数据，例如探测数量。❹最后，探测后的内核和元数据将被添加到内核存储。失败的内核也将以status=false添加到内核存储以避免重复。

其他函数未钩住，我们通过解析头cuda.h/hip_runtime.h和共享库的符号表自动生成它们，如libcuda.so。

4.2 探测引擎

探测引擎：如图9所示，NEUTRINO探测引擎首先objdump转储的GPU二进制文件，以提取文本格式的并行汇编。然后，它将使用内核名称匹配并修剪多内核原始汇编成单内核汇编，同时保留全局定义和设备函数。

接下来，NEUTRINO将处理并添加从环境变量读取的探测，涉及以下步骤：❶它规划探测映射，根据映射定义（§3.3，level/type/size/cap）和线程索引（如threadIdx）指导每个warp/线程到其映射段，如附录C中详细说明。❷它粗略解析内核汇编成参数、寄存器声明和指令。然后匹配跟踪点到特定指令。❸它彻底解析每个匹配的行（例如，ld.global.u64 %rd1, [%rd2];//%rd1=*%rd2）成令牌，如操作码（ld.global.u64）和操作数%rd1, %rd2。然后填充片段中的助手，如ADDR，到真实寄存器%rd2。最后，它在匹配指令之前/之后放置片段，以及内核参数末尾的映射地址，以及内核开头的映射规划汇编。

探测后，它通过汇编器如ptxas【61，PTXAS, 2024】将探测后的汇编转换为机器码。探测引擎还将保存对钩子驱动程序有用的内核元数据，如探测、映射、回调等。

4.3 探测DSL和编译器

探测DSL和编译器：探测引擎的一个实际问题是探测是汇编，这是一种低级且硬件依赖的语言。直接汇编编程对一般开发者可能不太友好。因此，为了增强NEUTRINO的硬件独立性和可用性，我们提出一个最小化的Python领域特定语言（DSL）作为NEUTRINO探测的高级接口，类似于eBPF的bpftrace【13，bpftrace: High-level tracing language for Linux, 2025】【24，eBPF, 2024】。注意DSL对NEUTRINO是可选的，有经验的开发者仍然可以手工制作汇编用于高级用法。

如图10所示，NEUTRINO DSL密切遵循Python语法，允许用户使用@probe装饰器声明探测，跟踪点指定为装饰器参数，片段作为函数体。类似地，映射可以用@Map装饰器声明，结构定义为类成员。而且，跨探测共享的上下文探测寄存器可以通过值赋值语法定义，类型在全局范围内注释。NEUTRINO探测不允许使用其他函数如open。相反，我们提供助手操作数如nl.addr用于读取寄存器，以及助手函数如nl.clock()和Map.save()用于获取设备侧时钟和保存结果。

这个DSL将被即时编译成平台特定的基于汇编的探测，分两个步骤，图10的编译示例在附录D中：❶它使用Python的ast模块解析并转换跟踪DSL到类似于eBPF ISA【90，BPF Instruction Set Architecture (ISA), 2024】的中间表示（IR）。❷IR将被翻译成平台特定的汇编，即CUDA的PTX Assembly和ROCm/HIP的GCNAsm，助手操作数将被保留供探测引擎使用。我们将IR设计成类似于eBPF的，可能启用重用成熟的eBPF工具链，如著名的eBPF验证器【28，Simple and precise static analysis of untrusted linux kernel extensions, 2019】。

from neutrino import probe, Map
import neutrino.language as nl
CALLBACK = "block_sched.py"  # for trace analysis
# declare maps for persistence
@Map(level="warp", type="array", size=16, cap=1)
class block_sched:
    start: nl.u64
    elapsed: nl.u32
    cuid: nl.u32
# declare probe registers shared across probes
start: nl.u64 = 0  # starting clock
elapsed: nl.u64 = 0  # elapsed time, initialized to 0
# define probes with decorator
@probe(pos="kernel", level="warp", before=True)
def thread_start():
    start = nl.clock()
@probe(pos="kernel", level="warp")
def thread_end():
    elapsed = nl.clock() - start
    block_sched.save(start, elapsed, nl.cuid())

4.4 实用工具

实用工具：生态系统集成：仅钩住驱动程序或在汇编上探测将缺少如张量形状的高级信息，这对分析有用。因此，我们还实现了如PyTorch【8，Pytorch 2: Faster machine learning through dynamic python bytecode transformation and graph compilation, 2024】的生态系统集成实用工具，通过Python sys.settrace暴露高级张量信息。

基准测试模式：为了评估系统开销并为重探测提供时间对齐，NEUTRINO提供基准测试模式，使用CUDA/HIP事件定时器启动探测后的内核和修剪后的内核（剥离探测并以相同配置组装），以基准测试探测引起的额外执行延迟。

分析代码生成和回调：为了促进分析，NEUTRINO支持基于映射定义（如图10）生成Python跟踪解析代码，以方便用户从跟踪中提取信息，附录E中有示例。而且，NEUTRINO支持添加回调（如图10的第3行）用于自动化后验分析。

源代码注释：为了提供更精确的控制，我们实现了类似于NVTX的API【59，NVTX: NVIDIA Tools Extension SDK, 2025】的源代码注释工具。探测引擎可以查找lineinfo（特殊注释，例如".file 1 example.py"和".loc 1 33 45"）并与源代码交叉引用（第1文件example.py的第33行）以包括或排除指令。

4.5 将NEUTRINO扩展到其他平台

将NEUTRINO扩展到其他平台：虽然当前实现仅支持NVIDIA和AMD GPU，但NEUTRINO可以扩展到其他平台，如Intel oneAPI【26，Unified Acceleration Foundation. oneAPI Programming Model, 2025】。NEUTRINO的硬件独立性源于其在跨平台两个共同组件的设计：并行汇编以适应快速的架构级演变，以及从主机OS控制执行的驱动程序。

在实践中，要将NEUTRINO扩展到其他平台，需要实现钩子驱动程序、探测引擎和（可选）DSL编译器后端。对于钩子驱动程序，由于大多数功能标准化为平台无关模块，我们预计大多数变化围绕API重命名和调试，例如，对于ROCm/HIP支持的cuLoadModule到hipLoadModule。至于探测引擎，需要一个新的解析器和匹配器用于不同的汇编语法（例如GCNAsm【4，GCN Assembly, 2024】），但整体基础设施（图9）保持不变。DSL编译器需要一个后端将我们的类似于eBPF的IR翻译成汇编。我们预计这类似于Triton【91，Triton: an intermediate language and compiler for tiled neural network computations, 2019】如何通过扩展代码生成（像我们的探测引擎和DSL编译器）和启动器（像我们的钩子驱动程序）支持新硬件。

4.6 使用：将所有整合在一起

使用：将所有整合在一起：上述组件形成了NEUTRINO的用户友好且易用的剖析工具。与许多框架如PyTorch【8，Pytorch 2: Faster machine learning through dynamic python bytecode transformation and graph compilation, 2024】、Triton【91，Triton: an intermediate language and compiler for tiled neural network computations, 2019】和JAX【27，Compiling machine learning programs via high-level tracing, 2018】兼容，NEUTRINO的使用与bpftrace一样简单，具有许多内置工具，如block_sched（图10）来检查内核的块调度成本。NEUTRINO的用户友好性通过一个简单示例最佳展示，其中我们尝试剖析以下PyTorch代码行并获得洞见。

要这样做，用户只需使用–probe/-p选项运行NEUTRINO CLI：

$ neutrino -p block_sched python -c "torch.zeros(... 然后，当完成时，跟踪将被放置在目录中，并从分析回调打印出消息如下： vectorized_elementwise: # kernel name, truncated No.block:32768 Exec:680869 Sched:142674 (cycle/SM) 这里vectorized_elementwise内核【88，PyTorch Vectorized Elementwise Kernel, 2024】，广泛用于一元张量操作，用于用零初始化分配的内存。我们通过模拟块分派到CU来测量调度时间。对于CU上的每个记录块，如果其开始时钟大于任何现有块的结束时钟（start + elapsed），则发生块替换，调度周期由下一个start - 前一个end估计，执行周期由前一个elapsed测量（附录E中的完整代码）。 剖析结果揭示了令人惊讶的∼20%时间花在调度块到物理SM上，因为内核启动了大量（32,768）块，每个块的执行时间相对较小。基于洞见，可以通过减少内核启动的块数来优化性能。要这样做，我们可以使用CUDA memset或指示LLM编写持久内核（附录F中的示例），将块数固定到硬件限制。要应用自定义初始化，我们用torch.empty替换torch.zeros，该函数分配内存而不初始化，并添加一行cuMemset或zero_persistent。此修改提供了∼28%加速： # Original kernel time: 34,493 ns torch.zeros((4096,4096),torch.float16,device="cuda") # Updated memset time: 24,630 ns t=torch.empty((4096,4096),torch.float16,device="cuda") driver.cuMemsetD16(t.data_ptr(), 0, 4096*4096) # Updated zero_persistent kernel: 24,891 ns zero_persistent(t) # code in Appendix F 上述一行调试示例展示了NEUTRINO pinpoint 内核内性能瓶颈的能力和简单性。注意NEUTRINO还支持超出上述示例的更高级用法，包括使用其他有趣工具剖析整个模型： $ neutrino -p tensorop_count # number of tensor op $ neutrino -p gmem_bytes # number of GEMM bytes used $ neutrino -p dmat # draw DMAT Plot $ neutrino -p

5 NEUTRINO可视化

DMAT介绍：在本节中，我们介绍稠密内存访问时间线（DMAT）图，这是一种对GPU运行时工作负载行为的洞见可视化，如图1和图11所示。

DMAT图受页面引用映射【11，An anomaly in space-time characteristics of certain programs running in a paging machine, 1969】【22，Working set analytics, 2021】启发，该映射显示虚拟时间作为x轴，页面访问作为y轴，其中一个点表示在时间对页面的访问。作为参考标准，页面引用映射已被证明是虚拟内存管理【11，An anomaly in space-time characteristics of certain programs running in a paging machine, 1969】和替换算法【10，A study of replacement algorithms for a virtual-storage computer, 1966】研究的有用工具。

为了适应GPU的大规模并行性，DMAT从两个视角扩展了原始页面引用映射：物理时间：页面引用映射【11，An anomaly in space-time characteristics of certain programs running in a paging machine, 1969】【22，Working set analytics, 2021】和一般内存跟踪器【66，Lightweight memory tracing, 2013】【74，Using valgrind to detect undefined value errors with BitPrecision, 2005】通常使用线程本地自增索引作为虚拟时间，表示访问顺序。然而，我们发现虚拟时间在并行性下不足，因为每个线程仅持有整个页面引用映射的一部分。由于线程的起始时间和执行步调不同，当聚合单个线程的页面引用映射以形成完整跟踪时，虚拟时间之间将不可避免地存在未对齐。

因此，对于DMAT，我们使用设备侧物理时间提供可靠聚合。具体而言，我们提供两种类型的DMAT：❶归一化到起始时钟（图1、图11）的未同步CU级时钟，适合分析算法行为；❷同步到不太准确（MHz）GPU本地定时器（图15），表示实际内存访问，用于硬件/缓存分析（附录G）。

页面访问密度：先前的页面引用映射是2D的，其中一个点表示在时间对页面的访问。然而，对于高度并行化的环境，可能有许多来自多个线程的并发访问同一页面在同一时间。我们记录这种并行访问强度作为密度，并将其标记为颜色深度，以将其与传统页面引用映射【11，An anomaly in space-time characteristics of certain programs running in a paging machine, 1969】【22，Working set analytics, 2021】中的时间频率区分开来，突出了用于分析并行性效果的新信息维度。

所提出的DMAT不仅促进了GPU上的传统内存分析（例如数据竞争、访问异常），还为GPU运行时分析带来了独特的好处。

■颜色深度：DMAT中的颜色深度表示并行化的密度。当对齐到GPU本地定时器时，它反映了真实内存负载，浅色区域 pinpoint 线程间未合并访问，过度强度表示潜在内存I/O竞争。当与起始时钟对齐时，颜色深度反映线程间的分歧，其中淡色模式通常表示分歧线程执行或不平衡工作负载，可能浪费计算能力。

■空洞：DMAT中的空洞表明页面在显著时间内保持未使用，这包括两种情况（图11b）：❶离散空洞通常反映CU内的计算。持续时间反映每个主循环的操作强度【95，Roofline: an insightful visual performance model for multicore architectures, 2009】，太长的空洞可能在流水线【20，Flashattention: Fast and memory-efficient exact attention with io-awareness, 2022】中低效；❷结构化空洞通常反映算法改进，而额外大的结构空洞可能反映时间碎片化和优化机会。

A4 实验环境

数据集：无明确数据集，但使用真实GPU工作负载，如Flash-Attn-v2【19，FlashAttention-2: Faster attention with better parallelism and work partitioning, 2024】、FlashAttn-v1【20，Flashattention: Fast and memory-efficient exact attention with io-awareness, 2022】、Memory Efficient Attention【69，Self-attention does not need o(n2) memory, 2022】、ResNet【34，Deep residual learning for image recognition, 2015】、Stable-Diffusion【71，High-resolution image synthesis with latent diffusion models, 2022】、Mamba-1.7B【31，Mamba: Linear-time sequence modeling with selective state spaces, 2024】、Llama3-1/3/8B【29，The llama 3 herd of models, 2024】等，用于剖析和基准测试。规模包括大批量（256）和模型如LLM，用于内存使用测试。用途：用于验证剖析准确性、开销和适用性，包括DMAT可视化。

模型架构关键参数：评估中使用了如GEMM (M=N=K=2048或4096)、SoftMax、Flash-Attn-v2等内核，参数包括tile大小（如128x128）、阶段（2或3）、warp数（4或8）等。

硬件配置：NVIDIA A100 80GB GPU和NVIDIA RTX4090 24GB GPU。连接关系：标准PCIe连接到主机。

软件配置：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。NEUTRINO使用C（≈2500行）和Python（≈3000行）实现。

A4 实验结果

正确性验证：执行正确性通过比较探测后和原始内核的输出差异验证，结果显示无显著差异。剖析准确性通过与Nsight Compute【58，NSight Compute System, 2024】比较重叠指标（如block_sched、gmem_bytes、tensorop_count）验证，结果一致。对于新指标如DMAT，通过设计微基准内核验证，包括线性、跨步、收集、散布和随机访问模式。表2显示DMAT正确捕获内存地址（汉明距离为零）和实现小于200周期时间分辨率（<循环时间的7%）。RMSE错误因未合并访问的动态内存延迟而较高。

剖析开销：性能开销：轻量探测平均1.04x减速，重DMAT平均7.12x。资源开销：轻量探测平均3.78额外寄存器，重DMAT 5.09。异常加速归因于汇编器优化（附录K）。

广泛研究：GMEM使用：在ResNet、Stable-Diffusion、Mamba-1.7B和Llama模型上测试端到端剖析，轻量探测内存使用至少比原始小一个数量级，重DMAT在原始内存使用内。随着模型缩放，NEUTRINO内存比例下降。暴露延迟：与Nsight Compute相比，NEUTRINO在重叠指标上减少了暴露延迟。

案例研究：在Flash-Attn-v2和GEMM上，独占块显示一致执行时间，共享块显示尾随效应（高达24.69%和50.93%）。DMAT显示共享块的非结构化访问，导致更高停顿周期（表4）。吞吐量时间线显示共享块从慢阶段（∼1.8 TFLOP/s）到快阶段（∼2.2 TFLOP/s）的转换，暗示FIFO-like调度。

A5 结论

本论文介绍了NEUTRINO，这是一个GPU汇编探测基础设施，通过其独特的片段、跟踪点和映射探测设计实现了细粒度、多功能和可编程的GPU内核运行时剖析。我们为CUDA和ROCm生态系统实现了NEUTRINO，包括钩子驱动程序、探测引擎和DSL编译器。我们引入了新型稠密内存访问时间线（DMAT）来有效可视化全面的GPU内存访问模式。我们进行了广泛实验，验证了NEUTRINO的可信度、低开销和适用性。此外，我们对同步影响进行了案例研究，使用NEUTRINO获得的新洞见成功 pinpoint 性能瓶颈。为了最大化NEUTRINO的潜力，我们已在https://github.com/open-neutrino/neutrino开源，并计划构建协作社区以支持其向统一GPU内核剖析框架的持续增长。

未来工作包括理解GPU调度、GPU共享影响、完整探测验证、异常加速的汇编器优化，以及整合硬件-软件剖析器。

A6 附录

B 钩子驱动程序：钩子驱动程序通过暴露相同签名的函数并使用dlfcn调用原始函数实现，钩子留给添加功能。钩子驱动程序通过LD_PRELOAD注入程序，用于动态加载，或LD_LIBRARY用于静态链接。这也可以扩展以过滤专有产品如cublas以确保法律安全。

C 探测映射规划：我们的探测引擎根据定义自动规划探测映射，例如thread:array:8:1。其机制类似于以下CUDA/HIP C++代码。

#define NO_BYTES ... // filled by datamodel, 8*1 here
__global__ void tmp(void* buff) {
  int thr_idx = (blockDim.y * threadIdx.z + threadIdx.y) * blockDim.x + threadIdx.x;
  int blk_idx = (gridDim.y * blockIdx.z + blockIdx.y) * gridDim.x + blockIdx.x;
  int blk_size = blockDim.x * blockDim.y * blockDim.z;
  int buf_idx = blk_idx * blk_size + thr_idx;
  void* buf_loc = buff + buf_idx * NO_BYTES;
}

在实践中，探测引擎通过格式化Python字符串的汇编与name和no_bytes实现它，如以下PTX示例。

D 探测DSL编译示例：我们的编译器首先将DSL翻译成类似于eBPF的IR。然后，这个中间表示被翻译成平台特定的汇编，如PTX和GCNAsm。在AMD GPU上，标量寄存器每个warp和向量寄存器每个线程是分离的，这反映在寄存器中的新type字段中。我们还为AMD平台实现了更多助手来覆盖缺失功能，如SUB64用于64位减法，因为AMD不支持标量寄存器的64位加/减【4，GCN Assembly, 2024】（我们通过s_sub_u32仿真它）。

E 跟踪分析代码示例：典型的NEUTRINO跟踪分析代码包括两部分：❶ NEUTRINO探测引擎自动生成的跟踪读取代码，用于通过struct将二进制跟踪读取到Python数据结构中。❷ 基于读取结构的跟踪分析代码。整个程序可以注册到钩子驱动程序作为CALLBACK，如以下示例。

F 用于torch.zeros的持久内核：以下代码由GPT4o生成，提示："Write a Persistent Kernel to initialize a torch tensor to 0"。

import torch
import triton
import triton.language as tl
@triton.jit
def zero_persistent_kernel(output_ptr, numel, BLOCK_SIZE: tl.constexpr, NUM_SMS: tl.constexpr):
    start_pid = tl.program_id(axis=0)
    num_blocks = tl.cdiv(numel, BLOCK_SIZE)
    blocks_per_sm = num_blocks // NUM_SMS
    if start_pid < num_blocks % NUM_SMS:
        blocks_per_sm += 1
    block_id = start_pid - NUM_SMS
    for _ in range(blocks_per_sm):
        block_id += NUM_SMS
        offsets = block_id * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
        mask = offsets < numel
        tl.store(output_ptr + offsets, tl.zeros([BLOCK_SIZE], dtype=tl.float16), mask)
def zero_persistent(x: torch.Tensor):
    numel = x.numel()
    NUM_SMS = torch.cuda.get_device_properties("cuda").multi_processor_count
    BLOCK_SIZE = 128
    grid = lambda META: (min(NUM_SMS, triton.cdiv(numel, META['BLOCK_SIZE'])),)
    zero_persistent_kernel[grid](x, numel, BLOCK_SIZE, NUM_SMS)
t = torch.empty((4096, 4096), torch.float16, device="cuda")
zero_persistent(t)

G 全局定时器对齐的DMAT：NEUTRINO DMAT也可以同步到GPU本地定时器，即不同计算单元共享同一同步定时器，而不是CU本地周期定时器。同一配置下的DMAT如图1在全局同步下成为图15。在全局同步下，DMAT变得阶梯状，每个"步"代表虚拟块的启动波，在这种情况下为32x128，到物理计算单元（A100上的108 SM）。我们推荐放大每个"步"，用于洞见分析。放大的视图（图15的右部分，注意y轴值差异）反映了应用于硬件的实际内存访问模式，并对硬件级分析和缓存模拟有洞见。

H 微基准测试细节：这里我们呈现理论估计背后的更多细节，以及时钟估计变化的原因。以最复杂的随机访问为例，我们使用Fisher–Yates shuffle在A中创建可重现的随机索引序列。在运行时，每个线程从A读取随机索引并相应写入B。然后，在验证中，我们重建相同的随机索引并模拟访问。我们将模拟地址与收集的地址对齐到基地址。然而，一个关键挑战是内存访问延迟（代码中的LATENCY），在实践中将因缓存命中/未命中而变化。为了公平，我们禁用L1（通过cg缓存操作符【60，PTX ISA 8.5, 2024】）并选择L1禁用延迟（A100上的570周期）作为LATENCY。因为不同模式有不同的L1命中率（线性88%和跨步0%），导致不可预测的估计错误。

I 评估NEUTRINO探测减速：如表3所示，NEUTRINO的dmat探测导致不同内核的不同减速程度，从Flash-Attn-v2的2.75x到SoftMax的13.163x。我们首先从内存I/O视角分析这种效果，通过采样DMAT访问添加采样计数器，使DMAT仅保存1/2、1/3、1/4、1/8，并测量内核减速。图16的结果表明DMAT的减速主要与采样频率成线性关系，突出了大多数DMAT成本来自I/O。而且，我们识别了两个主要因素贡献于确切减速比率。首先，更小的内核（即较低块时间）遭受更大减速，因为dmat的内存I/O比例相对于原始块时间更大（pool，∼5555周期13x减速和attn，∼205101周期2.75x减速）。其次，对于类似块时间的内核，内存指令比例显著影响减速（attn，11.03%内存指令2.75x减速，和gemm，17.57%内存指令6.55x减速）。

J GEMM内核案例研究：这里，我们还制定了GEMM内核（M=N=K=4096）的两个配置，对物理SM施加相同负担：❶ 128x256x64 tile，3阶段，8 warps，A100上每个SM 1块；❷ 128x128x64 tile，3阶段，4 warps，A100上每个SM 2块。我们呈现块完成时间的累积分布函数分别在图17A和图17B。我们可以识别与§7相同的模式，即共享块遭受尾随块。类似地，我们绘制了块内级吞吐量时间线和warp执行时间线在图17A和图17B，产生与§7类似的发现。但不同于前者，这里BLOCK_N，128或256，是并行维度。图14和图17一起表明，由于GPU同步和调度（§7），共享块中的尾随效应存在是通用的，无论内核或tile组织（并行或顺序）。

K 分析异常加速：如表3所示，NEUTRINO探测可能导致异常加速。这样的加速的一个显著案例是GEMM与M=N=K=2048，其中block_sched可以导致CUTLASS实现高达0.94x加速和Triton实现的0.96x加速。我们挑选Triton实现（BLOCK_M=128, BLOCK_N=128, BLOCK_K=32, STAGES=3, WARPS=4），和block_sched探测，因为它对程序的最低中断（仅线程开始和结束）。我们使用硬件剖析器Nsight Compute【58，NSight Compute System, 2024】剖析了探测后和原始内核（相同-O3），结果表明戏剧性变化。剖析指标突出了5.88% IPC改进和4.97%更忙的指令调度。而且，我们识别出NEUTRINO探测内核有10.22%更少的全局内存等待时间、32.8%更少的内存指令队列等待时间和17.32%更少的指令或寄存器缓存等待时间。为了进一步解决改进，我们检查了组装的机器码，特别是带有张量核心（HMMA）和异步复制（LDGSTS）的主循环。从截断和优化的机器码中，我们可以发现探测内核的HMMA指令更连续，有更多机会重用寄存器缓存（因为结果沿K累积）。而且，内存指令LDGSTS更分布式，可能导致内存系统压力更平衡。这个实验显示来自NEUTRINO探测的新指令可能导致汇编器【61，PTXAS, 2024】更好的执行流。虽然当前实验不全面（仅一个内核），我们相信这种异常加速和我们的分析突出了低级汇编器优化中仍有未探索的机会。

A Artifact Appendix：NEUTRINO的工件托管在GitHub（artifact分支），包含源代码、安装/收集/分析脚本、收集的跟踪，用于重现论文中的所有评估结果。我们还将工件评估打包为托管在Google Colab的Jupyter Notebook，提供一键结果重现，而无需本地运行时设置。此外，我们还维护NEUTRINO的在线文档，包含项目亮点、用户指南、路线图和评估功能的参考。工件声明：收集的跟踪和代码与论文相应描述相同。您可以使用我们提供的跟踪和分析代码复制所有主要结果（下面预期结果部分有细节）。我们还提供跟踪收集代码，以便您在自己的设备上收集自己的跟踪。值得注意的是，自定义跟踪，特别是DMAT，只能产生类似结果，由于硬件和运行时动态。

范围（元信息）：设计：NEUTRINO是一个GPU汇编探测工具，设计用于在运行时附加小代码片段（探测）到GPU内核，以暴露运行时执行细节（剖析）。系统：NEUTRINO系统由两部分组成，探测引擎附加代码片段，钩子驱动程序捕获运行时启动的GPU内核。源代码在GitHub可用，并可作为Python包安装。探测：NEUTRINO探测是小TOML文件，通过片段、数据模型、位置和回调定义剖析任务。论文中使用的探测在Github可用。输出：论文中的图1、图10 (A/B/C)、图11、图12、图13和表2。评估：我们将评估安排在线性结构的笔记本中，允许简单点击Runtime -> Run All执行。请参考README和每个Jupyter Notebook (Colab)中的说明。特殊要求：静态跟踪分析无特殊要求，对于动态跟踪收集，需要NVIDIA GPU，例如A100，以及PyTorch v2.5.0和CUTLASS v3.5.0的包含PTX的构建。磁盘空间要求：Google Colab评估不需要任何磁盘空间。对于本地评估，请安排3GB用于静态跟踪和至少10GB用于收集动态跟踪。实验时间：静态评估分析收集的跟踪在CPU上少于30分钟，动态评估在GPU上收集跟踪≈10小时。环境设置时间：静态评估下载跟踪≈2分钟。动态评估构建NEUTRINO≈15秒。设置PyTorch和CUTLASS可能≈3分钟。公开可用：是的。许可证：我们对系统源代码使用Apache License, Version 2.0，对论文中使用的探测使用CC BY 4.0许可证。

内容：NEUTRINO的工件评估安排在6部分，对应论文中的不同图或表：1. block_sched: §4.6 2. dmat: 图1、图11 3. kernel_overhead: 表3 4. max_mem: 图12 5. exposed_latency: 图13 6. warp_sched: 图14。我们将每个部分对应于Jupyter Notebook中的一个节。而且，每个评估提供两种模式：静态模式解析收集的跟踪，适合在无特殊硬件/软件要求的本地CPU-only设备上入门，以及动态模式在真实GPU启用环境中收集跟踪，适合完整评估。

托管和要求：如何访问：请从以下选择访问工件。Github：1. 静态评估: artifact/static.ipynb 2. 动态评估: artifact/dynamic.ipynb Colab：1. 静态评估 (使用CPU作为Runtime) 在这里。2. 动态评估 (使用GPU作为Runtime) 在这里。硬件要求：静态评估仅需要带有Python 3运行时的CPU机器。您不需要为静态评估安装NEUTRINO。对于动态评估，您需要安装CUDA驱动程序的NVIDIA GPU。请注意：1. 硬件选择将显著影响结果：• 请使用RTX3080用于所有DMAT图 (Part. 2)。• 请使用A100用于所有其余 (Part. 1, Part. 3-6)。2. 请确保同一GPU上没有其他工作负载执行。3. 请安排足够磁盘空间，至少10GB，用于动态跟踪收集。软件要求：NEUTRINO系统仅依赖GNU工具链 (gcc, file, git, nm)、CUDA工具链 (cuobjdump, ptxas) 和 Python 3.12 (pip, toml)。但评估工作负载需要包含PTX的PyTorch和CUTLASS构建。我们将依赖检查和安装打包在prepare_env.py中用于一键安装。安装：推荐使用虚拟环境，例如conda create -y -n ae_env python=3.11 && conda activate ae_env，当不使用Colab时用于安装。自动安装：我们提供助手脚本prepare_env.py，可以python prepare_env.py安装所有依赖和neutrino。Jupyter Notebook (也Google Colab) 使用这种方式。手动安装：1. 克隆Github仓库: git clone -b artifact https://github.com/open-neutrino/neutrino.git 2. 构建并安装neutrino: cd neutrino && python setup.py install && cd .. 3. 使用neutrino –help测试安装 请参考README文件了解安装包含PTX的PyTorch和CUTLASS的详细描述。

评估工作流：入门说明：入门说明，需要<30 min，包括：1. 所有基于收集跟踪重现论文中所有图和表的静态评估。2. 动态评估的block_sched部分 (第1部分)，收集并分析块调度跟踪。这一部件需要<1分钟，并帮助证明正确环境设置用于详细说明。您可以使用Colab执行评估脚本。要这样做，首先选择正确的Runtime (如上所述的CPU或GPU)，然后点击Colab网页顶部的Runtime按钮，并在下拉菜单中点击Run All按钮执行脚本。每个节 (几个块) 可以独立执行。统计或图将在执行完成时显示在每个单元格下方。如果您选择本地评估，请下载Jupyter Notebook并遵循与Colab执行说明相同的步骤。详细说明：详细说明覆盖动态评估的其余五个节。它们也被打包在Jupyter Notebook中 (也在Colab可用)，允许一键 (Runtime -> Run All) 执行和评估。每个节也可以独立执行。所以您可以在每个节后清理跟踪以节省磁盘空间。预期结果：收集跟踪的静态评估预期密切符合论文中呈现的图和表。为了节省磁盘空间，我们错误删除了这些结果的原始跟踪。而且因为这些结果捕获GPU的最细运行时动态，精确重现将不可能。我们后来的实验只能重现类似结果。请接受我们的道歉带来的不便，我们将更新修订论文以包括最新结果。自定义跟踪的动态评估预期产生类似结果，即类似数字或图形状。进一步评估：在完成以上评估并阅读文档后，我们推荐几种进一步评估方式：1. 测试您的工作负载：NEUTRINO支持大多数GPU工作负载。您可以导入您的GPU内核 (CUDA C++, Triton等) 并通过neutrino 测试它们。在此查看NEUTRINO支持的更多信息。2. 测试您的工作负载：首先阅读可编程探测指南，在.toml中编写并本地保存您的探测，并使用neutrino -p 应用它。3. 调查实现：NEUTRINO的实现组织良好，是理解GPU代码如何从OS分派的良好入口。您可以在neutrino/src/中找到钩子驱动程序的实现，以及在neutrino/probe/中的探测引擎。牧羊过程中添加的实验：在牧羊过程中，我们添加了几个更多实验来解决审稿人的技术评论。虽然AE不要求，我们也准备了重现代码：• 微基准 (表2): microbench.ipynb • 全局DMAT (图15): dmat_global.ipynb • DMAT减速 (图16): dmat_slowdown.ipynb • 异常加速 (附录K): speedup.ipynb

引用汇总

• [4]：GCN Assembly, 2024. 在背景章节中引用，描述AMD的汇编语言。
• [5]：GPU Performance API, 2024. 在引言中引用，作为硬件依赖剖析器示例。
• [6]：HCC Compiler for ROCm, 2024. 在背景中引用，作为AOT编译器。
• [7]：Radeon GPU Profiler, 2024. 在引言中引用，作为AMD硬件剖析器。
• [8]：Pytorch 2: Faster machine learning through dynamic python bytecode transformation and graph compilation, 2024, ASPLOS. 在背景和实用工具中多次引用。
• [10]：A study of replacement algorithms for a virtual-storage computer, 1966, IBM Systems Journal.
• [11]：An anomaly in space-time characteristics of certain programs running in a paging machine, 1969, Commun. ACM.
• [12]：Triangulating python performance issues with Scalene, 2023, OSDI.
• [13]：bpftrace: High-level tracing language for Linux, 2025.
• [14]：Cuda flux: A lightweight instruction profiler for cuda applications, 2019, PMBS.
• [18]：LDB: An efficient latency profiling tool for multithreaded applications, 2024, NSDI.
• [19]：FlashAttention-2: Faster attention with better parallelism and work partitioning, 2024, ICLR. 在引言和DMAT中多次引用。
• [20]：Flashattention: Fast and memory-efficient exact attention with io-awareness, 2022. 在引言和DMAT中引用。
• [21]：Low-overhead trace collection and profiling on gpu compute kernels, 2024, ACM Trans. Parallel Comput. 在方法中多次引用。
• [22]：Working set analytics, 2021, ACM Comput. Surv. 在引言和DMAT中多次引用。
• [24]：eBPF, 2024. 在多个章节中引用，作为灵感来源。
• [25]：Executable and Linkable Format, 2024.
• [26]：Unified Acceleration Foundation. oneAPI Programming Model, 2025.
• [27]：Compiling machine learning programs via high-level tracing, 2018.
• [28]：Simple and precise static analysis of untrusted linux kernel extensions, 2019, PLDI.
• [29]：The llama 3 herd of models, 2024.
• [30]：SPIR: The Standard IR for Parallel Compute and Graphics, 2024.
• [31]：Mamba: Linear-time sequence modeling with selective state spaces, 2024.
• [34]：Deep residual learning for image recognition, 2015.
• [39]：A novel renaming scheme to exploit value temporal locality through physical register reuse and unification, 1998, MICRO.
• [44]：LLVM: A compilation framework for lifelong program analysis & transformation, 2004, CGO.
• [45]：MLIR: Scaling compiler infrastructure for domain specific computation, 2021, CGO.
• [48]：Understanding performance of ebpf maps, 2024, eBPF.
• [53]：CUDA C++, 2024.
• [54]：CUDA Compiler Driver NVCC, 2024.
• [55]：Cuda profiling tools interface, 2024.
• [56]：Fat Binaries, 2024.
• [57]：Instruction Set Reference, 2024.
• [58]：NSight Compute System, 2024. 在评估中多次引用。
• [59]：NVTX: NVIDIA Tools Extension SDK, 2025.
• [60]：PTX ISA 8.5, 2024. 在方法中多次引用。
• [61]：PTXAS, 2024. 在评估中引用。
• [62]：TensorRT, 2024.
• [66]：Lightweight memory tracing, 2013, ATC.
• [68]：TOML, 2024.
• [69]：Self-attention does not need o(n2) memory, 2022.
• [71]：High-resolution image synthesis with latent diffusion models, 2022.
• [72]：Register integration: a simple and efficient implementation of squash reuse, 2000, MICRO.
• [74]：Using valgrind to detect undefined value errors with BitPrecision, 2005, ATC.
• [77]：Cudaadvisor: Llvm-based runtime profiling for modern gpus, 2018, CGO.
• [79]：Megatron-lm: Training multi-billion parameter language models using model parallelism, 2020.
• [80]：Flexible binary instrumentation framework to profile code running on intel gpus, 2022, ISPASS.
• [82]：Flexible software profiling of gpu architectures, 2015, ISCA.
• [84]：Validating the eBPF verifier via state embedding, 2024, OSDI.
• [86]：LLVM PTX Backend, 2024.
• [87]：PyTorch Profiler, 2024.
• [88]：PyTorch Vectorized Elementwise Kernel, 2024.
• [89]：CUTLASS, 2023.
• [90]：BPF Instruction Set Architecture (ISA), 2024.
• [91]：Triton: an intermediate language and compiler for tiled neural network computations, 2019, MAPL.
• [93]：Nvbit: A dynamic binary instrumentation framework for nvidia gpus, 2019, MICRO.
• [94]：Diffusers: State-of-the-art diffusion models, 2022.
• [95]：Roofline: an insightful visual performance model for multicore architectures, 2009, Commun. ACM.
• [96]：Transformers: State-of-the-art natural language processing, 2020, EMNLP.