✨ feat: release v0.4.0 profiling primitives

.github/workflows/hardware.yml

@@ -24,3 +24,5 @@ jobs:
       - run: just test-injection-live
       - run: just test-multisource-live
       - run: just benchmark-gpu
+      - run: just benchmark-aggregate
+      - run: just benchmark-multiprocess

AGENTS.md

@@ -19,6 +19,9 @@
 - Do not add a daemon or persistent session service without a concrete workflow
   that cannot be expressed by the four bounded commands.
 - Keep eBPF and CUPTI callback hot paths bounded and free of blocking I/O.
+- Resolve architecture-specific syscall names in core. In BPF, filter process
+  identity and syscall number before reading scalar registers or reserving a
+  record. Arm multi-link collectors only after every link is attached.
 
 ## Failures and safety
 
@@ -30,8 +33,16 @@
 - `validate` is read-only. `measure` may inject CUPTI only after validation
   reports `injection_required`; log the mutation and include a JSON warning.
 - Stop CUPTI logically after collection. Do not `dlclose` the injected agent.
+- NVTX callback routing must be initialized before the target's first NVTX API
+  call. Once initialized, keep the CUPTI subscriber mapped and registered;
+  logical stop disables callback domains instead of unsubscribing it.
+- Correlate NVTX IDs as ordered lifecycles. Thread range keys reuse push levels,
+  and a bounded capture may end with one range open; retain earlier closed
+  pairs and report only the trailing start as unmatched.
 - Never collect pointer-referenced payloads, environments, or GPU buffer data by
-  default, and never describe temporal correlation as exact causality.
+  default. Named tracepoints retain identity and timestamps unless a versioned
+  scalar payload is explicitly designed. Never describe temporal correlation
+  as exact causality.
 
 ## Agent workflow
 
@@ -65,10 +76,14 @@
 - Run `just fmt-check`, `just lint`, and `just test` for Rust changes.
 - Run `just test-bpf-live` for BPF attachment changes.
 - Run `just test-cupti-live` for CUPTI ABI or callback changes.
+- Run `just test-nvtx-live` for NVTX callback, filtering, or lifecycle changes.
 - Run `just test-injection-live` for injection or agent lifecycle changes.
 - Run `just test-multisource-live` for host/GPU orchestration changes.
 - Run `just test-agent-contract` for CLI, schema, docs, or Skill changes.
 - Test bundled Skill scripts with deterministic fixtures and include them in
   `just test-agent-contract`.
 - Run `just benchmark-gpu` for callback hot-path changes.
+- Run `just benchmark-aggregate` for aggregate inventory hot-path or capacity
+  changes.
+- Run `just benchmark-multiprocess` for concurrent worker orchestration changes.
 - Use emoji conventional commits, for example `🐛 fix: restore target registers`.

CMakeLists.txt

@@ -1,6 +1,6 @@
 cmake_minimum_required(VERSION 3.20)
 
-project(xprobe VERSION 0.3.3 LANGUAGES C)
+project(xprobe VERSION 0.4.0 LANGUAGES C)
 
 include(CTest)
 

Cargo.toml

