Repository navigation
Conversation
Co-authored-by: Vivek Kale <vivek.kale@amd.com> (cherry picked from commit c61850d)
…apshot test kernels Co-authored-by: Vivek Kale <vivek.kale@amd.com>
github-actions
Bot
added
documentation
❌ PR Check — Action Required
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🚫 Please fix the failed policies before requesting reviews. The following policy checks failed:
The |
This was referenced Oct 1, 2026
Draft
amd-vkale
marked this pull request as ready for review
October 5, 2026 05:42
Contributor
There was a problem hiding this comment.
Copilot review overview
🟡 Changes recommended
The repository’s ROCr reports the constant attribute with inverted semantics, causing the filter to skip writable globals while retaining non-writable symbols.
Review effort: Balanced
Findings: 1
What changed in this PR
Updates kernel replay snapshots to exclude constant symbols, aiming to avoid restoring data to protected GPU pages.
Changes:
- Adds filtering based on the HSA constant-symbol attribute.
- Adds a constant-symbol fixture and exclusion regression test.
- Updates snapshot documentation and API guidance.
| File | Description |
|---|---|
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/tests/snap_restore.cpp | Tests constant exclusion and restore. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/tests/snap_kernels.hpp | Declares the constant-address helper. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/tests/snap_kernels.cpp | Adds the constant fixture and helper. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/memory_snapshot.hpp | Documents writable-only module capture. |
| projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_replay/memory_snapshot.cpp | Filters symbols using constant metadata. |
| projects/rocprofiler-sdk/source/include/rocprofiler-sdk/experimental/kernel_replay.h | Updates public snapshot guidance. |
| projects/rocprofiler-sdk/source/docs/how-to/using-rocprofv3.rst | Revises CLI snapshot limitations. |
| projects/rocprofiler-sdk/source/docs/how-to/using-kernel-replay.rst | Clarifies restored memory categories. |
| projects/rocprofiler-sdk/source/docs/how-to/using-kernel-replay-rocprofv3.rst | Documents constant exclusion. |
| projects/rocprofiler-sdk/source/docs/conceptual/kernel_replay/kernel_replay_memory_snapshot.md | Explains symbol filtering. |
| projects/rocprofiler-sdk/CHANGELOG.md | Updates snapshot coverage notes. |
💡 Add a code-review agent skill or configure MCP servers for context-aware, tailored reviews. Learn more in the docs.
Revert the HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_IS_CONST filter from 02c1fdd, and the test-kernel alignment from 08b2ce0 that came with it. ROCr reports that attribute inverted: Symbol::IsConst() returns true when the ELF section has SHF_WRITE. HIP places __device__ globals in .data/.bss and __constant__ ones in .rodata, so the filter dropped the writable globals that replay must restore and kept the constants. CI failed the four module-variable snapshot unit tests and both kernel replay generate tests in every job. Restoring constants is harmless: on MI325 the exclusion test's restore() succeeded with the __constant__ fixture still in the snapshot. Leave a comment so the attribute is not used for this again. Co-authored-by: Vivek Kale <vivek.kale@amd.com>
cursor
Bot
changed the title
Contributor
Author
|
@copilot Fix the code for all comments in this review thread. When a review comment includes a suggested change, apply the suggestion exactly. Do not make changes beyond what is described in the linked review thread. |
Co-authored-by: amd-vkale <291155109+amd-vkale@users.noreply.github.com>
Contributor
Updated module-variable filtering to use HSA_VARIABLE_SEGMENT_READONLY instead of VARIABLE_IS_CONST, and fail closed if segment data is unavailable or unrecognized. Commit: |
|
|
||
| if(kind != HSA_SYMBOL_KIND_VARIABLE) return HSA_STATUS_SUCCESS; | ||
|
|
||
| // __constant__ variables are captured too; restoring them is harmless. Do not skip them via |
cursor Bot
pushed a commit
that referenced
this pull request
Oct 6, 2026
…constant comment Where __constant__ variables land depends on the compiler. Ubuntu clang 18 puts a __constant__ int in .data, but the amdclang++ 23 that CI builds with puts it in a non-writable section: on #12605 the IS_CONST filter left the excludes_read_only_module_constant fixture in the snapshot, which only happens when its section lacks SHF_WRITE. Neither placement sets SHF_AMDGPU_HSA_READONLY, so the comment's reasons for not filtering on IS_CONST or VARIABLE_SEGMENT hold either way. State those reasons without the placement claim. Co-authored-by: Vivek Kale <vivek.kale@amd.com>
This was referenced Oct 6, 2026
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.
First PR of the kernel replay host-side perf stack: leave constant symbols out of the snapshot, since kernels cannot write them and restoring them can fault read-only pages. JIRA AIPROFSDK-68.