bpftime for GPU 深度实战：将 eBPF 带进 GPU Kernel 内部——从 PTX 级插桩到线程级可观测性的全链路架构解析

在 AI 训练、科学计算、高性能推理快速发展的今天，GPU 已经成为现代算力基础设施的核心。但一个长期存在的问题是：我们知道 GPU 很重要，却始终很难真正"看清"它内部发生了什么。2026 年 4 月 18 日，在第四届 eBPF 开发者大会上，bpftime for GPU 给出了答案——用 eBPF 的方式深入 GPU Kernel 内部。

一、背景：GPU 可观测性的"黑箱"困境

1.1 GPU 为什么难以观测？

GPU 与 CPU 的架构差异决定了它们的可观测性路径截然不同。

CPU 侧的可观测性已经相当成熟：

• perf 可以采集硬件性能计数器
• strace 可以追踪系统调用
• ftrace 可以追踪内核函数
• eBPF 可以在内核和用户态任意位置插桩

这些工具的共同特点是：它们可以深入到指令级别，观察每一个线程的执行细节。

GPU 侧的可观测性则完全不同：

• GPU Kernel 是大规模并行的，一次启动可能涉及数万个线程
• GPU 的执行模型是 SIMT（单指令多线程），一个 warp 中的 32 个线程同时执行相同指令
• GPU 没有操作系统，没有系统调用，没有特权级切换
• GPU Kernel 执行期间，CPU 几乎无法干预

这意味着，传统的 CPU 观测工具在 GPU 上完全失效。

1.2 现有 GPU 观测工具的局限

NVIDIA 提供的 Nsight 工具链是目前最强大的 GPU 观测方案：

这些工具可以回答以下问题：

• 某个 Kernel 总共运行了多久
• 某个阶段的平均表现如何
• 哪些硬件资源使用较多

但它们很难回答：

• 某一个线程究竟卡在了哪里？
• Kernel 内部的某个分支路径执行了多少次？
• 为什么某个线程的执行时间比其他线程长 10 倍？

1.3 为什么需要 eBPF 的能力？

eBPF 在 CPU 世界中已经证明了它的价值。它让开发者能够：

1. 动态插桩：无需重启应用，无需修改源码
2. 可编程观测：用 C 语言编写自定义观测逻辑
3. 低开销：在内核沙盒中运行，性能影响可控
4. 安全可靠：验证器确保程序不会崩溃系统

如果这些能力能扩展到 GPU，会带来什么？

• Per-thread 级别的时间戳采集：不再只是聚合统计
• 任意位置的插桩：在 Kernel 内部任意指令处设置观测点
• 可编程的观测逻辑：根据条件选择性采集数据
• 动态注入：对运行中的 CUDA 应用进行观测

这正是 bpftime for GPU 想要解决的问题。

二、bpftime 项目概览：用户态 eBPF 运行时

2.1 bpftime 是什么？

bpftime 是一个高性能的用户态 eBPF 运行时框架，它的核心目标是：让 eBPF 程序能够在用户态运行，同时保持与内核 eBPF 的兼容性。

项目架构：

┌─────────────────────────────────────────────────────────────┐
│                      Application Layer                       │
├─────────────────────────────────────────────────────────────┤
│  bpftrace  │  libbpf  │  CO-RE BPF  │  Custom Tools         │
├─────────────────────────────────────────────────────────────┤
│                     bpftime Runtime                         │
│  ┌─────────────┐  ┌─────────────┐  ┌─────────────┐         │
│  │   Loader    │  │  Verifier   │  │    Maps     │         │
│  └─────────────┘  └─────────────┘  └─────────────┘         │
│  ┌─────────────┐  ┌─────────────┐  ┌─────────────┐         │
│  │  VM (LLVM)  │  │   Helpers   │  │   Attach    │         │
│  └─────────────┘  └─────────────┘  └─────────────┘         │
├─────────────────────────────────────────────────────────────┤
│                     Event Sources                            │
│  Uprobe  │  Syscall  │  XDP  │  GPU Kernel  │  ...         │
└─────────────────────────────────────────────────────────────┘

