rocjitsu: Harden ISA layer, virtual machine, and adds initial DBT by atgutier · Pull Request #5512 · ROCm/rocm-systems

atgutier · 2026-04-27T21:39:09Z

Motivation

Provides many fixes and enhancements to establish our MVP:

Opt125 runs a forward pass on the simulator in the vllm-openai-rocm container for the cdna4 KMD config and produces the correct token:

Final logits: NaN=0 Inf=0 total=50272
GPU top-5: [(' the', '9.03'), (' is', '7.57'), (' a', '6.89'), (' has', '6.19'), (' "', '5.50')]

Loading model on CPU for comparison...
CPU top-5: [(' the', '7.66'), (' a', '7.50'), (' not', '6.82'), (' in', '6.61'), (' being', '6.24')]
Logit max_diff=9.1250 mean_diff=1.3923

DBT layer demonstrates capabilities to translate between CDNA4 native source binary and RDNA4 Wave64 mode target with correct results against a fuzzing test of random inputs.

Technical Details

Many fixes to the ISA layer, pipeline, and KMD. Initial architecture in place for DBT.

JIRA ID

Test Plan

Run all tests in our suite and validate Opt125 run in the simulator on the vllm-openai-rocm container.

Test Result

Tests all pass. Real GFX12 HW runs of translated binaries pass (requires GFX12 HW), Opt125 forward pass produces good token in Docker when run on CDNA4 KMD config.

Submission Checklist

Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

Copilot

Copilot wasn't able to review this pull request because it exceeds the maximum number of files (300). Try reducing the number of changed files and requesting a review from Copilot again.

MaheshRavishankar

Could you please clean up the commit title...

This is experimental and reviewed through agents. I am fine with this landing and fixing forward any issues since this is MVP.

Groverkss · 2026-04-28T13:51:04Z

I left some comments on some potential bugs that I fixed locally, feel free to land without them and I can send a follow up patch for them.

accauble

Not done reviewing but since you're making changes, you may want to address this sooner rather than later.

…ning - Fix MUBUF address calculation (buffer bounds checking, OOB modes, stride) - Fix v_cmpx to write VCC on CDNA (CDNA writes VCC+EXEC, RDNA only EXEC) - Fix v_div_fmas scaling exponents per ISA spec (F32 and F64) - Rewrite v_div_scale from ISA pseudocode for correct special-case handling - Add MTBUF idxen support with full buffer addressing - Throw exception on unsupported GDS bit - Add SDWA float clamp bit support for VOP1/VOP2 - Fix MFMA F64 const_acc and output_loc_64 register convention - Implement FP8 MFMA X2 scaling and remove mfma namespace wrapper - Add rocBLAS GEMM test suite (2x2 through 2048x2048) - Fix codegen build_modifiers to skip fields not in ISA (lds on rdna3/rdna3_5) - Fix codegen to emit log/format includes for saveexec instructions

…ked ISA - Enable SDMA queue processing and fix 2D/3D workgroup dispatch - Fix LDS allocation collision and L2 write-through behavior - Fix FP16 ISA correctness (v_cvt, packed ops) for vLLM kernel support - Fix dot product, packed F32, and SDWA instruction implementations - Implement v_pack_b32_f16 and fix cross-ISA sharing analysis - Synchronous memory pipeline, ABI fixes, and allocation zeroing - PackedTID workitem ID encoding for CDNA3/4 - CU count fix and code object metadata extraction - SRAM ECC D16 zeroing behavior - Add tmp/ to gitignore

… and legalization - Implement v_bitop3_b32 (3-input boolean) and v_permlane_swap instructions - Add amdisa legalization extension for cross-ISA instruction mapping (identity, substitute, lower, expand, illegal) - Fix ds_read_b64_tr_b16 transpose to use word-level stride-4 pairing - Misc instruction hardening and edge-case fixes

