How to Profile GPU Kernels with PTX Probes: From CUDA Basics to Custom Instrumentation

Introduction

The author explores GPU kernel performance analysis, beginning with a motivation to achieve fine‑grained, user‑defined profiling similar to eBPF on CPUs. Using a CUDA matrix‑multiplication example, the article demonstrates how to identify bottlenecks such as non‑coalesced memory accesses and high global‑memory latency.

CUDA Architecture Basics

CUDA kernels execute on a hierarchy of SMs, warps, and threads. Each warp consists of 32 threads that share a common instruction stream. Memory is organized into global, shared, registers, and cache layers, each with different latency and bandwidth characteristics. Understanding this hierarchy is essential for effective optimization.

Performance Optimization Strategies

The article outlines four main strategies: maximizing parallel execution, improving instruction‑level parallelism, optimizing memory usage (including reducing non‑coalesced accesses and bank conflicts), and minimizing host‑device data transfers. A tiled matrix multiplication implementation using shared memory is presented to illustrate these concepts.

#include <iostream>
#include <vector>
#define TILE_SIZE 32
__global__ void gemm_gpu(const float* A, const float* B, float* C, int M, int N, int K) {
    __shared__ float sA[TILE_SIZE][TILE_SIZE];
    __shared__ float sB[TILE_SIZE][TILE_SIZE];
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float sum = 0.0f;
    for (int i = 0; i < (K + TILE_SIZE - 1) / TILE_SIZE; ++i) {
        sA[threadIdx.y][threadIdx.x] = (row < M && i * TILE_SIZE + threadIdx.x < K) ? A[row * K + i * TILE_SIZE + threadIdx.x] : 0.0f;
        sB[threadIdx.y][threadIdx.x] = (i * TILE_SIZE + threadIdx.y < K && col < N) ? B[(i * TILE_SIZE + threadIdx.y) * N + col] : 0.0f;
        __syncthreads();
        for (int k = 0; k < TILE_SIZE; ++k) sum += sA[threadIdx.y][k] * sB[k][threadIdx.x];
        __syncthreads();
    }
    if (row < M && col < N) C[row * N + col] = sum;
}

PTX Overview

PTX is an intermediate assembly language between CUDA C++ and the GPU’s native SASS. It provides a stable, forward‑compatible representation of kernel code. The article shows how to generate PTX with nvcc --keep -c and inspect it using cuobjdump -ptx. A snippet of the generated PTX for the matrix multiplication kernel is annotated to explain parameter loading, thread indexing, and memory accesses.

.visible .entry _Z8gemm_gpuPKfS0_Pfiii(
    .param .u64 _Z8gemm_gpuPKfS0_Pfiii_param_0,
    .param .u64 _Z8gemm_gpuPKfS0_Pfiii_param_1,
    .param .u64 _Z8gemm_gpuPKfS0_Pfiii_param_2,
    .param .u32 _Z8gemm_gpuPKfS0_Pfiii_param_3,
    .param .u32 _Z8gemm_gpuPKfS0_Pfiii_param_4,
    .param .u32 _Z8gemm_gpuPKfS0_Pfiii_param_5)
{
    .reg .pred %p9;
    .reg .f32 %f30;
    .reg .b32 %r14;
    ld.param.u64 %rd18, [_Z8gemm_gpuPKfS0_Pfiii_param_0];
    // ... (omitted for brevity) ...
    ret;
}

Neutrino Probe Framework

Neutrino brings an eBPF‑like programmable probing model to GPUs by inserting custom PTX snippets (Snippets) at specified Tracepoints. Probes consist of a Snippet, a Tracepoint (e.g., before or after a kernel), and a Map for storing results. The framework isolates probe registers and memory, ensuring non‑intrusive instrumentation.

A Python DSL defines a warp‑level timing probe that records the start clock, computes elapsed cycles, and saves the data to a map:

from neutrino import probe, Map
import neutrino.language as nl

@Map(level="warp", type="array", size=16, cap=1)
class block_sched:
    start: nl.u64
    elapsed: nl.u32
    cuid: nl.u32

start = nl.u64 = 0
elapsed = nl.u64 = 0

@probe(pos="kernel", level="warp", before=True)
def thread_start():
    start = nl.clock()

@probe(pos="kernel", level="warp")
def thread_end():
    elapsed = nl.clock() - start
    block_sched.save(start, elapsed, nl.cuid())

When executed with neutrino -p dmat python main.py, Neutrino hooks the CUDA driver, injects the probe into the kernel’s PTX, recompiles, and runs the instrumented binary. Results are stored in a trace directory containing original and probed binaries, PTX files, and binary map data.

Data Collection and Analysis

The binary map files follow a custom layout: a header, section descriptors, and per‑warp records (start time, elapsed cycles, SM ID). Python scripts parse these files using struct and reconstruct structured records. The article provides a full parsing script that aggregates per‑warp timings, computes average running and idle times, and prints a summary.

import struct
from neutrino import TraceHeader, TraceSection

class warp_duration(NamedTuple):
    start: int
    elapsed: int
    warpid: int

# Parsing logic omitted for brevity

Analysis of the example shows that warps are not launched in sequential order and that later‑launched warps experience higher latency due to resource contention. After applying tiling and reducing TILE_SIZE from 32 to 16, the warp execution times decrease noticeably, confirming the effectiveness of the optimization.

Conclusion

While Nsight Compute suffices for many profiling tasks, PTX‑level instrumentation via Neutrino offers deeper insight into kernel behavior, enabling custom metrics and fine‑grained analysis that are valuable for advanced GPU performance tuning.