GPU可观测性差距：为什么我们需要卸载到GPU设备上的eBPF

Yusheng Zheng=, Tong Yu=, Yiwei Yang=

eBPF 作为一种在内核态提供可编程能力的革命性技术，已经在 CPU 的可观测性、网络和安全领域取得了巨大成功。然而，对于日益重要的 GPU 计算领域，我们同样需要一种灵活、高效的观测手段。目前，大多数 GPU 性能分析工具都局限于在 CPU 侧通过驱动/用户态 API 或厂商的性能分析接口（CUPTI）进行观测，难以深入了解 GPU 内部的执行细节。为了解决这一问题，bpftime 通过其 CUDA/SYCL 附加实现提供 GPU 支持，使 eBPF 程序能够在 NVIDIA 和 AMD GPU 上的 GPU 内核内部执行。这将 eBPF 的可编程性、可观测性和定制能力带到 GPU 计算工作负载中。通过这种方式，它实现了对 GPU 应用程序的实时性能分析、调试和运行时扩展，无需修改源代码，填补了当前可观测性领域的空白。

注意： GPU 支持仍处于实验阶段。如有问题或建议，提交问题或联系我们。

问题：GPU 可观测性挑战

核心挑战：不透明的 GPU 执行模型

GPU 已成为机器学习、科学计算和高性能计算工作负载的主流加速器，但其 SIMT(单指令多线程)执行模型带来了重大的可观测性和扩展性挑战。现代 GPU 将数千个线程组织成线程束(通常 32 个线程)，这些线程束在流式多处理器(SM)上同步执行，内核从主机异步启动。这些线程在复杂的多级内存层次结构中导航，从快速但有限的每线程寄存器，到线程块内的共享内存/LDS,通过 L1/L2 缓存，再到较慢但丰富的设备内存，同时面临有限的抢占能力，这使得内核执行难以中断、检查或扩展。这种架构复杂性创造了丰富的性能特征，包括线程束分化、内存合并模式、存储体冲突和占用率变化，这些都直接影响吞吐量，但传统可观测性工具很难看清这些行为。理解和优化内核停滞、内存瓶颈、低效同步或次优 SM 利用率等问题，需要对 GPU 深处发生的执行流、内存访问模式和线程束间协调进行细粒度的可见性，以及动态注入自定义逻辑的能力。而这些正是现有工具难以以灵活、可编程的方式提供的能力。

时间线可见性差距：什么可以观察，什么不可以

为了具体说明这一挑战，考虑一个常见的调试场景。"我的 CUDA 应用程序需要 500ms 完成，但我不知道时间花在哪里。是内存传输、内核执行还是 API 开销？"答案关键取决于应用程序是使用同步还是异步 CUDA API,并揭示了 CPU 端可观测性的基本限制。

在同步模式下，CUDA API 调用会阻塞，直到 GPU 完成每个操作，在 CPU 和 GPU 时间线之间创建紧密耦合。考虑一个典型的工作流：分配设备内存、将数据传输到 GPU(主机到设备)、执行内核并等待完成。CPU 端性能分析器可以测量每个阻塞 API 调用的挂钟时间，提供有用的高级见解。例如，如果 cudaMemcpy() 需要 200μs 而 cudaDeviceSynchronize()(等待内核完成)只需要 115μs，开发人员可以快速识别数据传输占主导地位而非计算，表明 PCIe 瓶颈可能通过使用固定内存、更大的批量大小或异步传输来解决。然而，当开发人员问"内核同步需要 115μs，但为什么我的内核这么慢？是启动开销、内存停滞、线程束分化还是低 SM 利用率？"时，CPU 端工具遇到了根本性的障碍。115μs 的同步时间是一个不透明的聚合，混合了多个隐藏的 GPU 端阶段，包括内核启动开销(约 5μs 在 SM 上调度工作)、实际内核执行(约 100μs 在流式多处理器上计算)和清理(约 10μs 排空管道和释放资源)。即使完美计时同步 API 调用，CPU 端工具也无法区分内核性能差是由于过度的启动开销、计算低效、导致停滞的内存访问模式，还是 SM 利用不足。这些需要从内核执行期间的 GPU 内部获得线程束级执行、内存事务统计和每线程行为的可见性。我们需要更细粒度、可编程的 GPU 内部视角来理解内核执行期间发生的事情，而不仅仅是它何时开始和结束。

