[stinkytofu] Expand RawAsmParser coverage for custom operands, MX modifiers, and DPP#7097
Open
darrenhsieh-amd wants to merge 3 commits into
Open
[stinkytofu] Expand RawAsmParser coverage for custom operands, MX modifiers, and DPP#7097darrenhsieh-amd wants to merge 3 commits into
darrenhsieh-amd wants to merge 3 commits into
Conversation
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
shared: stinkytofu
labels
…tions Instructions with custom operand syntax (s_delay_alu, s_wait_alu, s_setreg_IMM32_b32) were falling to TEXTBLOCK because the operand parser used a generic while-loop that called parseOneRegister on every operand, regardless of FieldType. Refactor the operand parsing from a while-loop to a per-field for-loop over operandFields, dispatching each field by FieldType via parseCustomOperand(): - FieldType::delay -> parseDelayAluSyntax() - FieldType::wait_alu -> parseWaitAluSyntax() - FieldType::hwreg -> parseHwregOperand() - register/immediate -> parseOneRegister with existing recovery logic Custom operand parsers are called directly in the operand loop, consistent with their role as operands rather than trailing modifiers. Uses field.isDest directly from operandFields instead of pre-counted numDest. This also enables mixed-type instructions like s_setreg_IMM32_b32 (hwreg + simm32) where custom and register operands coexist.
v_wmma_scale_* instructions with matrix_a_fmt/matrix_b_fmt modifiers were falling to TEXTBLOCK because their microcode format (MC_VOP3PX2/ MC_VOP3PX3) had no modKey mapping, and key:Identifier modifier values (e.g. matrix_a_fmt:MATRIX_FMT_FP8) were unconditionally marked as sawUnrepresentable. Add MC_VOP3PX2/MC_VOP3PX3 to the parseModifiers microcode switch, mapping to "mod.mfma" (same modifier path as the rocisa pipeline). Store key:Identifier modifier values in the generic fields collection instead of marking unrepresentable — formats without a modifier namespace are caught by the existing modKey.empty() check. Add mod.mfma handler for matrix_a_fmt, matrix_b_fmt, matrix_a_reuse, matrix_b_reuse. Build inputPermute from individual fields in ModifierSerializer::deserialize, consistent with the rocisa pipeline path (ToStinkyTofuUtils::extractMatrixFormatStr).
VOP instructions with DPP modifiers (row_xmask, row_shl, row_ror, etc.) were falling to TEXTBLOCK because VOP microcode formats have no modKey mapping. DPP is an add-on encoding variant layered on top of the base VOP format — the assembler picks a wider encoding when DPP modifiers are present. Add field-based DPP detection after modifier token collection via hasDPPFields() helper. When modKey is empty and DPP fields are present, assign modKey = "mod.dpp" post-hoc. This avoids mapping all VOP formats to DPP which would interfere with non-DPP VOP modifiers. Add dppCtrl string field to DPPModifiers for unmodelled DPP modes (row_xmask, row_shl, etc.). Existing typed fields (row_shr, row_bcast, bound_ctrl) preserved for rocisa pipeline compatibility. Add mod.dpp deserialize support in ModifierSerializer.
darrenhsieh-amd
force-pushed
the
users/darrenh/rawasmparser-coverage
branch
from
5be5e27 to
fbb7fa2
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 #7097 +/- ##
========================================
Coverage 65.26% 65.26%
========================================
Files 2083 2083
Lines 323506 323506
Branches 42438 42438
========================================
Hits 211111 211111
Misses 94821 94821
Partials 17574 17574
*This pull request uses carry forward flags. Click here to find out more. 🚀 New features to boost your workflow:
|
Contributor
Author
|
testing asm for the stinkytofu-opt round-trip: |
darrenhsieh-amd
commented
May 6, 2026
darrenhsieh-amd
commented
May 6, 2026
cycheng
approved these changes
May 7, 2026
| if (fields.count("glc")) modFields["glc"] = "true"; | ||
| if (fields.count("nv")) modFields["nv"] = "true"; | ||
|
|
||
| } else if (modKey == "mod.dpp") { |
Contributor
There was a problem hiding this comment.
I feel there is a future TODO:
We now have 4 conversions for modifier:
- StinkyModifier to assembly string (asm emitter)
- StinkyModifier to stinkytofu IR string (modifier serializer)
- assembly string to StinkyModifier (raw asm parser)
- stinkytofu IR string to StinkyModifier (modifier de-serializer)
4 for instruction:
- StinkyInstruction to assembly string (asm emitter)
- StinkyInstruction to stinkytofu IR string (IRConverter)
- assembly string to StinkyInstruction (raw asm parser)
- stinkytofu IR string to StinkyInstruction (IRConverter)
Maybe we could define a way that can handle serializer/de-serializer in a unified way.
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
3 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
Instructions with custom operand syntax, MX matrix format modifiers, or DPP encodings silently fall through to TEXTBLOCK in RawAsmParser. As TEXTBLOCK, they are opaque to all optimization passes. This PR fixes the parser to recognize these instruction patterns so they can be properly analyzed and transformed.
Technical Details
Per-field operand dispatch for custom syntax instructions
Wire MX matrix format modifiers
Add DPP modifier support
Test Plan
ctest pass
filechecks pass
stinkytofu-opt --arch gfx1250 --emit-asm round-trip correct
Submission Checklist