- DRM interposer fixes and vLLM platform plugin for KFD emulation - Cross-ISA encoding translators and legalization tables (28 ISA pairs, auto-generated) - BinaryTranslator, CodeObjectPatcher, and rj_code_translate public C API - Two-tier DBT architecture: semantic legalization + encoding field remapping - CDNA4→RDNA4 binary translation pipeline with hardware dispatch test - E2E test: translate gfx950 vector_add → gfx1201, run on real RDNA4 GPU

…up, buffer_wbl2 - Add LDS capacity check to workgroup scheduling (prevents OOM: 4 WGs × 42.5KB > 160KB) - Add operand delegate pattern for DPP/SDWA transparency (~1095 shared execute sites) - Fix packed F32 inline constant interpretation (F32 replicated, not F64 encoding) - Fix ds_write2/ds_read2 missing lds_base offset (per-workgroup LDS addressing) - Implement buffer_wbl2 and buffer_inv cache control instructions (Triton flash attention) - Remove instruction-count safety valve and add SIGSEGV handler - Fix v_div_fixup and v_div_scale special-case handling (zero/inf numerator/denominator) - MFMA accumulation unit tests These fixes enable OPT-125M to produce correct next-token predictions on the rocjitsu simulator with eager attention.

Add per-basic-block VGPR liveness analysis in a new top-level analysis/ directory. The RegisterLiveness class provides per-instruction live set queries and free register run search, used by the semantic translator for safe register remapping. Register classification uses auto-generated is_vgpr_operand_type() functions (one per ISA) that classify operands by their OperandType enum, avoiding hardcoded encoding value ranges. The Operand base class exposes is_vgpr(), unified_vgpr_index(), and vgpr_count() accessors set at construction time by ISA-specific subclasses. Also adds AccVGPR elimination to the semantic translator: v_accvgpr_read and v_accvgpr_write are lowered to NOPs or v_mov_b32 on RDNA4 (where AccVGPRs map to the unified VGPR file). The KD patcher now accounts for AccVGPR allocation from compute_pgm_rsrc3 ACCUM_OFFSET when computing the RDNA4 VGPR granularity. Signed-off-by: Tony Gutierrez <anthony.gutierrez@amd.com>

Translates v_mfma_f32_16x16x16_f16 (CDNA4) to v_wmma_f32_16x16x16_f16 (RDNA4) in Wave64 mode with correct lane layout remapping. WMMA swaps output lane groups 1 and 2 vs MFMA; ds_bpermute_b32 with XOR-48 byte addresses at lanes 16-47 corrects this in a single post-WMMA pass. Translation pipeline: - v_mbcnt_lo/hi computes lane_id (uses HW thread mask, not EXEC) - v_lshlrev_b32 converts to byte address, v_xor_b32 applies XOR-48 at lanes 16-47 via EXEC mask - v_wmma_f32_16x16x16_f16 executes (single pass, all 64 lanes) - s_wait_idle drains pipelines for DS access to WMMA output - ds_bpermute_b32 x4 remaps lanes, s_wait_dscnt 0 completes Address VGPR selected by RegisterLiveness::find_free_run() to avoid clobbering live kernel state. GFX12 hazards handled via s_delay_alu (SALU_CYCLE_1, VALU_DEP_1) per LLVM backend machine model. Other changes: - translation_rule.h: three-tier framework types with HazardPipeline enum, updated ExpandFn signature and LaneLayout constants - binary_translator: pass RegisterLiveness to try_lower_expand, literal constant passthrough for trailing instruction words - code_object_patcher: VGPR floor 128 for temp register headroom, KD symbol matching when strtab is empty (GFX950 COs) - hsa_translate_test: MFMA 16x16 fuzzing test (10 iterations, random FP16 inputs, CPU golden reference, 0 mismatches) - dbt-design.md: updated MFMA/WMMA and testing sections Tested: vector_add + matmul_mfma_16x16 pass on GFX1201. Signed-off-by: Tony Gutierrez <anthony.gutierrez@amd.com>

