将 eBPF 扩展到 GPU 设备上下文：bpftime 的 GPU 可观测性与调度实践- 大数跨境

首页

将 eBPF 扩展到 GPU 设备上下文：bpftime 的 GPU 可观测性与调度实践

GTOC

2025-12-09

导读：本文为作者在OSDT 2025-第17届开源开发工具大会上分享的 bpftime 的 GPU 可观测性与调度实践的文稿。

本文为作者在OSDT 2025-第17届开源开发工具大会上分享的 bpftime 的 GPU 可观测性与调度实践的文稿，首发于知乎（文末阅读原文可访问）。

另外，特别感谢本项目的 maintainer，包括但不仅限于@云微，Officeyutong 和 @vickieGPT(Yiwei Yang) 。他们对本文亦有贡献。

在当前的大规模计算场景中，GPU 已经从通用加速设备演变为整机中最关键的算力资源之一。

在大模型训练、在线推理、图计算、向量检索等工作负载下，整体性能和稳定性往往直接由 GPU 的利用效率和运行状态决定。

如果希望系统地回答以下问题：

给定一块 GPU，在某个时间区间内具体执行了哪些 kernel、耗时如何；某些 step 尾延迟升高时，瓶颈出现在数据准备、计算还是通信阶段；多机训练中，哪些节点在迭代过程里持续或间歇性地变慢。

仅依赖传统 CPU 侧监控和现有 GPU profiler 通常难以获得足够细粒度、且易于与 CPU 事件对齐的数据。在这种观测体系下，GPU 在多数情况下仍然缺乏足够的可观测性。

本文围绕我们在 bpftime 中扩展 GPU 观测与策略控制能力所做的工作，介绍以下几个方面：

从用户态 eBPF 框架 bpftime 出发，如何把 eBPF 的可编程观测能力扩展到 GPU 设备端；如何在此基础上构建统一的 CPU–GPU 观测视角；

这套能力在大模型训练慢节点检测、内存管理和调度策略等场景中的初步尝试。

GPU 性能分析背景

1.1 eBPF 在 CPU 侧的成功

随着 Linux 内核版本的持续升级，eBPF 已成为 Linux 内核的核心可观测性基础设施：

· Linux 3.15+ 引入 eBPF
· Linux 4.x 持续增强 kprobe/uprobe/tracepoint
· Linux 5.x+ 支持 BTF、CO-RE、更多 helper

基于 eBPF，运维和开发团队可以：

持续采样 CPU 栈，生成进程级、容器级的火焰图；按需挂载 kprobe/tracepoint，捕获内核里的调度、IO、内存事件；在用户态函数上用 uprobe 做细粒度的调用跟踪。

对传统的 Web 服务、数据库、中间件来说，从内核到应用栈，可以通过统一的 eBPF 接口获取较完整的性能与行为信息。

1.2 GPU 侧的观测困境

在 AI 场景下，情况明显不同。当我们将这套能力迁移到大模型训练和推理集群时，可以观察到：

在这种状态下，典型问题往往只能做出概率性的推断，例如：

推理尾延迟升高”可能”与某个 kernel 相关；训练 step time 偶尔变长”可能”由 all-reduce 排队导致。

这类结论难以直接用于驱动可靠的调优和自动化策略。

1.3 现有 GPU 观测工具的局限

现有 GPU 观测工具大致可以分为两类：

CPU–GPU 边界上的跟踪工具。通过拦截 CUDA /ROCm 用户态库或在内核驱动层插入探针，可以观测：

进程何时调用 cudaLaunchKernel 等 API；API 调用耗时、传输数据量等信息。

这类工具通常将 GPU 视为整体设备，无法深入到 kernel 内部执行。

GPU 厂商的性能分析工具（CUPTI、Nsight、GTPin、NVBit、Neutrino 等）这类工具能在 GPU 中插桩，获取 warp 级甚至指令级数据，但存在一些限制：

与 Linux 内核 eBPF 生态相对独立，数据管道和分析流程不统一与通用监控平台的集成成本较高；细粒度模式下开销较大（NVBit 30-45%，CUPTI 15-25%），难以长期部署在生产环境。

在这两类工具之间存在明显空白：要么获得设备端细节，代价是较高开销和与现有观测体系割裂；要么在现有 eBPF 体系下保持良好一致性，但 GPU 内部行为仍缺乏细粒度可见性。

GPU 可观测性为何困难？

GPU 之所以难以观测，并不仅仅是”缺少工具”，更主要的原因在于其执行模型和内存体系与 CPU 有本质差异。

2.1 SIMT 执行模型

以 NVIDIA 为例，今天的主流 GPU 基本都遵循 SIMT（Single Instruction, Multiple Thread）模型：

线程层次结构：Thread → Warp (32) → Block → Grid → SM

在这样的执行环境下，较小的代码或数据模式变化就可能导致显著的性能差异：

Warp Divergence：warp 内分支条件稍有变化，就产生明显的线程束分化；

访存合并失败：访存对齐或布局改变，引起访存合并失败、bank conflict 或 cache 行为变化；

Occupancy 下降：共享内存使用方式不同，可能导致 occupancy 下降；

时间线复杂化：持久 kernel、异步流、多 stream 之间的依赖关系，使时间线结构更加复杂。

2.2 复杂的内存层次结构

GPU 的内存层次比 CPU 更复杂，性能对访存模式极其敏感：

2.3 异步执行：CPU 视角的天然盲区

很多人第一次调 CUDA 性能，是在 CPU 端给 cudaMemcpy、cudaLaunchKernel、cudaDeviceSynchronize 打点。

