5 — Where eBPF meets CUDA

Requirements: chapters 2 to 4, and ideally Linux, Observed. For the experiments: root on a Linux test machine, an NVIDIA GPU and bpftrace. Never run these on a production or shared host.

Linux, Observed rests on one habit: don’t guess, attach an instrument. This chapter applies it to the GPU.

The blind spot

When a PyTorch script stalls, top shows a CPU at 100%. nvidia-smi shows utilization sampled over a period between 1 second and 1/6 second, depending on the product (nvidia-smi docs). Neither tells you why.

Here’s the part that makes it hard. libcuda.so writes GPU commands (methods) into a pushbuffer in memory, adds a GPFIFO entry pointing at them, advances GP_PUT, and rings a doorbell with a single memory-mapped write. The evidence comes from the kernel’s own channels in uvm_pushbuffer.h, because libcuda.so itself is closed source. The common path doesn’t enter the kernel at all, so Linux doesn’t see it happen.

+-------------------------------------------------------------------------+
|                        USER SPACE (PyTorch / vLLM)                      |
|  [Python] -> [libtorch_cuda.so] -> [libcudart.so] -> [libcuda.so]       |
|                         (Layer 1: eBPF uprobes)                         |
+------------------------------------------|----------------|-------------+
                                           | ioctl()        | doorbell
+------------------------------------------|----------------|-------------+
|                         LINUX KERNEL SPACE                              |
|  [nvidia.ko] -> [nvidia-uvm.ko] -> [PCIe/NVLink]                        |
|                  (Layer 2: eBPF kprobes)                                |
+-------------------------------------------------------------------------+
                                           |
                                           v
+-------------------------------------------------------------------------+
|                           GPU SILICON (SMs)                             |
|                  (Layer 3: research instrumentation)                    |
+-------------------------------------------------------------------------+

Tracing the runtime with uprobes

A uprobe fires when a program reaches a function in a library, with no code changes. Two useful targets:

• cudaLaunchKernel shows each launch’s grid and block size.
• cudaMalloc forces the GPU to synchronize across streams (NVIDIA), so it’s a good thing to time. NVIDIA doesn’t publish a duration, so measure yours.

Three cautions from real use:

• The library is versioned and may be bundled, so find it with grep libcudart /proc/$PID/maps.
• Whether <<<>>> reaches cudaLaunchKernel or __cudaLaunchKernel is unsettled, so count *cudaLaunch* first.
• A uretprobe on cudaStreamSynchronize gives only an upper bound on kernel time, if that’s the wait path used. For real timestamps, use CUPTI (Polar Signals).

How a launch’s arguments arrive

dim3 is a 12-byte struct, and on x86-64 it’s split across registers. The grid’s x and y share one register and z sits in the next. So arg1 holds both x and y, in its low and high halves. This was checked by compiling a matching declaration, repeated in Experiment 1. Only x86-64 has been checked; Grace (Arm) needs separate analysis.

Tracing the driver with kprobes

NVIDIA’s open kernel modules let you read real function names. cudaMallocManaged memory uses Unified Memory, so oversubscribing VRAM triggers page faults serviced by nvidia-uvm.ko. Useful targets:

• uvm_parent_gpu_service_replayable_faults
• uvm_parent_gpu_service_non_replayable_fault_buffer
• uvm_tools_record_thrashing, whose second argument is the address

Names change between drivers, so always check them first (nvidia-uvm source).

What the research tools do

The boundary is moving fast. Keep these three projects separate:

• eGPU is a paper presented at HCDS ‘25.
• bpftime is a runtime whose GPU support is “still experimental”. It JIT-compiles eBPF to PTX and rewrites a program’s PTX to add probes. SPIR-V is a stated direction, but PTX is the confirmed path. It uses per-thread ring buffers.
• gpu_ext (paper) brings struct_ops-style policy to the GPU driver. It ships a modified NVIDIA module (v575.57.08), though policies can then load without further rebuilds. It sets attributes through kfuncs, and its scheduling ops are limited to task_init and task_destroy.