- Don't unmap host pages on free_memory_ioctl for FMM (user_va) allocations. ROCR's caching allocator and PyTorch's block pool reuse freed handles without unmapping, so premature unmap causes the GPU to read zeros on reuse. Wait until the actual munmap instead. - Fix SDMA COPY_LINEAR, CONST_FILL, and WRITE_DATA count masks from 0x3FFFFF (22-bit, 4MB max) to 0x3FFFFFF (26-bit, 64MB max) per the GFX9+ SDMA spec. Signed-off-by: Tony Gutierrez <anthony.gutierrez@amd.com>

- Handle SRC_POPS_EXITING_WAVE_ID (249) and NULL (250) in the scalar operand read path. Both return 0 in compute shaders. Previously triggered 'Unsupported encoding value for scalar read: 249' abort during OPT-125M inference. Signed-off-by: Tony Gutierrez <anthony.gutierrez@amd.com>

Signed-off-by: Tony Gutierrez <anthony.gutierrez@amd.com>

newling

I haven't looked at this, but after rebasing onto this, issues I have on my race detection plugin development branch are resolved. So LGTM!

* Added tests to make sure that if there is a refactor of the cmpx_writes_vcc function (for instance returns false if the instruction writes to scalar SGPR pair instead of VCC) that the refactor doesn't siliently generate the wrong code. * Made a minor fix to `test_v_cmpx_instructions_cannot_share_execute` in `test_cross_isa_analyzer.py` which was passing the encoding to `_can_share_execute`. This was changed in PR #5512

## Motivation [rocjitsu] Two small nit picks from PR #5512: * Added tests to make sure that if there is a refactor of the cmpx_writes_vcc function (for instance returns false if the instruction writes to scalar SGPR pair instead of VCC) that the refactor doesn't siliently generate the wrong code. * Made a minor fix to `test_v_cmpx_instructions_cannot_share_execute` in `test_cross_isa_analyzer.py` which was passing the encoding to `_can_share_execute`. This was changed in PR #5512 ## JIRA ID None ## Test Result Tests pass! ## Submission Checklist - [ X] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

The CDNA5 descriptor tests in triton-lang/triton#12512 pass on hardware but fail under Rocjitsu. Crossing raw B128 accesses load and store one extra DWORD near NUM_RECORDS. - CDNA5 buffer bounds (PR #10362, f81f9d4): retain complete accesses, including exact fits, and apply the PR's stated end + 4 <= NUM_RECORDS rule to crossing raw vector accesses. Keep SOFFSET in the check and preserve structured and independent stride bounds. - GFX12 atomic returns (PR #7336, 3a4e98a): test TH[0] independently of the non-temporal and cascade bits. Returning modes 3, 5 and 7 otherwise omit the old value and select non-returning metadata and wait counters. - Signed scratch operands (CDNA5 PR #6673, f6aec50; RDNA4 PR #5053, ca4b97e): sign-extend scalar and vector offsets before combining them and before CDNA5 payload scaling. Compensated negative operands otherwise become large positive private offsets and fault. - CDNA GLOBAL operands (PR #5512, 5a107ce): zero-extend the 32-bit VGPR offset when a scalar base is present. High-bit offsets otherwise resolve four GiB below the intended address. Preserve absolute VGPR pairs and CDNA5 signed addressing. CDNA1-4 manuals and LLVM agree on the unsigned rule. For buffer bounds, the Triton PR measures aligned B128 crossings with four or eight bytes remaining. Other widths and byte-granular bounds follow its stated rule as a model choice pending gfx1250 qualification. Record that limit beside the implementation and regression expectations. Cover bounds, inactive lanes, signed and unsigned operand combinations, atomic writeback and wait counters, including image metadata. Validate 49 Gluon cases with matching ISA and 172 hardware comparisons each on gfx1201 and gfx1100. Audit gfx942/gfx950 buffer, scratch and atomic contracts against the manuals. The standard CTest sweep retains the same 31 baseline failures.

