perf with gpu

@amitmund September 10, 2026

Linux perf on GPU — Complete Learning Notes & Output Guide

While perf is primarily known for CPU profiling, it interfaces directly with the Linux kernel's Direct Rendering Manager (DRM), DMA fence subsystems, and vendor GPU PMU drivers (such as Intel i915/xe and AMD amdgpu). It profiles the CPU-GPU boundary, command submission queues, rendering engines, and synchronization latency.

1. How perf Interacts with GPUs

Linux perf operates at the kernel-driver boundary. It does not replace vendor-specific shader profiling tools (like NVIDIA Nsight or AMD ROCm), but it is the primary native tool for diagnosing host-to-device bottlenecks, kernel submission stalls, and GPU engine utilization.

+-------------------------------------------------------------+
|                        USER SPACE                           |
|   App / ML Framework (PyTorch, Vulkan, OpenGL, OpenCL)      |
+-------------------------------------------------------------+
                               |
                               | ioctl / memory mappings
                               v
+-------------------------------------------------------------+
|                       KERNEL SPACE                          |
|   [DRM Core] <---> [DMA Fence API] <---> [GPU Driver Module]|
|                                                 |           |
|   ===========================================   |           |
|   ==> [ perf intercepts tracepoints / PMU ] =   |           |
|   ===========================================   |           |
+-------------------------------------------------|-----------+
                                                  v
+-------------------------------------------------------------+
|                      HARDWARE LAYER                         |
|   PCIe Bus <-----> Command Ring Buffers <-----> GPU Engines |
+-------------------------------------------------------------+

perf monitors GPU activity through two mechanisms:

  1. GPU Hardware PMUs (Performance Monitoring Units): Expose hardware engine busyness (Render, Compute, Video, Copy engines) and dynamic clock frequencies.
  2. Kernel DRM/DMA Tracepoints: Track command buffer submissions, page flips, and synchronization fences (dma_fence).

2. Discovering Available GPU Events

To see what GPU events your hardware and kernel expose, use perf list:

# List all DRM and DMA fence subsystem tracepoints
perf list 'drm:*' 'dma_fence:*'

# List Intel GPU PMU events (if running Intel integrated/discrete GPU)
perf list 'i915/*' 'xe/*'

# List AMD GPU tracepoints
perf list 'amdgpu:*'

3. Example 1: GPU Engine Monitoring with perf stat (Intel / DRM PMU)

On systems with supported DRM drivers (e.g., Intel i915), the kernel exposes GPU engine utilization as hardware PMU counters.

Command

sudo perf stat -e \
  i915/rcs0-busy/,i915/bcs0-busy/,i915/vcs0-busy/,i915/actual-frequency/ \
  -I 1000

Raw Output

#           time             counts unit events
     1.001045120         450.210 ms      i915/rcs0-busy/           #   45.02% busyness
     1.001045120          12.105 ms      i915/bcs0-busy/           #    1.21% busyness
     1.001045120           0.000 ms      i915/vcs0-busy/           #    0.00% busyness
     1.001045120            1150 MHz     i915/actual-frequency/
     2.002150340         980.450 ms      i915/rcs0-busy/           #   98.05% busyness
     2.002150340           5.400 ms      i915/bcs0-busy/           #    0.54% busyness
     2.002150340           0.000 ms      i915/vcs0-busy/           #    0.00% busyness
     2.002150340            1300 MHz     i915/actual-frequency/

Explanation of Output Columns & Metrics

Column / Field Example Technical Meaning
# time 1.001045120 Interval timestamp (in seconds) since perf was initiated (via -I 1000 for 1-second intervals).
counts 450.210 / 1150 The raw metric value recorded during that 1-second sampling window.
unit ms / MHz The measurement unit (ms of busy engine time per elapsed second, or hardware clock frequency in MHz).
events i915/rcs0-busy/ The kernel PMU counter representing a specific GPU engine.
# % busyness 45.02% Computed utilization ratio: $\frac{\text{Busy Time (ms)}}{\text{Interval Duration (1000ms)}} \times 100$.

Breakdown of Specific GPU PMU Engines

  • rcs0-busy (Render / Compute Engine): Time the primary 3D rendering and compute shader pipeline was executing commands. High numbers indicate compute-bound rendering or ML workloads.
  • bcs0-busy (Blitter / Copy Engine): Time the DMA copy engine spent transferring buffers between system RAM and GPU VRAM over PCIe.
  • vcs0-busy (Video Engine): Hardware video decode/encode engine activity (NVDEC/QuickSync equivalent).
  • actual-frequency: Real-time operating clock speed of the GPU core, exposing whether the chip is throttling due to thermals or power limits.

4. Example 2: Tracing GPU Synchronization Stalls with dma_fence

A common GPU performance issue occurs when the CPU sits idle waiting for the GPU to finish rendering (or vice-versa). Linux manages this via DMA fences.

Command

Record DMA fence latency across all processes for 5 seconds:

sudo perf record -e \
  dma_fence:dma_fence_init,\
  dma_fence:dma_fence_wait_start,\
  dma_fence:dma_fence_wait_end \
  -a sleep 5

Analyze the results:

sudo perf report --stdio

Raw Output

# Overhead  Command          Shared Object        Symbol
# ........  ...............  ...................  ......................................
#
    62.40%  game_engine      [kernel.kallsyms]    [k] dma_fence_wait_start
    25.10%  game_engine      [kernel.kallsyms]    [k] dma_fence_wait_end
     8.20%  Xorg             [kernel.kallsyms]    [k] dma_fence_init
     4.30%  kworker/u16:2    [kernel.kallsyms]    [k] dma_fence_signal