现代 CUDA 应用程序更多地使用异步 API(cudaMemcpyAsync()、带流的 cudaLaunchKernel())来通过重叠 CPU 工作和 GPU 执行来最大化硬件利用率。这引入了时间解耦，其中 API 调用在将工作排入流后立即返回，允许 CPU 继续执行，而 GPU 在后台按顺序处理操作。这打破了 CPU 端工具在同步模式下的可观测性。

考虑现在异步执行的相同工作流。开发人员将主机到设备传输(200μs)、内核执行(100μs)和设备到主机传输(150μs)排队，然后继续 CPU 工作，最后调用 cudaStreamSynchronize() 等待所有 GPU 操作完成。从 CPU 的角度来看，所有排队操作都在几微秒内返回，只有最后的同步点阻塞，报告总共 456μs(1 + 200 + 5 + 100 + 150 μs 的顺序 GPU 工作)。

CPU 时间线(传统工具看到的):
─────────────────────────────────────────────────────────────────────────────────
cudaMallocAsync() cudaMemcpyAsync() cudaLaunchKernel() cudaMemcpyAsync()      cudaStreamSync()
●─────●─────●─────●─────────────────────────────────────────────────────────────────●────
1μs  1μs  1μs  1μs         CPU 继续做其他工作...               456μs 等待
(分配)(H→D)(内核)(D→H)                                                       (全部完成)

GPU 时间线(实际执行 - 流中顺序):
─────────────────────────────────────────────────────────────────────────────────
◄─分配─►◄───────H→D DMA────────►◄─启动─►◄────内核执行────►◄────D→H DMA────►
│ ~1μs  │       200μs            │   5μs   │       100μs        │      150μs     │
┴────────┴────────────────────────┴─────────┴────────────────────┴────────────────┴─────
         ↑                                  ↑                                     ↑
    CPU 已继续                          GPU 仍在工作                    同步返回

在同步执行中，测量各个 API 调用持续时间允许开发人员识别传输或计算是否占主导地位。在异步模式下，这种能力完全消失，因为所有时间信息都折叠到同步点的单个 456μs 聚合中。问题"我的瓶颈是内存传输还是内核执行？"从 CPU 端变得无法回答。如果第一次传输由于未固定内存而需要两倍时间(400μs 而不是 200μs)，将所有后续操作延迟 200μs，开发人员只会看到总时间从 456μs 增加到 656μs，完全不知道是哪个操作导致了延迟、何时发生或是否传播到下游操作。

异步执行不仅隐藏了在同步模式下已经不可见的 GPU 内部细节(线程束分化、内存停滞、SM 利用率)，还消除了 CPU 工具以前可以提供的粗粒度阶段时序。开发人员失去了进行基本分类的能力。他们无法确定是将优化工作集中在内存传输、内核逻辑还是 API 使用模式上，除非(1)恢复到缓慢的同步执行进行调试(违背异步的目的)，或(2)添加需要重新编译的手动检测，只提供静态测量点。

像 LLM 服务这样的现代 GPU 应用程序通过高级优化技术进一步复杂化了这种情况。批处理策略组合多个操作以最大化吞吐量，就像管道一样，但使得更难识别哪些单个操作速度慢。持久内核驻留在 GPU 上处理多个工作批次，消除启动开销但模糊了阶段边界。具有流间复杂依赖关系的多流执行创建了复杂的执行图，其中来自不同流的操作不可预测地交错。每个线程块的共享内存使用限制了占用率并限制了并发线程束执行，创建了根据内核配置变化的微妙资源争用。这些优化显著提高了吞吐量，但使已经不透明的异步执行模型从 CPU 端观察和调试变得更加困难。

然而，即使我们有完美的 GPU 内部可见性显示高内存停滞周期，我们仍然无法确定根本原因：是线程束分化导致未合并访问？主机线程取消调度延迟异步内存复制？并发 RDMA 操作的 PCIe 拥塞？还是来自内核的按需页面迁移延迟？同样，当观察 SM 利用不足时，仅设备指标无法区分是网格太小、启动被用户空间互斥锁序列化，还是驱动程序在 ECC 错误或电源事件后进行了限流。