2.2 bpftime 的核心特性

性能优势：相比内核 eBPF，bpftime 在用户态避免了内核态切换的开销：

跨平台能力：bpftime 不依赖内核 eBPF 支持，可以在以下环境运行：

• 旧版本 Linux 内核
• 非 Linux 操作系统（FreeBSD、macOS）
• 容器环境（无需特权）

兼容性：使用标准 eBPF 工具链开发：

• clang 编译 BPF 程序
• libbpf 加载和管理程序
• bpftrace 一行命令

2.3 bpftime 的架构组件

bpftime/
├── vm/                 # eBPF 虚拟机和 JIT 编译器
│   ├── llvmbpf/       # 基于 LLVM 的高性能 JIT
│   └── ubpf/          # 基于 ubpf 的解释器/JIT
├── runtime/           # 用户态运行时
│   ├── maps/          # eBPF Map 实现
│   ├── helpers/       # Helper 函数实现
│   └── syscall-server/# 加载器
├── attach/            # 事件附加机制
│   ├── frida/         # 用户态函数 hook
│   ├── syscall/       # 系统调用 hook
│   └── nv_attach_impl/# GPU Kernel 插桩
├── verifier/          # 用户态验证器
└── daemon/            # 与内核 eBPF 协作的守护进程

三、bpftime for GPU：核心技术架构

3.1 整体设计思路

bpftime for GPU 的核心目标是将 eBPF 的可编程观测能力带入 GPU Kernel 内部。实现这一目标需要解决三个关键问题：

1. 如何进入 GPU Kernel？

   • GPU Kernel 在 GPU 上执行，CPU 无法直接干预
   • 需要在 Kernel 加载和编译阶段介入

2. 如何编写 GPU 侧的观测逻辑？

   • GPU 的编程模型与 CPU 完全不同
   • 需要将 eBPF 程序转换为 GPU 可执行的形式

3. 如何回传观测数据？

   • GPU 和 CPU 有独立的内存空间
   • 需要高效的数据传输机制

bpftime 的解决方案是：PTX 级插桩。

3.2 PTX：GPU 的中间表示

PTX（Parallel Thread Execution）是 NVIDIA GPU 的中间指令集。它的定位类似于 CPU 的 LLVM IR：

CUDA C 代码
    ↓ nvcc 前端
PTX 代码（中间表示）
    ↓ PTX 汇编器
SASS 代码（GPU 机器码）
    ↓ GPU 执行

选择 PTX 级插桩的关键优势：

1. 可移植性：PTX 是跨架构的，同一份代码可以在 sm_75、sm_80、sm_90 等不同架构上运行
2. 可读性：PTX 是文本格式，便于分析和修改
3. 工具链成熟：NVIDIA 提供了完整的 PTX 工具链

3.3 nv_attach_impl：GPU 动态插桩流水线

bpftime GPU 的核心实现是 nv_attach_impl，它构建了一条完整的动态插桩流水线：

┌──────────────────────────────────────────────────────────────────┐
│                    nv_attach_impl 流水线                          │
├──────────────────────────────────────────────────────────────────┤
│                                                                  │
│  1. 拦截 CUDA 模块加载                                            │
│     ↓ Frida-gum hook __cudaRegisterFatBinary                    │
│                                                                  │
│  2. 提取 PTX                                                      │
│     ↓ cuobjdump --extract-ptx                                   │
│                                                                  │
│  3. 执行 PTX Pass                                                 │
│     ↓ kprobe / kretprobe / memcapture                           │
│                                                                  │
│  4. 编译 eBPF 探针为 PTX                                          │
│     ↓ bpftime LLVM JIT → PTX                                    │
│                                                                  │
│  5. 自动寄存器保护                                                │
│     ↓ Register Guard                                            │
│                                                                  │
│  6. 重新编译为 cubin                                              │
│     ↓ nvPTXCompiler                                             │
│                                                                  │
│  7. 替换原始 GPU Module                                           │
│     ↓ 透明替换，应用无感知                                         │
│                                                                  │
└──────────────────────────────────────────────────────────────────┘

