rccl-tests/gin-sdma: in-device (wall_clock64) A2A timing mechanism by dlamd1dai · Pull Request #1 · dlamd1dai/rocm-systems

dlamd1dai · 2026-08-12T19:51:19Z

Summary

Adds device-side (in-kernel wall_clock64) timing for the GIN-SDMA AllToAll collective, measuring pure GPU device-function latency (excludes host launch / graph capture / teardown overhead).

gin_sdma_devtime.h — portable gin_devtime::measure() scaffold: queries the GPU wall-clock rate, allocates per-CTA start/end stamps, launches a persistent skip+loop timed kernel, reduces min(start)..max(end) over CTAs, takes the slowest local GPU, then MPI_MAX across ranks. Fails if wall-clock rate query fails (no silent fallback).
alltoall.cu — factors the GIN/Hybrid collective bodies into reusable __device__ functions, adds persistent GinAlltoAllTimedKernel/HybridAlltoAllTimedKernel loop variants, and implements the AlltoAllDeviceTime hook (CLI-controlled loop/skip/tier overrides; selects the timed kernel via deviceImpl).
common.h / common.cu — adds the optional testColl::deviceTime hook and exposes deviceImpl (non-static) so per-collective hooks can pick the matching timed kernel. Integrates device-time modes into BenchTime:
- --device_timing 1 (-B 1) — augment: normal run plus an extra #[a2a-devtime] line.
- --device_timing 2 (-B 2) — device-time-only: reports in-kernel device latency as the metric (skips the host graph loop for the out-of-place pass).
CMakeLists.txt — hipify/copy gin_sdma_devtime.h with the collective tests.

CLI flags (replaces env vars; per @abouteiller review)

Flag	Purpose	Default
`-B` / `--device_timing`	0=off, 1=augment, 2=device-only	0
`-L` / `--devtime_loop`	Timed-kernel loop iterations	10
`-P` / `--devtime_skip`	Timed-kernel warmup skip iterations	10
`-W` / `--devtime_loop_mid`	Loop at per-peer ≥ 8 MiB (0=use `-L`)	0
`-Q` / `--devtime_loop_large`	Loop at per-peer ≥ 64 MiB (0=use tier below)	0
`-j` / `--devtime_skip_mid`	Skip at mid tier (-1 → min(`-P`, 2))	-1
`-k` / `--devtime_skip_large`	Skip at large tier (-1 → min(`-P`, 1))	-1
`-H` / `--devtime_check`	Validate timed-kernel output before datacheck	0

Example (mode 1 augment, MI355 sweep):

alltoall_perf -D 3 -V 8 -g 1 -R 2 -B 1 -L 20 -P 5 -b 128 -e 4G -f 2 -w 5 -n 20

JIRA ID: AICOMRCCL-1458

AICOMRCCL-1459, AICOMRCCL-1460

Stacking

Base branch: develop. Depends on upstream ROCm/rocm-systems#10420 (128 MiB Anvil-SDMA put cap), merged 2026-08-24. This PR adds the devtime commits on top of develop; ready for upstream PR to ROCm/rocm-systems.

Rebased onto upstream develop @ d787b5f (2026-08-24). Devtime tip: 82460f9. Zero diff in projects/rccl/; +478 lines in 5 rccl-tests files only.

Build gate

The timed path (AlltoAllDeviceTime + the *TimedKernel templates) is compiled only under ENABLE_DEVICE_API and NCCL_VERSION_CODE >= NCCL_VERSION(2,28,7); otherwise a #else stub is compiled instead. ENABLE_DEVICE_API defaults OFF in install.sh/Makefile, so a default build does not compile this code. Build with --enable-device-api (cmake -DENABLE_DEVICE_API=ON) against RCCL >= 2.28.7 to exercise it; CI covering this collective should do the same.

Compile-verified with -DENABLE_DEVICE_API=ON:

gfx942 (MI300X): alltoall_perf builds & links; AlltoAllDeviceTime + GinAlltoAllTimedKernel/HybridAlltoAllTimedKernel present; Allreduce<T> cross-thread reduction resolves at link time.
gfx950 (MI355X, ROCm 7.13 docker): BUILD SANITY OK — GIN host plugins in librccl.so, Anvil SDMA factory linked, AlltoAllDeviceTime present in alltoall_perf. GIN Anvil-SDMA runtime smoke OK. Use the default device linker (do not pass --no-device-linker to RCCL install.sh on ROCm 7.13; otherwise cooperative_groups::this_cluster() duplicate symbols fail the link).

Test plan

Compile-verified on MI355X / gfx950 (docker, ROCm 7.13): alltoall_perf builds & links (AlltoAllDeviceTime + timed kernels present), GIN Anvil-SDMA smoke OK.
Mode 1 (-B 1 -L 20 -P 5, augment) on MI355X: normal numbers plus #[a2a-devtime] line per size (128 B–4 GB); datacheck Out of bounds values : 0 OK.
Mode 2 (-B 2 -L 20 -P 5, device-time-only) on MI355X: reported metric is in-kernel device latency; #wrong=0 across 128 B–4 GB.
Mode 2 + -H 1 on MI355X: timed-kernel output validated against expected — 0 wrong elements across all sizes through 4 GB; datacheck clean.
Mode 1/2 + -H 1 validated earlier on MI300X (see below; pre-rebase stack).

Review updates

@pvallem

measure() split into launch-all then sync-all passes to avoid the -g N>1 device-barrier deadlock.
Cross-rank/thread reduction routed through the harness Allreduce() helper (MPI serialized to the last thread; also combines across threads) instead of a direct MPI_Allreduce on MPI_COMM_WORLD.
AlltoAllDeviceTime gated to deviceImpl 3/4 (GIN/Hybrid) only.
skip clamped to >= 0 so the start stamp is always written.
Opt-in --devtime_check validates the timed kernel's output against expected before the datacheck loop re-inits.
Per-GPU wall-clock rate query; fail on bad rate (no khz=1 fallback).

@abouteiller

Device timing controlled via CLI flags (generic names, not ANVIL-specific env vars).
Removed legacy NCCL_GIN_ANVIL_A2A_DEVICE_TIMING and all other devtime env vars.
Replaced AUTOREDUCE with explicit --devtime_loop_mid / --devtime_loop_large tier knobs.

Runtime validation (MI355X / gfx950, GIN Anvil-SDMA, CLI flags)