这个挑战在生产环境中变得严峻，那里尾延迟峰值间歇性发生。它们是由 GPU 缓存效应、主机生产者线程上的 cgroup 限流，还是另一个容器发出大型 DMA 传输的干扰引起的？仅设备工具报告"GPU 上发生了什么"，但不报告"为什么现在在这个异构系统中发生"。如果没有与用户空间行为(线程、系统调用、分配)、内核事件(页面错误、调度器上下文切换、块 I/O)和驱动程序决策(流依赖解析、内存注册、电源状态转换)的时间对齐关联，工程师必须缓慢迭代：观察 GPU 症状 → 猜测主机原因 → 使用临时检测重建 → 重复。这个反馈循环代价高昂，在延迟敏感的部署中通常不可行。为了解决生产环境中的这些挑战，业界开始转向持续性能分析（Continuous Profiling），其目标是在不显著影响性能的前提下，实现对线上服务的“永远在线”监控。然而，现有的 GPU 持续性能分析方案通常依赖于采样高层次指标（如 GPU 利用率或功耗），这无法解释“为什么”内核执行缓慢；或者依赖于定期运行重量级的供应商工具，但这会带来不可接受的性能开销。这些方案都无法在低开销、高保真和细粒度可见性之间取得平衡，也无法解决跨层关联的根本难题。

关键见解：有效的 GPU 可观测性和扩展性需要一个跨越异构计算堆栈多个层次的统一解决方案：从进行 CUDA API 调用的用户空间应用程序，通过管理设备资源的操作系统内核驱动程序，到在 GPU 硬件上执行的设备代码。传统工具在这些层之间是分散的，仅在 CPU-GPU 边界或单独在 GPU 内核内提供隔离的可见性，但缺乏理解一个层级的决策和事件如何影响另一个层级的性能和行为所需的跨层关联。

现有工具的局限性

现有的 GPU 跟踪和性能分析工具大致可分为三类，但每一类都有其固有的局限性，无法提供一个完整、统一的解决方案。

1. 仅限 CPU-GPU 边界的跟踪工具

许多跟踪工具仅在 CPU-GPU 边界运行，通过拦截 CUDA/SYCL 用户空间库调用(例如，通过对 libcuda.so 的 LD_PRELOAD 钩子)或在系统调用层检测内核驱动程序。

优势: 可以捕获主机端事件，如内核启动、内存传输和 API 时序。
局限性: 这种方法从根本上将 GPU 视为一个“黑盒”。它无法深入观察内核执行期间发生的具体事件，无法将性能问题与特定的线程束行为或内存停滞相关联，也无法根据设备内部的运行时条件自适应地修改行为。

2. 供应商特定的重量级性能分析器: 以 Nsight 为例

NVIDIA 的 Nsight 套件试图解决一些跨域可见性挑战。Nsight Systems 提供系统范围的时间线，在统一视图中显示 CPU 线程、CUDA API 调用、内核启动和内存传输。Nsight Compute 通过 CUPTI 计数器和基于重放的分析提供深度内核级微架构指标，具有引导式优化规则。这些工具可以将启动与 CPU 线程调度关联，并提供丰富的每内核停滞原因和内存指标。

然而，与统一的 eBPF 方法相比，Nsight 和类似的供应商工具存在基本限制。 - 封闭的事件模型：Nsight 提供具有固定事件集的封闭事件模型，在附加点没有任意可编程逻辑——您无法在不重新编译应用程序的情况下动态添加自定义检测，也无法编写过滤谓词，如"仅在内核执行超过 100ms 时收集数据"。 - 干扰性的性能分析会话：这些工具需要特殊的性能分析会话，通过计数器复用和重放机制扰乱工作负载行为，使它们不适合生产中的始终在线连续遥测，并且基于重放的收集会错过瞬态异常和罕见事件。 - 缺乏在设备端的过滤和聚合：Nsight 缺乏原地过滤和聚合，强制导出所有原始数据然后进行后处理，这从具有大量异步管道的大型应用程序创建多 GB 跟踪，没有可编程的自适应响应来根据观察到的状态动态更改采样逻辑。 - 有限的系统集成：Nsight 无法在不重启的情况下将动态探针附加到持久内核，缺乏与 Linux eBPF 基础设施(kprobes/uprobes/tracepoints)的集成，并且无法跨 CPU 和 GPU 检测共享数据结构(映射)，这使得非常难以拼接因果链，如页面错误(主机)→ 延迟启动排队 → 线程束停滞峰值。 - 供应商锁定：这些是 NVIDIA 专用工具，在异构系统中跨 AMD、Intel 或其他加速器的供应商中立部署没有明确路径。

