GPU Probes

Probes are in-process metric instruments for GPU and host code — added by hand, or by an AI agent instrumenting the code it optimizes. Register an instrument once, record values on the hot path, and when the process runs under graphsignal-run, the profiler discovers the registry, reads the instruments every second, and publishes them to the local /signals endpoint — next to the built-in cuda_*, gpu_*, and engine metrics.

The record path is lock-free — a handful of relaxed 64-bit atomics, no allocation — and probes are inert when nothing reads them, so instrumented code ships to production unchanged.

When to use probes

The built-in cuda_kernels_nanoseconds profile already ranks GPU time per kernel symbol for free, and graphsignal-run --cuda-graph-trace node extends that ranking to the kernels inside a replayed CUDA graph — so for “which kernel is hot”, in any engine, reach for the profiles first and add no code at all.

Probes are for the two things CUPTI structurally cannot see:

• Phases inside one kernel — the stages of a megakernel, the steps of a fused op. CUPTI times a kernel launch; only code inside the kernel can time a part of it.
• Concepts that are not linker symbols — per layer, per op, per request stage, queue depths, batch sizes. A profile frame can be named whatever the question is.

The usual order is: rank kernels from the built-in profile (node mode if the engine replays graphs), then probe the kernel that dominates.

Probes are meant to ship: leaving them in production code is the intended use, since the record path is lock-free and allocation-free and the instruments are inert when nothing reads them. Cost follows how densely you record, not whether probes are present — a few instruments per request or per iteration is not measurable, while instrumentation dense enough to break one kernel into its phases (hundreds to thousands of records per iteration) costs a few percent. So register instruments once at init, probe the place you have already narrowed down to, and when the instrumentation is that dense, take headline numbers from an unprobed build. See Profiler Overhead for measured figures.

Vendor the header

The probe API is a single self-contained C++17 header with no dependencies beyond libc (Apache-2.0), designed to be vendored into your project:

curl -fsSL https://raw.githubusercontent.com/graphsignal/graphsignal/main/include/graphsignal/probe.h \
  -o third_party/graphsignal/probe.h --create-dirs

For device-side recording, also vendor the CUDA or HIP companion header:

curl -fsSL https://raw.githubusercontent.com/graphsignal/graphsignal/main/include/graphsignal/probe_cuda.h \
  -o third_party/graphsignal/probe_cuda.h
# or
curl -fsSL https://raw.githubusercontent.com/graphsignal/graphsignal/main/include/graphsignal/probe_rocm.h \
  -o third_party/graphsignal/probe_rocm.h

Compile with -std=c++17 or later (g++, clang, nvcc, hipcc) and add the vendored directory to your include path.

Instrument types

Type	Records	Appears in /signals as
GRAPHSIGNAL_GAUGE	the current value (graphsignal_set)	gauge: {"value": ...}
GRAPHSIGNAL_COUNTER	cumulative additions (graphsignal_add)	counter: {"total": ...}
GRAPHSIGNAL_HISTOGRAM	observed values (graphsignal_record)	histogram: exact count/sum/min/max, plus mean/p50/p95
GRAPHSIGNAL_PROFILE	cumulative counters per named frame	profile: frames sorted by value descending, each with its sample count

Values are unsigned 64-bit integers (nanoseconds, bytes, counts); gauges take a double. Instruments are cumulative forever — the profiler snapshots and diffs, so you never reset anything. Name metrics with underscores and a unit suffix, like the built-ins: myengine_batch_duration_nanoseconds, myengine_bytes_in.

Histogram bins are log-linear with 4 sub-bins per power of two (worst-case bucket width 25%), which is what bounds p50/p95. The count, sum, min and max beside them are exact, so sum / <iterations> is the true time per iteration of a probed region and count is how often it ran. Limits: 4096 instruments per process, 4 tags per instrument, 250 frames per profile.

Host code

Register instruments at startup and record from anywhere:

#include <graphsignal/probe.h>

static graphsignal_probe_entry* g_batch_ns;
static graphsignal_probe_entry* g_bytes_in;
static graphsignal_probe_entry* g_queue_depth;

void init() {
    g_batch_ns = graphsignal_probe_register(
        "myengine_batch_duration_nanoseconds", GRAPHSIGNAL_HISTOGRAM, NULL, NULL, 0);
    g_bytes_in = graphsignal_probe_register(
        "myengine_bytes_in", GRAPHSIGNAL_COUNTER, NULL, NULL, 0);
    g_queue_depth = graphsignal_probe_register(
        "myengine_queue_depth", GRAPHSIGNAL_GAUGE, NULL, NULL, 0);
}