8-GPU MI355X node (smci355-ccs-aus-m03-17), docker image rccl-gin-devtime-pr9927:latest (ROCm 7.13, -DENABLE_DEVICE_API=ON). Branch: users/dondai/gin-sdma-a2a-devtime @ 82460f9 (rebased onto upstream develop / ROCm#10420).

Config: NCCL_GIN_TYPE=6, -D 3 -V 8 -g 1 -R 2 -w 5 -n 20, sweep 128 B–4 GB (-b 128 -e 4G -f 2).

Gate	CLI	Result
Docker build + smoke	—	PASS — GIN smoke OK (`devtime-docker-build-20260821-*.log`)
Test#5 mode 1 (augment)	`-B 1 -L 20 -P 5`	PASS — `#wrong=0`, 26 `#[a2a-devtime]` lines; avg busbw 142.9 GB/s (host)
Test#5 mode 2 (device-only)	`-B 2 -L 20 -P 5`	PASS — `#wrong=0`; avg busbw 131.8 GB/s (device-time metric)
Test#5 mode 2 + datacheck	`-B 2 -L 20 -P 5 -H 1`	PASS — timed-kernel check 0 wrong; final datacheck `#wrong=0`
4 GB peak (mode 1)	—	host 419.7 GB/s busbw; `#[a2a-devtime]` 349.9 GB/s busbw @ 10.74 ms devtime

Log: ~/rocm-systems/ddai-artifacts/logs/devtime-a2a-8g-128B-4G-20260821-152042.log

Runtime validation (MI300X / gfx942, earlier stack)

Ran on 8-GPU MI300X with alltoall_perf -D3 (NCCL_GIN_TYPE=5 on the pre-pr9927 stack; current stack uses NCCL_GIN_TYPE=6), -n 8 -g 1, 128 B–1 MiB:

Mode 1 (-B 1, augment): normal numbers plus a #[a2a-devtime] line per size; datacheck Out of bounds values : 0 OK.
Mode 2 (-B 2, device-time-only): reported metric is the in-kernel device latency and tracks the mode-1 #[a2a-devtime] value (e.g. 1 MiB: device ~40.9 µs vs host-graph ~28.7 µs); datacheck clean.
Mode 2 + -H 1: the timed-kernel output is validated against expected — 0 wrong elements across all sizes; datacheck clean.

Made with Cursor

dlamd1dai · 2026-08-12T19:52:55Z

Stacked on ROCm#9927 (base branch users/dondai/gin-sdma-a2a-meta-repro is the head of that PR). This PR contains only the single in-device timing commit on top of ROCm#9927; that PR should merge first. Once ROCm#9927 lands in users/abouteiller/rccl-gin/installsh, this PR's base will be retargeted accordingly.

Fixes from PR #1 review of the in-device (wall_clock64) A2A timing path: - gin_sdma_devtime.h: split measure() into launch-all then sync-all passes. Synchronizing GPU i before launching GPU i+1 deadlocks under -g N>1 because the collective bodies do an all-ranks device barrier and each local GPU is its own rank (GPU 0 waited on GPUs never launched). Mirrors the startColl()/completeColl() launch-then-complete split. - gin_sdma_devtime.h: reduce via the harness Allreduce() helper instead of a direct MPI_Allreduce on MPI_COMM_WORLD. MPI is MPI_THREAD_SINGLE, so under -t N>1 all threads would issue concurrent collectives (can hang); Allreduce() combines across threads, serializes MPI to the last thread, and broadcasts back -- also fixing localMaxUs not being combined across threads. - common.h/common.cu: declare Allreduce<T> and add double/long long explicit instantiations so gin_sdma_devtime.h (via alltoall.cu) can call it. - alltoall.cu: gate AlltoAllDeviceTime to deviceImpl 3/4 (GIN/Hybrid) only; -D1/-D2 (Nvl LSA) lack GIN provisioning and would fault/hang. - alltoall.cu: clamp skip to >= 0 so the start stamp is always written (a negative skip left start_time[] uninitialized, skewing the mode-2 metric). - common.cu: document that mode-2 datacheck validates the production path; the timed kernels reuse the same __device__ collective functions, so only the loop/skip/stamp wrapper is not exercised by datacheck. Co-authored-by: Cursor <cursoragent@cursor.com>

Addresses PR #1 review (#5): in device-time-only mode the datacheck loop re-inits buffers and validates the PRODUCTION kernel, not the timed kernel. Add NCCL_GIN_ANVIL_DEVTIME_CHECK=1 (default off): after the deviceTime hook runs skip+loop back-to-back collectives into recvbuff, validate that output against `expected` before the datacheck loop overwrites it. Each timed iteration is a full deterministic collective over the unchanged sendbuff, so the final recvbuff equals `expected`. Wrong-element counts are summed across threads/ranks; any mismatch fails the run. Requires datacheck on (populates `expected`). Default behavior is unchanged. Co-authored-by: Cursor <cursoragent@cursor.com>

dlamd1dai · 2026-08-13T08:14:35Z

@pvallem — thanks for the thorough review. All six threads are addressed (and resolved above); summary:

Deadlock (measure()): split into launch-all then sync-all passes (mirrors startColl/completeColl).
MPI on MPI_COMM_WORLD: routed through the harness Allreduce() helper (max, last-thread MPI, cross-thread combine); added Allreduce<T> decl + explicit instantiations.
deviceImpl gating: hook now runs only for 3 (GIN) / 4 (Hybrid).
skip clamp: clamped to >= 0.
Timed-path datacheck: documented the default rationale and added opt-in NCCL_GIN_ANVIL_DEVTIME_CHECK=1 that validates the timed kernel's output against expected.
Build gate: called out ENABLE_DEVICE_API + NCCL >= 2.28.7 in the description; compile-verified with -DENABLE_DEVICE_API=ON.

Commits: b9eb26f (fixes) and 9cba85a (opt-in check).

I also added a Runtime validation (MI300X / gfx942, GIN Anvil-SDMA) section to the PR description. Ran alltoall_perf -D3 (NCCL_GIN_TYPE=5, -n 8 -g 1, 128 B–1 MiB):

Mode 1 (=1): #[a2a-devtime] line per size; datacheck Out of bounds values : 0 OK.
Mode 2 (=2): reported metric = in-kernel device latency, tracks the mode-1 value (e.g. 1 MiB: ~40.9 µs device vs ~28.7 µs host-graph); datacheck clean.
Mode 2 + NCCL_GIN_ANVIL_DEVTIME_CHECK=1: timed-kernel output validated — 0 wrong elements across all sizes.

PTAL when you get a chance.

pvallem · 2026-08-13T16:38:25Z

+  // signals/barriers the timed bodies rely on. The Nvl LSA impls (-D1/-D2) only
+  // provision the LSA barrier, so running the GIN timed kernel there would fault
+  // on an out-of-range signal index or wait forever -- skip them.
+  if (deviceImpl != 3 && deviceImpl != 4) return testSuccess;


This early-return correctly stops -D1/-D2 from running the GIN body on a comm with no GIN provisioning. But in mode 2 (NCCL_GIN_ANVIL_DEVICE_TIMING=2) deviceOnly is already true for these impls, so the host timed loop is skipped and this return leaves deviceDeltaSec at its initial 0.0. That flows into getBw as the divisor, so the row reports 0.00 us and inf GB/s with no warning — the same silent zero also happens for op != ncclSum, bf16/fp8, count == 0, or LOOP < 1. Could we emit a WARN and skip the row (or fail) when the hook can't produce a real measurement, rather than reporting a 0.0/inf metric?

Fixed in 8147eac. AlltoAllDeviceTime now returns a negative sentinel (-1) whenever the timed grid produced no positive busy window, and the mode-2 path in BenchTime initializes deviceDeltaSec = -1.0 and checks it before use:

double deviceDeltaSec = -1.0; TESTCHECK(args->collTest->deviceTime(args, type, op, root, in_place, &deviceDeltaSec)); if (deviceDeltaSec <= 0.0) { if (args->proc == 0 && args->thread == 0) { fprintf(stderr, "# WARN NCCL_GIN_ANVIL_DEVICE_TIMING=2: no in-kernel device-time " "measurement for this size (timed kernel did not run: non-GIN impl, " "unsupported op/type, count==0, or loop<1); skipping row " "(size %ld bytes)\n", args->expectedBytes); } return testSuccess; }

So instead of feeding 0.0 into getBw, we emit a one-line WARN and skip the row. This covers every "no measurement" case you listed (non-GIN -D1/-D2, op != ncclSum, bf16/fp8, count == 0, loop < 1) because they all early-return before writing *outDeltaSec, and the mode-2 branch itself now emits -1 when devUs <= 0.

Verified on MI355X (gfx950), 8-GPU GIN-SDMA A2A: -D1 in mode 2 now prints the WARN+skip for all 14 sizes with zero inf and zero 0.00 us rows; the in-place pass (normal timing) still prints as expected.

pvallem · 2026-08-13T16:56:02Z

+      const char* e = getenv("NCCL_GIN_ANVIL_DEVTIME_CHECK");
+      return (e && *e) ? atoi(e) : 0;
+    }();
+    if (devTimeCheck && datacheck) {


This validates the timed kernel's output only by comparing recvbuff to expected, without checking that the timed kernel actually ran. When the hook early-returns (deviceImpl 1/2, op != ncclSum, bf16/fp8, count == 0, LOOP < 1) the kernel never executes, but recvbuff still holds the warmup production result, which equals expected — so the check passes green and validates the production path, not the timed one it's meant to cover. The inverse also bites: with -w 0 and an early-returning hook, there's no warmup output left, so recvbuff is the raw initData scratch and this reports "timed-kernel output datacheck failed" for a kernel that never ran. Could we gate the check on the timed kernel having actually run — e.g. loop > 0 && kernel != nullptr && (deviceImpl == 3 || deviceImpl == 4) — before comparing?

Fixed in 8147eac. The NCCL_GIN_ANVIL_DEVTIME_CHECK block now sits after the WARN+skip guard added for the 0.0/inf issue, so it is reached only when the timed kernel actually ran. Concretely, deviceDeltaSec > 0 is now a precondition — which is exactly equivalent to your suggested loop > 0 && kernel != nullptr && (deviceImpl == 3 || deviceImpl == 4), since those are the only conditions under which the hook launches the kernel and returns a positive measurement. When the hook early-returns (deviceImpl 1/2, op != ncclSum, bf16/fp8, count == 0, loop < 1), deviceDeltaSec <= 0 triggers WARN+skip and the check never runs — eliminating both the false-green (stale warmup production result) and the false-red (-w 0 raw initData scratch) cases.

Verified: mode 2 + NCCL_GIN_ANVIL_DEVTIME_CHECK=1 on gfx950 passes with 0 wrong elements across 128 B–1 MiB; the -D1 case is now preempted by WARN+skip rather than reporting a spurious verdict.

pvallem · 2026-08-13T17:27:45Z

+  if (khz == 0) {
+    int dev = 0;
+    if (cudaGetDevice(&dev) != cudaSuccess) return 1;
+    if (hipDeviceGetAttribute(&khz, hipDeviceAttributeWallClockRate, dev) != hipSuccess || khz <= 0)


Nit: this rate is sampled once from whatever device is current and then used to convert every GPU's cycle delta, so it's slightly off for the non-reference GPUs on a mixed-clock node — could we sample it per-GPU in the Pass-2 loop?

Fixed in 8147eac. wallClockRateKhz() now takes a device ordinal and is sampled per-GPU inside the Pass-2 stamp-reduction loop of measure():

for (int i = 0; i < args->nGpus; i++) { CUDACHECK(cudaSetDevice(args->gpus[i])); CUDACHECK(cudaStreamSynchronize(args->streams[i])); const int rate = wallClockRateKhz(args->gpus[i]); // this GPU's own rate ... double totalUs = (double)(mx - mn) / (double)rate * 1.0e3; }

So each GPU's cycle delta is converted with its own fixed-frequency rate; no more single-reference-rate skew on mixed-clock nodes. On the (homogeneous) MI355X node the mode 1 and mode 2 device times stay consistent (~40.9 us @ 1 MiB).

@1

Addresses three reviewer (pvallem) findings on the in-device timing path: 1. Silent 0.00 us / inf GB/s in device-time-only mode (mode 2). When the timing hook produced no measurement -- the timed kernel never ran (non-GIN impl -D1/-D2, unsupported op/type, count==0, or loop<1) -- deviceDeltaSec stayed 0 and was reported as THE metric, so getBw divided by zero and printed inf GB/s with no warning. AlltoAllDeviceTime now returns a negative sentinel (-1) when the grid produced no positive busy window, and BenchTime emits a one-line WARN and skips the row instead of a bogus metric. 2. NCCL_GIN_ANVIL_DEVTIME_CHECK validated the wrong buffer when the kernel did not run: on early-return the stale warmup/production result passed (false green), and under -w 0 the raw init scratch failed (false red). The check is now reached only after the WARN+skip guard above, i.e. only when the timed kernel actually ran and recvbuff holds its output. 3. Single wall-clock rate applied to all GPUs. wallClockRateKhz() sampled the rate once from the current device; on a mixed-clock node that skews the non-reference GPUs' cycle->us conversion. It now takes a device ordinal and is sampled per-GPU inside the Pass-2 stamp-reduction loop of measure(). Verified on smci355 (MI355X/gfx950), -DENABLE_DEVICE_API=ON -DENABLE_ROCSHMEM=ON, 8-GPU GIN-SDMA AllToAll (128 B..1 MiB): mode 1 augment lines and mode 2 device metric agree (~40.9 us @1 MiB); mode 2 + DEVTIME_CHECK passes with 0 wrong elements; mode 2 with -D1 now emits the WARN+skip for every size with zero inf and zero bogus 0.00 us rows. Co-authored-by: Cursor <cursoragent@cursor.com>

abouteiller · 2026-08-19T00:44:51Z

+//
+// skip warmup iterations run first (steady-state caches/queues, discarded).
+// At i == skip every CTA records its start stamp; after the final iteration
+// every CTA records its end stamp. The host reduces min(start)/max(end) across


Yes — the timed kernels use per-CTA start/end wall-clock stamps (rocSHMEM-style): CTA c writes start_time[c] at iteration skip and end_time[c] after the final skip+loop iteration, then gin_devtime::measure() reduces min(start)..max(end) over the grid and divides by loop for per-iteration device latency. The stamp buffers are sized to gridCtas (one slot per CTA). See GinAlltoAllTimedKernel / HybridAlltoAllTimedKernel in alltoall.cu and the Pass-2 reduction in gin_sdma_devtime.h.

abouteiller · 2026-08-19T00:46:40Z

 }

+// Device-side (in-kernel wall_clock64) timing for AllToAll (GIN and Hybrid
+// tiers). Opt-in via NCCL_GIN_ANVIL_DEVICE_TIMING (legacy


Not specific to ANVIL, variable should be generic, or better, an command line parameter rather than an envvar.

abouteiller · 2026-08-19T00:49:16Z

+testResult_t AlltoAllDeviceTime(struct threadArgs* args, ncclDataType_t type, ncclRedOp_t op, int root, int in_place, double* outDeltaSec) {
+  static const int devTiming = []() {
+    const char* e = getenv("NCCL_GIN_ANVIL_DEVICE_TIMING");          // generic (all collectives)
+    if (!(e && *e)) e = getenv("NCCL_GIN_ANVIL_A2A_DEVICE_TIMING");  // legacy A2A-specific name


I would remove it altogether this has never been part of a release safe to remove without first deprecating.

abouteiller · 2026-08-19T14:40:59Z

+  // window matches the host tests' -w/-n. Opt in to auto-reducing iteration
+  // counts for very large chunks (bounded runtime; per-call copy dominates) via
+  // NCCL_GIN_ANVIL_A2A_DEVTIME_AUTOREDUCE=1.
+  static const int autoReduce = gin_devtime::envInt("NCCL_GIN_ANVIL_A2A_DEVTIME_AUTOREDUCE", 0);


Prefer a command line arg again, autoreduce is a confusing name, maybe loop-large, loop-mid?

abouteiller · 2026-08-19T14:47:42Z

+  int devTimeMode = 0;
+  if (deviceImpl != 0 && args->collTest->deviceTime != nullptr) {
+    const char* e = getenv("NCCL_GIN_ANVIL_DEVICE_TIMING");
+    if (!(e && *e)) e = getenv("NCCL_GIN_ANVIL_A2A_DEVICE_TIMING");


remove duplicate/legacy

Removed in 1df658c: the duplicate/legacy devtime env plumbing in common.cu / alltoall.cu (including NCCL_GIN_ANVIL_A2A_DEVICE_TIMING and the other ANVIL-specific devtime env vars). Everything now goes through the shared rccl-tests CLI flags (-B/-L/-P/-W/-Q/-j/-k/-H) declared once in common.cu / common.h.

abouteiller · 2026-08-19T14:48:50Z

-    if (agg_iters>1) NCCLCHECK(ncclGroupStart());
-    for (int aiter = 0; aiter < agg_iters; aiter++) {
-      TESTCHECK(startColl(args, type, op, root, in_place, iter*agg_iters+aiter));
+  if (!deviceOnly) {


More a question than a request for change: can we have both modes active (i.e. launch iter external kernels and have each kernel perform loop in-device timing loops)?

Not simultaneously in one metric today — by design the modes are sequential:

Mode 0: off.

Mode 1 (-B 1): the normal host/graph timed loop runs first (iters × production kernel via startColl/completeColl), then the collective's deviceTime hook launches one persistent timed kernel (skip+loop in-device iterations) and prints an extra #[*-devtime] line. You get both numbers, but they are separate measurements.

Mode 2 (-B 2, out-of-place only): the host/graph loop is skipped (deviceOnly); the hook's in-kernel measurement is the reported metric.

Nesting host iters with an in-kernel loop inside each host iteration would mix launch/capture overhead with device time and complicate reduction (which iteration is the metric?). Happy to explore that if you want it for a specific use case, but the current split keeps host bandwidth numbers and pure device-function latency cleanly separated.

abouteiller · 2026-08-19T14:56:28Z

+// fails, which yields a harmless raw-cycle magnitude instead of a divide-by-zero.
+static inline int wallClockRateKhz(int dev) {
+  int khz = 0;
+  if (hipDeviceGetAttribute(&khz, hipDeviceAttributeWallClockRate, dev) != hipSuccess || khz <= 0)


should fail rather than report nonsense?

abouteiller · 2026-08-19T14:58:37Z

+    double perIterUs = (loop > 0) ? totalUs / (double)loop : totalUs;
+    if (perIterUs > localMaxUs) localMaxUs = perIterUs;
+  }
+  Barrier(args);


Is this barrier needed?

Yes — kept intentionally. After Pass 2 finishes copying/reducing stamps on every local GPU, Barrier(args) ensures all ranks (and threads) have completed the timing run before Allreduce() publishes the result and before BenchTime moves on to datacheck / the next size. Without it, a fast rank could re-init buffers or start the production datacheck kernel while a slow rank is still draining stamps, racing the timed kernel's outputs. The pre-measurement Barrier() at the top of measure() serves the same isolation role on entry.

Fixes from PR #1 review of the in-device (wall_clock64) A2A timing path: - gin_sdma_devtime.h: split measure() into launch-all then sync-all passes. Synchronizing GPU i before launching GPU i+1 deadlocks under -g N>1 because the collective bodies do an all-ranks device barrier and each local GPU is its own rank (GPU 0 waited on GPUs never launched). Mirrors the startColl()/completeColl() launch-then-complete split. - gin_sdma_devtime.h: reduce via the harness Allreduce() helper instead of a direct MPI_Allreduce on MPI_COMM_WORLD. MPI is MPI_THREAD_SINGLE, so under -t N>1 all threads would issue concurrent collectives (can hang); Allreduce() combines across threads, serializes MPI to the last thread, and broadcasts back -- also fixing localMaxUs not being combined across threads. - common.h/common.cu: declare Allreduce<T> and add double/long long explicit instantiations so gin_sdma_devtime.h (via alltoall.cu) can call it. - alltoall.cu: gate AlltoAllDeviceTime to deviceImpl 3/4 (GIN/Hybrid) only; -D1/-D2 (Nvl LSA) lack GIN provisioning and would fault/hang. - alltoall.cu: clamp skip to >= 0 so the start stamp is always written (a negative skip left start_time[] uninitialized, skewing the mode-2 metric). - common.cu: document that mode-2 datacheck validates the production path; the timed kernels reuse the same __device__ collective functions, so only the loop/skip/stamp wrapper is not exercised by datacheck. Co-authored-by: Cursor <cursoragent@cursor.com> modified: projects/rccl-tests/src/common.cu

Addresses PR #1 review (#5): in device-time-only mode the datacheck loop re-inits buffers and validates the PRODUCTION kernel, not the timed kernel. Add NCCL_GIN_ANVIL_DEVTIME_CHECK=1 (default off): after the deviceTime hook runs skip+loop back-to-back collectives into recvbuff, validate that output against `expected` before the datacheck loop overwrites it. Each timed iteration is a full deterministic collective over the unchanged sendbuff, so the final recvbuff equals `expected`. Wrong-element counts are summed across threads/ranks; any mismatch fails the run. Requires datacheck on (populates `expected`). Default behavior is unchanged. Co-authored-by: Cursor <cursoragent@cursor.com>

@1

Addresses three reviewer (pvallem) findings on the in-device timing path: 1. Silent 0.00 us / inf GB/s in device-time-only mode (mode 2). When the timing hook produced no measurement -- the timed kernel never ran (non-GIN impl -D1/-D2, unsupported op/type, count==0, or loop<1) -- deviceDeltaSec stayed 0 and was reported as THE metric, so getBw divided by zero and printed inf GB/s with no warning. AlltoAllDeviceTime now returns a negative sentinel (-1) when the grid produced no positive busy window, and BenchTime emits a one-line WARN and skips the row instead of a bogus metric. 2. NCCL_GIN_ANVIL_DEVTIME_CHECK validated the wrong buffer when the kernel did not run: on early-return the stale warmup/production result passed (false green), and under -w 0 the raw init scratch failed (false red). The check is now reached only after the WARN+skip guard above, i.e. only when the timed kernel actually ran and recvbuff holds its output. 3. Single wall-clock rate applied to all GPUs. wallClockRateKhz() sampled the rate once from the current device; on a mixed-clock node that skews the non-reference GPUs' cycle->us conversion. It now takes a device ordinal and is sampled per-GPU inside the Pass-2 stamp-reduction loop of measure(). Verified on smci355 (MI355X/gfx950), -DENABLE_DEVICE_API=ON -DENABLE_ROCSHMEM=ON, 8-GPU GIN-SDMA AllToAll (128 B..1 MiB): mode 1 augment lines and mode 2 device metric agree (~40.9 us @1 MiB); mode 2 + DEVTIME_CHECK passes with 0 wrong elements; mode 2 with -D1 now emits the WARN+skip for every size with zero inf and zero bogus 0.00 us rows. Co-authored-by: Cursor <cursoragent@cursor.com>

dlamd1dai · 2026-08-24T19:08:31Z

@abouteiller — replied inline on the four remaining threads:

start-end (alltoall.cu): per-CTA start/end stamps + min(start)..max(end) reduction in measure().
remove duplicate/legacy (common.cu): legacy ANVIL devtime env vars removed in 1df658c; shared CLI flags only.
dual modes (common.cu): modes are sequential by design (mode 1 = host loop then timed kernel augment; mode 2 = device-only metric).
barrier (gin_sdma_devtime.h): post-measurement Barrier() isolates the timing run before datacheck / next size.

Re-triggering RCCL cluster CI (device-API + GIN) now — prior run timed out/cancelled at 24h.

PTAL when you have a moment; happy to adjust if the dual-mode sequencing should change.

Fixes from PR #1 review of the in-device (wall_clock64) A2A timing path: - gin_sdma_devtime.h: split measure() into launch-all then sync-all passes. Synchronizing GPU i before launching GPU i+1 deadlocks under -g N>1 because the collective bodies do an all-ranks device barrier and each local GPU is its own rank (GPU 0 waited on GPUs never launched). Mirrors the startColl()/completeColl() launch-then-complete split. - gin_sdma_devtime.h: reduce via the harness Allreduce() helper instead of a direct MPI_Allreduce on MPI_COMM_WORLD. MPI is MPI_THREAD_SINGLE, so under -t N>1 all threads would issue concurrent collectives (can hang); Allreduce() combines across threads, serializes MPI to the last thread, and broadcasts back -- also fixing localMaxUs not being combined across threads. - common.h/common.cu: declare Allreduce<T> and add double/long long explicit instantiations so gin_sdma_devtime.h (via alltoall.cu) can call it. - alltoall.cu: gate AlltoAllDeviceTime to deviceImpl 3/4 (GIN/Hybrid) only; -D1/-D2 (Nvl LSA) lack GIN provisioning and would fault/hang. - alltoall.cu: clamp skip to >= 0 so the start stamp is always written (a negative skip left start_time[] uninitialized, skewing the mode-2 metric). - common.cu: document that mode-2 datacheck validates the production path; the timed kernels reuse the same __device__ collective functions, so only the loop/skip/stamp wrapper is not exercised by datacheck. Co-authored-by: Cursor <cursoragent@cursor.com> modified: projects/rccl-tests/src/common.cu

Addresses PR #1 review (#5): in device-time-only mode the datacheck loop re-inits buffers and validates the PRODUCTION kernel, not the timed kernel. Add NCCL_GIN_ANVIL_DEVTIME_CHECK=1 (default off): after the deviceTime hook runs skip+loop back-to-back collectives into recvbuff, validate that output against `expected` before the datacheck loop overwrites it. Each timed iteration is a full deterministic collective over the unchanged sendbuff, so the final recvbuff equals `expected`. Wrong-element counts are summed across threads/ranks; any mismatch fails the run. Requires datacheck on (populates `expected`). Default behavior is unchanged. Co-authored-by: Cursor <cursoragent@cursor.com>

@1

Addresses three reviewer (pvallem) findings on the in-device timing path: 1. Silent 0.00 us / inf GB/s in device-time-only mode (mode 2). When the timing hook produced no measurement -- the timed kernel never ran (non-GIN impl -D1/-D2, unsupported op/type, count==0, or loop<1) -- deviceDeltaSec stayed 0 and was reported as THE metric, so getBw divided by zero and printed inf GB/s with no warning. AlltoAllDeviceTime now returns a negative sentinel (-1) when the grid produced no positive busy window, and BenchTime emits a one-line WARN and skips the row instead of a bogus metric. 2. NCCL_GIN_ANVIL_DEVTIME_CHECK validated the wrong buffer when the kernel did not run: on early-return the stale warmup/production result passed (false green), and under -w 0 the raw init scratch failed (false red). The check is now reached only after the WARN+skip guard above, i.e. only when the timed kernel actually ran and recvbuff holds its output. 3. Single wall-clock rate applied to all GPUs. wallClockRateKhz() sampled the rate once from the current device; on a mixed-clock node that skews the non-reference GPUs' cycle->us conversion. It now takes a device ordinal and is sampled per-GPU inside the Pass-2 stamp-reduction loop of measure(). Verified on smci355 (MI355X/gfx950), -DENABLE_DEVICE_API=ON -DENABLE_ROCSHMEM=ON, 8-GPU GIN-SDMA AllToAll (128 B..1 MiB): mode 1 augment lines and mode 2 device metric agree (~40.9 us @1 MiB); mode 2 + DEVTIME_CHECK passes with 0 wrong elements; mode 2 with -D1 now emits the WARN+skip for every size with zero inf and zero bogus 0.00 us rows. Co-authored-by: Cursor <cursoragent@cursor.com>

github-actions · 2026-08-24T19:47:25Z

rocprofiler-systems CI Summary

Test Results (for commit `96f6057`)

Workflow jobs covered: 36

Workflow	Runs	Avg Test Time	Tests	Skipped
Ubuntu 22.04	5	27m 18s	2,590	83
Ubuntu 24.04	5	32m 01s	2,610	83
Debian	4	34m 10s	2,184	57
RHEL	12	32m 17s	6,216	230
System deps	10	27m 54s	5,404	168
Total	36	30m 33s	19,004	621

Failed Tests

No failed tests reported.

Sanitizers

Thread	Undefined	Address
Passed	Passed	Passed

## Motivation JIRA ID -AIRUNTIME-2710, to remove the HRR support in the changelog for ROCm 10.0 ## Technical Details Updated the changelog in HIP 10.0.0 section, remove the HRR feature support. ## Issue Tracking AIRUNTIME-2710 ## Test Plan N/A ## Test Result N/A ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.

## Summary JIRA ID: AIESW-38547 Approval via - DEGFX11-14512 DEGFX11-14513 DEGFX11-14514 Add approved gfx11.5 (RDNA 3.5) performance counters in `config.yaml` for gfx1150-1153: - 10 SQ (VALU/VMEM stall, wait, bus) + 4 GL1C (read-size, arb level) - 1 GRBM (TCP_BUSY) + 11 TA (addresser/aligner/write-path/flat/store/vmem) - TCP_TA_REQ (new) + gfx11.5 definition added to TCP_PENDING_STALL_CYCLES ## Validation (gfx1151) - All event IDs cross-checked against the gfx11.5 RDL (23/23 exact where the RDL is explicit; GL1C/GRBM corroborated via soc21 + hardware). - Collected via rocprofv3 on synthetic HIP kernels and a real vLLM (Qwen2.5-0.5B AWQ) inference run. Quantitative check: GL1C_REQ_READ_128B measured 4,000,016 vs 4,000,000 expected. - TCP events pinned empirically (TA_REQ=2, GL1_PENDING_STALL=54); soc21_enum.h is off-by-one from the RDL/HW in the TCP high region. ## Notes - gfx1151 limitation: only one TCP-block counter can be collected per pass. --------- Signed-off-by: Vinod Pangul <146476973+vipangul@users.noreply.github.com> Co-authored-by: Cursor <cursoragent@cursor.com>

Adapt the in-kernel device-side timing instrumentation from gin-stage2j-sdma-A1457-a2av-wip to this branch's simpler two-kernel (GIN impl 3 / Hybrid impl 4) A2A architecture. Measures pure GPU device-function latency, excluding host launch/graph/teardown overhead. - gin_sdma_devtime.h: portable gin_devtime::measure() scaffold that queries the GPU wall-clock rate, allocates per-CTA start/end stamps, launches a persistent skip+loop timed kernel, reduces min(start)..max(end) over CTAs, takes the slowest local GPU, then MPI_MAX across ranks. - alltoall.cu: factor GIN/Hybrid collective bodies into reusable __device__ functions, add persistent GinAlltoAllTimedKernel/HybridAlltoAllTimedKernel loop variants, and implement AlltoAllDeviceTime hook (env-controlled loop/skip/autoreduce; selects timed kernel via deviceImpl). - common.h/common.cu: add optional testColl::deviceTime hook and expose deviceImpl (non-static) so per-collective hooks can pick the matching timed kernel. Integrate device-time-only mode into BenchTime: NCCL_GIN_ANVIL_DEVICE_TIMING=1 prints an extra #[a2a-devtime] line; =2 reports in-kernel device latency as the metric (skips host graph loop). - CMakeLists.txt: hipify/copy gin_sdma_devtime.h with the collective tests. Compile-verified on MI355X/gfx950: alltoall_perf builds and links (timed kernels + hook present), BUILD SANITY OK, GIN Anvil-SDMA smoke OK. Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit 540956b) modified: projects/rccl-tests/src/CMakeLists.txt modified: projects/rccl-tests/src/alltoall.cu modified: projects/rccl-tests/src/common.cu modified: projects/rccl-tests/src/common.h new file: projects/rccl-tests/src/gin_sdma_devtime.h

Fixes from PR #1 review of the in-device (wall_clock64) A2A timing path: - gin_sdma_devtime.h: split measure() into launch-all then sync-all passes. Synchronizing GPU i before launching GPU i+1 deadlocks under -g N>1 because the collective bodies do an all-ranks device barrier and each local GPU is its own rank (GPU 0 waited on GPUs never launched). Mirrors the startColl()/completeColl() launch-then-complete split. - gin_sdma_devtime.h: reduce via the harness Allreduce() helper instead of a direct MPI_Allreduce on MPI_COMM_WORLD. MPI is MPI_THREAD_SINGLE, so under -t N>1 all threads would issue concurrent collectives (can hang); Allreduce() combines across threads, serializes MPI to the last thread, and broadcasts back -- also fixing localMaxUs not being combined across threads. - common.h/common.cu: declare Allreduce<T> and add double/long long explicit instantiations so gin_sdma_devtime.h (via alltoall.cu) can call it. - alltoall.cu: gate AlltoAllDeviceTime to deviceImpl 3/4 (GIN/Hybrid) only; -D1/-D2 (Nvl LSA) lack GIN provisioning and would fault/hang. - alltoall.cu: clamp skip to >= 0 so the start stamp is always written (a negative skip left start_time[] uninitialized, skewing the mode-2 metric). - common.cu: document that mode-2 datacheck validates the production path; the timed kernels reuse the same __device__ collective functions, so only the loop/skip/stamp wrapper is not exercised by datacheck. Co-authored-by: Cursor <cursoragent@cursor.com> modified: projects/rccl-tests/src/common.cu

Addresses PR #1 review (#5): in device-time-only mode the datacheck loop re-inits buffers and validates the PRODUCTION kernel, not the timed kernel. Add NCCL_GIN_ANVIL_DEVTIME_CHECK=1 (default off): after the deviceTime hook runs skip+loop back-to-back collectives into recvbuff, validate that output against `expected` before the datacheck loop overwrites it. Each timed iteration is a full deterministic collective over the unchanged sendbuff, so the final recvbuff equals `expected`. Wrong-element counts are summed across threads/ranks; any mismatch fails the run. Requires datacheck on (populates `expected`). Default behavior is unchanged. Co-authored-by: Cursor <cursoragent@cursor.com>

@1

Addresses three reviewer (pvallem) findings on the in-device timing path: 1. Silent 0.00 us / inf GB/s in device-time-only mode (mode 2). When the timing hook produced no measurement -- the timed kernel never ran (non-GIN impl -D1/-D2, unsupported op/type, count==0, or loop<1) -- deviceDeltaSec stayed 0 and was reported as THE metric, so getBw divided by zero and printed inf GB/s with no warning. AlltoAllDeviceTime now returns a negative sentinel (-1) when the grid produced no positive busy window, and BenchTime emits a one-line WARN and skips the row instead of a bogus metric. 2. NCCL_GIN_ANVIL_DEVTIME_CHECK validated the wrong buffer when the kernel did not run: on early-return the stale warmup/production result passed (false green), and under -w 0 the raw init scratch failed (false red). The check is now reached only after the WARN+skip guard above, i.e. only when the timed kernel actually ran and recvbuff holds its output. 3. Single wall-clock rate applied to all GPUs. wallClockRateKhz() sampled the rate once from the current device; on a mixed-clock node that skews the non-reference GPUs' cycle->us conversion. It now takes a device ordinal and is sampled per-GPU inside the Pass-2 stamp-reduction loop of measure(). Verified on smci355 (MI355X/gfx950), -DENABLE_DEVICE_API=ON -DENABLE_ROCSHMEM=ON, 8-GPU GIN-SDMA AllToAll (128 B..1 MiB): mode 1 augment lines and mode 2 device metric agree (~40.9 us @1 MiB); mode 2 + DEVTIME_CHECK passes with 0 wrong elements; mode 2 with -D1 now emits the WARN+skip for every size with zero inf and zero bogus 0.00 us rows. Co-authored-by: Cursor <cursoragent@cursor.com>

Replace ANVIL-specific devtime env vars with rccl-tests CLI flags (-B/-L/-P tier overrides and -H datacheck) per review feedback, remove the legacy A2A env alias, and fail when wall-clock rate query fails. Co-authored-by: Cursor <cursoragent@cursor.com>

dlamd1dai · 2026-08-24T20:23:37Z

Superseded by upstream PR ROCm/rocm-systems#10658. Branch users/dondai/gin-sdma-a2a-devtime pushed to ROCm/rocm-systems (5 commits, rebased onto upstream develop @ 60ca803). Closing this fork PR.

github-actions · 2026-08-24T21:52:37Z

Code Coverage Report

Unit Tests

Code Coverage: unit-tests

Lines: 🔴 19.5% — 6,951/35,582 across 320 files
Functions: 🔴 19.6% — 717/3,654

🔴 0-20% (201 files)

Lines	Functions	File
0.0% 0/10	0/2	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/byte_container_with_subtype.hpp`
0.0% 0/1	0/1	`build/coverage/_deps/spdlog-src/include/spdlog/details/log_msg_buffer.h`
0.0% 0/1	0/1	`build/coverage/_deps/spdlog-src/include/spdlog/sinks/ansicolor_sink.h`
0.0% 0/2	0/1	`build/coverage/include/profiler-hub/shared_types.hpp`
0.0% 0/105	0/9	`source/bin/common/argument_registration.hpp`
0.0% 0/254	0/28	`source/bin/common/tool_runner.cpp`
0.0% 0/941	0/48	`source/bin/rocprof-sys-avail/avail.cpp`
0.0% 0/61	0/6	`source/bin/rocprof-sys-avail/avail.hpp`
0.0% 0/227	0/25	`source/bin/rocprof-sys-avail/common.cpp`
0.0% 0/17	0/3	`source/bin/rocprof-sys-avail/component_categories.hpp`
0.0% 0/258	0/12	`source/bin/rocprof-sys-avail/generate_config.cpp`
0.0% 0/39	0/2	`source/bin/rocprof-sys-avail/get_availability.hpp`
0.0% 0/10	0/1	`source/bin/rocprof-sys-avail/get_categories.hpp`
0.0% 0/102	0/6	`source/bin/rocprof-sys-avail/hw_counter_query.cpp`
0.0% 0/4	0/1	`source/bin/rocprof-sys-avail/info_type.cpp`
0.0% 0/9	0/1	`source/bin/rocprof-sys-avail/info_type.hpp`
0.0% 0/446	0/47	`source/bin/rocprof-sys-causal/impl.cpp`
0.0% 0/80	0/2	`source/bin/rocprof-sys-causal/rocprof-sys-causal.cpp`
0.0% 0/673	0/46	`source/bin/rocprof-sys-instrument/details.cpp`
0.0% 0/91	0/7	`source/bin/rocprof-sys-instrument/function_signature.cpp`
0.0% 0/20	0/5	`source/bin/rocprof-sys-instrument/function_signature.hpp`
0.0% 0/158	0/7	`source/bin/rocprof-sys-instrument/info.hpp`
0.0% 0/264	0/22	`source/bin/rocprof-sys-instrument/internal_libs.cpp`
0.0% 0/64	0/6	`source/bin/rocprof-sys-instrument/log.cpp`
0.0% 0/5	0/1	`source/bin/rocprof-sys-instrument/log.hpp`
0.0% 0/533	0/41	`source/bin/rocprof-sys-instrument/module_function.cpp`
0.0% 0/93	0/7	`source/bin/rocprof-sys-instrument/module_function.hpp`
0.0% 0/1541	0/87	`source/bin/rocprof-sys-instrument/rocprof-sys-instrument.cpp`
0.0% 0/156	0/12	`source/bin/rocprof-sys-instrument/rocprof-sys-instrument.hpp`
0.0% 0/2	0/1	`source/bin/rocprof-sys-run/rocprof-sys-run.cpp`
0.0% 0/2	0/1	`source/bin/rocprof-sys-sample/rocprof-sys-sample.cpp`
0.0% 0/24	0/1	`source/lib/backends/amd_smi/wrapper.hpp`
0.0% 0/21	0/3	`source/lib/backends/rocprofiler_sdk/wrapper.hpp`
0.0% 0/16	0/4	`source/lib/binary/address_multirange.cpp`
0.0% 0/9	0/1	`source/lib/binary/address_multirange.hpp`
0.0% 0/117	0/5	`source/lib/binary/analysis.cpp`
0.0% 0/7	0/1	`source/lib/binary/binary_info.hpp`
0.0% 0/102	0/9	`source/lib/binary/dwarf_entry.cpp`
0.0% 0/1	0/1	`source/lib/binary/dwarf_entry.hpp`
0.0% 0/69	0/6	`source/lib/binary/link_map.cpp`
0.0% 0/3	0/1	`source/lib/binary/link_map.hpp`
0.0% 0/10	0/1	`source/lib/binary/scope_filter.cpp`
0.0% 0/5	0/2	`source/lib/binary/scope_filter.hpp`
0.0% 0/158	0/18	`source/lib/binary/symbol.cpp`
0.0% 0/8	0/8	`source/lib/binary/symbol.hpp`
0.0% 0/9	0/2	`source/lib/common/env_vars.hpp`
0.0% 0/40	0/3	`source/lib/common/invoke.hpp`
0.0% 0/150	0/10	`source/lib/common/md5sum.hpp`
0.0% 0/22	0/2	`source/lib/common/setup.hpp`
0.0% 0/3	0/1	`source/lib/common/units.hpp`
0.0% 0/166	0/2	`source/lib/core/agent_info.hpp`
0.0% 0/33	0/3	`source/lib/core/amd_smi.cpp`
0.0% 0/80	0/18	`source/lib/core/binary/address_range.cpp`
0.0% 0/9	0/3	`source/lib/core/binary/address_range.hpp`
0.0% 0/49	0/10	`source/lib/core/categories.cpp`
0.0% 0/11	0/11	`source/lib/core/components/fwd.hpp`
0.0% 0/26	0/3	`source/lib/core/config.hpp`
0.0% 0/14	0/2	`source/lib/core/containers/aligned_static_vector.hpp`
0.0% 0/11	0/3	`source/lib/core/containers/c_array.hpp`
0.0% 0/1	0/1	`source/lib/core/containers/operators.hpp`
0.0% 0/46	0/12	`source/lib/core/containers/stable_vector.hpp`
0.0% 0/23	0/5	`source/lib/core/containers/static_vector.hpp`
0.0% 0/89	0/6	`source/lib/core/cpu.cpp`
0.0% 0/1	0/1	`source/lib/core/cpu.hpp`
0.0% 0/53	0/6	`source/lib/core/dynamic_library.cpp`
0.0% 0/28	0/7	`source/lib/core/exception.cpp`
0.0% 0/178	0/23	`source/lib/core/gpu.cpp`
0.0% 0/160	0/12	`source/lib/core/gpu_metrics.cpp`
0.0% 0/33	0/6	`source/lib/core/mpi.hpp`
0.0% 0/56	0/4	`source/lib/core/mproc.cpp`
0.0% 0/24	0/2	`source/lib/core/node_info.cpp`
0.0% 0/2	0/2	`source/lib/core/node_info.hpp`
0.0% 0/79	0/5	`source/lib/core/perf.cpp`
0.0% 0/51	0/5	`source/lib/core/perfetto/cached_perfetto_session.cpp`
0.0% 0/33	0/4	`source/lib/core/perfetto/counter_track.hpp`
0.0% 0/29	0/7	`source/lib/core/perfetto/emitter.cpp`
0.0% 0/48	0/9	`source/lib/core/perfetto/emitter.hpp`
0.0% 0/7	0/1	`source/lib/core/perfetto/sinks/append_mode.hpp`
0.0% 0/13	0/3	`source/lib/core/perfetto/sinks/tee_sink.cpp`
0.0% 0/1	0/1	`source/lib/core/perfetto/sinks/tee_sink.hpp`
0.0% 0/1	0/1	`source/lib/core/progress/bar.hpp`
0.0% 0/15	0/4	`source/lib/core/sdk-tracing-config-deps.hpp`
0.0% 0/54	0/3	`source/lib/core/trace_cache/cache_manager.cpp`
0.0% 0/7	0/3	`source/lib/core/trace_cache/cache_manager.hpp`
0.0% 0/516	0/71	`source/lib/core/trace_cache/metadata_registry.cpp`
0.0% 0/764	0/103	`source/lib/core/trace_cache/perfetto_processor.cpp`
0.0% 0/119	0/12	`source/lib/core/trace_cache/post_processor.cpp`
0.0% 0/749	0/31	`source/lib/core/trace_cache/rocpd_processor.cpp`
0.0% 0/77	0/1	`source/lib/core/trace_cache/sample_processor.hpp`
0.0% 0/547	0/90	`source/lib/rocprof-sys-dl/dl.cpp`
0.0% 0/16	0/1	`source/lib/rocprof-sys-dl/main.c`
0.0% 0/280	0/31	`source/lib/rocprof-sys-rt/src/RTcommon.c`
0.0% 0/29	0/4	`source/lib/rocprof-sys-rt/src/RTheap-linux.c`
0.0% 0/65	0/4	`source/lib/rocprof-sys-rt/src/RTheap.c`
0.0% 0/212	0/22	`source/lib/rocprof-sys-rt/src/RTlinux.c`
0.0% 0/97	0/8	`source/lib/rocprof-sys-rt/src/RTposix.c`
0.0% 0/6	0/2	`source/lib/rocprof-sys-rt/src/RTsignal.c`
0.0% 0/11	0/2	`source/lib/rocprof-sys-rt/src/RTthread-x86-64.c`
0.0% 0/3	0/1	`source/lib/rocprof-sys-rt/src/RTthread.c`
0.0% 0/79	0/13	`source/lib/rocprof-sys-user/user.cpp`
0.0% 0/70	0/18	`source/lib/rocprof-sys/api.cpp`
0.0% 0/620	0/36	`source/lib/rocprof-sys/library.cpp`
0.0% 0/77	0/6	`source/lib/rocprof-sys/library/causal/components/backtrace.cpp`
0.0% 0/7	0/2	`source/lib/rocprof-sys/library/causal/components/backtrace.hpp`
0.0% 0/93	0/11	`source/lib/rocprof-sys/library/causal/components/blocking_gotcha.cpp`
0.0% 0/33	0/7	`source/lib/rocprof-sys/library/causal/components/causal_gotcha.cpp`
0.0% 0/85	0/22	`source/lib/rocprof-sys/library/causal/components/progress_point.cpp`
0.0% 0/23	0/3	`source/lib/rocprof-sys/library/causal/components/progress_point.hpp`
0.0% 0/44	0/8	`source/lib/rocprof-sys/library/causal/components/unblocking_gotcha.cpp`
0.0% 0/553	0/33	`source/lib/rocprof-sys/library/causal/data.cpp`
0.0% 0/87	0/16	`source/lib/rocprof-sys/library/causal/delay.cpp`
0.0% 0/373	0/29	`source/lib/rocprof-sys/library/causal/experiment.cpp`
0.0% 0/3	0/2	`source/lib/rocprof-sys/library/causal/experiment.hpp`
0.0% 0/24	0/5	`source/lib/rocprof-sys/library/causal/sample_data.cpp`
0.0% 0/268	0/25	`source/lib/rocprof-sys/library/causal/sampling.cpp`
0.0% 0/4	0/1	`source/lib/rocprof-sys/library/causal/selected_entry.cpp`
0.0% 0/5	0/3	`source/lib/rocprof-sys/library/causal/selected_entry.hpp`
0.0% 0/71	0/11	`source/lib/rocprof-sys/library/components/backtrace.cpp`
0.0% 0/1	0/1	`source/lib/rocprof-sys/library/components/backtrace.hpp`
0.0% 0/215	0/30	`source/lib/rocprof-sys/library/components/backtrace_metrics.cpp`
0.0% 0/11	0/2	`source/lib/rocprof-sys/library/components/backtrace_metrics.hpp`
0.0% 0/9	0/3	`source/lib/rocprof-sys/library/components/backtrace_timestamp.cpp`
0.0% 0/4	0/2	`source/lib/rocprof-sys/library/components/backtrace_timestamp.hpp`
0.0% 0/71	0/10	`source/lib/rocprof-sys/library/components/callchain.cpp`
0.0% 0/1	0/1	`source/lib/rocprof-sys/library/components/callchain.hpp`
0.0% 0/367	0/30	`source/lib/rocprof-sys/library/components/comm_data.cpp`
0.0% 0/19	0/5	`source/lib/rocprof-sys/library/components/comm_data.hpp`
0.0% 0/8	0/2	`source/lib/rocprof-sys/library/components/ensure_storage.hpp`
0.0% 0/28	0/6	`source/lib/rocprof-sys/library/components/exit_gotcha.cpp`
0.0% 0/1	0/1	`source/lib/rocprof-sys/library/components/exit_gotcha.hpp`
0.0% 0/68	0/8	`source/lib/rocprof-sys/library/components/fork_gotcha.cpp`
0.0% 0/1	0/1	`source/lib/rocprof-sys/library/components/fork_gotcha.hpp`
0.0% 0/13	0/3	`source/lib/rocprof-sys/library/components/kill_gotcha.cpp`
0.0% 0/1	0/1	`source/lib/rocprof-sys/library/components/kill_gotcha.hpp`
0.0% 0/158	0/22	`source/lib/rocprof-sys/library/components/mpi_gotcha.cpp`
0.0% 0/2	0/1	`source/lib/rocprof-sys/library/components/mpi_gotcha.hpp`
0.0% 0/533	0/12	`source/lib/rocprof-sys/library/components/mpip.hpp`
0.0% 0/78	0/19	`source/lib/rocprof-sys/library/components/numa_gotcha.cpp`
0.0% 0/2	0/1	`source/lib/rocprof-sys/library/components/numa_gotcha.hpp`
0.0% 0/336	0/20	`source/lib/rocprof-sys/library/components/pthread_create_gotcha.cpp`
0.0% 0/3	0/1	`source/lib/rocprof-sys/library/components/pthread_create_gotcha.hpp`
0.0% 0/34	0/9	`source/lib/rocprof-sys/library/components/pthread_gotcha.cpp`
0.0% 0/24	0/2	`source/lib/rocprof-sys/library/components/pthread_mutex_gotcha_policy.hpp`
0.0% 0/164	0/31	`source/lib/rocprof-sys/library/components/vaapi_gotcha.cpp`
0.0% 0/2	0/1	`source/lib/rocprof-sys/library/components/vaapi_gotcha.hpp`
0.0% 0/145	0/10	`source/lib/rocprof-sys/library/coverage.cpp`
0.0% 0/29	0/5	`source/lib/rocprof-sys/library/coverage.hpp`
0.0% 0/60	0/16	`source/lib/rocprof-sys/library/coverage/impl.hpp`
0.0% 0/252	0/37	`source/lib/rocprof-sys/library/kokkosp.cpp`
0.0% 0/245	0/31	`source/lib/rocprof-sys/library/perf.cpp`
0.0% 0/11	0/3	`source/lib/rocprof-sys/library/perf.hpp`
0.0% 0/186	0/12	`source/lib/rocprof-sys/library/pmc/collectors/common/settings.hpp`
0.0% 0/66	0/5	`source/lib/rocprof-sys/library/pmc/collectors/cpu/cache_policy.hpp`
0.0% 0/47	0/6	`source/lib/rocprof-sys/library/pmc/collectors/cpu/cpu_traits.hpp`
0.0% 0/52	0/9	`source/lib/rocprof-sys/library/pmc/collectors/cpu/perfetto_policy.hpp`
0.0% 0/100	0/6	`source/lib/rocprof-sys/library/pmc/collectors/gpu/cache_policy.hpp`
0.0% 0/222	0/36	`source/lib/rocprof-sys/library/pmc/collectors/gpu/perfetto_policy.hpp`
0.0% 0/25	0/3	`source/lib/rocprof-sys/library/pmc/collectors/gpu/sample.hpp`
0.0% 0/32	0/3	`source/lib/rocprof-sys/library/pmc/collectors/gpu_perf_counter/cache_policy.hpp`
0.0% 0/39	0/7	`source/lib/rocprof-sys/library/pmc/collectors/gpu_perf_counter/perfetto_policy.hpp`
0.0% 0/16	0/2	`source/lib/rocprof-sys/library/pmc/collectors/nic/sample.hpp`
0.0% 0/14	0/2	`source/lib/rocprof-sys/library/pmc/device_providers/procfs/provider.hpp`
0.0% 0/148	0/18	`source/lib/rocprof-sys/library/pmc/sampler.cpp`
0.0% 0/113	0/11	`source/lib/rocprof-sys/library/process_sampler.cpp`
0.0% 0/13	0/1	`source/lib/rocprof-sys/library/process_sampler.hpp`
0.0% 0/2	0/1	`source/lib/rocprof-sys/library/rocm.cpp`
0.0% 0/1316	0/88	`source/lib/rocprof-sys/library/rocprofiler-sdk.cpp`
0.0% 0/111	0/13	`source/lib/rocprof-sys/library/rocprofiler-sdk/counters.cpp`
0.0% 0/17	0/6	`source/lib/rocprof-sys/library/rocprofiler-sdk/counters.hpp`
0.0% 0/129	0/9	`source/lib/rocprof-sys/library/rocprofiler-sdk/fwd.cpp`
0.0% 0/34	0/7	`source/lib/rocprof-sys/library/rocprofiler-sdk/fwd.hpp`
0.0% 0/265	0/25	`source/lib/rocprof-sys/library/rocprofiler-sdk/kfd_events.cpp`
0.0% 0/36	0/9	`source/lib/rocprof-sys/library/rocprofiler-sdk/marker_writer.cpp`
0.0% 0/136	0/13	`source/lib/rocprof-sys/library/rocprofiler-sdk/rccl.cpp`
0.0% 0/112	0/10	`source/lib/rocprof-sys/library/rocprofiler-sdk/roctx_client.cpp`
0.0% 0/118	0/20	`source/lib/rocprof-sys/library/runtime.cpp`
0.0% 0/2	0/1	`source/lib/rocprof-sys/library/runtime.hpp`
0.0% 0/994	0/65	`source/lib/rocprof-sys/library/sampling.cpp`
0.0% 0/145	0/31	`source/lib/rocprof-sys/library/thread_data.hpp`
0.0% 0/10	0/3	`source/lib/rocprof-sys/library/thread_data_growth.hpp`
0.0% 0/10	0/1	`source/lib/rocprof-sys/library/thread_deleter.cpp`
0.0% 0/5	0/1	`source/lib/rocprof-sys/library/thread_deleter.hpp`
0.0% 0/161	0/27	`source/lib/rocprof-sys/library/thread_info.cpp`
0.0% 0/2	0/1	`source/lib/rocprof-sys/library/thread_info.hpp`
0.0% 0/113	0/11	`source/lib/rocprof-sys/library/tracing.cpp`
0.0% 0/109	0/30	`source/lib/rocprof-sys/library/tracing.hpp`
0.0% 0/3	0/1	`source/lib/rocprof-sys/library/tracing/annotation.cpp`
0.0% 0/45	0/4	`source/lib/rocprof-sys/library/tracing/annotation.hpp`
0.0% 0/8	0/2	`source/lib/rocprof-sys/progress.cpp`
0.0% 0/68	0/14	`source/lib/rocprof-sys/regions.cpp`
0.0% 0/82	0/4	`source/lib/rocprof-sys/timeout.cpp`
0.0% 0/454	0/46	`source/python/libpyrocprofsys.cpp`
0.6% 1/166	1/13	`source/lib/core/perfetto.cpp`
0.8% 7/859	1/74	`source/lib/core/argparse.cpp`
2.5% 43/1694	7/289	`source/lib/core/config.cpp`
7.1% 1/14	0/1	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/iterators/primitive_iterator.hpp`
8.4% 15/178	3/25	`source/lib/core/constraint.cpp`
15.4% 4/26	0/6	`source/lib/core/trace_cache/metadata_registry.hpp`
18.2% 4/22	1/3	`source/lib/common/static_object.hpp`
18.8% 9/48	2/5	`source/lib/core/output_file_registry.cpp`
19.4% 6/31	2/5	`source/bin/common/common_utils.hpp`

🟠 20-50% (11 files)

Lines	Functions	File
20.0% 1/5	0/2	`source/lib/core/argparse.hpp`
21.3% 13/61	3/9	`source/lib/core/utility.hpp`
22.2% 2/9	0/6	`source/lib/core/constraint.hpp`
23.8% 15/63	3/12	`source/lib/core/agent_manager.cpp`
30.0% 127/423	12/46	`build/coverage/_deps/fmt-src/include/fmt/format.h`
32.1% 52/162	7/17	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/input/json_sax.hpp`
38.2% 13/34	1/4	`source/lib/core/perfetto/sinks/io_helpers.hpp`
40.3% 195/484	14/17	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/input/lexer.hpp`
40.5% 17/42	5/12	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/exceptions.hpp`
47.2% 221/468	17/38	`build/coverage/_deps/fmt-src/include/fmt/base.h`
49.6% 127/256	14/32	`source/lib/rocprof-sys/library/components/category_region.hpp`

🟡 50-80% (33 files)

Lines	Functions	File
50.0% 153/306	9/11	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/output/serializer.hpp`
50.0% 1/2	0/1	`source/lib/rocprof-sys/library/pmc/collectors/nic/types.hpp`
52.0% 13/25	2/3	`source/lib/rocprof-sys/library/pmc/collectors/gpu_perf_counter/gpu_perf_counter_traits.hpp`
53.0% 220/415	25/37	`source/bin/common/common_utils.cpp`
53.3% 8/15	3/3	`source/lib/core/perfetto/sinks/per_pid_file_sink.cpp`
55.0% 22/40	4/6	`build/coverage/_deps/fmt-src/include/fmt/ranges.h`
55.6% 5/9	1/3	`build/coverage/_deps/fmt-src/include/fmt/chrono.h`
55.7% 34/61	5/8	`source/lib/rocprof-sys/library/pmc/collectors/cpu/sample.hpp`
58.3% 14/24	3/8	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/output/output_adapters.hpp`
58.3% 7/12	0/4	`build/coverage/_deps/spdlog-src/include/spdlog/common.h`
60.0% 3/5	1/1	`source/lib/core/agent.hpp`
62.1% 18/29	1/1	`source/lib/core/perfetto/locked_file_append.cpp`
62.1% 41/66	7/8	`source/lib/rocprof-sys/library/pmc/collectors/base/collector.hpp`
62.8% 49/78	4/5	`source/lib/core/trace_cache/discovery.cpp`
63.1% 70/111	7/10	`source/lib/core/progress/bar.cpp`
63.4% 272/429	33/40	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/json.hpp`
64.1% 336/524	26/40	`source/bin/common/json_config.cpp`
65.4% 34/52	4/4	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/conversions/from_json.hpp`
66.2% 47/71	5/5	`source/lib/rocprof-sys/library/pmc/collectors/gpu/gpu_traits.hpp`
66.4% 87/131	4/5	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/input/parser.hpp`
66.7% 2/3	0/2	`source/lib/rocprof-sys/library/pmc/collectors/cpu/types.hpp`
67.2% 43/64	11/14	`source/lib/core/trace_cache/data_types.cpp`
68.4% 39/57	4/6	`source/lib/core/perfetto.hpp`
68.8% 22/32	3/8	`source/lib/rocprof-sys/library/pmc/collectors/gpu_perf_counter/types.hpp`
71.1% 69/97	8/8	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/iterators/iter_impl.hpp`
71.6% 212/296	14/16	`source/bin/common/preset_registry.cpp`
72.3% 86/119	14/17	`source/lib/core/perfetto/engine.cpp`
73.7% 28/38	6/7	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/iterators/iteration_proxy.hpp`
75.0% 3/4	0/1	`source/lib/backends/amd_smi/nic_types.hpp`
75.0% 3/4	1/2	`source/lib/core/perfetto/sinks/single_file_sink.hpp`
76.5% 78/102	4/4	`source/lib/core/perfetto/sinks/single_file_sink.cpp`
77.8% 77/99	10/11	`source/lib/core/perfetto/engine_impl.hpp`
79.3% 88/111	13/15	`source/lib/common/path.hpp`

🟢 80-100% (75 files)

Lines	Functions	File
80.0% 8/10	1/2	`build/coverage/include/profiler-hub/writer_types.hpp`
80.5% 62/77	9/10	`source/lib/rocprof-sys/library/pmc/device_providers/rocprofiler_sdk/provider.hpp`
81.2% 39/48	3/4	`source/lib/core/utility.cpp`
82.8% 48/58	6/7	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/conversions/to_json.hpp`
83.3% 279/335	17/23	`source/lib/core/sdk-tracing-config.hpp`
83.9% 135/161	12/16	`source/lib/backends/procfs/backend.hpp`
84.6% 11/13	2/4	`source/lib/rocprof-sys/library/rocprofiler-sdk/roctx_client.hpp`
85.0% 85/100	11/15	`source/lib/logger/logger.hpp`
85.1% 40/47	2/2	`source/lib/core/trace_cache/storage_parser.hpp`
86.9% 213/245	15/16	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/conversions/to_chars.hpp`
88.1% 273/310	29/35	`source/lib/common/environment.hpp`
89.1% 49/55	4/4	`source/lib/rocprof-sys/library/pmc/collectors/gpu_perf_counter/device.hpp`
89.5% 34/38	5/8	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/input/input_adapters.hpp`
90.0% 54/60	4/4	`source/lib/core/trace_cache/buffer_storage.cpp`
90.1% 100/111	8/8	`source/lib/rocprof-sys/library/pmc/collectors/cpu/device.hpp`
90.7% 303/334	28/31	`source/lib/core/trace_cache/sample_type.hpp`
90.9% 10/11	1/2	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/json_ref.hpp`
91.0% 243/267	19/20	`source/lib/core/trace_cache/unified_memory_processor.cpp`
91.3% 178/195	14/18	`source/lib/rocprof-sys/library/components/shmem_gotcha.hpp`
91.7% 11/12	2/2	`source/lib/core/perfetto/session_backend.hpp`
92.6% 163/176	9/9	`source/lib/rocprof-sys/library/pmc/collectors/gpu/device.hpp`
92.9% 13/14	0/1	`source/lib/rocprof-sys/library/pmc/collectors/common/collector_slice.hpp`
93.0% 40/43	6/11	`source/lib/core/demangler.hpp`
93.2% 68/73	7/8	`source/lib/core/trace_cache/buffer_storage.hpp`
94.7% 36/38	7/7	`build/coverage/_deps/spdlog-src/include/spdlog/logger.h`
94.7% 54/57	6/8	`source/lib/rocprof-sys/library/components/pthread_mutex_gotcha.hpp`
94.8% 55/58	1/5	`source/lib/core/trace_cache/cacheable.hpp`
94.8% 92/97	9/10	`source/lib/rocprof-sys/library/rocprofiler-sdk/trace_control.cpp`
95.0% 19/20	4/6	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/string_concat.hpp`
95.7% 22/23	4/6	`source/lib/common/synchronized.hpp`
95.8% 23/24	2/4	`source/lib/common/pci_bdf.hpp`
95.9% 162/169	19/22	`source/lib/rocprof-sys/library/components/ucx_gotcha.hpp`
96.2% 25/26	1/2	`source/lib/common/delimit.hpp`
96.5% 82/85	11/12	`source/lib/core/state.hpp`
96.8% 30/31	3/4	`source/lib/rocprof-sys/library/pmc/collectors/gpu_perf_counter/sample.hpp`
97.3% 73/75	3/3	`source/lib/core/perfetto/packet_framing.hpp`
97.4% 74/76	6/6	`source/lib/core/common_types.hpp`
97.5% 39/40	6/7	`source/lib/core/trace_cache/rocpd_helpers.hpp`
97.8% 45/46	7/7	`source/lib/rocprof-sys/library/rocprofiler-sdk/rccl_internal.hpp`
97.9% 47/48	3/3	`source/lib/rocprof-sys/library/pmc/collectors/nic/device.hpp`
98.1% 51/52	7/7	`source/lib/rocprof-sys/library/pmc/device_providers/amd_smi/provider.hpp`
98.7% 78/79	15/16	`source/lib/backends/amd_smi/backend.hpp`
99.1% 110/111	25/26	`source/lib/backends/rocprofiler_sdk/backend.hpp`
100.0% 14/14	1/1	`build/coverage/_deps/fmt-src/include/fmt/ostream.h`
100.0% 6/6	0/3	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/adl_serializer.hpp`
100.0% 2/2	0/1	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/input/position_t.hpp`
100.0% 2/2	0/1	`build/coverage/_deps/spdlog-src/include/spdlog/details/backtracer.h`
100.0% 2/2	0/1	`build/coverage/_deps/spdlog-src/include/spdlog/details/circular_q.h`
100.0% 1/1	0/1	`build/coverage/include/profiler-hub/reader_types.hpp`
100.0% 1/1	0/1	`source/bin/common/json_config.hpp`
100.0% 107/107	13/14	`source/lib/backends/amd_smi/device.hpp`
100.0% 7/7	0/1	`source/lib/backends/amd_smi/gpu_types.hpp`
100.0% 1/1	0/1	`source/lib/backends/rocprofiler_sdk/types.hpp`
100.0% 12/12	1/1	`source/lib/common/join.hpp`
100.0% 2/2	0/2	`source/lib/core/agent_manager.hpp`
100.0% 1/1	0/1	`source/lib/core/output_file_registry.hpp`
100.0% 1/1	0/1	`source/lib/core/perfetto/category_registry.hpp`
100.0% 1/1	0/1	`source/lib/core/perfetto/engine.hpp`
100.0% 7/7	0/1	`source/lib/core/perfetto/locked_file_append.hpp`
100.0% 9/9	2/3	`source/lib/core/perfetto/sinks/polymorphic_sink_view.hpp`
100.0% 6/6	2/2	`source/lib/core/perfetto/sinks/recording_sink.cpp`
100.0% 5/5	0/2	`source/lib/core/perfetto/sinks/recording_sink.hpp`
100.0% 6/6	2/2	`source/lib/core/progress/tracker.cpp`
100.0% 1/1	0/1	`source/lib/core/progress/tracker.hpp`
100.0% 4/4	0/3	`source/lib/core/trace_cache/data_types.hpp`
100.0% 9/9	3/3	`source/lib/core/trace_cache/type_registry.hpp`
100.0% 24/24	4/5	`source/lib/core/trace_cache/unified_memory_processor.hpp`
100.0% 5/5	2/2	`source/lib/core/track_registry.cpp`
100.0% 4/4	0/2	`source/lib/core/track_registry.hpp`
100.0% 10/10	1/1	`source/lib/logger/logger.cpp`
100.0% 9/9	0/2	`source/lib/rocprof-sys/library/pmc/collectors/gpu/types.hpp`
100.0% 20/20	4/4	`source/lib/rocprof-sys/library/pmc/collectors/nic/perfetto_policy.hpp`
100.0% 2/2	0/2	`source/lib/rocprof-sys/library/pmc/common/types.hpp`
100.0% 34/34	3/4	`source/lib/rocprof-sys/library/rocprofiler-sdk/marker_writer.hpp`
100.0% 3/3	1/2	`source/lib/rocprof-sys/library/rocprofiler-sdk/trace_control.hpp`

Distribution

Range	Files
0-20%	201
20-50%	11
50-80%	33
80-100%	75

E2E Tests

Code Coverage: e2e-tests

Lines: 🟡 62.4% — 22,219/35,582 across 320 files
Functions: 🟡 62.8% — 2,296/3,654

🔴 0-20% (54 files)

Lines	Functions	File
0.0% 0/10	0/2	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/byte_container_with_subtype.hpp`
0.0% 0/1	0/1	`build/coverage/_deps/spdlog-src/include/spdlog/details/log_msg_buffer.h`
0.0% 0/1	0/1	`build/coverage/_deps/spdlog-src/include/spdlog/sinks/ansicolor_sink.h`
0.0% 0/2	0/1	`build/coverage/include/profiler-hub/shared_types.hpp`
0.0% 0/102	0/6	`source/bin/rocprof-sys-avail/hw_counter_query.cpp`
0.0% 0/24	0/1	`source/lib/backends/amd_smi/wrapper.hpp`
0.0% 0/150	0/10	`source/lib/common/md5sum.hpp`
0.0% 0/166	0/2	`source/lib/core/agent_info.hpp`
0.0% 0/89	0/6	`source/lib/core/cpu.cpp`
0.0% 0/1	0/1	`source/lib/core/cpu.hpp`
0.0% 0/53	0/6	`source/lib/core/dynamic_library.cpp`
0.0% 0/28	0/7	`source/lib/core/exception.cpp`
0.0% 0/160	0/12	`source/lib/core/gpu_metrics.cpp`
0.0% 0/79	0/5	`source/lib/core/perf.cpp`
0.0% 0/13	0/3	`source/lib/core/perfetto/sinks/tee_sink.cpp`
0.0% 0/1	0/1	`source/lib/core/perfetto/sinks/tee_sink.hpp`
0.0% 0/15	0/4	`source/lib/core/sdk-tracing-config-deps.hpp`
0.0% 0/749	0/31	`source/lib/core/trace_cache/rocpd_processor.cpp`
0.0% 0/6	0/2	`source/lib/rocprof-sys-rt/src/RTsignal.c`
0.0% 0/11	0/2	`source/lib/rocprof-sys-rt/src/RTthread-x86-64.c`
0.0% 0/3	0/1	`source/lib/rocprof-sys-rt/src/RTthread.c`
0.0% 0/1	0/1	`source/lib/rocprof-sys/library/components/exit_gotcha.hpp`
0.0% 0/1	0/1	`source/lib/rocprof-sys/library/components/kill_gotcha.hpp`
0.0% 0/164	0/31	`source/lib/rocprof-sys/library/components/vaapi_gotcha.cpp`
0.0% 0/52	0/9	`source/lib/rocprof-sys/library/pmc/collectors/cpu/perfetto_policy.hpp`
0.0% 0/100	0/6	`source/lib/rocprof-sys/library/pmc/collectors/gpu/cache_policy.hpp`
0.0% 0/222	0/36	`source/lib/rocprof-sys/library/pmc/collectors/gpu/perfetto_policy.hpp`
0.0% 0/25	0/3	`source/lib/rocprof-sys/library/pmc/collectors/gpu/sample.hpp`
0.0% 0/32	0/3	`source/lib/rocprof-sys/library/pmc/collectors/gpu_perf_counter/cache_policy.hpp`
0.0% 0/39	0/7	`source/lib/rocprof-sys/library/pmc/collectors/gpu_perf_counter/perfetto_policy.hpp`
0.0% 0/16	0/2	`source/lib/rocprof-sys/library/pmc/collectors/nic/sample.hpp`
0.0% 0/2	0/1	`source/lib/rocprof-sys/library/rocm.cpp`
0.0% 0/111	0/13	`source/lib/rocprof-sys/library/rocprofiler-sdk/counters.cpp`
0.0% 0/17	0/6	`source/lib/rocprof-sys/library/rocprofiler-sdk/counters.hpp`
0.0% 0/129	0/9	`source/lib/rocprof-sys/library/rocprofiler-sdk/fwd.cpp`
0.0% 0/34	0/7	`source/lib/rocprof-sys/library/rocprofiler-sdk/fwd.hpp`
0.0% 0/265	0/25	`source/lib/rocprof-sys/library/rocprofiler-sdk/kfd_events.cpp`
0.0% 0/136	0/13	`source/lib/rocprof-sys/library/rocprofiler-sdk/rccl.cpp`
0.0% 0/112	0/10	`source/lib/rocprof-sys/library/rocprofiler-sdk/roctx_client.cpp`
0.0% 0/3	0/1	`source/lib/rocprof-sys/library/tracing/annotation.cpp`
0.0% 0/45	0/4	`source/lib/rocprof-sys/library/tracing/annotation.hpp`
0.0% 0/454	0/46	`source/python/libpyrocprofsys.cpp`
0.6% 1/166	1/13	`source/lib/core/perfetto.cpp`
2.5% 33/1316	3/88	`source/lib/rocprof-sys/library/rocprofiler-sdk.cpp`
6.1% 2/33	1/3	`source/lib/core/amd_smi.cpp`
7.1% 1/14	0/1	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/iterators/primitive_iterator.hpp`
8.3% 3/36	1/9	`source/lib/rocprof-sys/library/rocprofiler-sdk/marker_writer.cpp`
11.3% 8/71	4/10	`source/lib/rocprof-sys/library/components/callchain.cpp`
13.6% 38/280	7/31	`source/lib/rocprof-sys-rt/src/RTcommon.c`
13.8% 4/29	1/4	`source/lib/rocprof-sys-rt/src/RTheap-linux.c`
14.2% 30/212	4/22	`source/lib/rocprof-sys-rt/src/RTlinux.c`
16.1% 83/516	13/71	`source/lib/core/trace_cache/metadata_registry.cpp`
18.0% 32/178	6/23	`source/lib/core/gpu.cpp`
19.0% 4/21	0/3	`source/lib/backends/rocprofiler_sdk/wrapper.hpp`

🟠 20-50% (22 files)

Lines	Functions	File
22.8% 49/215	5/30	`source/lib/rocprof-sys/library/components/backtrace_metrics.cpp`
27.8% 27/97	4/8	`source/lib/rocprof-sys-rt/src/RTposix.c`
29.3% 224/764	26/103	`source/lib/core/trace_cache/perfetto_processor.cpp`
31.7% 59/186	6/12	`source/lib/rocprof-sys/library/pmc/collectors/common/settings.hpp`
32.1% 52/162	7/17	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/input/json_sax.hpp`
35.7% 10/28	3/6	`source/lib/rocprof-sys/library/components/exit_gotcha.cpp`
38.5% 5/13	2/3	`source/lib/rocprof-sys/library/components/kill_gotcha.cpp`
38.8% 31/80	1/2	`source/bin/rocprof-sys-causal/rocprof-sys-causal.cpp`
39.8% 37/93	5/11	`source/lib/rocprof-sys/library/causal/components/blocking_gotcha.cpp`
40.3% 195/484	14/17	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/input/lexer.hpp`
40.5% 17/42	5/12	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/exceptions.hpp`
41.3% 26/63	7/12	`source/lib/core/agent_manager.cpp`
43.8% 21/48	6/9	`source/lib/core/perfetto/emitter.hpp`
43.9% 36/82	2/4	`source/lib/rocprof-sys/timeout.cpp`
44.4% 4/9	1/3	`source/lib/rocprof-sys/library/components/backtrace_timestamp.cpp`
45.0% 27/60	6/16	`source/lib/rocprof-sys/library/coverage/impl.hpp`
45.6% 72/158	3/7	`source/bin/rocprof-sys-instrument/info.hpp`
46.2% 36/78	6/19	`source/lib/rocprof-sys/library/components/numa_gotcha.cpp`
47.7% 31/65	3/4	`source/lib/rocprof-sys-rt/src/RTheap.c`
48.3% 14/29	5/7	`source/lib/core/perfetto/emitter.cpp`
48.5% 178/367	16/30	`source/lib/rocprof-sys/library/components/comm_data.cpp`
49.4% 39/79	6/13	`source/lib/rocprof-sys-user/user.cpp`

🟡 50-80% (70 files)

Lines	Functions	File
50.0% 8/16	2/4	`source/lib/binary/address_multirange.cpp`
50.0% 1/2	0/1	`source/lib/rocprof-sys/library/components/numa_gotcha.hpp`
50.0% 1/2	0/1	`source/lib/rocprof-sys/library/components/vaapi_gotcha.hpp`
50.0% 1/2	0/1	`source/lib/rocprof-sys/library/pmc/collectors/nic/types.hpp`
50.0% 1/2	0/1	`source/lib/rocprof-sys/library/thread_info.hpp`
50.0% 4/8	1/2	`source/lib/rocprof-sys/progress.cpp`
50.5% 55/109	12/30	`source/lib/rocprof-sys/library/tracing.hpp`
50.8% 128/252	25/37	`source/lib/rocprof-sys/library/kokkosp.cpp`
50.8% 215/423	24/46	`build/coverage/_deps/fmt-src/include/fmt/format.h`
52.0% 13/25	2/3	`source/lib/rocprof-sys/library/pmc/collectors/gpu_perf_counter/gpu_perf_counter_traits.hpp`
53.2% 41/77	0/1	`source/lib/core/trace_cache/sample_processor.hpp`
53.4% 79/148	12/18	`source/lib/rocprof-sys/library/pmc/sampler.cpp`
55.0% 22/40	4/6	`build/coverage/_deps/fmt-src/include/fmt/ranges.h`
55.6% 5/9	1/3	`build/coverage/_deps/fmt-src/include/fmt/chrono.h`
55.7% 39/70	13/18	`source/lib/rocprof-sys/api.cpp`
56.2% 172/306	9/11	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/output/serializer.hpp`
57.5% 23/40	1/3	`source/lib/common/invoke.hpp`
58.0% 142/245	20/31	`source/lib/rocprof-sys/library/perf.cpp`
58.3% 14/24	3/8	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/output/output_adapters.hpp`
58.3% 7/12	0/4	`build/coverage/_deps/spdlog-src/include/spdlog/common.h`
60.0% 63/105	5/9	`source/bin/common/argument_registration.hpp`
60.0% 3/5	1/1	`source/lib/core/agent.hpp`
61.5% 288/468	21/38	`build/coverage/_deps/fmt-src/include/fmt/base.h`
61.8% 42/68	5/8	`source/lib/rocprof-sys/library/components/fork_gotcha.cpp`
62.1% 18/29	1/1	`source/lib/core/perfetto/locked_file_append.cpp`
62.3% 619/994	51/65	`source/lib/rocprof-sys/library/sampling.cpp`
62.5% 50/80	13/18	`source/lib/core/binary/address_range.cpp`
62.8% 49/78	4/5	`source/lib/core/trace_cache/discovery.cpp`
63.1% 70/111	7/10	`source/lib/core/progress/bar.cpp`
64.7% 22/34	3/4	`source/lib/core/perfetto/sinks/io_helpers.hpp`
65.4% 34/52	4/4	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/conversions/from_json.hpp`
65.5% 281/429	33/40	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/json.hpp`
66.1% 37/56	3/4	`source/lib/core/mproc.cpp`
66.2% 47/71	5/5	`source/lib/rocprof-sys/library/pmc/collectors/gpu/gpu_traits.hpp`
66.4% 87/131	4/5	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/input/parser.hpp`
66.9% 79/118	15/20	`source/lib/rocprof-sys/library/runtime.cpp`
67.1% 106/158	14/22	`source/lib/rocprof-sys/library/components/mpi_gotcha.cpp`
68.8% 22/32	3/8	`source/lib/rocprof-sys/library/pmc/collectors/gpu_perf_counter/types.hpp`
69.2% 18/26	2/6	`source/lib/core/trace_cache/metadata_registry.hpp`
69.2% 9/13	0/1	`source/lib/rocprof-sys/library/process_sampler.hpp`
69.7% 46/66	7/8	`source/lib/rocprof-sys/library/pmc/collectors/base/collector.hpp`
70.2% 40/57	4/6	`source/lib/core/perfetto.hpp`
70.6% 24/34	5/9	`source/lib/rocprof-sys/library/components/pthread_gotcha.cpp`
71.1% 69/97	8/8	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/iterators/iter_impl.hpp`
71.4% 10/14	1/2	`source/lib/core/containers/aligned_static_vector.hpp`
71.8% 393/547	46/90	`source/lib/rocprof-sys-dl/dl.cpp`
72.3% 86/119	11/12	`source/lib/core/trace_cache/post_processor.cpp`
72.7% 24/33	5/7	`source/lib/rocprof-sys/library/causal/components/causal_gotcha.cpp`
72.7% 32/44	5/8	`source/lib/rocprof-sys/library/causal/components/unblocking_gotcha.cpp`
73.7% 28/38	6/7	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/iterators/iteration_proxy.hpp`
74.0% 409/553	30/33	`source/lib/rocprof-sys/library/causal/data.cpp`
74.0% 219/296	15/16	`source/bin/common/preset_registry.cpp`
74.5% 35/47	4/6	`source/lib/rocprof-sys/library/pmc/collectors/cpu/cpu_traits.hpp`
74.5% 120/161	23/27	`source/lib/rocprof-sys/library/thread_info.cpp`
75.0% 3/4	0/1	`source/lib/backends/amd_smi/nic_types.hpp`
75.0% 18/24	4/5	`source/lib/rocprof-sys/library/causal/sample_data.cpp`
75.2% 394/524	30/40	`source/bin/common/json_config.cpp`
75.9% 511/673	35/46	`source/bin/rocprof-sys-instrument/details.cpp`
76.1% 1172/1541	52/87	`source/bin/rocprof-sys-instrument/rocprof-sys-instrument.cpp`
76.6% 49/64	5/6	`source/bin/rocprof-sys-instrument/log.cpp`
76.8% 53/69	5/6	`source/lib/binary/link_map.cpp`
77.3% 17/22	3/3	`source/lib/common/static_object.hpp`
77.7% 482/620	27/36	`source/lib/rocprof-sys/library.cpp`
78.0% 209/268	23/25	`source/lib/rocprof-sys/library/causal/sampling.cpp`
78.1% 139/178	19/25	`source/lib/core/constraint.cpp`
78.6% 114/145	26/31	`source/lib/rocprof-sys/library/thread_data.hpp`
79.0% 328/415	35/37	`source/bin/common/common_utils.cpp`
79.1% 201/254	24/28	`source/bin/common/tool_runner.cpp`
79.2% 19/24	2/2	`source/lib/core/node_info.cpp`
79.9% 211/264	20/22	`source/bin/rocprof-sys-instrument/internal_libs.cpp`

🟢 80-100% (174 files)

Lines	Functions	File
80.0% 8/10	1/2	`build/coverage/include/profiler-hub/writer_types.hpp`
80.0% 4/5	0/1	`source/bin/rocprof-sys-instrument/log.hpp`
80.5% 1363/1694	262/289	`source/lib/core/config.cpp`
80.5% 62/77	9/10	`source/lib/rocprof-sys/library/pmc/device_providers/rocprofiler_sdk/provider.hpp`
80.9% 55/68	12/14	`source/lib/rocprof-sys/regions.cpp`
81.2% 69/85	17/22	`source/lib/rocprof-sys/library/causal/components/progress_point.cpp`
81.2% 39/48	4/5	`source/lib/core/output_file_registry.cpp`
81.2% 13/16	1/1	`source/lib/rocprof-sys-dl/main.c`
81.8% 81/99	11/11	`source/lib/core/perfetto/engine_impl.hpp`
82.2% 438/533	41/41	`source/bin/rocprof-sys-instrument/module_function.cpp`
82.4% 42/51	5/5	`source/lib/core/perfetto/cached_perfetto_session.cpp`
82.4% 187/227	22/25	`source/bin/rocprof-sys-avail/common.cpp`
83.3% 85/102	4/4	`source/lib/core/perfetto/sinks/single_file_sink.cpp`
84.2% 314/373	18/29	`source/lib/rocprof-sys/library/causal/experiment.cpp`
84.2% 16/19	3/5	`source/lib/rocprof-sys/library/components/comm_data.hpp`
84.3% 86/102	8/9	`source/lib/binary/dwarf_entry.cpp`
84.6% 11/13	2/4	`source/lib/rocprof-sys/library/rocprofiler-sdk/roctx_client.hpp`
84.8% 134/158	17/18	`source/lib/binary/symbol.cpp`
85.1% 74/87	14/16	`source/lib/rocprof-sys/library/causal/delay.cpp`
85.9% 134/156	12/12	`source/bin/rocprof-sys-instrument/rocprof-sys-instrument.hpp`
86.0% 86/100	11/15	`source/lib/logger/logger.hpp`
86.0% 80/93	5/7	`source/bin/rocprof-sys-instrument/module_function.hpp`
86.3% 290/336	16/20	`source/lib/rocprof-sys/library/components/pthread_create_gotcha.cpp`
86.7% 98/113	9/11	`source/lib/rocprof-sys/library/process_sampler.cpp`
87.0% 20/23	4/5	`source/lib/core/containers/static_vector.hpp`
87.0% 47/54	3/3	`source/lib/core/trace_cache/cache_manager.cpp`
87.4% 390/446	35/47	`source/bin/rocprof-sys-causal/impl.cpp`
87.6% 99/113	10/11	`source/lib/rocprof-sys/library/tracing.cpp`
88.9% 8/9	5/6	`source/lib/core/constraint.hpp`
89.0% 129/145	10/10	`source/lib/rocprof-sys/library/coverage.cpp`
89.0% 276/310	31/35	`source/lib/common/environment.hpp`
89.1% 49/55	4/4	`source/lib/rocprof-sys/library/pmc/collectors/gpu_perf_counter/device.hpp`
89.4% 768/859	59/74	`source/lib/core/argparse.cpp`
89.8% 220/245	15/16	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/conversions/to_chars.hpp`
90.0% 18/20	5/5	`source/bin/rocprof-sys-instrument/function_signature.hpp`
90.0% 9/10	1/1	`source/lib/binary/scope_filter.cpp`
90.0% 54/60	4/4	`source/lib/core/trace_cache/buffer_storage.cpp`
90.0% 9/10	1/1	`source/lib/rocprof-sys/library/thread_deleter.cpp`
90.1% 64/71	9/11	`source/lib/rocprof-sys/library/components/backtrace.cpp`
90.6% 106/117	5/5	`source/lib/binary/analysis.cpp`
90.7% 234/258	11/12	`source/bin/rocprof-sys-avail/generate_config.cpp`
90.7% 303/334	28/31	`source/lib/core/trace_cache/sample_type.hpp`
90.9% 10/11	1/2	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/json_ref.hpp`
90.9% 20/22	2/2	`source/lib/common/setup.hpp`
91.0% 243/267	19/20	`source/lib/core/trace_cache/unified_memory_processor.cpp`
91.2% 83/91	6/7	`source/bin/rocprof-sys-instrument/function_signature.cpp`
91.3% 178/195	14/18	`source/lib/rocprof-sys/library/components/shmem_gotcha.hpp`
91.7% 11/12	2/2	`source/lib/core/perfetto/session_backend.hpp`
91.9% 102/111	8/8	`source/lib/rocprof-sys/library/pmc/collectors/cpu/device.hpp`
92.6% 163/176	9/9	`source/lib/rocprof-sys/library/pmc/collectors/gpu/device.hpp`
92.8% 873/941	41/48	`source/bin/rocprof-sys-avail/avail.cpp`
92.9% 13/14	0/1	`source/lib/rocprof-sys/library/pmc/collectors/common/collector_slice.hpp`
92.9% 13/14	2/2	`source/lib/rocprof-sys/library/pmc/device_providers/procfs/provider.hpp`
93.1% 27/29	3/5	`source/lib/rocprof-sys/library/coverage.hpp`
93.2% 68/73	7/8	`source/lib/core/trace_cache/buffer_storage.hpp`
93.3% 111/119	16/17	`source/lib/core/perfetto/engine.cpp`
93.3% 14/15	3/3	`source/lib/core/perfetto/sinks/per_pid_file_sink.cpp`
93.7% 314/335	19/23	`source/lib/core/sdk-tracing-config.hpp`
93.8% 240/256	28/32	`source/lib/rocprof-sys/library/components/category_region.hpp`
93.9% 31/33	6/6	`source/lib/core/mpi.hpp`
94.1% 16/17	3/3	`source/bin/rocprof-sys-avail/component_categories.hpp`
94.4% 152/161	14/16	`source/lib/backends/procfs/backend.hpp`
94.6% 105/111	14/15	`source/lib/common/path.hpp`
94.7% 36/38	7/7	`build/coverage/_deps/spdlog-src/include/spdlog/logger.h`
94.8% 55/58	1/5	`source/lib/core/trace_cache/cacheable.hpp`
95.0% 19/20	4/6	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/string_concat.hpp`
95.1% 58/61	7/8	`source/lib/rocprof-sys/library/pmc/collectors/cpu/sample.hpp`
95.7% 44/46	10/12	`source/lib/core/containers/stable_vector.hpp`
95.7% 45/47	2/2	`source/lib/core/trace_cache/storage_parser.hpp`
95.8% 23/24	2/4	`source/lib/common/pci_bdf.hpp`
95.9% 162/169	19/22	`source/lib/rocprof-sys/library/components/ucx_gotcha.hpp`
96.1% 74/77	6/6	`source/lib/rocprof-sys/library/causal/components/backtrace.cpp`
96.2% 25/26	1/2	`source/lib/common/delimit.hpp`
96.5% 82/85	11/12	`source/lib/core/state.hpp`
96.6% 56/58	7/7	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/conversions/to_json.hpp`
96.7% 59/61	9/9	`source/lib/core/utility.hpp`
96.8% 30/31	3/4	`source/lib/rocprof-sys/library/pmc/collectors/gpu_perf_counter/sample.hpp`
96.9% 62/64	14/14	`source/lib/core/trace_cache/data_types.cpp`
96.9% 94/97	10/10	`source/lib/rocprof-sys/library/rocprofiler-sdk/trace_control.cpp`
97.0% 32/33	4/4	`source/lib/core/perfetto/counter_track.hpp`
97.3% 73/75	3/3	`source/lib/core/perfetto/packet_framing.hpp`
97.4% 74/76	6/6	`source/lib/core/common_types.hpp`
97.5% 39/40	6/7	`source/lib/core/trace_cache/rocpd_helpers.hpp`
97.8% 45/46	7/7	`source/lib/rocprof-sys/library/rocprofiler-sdk/rccl_internal.hpp`
97.9% 47/48	4/4	`source/lib/core/utility.cpp`
97.9% 47/48	3/3	`source/lib/rocprof-sys/library/pmc/collectors/nic/device.hpp`
98.1% 51/52	7/7	`source/lib/rocprof-sys/library/pmc/device_providers/amd_smi/provider.hpp`
98.2% 56/57	6/8	`source/lib/rocprof-sys/library/components/pthread_mutex_gotcha.hpp`
98.7% 78/79	15/16	`source/lib/backends/amd_smi/backend.hpp`
99.1% 110/111	25/26	`source/lib/backends/rocprofiler_sdk/backend.hpp`
100.0% 14/14	1/1	`build/coverage/_deps/fmt-src/include/fmt/ostream.h`
100.0% 6/6	0/3	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/adl_serializer.hpp`
100.0% 38/38	7/8	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/input/input_adapters.hpp`
100.0% 2/2	0/1	`build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/input/position_t.hpp`
100.0% 2/2	0/1	`build/coverage/_deps/spdlog-src/include/spdlog/details/backtracer.h`
100.0% 2/2	0/1	`build/coverage/_deps/spdlog-src/include/spdlog/details/circular_q.h`
100.0% 1/1	0/1	`build/coverage/include/profiler-hub/reader_types.hpp`
100.0% 31/31	4/5	`source/bin/common/common_utils.hpp`
100.0% 1/1	0/1	`source/bin/common/json_config.hpp`
100.0% 61/61	6/6	`source/bin/rocprof-sys-avail/avail.hpp`
100.0% 39/39	2/2	`source/bin/rocprof-sys-avail/get_availability.hpp`
100.0% 10/10	1/1	`source/bin/rocprof-sys-avail/get_categories.hpp`
100.0% 4/4	1/1	`source/bin/rocprof-sys-avail/info_type.cpp`
100.0% 9/9	0/1	`source/bin/rocprof-sys-avail/info_type.hpp`
100.0% 2/2	1/1	`source/bin/rocprof-sys-run/rocprof-sys-run.cpp`
100.0% 2/2	1/1	`source/bin/rocprof-sys-sample/rocprof-sys-sample.cpp`
100.0% 107/107	13/14	`source/lib/backends/amd_smi/device.hpp`
100.0% 7/7	0/1	`source/lib/backends/amd_smi/gpu_types.hpp`
100.0% 1/1	0/1	`source/lib/backends/rocprofiler_sdk/types.hpp`
100.0% 9/9	0/1	`source/lib/binary/address_multirange.hpp`
100.0% 7/7	1/1	`source/lib/binary/binary_info.hpp`
100.0% 1/1	0/1	`source/lib/binary/dwarf_entry.hpp`
100.0% 3/3	0/1	`source/lib/binary/link_map.hpp`
100.0% 5/5	1/2	`source/lib/binary/scope_filter.hpp`
100.0% 8/8	7/8	`source/lib/binary/symbol.hpp`
100.0% 9/9	2/2	`source/lib/common/env_vars.hpp`
100.0% 12/12	1/1	`source/lib/common/join.hpp`
100.0% 23/23	5/6	`source/lib/common/synchronized.hpp`
100.0% 3/3	1/1	`source/lib/common/units.hpp`
100.0% 2/2	1/2	`source/lib/core/agent_manager.hpp`
100.0% 5/5	0/2	`source/lib/core/argparse.hpp`
100.0% 9/9	1/3	`source/lib/core/binary/address_range.hpp`
100.0% 49/49	10/10	`source/lib/core/categories.cpp`
100.0% 11/11	11/11	`source/lib/core/components/fwd.hpp`
100.0% 26/26	3/3	`source/lib/core/config.hpp`
100.0% 11/11	1/3	`source/lib/core/containers/c_array.hpp`
100.0% 1/1	0/1	`source/lib/core/containers/operators.hpp`
100.0% 43/43	7/11	`source/lib/core/demangler.hpp`
100.0% 2/2	2/2	`source/lib/core/node_info.hpp`
100.0% 1/1	0/1	`source/lib/core/output_file_registry.hpp`
100.0% 1/1	0/1	`source/lib/core/perfetto/category_registry.hpp`
100.0% 1/1	0/1	`source/lib/core/perfetto/engine.hpp`
100.0% 7/7	0/1	`source/lib/core/perfetto/locked_file_append.hpp`
100.0% 7/7	1/1	`source/lib/core/perfetto/sinks/append_mode.hpp`
100.0% 9/9	2/3	`source/lib/core/perfetto/sinks/polymorphic_sink_view.hpp`
100.0% 6/6	2/2	`source/lib/core/perfetto/sinks/recording_sink.cpp`
100.0% 5/5	0/2	`source/lib/core/perfetto/sinks/recording_sink.hpp`
100.0% 4/4	2/2	`source/lib/core/perfetto/sinks/single_file_sink.hpp`
100.0% 1/1	0/1	`source/lib/core/progress/bar.hpp`
100.0% 6/6	2/2	`source/lib/core/progress/tracker.cpp`
100.0% 1/1	0/1	`source/lib/core/progress/tracker.hpp`
100.0% 7/7	1/3	`source/lib/core/trace_cache/cache_manager.hpp`
100.0% 4/4	0/3	`source/lib/core/trace_cache/data_types.hpp`
100.0% 9/9	3/3	`source/lib/core/trace_cache/type_registry.hpp`
100.0% 24/24	4/5	`source/lib/core/trace_cache/unified_memory_processor.hpp`
100.0% 5/5	2/2	`source/lib/core/track_registry.cpp`
100.0% 4/4	0/2	`source/lib/core/track_registry.hpp`
100.0% 10/10	1/1	`source/lib/logger/logger.cpp`
100.0% 7/7	1/2	`source/lib/rocprof-sys/library/causal/components/backtrace.hpp`
100.0% 23/23	2/3	`source/lib/rocprof-sys/library/causal/components/progress_point.hpp`
100.0% 3/3	1/2	`source/lib/rocprof-sys/library/causal/experiment.hpp`
100.0% 4/4	1/1	`source/lib/rocprof-sys/library/causal/selected_entry.cpp`
100.0% 5/5	2/3	`source/lib/rocprof-sys/library/causal/selected_entry.hpp`
100.0% 1/1	1/1	`source/lib/rocprof-sys/library/components/backtrace.hpp`
100.0% 11/11	1/2	`source/lib/rocprof-sys/library/components/backtrace_metrics.hpp`
100.0% 4/4	0/2	`source/lib/rocprof-sys/library/components/backtrace_timestamp.hpp`
100.0% 1/1	1/1	`source/lib/rocprof-sys/library/components/callchain.hpp`
100.0% 8/8	2/2	`source/lib/rocprof-sys/library/components/ensure_storage.hpp`
100.0% 1/1	0/1	`source/lib/rocprof-sys/library/components/fork_gotcha.hpp`
100.0% 2/2	0/1	`source/lib/rocprof-sys/library/components/mpi_gotcha.hpp`
100.0% 533/533	12/12	`source/lib/rocprof-sys/library/components/mpip.hpp`
100.0% 3/3	0/1	`source/lib/rocprof-sys/library/components/pthread_create_gotcha.hpp`
100.0% 24/24	1/2	`source/lib/rocprof-sys/library/components/pthread_mutex_gotcha_policy.hpp`
100.0% 11/11	0/3	`source/lib/rocprof-sys/library/perf.hpp`
100.0% 66/66	5/5	`source/lib/rocprof-sys/library/pmc/collectors/cpu/cache_policy.hpp`
100.0% 3/3	0/2	`source/lib/rocprof-sys/library/pmc/collectors/cpu/types.hpp`
100.0% 9/9	0/2	`source/lib/rocprof-sys/library/pmc/collectors/gpu/types.hpp`
100.0% 20/20	4/4	`source/lib/rocprof-sys/library/pmc/collectors/nic/perfetto_policy.hpp`
100.0% 2/2	0/2	`source/lib/rocprof-sys/library/pmc/common/types.hpp`
100.0% 34/34	3/4	`source/lib/rocprof-sys/library/rocprofiler-sdk/marker_writer.hpp`
100.0% 3/3	1/2	`source/lib/rocprof-sys/library/rocprofiler-sdk/trace_control.hpp`
100.0% 2/2	0/1	`source/lib/rocprof-sys/library/runtime.hpp`
100.0% 10/10	3/3	`source/lib/rocprof-sys/library/thread_data_growth.hpp`
100.0% 5/5	1/1	`source/lib/rocprof-sys/library/thread_deleter.hpp`

Distribution

Range	Files
0-20%	54
20-50%	22
50-80%	70
80-100%	174

Fixes from PR #1 review of the in-device (wall_clock64) A2A timing path: - gin_sdma_devtime.h: split measure() into launch-all then sync-all passes. Synchronizing GPU i before launching GPU i+1 deadlocks under -g N>1 because the collective bodies do an all-ranks device barrier and each local GPU is its own rank (GPU 0 waited on GPUs never launched). Mirrors the startColl()/completeColl() launch-then-complete split. - gin_sdma_devtime.h: reduce via the harness Allreduce() helper instead of a direct MPI_Allreduce on MPI_COMM_WORLD. MPI is MPI_THREAD_SINGLE, so under -t N>1 all threads would issue concurrent collectives (can hang); Allreduce() combines across threads, serializes MPI to the last thread, and broadcasts back -- also fixing localMaxUs not being combined across threads. - common.h/common.cu: declare Allreduce<T> and add double/long long explicit instantiations so gin_sdma_devtime.h (via alltoall.cu) can call it. - alltoall.cu: gate AlltoAllDeviceTime to deviceImpl 3/4 (GIN/Hybrid) only; -D1/-D2 (Nvl LSA) lack GIN provisioning and would fault/hang. - alltoall.cu: clamp skip to >= 0 so the start stamp is always written (a negative skip left start_time[] uninitialized, skewing the mode-2 metric). - common.cu: document that mode-2 datacheck validates the production path; the timed kernels reuse the same __device__ collective functions, so only the loop/skip/stamp wrapper is not exercised by datacheck. Co-authored-by: Cursor <cursoragent@cursor.com> modified: projects/rccl-tests/src/common.cu

Addresses PR #1 review (#5): in device-time-only mode the datacheck loop re-inits buffers and validates the PRODUCTION kernel, not the timed kernel. Add NCCL_GIN_ANVIL_DEVTIME_CHECK=1 (default off): after the deviceTime hook runs skip+loop back-to-back collectives into recvbuff, validate that output against `expected` before the datacheck loop overwrites it. Each timed iteration is a full deterministic collective over the unchanged sendbuff, so the final recvbuff equals `expected`. Wrong-element counts are summed across threads/ranks; any mismatch fails the run. Requires datacheck on (populates `expected`). Default behavior is unchanged. Co-authored-by: Cursor <cursoragent@cursor.com>

@1

Addresses three reviewer (pvallem) findings on the in-device timing path: 1. Silent 0.00 us / inf GB/s in device-time-only mode (mode 2). When the timing hook produced no measurement -- the timed kernel never ran (non-GIN impl -D1/-D2, unsupported op/type, count==0, or loop<1) -- deviceDeltaSec stayed 0 and was reported as THE metric, so getBw divided by zero and printed inf GB/s with no warning. AlltoAllDeviceTime now returns a negative sentinel (-1) when the grid produced no positive busy window, and BenchTime emits a one-line WARN and skips the row instead of a bogus metric. 2. NCCL_GIN_ANVIL_DEVTIME_CHECK validated the wrong buffer when the kernel did not run: on early-return the stale warmup/production result passed (false green), and under -w 0 the raw init scratch failed (false red). The check is now reached only after the WARN+skip guard above, i.e. only when the timed kernel actually ran and recvbuff holds its output. 3. Single wall-clock rate applied to all GPUs. wallClockRateKhz() sampled the rate once from the current device; on a mixed-clock node that skews the non-reference GPUs' cycle->us conversion. It now takes a device ordinal and is sampled per-GPU inside the Pass-2 stamp-reduction loop of measure(). Verified on smci355 (MI355X/gfx950), -DENABLE_DEVICE_API=ON -DENABLE_ROCSHMEM=ON, 8-GPU GIN-SDMA AllToAll (128 B..1 MiB): mode 1 augment lines and mode 2 device metric agree (~40.9 us @1 MiB); mode 2 + DEVTIME_CHECK passes with 0 wrong elements; mode 2 with -D1 now emits the WARN+skip for every size with zero inf and zero bogus 0.00 us rows. Co-authored-by: Cursor <cursoragent@cursor.com>

…OCm#11580) ## Motivation `ihipLaunchKernel` could be made to dereference an arbitrary host pointer as an `amd::Kernel*`, killing the process with `SIGSEGV` inside `hipLaunchKernel` instead of returning an error. `StatCO::GetFunc()` returns `hipErrorInvalidSymbol` for any pointer that is not in the registered function table. `ihipLaunchKernel` took that as a cue to `reinterpret_cast` the pointer to `hipFunction_t`, and `hip::asKernel()` is a raw `reinterpret_cast<amd::Kernel*>`, so `ihipLaunchKernel_validate()` goes on to read the caller's pointer as a C++ object. The `if (!kernel)` check there does not help, since non-null garbage passes it. The same cast has crashed before through a different error code: kpack fat binaries with `ROCM_KPACK_ENABLED=OFF` made `GetFunc()` fail with `hipErrorNotSupported`, which was fixed by propagating that error in e14bd06 (ROCm#10150). `hipErrorInvalidSymbol` was the last code still reaching the cast. ## Technical Details The cast cannot simply be dropped. A caller may legitimately pass a handle from `hipModuleGetFunction` or `hipGetFuncBySymbol` straight into `hipLaunchKernel`, and that only works because of it — both confirmed against an unpatched build. So the cast stays, but is taken only for a pointer the runtime actually produced; anything else returns `hipErrorInvalidDeviceFunction`. Recognising a kernel handle needed a registry. `PlatformState` now keeps one in a set, populated in `Function::BuildKernel()` and cleared in `~Function()`. `BuildKernel()` is the only place an `amd::Kernel` is constructed, so the set covers every handle the runtime hands out, from statically registered fat binaries and dynamically loaded modules alike. This replaces `IsValidDynFunc()`, which scanned every function of every loaded module on each call and only ever matched dynamic modules. The new lookup is O(1), and is renamed `IsValidFuncHandle()` because it now answers for static handles too. `DynCO::isValidDynFunc()` and `Function::IsValidDynFunc()` become unused and are removed. The set takes its own leaf mutex rather than `PlatformState::lock_`. Kernels are built with `sclock_`, `dclock_` or `FatBinaryLock` already held, and other paths run `lock_ -> sclock_`, so reusing `lock_` here would invert that order. Two incidental improvements fall out: - `Function::IsValidDynFunc()` compared only against `dFunc_[ihipGetDevice()]`, while `ihipLaunchKernel` resolves against the stream's device, so a valid handle could be rejected when the two differed. A flat set has no such mismatch. - A handle used after `hipModuleUnload` is now rejected instead of dereferencing the released kernel. ## Issue Tracking GitHub issue: ROCm#11563 This covers that issue's request that "HIP validates this condition and returns an error rather than dereferencing invalid state in `hipLaunchKernel`". It does **not** explain the reporter's gfx950 VF `hipErrorInvalidValue` / `hipErrorOutOfMemory`: their reproducer passes on gfx942 with ROCm 10.1, including with `TMPDIR` on a full `/dev/shm`, and neither HIP nor ROCr reads `TMPDIR`. Those errors look like device state specific to that VF and want separate investigation. ## Test Plan gfx942 (MI300X). Every case was run against a baseline library built from the same tree, so this change is the only difference: 1. `hipLaunchKernel` with an unregistered host function address, with a zeroed heap buffer, and with an unmapped address. 2. A `hipModuleGetFunction` handle passed into `hipLaunchKernel`. 3. A `hipGetFuncBySymbol` handle passed into `hipLaunchKernel`. 4. `hipLaunchKernel` with a module's handle after `hipModuleUnload`. ## Test Result | case | before | after | | --- | --- | --- | | unregistered host function address | SIGSEGV | `hipErrorInvalidDeviceFunction` | | zeroed heap buffer | SIGSEGV | `hipErrorInvalidDeviceFunction` | | unmapped address | SIGSEGV | `hipErrorInvalidDeviceFunction` | | `hipModuleGetFunction` handle | launches, correct result | unchanged | | `hipGetFuncBySymbol` handle | launches, correct result | unchanged | | after `hipModuleUnload` | `hipErrorInvalidDeviceOrdinal`, released kernel dereferenced | `hipErrorInvalidDeviceFunction` | For reference, the crash signature this removes, with the host stub address reaching `getDeviceKernel` as `this`: ``` #0 amd::Symbol::getDeviceKernel hashtable.h:1954 <-- deviceKernels_.find() faults #1 amd::Kernel::getDeviceKernel rocclr/platform/kernel.cpp:35 #2 hip::ihipLaunchKernel_validate hipamd/src/hip_module.cpp:356 #3 hip::ihipModuleLaunchKernel hipamd/src/hip_module.cpp:511 #4 hip::ihipLaunchKernel hipamd/src/hip_platform.cpp #5 hip::hipLaunchKernel_common hipamd/src/hip_module.cpp:903 #6 hipLaunchKernel ``` ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md.

…l api tracing (ROCm#11798) ## Motivation ML API tracing emits one ROCTX range per call, with operator arguments and thread ids on the marker, instead of stacked marker paths. Marker `Function` text changes from stacked `nn.Module.Net.forward/nn.Module.Linear.forward:#1@train.py:27/#2@train.py:19` paths to one-level `nn.Module.Linear.forward:train.py:19|seqNr=n/a|tid=n/a|ftid=n/a|ltid=n/a|scope=n/a|args=n/a|torch`. Parsing the values and correlation is done in `analyze `phase. ## Technical Details - Marker: `encode(name):<location>|seqNr=<v>|tid=<v>|ftid=<v>|ltid=<v>|scope=<v>|args=<v>|<backend>`. `/` and `%` in the name are encoded. Backend is a raw trailing `|torch` or `|triton`. Python wraps use user `file:line` (or `n/a`) and leave seqNr/tid/ftid/ltid/scope as `n/a`. RecordFunction uses `n/a` as location and fills the other fields when available. - `src/lib/torch_trace_collector`: RecordFunction callback pushes that range. Args from inputs and schema names, capped and encoded. Library: `torch_trace_collector.so`. - Removed stacked collector: snapshot store, marker stack, leaf context, pybind, C++ collector tests. - `src/utils/inject_roctx/_backends/torch_trace_collector.py`: ctypes-load the `.so`, call `torch_trace_collector_install`. Push/pop launcher tid around selected autograd wraps. - `src/utils/inject_roctx`: torch and triton wraps use the same wire; each call pushes and pops its own range. A Triton `JITFunction.run` that calls `CompiledKernel.run` emits both markers. - `src/utils/utils_profile.py`: `save_ml_api_trace_inputs` copies marker and counter CSVs unchanged. - Removed tests: `test_torch_cpp_loader`, `test-torch-trace-collector`, `test_torch_trace_coverage`, `test_profile_triton_trace`, `test_triton_backend_reentrancy_dedups_nested_launch`. - This wire no longer matches the old stacked `:#` shape. Analyze does not work on this output; that is ROCm#11830. - Restore Torch-independent collector builds with a generic `torch_trace_collector.so`, using the application's PyTorch at runtime. - Fix missing tensor metadata and stale operator argument names, limit tensor-list traversal, and preserve the correct subclass binding for inherited `Function.apply`. - Regression tests for these fixes are in ROCm#12809, stacked on ROCm#11834. ## Issue Tracking JIRA ID: AIPROFCOMP-860 ## Test Plan `rocprof-compute profile --experimental -b 2.1.1 --torch-trace -n simpleCode -- python3 simple_net.py --dry-run --epochs 1 --batch-size 1` ## Test Result From marker_trace.csv, `Domain,Function,Process_Id,Thread_Id,Correlation_Id,GUID,Start_Timestamp,End_Timestamp MARKER_CORE_RANGE_API,nn.Module.SimpleNet.__init__:simple_net.py:30|seqNr=n/a|tid=n/a|ftid=n/a|ltid=n/a|scope=n/a|args=n/a|torch,6921,6921,1,000084b2-a044-7044-b2bd-c1a4d65f6110,2226300689148270,2226300689187709 MARKER_CORE_RANGE_API,nn.Module.Linear.__init__:simple_net.py:31|seqNr=n/a|tid=n/a|ftid=n/a|ltid=n/a|scope=n/a|args=n/a|torch,6921,6921,2,000084b2-a044-7044-b2bd-c1a4d65f6110,2226300690158420,2226300690174264 MARKER_CORE_RANGE_API,torch.empty:simple_net.py:31|seqNr=n/a|tid=n/a|ftid=n/a|ltid=n/a|scope=n/a|args=n/a|torch,6921,6921,3,000084b2-a044-7044-b2bd-c1a4d65f6110,2226300690210037,2226300690347283 MARKER_CORE_RANGE_API,"aten::empty:n/a|seqNr=n/a|tid=1|ftid=0|ltid=n/a|scope=FUNCTION|args=(size=GenericList, dtype=None, layout=None, device=Device, pin_memory=Bool, memory_format=None)|torch",6921,6921,4,000084b2-a044-7044-b2bd-c1a4d65f6110,2226300690316316,2226300690342836 MARKER_CORE_RANGE_API,aten::detach:n/a|seqNr=0|tid=1|ftid=0|ltid=n/a|scope=FUNCTION|args=(self=float32[20x10])|torch,6921,6921,5,000084b2-a044-7044-b2bd-c1a4d65f6110,2226300690389606,2226300690411058 ` ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/rocm-systems/blob/develop/CONTRIBUTING.md. --------- Co-authored-by: Harsha Bajpai <habajpai@amd.com> Co-authored-by: Vignesh Edithal <Vignesh.Edithal@amd.com> Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>

github-actions Bot added the project: rccl-tests label Aug 12, 2026

dlamd1dai mentioned this pull request Aug 12, 2026

rccl-tests/gin-sdma: cap A2A GIN put to 128 MiB to fix >1 GiB SDMA hang ROCm/rocm-systems#9927

Closed

2 of 3 tasks