让我们逐一分析每个步骤。

步骤 1：拦截 CUDA 模块加载

CUDA 程序在运行时需要加载 GPU 代码模块（fat binary）。这个模块包含了编译好的 Kernel 代码。bpftime 使用 Frida-gum 框架 hook __cudaRegisterFatBinary 函数：

// 原始函数签名
void** __cudaRegisterFatBinary(void *fatBin);

// Hook 逻辑
void** hooked_cudaRegisterFatBinary(void *fatBin) {
    // 1. 提取 fat binary 中的信息
    fat_binary_header_t *header = (fat_binary_header_t *)fatBin;
    
    // 2. 遍历所有 PTX 模块
    for (each_ptx_module in fatBin) {
        // 3. 对目标 Kernel 进行插桩
        if (should_instrument(kernel_name)) {
            instrumented_ptx = apply_ptx_pass(original_ptx);
        }
    }
    
    // 4. 调用原始注册函数
    return original_cudaRegisterFatBinary(modified_fatBin);
}

关键点：

• 透明性：应用完全无感知，不需要修改任何代码
• 动态性：可以在应用运行时注入，不需要重启

步骤 2：提取 PTX

从 fat binary 中提取 PTX 使用 NVIDIA 提供的工具：

# 从 fat binary 提取 PTX
cuobjdump --extract-ptx app.cubin -output app.ptx

提取出的 PTX 代码示例：

.version 8.0
.target sm_90
.address_size 64

.visible .entry vector_add(
    .param .u64 vector_add_param_0,
    .param .u64 vector_add_param_1,
    .param .u64 vector_add_param_2,
    .param .u32 vector_add_param_3
)
{
    .reg .u64 %rd<10>;
    .reg .u32 %r<10>;
    .reg .f32 %f<10>;

    // Kernel 入口
    ld.param.u64 %rd1, [vector_add_param_0];
    ld.param.u64 %rd2, [vector_add_param_1];
    ld.param.u64 %rd3, [vector_add_param_2];
    ld.param.u32 %r1, [vector_add_param_3];

    // 计算线程索引
    cvta.to.global.u64 %rd4, %rd1;
    cvta.to.global.u64 %rd5, %rd2;
    cvta.to.global.u64 %rd6, %rd3;
    
    mad.wide.u32 %rd7, %r1, 4, %rd4;
    mad.wide.u32 %rd8, %r1, 4, %rd5;
    mad.wide.u32 %rd9, %r1, 4, %rd6;
    
    ld.global.f32 %f1, [%rd7];
    ld.global.f32 %f2, [%rd8];
    add.f32 %f3, %f1, %f2;
    st.global.f32 [%rd9], %f3;
    
    ret;
}

步骤 3：执行 PTX Pass

PTX Pass 是对 PTX 代码进行改写的核心逻辑。bpftime 支持三种类型的 Pass：

kprobe Pass：在 Kernel 入口插桩

// 原始 PTX
.visible .entry vector_add(...) {
    // 原始代码
    ld.param.u64 %rd1, [vector_add_param_0];
    ...
    ret;
}

// 插桩后 PTX
.visible .entry vector_add(...) {
    // ===== 插桩代码开始 =====
    {
        // 获取全局时间戳
        mov.u64 %clock, %clock64;
        // 获取线程索引
        mov.u32 %tid, %tid.x;
        // 写入 ring buffer
        call _bpf_kprobe_entry, (%clock, %tid);
    }
    // ===== 插桩代码结束 =====
    
    // 原始代码
    ld.param.u64 %rd1, [vector_add_param_0];
    ...
    ret;
}

kretprobe Pass：在每个 ret 指令前插桩

// 原始 PTX
.visible .entry vector_add(...) {
    ...
    ret;  // 原始返回
}