在实践中，开发人员面临大型跟踪的迭代根本原因分析速度变慢，错过在性能分析开销下不重现的生产问题，并且无法将 GPU 事件与现有生产可观测性堆栈(perf、bpftrace、自定义 eBPF 代理)关联，而无需复杂的模式切换到特殊的"性能分析会话"。

3. 用于细粒度分析的研究工具与接口

当需要比 Nsight 更深层次的可见性时，业界探索了基于二进制插桩或更底层接口的工具，如 NVIDIA CUPTI、NVBit 和 NEUTRINO。

CUPTI (CUDA Profiling Tools Interface): 作为一个成熟的接口，CUPTI 非常适合获取内核级的宏观指标（如开始/结束时间）和硬件性能计数器。项目如 xpu-perf 已证明其在关联 CPU-GPU 数据方面的有效性。然而，当需要理解“为什么”一个内核缓慢时，CUPTI 提供的宏观指标往往不足。
二进制插桩工具 (NVBit, NEUTRINO): 这些工具通过在汇编或 PTX（NVIDIA GPU 的中间语言）层面进行插桩，实现了对指令级的细粒度观测。例如，在同一时间出现的 NEUTRINO 通过汇编层探测来获取数据，但它通常需要直接使用汇编语言进行编程，这不仅复杂，而且缺乏 eBPF 所提供的安全性和可移植性，它们也通常是独立于 CPU 分析器的，难以提供统一的跨层、跨多设备的可见性，关联事件逻辑并处理时钟漂移非常困难，在关联事件的过程中也可能造成多次拷贝，导致额外的性能开销。Neutrino 并非为持续在线（always-on）监控而设计，它也是基于会话（session-based）的，会生成大量信息并等待后续处理。

总而言之，尽管现有工具在某些方面很强大，但它们要么过于宏观，要么过于笨重和孤立。开发者面临着在迭代速度、生产安全和问题定位深度之间的艰难权衡，并且无法将 GPU 事件与现有的 CPU 可观测性堆栈（如 perf、bpftrace）轻松关联。

解决方案：将 eBPF 扩展到 GPU

为了克服现有工具的局限性，我们需要一种能够统一 CPU 和 GPU，并提供可编程、低开销观测能力的解决方案。eBPF 技术及其在 bpftime 项目中的实现，为此提供了可能。

为什么选择 eBPF？

要理解为什么 eBPF 是应对这一挑战的正确工具，查看它在 CPU 世界的影响是有帮助的。eBPF(扩展的伯克利包过滤器)是 Linux 内核中的一项革命性技术，允许沙箱程序动态加载以安全地扩展内核能力。在 CPU 端，eBPF 已成为现代可观测性、网络和安全的基石，因为它独特地结合了可编程性、安全性和性能。它使开发人员能够将自定义逻辑附加到数千个钩子点，以最小的开销收集深度、定制的遥测。bpftime 背后的核心思想是将这种相同的变革力量带到传统上不透明的 GPU 计算世界。

通过在 GPU 内核内原生运行 eBPF 程序，bpftime 在整个堆栈中提供安全、可编程、统一的可观测性和扩展性。 - 统一的跨层可观测性：该架构将 CPU 和 GPU 探针视为统一控制平面中的对等体。共享的 BPF 映射和环形缓冲区实现直接数据交换，动态检测无需重新编译或重启即可工作。与现有 eBPF 基础设施(perf、bpftrace、自定义代理)的集成不需要模式切换。开发人员可以同时通过 uprobes 跟踪 CPU 端 CUDA API 调用、通过 kprobes 跟踪内核驱动程序交互、通过 CUDA 探针跟踪 GPU 端内核执行，所有这些都使用相同的 eBPF 工具链，并关联跨主机-设备边界的事件。示例问题现在变得可以回答："T+50μs 的 CPU 系统调用延迟是否导致 T+150μs 的 GPU 内核停滞？"或"哪些 CPU 线程正在启动表现出高线程束分化的内核？"这种跨层可见性实现了跨越整个异构执行堆栈的根本原因分析，从用户空间应用程序逻辑通过内核驱动程序到 GPU 硬件行为，而无需离开生产可观测性工作流。 - 低开销的生产环境监控：与基于会话的性能分析器不同，它通过动态加载/卸载探针和设备端谓词过滤来实现始终在线的生产监控，以减少开销。 - 恢复异步可见性：它通过每阶段时间戳(H→D 在 T+200μs，内核在 T+206μs，D→H 在 T+456μs)恢复异步模式可见性，通过纳秒粒度的线程束执行和内存模式遥测暴露 GPU 内部细节，并关联 CPU 和 GPU 事件，而无需传统单独性能分析器的重量级开销。