void step(size_t request_bytes) {
    uint64_t t0 = graphsignal_now_ns();
    // ... work ...
    graphsignal_record(g_batch_ns, graphsignal_now_ns() - t0);  // histogram
    graphsignal_add(g_bytes_in, request_bytes);                 // counter
    graphsignal_set(g_queue_depth, queue.size());               // gauge
}

Registration is idempotent by (name, tags): registering the same instrument twice returns the same entry. Every record function is a safe no-op on NULL, so a failed registration never needs a branch at the call site.

Tags distinguish series of one name (up to 4 per instrument):

const char* keys[] = {"stage"};
const char* vals[] = {"prefill"};
g_prefill_ns = graphsignal_probe_register(
    "myengine_stage_duration_nanoseconds", GRAPHSIGNAL_HISTOGRAM, keys, vals, 1);

Profile instruments

A profile aggregates cumulative counters per named frame — time per operation, per layer, per anything you can name (up to 250 frames):

static graphsignal_probe_entry* g_ops;

void init() {
    g_ops = graphsignal_probe_register(
        "myengine_ops_nanoseconds", GRAPHSIGNAL_PROFILE, NULL, NULL, 0);
}

void run_op(const char* op_name) {
    uint64_t t0 = graphsignal_now_ns();
    // ... op ...
    graphsignal_profile_add_by_name(g_ops, op_name, graphsignal_now_ns() - t0);
}

graphsignal_profile_add_by_name does a linear name scan per call — fine for moderate rates. On hot paths, resolve the frame once and cache it:

static graphsignal_profile_frame* g_attn_frame;

void init() {
    g_attn_frame = graphsignal_profile_frame_get(g_ops, "attention");
}

void attention_op() {
    uint64_t t0 = graphsignal_now_ns();
    // ... op ...
    graphsignal_profile_add(g_attn_frame, graphsignal_now_ns() - t0);
}

In /signals, the profile reports its frames sorted by value descending, each with the number of recordings behind it — the same shape as the built-in cuda_kernels_nanoseconds profile.

CUDA kernels

Device-storage instruments live in GPU memory and are recorded from inside kernels. Register with <graphsignal/probe_cuda.h> and pass the device data pointer to your kernel:

#include <graphsignal/probe_cuda.h>

static graphsignal_probe_entry* g_tile_ns;

void init() {
    g_tile_ns = graphsignal_probe_register_cuda(
        "mykernels_tile_duration_nanoseconds", NULL, NULL, 0);
}

__global__ void my_kernel(..., graphsignal_instrument_data* tile_ns) {
    unsigned long long t0 = graphsignal_gtimer();
    // ... tile work ...
    if (threadIdx.x == 0) {
        GRAPHSIGNAL_RECORD(tile_ns, graphsignal_gtimer() - t0);
    }
}

// launch:
my_kernel<<<grid, block>>>(..., graphsignal_probe_device_data(g_tile_ns));

Device instruments are histograms: GRAPHSIGNAL_RECORD observes one value with the same relaxed-atomics path as the host API. graphsignal_gtimer() reads the device’s globaltimer register in nanoseconds. Record once per block (e.g. from threadIdx.x == 0), not per thread, unless you mean to observe per-thread values.

Registration allocates the data block with cudaMalloc and returns NULL on failure; graphsignal_probe_device_data(NULL) is NULL, and recording on NULL is a no-op in the kernel too — instrumented kernels are safe even when registration failed.

The profiler reads device instruments through the CUDA driver API from its injected library; nothing in your code copies them back.

HIP kernels

<graphsignal/probe_rocm.h> is the HIP analog — graphsignal_probe_register_rocm, graphsignal_probe_device_data_rocm, and the same GRAPHSIGNAL_RECORD/graphsignal_gtimer device API. Compile with hipcc. One difference: on HIP, graphsignal_gtimer() returns wall_clock64() ticks, not nanoseconds — convert with the device’s wall clock frequency (hipDeviceAttributeWallClockRate, kHz) or record raw ticks consistently and interpret downstream.

How discovery works

The header keeps a process-global registry exported as a C symbol (__graphsignal_probe_registry_v1). When the process runs under graphsignal-run, the injected profiling library finds it via dlsym, snapshots every instrument once per second, and publishes the values. Multiple libraries in one process share the registry; without the profiler, the registry is just a few pages of memory that nothing reads.

Verify

Run your instrumented workload under the profiler and read your instruments back:

graphsignal-run ./my_app
curl -s http://127.0.0.1:18259/signals | grep myengine

Probe metrics appear under their registered names, tagged with the process that recorded them.