// 插桩后 PTX
.visible .entry vector_add(...) {
    ...
    // ===== 插桩代码开始 =====
    {
        // 获取全局时间戳
        mov.u64 %clock, %clock64;
        // 获取线程索引
        mov.u32 %tid, %tid.x;
        // 写入 ring buffer
        call _bpf_kretprobe_exit, (%clock, %tid);
    }
    // ===== 插桩代码结束 =====
    ret;
}

memcapture Pass：追踪内存访问

// 原始 PTX
ld.global.f32 %f1, [%rd7];
st.global.f32 [%rd9], %f3;

// 插桩后 PTX
{
    // 记录 load 地址和值
    mov.u64 %addr, %rd7;
    call _bpf_memcapture_load, (%addr);
}
ld.global.f32 %f1, [%rd7];

{
    // 记录 store 地址和值
    mov.u64 %addr, %rd9;
    mov.f32 %val, %f3;
    call _bpf_memcapture_store, (%addr, %val);
}
st.global.f32 [%rd9], %f3;

步骤 4：编译 eBPF 探针为 PTX

这是 bpftime 最核心的创新：将用户编写的 eBPF C 程序编译为 PTX 代码。

编译链路：

eBPF C 程序
    ↓ clang -target bpf
eBPF 字节码
    ↓ bpftime LLVM JIT
LLVM IR
    ↓ LLVM PTX 后端
PTX 代码

示例：编写一个简单的 GPU 探针

// gpu_probe.c
#include <bpf/bpf_helpers.h>
#include <bpf/bpf_tracing.h>

struct event_t {
    u64 timestamp;
    u32 thread_idx;
    u32 sm_id;
    u64 kernel_id;
};

struct {
    __uint(type, BPF_MAP_TYPE_RINGBUF);
    __uint(max_entries, 1 << 24);
} events SEC(".maps");

SEC("gpu_kprobe")
int trace_kernel_entry(struct pt_regs *ctx) {
    struct event_t *event;
    
    event = bpf_ringbuf_reserve(&events, sizeof(*event), 0);
    if (!event)
        return 0;
    
    event->timestamp = bpf_get_globaltimer();
    event->thread_idx = bpf_get_thread_idx();
    event->sm_id = bpf_get_sm_id();
    event->kernel_id = ctx->kernel_id;
    
    bpf_ringbuf_submit(event, 0);
    return 0;
}

char LICENSE[] SEC("license") = "GPL";

编译和加载：

# 编译 eBPF 程序
clang -target bpf -g -O2 -c gpu_probe.c -o gpu_probe.o

# 使用 bpftime 加载到 GPU Kernel
bpftime gpu-attach --pid <cuda_app_pid> --kernel vector_add gpu_probe.o

步骤 5：自动寄存器保护

在 GPU Kernel 中插桩面临一个严峻挑战：寄存器冲突。GPU 的寄存器资源非常宝贵，如果探针代码使用了原始 Kernel 已经使用的寄存器，会导致数据损坏。

bpftime 实现了 Register Guard 机制：

1. 分析原始 PTX：识别所有已使用的寄存器
2. 分析探针 PTX：识别探针需要的寄存器
3. 分配空闲寄存器：如果没有空闲寄存器，则保存/恢复

// Register Guard 生成的保护代码
{
    // 保存寄存器
    mov.u64 %saved_r10, %r10;
    mov.u64 %saved_r11, %r11;
    mov.f32 %saved_f10, %f10;
    
    // 探针代码（使用 %r10, %r11, %f10）
    ...
    
    // 恢复寄存器
    mov.u64 %r10, %saved_r10;
    mov.u64 %r11, %saved_r11;
    mov.f32 %f10, %saved_f10;
}

步骤 6：重新编译为 cubin

使用 NVIDIA 的 nvPTXCompiler 库将修改后的 PTX 编译为 cubin：

#include <nvPTXCompiler.h>

nvPTXCompilerHandle compiler;
nvPTXCompilerResult result;

// 创建编译器句柄
nvPTXCompilerCreate(&compiler, ptx_source, ptx_size);