`bpftime`：在 GPU 上原生运行 eBPF

bpftime 的方法通过将 eBPF 的可编程性和定制模型直接扩展到 GPU 执行上下文来弥合这一差距，使 eBPF 程序能够在 GPU 内核内与应用程序工作负载一起原生运行。它采用的 PTX 注入技术，通过向 NVIDIA GPU 的中间汇编语言（PTX）中动态注入 eBPF 程序，实现了对 GPU 线程的直接挂钩（hook）。

这种方法让我们能够获取极其细粒度的、内核内部的运行时信息，而这些信息是 CUPTI 等高层 API 难以企及的。例如，通过 bpftime 我们可以：

跟踪单个线程或线程块（Thread Block）的内存访问模式：精确了解访存指令如何执行，是否存在非合并访问等问题。
观察 SM（Streaming Multiprocessor）的调度行为：了解线程束（Warp）如何在 SM 上被调度和执行。
分析内核内部的控制流：识别导致线程束分化（Warp Divergence）的具体分支，并量化其影响。

bpftime 的 PTX 注入并非要替代 CUPTI，而是作为其能力的延伸和补充。当开发者需要深入到 GPU 内核执行的微观层面，去定位那些由线程行为、内存访问或调度策略引发的复杂问题时，bpftime 填补了当前工具链在细粒度、可编程观测方面的空白。相比其他二进制插桩工具，bpftime 采用的 eBPF 方案具备几个独特优势：

更高的安全性：eBPF 的验证器（Verifier）提供了一个安全的沙箱环境，防止了不安全的指令执行。
便捷的跨层关联：卸载到 GPU 上的 eBPF 本质上是 CPU 上 eBPF 的一个扩展，能够方便地与 CPU 端用户态或内核态的 eBPF 进行统一插桩和编写，更容易实现跨 CPU-GPU-网卡、跨用户态、内核态的事件关联。
简化的编程模型：eBPF 使用受限的 C 语言进行编程，屏蔽了底层汇编的复杂细节，使得开发者可以更高效地编写功能强大的观测和分析程序。

工具类型	代表工具	优势	局限	适用场景
厂商性能分析器	NVIDIA Nsight, AMD ROCProfiler	全面的硬件指标，可视化强	重量级，需要特殊环境，跨平台困难	开发阶段深度优化
高级 API	CUPTI, ROCTracer	稳定接口，支持追踪和采样	粒度粗，无法观测内核内部	应用级性能监控
二进制插桩	NVBit, SASSI, NEUTRINO	细粒度观测，指令级控制	性能开销大（3-10×），编程受限	离线深度分析
eBPF 扩展	bpftime	统一 CPU/GPU 观测，低开销，安全可编程，跨平台	需要运行时支持，功能仍在发展	生产环境实时监控，跨层关联分析

`bpftime` 的架构与优势

该系统定义了一套全面的 GPU 端附加点，反映了 CPU 端 kprobes/uprobes 的灵活性。开发人员可以检测 CUDA/SYCL 设备函数入口和出口点(类似于函数探针)、线程块生命周期事件(块开始/结束)、同步原语(屏障、原子操作)、内存操作(加载、存储、传输)和流/事件操作。用受限 C 编写的 eBPF 程序通过 LLVM 编译成设备原生字节码(用于 NVIDIA GPU 的 PTX(并行线程执行)汇编或用于 AMD/Intel 的 SPIR-V)，并通过二进制检测在运行时动态注入到目标内核中，无需修改源代码或重新编译。

运行时在 GPU 上提供完整的 eBPF 执行环境，包括(1)安全验证器，以确保在 SIMT 上下文中的有界执行和内存安全，(2)一组丰富的 GPU 感知辅助函数，用于访问线程/块/网格上下文、计时、同步和格式化输出，(3)专用的 BPF 映射类型，位于 GPU 内存中，用于高吞吐量每线程数据收集(GPU 数组映射)和事件流(GPU 环形缓冲区映射)，以及(4)主机-GPU 通信协议，使用共享内存和自旋锁，在需要时安全地调用主机端辅助函数。