在同步 API 模式下，这种办法多少还能用：每个 API 调用阻塞，CPU 侧的时间间隔和 GPU 上的执行时间勉强能对上。

问题在于，真实的深度学习训练与推理服务几乎不会采用这种同步方式：

用 cudaMemcpyAsync + 多 stream 做传输和计算重叠；使用持久 kernel 避免频繁 launch；甚至直接交由框架（vLLM、TensorRT 等）统一组织 pipeline。

在异步模式下，CPU 端能观测到的主要是：

一系列几乎瞬间返回的 API 调用（工作被排队到 GPU）；若干 cudaStreamSynchronize 或事件等待（聚合了排队与执行时间）。

CPU 时间线 (传统工具看到的):──────────────────────────────────────────────────────────────── LaunchAsync    LaunchAsync    LaunchAsync       Sync────●─────────────●─────────────●────────────────●────────────   1μs           1μs           1μs             455μs                                            (所有时间折叠到这里)GPU 时间线 (实际执行):────────────────────────────────────────────────────────────────    ◄─排队─►◄──Kernel 1──►◄──Kernel 2──►◄──Kernel 3──►────────────────────────────────────────────────────────────────

此时，CPU 视角下的时间线信息不足以区分：

排队时间和 GPU 实际执行时间的比例；单个 kernel 自身执行特征与其前后数据传输、调度行为之间的关系。

这也是 CPU 侧监控在 GPU 场景中的典型限制：可以发现问题存在，但难以给出足够精确的空间位置和原因定位。

目标：统一的 GPU 可观测与策略控制

3.1 设计目标

基于已有的 bpftime 能力，我们希望：

不额外为 GPU 再造一套独立的 profiling 框架；而是将 GPU 纳入与 CPU、内核一致的 eBPF 编程与执行框架中。

3.2 核心挑战

将 eBPF 扩展到 GPU 面临以下挑战：

挑战一：GPU 驱动闭源，缺乏扩展点。

GPU 驱动把硬件控制逻辑和策略逻辑紧密耦合在一起，没有类似 CPU 内核那样专门为 eBPF 预留的 attach 点。

挑战二：eBPF 语义和 GPU SIMT 执行模型不匹配。

eBPF 默认假设单线程、顺序执行、有强内存一致性。GPU 是 SIMT 模型：

一个 warp 32 线程要走相对统一的控制流；如果每个线程都独自跑 eBPF，很容易造成 warp divergence 和性能问题。

挑战三：缺乏高效的 host–device 共享状态机制。

CPU 内存大但延迟高，GPU 显存小但带宽高，二者通过 PCIe / NVLink 交互，时序差异大。策略往往需要同时访问 host 和 device 的状态。

3.3 bpftime for GPU 的定位

bpftime for GPU 正是在这一设计目标下逐步形成的方案，它要解决三件事：

把 eBPF 程序变成 GPU 能跑的代码（eBPF → PTX）；

在不改业务代码的前提下，把这些代码”缝”到已有 kernel 里（PTX 解包与插桩）；

让 GPU 设备端的 eBPF 与 CPU 端的 eBPF 共享同一套控制面和数据面（map 与 helper）。

bpftime 用户态 eBPF 运行时架构

在介绍 GPU 扩展之前，有必要先了解 bpftime 本身的架构。

bpftime 是一个用户态 eBPF 运行时，它的核心思路是：不依赖内核 eBPF 子系统，而是在用户空间完整实现 eBPF 的编译、验证和执行能力。

从架构图可以看到，bpftime 分为三个主要层次：

Userspace 层

左侧是 eBPF 程序源码和现有工具链（clang/llvm/bpftrace）；

中间是 bpftime 的核心组件：bpftime-syscall.so 负责拦截 bpf 系统调用，verifier 做程序验证，AOT compiler 做编译。

右侧是 Target Process，通过 uprobe 和 function hook 实现对用户态程序的插桩。

eBPF userspace applications 层

这里展示了两条并行路径：普通的 Target Process 和 Target CUDA Process。

对于 GPU 场景，bpftime-agent.so 负责 PTX injection，将 eBPF 程序编译为 GPU 可执行的 PTX 代码。

通过 Shared memory 和 eBPF maps 实现 CPU 与 GPU 之间的数据共享。

Kernel space 层

与内核 eBPF 子系统的交互：verifier、JIT compiler、eBPF maps、eBPF programs；支持 kprobe、tracepoint、socket 等多种 attach 类型。

这种架构带来几个关键优势：

无需 root 权限：用户态 eBPF 不需要加载内核模块；

跨平台：不依赖特定内核版本，可以在容器、虚拟机等环境中运行；

可扩展：正是这种架构使得 GPU 扩展成为可能——只需要在用户态增加对 CUDA Runtime 的 hook 和 PTX 编译能力。

透明插装：从 fatbin+ 到 PTX 注入

我们首先要解决的问题是：如何在不修改业务代码的前提下，对现有 kernel 进行透明插桩。

5.1 CUDA 程序编译流程与 fatbin

CUDA 程序包含两部分代码：

CUDA 提供了一种叫 fatbin 的打包格式，可以把针对不同架构的 PTX、SASS 统统打在一个包里。

在应用启动时，CUDA 运行时会调用 __cudaRegisterFatBinary 之类的函数，把这个包注册给驱动。

fatbin 结构：

+----------------+
| FatbinHeader   |  magic, version, header_size, size
+----------------+
| section header |  kind=1 (PTX), size
+----------------+
| payload        |  PTX assembly (文本)
+----------------+
| section header |  kind=2 (ELF), sm_62
+----------------+
| payload        |  Cubin sm_62 (二进制)
+----------------+
| section header |  kind=2 (ELF), sm_70
+----------------+
| payload        |  Cubin sm_70 (二进制)
+----------------+

