Repository navigation
TheRock : Code Coverage - #7783
Open
jainprad wants to merge 82 commits into
Open
TheRock : Code Coverage #7783jainprad wants to merge 82 commits into
jainprad wants to merge 82 commits into
Conversation
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
therock-pr-bot
Bot
added
the
Not ready to Review
|
🎉 All checks passed! This PR is ready for review. |
jainprad
marked this pull request as draft
August 31, 2026 19:42
![github-advanced-security[bot]](https://avatars.githubusercontent.com/in/57789?s=60&v=4){kind=small}
Comment on lines
+121
to
+128
| - name: Checking out repository | ||
| uses: actions/checkout@3d3c42e5aac5ba805825da76410c181273ba90b1 # v7.0.1 | ||
| timeout-minutes: 15 | ||
| with: | ||
| repository: ${{ inputs.repository || github.repository }} | ||
| ref: ${{ inputs.ref }} | ||
|
|
||
| - name: Checkout CI config |
Comment on lines
+117
to
+123
| if: >- | ||
| ${{ | ||
| !cancelled() && !failure() && | ||
| !contains(inputs.prebuilt_stages, 'compiler-runtime') && | ||
| (inputs.build_stages == '' || | ||
| contains(format(',{0},', inputs.build_stages), ',compiler-runtime,')) | ||
| }} |
Contributor
There was a problem hiding this comment.
Shouldn't we have only one if block? AFAIK, GitHub Actions rejects duplicate keys in a mapping. Please check and merge the if block conditions into one block, delete the duplicate.
Member
There was a problem hiding this comment.
Yes, this is invalid yaml and the workflow does not run: https://github.com/ROCm/TheRock/actions/runs/33431950362
Invalid workflow file: .github/workflows/multi_arch_ci.yml#L139
error parsing called workflow
".github/workflows/multi_arch_ci.yml"
-> "./.github/workflows/multi_arch_ci_linux.yml"
--> "./.github/workflows/multi_arch_build_portable_linux.yml"
: (Line: 124, Col: 5): 'if' is already defined, (Line: 158, Col: 5): 'if' is already defined, (Line: 192, Col: 5): 'if' is already defined, (Line: 223, Col: 5): 'if' is already defined, (Line: 270, Col: 5): 'if' is already defined, (Line: 306, Col: 5): 'if' is already defined, (Line: 341, Col: 5): 'if' is already defined, (Line: 375, Col: 5): 'if' is already defined, (Line: 409, Col: 5): 'if' is already defined, (Line: 443, Col: 5): 'if' is already defined
Member
There was a problem hiding this comment.
I think you can probably undo the changes in this file, and then call multi_arch_build_portable_linux_artifacts.yml in your workflows as follows:
jobs:
build_instrumented_compiler_runtime:
uses: ./.github/workflows/multi_arch_build_portable_linux_artifacts.yml
with:
stage_name: compiler-runtime
extra_cmake_options: -DCOMPILER_RT_BUILD_PROFILE_ROCM=ON
build_instrumented_math_libs:
needs: build_instrumented_compiler_runtime
uses: ./.github/workflows/multi_arch_build_portable_linux_artifacts.yml
with:
stage_name: math-libs
stage_projects: "hiprand;rocrand;hipfft;rocfft" # All projects to instrument
extra_cmake_options: -DTHEROCK_COVERAGE_ROCM_LIBRARIES_ALL=ON
test_coverage:
needs: build_instrumented_math_libs
strategy:
matrix:
project: [hiprand, rocrand, hipfft, rocfft]
uses: ./.github/workflows/test_component.yml
with:
component: ${{ matrix.project }}
coverage_enabled: trueReally all you seem to be doing with these changes is adding a bunch of conditionals so that we can filter everything except for compiler-runtime and math-libs, and eventually rocm-systems, so this might be a more ideal approach.
Author
There was a problem hiding this comment.
Done. Removed all changes from this file. Please review
Comment on lines
151
to
158
| if: >- | ||
| ${{ | ||
| !cancelled() && !failure() && | ||
| !contains(inputs.prebuilt_stages, 'runtime-tests') && | ||
| (inputs.build_stages == '' || | ||
| contains(format(',{0},', inputs.build_stages), ',runtime-tests,')) | ||
| }} | ||
| if: ${{ !cancelled() && !failure() && !contains(inputs.prebuilt_stages, 'runtime-tests') && !contains(inputs.skip_stages, 'runtime-tests') }} |
Contributor
There was a problem hiding this comment.
Same as above - Shouldn't we have only one if block? AFAIK, GitHub Actions rejects duplicate keys in a mapping. Please check and merge the if block conditions into one block, delete the duplicate. This is also applicable to multiple other instances in this yml.
…verage-workflows-therock # Conflicts: # .github/workflows/multi_arch_build_portable_linux.yml
reboss
suggested changes
Sep 2, 2026
reboss
left a comment
reboss
left a comment
Member
There was a problem hiding this comment.
We should try and reduce our footprint for code coverage. We want it to be an add-on that reuses much of the existing process, not a tightly coupled feature that requires major refactoring of existing code.
Comment on lines
+45
to
+57
| build_stages: | ||
| type: string | ||
| default: "" | ||
| description: >- | ||
| Comma-separated allowlist of build stages to run. Empty runs every | ||
| stage. Variant pipelines such as coverage use this to build only the | ||
| stages they need. | ||
| stage_projects: | ||
| type: string | ||
| default: "" | ||
| description: >- | ||
| Space-separated project names to narrow per-arch stages down to | ||
| (e.g. "hiprand"). Empty builds every project in the stage. |
Member
There was a problem hiding this comment.
Do we really need both of these variables for a nightly code coverage run?
Member
There was a problem hiding this comment.
I don't think we do. Please revert the changes in this file, since they are not needed. You can localize your changes to .github/workflows/multi_arch_build_portable_linux_artifacts.yml.
Also, please make sure you are fetching the most recent changes to the RFC and confirming that your branch remains in alignment with what is documented.
Comment on lines
+78
to
+85
| # Coverage instrumented builds need the ROCm profiling runtime | ||
| # (libclang_rt.profile_rocm) so that instrumented device code can write | ||
| # profraw files. The flag is only meaningful to compiler-rt, so forward it. | ||
| if(COMPILER_RT_BUILD_PROFILE_ROCM) | ||
| message(STATUS "COMPILER_RT_BUILD_PROFILE_ROCM enabled - building the ROCm profiling runtime") | ||
| list(APPEND _extra_llvm_cmake_args "-DCOMPILER_RT_BUILD_PROFILE_ROCM=ON") | ||
| endif() | ||
|
|
Member
There was a problem hiding this comment.
This likely is no longer needed. We're doing some testing without this option atm, but feel free to experiment without it on your own
Comment on lines
+117
to
+123
| if: >- | ||
| ${{ | ||
| !cancelled() && !failure() && | ||
| !contains(inputs.prebuilt_stages, 'compiler-runtime') && | ||
| (inputs.build_stages == '' || | ||
| contains(format(',{0},', inputs.build_stages), ',compiler-runtime,')) | ||
| }} |
Member
There was a problem hiding this comment.
I think you can probably undo the changes in this file, and then call multi_arch_build_portable_linux_artifacts.yml in your workflows as follows:
jobs:
build_instrumented_compiler_runtime:
uses: ./.github/workflows/multi_arch_build_portable_linux_artifacts.yml
with:
stage_name: compiler-runtime
extra_cmake_options: -DCOMPILER_RT_BUILD_PROFILE_ROCM=ON
build_instrumented_math_libs:
needs: build_instrumented_compiler_runtime
uses: ./.github/workflows/multi_arch_build_portable_linux_artifacts.yml
with:
stage_name: math-libs
stage_projects: "hiprand;rocrand;hipfft;rocfft" # All projects to instrument
extra_cmake_options: -DTHEROCK_COVERAGE_ROCM_LIBRARIES_ALL=ON
test_coverage:
needs: build_instrumented_math_libs
strategy:
matrix:
project: [hiprand, rocrand, hipfft, rocfft]
uses: ./.github/workflows/test_component.yml
with:
component: ${{ matrix.project }}
coverage_enabled: trueReally all you seem to be doing with these changes is adding a bunch of conditionals so that we can filter everything except for compiler-runtime and math-libs, and eventually rocm-systems, so this might be a more ideal approach.
…verage-workflows-therock # Conflicts: # .github/workflows/multi_arch_build_portable_linux.yml
Review feedback asked for coverage to be an add-on that reuses the existing build process, rather than a feature the shared pipeline has to know about. multi_arch_build_portable_linux.yml is restored to its state on main, dropping the build_stages, stage_projects, and extra_cmake_options inputs along with the per-stage `if` rework that left the file with duplicate `if` keys and made it invalid YAML. The coverage workflows now call multi_arch_build_portable_linux_artifacts.yml once per stage they actually need, so no stage filtering is required anywhere. The nightly builds compiler-runtime and then math-libs as explicit jobs, which makes configure_coverage_ci.py's build_stages output unnecessary. The math-libs stage is deliberately not narrowed with stage_projects: its test jobs fetch artifacts by this run's id and so need the stage's full dependency set, and coverage_cmake_options already decides which projects get instrumented. COMPILER_RT_BUILD_PROFILE_ROCM is dropped from compiler/CMakeLists.txt, from the coverage workflows, and from the docs, since the ROCm profiling runtime is expected to come from the compiler build rather than from a TheRock level option. Co-authored-by: Cursor <cursoragent@cursor.com>
jainprad
commented
Sep 2, 2026
jainprad
left a comment
jainprad
left a comment
Author
There was a problem hiding this comment.
@reboss @endurthiabhilash , Thanks for the review.
I have addressed most of the the review comments. Please check
Implements the nightly half of RFC0014: option C for scheduling and option B for artifacts, on the phase 1 scope of every onboarded project on the single default architecture (gfx942 / gfx94X-dcgpu). The regular nightly now dispatches multi_arch_ci_coverage_nightly.yml once its build stage finishes, passing its own run id as baseline_run_id. The cron on the coverage workflow is gone: it would have produced runs with no baseline. Coverage keeps its own run id, so a coverage failure leaves the nightly's status alone and the nightly's unstable test jobs do not gate it, while the baseline it tests against is still fresh from the same commit. The dispatch is gated on the COVERAGE_NIGHTLY_ENABLED repository variable and on the nightly channel. Instrumenting the whole stack in one build and then reporting per project means a project's dependencies would otherwise be instrumented too, which both writes stray profiles and moves the coverage denominator whenever a dependency changes. Test jobs therefore install the baseline run in full and overlay only the project under test from the coverage run. overlay_coverage_artifacts.py does the overlay; it works at the subproject stage directory level rather than the artifact level because TheRock's artifacts are grouped (rand carries both rocRAND and hipRAND), and fails rather than silently reporting coverage for a non-instrumented build. The baseline is read from its own release channel, since artifacts are bucketed per channel. configure_coverage_ci.py grows the artifact_names and artifact_relpaths the overlay needs, and now collapses a whole-group selection to THEROCK_COVERAGE_ROCM_LIBRARIES_ALL / _ROCM_SYSTEMS_ALL / THEROCK_COVERAGE_ALL, which is how the RFC describes the nightly full-stack build. quartz_tracking_id is threaded through the coverage workflows, which the new dispatch edge pulls into the release graph.
Coverage now has a single entry point, the nightly workflow dispatched downstream from the regular nightly. The label-triggered PR workflow and the per-project instrumented rebuild it needed are removed, which also takes the stage_projects hook back out of the shared per-stage build workflow. multi_arch_ci_coverage_linux.yml keeps only the test, merge, and report half; the instrumented stack is always built by its caller. The registry fields that existed solely to drive the per-project rebuild (stage_project, and the coverage_flag / build_stage matrix keys) go with it. Co-authored-by: Cursor <cursoragent@cursor.com>
code_coverage.md covers the CMake options and onboarding a project, but not how a nightly report is actually produced. This adds the CI walkthrough: the dispatch and its gates, the instrumented build, how the baseline run id reaches the test jobs, the artifact overlay, and report generation. Linked from code_coverage.md, ci_overview.md, and the docs index. Co-authored-by: Cursor <cursoragent@cursor.com>
reboss
reviewed
Sep 3, 2026
The RFC's option C, where the regular nightly dispatches coverage and hands over its run id, lands separately. This drops the trigger_coverage_nightly job, so multi_arch_release_linux.yml is untouched by this PR and coverage is started by hand with a baseline_run_id from a recent nightly (the RFC's option B1). Nothing downstream changes: baseline_run_id and baseline_release_type are already inputs, so adding the trigger later only means adding the dispatching job. Co-authored-by: Cursor <cursoragent@cursor.com>
The phase 1 work exists to show that an instrumented build and a coverage report work end to end, and it currently cannot show that. multi_arch_ci_coverage_nightly.yml is triggered by hand, but a workflow_dispatch is only offered once the file is on the default branch, so there is no way to run it before it merges and no way to see it in this PR's checks. Merging it unvalidated to make it validatable is the wrong order. A caller is not subject to that restriction. Paths under ./.github/workflows resolve against the caller's own commit, so a pull_request-triggered caller runs this branch's copy of the coverage workflow and reports as jobs of this PR. The coverage workflow therefore gains a workflow_call trigger, duplicating the dispatch inputs since Actions does not expand YAML anchors. The baseline is pinned rather than resolved, so that a change in the report means a change in this branch: Multi-Arch CI run 33990288042, whose Stage - Compiler Runtime and Stage - Math Libs (gfx94X-dcgpu) both succeeded. That run's own test jobs failed, which does not matter here, as only its build stages are consumed. It published to therock-ci-artifacts rather than the nightly bucket, so baseline_release_type is ci and not the default; the bucket for a run follows the release type it published under. A ROCm/TheRock run rather than a rockrel release nightly. This workflow only installs the baseline at test time and never configures against it, so the embedded artifact workspace paths that rule out a cross-repository baseline elsewhere do not apply, but keeping the baseline in this repository is one fewer variable while the instrumented build is still unproven. The caller also has to grant id-token and actions:read itself, because a called workflow cannot hold permissions its caller does not. The paths filter is deliberately narrow and concurrency cancels in progress: each run is a multi-hour instrumented build plus a full test suite, and it is not worth spending one on code that has already been superseded. This file is scaffolding and is not part of the design. Coverage on presubmit is precisely what phase 1 rejected, and leaving this in place would charge that cost to every PR touching coverage. Delete it, and the workflow_call trigger, once the run has gone green. Co-authored-by: Cursor <cursoragent@cursor.com>
The nightly coverage run built an instrumented stack and then, at test time, threw most of it away: each job installed a regular nightly's artifacts in full and overlaid only the project under test, so the project's dependencies stayed non-instrumented. That is unnecessary. A coverage run builds every stage its selection needs and instruments only the projects the selection names, so everything those projects depend on is already built normally in the same run. Its own artifacts are the stack to test against. The isolation the overlay was protecting comes from somewhere else anyway. merge_coverage_report.py hands llvm-cov an explicit -object per binary, and llvm-cov only reports functions found in those objects' coverage mappings, so counters written by a sibling project are never looked up. A selection covering several projects still yields one report per project. Removing it also removes the failure mode it introduced. Run 34098007786 died in the overlay step, where shutil.copytree could not replace an existing libhiprand.so -> .so.1 -> .so.1.1 chain that the baseline install had already put in place. Drops the script and its tests, the baseline_run_id and baseline_release_type inputs threaded through four workflows, and the --artifact-names flag added to artifact_manager.py for the overlay's sake alone. Co-authored-by: Cursor <cursoragent@cursor.com>
THEROCK_COVERAGE_* set <PROJECT>_ENABLE_COVERAGE and therock_subproject.cmake passed that name straight through to the subproject. Almost no project implements it. Of the projects registered here only hipBLASLt, hipSPARSE and hipDNN use that spelling; the rest use BUILD_CODE_COVERAGE, CODE_COVERAGE or ENABLE_CODE_COVERAGE. hipRAND, the only project the registry had until now, gates its instrumentation on BUILD_CODE_COVERAGE. So -DHIPRAND_ENABLE_COVERAGE=ON was forwarded, ignored, and the build produced uninstrumented binaries. The tests would have written no profraw files and the report job would have failed on finding none. Run 34098007786 never got that far. Registry entries now name the option their project actually implements, and therock_subproject.cmake translates TheRock's flag into it. The translation is per subproject on purpose: BUILD_CODE_COVERAGE and CODE_COVERAGE are generic names, so defining either globally would instrument everything in the build that understands it, which is what per-project coverage exists to avoid. Also expands the registry to the 27 projects across rocm-libraries and rocm-systems, so a run can measure a chosen subset, and makes the nightly build only the stages that subset needs. Eight of those cannot be measured by this pipeline. hipSPARSE, rocALUTION and rocprofiler-sdk instrument with gcov, which writes .gcda files rather than the .profraw llvm-profdata merges; MIOpen, aqlprofile, amdsmi, rocprofiler-compute and rocprofiler-systems have no build-time coverage option at all. They stay registered so the gap is recorded, but carry an unsupported_reason instead of an option name, stay out of the group aliases, and are rejected with that reason when named rather than failing hours later with an empty report. Option names verified against the pinned rocm-libraries and rocm-systems commits. Co-authored-by: Cursor <cursoragent@cursor.com>
Phase 1 instruments the whole stack in one build but has to report per project, so a test job installs a non-instrumented baseline run and overlays only the project under test. That path was removed in 0585dc3 on the argument that object_globs already scope the report. They scope the export, but not what gets measured: inline and template code from headers is emitted linkonce_odr into every binary that uses it, so an instrumented sibling writes counters that llvm-profdata merges into this project's and no -object choice separates again. It is worst for the header-only projects (rocPRIM, hipCUB, rocThrust, rocWMMA), where nearly all the measured code lives in headers. This reverts that commit and restores baseline_run_id as the RFC specifies. The overlay crashed in run 34098007786 because shutil.copytree recreates symlinks with os.symlink(), which raises once the baseline install has already placed the libfoo.so -> libfoo.so.1 chain -- dirs_exist_ok forgives directory collisions, not symlink ones. Replaced with a merge_tree() that unlinks what it is about to replace, covered by two regression tests. The PR check now pins ROCm/rockrel 34068476531, a real Multi-Arch Release (nightly), and selects every project rather than one. Unlike PR #7944, which copies baseline stages into its own run and builds on top of them, nothing here configures CMake against the baseline, so the embedded /__w/rockrel/rockrel paths that rule out release nightlies there do not matter: the baseline is only installed on the test runner, where prebuilt binaries run and no package config is read. Selecting everything exposed two gaps a single-project run could not: - comm-libs takes hipify and rocjitsu from emulation, which the coverage workflow never built, so selecting rccl or rocshmem would have failed on a missing inbound artifact. build_stage_plan now derives the required stages transitively from the build topology and orders them into waves, and the workflow builds one wave per job. - The artifacts bucket lookup resolves the run through the GitHub API against this repository, which 404s for a rockrel id. Plumbed a RUN_GITHUB_REPO passthrough down to install_rocm_from_artifacts.py. Also added a test asserting every registry entry names the stage the topology says builds its artifacts. The registry duplicates that fact with nothing keeping the two in step, and a project moving upstream would otherwise surface as a stage never built and an overlay that finds nothing, hours into a run. Co-authored-by: Cursor <cursoragent@cursor.com>
Run 34119302776 selected every onboarded project and found the stack is
not ready to be instrumented all at once:
- libclang_rt.profile_rocm.a is linked without -ldl or -lpthread, so
an instrumented executable cannot resolve dlsym, dladdr or
pthread_once. Shared libraries tolerate the undefined symbols, which
is why this stayed hidden until projects that build their own
binaries were selected: rocRAND's tools, hipCUB's and rocPRIM's
tests.
- rccl's CMake refuses coverage outside a Debug build, and TheRock
builds Release.
- rocshmem does not survive coverage together with HIP relocatable
device code: "unknown special variable with appending linkage:
llvm.global_ctors".
Narrow to hipRAND, which built cleanly instrumented in 34098007786, so
the baseline install and the per-project overlay can be proven on their
own. The rocm-systems projects wait on the second and third problem, the
rest of math-libs on the first; the selection widens once they are
fixed. Nothing about the workflow changes -- it still resolves the whole
registry when asked, and the emulation/comm-libs wave jobs stay in place
for when rccl and rocshmem come back.
Co-authored-by: Cursor <cursoragent@cursor.com>
-fprofile-instr-generate links libclang_rt.profile_rocm.a, whose InstrProfilingPlatformROCm.cpp calls dlsym, dladdr and pthread_once. The driver puts neither -ldl nor -lpthread on the link line, and the manylinux base still keeps those symbols in libdl and libpthread rather than in libc, so an instrumented shared library ends up with undefined references. The library itself links, since shared objects tolerate that, and the failure surfaces one step later when an executable links against it: ld.lld: error: undefined reference: dlsym >>> referenced by library/libhiprand.so.1.1 >>> (disallowed by --no-allow-shlib-undefined) This blocked every instrumented project, not just the ones building their own executables, and is why run 34126836356 failed on hipRAND alone. It stayed hidden until 1fa0945 made the coverage flag actually reach the subproject; before that hipRAND was building uninstrumented. Add the two libraries to the subproject's linker flags alongside the coverage option. The flags go through the project init file, which appends to CMAKE_<TYPE>_LINKER_FLAGS and so lands ahead of the objects that need the symbols; push-state/--no-as-needed keeps the DT_NEEDED entries from being dropped on that account, and pop-state keeps the relaxation from applying to anything else on the line. Verified the mechanism outside ROCm, reproducing the shape of the failure: a static archive with an unresolved reference linked into a shared library, an executable linking against it under --no-allow-shlib-undefined, and the same init-file flags fixing it and recording the expected DT_NEEDED. Co-authored-by: Cursor <cursoragent@cursor.com>
The staged build jobs existed for one dependency edge: comm-libs takes hipify and rocjitsu from emulation, so selecting rccl or rocshmem had to build emulation first. Those two projects are the only consumers, and both are blocked upstream -- rccl's CMake refuses coverage outside a Debug build, and rocshmem fails on coverage plus HIP relocatable device code. The other 17 measurable projects are all in math-libs, so the wave jobs were carrying no traffic. Drop them, along with the topology walk and wave ordering in configure_coverage_ci.py. build_stage_plan() becomes resolve_build_stages(), which checks the selection against BUILDABLE_STAGES and reports whether math-libs is needed. Naming an unbuildable project is now an error at matrix time. Without it nothing would be built for that project and the failure would not surface until the test job found no instrumented files to overlay, hours into the run. An empty selection meant every measurable project, which would now reject itself, so it narrows to the buildable set. That applies to the default only; an explicit request is refused rather than reduced. The ordering constraint is recorded under "Adding a stage" in docs/development/code_coverage.md for whoever restores comm-libs. Co-authored-by: Cursor <cursoragent@cursor.com>
This was referenced Sep 29, 2026
…mespace Fixes the failing Linux CodeCoverage jobs in the rocm-libraries nightly (run 36913152797): - Instrumented subprojects no longer compile through ccache. ccache drops every -Xarch_<arch> <arg> pair, which removed the device-side negation (kernels were instrumented, and their colon-named profiles broke the profraw upload) and rocRAND's and hipRAND's -Xarch_host coverage flags (those libraries were not instrumented at all). - The instrumented stack publishes under <run_id>-coverage. The nightly's regular build shares the workflow run, and the two overwrote each other's math-libs artifacts. WorkflowOutputRoot looks a suffixed run id up by its numeric part. - external_repo_config is forwarded to the coverage builds, which otherwise measured TheRock's pinned submodules rather than the commit under test. - rocWMMA's coverage option is ROCWMMA_CODE_COVERAGE. rocPRIM, hipCUB and rocThrust report on their own test binaries; a "!" glob excludes a sibling project's. - Coverage test timeouts are 4x the release ones, capped at 180 minutes. - Device-side profiles are dropped, with a warning, before upload. - merge_coverage_report.py fails a report whose profiles hold no counters, whose objects carry no coverage mapping, or in which no line ran, instead of publishing 0%. Co-authored-by: Cursor <cursoragent@cursor.com>
1 task done
The rocm-libraries nightly starts coverage only after its whole regular Linux build and test job finishes, though only the coverage tests need the baseline, and only its build stages. A reusable workflow cannot `needs:` a job inside another one, so wait_for_baseline polls the run's jobs instead and holds coverage_report until the regular build's math-libs stage has succeeded for every coverage family (compiler-runtime when no math-libs project is selected). The gate runs only when baseline_run_id is this run's id; a baseline from an earlier run skips it. It fails, skipping the tests, when the stage job fails, is skipped, or never runs, when the run builds the stage only for other families, or after 300 minutes. Jobs are matched by display name, ignoring the Windows build, which reuses the names; a test checks the names against multi_arch_build_portable_linux.yml. Listing a public repository's jobs needs no `actions` permission (the nightly's Coverage Summary already does it with contents: read), so callers need no new grant. With the nightly as it is, the gate finds the baseline already built and passes on its first poll. Co-authored-by: Cursor <cursoragent@cursor.com>
…d links Instrumented tests in the rocm-libraries nightly (run 37288017642) abort at exit with "Cannot create GlobalVar Obj for symbol: __llvm_profile_sections_<cuid>". Keeping kernels uninstrumented does not stop clang from registering a host-side shadow of every HIP translation unit's device counters, or the driver from force-linking InstrProfilingPlatformROCm.cpp to take the registrations. At exit that file asks the HIP runtime for each shadow's device side, which only exists when the device code has profile data, and the HIP runtime aborts once the first module has written its profile. Every other module loses its profile; with an instrumented test executable, that includes the library under test. Coverage subprojects now compile cmake/therock_coverage_profile_stub.c at configure time and link it into every executable and shared library. It defines the file's five entry points as hidden no-ops, so the real one, with its HIP interceptors and HSA drain, is never pulled in. A compiler that adds an entry point the stub lacks fails the link on duplicate symbols. -fuse-cuid=none would stop the shadows instead, but every translation unit then defines the same __hip_cuid_, and libraries with two HIP sources fail to link. Co-authored-by: Cursor <cursoragent@cursor.com>
…ed libraries kpack moves the program header table of each library it rewrites to the end of the file, mapped by a PT_LOAD of its own at an address only page-congruent with the file offset. compiler-rt's profile writer reads the table at __ehdr_start + e_phoff to record binary IDs, so instrumented rocBLAS and rocSPARSE segfault at exit before writing a counter, and rocRAND reads another segment's bytes instead. kpack already pins p_vaddr == p_offset for executables (normalize_phdr_vaddr). Until it does the same for shared libraries, install_rocm_code_coverage_build.py applies that edit to each replaced file with a __llvm_prf_cnts section once it is extracted (_therock_utils/elf_phdr.py). Uninstrumented libraries and the baseline install are left as they are. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
This PR has been split into:
Could you please review those PRs? |
Member
|
I can perform second level reviews once they are reviewed by others working more directly on the features. Please find enough reviewers/maintainers for these features, we're over capacity for what core project admins can review. |
Contributor
|
Sure, I have done a first pass on #8598. I will find few other reviewers for the PRs. |
…rchives The coverage math-libs build in rocm-libraries run 37441751375 fails to link librocsparse.so on duplicate symbols between the stub and InstrProfilingPlatformROCm.cpp, the GPU half of the profile runtime it stands in for. rocSPARSE names libclang_rt.profile.a and libclang_rt.profile_rocm.a on its own link line, ahead of the HIP runtime. The GPU half also defines weak hipLaunchKernel, hipModuleLoad and similar wrappers, so lld extracts it from libclang_rt.profile.a for the hipLaunchKernel references rocSPARSE's objects still have at that point. The generated coverage include now puts the toolchain's libamdhip64.so right after the stub, under --as-needed, so those names already resolve to the HIP runtime when any archive is scanned. Links that don't call into HIP gain no dependency on it, and subprojects without THEROCK_TOOLCHAIN_ROOT get the stub alone, as before. Co-authored-by: Cursor <cursoragent@cursor.com>
rocThrust's coverage report in rocm-libraries run 37803818804 fails with none of its 67275 instrumented lines run. Its test job installed the baseline and replaced only files whose path names the project, which for rocThrust's tests left 2 of 372: bin/rocthrust/generate_resource_spec and CTestTestfile.cmake. ctest then ran the baseline's uninstrumented bin/*.hip, and the one profraw uploaded came from generate_resource_spec. rocPRIM's tests (bin/test_block_scan) and rocWMMA's (bin/gemm_*-validate) are missed the same way; their reports pass on the 1.72% and 2.20% that their few swapped binaries happen to share with them. A file under the project's own stage directory in the artifact manifest, math-libs/rocThrust/stage or math-libs/rocPRIM_tests/stage, is now replaced whatever its name. Paths elsewhere still have to name the project, and so does the device code under .kpack/, which holds every project's kernels in a multi-project artifact like prim; rocWMMA's, named after it, is replaced as before. Co-authored-by: Cursor <cursoragent@cursor.com>
rocWMMA's report job died in every attempt of rocm-libraries run 37912746375: the hosted runner got a shutdown signal two to three minutes into llvm-cov show, after the lcov and the summary were written. llvm-cov's HTML repeats a header's source once per template instantiation, so rocWMMA's report came to 3.3GB with single pages near 1GB, and rendering it peaked at 17.9GB on four cores, past the 16GB of an ubuntu-24.04 runner. Most of that is the coverage mappings of the 29 test binaries, which the lcov export holds too (14.6GB); demangling every function name up front added 3GB, and rendering files in parallel another 2GB. Show each file once, its counts summed over instantiations as in the lcov. No function name is printed then, so drop the demangler, and render one file at a time. On the same inputs the report is 19MB, peaks at 14.1GB, and its totals are unchanged. Co-authored-by: Cursor <cursoragent@cursor.com>
jainprad
added a commit
that referenced
this pull request
Oct 9, 2026
From runs 37756059801 and 37756059792: - rpp: librpp.so failed its load check on an undefined __sanitizer_internal_memcpy, which skipped every phase 2 test. rpp links HIP only as hip::host, so the shadow registrations of its HIP translation units pull the GPU half of the profile runtime out of libclang_rt.profile.a, and only libclang_rt.profile_rocm.a defines the sanitizer helpers that half calls. The library also exported the hipLaunchKernel interceptors. Projects instrumented for host coverage only now link #7783's stub of the GPU half's entry points, and every coverage link puts the HIP runtime ahead of its objects, as #7783 does. A CMake test configures a host-only and a device subproject and checks their link flags. - rocALUTION: no instrumented device code. Its HIP backend is built with FindHIP's hip_add_library(), whose hipcc commands never see the compile rule. cmake/coverage/rocALUTION.cmake seeds HIP_CLANG_FLAGS with the flags the rule gets, now also set as THEROCK_COVERAGE_COMPILE_FLAGS. A test checks they survive FindHIP creating the cache entry under rocALUTION's policy settings. - hip-kernel-provider: the coverage test job had no uv for install_additional_requirements.py. It now installs uv and points it at the setup-python interpreter. - rocThrust: timed out after 125 of 167 tests. Coverage runs can now set coverage_total_shards_dict; rocThrust uses 3 shards, split by gtest case alone so that every shard still runs each test binary. - rocPRIM's test_config_dispatch and rocThrust's scan_by_key.inclusive and reproducibility fail only because of the instrumentation, and are skipped in coverage runs (COVERAGE_CTEST_EXCLUSIONS in test_runner.py). test_config_dispatch launches on hipStreamLegacy, which the profile runtime's hipLaunchKernel interceptor passes to hipStreamGetDevice, and that dereferences it. The scan-by-key binaries return wrong results with instrumented kernels and pass with the same host binaries and the regular build's kernels. Co-authored-by: Cursor <cursoragent@cursor.com>
ScottTodd
removed their request for review
October 9, 2026 19:36
1 task done
jainprad
added a commit
that referenced
this pull request
Oct 19, 2026
…mespace Fixes the failing Linux CodeCoverage jobs in the rocm-libraries nightly (run 36913152797): - Instrumented subprojects no longer compile through ccache. ccache drops every -Xarch_<arch> <arg> pair, which removed the device-side negation (kernels were instrumented, and their colon-named profiles broke the profraw upload) and rocRAND's and hipRAND's -Xarch_host coverage flags (those libraries were not instrumented at all). - The instrumented stack publishes under <run_id>-coverage. The nightly's regular build shares the workflow run, and the two overwrote each other's math-libs artifacts. WorkflowOutputRoot looks a suffixed run id up by its numeric part. - external_repo_config is forwarded to the coverage builds, which otherwise measured TheRock's pinned submodules rather than the commit under test. - rocWMMA's coverage option is ROCWMMA_CODE_COVERAGE. rocPRIM, hipCUB and rocThrust report on their own test binaries; a "!" glob excludes a sibling project's. - Coverage test timeouts are 4x the release ones, capped at 180 minutes. - Device-side profiles are dropped, with a warning, before upload. - merge_coverage_report.py fails a report whose profiles hold no counters, whose objects carry no coverage mapping, or in which no line ran, instead of publishing 0%. Ported from #7783 onto the split #8598/#8602/#8605 stack: the registry entries keep #8598's derived stage/test_component/coverage_config, the report's run id goes through env like the rest of the step, the external_repo_config input this PR had already added is merged with this one, and copy_compiler_runtime writes into the same <run_id>-coverage namespace. (cherry picked from commit 8523c5a) Co-authored-by: Cursor <cursoragent@cursor.com>
## Motivation hipCUB coverage tests in rocm-libraries run 41136053832 (gfx94X-dcgpu) failed: 33 of 49 tests failed, 32 of them aborting with `Cannot find Symbol` for `hipcub::HIPCUB_400700_NS` kernels, and `GetVersionOnDevice` read 400701 instead of 400700. Since 80e80ed, every file under the project's own stage directories is replaced, except the device code under `.kpack/` unless its name names the project. prim_test's kpack is `prim_test_gfx942.kpack`, shared by hipCUB, rocPRIM and rocThrust, so it was never replaced. The instrumented test binaries then looked up their kernels by name in the baseline's kpack, which was built at a different hipCUB version. ## Technical Details - `install_rocm_code_coverage_build.py`: replace the archive's `.kpack/` device code whenever the component has its own stage directory in the artifact's manifest (`owns_archive`), so the host binaries and their kernels come from the same build. Archives where the component has no stage directory still keep the baseline's kpack. Split the regexes into `_names_folder` and `_is_own_stage`. - Trade-off: other projects' baseline binaries in the same artifact (e.g. rocPRIM and rocThrust in prim_test) now load the instrumented kpack. A coverage job only runs the swapped project's tests. - Tests: the prim_test case now expects the shared kpack to be replaced for each project; added a case where the component has no stage directory in the artifact (rocBLAS vs a hipBLASLt-only blas_lib), which keeps the baseline kpack. - `docs/development/code_coverage.md`: describe the new `.kpack/` rule and the trade-off.
…verage-workflows-therock # Conflicts: # .github/workflows/test_code_coverage_component.yml
## Motivation The coverage test job installs the baseline stack and overlays only the selected project's instrumented files. hip* wrappers call a backend at runtime that ships in the same artifact, so the instrumented wrapper ran against the baseline's backend, which can be a different version than the one it was built against. hipRAND's linkage test failed this way: hiprandGetVersion() (answered by the baseline librocrand, 500300) did not match the ROCRAND_VERSION it was compiled with (500100). ## Technical Details - Add SAME_ARTIFACT_BACKENDS to install_rocm_code_coverage_build.py: hipblas->rocBLAS, hipfft->rocFFT, hiprand->rocRAND, hipsolver->rocSOLVER, hipsparse->rocSPARSE. --replace-<wrapper> now also extracts the backend's folder from the same instrumented artifact. - Move folder collection out of main() into replacement_folders(), which de-duplicates folders when wrapper and backend are both selected. - Add tests for the mapping, for backends living in their wrapper's artifact, and for a rand archive swap installing both libhiprand and librocrand. - Document the behavior in docs/development/code_coverage.md. The backend's objects stay out of the wrapper's --object-globs, but header templates both instantiate share profile counters. Dependencies in other artifacts (fftw3, rocBLAS for hipSOLVER) still come from the baseline.
Comment on lines
+157
to
+165
| - name: Generate coverage matrix | ||
| id: matrix | ||
| env: | ||
| PROJECTS_TO_TEST: ${{ inputs.projects_to_test || '' }} | ||
| AMDGPU_FAMILIES: ${{ inputs.amdgpu_families || 'gfx94X-dcgpu' }} | ||
| COVERAGE_CONFIG_SOURCE: ${{ inputs.coverage_config_source || 'ROCm/rocm-libraries@main' }} | ||
| run: python ./build_tools/github_actions/configure_coverage_ci.py | ||
|
|
||
| - name: Compute package version |
Comment on lines
+165
to
+174
| - name: Compute package version | ||
| id: rocm_package_version | ||
| run: | | ||
| python ./build_tools/compute_rocm_package_version.py \ | ||
| --release-type=ci \ | ||
| --override-git-sha=${{ steps.checkout.outputs.commit }} | ||
|
|
||
| # Built first so later stages have a compiler-runtime to link against and | ||
| # fetch as an inbound artifact. | ||
| build_compiler_runtime: |
Comment on lines
+258
to
+271
| - name: Wait for the baseline stage | ||
| env: | ||
| GITHUB_TOKEN: ${{ github.token }} | ||
| BASELINE_STAGE: ${{ needs.setup_coverage_matrix.outputs.needs_math_libs == 'true' && 'math-libs' || 'compiler-runtime' }} | ||
| AMDGPU_FAMILIES: ${{ needs.setup_coverage_matrix.outputs.dist_amdgpu_families }} | ||
| run: | | ||
| python build_tools/github_actions/wait_for_baseline_stage.py \ | ||
| --stage="${BASELINE_STAGE}" \ | ||
| --amdgpu-families="${AMDGPU_FAMILIES}" | ||
|
|
||
| # One test + report job per project. Each installs baseline_run_id and replaces | ||
| # only its own project from this run's instrumented stack, so dependencies stay | ||
| # non-instrumented even though the build instrumented them all at once. | ||
| coverage_report: |
Comment on lines
+333
to
+338
| - name: Evaluate workflow results | ||
| env: | ||
| GITHUB_TOKEN: ${{ github.token }} | ||
| run: | | ||
| python build_tools/github_actions/workflow_summary.py \ | ||
| --needs-json '${{ toJSON(needs) }}' |
Comment on lines
+124
to
+140
| - name: Configuring coverage test options | ||
| id: configure | ||
| env: | ||
| ARTIFACT_GROUP: multi-arch-coverage | ||
| AMDGPU_FAMILIES: ${{ inputs.amdgpu_families }} | ||
| TEST_TYPE: ${{ inputs.test_type }} | ||
| TEST_LABELS: "[]" | ||
| # One project per report, so only its tests run here; other projects | ||
| # get their own job. | ||
| PROJECTS_TO_TEST: ${{ inputs.test_component }} | ||
| CI_CONFIG_PATH: ci-config | ||
| BUILD_VARIANT: coverage | ||
| run: | | ||
| python ./build_tools/github_actions/fetch_test_configurations.py \ | ||
| --platform=linux | ||
|
|
||
| test_coverage: |
Comment on lines
+326
to
+333
| - name: Checking out repository | ||
| uses: actions/checkout@3d3c42e5aac5ba805825da76410c181273ba90b1 # v7.0.1 | ||
| timeout-minutes: 15 | ||
| with: | ||
| # Cross-repo callers pass ROCm/TheRock + ref; same-repo defaults to github.repository. | ||
| repository: ${{ inputs.repository || github.repository }} | ||
| ref: ${{ inputs.ref }} | ||
| - name: Evaluate workflow results |
This branch has not been deployed
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.
Motivation
GitHub issue: #6572
Adds an end-to-end code coverage pipeline to TheRock: propagate per-project -D_ENABLE_COVERAGE flags from the superbuild down to subprojects, and build/test/report coverage in CI.
Technical Details
Test Plan
Test Result
Submission Checklist