这种架构不仅可以以纳秒粒度收集细粒度遥测(每线程束时序、内存访问模式、控制流分化)，还可以根据运行时条件自适应地修改内核行为，构建自定义扩展和优化，并将 GPU 可观测性与现有 CPU 端 eBPF 程序统一到单个分析管道中，同时保持生产就绪的开销特性。这使得：

比 NVBit 等工具快 3-10 倍的检测性能
供应商中立设计，适用于 NVIDIA、AMD 和 Intel GPU
统一可观测性和控制，与 Linux 内核 eBPF 程序(kprobes、uprobes)一起
细粒度性能分析和运行时定制，在线程束或指令级别
自适应 GPU 内核内存优化和跨 SM 的可编程调度
动态扩展，用于 GPU 工作负载，无需重新编译
通过利用 GPU 计算能力加速 eBPF 应用程序

该架构旨在实现四个核心目标：(1)提供统一的基于 eBPF 的接口，可无缝跨用户空间、内核、来自不同供应商的多个 CPU 和 GPU 上下文工作，(2)实现动态、运行时检测，无需修改源代码或重新编译，(3)在 GPU 硬件和 SIMT 执行模型的约束内保持安全高效的执行。(4)依赖更少且易于部署，建立在现有 CUDA/SYCL/OpenGL 运行时之上，无需自定义内核驱动程序、固件修改或像记录和重放系统这样的重量级运行时。

但空谈架构是一回事，亲眼见证其性能表现则是另一回事。它到底有多快？我们将在下一篇博客中深入探讨性能基准测试，敬请期待！

架构

CUDA 附加管道

GPU 支持构建为动态将 eBPF 程序注入 GPU 内核的检测管道。关键组件是：

CUDA/OpenCL 运行时钩子：使用 LD_PRELOAD,bpftime 拦截对 CUDA/SYCL 运行时库的调用。这允许它控制内核启动和其他 GPU 相关操作。
eBPF 到 PTX/SPIR-V JIT 编译：当内核启动时，bpftime 获取用于 GPU 探针的 eBPF 字节码，并即时(JIT)编译成目标 GPU 的指令集架构——用于 NVIDIA 的 PTX 或用于 AMD/Intel 的 SPIR-V。
二进制检测和注入：编译后的 eBPF 代码在加载到 GPU 之前注入到目标内核的二进制文件(例如 PTX 代码)中。这种运行时修改允许 eBPF 程序在内核上下文中原生执行。
辅助函数蹦床:bpftime 提供一组可从 GPU 访问的 eBPF 辅助函数。这些辅助函数实现为蹦床，执行诸如访问映射、获取时间戳或通过环形缓冲区提交数据等任务。
共享数据结构:BPF 映射和环形缓冲区通过共享内存(固定主机内存)或 CPU 和 GPU 都可以访问的设备内存实现，实现主机和设备之间的高效数据交换。

┌─────────────────────────────────────────────────────────────────┐
│                      应用程序进程                                │
│              (LD_PRELOAD)              JIT 编译                 │
│  ┌──────────────┐     ┌──────────────┐     ┌──────────────┐     │
│  │     PTX/     │───▶│  bpftime     │────▶│  GPU 内核    │     │
│  │    SPIR-V    │     │  运行时       │     │  带 eBPF     │     │
│  └──────────────┘     └──────────────┘     └──────────────┘     │
│                     管理     │                      │           │
│                              ▼                      ▼           │
│                       ┌────────────────────────────────────┐    │
│                       │ 共享映射                            │    │
│                       │  (主机-GPU)                        │    │
│                       └────────────────────────────────────┘    │
└─────────────────────────────────────────────────────────────────┘

示例

我们通过几个 bcc 风格的工具演示 GPU eBPF 能力：

kernelretsnoop - 每线程退出时间戳跟踪器

附加到 CUDA 内核退出并记录每个 GPU 线程完成执行时的精确纳秒时间戳。这揭示了传统性能分析器看不到的线程分化、内存访问模式和线程束调度问题。