// 编译
nvPTXCompilerCompile(compiler, options, num_options);

// 获取 cubin
void *cubin_data;
size_t cubin_size;
nvPTXCompilerGetCompiledProgram(compiler, &cubin_data, &cubin_size);

// 销毁编译器
nvPTXCompilerDestroy(&compiler);

步骤 7：替换原始 GPU Module

最后一步是将修改后的 cubin 替换原始模块：

// 创建新的 fat binary
void *new_fatbin = create_modified_fatbinary(
    original_fatbin,
    instrumented_cubin,
    kernel_name
);

// 替换
original_cudaRegisterFatBinary(new_fatbin);

四、GPU 专用 Map 和 Helper

4.1 GPU Map 类型

为了让 eBPF 在 GPU 内部真正可用，bpftime 扩展了一套 GPU 专用的 Map 类型：

GPU_RINGBUF_MAP 是最关键的 Map 类型，它实现了高效的 GPU → CPU 数据传输：

┌─────────────────────────────────────────────────────────┐
│                    GPU Ring Buffer                       │
├─────────────────────────────────────────────────────────┤
│                                                         │
│  GPU 线程写入                   CPU 进程读取             │
│  ┌──────────┐                  ┌──────────┐            │
│  │ Thread 0 │ ──────┐          │ Reader   │            │
│  └──────────┘       │          └──────────┘            │
│  ┌──────────┐       │                   ▲              │
│  │ Thread 1 │ ──────┼──► Ring Buffer ───┘              │
│  └──────────┘       │    (UVA 零拷贝)                   │
│  ┌──────────┐       │                                   │
│  │ Thread N │ ──────┘                                   │
│  └──────────┘                                           │
│                                                         │
└─────────────────────────────────────────────────────────┘

UVA（Unified Virtual Addressing）技术使得 GPU 和 CPU 可以共享同一块物理内存，从而实现零拷贝数据回传。

4.2 GPU Helper 函数

bpftime 为 GPU 扩展了专用的 Helper 函数：

// 获取纳秒级全局时钟
u64 bpf_get_globaltimer(void);

// 获取线程索引
u32 bpf_get_thread_idx(void);

// 获取 SM ID
u32 bpf_get_sm_id(void);

// 获取 Warp ID
u32 bpf_get_warp_id(void);

// 获取 Lane ID
u32 bpf_get_lane_id(void);

// GPU 内打印信息到 Host
void ebpf_puts(const char *fmt, ...);

// 插入内存屏障
void bpf_gpu_membar(void);

// 终止当前线程
void bpf_gpu_exit(void);

使用示例：统计每个 SM 的 Kernel 执行次数

struct {
    __uint(type, BPF_MAP_TYPE_HASH);
    __type(key, u32);
    __type(value, u64);
    __uint(max_entries, 1024);
} sm_counter SEC(".maps");

SEC("gpu_kprobe")
int count_sm_usage(struct pt_regs *ctx) {
    u32 sm_id = bpf_get_sm_id();
    u64 *count = bpf_map_lookup_elem(&sm_counter, &sm_id);
    
    if (count) {
        __sync_fetch_and_add(count, 1);
    } else {
        u64 init = 1;
        bpf_map_update_elem(&sm_counter, &sm_id, &init, BPF_ANY);
    }
    
    return 0;
}

五、性能对比与基准测试

5.1 CPU 侧性能：延续 bpftime 一贯优势

在 uprobe/uretprobe 场景下，bpftime 相比内核 eBPF 展现出明显优势：

根本原因在于：bpftime 避免了内核态切换带来的额外成本。

5.2 GPU 侧性能：显著优于 NVBit

NVBit 是 NVIDIA 官方的动态二进制插桩工具，工作在 SASS 层。bpftime 工作在 PTX 层，带来了性能优势。

基准测试配置：

• 测试用例：10,000 次 vector addition
• GPU 1：NVIDIA P40 (Pascal 架构)
• GPU 2：NVIDIA RTX 5090 (Blackwell 架构)