5.2 Hook 机制：拦截 fatbin 注册

bpftime 做的第一件事，就是在用户态 hook 住 fatbin 注册过程：

这样，我们就获得了一个不需要项目改代码、不需要重新编译的插桩入口。

5.3 eBPF → PTX：让 probe 程序真正跑进 GPU

有了要插桩的 PTX，接下来需要把 eBPF 程序也编译成 PTX 格式，这样才能注入到 GPU kernel 中执行。

得益于 LLVM 的架构设计，eBPF 和 PTX 都是 LLVM 支持的编译目标。我们可以把 eBPF 程序先编译成 LLVM IR（与目标无关的中间表示），然后切换到 NVPTX 后端，生成 GPU 可执行的 PTX 代码。

这个过程我们内部叫 llvmbpf，本质上就是一个 eBPF 到 PTX 的编译器。在转换过程中，eBPF 的 r0-r10 寄存器映射到 PTX 虚拟寄存器，BPF_CALL helper 调用变成 PTX 函数调用，内存访问指令转换为 ld.global / st.global 等。

生成的 PTX 代码大致如下：

.visible .func __retprobe_func_xxx()
{
    // 获取线程索引
    mov.u32 %r1, %tid.x;
    mov.u32 %r2, %tid.y;
    mov.u32 %r3, %tid.z;

    // 读取 GPU 全局计时器 (纳秒级精度)
    mov.u64 %rd1, %globaltimer;

    // 调用 helper 更新 map
    call.uni _bpf_helper_ext_map_update;

    ret;
}

5.4 PTX 注入与插桩位置

有了原始 PTX 和 probe 的 PTX 函数，下一步就是把它们合并在一起。ptxpass 会解析原始 PTX 的结构，在指定位置插入对 probe 函数的调用（call.uni），然后声明必要的符号，重新打包后交给 CUDA 驱动加载。

默认的插桩位置包括：函数入口（用于调用计数、参数采样）、ret 指令前（用于耗时统计）、以及 load/store 指令前（用于访存跟踪）。以一个简单的 kernel 为例，原始 PTX 可能是：

.entry _Z9myKernelPf(...){    ld.global.f32 %f1, [%rd1];    mul.f32 %f2, %f1, %f3;    st.global.f32 [%rd2], %f2;    ret;}

注入后变成：