We found no independent replication of the performance figures for bpftime or gpu_ext, so treat them as the authors’ claims.

How much does instrumentation cost? NVIDIA’s NVBit has to save and restore application state. Its cost depends on what you measure: light timing tools can cost under 2× in one microbenchmark, while per-instruction tools averaged 36× in NVBit’s own paper and reached 112× at worst.

Don’t build on these. %globaltimer is “intended for use by NVIDIA tools” and “may change or be removed”, and %smid can change during execution (PTX ISA).

The verifier, the JIT and the ring buffer here are the same ideas Kernel, Traced follows through the Linux source.

Hands-on

Experiment 1: prove the dim3 layout

Needs: clang, no GPU.

1. Create t.cpp:

   struct dim3 {
       unsigned x, y, z;
       dim3(unsigned a, unsigned b, unsigned c) : x(a), y(b), z(c) {}
   };
   extern "C" int cudaLaunchKernel(const void *f, dim3 g, dim3 b, void **a,
                                   unsigned long sm, void *s);
   int go() { return cudaLaunchKernel(0, dim3(11, 22, 33), dim3(44, 55, 66), 0, 77, 0); }

2. Generate assembly:

   clang++ -target x86_64-linux-gnu -O1 -S -o - t.cpp -nostdinc++ -fno-exceptions -masm=intel

3. Expected: edx should hold 33 and r8d 66, and two 64-bit constants should decode to (y<<32)|x. 77 goes on the stack.

Experiment 2: count and decode launches

Needs: a GPU, root and bpftrace.

1. Start a CUDA program that keeps running, then find its library:

   grep libcudart /proc/$PID/maps

2. List the launch entry points, using your real path:

   sudo bpftrace -l 'uprobe:/full/path/libcudart.so.NN:*cudaLaunch*'

3. Count across all of them:

   sudo bpftrace -e 'uprobe:/full/path/libcudart.so.NN:*cudaLaunch* { @[probe] = count(); }'

4. Decode the plain symbol:

   sudo bpftrace -e 'uprobe:/full/path/libcudart.so.NN:cudaLaunchKernel {
     printf("grid=%u,%u,%u block=%u,%u,%u\n",
       arg1 & 0xffffffff, arg1 >> 32, arg2 & 0xffffffff,
       arg3 & 0xffffffff, arg3 >> 32, arg4 & 0xffffffff); }'

5. Expected: launch a kernel with a known grid and block size a known number of times. Counts and values must match what you wrote.

If you see zero: check the other symbols from step 3. Under vLLM, zero may mean CUDA Graphs replay rather than a broken probe, so compare with cudaGraphLaunch.

Experiment 3: time cudaMalloc

Needs: a GPU, root and bpftrace.

1. Write a small CUDA loop that allocates and frees repeatedly. PyTorch’s allocator reuses memory, so don’t start there.

2. Attach a timer:

   sudo bpftrace -e 'uprobe:/full/path/libcudart.so.NN:cudaMalloc { @s[tid] = nsecs; }
   uretprobe:/full/path/libcudart.so.NN:cudaMalloc /@s[tid]/ {
     @us = hist((nsecs-@s[tid])/1000); delete(@s[tid]); }'

3. Look for: the shape of the distribution on your machine. Record your driver version with the result.

Experiment 4: watch page faults

Needs: a GPU, root and the nvidia-uvm module.

1. Confirm the names on your driver:

   sudo bpftrace -l 'kprobe:*' | grep -E 'uvm_parent_gpu_service|uvm_tools_record_thrashing'

2. Count fault servicing, then capture thrashing addresses:

   sudo bpftrace -e 'kprobe:uvm_parent_gpu_service_replayable_faults { @ = count(); }'
   sudo bpftrace -e 'kprobe:uvm_tools_record_thrashing { @[arg1] = count(); }'

3. Allocate with cudaMallocManaged and touch the memory alternately from CPU and GPU.

4. Expected: counts should stay near zero with plain cudaMalloc. Thrashing addresses should fall in your managed allocation. No reliable workload for triggering thrashing was found, so this step may take some experimenting.