NVIDIA P40 结果：

NVIDIA RTX 5090 结果：

关键发现：

1. bpftime GPU 的探针开销显著低于 NVBit
2. 在新架构 GPU 上，这种优势更加明显
3. RTX 5090 上，bpftime 比 NVBit 快 6.8 倍

5.3 Attach 开销分析

首次挂载的时间开销：

关键点：

• 首次挂载需要约 100ms，主要耗时在 PTX 改写和重编译
• 完成后，后续 Kernel launch 可以直接复用已加载结果
• 不会再引入额外 attach 开销

六、实战案例：从入门到进阶

6.1 案例一：Kernel 执行时间追踪

目标：追踪每个 Kernel 的执行时间，找出性能瓶颈。

Step 1：编写 eBPF 程序

// kernel_time.c
#include <bpf/bpf_helpers.h>
#include <bpf/bpf_tracing.h>

struct start_event {
    u64 timestamp;
    u32 kernel_id;
};

struct end_event {
    u64 timestamp;
    u32 kernel_id;
    u64 duration_ns;
};

struct {
    __uint(type, BPF_MAP_TYPE_HASH);
    __type(key, u32);
    __type(value, u64);
    __uint(max_entries, 1024);
} start_times SEC(".maps");

struct {
    __uint(type, BPF_MAP_TYPE_RINGBUF);
    __uint(max_entries, 1 << 24);
} events SEC(".maps");

SEC("gpu_kprobe")
int trace_start(struct pt_regs *ctx) {
    u32 kernel_id = ctx->kernel_id;
    u64 ts = bpf_get_globaltimer();
    
    bpf_map_update_elem(&start_times, &kernel_id, &ts, BPF_ANY);
    return 0;
}

SEC("gpu_kretprobe")
int trace_end(struct pt_regs *ctx) {
    u32 kernel_id = ctx->kernel_id;
    u64 end_ts = bpf_get_globaltimer();
    
    u64 *start_ts = bpf_map_lookup_elem(&start_times, &kernel_id);
    if (!start_ts)
        return 0;
    
    struct end_event *event;
    event = bpf_ringbuf_reserve(&events, sizeof(*event), 0);
    if (!event)
        return 0;
    
    event->timestamp = end_ts;
    event->kernel_id = kernel_id;
    event->duration_ns = end_ts - *start_ts;
    
    bpf_ringbuf_submit(event, 0);
    return 0;
}

char LICENSE[] SEC("license") = "GPL";

Step 2：编译和加载

# 编译
clang -target bpf -g -O2 -c kernel_time.c -o kernel_time.o

# 找到目标 CUDA 应用
CUDA_PID=$(pgrep -f "my_cuda_app")

# 加载到所有 Kernel
bpftime gpu-attach --pid $CUDA_PID kernel_time.o

# 或者只加载到特定 Kernel
bpftime gpu-attach --pid $CUDA_PID --kernel "vector_add" kernel_time.o

Step 3：查看结果

bpftime cat-events kernel_time.o

# 输出示例
TIME         KERNEL_ID   DURATION_NS
10:23:45.123 0           15234
10:23:45.125 1           8721
10:23:45.128 2           156789
...

6.2 案例二：线程级热点分析

目标：找出哪些线程执行时间最长，分析负载不均衡问题。

// thread_histogram.c
#include <bpf/bpf_helpers.h>

struct {
    __uint(type, BPF_MAP_TYPE_HASH);
    __type(key, u32);  // thread_idx
    __type(value, u64); // execution_time
    __uint(max_entries, 65536);
} thread_times SEC(".maps");

struct {
    __uint(type, BPF_MAP_TYPE_HASH);
    __type(key, u64);  // time bucket (microseconds)
    __type(value, u32); // count
    __uint(max_entries, 1000);
} histogram SEC(".maps");