.entry _Z9myKernelPf(...){    call.uni __uprobe_func_myKernel;    // 入口 probe
    ld.global.f32 %f1, [%rd1];    mul.f32 %f2, %f1, %f3;    st.global.f32 [%rd2], %f2;
    call.uni __uretprobe_func_myKernel; // 出口 probe    ret;}

修改完成后，我们把这个带 probe 的 PTX 模块重新交给 CUDA 驱动加载，并借助 hook 的 cudaRegisterFunction / cudaLaunchKernel 等 API，确保后续真正运行的是修改过的版本。

对上层业务来说，这个过程是透明的：不需要改一行 CUDA 代码、不需要重新链接，只要通过 LD_PRELOAD 把 bpftime 的 agent 拉起来即可。

5.5 任意位置插桩：基于 Stub 的挂载点

上面介绍的插桩方式都是在 PTX 层面自动识别的固定位置。

但在实际调试和优化中，我们经常需要在 kernel 内部的特定位置进行观测，比如在一个复杂循环的每次迭代中采样，或者在多阶段计算的不同阶段之间测量耗时。这些需求无法通过固定位置的自动插桩来满足。

为此，bpftime 提供了基于 stub 的任意位置插桩能力。开发者在 CUDA kernel 中定义一个空的 device 函数作为”桩”，然后在需要观测的位置调用它：

__device__ __noinline__ void __bpftime_cuda__kernel_trace() {}__global__ void complexKernel(float *data, int n) {    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float val = preprocess(data[idx]);    __bpftime_cuda__kernel_trace();  // 插桩点 1：预处理完成
    for (int i = 0; i < n; i++) {        val = compute(val, i);        if (i % 100 == 0) {            __bpftime_cuda__kernel_trace();  // 插桩点 2：每 100 次迭代采样        }    }
    data[idx] = val;    __bpftime_cuda__kernel_trace();  // 插桩点 3：写回完成}

在编译时，这个 stub 函数会被编译成 PTX 中的一个 call 指令。bpftime 的 ptxpass 在处理 PTX 时，会识别所有对 stub 函数的调用，并将其重定向到 eBPF 生成的 probe 函数。

如果没有找到 stub 调用，则回退到默认的入口插桩。stub 名称可以通过配置文件自定义。

这个能力使得 bpftime 不仅可以做黑盒观测（不修改代码、只看入口出口），还可以做白盒调试（在代码中精确标记观测点）。生产环境可以使用透明的固定位置插桩，开发调试时则可以使用 stub 插桩获得更精细的观测粒度。

从更广的视角看，stub 机制实际上是在 GPU kernel 内部建立了一套可编程的观测点，类似于 CPU 侧的 tracepoint——在代码中预埋观测点，运行时决定是否启用以及执行什么逻辑。

开发者可以把 stub 调用理解为一种”GPU 端的 tracepoint”——在代码中预埋观测点，运行时决定是否启用以及执行什么逻辑。

GPU 上的 Map 与 Helper：设备端状态管理

在 eBPF 的世界里，map 是一个非常核心的概念：它既承载状态，又承担着用户态/内核态之间的通信通道。要把 eBPF 搬到 GPU 上，这个抽象也必须随之迁移。

6.1 为什么 map 必须放在 GPU 本地

一个直观的想法是：让 GPU 上的 probe 每次需要访问 map 时，通过 PCIe 去访问 host 内存。但这种方案在实际中性能很差。

在早期的实验版本里，我们测试过这种方式，很快就发现两个问题：

延迟过高：一次跨 PCIe 访问动辄几十微秒，远超可接受范围

锁竞争严重：SIMT 模式下成千上万个线程同时访问，锁竞争导致性能急剧下降

因此我们得出结论：GPU 上的多数 map 操作，必须在设备本地完成。

6.2 GPU Map 类型设计

针对不同场景，我们设计了几类 map：

在这些 map 之上，我们还提供了一套 GPU 端 helper：

// 获取线程索引u32 idx = bpf_get_thread_idx();u32 blk = bpf_get_block_idx();// 读取 GPU 全局计时器 (纳秒级)u64 ts = bpf_get_globaltimer();// 获取硬件拓扑信息u32 sm_id   = bpf_get_sm_id();u32 warp_id = bpf_get_warp_id();u32 lane_id = bpf_get_lane_id();// map 操作bpf_map_update_elem(&my_map, &key, &value, BPF_ANY);

6.3 SIMT-aware Verifier：让 GPU 上的 eBPF 安全可控

把 eBPF 程序直接搬到 GPU 上，最容易低估的一点是：CPU 世界里的 verifier 假定的是单线程语义，而 GPU 世界是成千上万个线程在 SIMT 模型下同时执行。

如果不加约束地让 probe 代码：

在每个线程里随意分支；用 lane-varying 的 key 去访问共享 map；写任意复杂、没有静态界限的循环。

那么在一个 warp 内就会产生严重的 divergence 与不可控的内存访问，轻则导致性能急剧下降，重则引发设备异常。

风险示例：

// 危险：lane-varying key → 写放大 32 倍int key = threadIdx.x;bpf_map_update(&map, &key, &val);// 危险：lane-varying 分支 → divergenceif (threadIdx.x % 2 == 0) { A } else { B }

SIMT-aware Verifier 规则：

与 CPU Verifier 的差异：

这些约束初看较为严格，但实践证明，这是让 GPU 上 eBPF 具备可预期性和可常驻性的前提。没有这些约束，难以判断性能问题是来自策略本身，还是某次执行路径导致的 SM 阻塞。

6.4 跨层 Map：host–device 状态统一管理

我们进一步探索了跨层 map 的设计：

逻辑上：BPF map 是一个 key-value store，host driver、GPU 设备 handler、用户态控制面都能访问；

物理上：runtime 将 map 分片布在不同位置：

一致性模型方面，每次 map 操作对单个 key 原子，但不提供全局顺序。以”快照”为主：GPU 侧频繁本地更新，runtime 在 kernel 完成或显式同步点，把局部 shard 合并回 canonical 实例。

对策略来说，这种弱一致 + 近似统计是足够的：即便看到稍旧的计数，只会影响决策质量，不会破坏内存映射正确性。

6.5 性能开销

在 GPU 上实现 eBPF 并不必然意味着要付出不可接受的 profiling 开销。

经过我们的测试，在典型场景下 bpftime 的开销在 2-3% 左右，而 NVBit 通常在 30-45%，CUPTI 在 15-25%。在 ResNet 推理等实际工作负载上，bpftime 的性能接近原生执行。

开销可控的原因主要有三点：SIMT-aware verifier 避免了 warp divergence 带来的性能损失；分层 map 设计让热数据保持在 GPU 本地访问；选择性插桩只在关键点进行观测。

更底层的 microbenchmark 表明，在引入 SIMT-aware 的 verifier 和 warp 级聚合执行之后，一个空 probe 的额外延迟可以压到非常接近 kernel baseline，简单的 array map lookup/update 相比 naive 的逐线程执行可以减少 30-40% 的开销，GPU 侧直接访问 map 比通过 PCIe 回到 CPU 去快了几个数量级。

这些数字说明，在合理选择插桩点、控制逻辑复杂度、谨慎使用跨域 map 的前提下，相关能力可以以常驻方式部署在生产环境。

GPU 观测工具实践

在运行时和安全机制基本成型之后，我们首先实现了一批结构上类似 bcc 工具的 GPU 侧示例程序。这些示例主要用于验证架构可行性，并帮助建立对 GPU 行为的直观认知。它们位于 bpftime/example/gpu/ 目录下。

7.1 kernelretsnoop：线程完成时间分布

传统 profiler 只能告诉你一个 kernel 总共跑了多久，但无法回答：这些时间是如何分布在各个线程上的？是否存在某些线程特别慢？

kernelretsnoop 能捕获每个 GPU 线程退出 kernel 的精确时刻，揭示传统 profiler 看不到的时序模式。

工作原理：

在 CUDA kernel 的退出点挂一个 GPU 端的 kretprobe；每个线程在返回前调用 GPU 专用 helper：

bpf_get_thread_idx()获取三维线程坐标

bpf_get_globaltimer()获取纳秒级 GPU 全局计时器

然后数据写入 GPU ringbuf map，host 侧进程持续消费。

核心 eBPF 代码：

#define BPF_MAP_TYPE_GPU_RINGBUF_MAP 1527struct {    __uint(type, BPF_MAP_TYPE_GPU_RINGBUF_MAP);    __uint(max_entries, 16);    __type(key, u32);    __type(value, struct big_struct);} rb SEC(".maps");// GPU 专用 helper 函数指针static const u64 (*bpf_get_globaltimer)(void) = (void *)502;static const u64 (*bpf_get_thread_idx)(u64 *x, u64 *y, u64 *z) = (void *)505;SEC("kretprobe/_Z9vectorAddPKfS0_Pf")int cuda__retprobe(){    struct data data;    bpf_get_thread_idx(&data.x, &data.y, &data.z);    data.timestamp = bpf_get_globaltimer();    bpf_perf_event_output(NULL, &rb, 0, &data, sizeof(struct data));    return 0;}

这段代码展示了几个关键点：

用 BPF_MAP_TYPE_GPU_RINGBUF_MAP (1527) 这个 GPU 专用的 map 类型；

helper 函数通过函数指针调用，ID 502 和 505 分别对应全局计时器和线程索引；

SEC 注解中的 _Z9vectorAddPKfS0_Pf 是 C++ mangled 后的 kernel 函数名。

典型输出：

Thread (0, 0, 0) timestamp: 1749147474550023136Thread (1, 0, 0) timestamp: 1749147474550023140  // +4nsThread (2, 0, 0) timestamp: 1749147474550023145  // +5ns...Thread (31, 0, 0) timestamp: 1749147474550023890 // +750ns 明显更晚!

线程 0 到 30 都在几纳秒内完成——这正是 warp 内线程同步执行的预期表现。但线程 31 晚了 750 纳秒，发生了什么？

诊断案例：Warp Divergence

检查 kernel 代码后发现，线程 31 恰好命中了边界条件——它处理数组的最后一个元素，触发了额外的边界检查或走了不同的代码分支。

由于 warp 内 32 个线程必须同步执行，当线程 31 走不同路径时，整个 warp 被迫串行化：先执行线程 0-30 的公共路径，再执行线程 31 的特殊路径。这就是 thread divergence，是影响 GPU 性能的常见因素。

诊断案例：内存访问模式问题

在另一个 kernel 中，分析时间戳发现：

Thread (0, 0, 0) timestamp: ...023140Thread (8, 0, 0) timestamp: ...023890  // 慢很多Thread (16, 0, 0) timestamp: ...023150Thread (24, 0, 0) timestamp: ...023900 // 又慢了

每隔 8 个线程就显著变慢。这指向内存访问模式问题：数据结构的布局导致每第 8 个线程访问不同的 memory bank，造成 bank conflict；或者这些线程触发了 cache miss 而其他线程命中了 cache。

通过关联线程索引和时序，可以精确定位哪些线程正在经历内存瓶颈，据此重构数据布局，确保连续线程访问连续内存地址——这是 GPU 最高效的访问模式。

可诊断的问题汇总：

7.2 threadhist：线程执行次数直方图

threadhist 关注的是在长期运行的 workload 中，各个线程累计执行 kernel 的次数。

实现上，它使用前文提到的 per-thread array map：每次 kernel 退出时，线程都会对自己的计数加一，host 侧定期将整张表导出并绘制直方图。

Thread 0: 210432Thread 1: 210432Thread 2: 210432Thread 3: 210432Thread 4: 158304  // 明显偏少 (只有 75%)

可诊断的问题：

在实际使用中，它主要用于识别两类典型问题：

grid-stride loop 里的工作不均衡：理论上每个线程应该处理差不多的迭代次数，但直方图可以清晰地显示出线程间的负载差异；

算法本身的负载不均：在一些稀疏图/稀疏矩阵场景下，不同 index 对应的工作量天差地别。

7.3 launchlate：把 CPU 与 GPU 之间的”排队空洞”照亮

launchlate 是一个实用的观测工具，专门测量：

从 CPU 调用 cudaLaunchKernel 到 GPU 真正开始执行 kernel 之间的时间。

它跨了 CPU 和 GPU 两个世界：

sequenceDiagram    participant CPU    participant GPU    CPU->>CPU: cudaLaunchKernel()    Note right of CPU: uprobe 记录 T1    CPU-->>GPU: 排队    GPU->>GPU: Kernel 开始    Note right of GPU: kprobe 记录 T2    Note over CPU,GPU: Latency = T2 - T1

典型输出：

12:34:56 Launch Latency Distribution:latency         : count    distribution100ns-1us       : 45       |********1-10us          : 234      |****************************************10-100us        : 167      |*****************************100us-1ms       : 89       |***************1-10ms          : 12       |**Total samples: 547

可诊断的问题：

这个直方图对我们有两个直接的帮助：

使大量小规模 kernel 中”launch 开销占比远大于计算本身”的情况变得可见，从而推动在合适场景下采用 kernel fusion、CUDA Graph 或持久 kernel 等手段；

在多租户、多进程场景下，帮助分析 P99 延迟偶发升高时，是否存在一小部分 kernel 受到上下文切换、PCIe 拥塞等因素影响，成为整体尾延迟的主导因素。

7.4 kernel_trace：stub 插桩示例

kernel_trace 是 5.5 节介绍的”任意位置插桩”能力的一个最小化示例。它展示了如何使用 stub 机制在 kernel 内部采集每个线程的执行信息。

示例代码：

__device__ __noinline__ void __bpftime_cuda__kernel_trace() {}__global__ void vectorAdd(const float *A, const float *B, float *C) {    int i = blockIdx.x * blockDim.x + threadIdx.x;    __bpftime_cuda__kernel_trace();  // stub 调用点    C[i] = A[i] + B[i];}

对应的 eBPF 程序会在每个线程执行到 stub 点时，采集 block/thread 索引和 GPU 时间戳，通过 ringbuf 传回 host。

输出示例：

[kernel_trace] ts=809635273301344 block=(0,0,0) thread=(0,0,0)[kernel_trace] ts=809635273301456 block=(0,0,0) thread=(1,0,0)[kernel_trace] ts=809635273301568 block=(0,0,0) thread=(2,0,0)...[kernel_trace] total events: 256

这个示例验证了 stub 机制的完整链路：从 CUDA 代码中的 stub 调用，到 PTX 层面的重定向，再到 eBPF 逻辑的执行和数据回传。

7.5 mem_trace：最小化的访存跟踪

mem_trace 是一个非常克制的例子：它没有试图做一个”GPU 上的 perf record”，而只是通过极少量的 probe 去统计：

某个 kernel 被调用了多少次；在插桩的几个 load/store 点上，各发生了多少次访问。

16:30:45
    pid=12345    mem_traces: 120

这类工具更多是在验证链路：从”GPU kernel 里的一条 ld.global 指令”，到”BPF 程序记录访问”，再到”host 进程看到统计并打印”，整个数据通路是否顺畅、开销是否可控。

7.6 SM/Warp/Lane 硬件映射（实验性）

传统 profiler 只能获取逻辑上的 CUDA 线程索引，无法回答：

我的线程块到底跑在哪个 SM 上？Warp 在 SM 内部是如何分布的？是否存在 SM 负载不均衡？

通过 PTX 特殊寄存器获取硬件拓扑信息：

// 新增 Helperu32 sm_id   = bpf_get_sm_id();    // %smidu32 warp_id = bpf_get_warp_id();  // %warpidu32 lane_id = bpf_get_lane_id();  // %laneid

我们解读一下输出的内容：

SM Utilization Histogram：展示每个 SM 上运行的线程数，可直观看出 SM 间负载是否均衡；

Load Balance Score: 43.9%：100% 表示完美均衡分布，43.9% 说明 SM 间负载严重不均；

Warp Distribution per SM：每个 SM 内的 Warp 分布情况；

Thread-to-Hardware Mapping：Block/Thread 到 SM/Warp/Lane 的映射。

这是从”逻辑追踪”迈向”物理/硬件感知追踪”的重要一步。

7.7 GPU L2 预取能力（实验性）

除了被动观测，bpftime for GPU 还可以用于主动优化。一个典型的例子是 UVM（Unified Virtual Memory）prefetch。

在 GPU 显存超量分配（oversubscription）的场景下，数据需要在 CPU 内存和 GPU 显存之间迁移。

通过 bpftime 挂载 eBPF 程序，我们可以在访存模式被识别后主动触发 prefetch，减少 page fault 带来的延迟。

// 新增 Helper (ID: 509)bpf_prefetch_l2(addr);// 对应 PTXprefetch.global.L2 [%0];

UVM 内存超额订阅测试：

从数据可以看到，通过 bpftime 的 prefetch 优化，执行时间从 107 秒降低到 70 秒，带宽提升约 52%。这说明 bpftime 不仅可以用于观测，还可以在观测的基础上实现主动的性能优化策略。

7.8 支持的工作负载

bpftime 的 GPU 插桩能力已经在多种实际工作负载上验证过。

对于深度学习框架，PyTorch 需要编译时包含 PTX 才能被插桩。

通过 TORCH_CUDA_ARCH_LIST=6.1+PTX 编译，之后可以观测其内部 CUDA kernel 的线程执行分布。

llama.cpp 可以直接用于分析 LLM 推理时的 GPU kernel 执行模式。vLLM 等其他推理框架也可以通过类似方式接入。

对于高性能计算库，Faiss（向量检索）和 CUTLASS（GEMM）都已经过验证。

在 Faiss 的测试中，我们可以观测到索引构建过程中各线程的执行次数分布，判断负载是否均衡。在 CUTLASS 的测试中，可以统计大矩阵乘法的 kernel launch 次数，验证插桩链路正常工作。

统一策略平面：从观测到控制

8.1 评估场景

我们基于同一套 eBPF 基础设施实现了驱动层和设备端统一的 GPU 策略平面，对以下几个问题进行了系统验证：

RQ1：单租户内存与调度策略

在显存超量分配的单租户场景下，利用 eBPF 在 UVM 和调度路径上的 hook，实现多种策略：

这些策略在不改应用代码的前提下，相比默认 UVM 和框架内置策略取得了显著的性能提升。

RQ2：多租户管理

在多租户环境中，通过驱动级调度 hook 和设备端 handler 组合，实现按租户区分的优先级和抢占策略：

相比仅依赖 CUDA stream 或静态优先级策略，具有更好的 tail latency 和吞吐折中。

RQ3：可编程性与通用性

例如，llama.cpp 的 region cache 策略，稍作修改就能用于 vLLM 的 KV-cache；顺序 prefetch 策略可以从 GNN adjacency block 迁移到 Faiss index segment。

RQ4：机制开销

即便叠加较复杂的内存与调度策略，总体开销仍保持在 5% 以内。

8.2 CLC + bpftime：GPU 内核级调度

CLC (Cooperative Launch Control) 是一种 GPU 内核级调度框架，用于解决 GPU 上的负载不均衡问题。

在很多 GPU kernel 中，不同 block 的工作量并不相同。例如在某些 GEMM 变体中，有些 tile 需要做大量的预处理（边界条件检查、稀疏结构展开等），有些则几乎不需要。

传统的 CUDA 启动方式采用静态分配策略，每个 SM 负责固定的 block 集合。这种方式在负载不均衡的场景下会导致部分 SM 提前完成后处于空闲状态，而另一部分 SM 因为承担了较重的任务而成为整体执行时间的瓶颈。

CLC 采用动态任务窃取（work-stealing）策略来解决这一问题。它将所有 block 组织为全局任务队列，当某个 SM 完成当前任务后，可以从队列中获取下一个待执行的 block，从而实现更均衡的负载分配，缩短整体执行时间。

bpftime 的集成方式：

整个流程分为控制面和数据面：

控制面：先在 host 起 eBPF 控制进程，向共享内存 map 写入策略配置（启用标志、目标 kernel 名、wrapper/policy PTX 路径），让 agent 知道这次要把哪个 kernel 换成 CLC 版本。

数据面：运行目标程序时，libbpftime-agent.so 通过 hook cudaLaunchKernel，在每次 launch 时：

1. 从当前 fatbin 中抽取 PTX；

2. 注入 eBPF pass（ptxpass 做插桩）；

3. 用 nvJitLink 把用户 kernel（重命名后的 baseline）与 wrapper/policy PTX 组合，生成 <kernel>_with_policy；

4. 如果策略配置命中，就改用 _with_policy kernel；否则继续走原始 baseline。

与静态编译的区别在于：bpftime 在运行期动态抽取 PTX、拼接 CLC wrapper，再决定是否重定向 kernel。这条”运行期注入 + 热切换”路径完全依赖 bpftime 才能成立。

性能验证

我们在 GEMM 负载不均衡场景下进行了测试。通过调整 prologue 参数（每个 block 开始前的额外计算量）来模拟不同程度的负载不均衡：

可以看到，负载越不均衡（prologue 越大），CLC 的收益越明显。在 prologue=1000 的极端不均衡场景下，性能提升超过 60%。

实际应用场景

在真实应用中，prologue 对应的是各种导致 block 间工作量差异的因素：

1. 每个 tile/block 进入主计算前的局部数据预处理、边界条件处理；

2. 稀疏结构展开/压缩；

3. 自适应算法中不同区域需要的迭代次数差异；

4. 不规则数据结构（图算法、稀疏矩阵）中不同 block 负责的数据量差异。

核心价值

bpftime 不仅是观测工具，开始具备通过动态注入来控制 GPU 调度行为的能力。整个闭环是：

观测 (bpftime) → 分析 (eBPF) → 决策 (Map) → 执行 (CLC)      ↑                                         ↓      └────────── 反馈指标 ──────────────────────┘

全链路观测与慢节点观测

有了 GPU 侧 eBPF 能力之后，我们很快意识到：仅观察 GPU 本身往往不足以解释真实问题。

例如，当通信相关 kernel 偶发性延迟时，它可能与 DataLoader 延迟、网络抖动或内核调度有关，仅靠设备端信息难以区分。

9.1 统一 CPU–GPU 观测架构

因此在架构上，我们刻意保持以下原则：bpftime for GPU 不被设计成独立的设备侧工具；而是与 CPU 侧、内核侧 eBPF 能力共同构成统一的观测和控制平面。

多层次 probe 部署：

所有这些事件最终通过共享的 map 和统一的时间基准被串联起来。

9.2 跨层关联分析

统一时间线示例：

CPU 侧:[0-50ms]   DataLoader 加载数据[50-51ms]  cudaLaunchKernel[51-200ms] cudaDeviceSynchronize
内核:[60-80ms]  page fault (UVM 迁移)[100-120ms] cgroup 内存限流
GPU 侧:[55-80ms]  排队等待[80-150ms] Kernel 执行[100-130ms] warp stall           ↑ 与 page fault 时间重叠!

9.3 大模型训练慢节点问题

在大规模分布式训练任务里，通常会有几十甚至上百块 GPU 参与同一轮 iteration。一旦其中某个节点在计算或通信阶段显著慢于其他节点，整体迭代时间就会被该节点决定。

木桶效应：在 all-reduce/barrier 模式下，所有 rank 必须等待最慢的那个完成，一个 rank 慢 = 全体等待。

根据实际排查经验，慢节点的原因大致可以归为三类：

9.4 分析流程：从原始信号到慢节点判别

结合 bpftime for GPU，我们大致建立了这样一条分析链路：

GPU 内核层（设备端 eBPF）

CPU 主机层（uprobes/kprobes）

跨节点聚合与对齐

基于这套剖面，我们可以给每一轮训练打上一些标签，例如：

哪些 rank 在 forward 上明显偏慢；哪些 rank 在 all-reduce 上经常是最后一个进入或最晚一个退出；哪些 rank 的 host 侧频繁出现 DataLoader 超时或 IO 抖动。

9.5 排查示例

在一次分布式训练任务中，我们遇到了 step time 抖动的问题：吞吐偏低，但粗粒度监控显示 GPU 利用率正常，难以定位是哪个节点、哪个阶段导致。

接入 bpftime 后，我们从 GPU 侧和 CPU 侧同时采集数据。GPU 侧通过 kernelretsnoop 观测到 rank 1 的 backward kernel 耗时在 50-85ms 之间波动，而 rank 0 稳定在 50-52ms。CPU 侧通过 kprobe 发现 rank 1 存在频繁的 page fault 和 cgroup 限流。

将两侧数据按时间线对齐后，可以清晰看到：rank 1 的 page fault spike 发生在 52-60ms 区间，紧接着 backward kernel 就出现了延长（60-85ms）。两个事件在时间上高度重叠。

最终确认根因：rank 1 节点上同时运行了 IO 密集型任务，触发频繁 page fault 并导致 cgroup 资源限流。CPU 侧算子下发延迟传导到 GPU。

解决方案：迁移 IO 任务到其他节点，step time 恢复稳定。关键价值：GPU 细粒度时间戳与 host 事件的明确关联。

传统方法只能看到 GPU 利用率正常、step time 偶尔变长，而 bpftime 能够定位具体哪个 rank、哪个阶段，并关联 CPU 侧事件。

工程实践中的典型问题

10.1 PTX 版本、架构与 fatbin 的”选择题”

在最初的实现中，我们假设”拿到 fatbin，解出一份 PTX 即可完成插桩”。实践表明，实际情况要复杂得多：

同一个 fatbin 里可能有多份 PTX/SASS，对应不同的 SM 架构；编译选项不同，PTX 暴露的符号、寄存器布局也有差异；有些应用只带 SASS，不带 PTX。

解决方案：

根据当前 GPU 的 SM 版本，精确选择匹配的 PTX；在 PTX 解析/重写时，尽量不依赖具体的 register 编号，而是构造自己的局部寄存器空间；对于 SASS-only 的场景，选择保守地降级，不强行插桩

10.2 Map 放哪儿：HBM、DRAM 还是 shared memory？

经验总结：

和策略紧密相关、更新频繁的”热状态”，尽量放在 GPU 本地（global memory / shared memory），按 epoch 或固定周期向 host 汇总；

低频的全局配置、统计信息放在 host DRAM，由 GPU 偶尔读一次；

对于需要 host 高频访问、GPU 高频更新的状态，尽量用分层 map + 批量同步，而不是每次操作都跨域。

这些经验后来被固化为 runtime 中的”层次化 map placement”策略。总体来说，在 GPU 上进行 eBPF 设计时，数据所在位置需要被作为一等设计维度对待。

10.3 与现有生态的融合

最后一个问题主要不在技术，而在生态。如果 bpftime for GPU 被设计为一个完全独立的 profiler，那么其与现有监控、调度、容量管理体系的融合成本会非常高。

因此在设计之初，我们就确定了以下原则：

尽量使用”标准 eBPF 程序 + map”的方式表达 probe；数据输出格式兼容现有的火焰图、热力图等可视化组件；在 API 层面贴近 bcc / libbpf 的使用习惯，降低用户迁移成本。

这部分工作虽然偏工程实现，但对系统能否落地到生产环境具有关键作用。

小结与展望

回顾这项工作，bpftime 在 GPU 上的主要目标是把原本存在于 CPU 和内核世界的 eBPF 可编程观测与策略能力，扩展到 GPU 驱动与设备端。

在技术实现上，我们通过 llvmbpf 编译器将 eBPF 程序编译为 PTX，通过 fatbin hook 和 PTX 注入实现透明插桩，设计了适配 GPU 的 per-thread array、ringbuf 和共享 map，并集成了 CLC 实现运行时策略注入与调度。

这套能力已经在 PyTorch、vLLM、llama.cpp、Faiss、CUTLASS 等工作负载上验证，典型开销在 2-3%，可以常驻生产环境。

从系统演进的角度看，bpftime for GPU 是一个以用户态观测和轻量策略验证为主的基础平台。

下一步可以与统一 GPU 策略平面方向形成完整闭环：

在 host 侧，以 Linux eBPF 为基础，在 GPU 驱动内部暴露稳定、安全的策略 hook，使内存放置、调度、抢占等成为可编程的内核机制；

在 device 侧，把经 verifier 约束的 eBPF 逻辑编译并注入到 GPU kernel 中，在 warp/block 粒度上实现设备端可编程策略；

通过跨 CPU–GPU 的层次化 BPF map，把 host 驱动、设备端 handler 以及用户态控制平面统一到一套状态视图下。

从更长远的视角看，这项工作的重点不在某一次具体的性能提升倍数，而在于观测与控制范围的扩展。

GPU 不再是 eBPF 体系之外的单独设备，而是可以与 CPU、内核、网络一起纳入统一的、可编程的观测与控制平面。

在整个异构计算栈上具备统一语言和工具之后，关于性能、可靠性和多租户公平性的一些长期问题，可以用更系统化、可验证的方式来分析和解决。

作者：蒋明辉(知乎 Sy11)

封面：蒋明辉(知乎 Sy11)

首图：主核(Kernyr)

审校：泽文(Zevorn)

图片、资料来源：

[1] bpftime . bpftime: Userspace eBPF runtime for uprobe, syscall and kernel-module tracing. https://github.com/eunomia-bpf/bpftime

[2] 尹瑞星. GPU Profiling: Extending eBPF to GPUs. 收录于 CLK2025 AI 基础设施与 eBPF 应用分论坛. https://github.com/ChinaLinuxKernel/CLK2025/tree/main/6%20AI%20基础设施与%20eBPF%20应用分论坛

[3] 程书意. 基于 eBPF 的大模型性能分析及慢节点检测. 收录于 CLK2025 AI 基础设施与 eBPF 应用分论坛. https://github.com/ChinaLinuxKernel/CLK2025/tree/main/6%20AI%20基础设施与%20eBPF%20应用分论坛

[4] llvmbpf 项目. llvmbpf: eBPF to various targets compiler based on LLVM. https://github.com/eunomia-bpf/llvmbpf

[5] NVIDIA. CUDA C++ Programming Guide - Cooperative Groups. https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#cooperative-groups

[6] Bpftime: Userspace eBPF runtime Linux Plumbers Conference Richmond, Virginia | November 13-15, 2023 https://lpc.events/event/17/contributions/1639/

[7] OSDI’25: Yusheng Zheng, Tong Yu, Yiwei Yang. bpftime: Userspace eBPF runtime for Observability, Network, GPU & General extensions Framework. https://www.usenix.org/conference/o