Detailed Breakdown of Headings

4.1 Overhead

  • Meaning: The percentage of captured tracepoint samples attributed to that specific event and code path.
  • Diagnostic Value: If dma_fence_wait_start commands more than 50% of trace samples, the application thread is blocked waiting for previously submitted GPU work to complete, indicating that the GPU is the bottleneck.

4.2 Command

  • Meaning: The name of the process/executable responsible for triggering the fence or wait.
  • Diagnostic Value: Identifies whether the client application (game_engine), the display server (Xorg / wayland), or kernel worker threads are stalling.

4.3 Shared Object

  • Meaning: The binary or library where the event occurred. [kernel.kallsyms] indicates an internal Linux kernel routine.

4.4 Symbol

  • Meaning: The specific kernel function or tracepoint hook executed.

5. Example 3: Tracing NVIDIA GPU Driver Calls

The proprietary NVIDIA Linux driver does not expose its internal chip counters to standard Linux kernel PMUs. However, perf can profile the user-to-kernel driver boundary (nvidia.ko and /dev/nvidia* system calls) to diagnose PCIe transfer overhead and driver stalls.

Command

Trace system calls and driver ioctl overhead for an NVIDIA workload:

sudo perf trace -e ioctl --filter 'fd == 3' python3 train_model.py

Or profile kernel-level CPU execution inside the NVIDIA driver:

sudo perf top --dsos=nvidia,nvidia_uvm

Raw Output

Overhead  Shared Object       Symbol
--------  ------------------  ----------------------------------------------
  48.12%  nvidia_uvm.ko       [k] _nvUvmCommitMemory
  28.45%  nvidia.ko           [k] os_flush_cpu_cache
  12.10%  nvidia_uvm.ko       [k] uvm_va_block_service_locked
   8.20%  nvidia.ko           [k] _nv012345rm

Explanation of Driver Symbols

  • nvidia_uvm.ko (_nvUvmCommitMemory / uvm_va_block_service_locked): Indicates Unified Virtual Memory (UVM) page faulting between CPU RAM and GPU VRAM. High values point to memory thrashing over the PCIe bus.
  • os_flush_cpu_cache: The CPU is flushing memory caches before DMA transfers can safely read host memory buffers.

6. perf vs. Vendor GPU Profilers

Diagnostic Capability Linux perf NVIDIA Nsight / NVPROF AMD ROCm rocprof
Primary Scope Linux Kernel, DRM, DMA fences, CPU context switches GPU SM warp occupancy, CUDA kernels, tensor cores AMD RDNA/CDNA wave execution, HIP kernels
Driver Support Open-source drivers (Intel i915/xe, AMDGPU, Nouveau) Proprietary NVIDIA Driver (nvidia.ko) Open-source ROCm driver stack
GPU Shader Counters No (engine level only) Yes (register usage, warp stalls) Yes (VALU/SALU utilization)
System-Wide Cost Negligible Moderate to High Low to Moderate
Host-to-Device Sync Excellent (dma_fence) Limited to CUDA API wrappers Limited to HIP/HSA events

7. Real-World Troubleshooting Scenarios

Scenario A: Diagnosing a "Stuttering" GPU Pipeline

An application maintains 60 FPS on average, but experiences micro-stutters:

sudo perf record -e drm:drm_vblank_event,dma_fence:dma_fence_wait_start -a sleep 10

Diagnosis: If dma_fence_wait_start timestamps regularly cross the 16.6ms boundary (the interval of drm_vblank_event at 60Hz), the GPU engine is failing to complete frame processing in time for the display refresh.

Scenario B: Verifying Hardware Video Encoding Acceleration

Confirming whether a transcoding daemon is utilizing the hardware media engine or falling back to CPU software rendering:

sudo perf stat -e i915/vcs0-busy/,i915/rcs0-busy/ -p $(pgrep ffmpeg)

Diagnosis:

  • vcs0-busy > 0%: Hardware video engine (QuickSync) is active.
  • vcs0-busy == 0% and rcs0-busy == 0%: FFMPEG is running pure CPU-based encoding (e.g., libx264).

8. Important Interview Questions & Answers

Q: Why can't Linux perf track NVIDIA streaming multiprocessor (SM) warp stalls directly?

Answer: Linux perf depends on standard Linux kernel hardware counter interfaces (perf_events) and open-source kernel tracepoints. NVIDIA's proprietary driver uses a closed user-space driver stack that interacts directly with GPU firmware registers, bypassing the Linux DRM subsystem and standard Linux PMUs. To profile internal SM warp states, memory pipelines, and tensor core utilization on NVIDIA hardware, you must use NVIDIA tools (such as Nsight Systems and Nsight Compute) via CUPTI (CUDA Profiling Tools Interface).

Q: What does high overhead in dma_fence_wait_start indicate in a GPU trace?

Answer: dma_fence_wait functions are kernel synchronization barriers where a CPU thread blocks waiting for a submitted GPU command buffer to finish execution. High overhead here indicates that the CPU has outrun the GPU and is idling at a synchronization barrier (e.g., glFinish(), vkQueueWaitIdle(), or cudaDeviceSynchronize()), demonstrating that the workload is GPU-bound.

0 Likes
2 Views
0 Comments

Filters

No filters available for this view.

Reset All