SEC("gpu_kprobe")
int record_start(struct pt_regs *ctx) {
    u32 tid = bpf_get_thread_idx();
    u64 ts = bpf_get_globaltimer();
    bpf_map_update_elem(&thread_times, &tid, &ts, BPF_ANY);
    return 0;
}

SEC("gpu_kretprobe")
int record_end(struct pt_regs *ctx) {
    u32 tid = bpf_get_thread_idx();
    u64 end_ts = bpf_get_globaltimer();
    
    u64 *start_ts = bpf_map_lookup_elem(&thread_times, &tid);
    if (!start_ts)
        return 0;
    
    u64 duration_us = (end_ts - *start_ts) / 1000;
    
    // 更新直方图
    u64 bucket = duration_us / 10 * 10;  // 10us 为一个桶
    u32 *count = bpf_map_lookup_elem(&histogram, &bucket);
    if (count) {
        (*count)++;
    } else {
        u32 init = 1;
        bpf_map_update_elem(&histogram, &bucket, &init, BPF_ANY);
    }
    
    return 0;
}

char LICENSE[] SEC("license") = "GPL";

分析结果：

bpftime map-dump histogram

# 输出示例
BUCKET(us)    COUNT
0-10          15234
10-20         8521
20-30         2341
30-40         567
40-50         123
50-60         45
>60           12

6.3 案例三：内存访问模式追踪

目标：分析 Kernel 的内存访问模式，优化缓存利用率。

// mem_trace.c
#include <bpf/bpf_helpers.h>

struct mem_event {
    u64 address;
    u32 thread_idx;
    u8 is_load;  // 1=load, 0=store
    u8 padding[3];
};

struct {
    __uint(type, BPF_MAP_TYPE_RINGBUF);
    __uint(max_entries, 1 << 24);
} mem_events SEC(".maps");

SEC("gpu_memcapture")
int trace_memory(struct pt_regs *ctx) {
    struct mem_event *event;
    event = bpf_ringbuf_reserve(&mem_events, sizeof(*event), 0);
    if (!event)
        return 0;
    
    event->address = ctx->mem_address;
    event->thread_idx = bpf_get_thread_idx();
    event->is_load = ctx->is_load;
    
    bpf_ringbuf_submit(event, 0);
    return 0;
}

char LICENSE[] SEC("license") = "GPL";

分析内存访问模式：

# analyze_memory.py
import pandas as pd
import matplotlib.pyplot as plt

# 读取 bpftime 导出的数据
events = pd.read_csv('mem_events.csv')

# 分析地址分布
plt.figure(figsize=(12, 6))
plt.subplot(1, 2, 1)
plt.hist(events['address'], bins=100)
plt.xlabel('Memory Address')
plt.ylabel('Access Count')
plt.title('Memory Access Distribution')

# 分析 Load/Store 比例
plt.subplot(1, 2, 2)
load_store = events['is_load'].value_counts()
plt.pie(load_store, labels=['Load', 'Store'], autopct='%1.1f%%')
plt.title('Load/Store Ratio')

plt.tight_layout()
plt.savefig('memory_analysis.png')

6.4 案例四：PyTorch 模型推理追踪

目标：追踪 PyTorch 模型推理过程中的 GPU Kernel 执行。

// pytorch_trace.c
#include <bpf/bpf_helpers.h>
#include <bpf/bpf_tracing.h>

struct kernel_event {
    u64 timestamp;
    char kernel_name[64];
    u32 grid_x, grid_y, grid_z;
    u32 block_x, block_y, block_z;
};

struct {
    __uint(type, BPF_MAP_TYPE_RINGBUF);
    __uint(max_entries, 1 << 26);
} events SEC(".maps");

SEC("gpu_kprobe")
int trace_pytorch_kernel(struct pt_regs *ctx) {
    struct kernel_event *event;
    event = bpf_ringbuf_reserve(&events, sizeof(*event), 0);
    if (!event)
        return 0;
    
    event->timestamp = bpf_get_globaltimer();
    bpf_probe_read_kernel_str(event->kernel_name, 64, ctx->kernel_name);
    event->grid_x = ctx->grid_dim_x;
    event->grid_y = ctx->grid_dim_y;
    event->grid_z = ctx->grid_dim_z;
    event->block_x = ctx->block_dim_x;
    event->block_y = ctx->block_dim_y;
    event->block_z = ctx->block_dim_z;
    
    bpf_ringbuf_submit(event, 0);
    return 0;
}