The CDNA5 descriptor tests in triton-lang/triton#12512 pass on hardware but fail under Rocjitsu. Crossing raw B128 accesses load and store one extra DWORD near NUM_RECORDS. - CDNA5 buffer bounds (PR #10362, f81f9d4): retain complete accesses, including exact fits, and apply the PR's stated end + 4 <= NUM_RECORDS rule to crossing raw vector accesses. Keep SOFFSET in the check and preserve structured and independent stride bounds. - GFX12 atomic returns (PR #7336, 3a4e98a): test TH[0] independently of the non-temporal and cascade bits. Returning modes 3, 5 and 7 otherwise omit the old value and select non-returning metadata and wait counters. - Signed scratch operands (CDNA5 PR #6673, f6aec50; RDNA4 PR #5053, ca4b97e): sign-extend scalar and vector offsets before combining them and before CDNA5 payload scaling. Compensated negative operands otherwise become large positive private offsets and fault. - CDNA GLOBAL operands (PR #5512, 5a107ce): zero-extend the 32-bit VGPR offset when a scalar base is present. High-bit offsets otherwise resolve four GiB below the intended address. Preserve absolute VGPR pairs and CDNA5 signed addressing. CDNA1-4 manuals and LLVM agree on the unsigned rule. - PC exports and targets (import handling #9844, e5109e3): apply the architecture's extension rule to GETPC, CALL and SWAPPC return addresses. RDNA4 zero-extends 48 bits; older targets sign-extend 48 bits and CDNA5 sign-extends 57 bits. Keep imported PCs as raw 48/57-bit addresses for instruction fetch, consistent with RFE, and preserve all CDNA5 target bits. Limit the positive linked-rel32 tag to 48-bit targets. Remove the negative tag on every target: it aliases ordinary sign-filled PC exports. Both forms otherwise corrupt valid 57-bit targets or older high-PC round trips when a load bias is present. For buffer bounds, the Triton PR measures aligned B128 crossings with four or eight bytes remaining. Other widths and byte-granular bounds follow its stated rule as a model choice pending gfx1250 qualification. Record that limit beside the implementation and regression expectations. Cover bounds, inactive lanes, signed and unsigned operand combinations, atomic writeback and wait counters, including image metadata. Validate 49 Gluon cases with matching ISA and 172 hardware comparisons each on gfx1201 and gfx1100. Audit gfx942/gfx950 buffer, scratch and atomic contracts against the manuals. The standard CTest sweep retains the same 31 baseline failures. Decode and execute PC regression cases on gfx1100, gfx1201, gfx950 and gfx1250, including high address bits, next-PC boundaries and operand aliasing. Follow SETPC/SWAPPC through real page tables to fetch a high mapped destination, including both relocation-tag collisions with a nonzero load bias. Older linked-rel32 debugger coverage is retained. These are manual/LLVM-backed simulation checks without physical high-address PC qualification. RDNA3/3.5 sign-filled exports follow LLVM against conflicting manual tables; record that choice beside the generator and test oracle. Add explicit test braces for GCC.

