[stinkytofu] Make hwreg() a typed StinkyRegister::HwReg operand variant#7729
Open
darrenhsieh-amd wants to merge 9 commits into
Open
[stinkytofu] Make hwreg() a typed StinkyRegister::HwReg operand variant#7729darrenhsieh-amd wants to merge 9 commits into
darrenhsieh-amd wants to merge 9 commits into
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
shared: stinkytofu
labels
darrenhsieh-amd
force-pushed
the
users/darrenh/stinkytofu-hwreg-tablegen-v2
branch
from
f0ecf57 to
39d34f6
Compare
Contributor
Author
darrenhsieh-amd
changed the title
darrenhsieh-amd
changed the title
darrenhsieh-amd
changed the title
darrenhsieh-amd
force-pushed
the
users/darrenh/stinkytofu-hwreg-tablegen-v2
branch
from
39d34f6 to
f263a8f
Compare
…tions - Operand loop walks operandFields and dispatches by FieldType. - s_delay_alu / s_wait_alu / s_setreg_IMM32_b32 etc. no longer fall to TEXTBLOCK. - Enables mixed-type instructions (hwreg + simm32 in s_setreg_IMM32_b32).
- Map MC_VOP3PX2/MC_VOP3PX3 → mod.matrix_fmt (typed MatrixFmtModifiers). - Store key:Identifier modifier values in fields instead of marking unrepresentable. - matrix_*_reuse cross-attaches to mod.mfma so one asm line yields both modifier types.
- inferModKeyFromFields detects dpp_ctrl/encoding-side keys → mod.dpp. - Reconstruct asm-form token (e.g. "row_xmask:8") and feed parseDppCtrlFromAsm. - DPP layers on many VOP formats, so MC-format alone can't pre-assign it.
Modifiers like `offset:8704*2` previously fell to TEXTBLOCK because evalExpr only knew `+`/`-`; passes saw no typed offset. Now evaluated to `17408`. Round-trip is no longer byte-identical for these cases (emitted in folded form) — passes need the int, and storing both original text and value would require dual-tracking through every modifier.
- Add entry to availablePasses. - Mark createCFGBuilderPass STINKYTOFU_EXPORT so the tool can link it.
- `--debug-pass FooPass` was treating "debug-pass" as a pass name. - Skip the value argument like `-o`'s two-arg form.
- Replace unordered_map with order-preserving wrapper (same operator[]/iter API). - Outer-key insertion order = canonical assembler trailing-modifier order. - Inner field map stays unordered (each serializeVisit emits in fixed order).
…enerator - DEF_HWREG(NAME, ID) blocks in each arch's Formats.def feed a constexpr NamedId table. - Consumer (HwReg.cpp) lands in the next commit.
- hardware/HwReg.{hpp,cpp}: Id enum, name<->id, SCHED_MODE sub-field accessors
- StinkyRegister::HwReg variant with {id, offset, size}
- HwRegHelpers.hpp: parseId, printOperand, isSetregTo
- RawAsmParser, IRParser, AsmEmitter, AsmPrinter, rocisa bridge: use the variant
darrenhsieh-amd
force-pushed
the
users/darrenh/stinkytofu-hwreg-tablegen-v2
branch
from
f263a8f to
7090df9
Compare
Codecov Report✅ All modified and coverable lines are covered by tests. ❌ Your project status has failed because the head coverage (77.83%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage. Additional details and impacted files@@ Coverage Diff @@
## develop #7729 +/- ##
===========================================
- Coverage 61.94% 61.87% -0.07%
===========================================
Files 2086 2086
Lines 357068 357219 +151
Branches 53779 53842 +63
===========================================
- Hits 221171 221013 -158
- Misses 117093 117410 +317
+ Partials 18804 18796 -8
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
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
2 participants
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
Insts s_setreg / s_getreg carry their hwreg(id, offset, size) operand as a plain LiteralString register, and that is opaque to
every pass. Any analysis that wants to reason about HWREG writes (e.g. detecting writers of a specific sub-field like
SCHED_MODE.DEP_MODE) has to pattern-match free-form text. This PR replaces the text blob with a typed StinkyRegister::HwReg operand variant and adds arch-aware name lookups so passes can interrogate the operand structurally.
This PR is on top of #7097
Technical Details
This branch stacks on top of
users/darrenh/rawasmparser-coverage(which lands the per-field operand dispatch the parser dispatch relies on).Test Plan
tests/filecheck/hwreg_operand_roundtrip.s(parse -> emit produces canonicalhwreg(...)form for each variant).test_gen_instructionsasserts the per-arch<Arch>_hwreg.incfile is emitted with the expected array + a known entry.ctest.Test Result
Test Plan all passed.
Submission Checklist