注意：这些示例中使用的 kprobe/kretprobe 命名约定是一个占位符，用于保持与 Linux 内核 eBPF 的概念相似性。在 bpftime 中，这些探针附加到设备端 GPU 内核函数，而不是 Linux 内核。这个命名可能会在未来修订，以更好地反映它们的范围。

用例：您发现内核比预期慢。kernelretsnoop 显示每个线程束中的线程 31 比线程 0-30 晚完成 750ns，暴露了导致分化的边界条件。您重构以消除分支，现在所有线程在几纳秒内完成。

// eBPF 程序在 GPU 内核退出时运行
SEC("kretprobe/_Z9vectorAddPKfS0_Pf")
int ret__cuda() {
    u64 tid_x, tid_y, tid_z;
    bpf_get_thread_idx(&tid_x, &tid_y, &tid_z);  // 我是哪个线程?
    u64 ts = bpf_get_globaltimer();               // 我何时完成?

    // 写入环形缓冲区供用户空间分析
    bpf_perf_event_output(ctx, &events, 0, &data, sizeof(data));
}

threadhist - 线程执行计数直方图

使用 GPU 数组映射来计算每个线程执行多少次。检测工作负载不平衡，其中一些线程做的工作远多于其他线程，浪费 GPU 计算能力。

用例：您的网格步进循环使用 5 个线程处理 100 万个元素。您期望工作负载平衡，但 threadhist 显示线程 4 的执行次数只有线程 0-3 的 75%。边界元素划分不均匀，线程 4 空闲而其他线程工作。您调整分布并实现平衡执行。

// eBPF 程序在 GPU 内核退出时运行
SEC("kretprobe/_Z9vectorAddPKfS0_Pf")
int ret__cuda() {
    u64 tid_x, tid_y, tid_z;
    bpf_get_thread_idx(&tid_x, &tid_y, &tid_z);

    // GPU 数组映射中的每线程计数器
    u64 *count = bpf_map_lookup_elem(&thread_counts, &tid_x);
    if (count) {
        __atomic_add_fetch(count, 1, __ATOMIC_SEQ_CST);  // 线程 N 又执行了一次
    }
}

launchlate - 内核启动延迟分析器

测量 CPU 上的 cudaLaunchKernel() 和 GPU 上实际内核执行之间的时间。揭示隐藏的队列延迟、流依赖关系和调度开销，这些使快速内核在生产中变慢。

用例：您的每个内核执行 100μs，但用户报告 50ms 延迟。launchlate 显示每个内核有 200-500μs 的启动延迟，因为每个都等待前一个和内存传输完成。总时间是 5ms，而不是 1ms。您切换到 CUDA 图，批处理所有启动，延迟降至 1.2ms。

BPF_MAP_DEF(BPF_MAP_TYPE_ARRAY, launch_time);

// CPU 端 uprobe 捕获启动时间
SEC("uprobe/app:cudaLaunchKernel")
int uprobe_launch(struct pt_regs *ctx) {
    u64 ts_cpu = bpf_ktime_get_ns();  // CPU 何时请求启动?
    bpf_map_update_elem(&launch_time, &key, &ts_cpu, BPF_ANY);
}

// GPU 端 kprobe 捕获执行开始
SEC("kprobe/_Z9vectorAddPKfS0_Pf")
int kprobe_exec() {
    u64 ts_gpu = bpf_get_globaltimer();  // GPU 实际何时开始?
    u64 *ts_cpu = bpf_map_lookup_elem(&launch_time, &key);

    u64 latency = ts_gpu - *ts_cpu;  // 内核在队列中等待多长时间?
    u32 bin = get_hist_bin(latency);
    // 更新直方图...
}

其他示例

cuda-counter：带时序测量的基本探针/返回探针
mem_trace：内存访问模式跟踪和分析
directly_run_on_gpu：直接在 GPU 上运行 eBPF 程序，无需附加到内核

关键组件

CUDA 运行时钩子：使用基于 Frida 的动态检测拦截 CUDA API 调用
PTX/ SPIR-V 修改：将 eBPF 字节码转换为 PTX（并行线程执行）或 SPIR-V 汇编并注入 GPU 内核
辅助蹦床：为映射操作、计时和上下文访问提供 GPU 可访问的辅助函数
主机-GPU 通信：通过固定共享内存实现从 GPU 到主机的同步调用

参考资料

Share on Share on