Bridging the GPU Observability Gap: Why eBPF on GPUs Matters

Why GPU Observability Is Hard

GPUs dominate machine‑learning, scientific, and high‑performance workloads, but their SIMT execution model and deep memory hierarchy make it difficult to observe and debug kernel behavior. Threads are grouped into warps, scheduled on streaming multiprocessors (SMs), and run asynchronously to the host, hiding launch overhead, memory stalls, warp divergence, and SM utilization from traditional CPU‑side tracing tools.

Limitations of Existing Tools

Current GPU profilers fall into two categories. First, CPU‑GPU boundary tools intercept CUDA/ROCm library calls (e.g., via LD_PRELOAD) and can only see host‑side events such as kernel launch and memory transfers, treating the GPU as a black box. Second, vendor‑specific analyzers (NVIDIA CUPTI, Nsight Compute, Intel GTPin, AMD ROCProfiler, NVBit, Neutrino) collect hardware counters and warp‑level traces on the device, but they operate in isolated ecosystems, lack integration with Linux eBPF probes, impose high overhead (10‑100× slower for fine‑grained tracing), and cannot dynamically filter or adapt probing logic without recompiling the application.

Timeline Visibility Gap

In a synchronous CUDA workflow, a developer can measure each API call’s wall‑clock time and infer whether data transfer or computation dominates. For example, cudaMemcpy() takes 200 µs while cudaDeviceSynchronize() takes 115 µs, suggesting a PCIe bottleneck. However, the 115 µs aggregate hides three hidden GPU phases: ~5 µs launch overhead, ~100 µs kernel execution, and ~10 µs cleanup. Without warp‑level visibility, developers cannot tell whether the slowdown is due to launch cost, low SM occupancy, memory divergence, or warp divergence.

In asynchronous execution, all API calls return immediately and only a final cudaStreamSynchronize() blocks, collapsing the entire timeline into a single 456 µs measurement. The loss of per‑stage timing makes it impossible to answer “Is the bottleneck memory transfer or kernel execution?” and obscures the impact of batch size, stream dependencies, or resource contention.

Why eBPF Is the Right Solution

eBPF has become the foundation of modern observability, networking, and security on Linux because it offers programmable, low‑overhead hooks that can be attached to thousands of kernel events. Extending this model to GPUs allows developers to inject custom logic directly into GPU kernels, collect per‑warp execution and memory‑access statistics, and dynamically adapt behavior at runtime without recompiling the application.

bpftime Architecture

bpftime implements a CUDA/ROCm attachment pipeline that injects eBPF programs into GPU kernels:

CUDA/OpenCL Runtime Hooks : Uses LD_PRELOAD and the bpftime runtime to intercept calls to the CUDA/ROCm libraries, gaining control over kernel launches and other GPU operations.

eBPF‑to‑PTX/SPIR‑V JIT : When a kernel starts, bpftime retrieves the eBPF bytecode, JIT‑compiles it to the target ISA (PTX for NVIDIA, SPIR‑V for AMD/Intel), and prepares it for injection.

Binary Detection and Injection : The compiled eBPF code is inserted into the kernel’s binary (e.g., PTX) before execution, enabling native execution inside the GPU kernel.

Helper‑Function Trampolines : Provides GPU‑accessible eBPF helper functions (e.g., map access, timestamp, ring‑buffer output) implemented as trampolines.

Shared Data Structures : BPF maps and ring buffers are allocated in memory accessible to both host and device, allowing efficient data exchange.

The design delivers:

3‑10× faster detection than NVBit.

Vendor‑neutral support for NVIDIA, AMD, and Intel GPUs.

Unified observability and control together with existing Linux eBPF probes (kprobes, uprobes).

Nanosecond‑scale per‑warp and per‑instruction performance data.

Dynamic, runtime‑adaptive instrumentation without source changes or recompilation.

Example Tools

kernelretsnoop attaches to CUDA kernel exit and records a nanosecond timestamp for each GPU thread, exposing warp divergence and memory‑access patterns that traditional profilers miss.

// eBPF program run on GPU kernel exit
SEC("kretprobe/_Z9vectorAddPKfS0_Pf")
int ret__cuda() {
    u64 tid_x, tid_y, tid_z;
    bpf_get_thread_idx(&tid_x, &tid_y, &tid_z); // which thread?
    u64 ts = bpf_get_globaltimer();           // when did it finish?
    bpf_perf_event_output(ctx, &events, 0, &data, sizeof(data));
}

threadhist uses a GPU array map to count how many times each thread executes, revealing load imbalance.

// eBPF program run on GPU kernel exit
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 *count = bpf_map_lookup_elem(&thread_counts, &tid_x);
    if (count) {
        __atomic_add_fetch(count, 1, __ATOMIC_SEQ_CST);
    }
}

launchlate measures the latency between the host’s cudaLaunchKernel() call and the actual start of kernel execution on the GPU, exposing queue delays and stream dependencies.

BPF_MAP_DEF(BPF_MAP_TYPE_ARRAY, launch_time);

// CPU‑side uprobe captures launch request
SEC("uprobe/app:cudaLaunchKernel")
int uprobe_launch(struct pt_regs *ctx) {
    u64 ts_cpu = bpf_ktime_get_ns(); // when host requested launch?
    bpf_map_update_elem(&launch_time, &key, &ts_cpu, BPF_ANY);
}

// GPU‑side kprobe captures execution start
SEC("kprobe/_Z9vectorAddPKfS0_Pf")
int kprobe_exec() {
    u64 ts_gpu = bpf_get_globaltimer(); // when GPU actually started?
    u64 *ts_cpu = bpf_map_lookup_elem(&launch_time, &key);
    u64 latency = ts_gpu - *ts_cpu; // queue wait time
    u32 bin = get_hist_bin(latency);
    // update histogram …
}

Key Components

CUDA runtime hooks (Frida‑based dynamic detection).

PTX modification to embed eBPF bytecode.

Helper trampolines for map access, timing, and context queries.

Host‑GPU communication via shared memory and spin‑locks.

References

bpftime OSDI ’25 paper.

CUDA Runtime API documentation.

PTX ISA specification.

eBPF documentation.

eGPU: Extending eBPF programmability to GPUs.

NVBit: Dynamic binary instrumentation framework for NVIDIA GPUs.