char LICENSE[] SEC("license") = "GPL";

加载到 PyTorch 应用：

# 启动 PyTorch 应用
python train.py &
PT_PID=$!

# 加载追踪程序
bpftime gpu-attach --pid $PT_PID pytorch_trace.o

# 运行一段时间后查看结果
bpftime cat-events pytorch_trace.o | head -20

# 输出示例
TIME         KERNEL_NAME              GRID         BLOCK
10:30:45.123 aten::addmm             (128,128,1)  (32,32,1)
10:30:45.124 aten::relu              (128,128,1)  (256,1,1)
10:30:45.125 aten::batch_norm        (64,1,1)     (1024,1,1)
...

七、与现有工具的对比

7.1 bpftime vs Nsight Systems

7.2 bpftime vs NVBit

7.3 bpftime 的独特价值

1. 可编程探针：不是预定义固定指标，而是让开发者按需编写观测逻辑
2. Per-thread 数据采集：从"聚合统计"走向"线程级诊断"
3. Kernel 内部插桩：不是只停留在 API 层或外部观测层
4. 内存访问模式追踪：分析复杂 GPU Kernel 的行为
5. 标准工具链兼容：使用 clang、libbpf、bpftrace

八、未来发展方向

8.1 短期目标

扩展 memcapture 能力：

• 增加寄存器值追踪
• 分析 shared memory 访问模式
• 支持 bank conflict 检测

AMD ROCm/HIP 后端：

• 覆盖 AMD GPU
• 支持 GCN/RDNA 指令集
• 实现跨厂商的统一 API

8.2 中期目标

实时流式 Profiling：

• 构建类似 Nsight 的在线可视化体验
• 支持远程监控
• 集成 Prometheus/Grafana

bpftrace 集成：

• 让 GPU 观测也具备 one-liner 的表达能力

# 未来可以这样使用
bpftrace -e 'gpu:kretprobe:* { @time[probe] = hist(duration); }'

8.3 学术成果

bpftime 的研究成果已经出现在多个顶级会议和期刊：

• OSDI '25: "Extending Applications Safely and Efficiently"
• Linux Plumbers 2023: Userspace eBPF Runtime
• USENIX ATC: uXDP 论文

这说明 bpftime 正在从工程实践走向系统化的方法论沉淀。

九、总结与展望

9.1 核心贡献

bpftime for GPU 代表的不仅仅是一个新工具，而是一种新的可能性：

如果 eBPF 可以深入 CPU 内核与用户态程序，那么它为什么不能继续进入 GPU Kernel？

通过 PTX 级插桩、per-thread 粒度观测、零拷贝数据回传以及标准 eBPF 工具链兼容，bpftime for GPU 正在把这个问题变成现实答案。

9.2 关键能力总结

9.3 对开发者的意义

1. AI 训练优化：深入理解 GPU Kernel 行为，优化训练效率
2. 科学计算调试：定位线程级问题，减少调试时间
3. 性能监控：构建实时的 GPU 可观测性系统
4. 故障诊断：在生产环境中动态注入探针

9.4 参与和贡献

bpftime 是开源项目，欢迎开发者参与：

• GitHub: github.com/eunomia-bpf/bpftime
• 文档: eunomia.dev/bpftime
• 社区: Linux eBPF 生态系统

随着 AI 时代 GPU 负载持续增长，这类能力的价值只会越来越大。

参考资料：

• bpftime 官方文档：https://eunomia.dev/bpftime
• eBPF 开发者大会 2026 演讲
• OSDI '25 论文：Extending Applications Safely and Efficiently
• NVIDIA PTX 文档
• bpftime GitHub 仓库