…14352) The CDNA5 descriptor tests in [triton-lang/triton#12512](triton-lang/triton#12512) pass on hardware but fail under Rocjitsu. Crossing raw B128 accesses load and store one extra DWORD near `NUM_RECORDS`. Auditing neighboring instructions and extension rules across gfx942/gfx950, gfx1100/gfx1201 and gfx1250 exposed further address and return defects. - **CDNA5 buffer bounds** (introduced by #10362, `f81f9d40db67`): keep complete accesses, including exact fits, and apply the Triton PR's stated `end + 4 <= NUM_RECORDS` rule to crossing raw vector accesses. Count `SOFFSET` and preserve structured and independent stride checks. - **GFX12 atomic returns** (introduced by #7336, `3a4e98a167e6`): test `TH[0]` independently of the non-temporal and cascade bits. Return modes 3, 5 and 7 otherwise omit the old value and select non-returning metadata and wait counters. - **GFX12 scratch operands** (CDNA5 #6673, `f6aec5064565`; RDNA4 #5053, `ca4b97ef2cc6`): sign-extend scalar/vector inputs before combining them and before CDNA5 payload scaling. Compensated negative inputs otherwise become large positive private offsets and fault. - **CDNA1–4 GLOBAL operands** (introduced by #5512, `5a107cebd601`): zero-extend the 32-bit VGPR offset when a scalar base is present. High-bit offsets otherwise address four GiB below the intended location. Preserve absolute VGPR pairs and the separate CDNA5 signed rule. - **PC extension and width**: GETPC, CALL and SWAPPC now export signed 48-bit PCs on older targets, zero-filled 48-bit PCs on RDNA4, and signed 57-bit PCs on CDNA5. SETPC/SWAPPC use raw masked 48/57-bit internal addresses so high targets remain fetchable through page tables. CDNA5 keeps all 57 target bits and bypasses linked-rel32 tags. The ambiguous negative tag is removed on every target so sign-filled PC exports round-trip with a load bias; 48-bit targets retain the positive tag covered by gfx950 debugger integration. Previously high PCs exposed the wrong upper bits or lost valid target bits. The generator and its 40 emitted PC handlers agree. RDNA3/3.5 sign-filled exports follow LLVM against conflicting manual tables; high-address PC behavior has no physical qualification. The gfx1250 buffer evidence covers aligned B128 crossings with four/eight bytes remaining. Other widths and byte-granular bounds follow the Triton PR's stated rule as model choices pending gfx1250 hardware qualification; the source and tests say so explicitly. The gfx942/gfx950 manual audit preserves the existing legacy scratch policy where no legal allocated-input mismatch was demonstrated. Local validation: - 49 Gluon cases pass through matching descriptor/instruction ISA. The installed Triton predates the new descriptor frontend API, so this is an ISA-equivalent replay, not a rebuild of that API. - 172 saved physical gfx1201 comparisons and 172 physical gfx1100 comparisons match simulation: buffer bounds, atomic return controls, legal scratch operands and GLOBAL offsets. The buffer-clamp rule remains CDNA5-specific. - gfx942 and gfx950 buffer/atomic/scratch/GLOBAL contracts were checked against their manuals and LLVM. All eight compiled HIP GLOBAL load/store simulation cases pass; no physical gfx942/gfx950 qualification is claimed. - 28 PC test groups cover gfx1100/gfx1201/gfx950/gfx1250, including next-PC boundaries, target width and SWAPPC operand aliasing. The original handlers fail 11 of the initial 16 groups. Page-table-backed fetch and nonzero-load-bias tests also reproduce the additional import/relocation gaps; all 28 final PC groups pass, including GETPC-to-SETPC/SWAPPC round trips with a nonzero load bias. Encodings were independently verified with LLVM. Four groups execute the destination through real page tables, including both 57-bit relocation-tag collisions. These are source-backed simulation checks without physical high-address PC qualification. - After rebasing onto current `develop`, the project builds and all 2,851 generator checks pass, with 46 skipped. Earlier GCC syntax checks cover the atomic and PC test translation units with `-Werror=dangling-else`; explicit braces fix the reported CI build error. Changed-file formatting hooks pass. - Rebase validation passes all 313 focused CTests, including PC, address, atomic and newly landed fork/join coverage. Standard CTest retains exactly the same 31 baseline failure names and classifications: 31/9181 after rebasing versus 31/9010 before rebasing. No new failure names.

Copilot AI review requested due to automatic review settings April 27, 2026 21:39

github-actions Bot added the documentation Improvements or additions to documentation label Apr 27, 2026

Copilot AI reviewed Apr 27, 2026

atgutier requested review from Groverkss, MaheshRavishankar, accauble, amd-arosa and newling April 27, 2026 21:39

MaheshRavishankar approved these changes Apr 27, 2026

View reviewed changes

systems-assistant Bot added the organization: ROCm label Apr 27, 2026

atgutier changed the title ~~Users/agutierr/rj opt125 dbt mvp staging~~ rocjitsu: Harden ISA layer, virtual machine, and adds initial DBT Apr 27, 2026

atgutier force-pushed the users/agutierr/rj-opt125-dbt-mvp-staging branch from 032216d to f3a8888 Compare April 28, 2026 00:32

amd-arosa approved these changes Apr 28, 2026

View reviewed changes

Comment thread experimental/rocjitsu/docs/dbt-design.md Resolved

Comment thread experimental/rocjitsu/tests/dbt/translate_test.cpp Resolved

Groverkss reviewed Apr 28, 2026

View reviewed changes

Comment thread experimental/rocjitsu/lib/rocjitsu/src/rocjitsu/code/patch/instruction_builder.h Resolved

Groverkss reviewed Apr 28, 2026

View reviewed changes

Comment thread experimental/rocjitsu/lib/rocjitsu/src/rocjitsu/code/dbt/translation_rule.h Resolved

Groverkss reviewed Apr 28, 2026

View reviewed changes

Comment thread experimental/rocjitsu/lib/rocjitsu/src/rocjitsu/code/dbt/generated/encoding_cdna4_to_rdna4.h Resolved

Groverkss reviewed Apr 28, 2026

View reviewed changes

Comment thread experimental/rocjitsu/lib/rocjitsu/src/rocjitsu/code/patch/code_object_patcher.cpp Resolved

accauble reviewed Apr 28, 2026

View reviewed changes

Comment thread experimental/rocjitsu/lib/python/amdisa/codegen.py Resolved

atgutier force-pushed the users/agutierr/rj-opt125-dbt-mvp-staging branch from 682c156 to 2a06e1e Compare April 28, 2026 19:57

accauble reviewed Apr 28, 2026

View reviewed changes

Comment thread experimental/rocjitsu/lib/python/amdisa/codegen.py Resolved

Comment thread experimental/rocjitsu/lib/rocjitsu/src/rocjitsu/isa/arch/amdgpu/shared/execute_shared.h Resolved

Comment thread experimental/rocjitsu/lib/python/amdisa/codegen.py Resolved

atgutier added 9 commits April 28, 2026 16:05

rocjitsu: Update reviewer feedback for PR#5512

aaca031

Signed-off-by: Tony Gutierrez <anthony.gutierrez@amd.com>

atgutier force-pushed the users/agutierr/rj-opt125-dbt-mvp-staging branch from 2a06e1e to aaca031 Compare April 28, 2026 23:05

newling approved these changes Apr 29, 2026

View reviewed changes

atgutier merged commit 5a107ce into develop Apr 29, 2026
21 checks passed

atgutier deleted the users/agutierr/rj-opt125-dbt-mvp-staging branch April 29, 2026 14:18

accauble mentioned this pull request Apr 29, 2026

[rocjitsu] 2 small testing improvements after PR #5512 #5609

Merged

kuhar mentioned this pull request Oct 27, 2026

[rocjitsu] Fix AMDGPU buffer bounds, atomic returns and addressing #14352

Merged

Navigation Menu

Repository navigation

rocjitsu: Harden ISA layer, virtual machine, and adds initial DBT - #5512

rocjitsu: Harden ISA layer, virtual machine, and adds initial DBT#5512
atgutier merged 10 commits into
developfrom
users/agutierr/rj-opt125-dbt-mvp-staging

atgutier commented Apr 27, 2026

Uh oh!

Copilot AI left a comment

Uh oh!

MaheshRavishankar left a comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Groverkss commented Apr 28, 2026

Uh oh!

Uh oh!

accauble left a comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

newling left a comment

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

7 participants

Conversation

atgutier commented Apr 27, 2026

Motivation

Technical Details

JIRA ID

Test Plan

Test Result

Submission Checklist

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Uh oh!

MaheshRavishankar left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Groverkss commented Apr 28, 2026

Uh oh!

Uh oh!

accauble left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

newling left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

7 participants