feat(rccl/gin-sdma): GIN SDMA Broadcast (-D3 hybrid Anvil-SDMA) - #5
Open
dlamd1dai wants to merge 18 commits into
Open
feat(rccl/gin-sdma): GIN SDMA Broadcast (-D3 hybrid Anvil-SDMA)#5dlamd1dai wants to merge 18 commits into
dlamd1dai wants to merge 18 commits into
Conversation
Port the GIN Anvil-SDMA hybrid AllGather kernel into stage3b, at parity with the AllReduce path: per-rank chunk size gates LSA-direct peer stores (<= sdmaThreshold) vs. all-peers GIN copy-engine puts (> threshold). - all_gather.cu: GinHybridAllGatherKernel (NCCL_GIN_TYPE=5) + host launcher; sizing/threshold/tier/bandwidth logic factored into a testable policy header. - gin_sdma_allgather_policy.h: pure host/device policy helpers (single source of truth for the LSA<->SDMA crossover) shared by kernel and unit tests. - rccl/test: host policy unit test + on-GPU tier-predicate/addressing test. - gin_plugin_anvil_sdma.cc: honor explicit NCCL_GIN_ANVIL_SDMA_THRESHOLD=0. - common.cu: don't report host cputime column for device-impl kernels. Validated on 8x MI355X across 128 B-128 MB: 0 wrong in both tiers, crossover lands exactly at 2 MiB/rank, SDMA tier ~390 GB/s bus bw. Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit ffe0d60)
Mirror the AllToAll GIN-SDMA >1 GiB hang-fix tests (PR ROCm#9927) for AllGather so the multi-channel/SDMA ordering concern is covered by an automated test, not just manual runs. Drives GinHybridAllGatherKernel (-D 3, NCCL_GIN_TYPE=5) with the SDMA tier forced (threshold=0) at per-rank chunks that cross the <=128 MiB segmentation boundary (256 MiB -> 2 seg, 2 GiB -> 16 seg), plus a 2 GiB-total completion guard with a subprocess timeout that turns a reintroduced SDMA hang into a failure. Exact-integer dtypes + built-in data check (-c 1) so a truncated /stale tail can't hide under float tolerance. Opt-in via RCCL_TESTS_GIN_SDMA_AG (needs an 8x MI355X-class node + MPI + a GIN-SDMA build); config mirrors test_AllToAll.py (NP/launcher/xenv/exe/timeout). Includes the harness's intermittent gfx950 cuMem-VMM connectivity-gate retry so the transient abort doesn't flake the test; a genuine data-check mismatch is never retried. Verified on 8x MI355X: 2 GiB-total hang guard passes, 0 wrong, ~425 GB/s (SDMA/SIMPLE tier, 2x128 MiB segments). Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit 38b741a)
Give the GIN Anvil-SDMA hybrid AllGather its own LSA<->SDMA crossover knob, NCCL_GIN_ANVIL_SDMA_THRESHOLD_ALLGATHER, decoupled from the backend gin.put inline-vs-copy-engine threshold (the global NCCL_GIN_ANVIL_SDMA_THRESHOLD / ctx->sdmaThreshold). The device API is untouched: the threshold is host-resolved and passed to GinHybridAllGatherKernel as a kernel argument, consumed only by the host-testable tier predicate (chunkUsesLsaTier); gin.put's signature is unchanged. - gin_sdma_allgather_policy.h: replace the context-based resolveSdmaThreshold() with a pure, precedence-based pickSdmaThreshold() (per-collective env > global env > compiled default) plus the 32 KiB/rank compiled default constant. - all_gather.cu: host AllGatherResolveSdmaThreshold() reads the per-collective env, falls back to the global env then the compiled default, parsed as 64-bit (explicit 0 honored -> all-SDMA tier; >2 GiB values do not wrap). Thread the resolved value through a new AllGather-specific launch helper into the kernel; guard the RunColl -D 3 case with ENABLE_DEVICE_API so a device-API-off build cleanly falls through to testNotImplemented (matches the shared launcher's prior behavior). - gin_sdma_allgather_policy_test.cpp: cover pickSdmaThreshold precedence, explicit-0, and the 64-bit no-wrap path. Validated: 13/13 host policy unit tests; docker device compile + A2A and AllGather GIN smoke gates pass on 8x MI355X (0 wrong). Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit 212d3fc)
AllGather launched a fixed deviceCtaCount (-V) grid for all sizes, which over-
subscribes the GIN-put/SDMA tier: only nRanks threads issue the puts (the copy
engines move the bytes), so extra CTAs are pure barrier/signal overhead. Mirror
the ReduceScatter design and self-select a size-adaptive CTA count decoupled from
-V, keyed off the same LSA<->SDMA tier predicate the kernel evaluates:
* LSA-direct tier (chunk <= threshold): 16 CTAs (grid-stride store scales with
threads; peaks ~16 on 8x MI355X, tiny sizes are latency-bound/CTA-indifferent).
* GIN-put/SDMA tier (chunk > threshold): 4 CTAs (stable near-peak 1 MiB-128 MiB).
Measured gains vs the old fixed -V 32 launch: 512 KiB +21%, 2 MiB +18%, 8 MiB
+13%, 32 MiB +5%; LSA tier unchanged. Datacheck bit-exact 128 B-256 MiB; A2A/RS
unaffected. Add allGatherCtas()/allGatherMaxCtas() to the policy header (host-
unit-testable), NCCL_GIN_ANVIL_AG_CTAS to pin a fixed count (diagnostic), an
explicit-grid launcher, and size the devComm barrier/signal pool to
max(deviceCtaCount, allGatherMaxCtas()) so the self-selected grid never exceeds it.
Co-authored-by: Cursor <cursoragent@cursor.com>
(cherry picked from commit e5ce620)
…tage3e Bring the in-kernel wall_clock64 device-timing hook (AllGatherDeviceTime, opt-in via NCCL_GIN_ANVIL_DEVICE_TIMING) to the AG branch, reconciled with the size- adaptive CTA selection just cherry-picked here. To share one implementation between the production and timed launches, factor the LSA/SDMA logic out of GinHybridAllGatherKernel into ginAllGatherBody<T>() and add GinHybridAllGather- TimedKernel that loops the body under one persistent launch. The devtime path self-selects the same allGatherCtas() count as the perf path (verified: the resolved per-collective sdmaThreshold so the timed tier matches the launched cfg. common.h already carries the testColl.deviceTime field and gin_sdma_devtime.h exists on this branch (used by AllToAll), so only all_gather.cu changes: include the scaffold, add the body/timed kernel, add AllGatherDeviceTime, and wire it as the testColl.deviceTime member. Builds clean; datacheck bit-exact; devtime prints. Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit 4435b34)
Reword the size-adaptive CTA comments so the AllGather design does not reference the ReduceScatter design (this is the AG branch). No functional change. Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit 3ae778c)
Cover the new allGatherCtas()/allGatherMaxCtas() policy helpers: LSA vs SDMA tier selection keyed off the crossover predicate, tracking of a moved threshold, env pin + 128 clamp + 0-falls-back-to-ladder, and max-covers-both-tiers. 18/18 host policy tests pass. Co-authored-by: Cursor <cursoragent@cursor.com> (cherry picked from commit e35adb7)
Resolves the seven review comments on #4 (GIN SDMA AllGather): - rccl-tests: clamp NCCL_GIN_ANVIL_AG_CTAS to the launched barrier/signal pool (allGatherPoolCtas = max(-V, ladder peak)); a pin above the pool previously launched more CTAs than slots -> OOB on the blockIdx.x-indexed pools. - rccl: fold Anvil-SDMA conn-check infra failures (write/verify kernel, D2H copy, slot-reset memset) into the collective decision instead of a per-rank goto, so a single-rank failure aborts all ranks together rather than leaving peers hung on bootstrapBarrier/allgather. - rccl: scope the gfx950 cuMem free-skip to the peer-aperture unmap in ncclCuMemFreeAddr only; ncclCuMemFree (physical handle) and ncclCudaFree (ordinary buffers) now free normally, so long-lived procs no longer leak on every ncclCommDestroy. Leak is confined to the peer VA aperture. - rccl: replace __builtin_unreachable() in the ANVIL_SDMA/ROCSHMEM_GDA get() specializations with __builtin_trap() so a future get() fails loudly. Uses the compiler builtin (no HIP runtime include) so it also compiles in the symmetric gensrc TU that pulls these headers via nccl_device.h. - rccl-tests: share allGatherRecvSliceOffset() between AllGatherLsaDirect and the GPU addressing test so a wrong slice formula is caught, not just self-consistent; scope the test comment to what it does/doesn't cover. - rccl: fix the rocshmem-gda template test stub to define atomic_add / atomic_add_single (was stale atomic_nofetch, broke --rocshmem-gin builds). - rccl-tests: replace tautological policy asserts with a behavioral "grid never exceeds pool" regression guarding the CTA-pin clamp. Validated on 8x MI355X (gfx950): --rocshmem-gin image builds clean; host policy unit tests pass (AllGather 19/19 incl. the new clamp+grid guards, SDMA 5/5); all_gather_perf -D3 (NCCL_GIN_TYPE=5) 128 B-128 MiB is bit-exact (#wrong=0) in both the LSA-direct and SDMA tiers with no hang. Co-authored-by: Cursor <cursoragent@cursor.com>
Replace stale NCCL_GIN_ANVIL_AG_DEVTIME_* env vars with the harness --device_timing/-L/-P tier knobs (matching A2A), gate to -D 3, and return the -1.0 mode-2 sentinel when no measurement is produced. Co-authored-by: Cursor <cursoragent@cursor.com>
…sync check Add LSA signal connectivity self-test (localFail collective abort) and the gin_anvil_sdma_oss7_device launchers so the PR branch matches the wip branch's plugin safety path. Include the AllGather manifest sync script used to keep the six core AG files identical across gin-stage3b-sdma-ag and the NCCL 2.30.7 wip. Co-authored-by: Cursor <cursoragent@cursor.com>
Remove the AllGather manifest and sync-check scripts; that validation tooling belongs on the NCCL 2.30.7 wip branch only. The conn-check plugin fix from the prior commit is unchanged. Co-authored-by: Cursor <cursoragent@cursor.com>
…ates Reconcile PR #4 with merged upstream ROCm#10658 (A2A devtime host helpers, BenchTime skip/augment/check path) and ROCm#10672 (GDA/SDMA device template tests + GIN CI wiring). AllGather devtime now buffers via devtimeAugmentLine to match the A2A augment pattern; AllGather-specific unit test targets remain. Co-authored-by: Cursor <cursoragent@cursor.com>
dlamd1dai
force-pushed
the
users/dondai/gin-stage3c-sdma-bc
branch
from
August 29, 2026 00:28
04450da to
abd91c5
Compare
dlamd1dai
changed the base branch from
develop
to
users/dondai/gin-stage3b-sdma-ag
August 29, 2026 00:29
rocprofiler-systems CI SummaryTest Results (for commit
|
| Workflow | Runs | Avg Test Time | Tests | Skipped | Failed |
|---|---|---|---|---|---|
| Ubuntu 22.04 | 5 | 30m 33s | 2,570 | 100 | 0 |
| Ubuntu 24.04 | 5 | 32m 23s | 2,590 | 110 | 0 |
| Debian | 4 | 37m 00s | 2,168 | 78 | 0 |
| RHEL | 12 | 36m 08s | 6,168 | 277 | 0 |
| System deps | 9 | 29m 19s | 4,822 | 197 | 0 |
| Total | 35 | 33m 09s | 18,318 | 762 | 0 |
Failed Tests
No failed tests reported.
Sanitizers
| Thread | Undefined | Address |
|---|---|---|
| Passed | Passed | Passed |
|
✅ MINOR ABI BREAKAGE resolved - ABI compliance check is now passing! |
|
| 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/61 | 0/6 | source/bin/rocprof-sys-avail/avail.hpp |
| 0.0% 0/258 | 0/12 | source/bin/rocprof-sys-avail/generate_config.cpp |
| 0.0% 0/102 | 0/6 | source/bin/rocprof-sys-avail/hw_counter_query.cpp |
| 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/1536 | 0/87 | source/bin/rocprof-sys-instrument/rocprof-sys-instrument.cpp |
| 0.0% 0/157 | 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/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/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/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/160 | 0/12 | source/lib/core/gpu_metrics.cpp |
| 0.0% 0/33 | 0/6 | source/lib/core/mpi.hpp |
| 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/8 | 0/3 | 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/817 | 0/32 | 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/14 | 0/6 | source/lib/rocprof-sys-causal-api/causal_api.cpp |
| 0.0% 0/539 | 0/85 | 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/84 | 0/20 | 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/1268 | 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/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/74 | 0/16 | source/lib/rocprof-sys/regions.cpp |
| 0.0% 0/82 | 0/4 | source/lib/rocprof-sys/timeout.cpp |
| 0.0% 0/440 | 0/45 | 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 |
| 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 |
| 9.6% 17/178 | 4/25 | source/lib/core/constraint.cpp |
| 13.2% 30/227 | 8/25 | source/bin/rocprof-sys-avail/common.cpp |
| 15.4% 4/26 | 0/6 | source/lib/core/trace_cache/metadata_registry.hpp |
| 18.0% 32/178 | 6/23 | source/lib/core/gpu.cpp |
| 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.0% 4/21 | 0/2 | source/lib/backends/rocprofiler_sdk/wrapper.hpp |
| 19.4% 6/31 | 2/5 | source/bin/common/common_utils.hpp |
🟠 20-50% (14 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.2% 13/56 | 1/4 | source/lib/core/mproc.cpp |
| 32.1% 52/162 | 7/17 | build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/input/json_sax.hpp |
| 34.8% 147/423 | 14/46 | build/coverage/_deps/fmt-src/include/fmt/format.h |
| 35.3% 612/1732 | 156/302 | source/lib/core/config.cpp |
| 36.5% 23/63 | 6/12 | source/lib/core/agent_manager.cpp |
| 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 |
| 46.0% 433/941 | 10/48 | source/bin/rocprof-sys-avail/avail.cpp |
| 47.4% 222/468 | 17/39 | 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% (35 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 |
| 51.9% 220/424 | 25/37 | source/bin/common/common_utils.cpp |
| 52.0% 13/25 | 2/3 | source/lib/rocprof-sys/library/pmc/collectors/gpu_perf_counter/gpu_perf_counter_traits.hpp |
| 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 |
| 65.4% 17/26 | 2/3 | source/lib/core/config.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.9% 220/306 | 15/17 | 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% 7/9 | 0/1 | source/bin/rocprof-sys-avail/info_type.hpp |
| 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% (81 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 |
| 82.8% 48/58 | 6/7 | build/coverage/_deps/nlohmann_json-src/include/nlohmann/detail/conversions/to_json.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 |
| 89.0% 276/310 | 30/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.1% 221/248 | 23/33 | source/lib/core/sdk/tracing-config.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 |
| 91.7% 44/48 | 4/4 | source/lib/core/utility.cpp |
| 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.1% 16/17 | 3/3 | source/bin/rocprof-sys-avail/component_categories.hpp |
| 94.2% 129/137 | 29/32 | source/lib/backends/rocprofiler_sdk/backend.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 |
| 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% 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% 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% 5/5 | 1/1 | source/lib/common/string_utility.hpp |
| 100.0% 8/8 | 0/2 | source/lib/common/version.hpp |
| 100.0% 2/2 | 1/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% | 191 |
| 20-50% | 14 |
| 50-80% | 35 |
| 80-100% | 81 |
E2E Tests
Code Coverage: e2e-tests
Lines: 🟡 64.3% — 22,674/35,268 across 321 files
Functions: 🟡 64.2% — 2,347/3,653
🔴 0-20% (49 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/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/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/8 | 0/3 | source/lib/core/sdk/tracing-config-deps.hpp |
| 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/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/440 | 0/45 | source/python/libpyrocprofsys.cpp |
| 0.6% 1/166 | 1/13 | source/lib/core/perfetto.cpp |
| 2.6% 33/1268 | 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 |
| 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.2% 12/74 | 4/16 | source/lib/rocprof-sys/regions.cpp |
| 18.0% 32/178 | 6/23 | source/lib/core/gpu.cpp |
| 19.0% 4/21 | 0/2 | source/lib/backends/rocprofiler_sdk/wrapper.hpp |
🟠 20-50% (25 files)
| Lines | Functions | File |
|---|---|---|
| 21.5% 17/79 | 3/5 | source/lib/core/perf.cpp |
| 23.3% 190/817 | 9/32 | source/lib/core/trace_cache/rocpd_processor.cpp |
| 23.6% 122/516 | 24/71 | source/lib/core/trace_cache/metadata_registry.cpp |
| 27.8% 27/97 | 4/8 | source/lib/rocprof-sys-rt/src/RTposix.c |
| 29.8% 228/764 | 32/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% 30/84 | 10/20 | source/lib/rocprof-sys/api.cpp |
| 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 |
| 47.7% 52/109 | 12/30 | source/lib/rocprof-sys/library/tracing.hpp |
| 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 |
🟡 50-80% (72 files)
| Lines | Functions | File |
|---|---|---|
| 50.0% 8/16 | 2/4 | source/lib/binary/address_multirange.cpp |
| 50.0% 7/14 | 2/6 | source/lib/rocprof-sys-causal-api/causal_api.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.4% 213/423 | 23/46 | build/coverage/_deps/fmt-src/include/fmt/format.h |
| 50.8% 128/252 | 25/37 | source/lib/rocprof-sys/library/kokkosp.cpp |
| 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 |
| 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.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 |
| 60.9% 131/215 | 12/30 | source/lib/rocprof-sys/library/components/backtrace_metrics.cpp |
| 61.8% 289/468 | 21/39 | 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.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 |
| 63.3% 155/245 | 23/31 | source/lib/rocprof-sys/library/perf.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.1% 465/673 | 34/46 | source/bin/rocprof-sys-instrument/details.cpp |
| 69.2% 373/539 | 40/85 | source/lib/rocprof-sys-dl/dl.cpp |
| 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.3% 64/91 | 5/7 | source/bin/rocprof-sys-instrument/function_signature.cpp |
| 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 |
| 72.6% 722/994 | 52/65 | source/lib/rocprof-sys/library/sampling.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.2% 227/306 | 16/17 | source/bin/common/preset_registry.cpp |
| 74.3% 411/553 | 30/33 | source/lib/rocprof-sys/library/causal/data.cpp |
| 74.5% 35/47 | 4/6 | source/lib/rocprof-sys/library/pmc/collectors/cpu/cpu_traits.hpp |
| 75.0% 15/20 | 5/5 | source/bin/rocprof-sys-instrument/function_signature.hpp |
| 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.5% 1159/1536 | 51/87 | source/bin/rocprof-sys-instrument/rocprof-sys-instrument.cpp |
| 75.6% 90/119 | 11/12 | source/lib/core/trace_cache/post_processor.cpp |
| 76.5% 408/533 | 40/41 | source/bin/rocprof-sys-instrument/module_function.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.4% 328/424 | 35/37 | source/bin/common/common_utils.cpp |
| 77.7% 482/620 | 27/36 | source/lib/rocprof-sys/library.cpp |
| 77.9% 113/145 | 26/31 | source/lib/rocprof-sys/library/thread_data.hpp |
| 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 |
| 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% (175 files)
| Lines | Functions | File |
|---|---|---|
| 80.0% 4/5 | 0/1 | source/bin/rocprof-sys-instrument/log.hpp |
| 80.5% 62/77 | 9/10 | source/lib/rocprof-sys/library/pmc/device_providers/rocprofiler_sdk/provider.hpp |
| 81.2% 69/85 | 17/22 | source/lib/rocprof-sys/library/causal/components/progress_point.cpp |
| 81.2% 13/16 | 1/1 | source/lib/rocprof-sys-dl/main.c |
| 81.4% 131/161 | 24/27 | source/lib/rocprof-sys/library/thread_info.cpp |
| 81.8% 1417/1732 | 277/302 | source/lib/core/config.cpp |
| 81.8% 81/99 | 11/11 | source/lib/core/perfetto/engine_impl.hpp |
| 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 |
| 83.4% 131/157 | 12/12 | source/bin/rocprof-sys-instrument/rocprof-sys-instrument.hpp |
| 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.0% 317/373 | 18/29 | source/lib/rocprof-sys/library/causal/experiment.cpp |
| 85.1% 74/87 | 14/16 | source/lib/rocprof-sys/library/causal/delay.cpp |
| 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.7% 98/113 | 9/11 | source/lib/rocprof-sys/library/process_sampler.cpp |
| 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.5% 42/48 | 4/5 | source/lib/core/output_file_registry.cpp |
| 87.6% 141/161 | 13/16 | source/lib/backends/procfs/backend.hpp |
| 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% 299/336 | 16/20 | source/lib/rocprof-sys/library/components/pthread_create_gotcha.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.1% 221/248 | 23/33 | source/lib/core/sdk/tracing-config.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% 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.1% 64/71 | 7/10 | source/lib/rocprof-sys/library/components/callchain.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.3% 178/195 | 14/18 | source/lib/rocprof-sys/library/components/shmem_gotcha.hpp |
| 91.3% 21/23 | 4/5 | source/lib/core/containers/static_vector.hpp |
| 91.5% 107/117 | 5/5 | source/lib/binary/analysis.cpp |
| 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.4% 239/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.2% 129/137 | 29/32 | source/lib/backends/rocprofiler_sdk/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.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.7% 42/43 | 7/11 | source/lib/core/demangler.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.0% 147/150 | 9/10 | source/lib/common/md5sum.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 |
| 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% 10/10 | 1/2 | build/coverage/include/profiler-hub/writer_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% 5/5 | 1/1 | source/lib/common/string_utility.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% 8/8 | 0/2 | source/lib/common/version.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% 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% 85/85 | 11/12 | source/lib/core/state.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% | 49 |
| 20-50% | 25 |
| 50-80% | 72 |
| 80-100% | 175 |
Remove a stray closing brace in alloc.h that broke RCCL host builds, and align AllGatherGetDevCommRequirements with common.h by taking ncclComm_t and querying comm properties internally (matching alltoall/all_reduce). Co-authored-by: Cursor <cursoragent@cursor.com>
Port the GIN Anvil-SDMA hybrid Broadcast design: policy header, tiered broadcast.cu kernels (LL/LSA/flat/scatter+AG/ring), devtime hook via the shared augment-line scaffold, broadcast_perf RDC link, and opt-in pytest multi-segment regression tests (-D 3). Co-authored-by: Cursor <cursoragent@cursor.com>
Wire rccl-UnitTestsGinSdmaBroadcastPolicy (host GTest for gin_sdma_broadcast_policy.h) and add broadcast_perf -D 3 to the GIN CI test matrix. Co-authored-by: Cursor <cursoragent@cursor.com>
… API. Fix GetBw/GetDevCommRequirements signatures and testColl init so broadcast_perf -D 3 builds against common.h on the wip branch. Co-authored-by: Cursor <cursoragent@cursor.com>
Mirrors the existing 2 GiB completion test with default tier selection for multi-GB ring/SAG coverage. Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
dlamd1dai
force-pushed
the
users/dondai/gin-stage3c-sdma-bc
branch
from
August 29, 2026 05:55
abd91c5 to
bcdd42d
Compare
dlamd1dai
force-pushed
the
users/dondai/gin-stage3b-sdma-ag
branch
from
August 29, 2026 19:23
cbdc463 to
b9b0cb4
Compare
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Stacked on #4 — adds GIN Anvil-SDMA hybrid Broadcast (
-D 3,NCCL_GIN_TYPE=5) on top of the AllGather work in #4.Broadcast
users/dondai/gin-stage3c-sdma-bcast:gin_sdma_broadcast_policy.h, tiered kernels (LL / LSA / flat / scatter+AG / ring), devtime hook via the shared augment-line scaffold, and broadcast RDC linkrccl-UnitTestsGinSdmaBroadcastPolicy) and opt-in pytest multi-segment regressions (test_Broadcast.py)broadcast_perf -D 3withNCCL_GIN_TYPE=5alloc.hhost-build fix + AllGatherGetDevCommRequirementsAPI alignment)broadcast.cualigned withcommon.htestEngine API (GetBwsize_t,ncclCommQueryProperties,getCollImplInfoslot)4 commits on top of #4 base (+ formatting), 7 files changed.
Depends on: #4
Test plan
rccl-UnitTestsGinSdmaBroadcastPolicy— broadcast tier policy (host GTest wired in CMakeLists; policy UT not in docker image — validated via tier logic in pytest + SUT gates)broadcast_perf -D 3withNCCL_GIN_ENABLE=1 NCCL_GIN_TYPE=5(128 B – multi-segment) — PASS on 8× MI355X (smci355-ccs-aus-m03-17): BC-C1 host baseline NP=8 128M; BC-C2 GIN NP=8 128M#wrong=0pytest projects/rccl-tests/test/test_Broadcast.py -k gin_sdma(opt-in) — includes 2 GiB + 4 GiB hang guards; SUT BC-C2-L PASS at 4G (~373 GB/s,#wrong=0, root=0)broadcast_perfentry ingin-tests.jsonviarun-gin-ci.sh— pending CI re-run after rebaseSUT evidence (MI355X, docker
rccl-gin-gda-sdma-713:latest, COLLECTIVE=bcast rebuild)