@@ -10,14 +10,15 @@ members = [
 resolver = "2"
 
 [workspace.package]
-version = "0.3.3"
+version = "0.4.0"
 edition = "2024"
 license = "Apache-2.0"
 repository = "https://github.com/itdevwu/xprobe"
 rust-version = "1.85"
 
 [workspace.dependencies]
 clap = { version = "4.5", features = ["derive"] }
+cpp_demangle = "0.4"
 libbpf-rs = { version = "0.26.2", features = ["vendored"] }
 nix = { version = "0.31.3", features = ["process", "ptrace", "signal", "uio"] }
 object = { version = "0.36", default-features = false, features = ["read", "std"] }

README.md

@@ -10,9 +10,10 @@
 </div>
 
 `xprobe` is an AI harness for measuring latency between two observable events
-in a process, on the CPU, NVIDIA GPU, or across both. Its bounded native profiler
-combines eBPF uprobes and NVIDIA CUPTI with an agent-friendly CLI, strict JSON
-contracts, explicit correlation quality, and no daemon or server lifecycle.
+in a process, on the CPU, NVIDIA GPU, or across both. Its bounded native
+profiler combines eBPF function, syscall, and tracepoint evidence with NVIDIA
+CUPTI, an agent-friendly CLI, strict JSON contracts, explicit correlation
+quality, and no daemon or server lifecycle.
 
 ## Install
 
@@ -21,7 +22,7 @@ only installation action required from the user:
 
 ```bash
 npx skills@1 add \
-  https://github.com/itdevwu/xprobe/tree/v0.3.3/skills/xprobe-measure-latency \
+  https://github.com/itdevwu/xprobe/tree/v0.4.0/skills/xprobe-measure-latency \
   --global
 ```
 
@@ -48,6 +49,11 @@ xprobe validate --pid 4242 \
   --to 'cuda:kernel_start:name~flash.*' \
   --match exact --json --non-interactive --no-color
 
+xprobe measure --pid 4242 \
+  --from 'cuda:kernel_start' --to 'cuda:kernel_end' \
+  --match exact --aggregate --duration-ms 1000 --max-groups 4096 \
+  --json --non-interactive --no-color
+
 xprobe measure --pid 4242 \
   --from 'cuda:runtime_api:cudaLaunchKernel:exit' \
   --to 'cuda:kernel_start:name~flash.*' \
@@ -57,8 +63,11 @@ xprobe measure --pid 4242 \
 ```
 
 Kernel launch latency is only one event pair. The same workflow measures host
-function spans, CUDA API calls, GPU operation durations, transfers, and paths
-across CPU and GPU events after selecting the correct CUDA worker.
+function spans, syscall latency, named Linux events, CUDA API calls, GPU
+operation durations, transfers, NVTX application ranges, and paths across CPU
+and GPU events after selecting the correct process. Aggregate mode provides a
+bounded coarse inventory of GPU operations before an exact evidence
+measurement narrows the question.
 
 `measure` also accepts completed `--input` captures and versioned live
 `--spec` files. Evidence can be exported as `jsonl` or `chrome`. JSON results
@@ -79,16 +88,19 @@ is also preserved when correlation or clock validation fails.
 `measure --pid` automatically loads the matching CUDA 12 or CUDA 13 CUPTI Agent
 when a selected endpoint requires it. It reports the target mutation on stderr
 and in JSON, disables collection afterward, and leaves the shared object mapped
-for safe reactivation.
+for safe reactivation. NVTX ranges are the exception: set
+`NVTX_INJECTION64_PATH` before the target's first NVTX call because online
+attach cannot retrofit an initialized NVTX dispatch.
 
 ## Support
 
-| Surface | 0.3.3 support |
+| Surface | Current support |
 | --- | --- |
 | OS/architecture | Linux x86_64, glibc 2.34 or newer |
-| Host events | ELF function entry/return through PID-scoped uprobes |
+| Host events | PID-scoped ELF function, named syscall, and tracepoint boundaries |
 | CUDA callbacks | Runtime and Driver API entry/exit |
 | GPU activity | Kernel, memcpy, and memset start/end |
+| Application ranges | Bounded ASCII NVTX thread and process ranges |
 | CUDA/CUPTI | 12.x and 13.x with automatic major selection |
 | Correlation | exact, first-after, nearest, stack-nested, stream-order |
 | Online injection | same mount namespace; ptrace permission required |

benchmarks/cuda-aggregate/cuda_aggregate_benchmark.cu

@@ -0,0 +1,63 @@
+#include <cuda_runtime.h>
+
+#include <cstdio>
+#include <unistd.h>
+
+__global__ void xprobe_aggregate_primary(int *output)
+{
+    *output += 1;
+}
+
+__global__ void xprobe_aggregate_secondary(int *output)
+{
+    *output += 2;
+}
+
+static int report_cuda_error(const char *operation, cudaError_t result)
+{
+    std::fprintf(stderr, "%s failed: %s\n", operation, cudaGetErrorString(result));
+    return 13;
+}
+
+int main(int argc, char **argv)
+{
+    if (argc != 3) {
+        std::fprintf(stderr, "usage: %s <ready-file> <stop-file>\n", argv[0]);
+        return 2;
+    }
+    cudaError_t result = cudaSetDevice(0);
+    if (result != cudaSuccess) {
+        return report_cuda_error("cudaSetDevice", result);
+    }
+    int *device_output = nullptr;
+    result = cudaMalloc(&device_output, sizeof(*device_output));
+    if (result != cudaSuccess) {
+        return report_cuda_error("cudaMalloc", result);
+    }
+
+    FILE *ready = std::fopen(argv[1], "w");
+    if (ready == nullptr || std::fclose(ready) != 0) {
+        std::perror("ready file");
+        return 5;
+    }
+    while (access(argv[2], F_OK) != 0) {
+        xprobe_aggregate_primary<<<1, 1>>>(device_output);
+        xprobe_aggregate_secondary<<<1, 1>>>(device_output);
+        result = cudaGetLastError();
+        if (result == cudaSuccess) {
+            result = cudaDeviceSynchronize();
+        }
+        if (result != cudaSuccess) {
+            return report_cuda_error("kernel workload", result);
+        }
+    }
+
+    result = cudaFree(device_output);
+    if (result == cudaSuccess) {
+        result = cudaDeviceReset();
+    }
+    if (result != cudaSuccess) {
+        return report_cuda_error("CUDA shutdown", result);
+    }
+    return 0;
+}