Repository navigation
feat(tensilelite): enable OutputAmaxD on gfx1250 via flat atomic fallback - #11525
Closed
geotseng-amd wants to merge 4 commits into
Closed
geotseng-amd wants to merge 4 commits into
geotseng-amd wants to merge 4 commits into
Conversation
…lel kernel assembly Convert three gfx1250 (ISA 12.5) hard crashes into clean solution-level rejections, so unsupported configurations are pruned instead of aborting the whole build: - OutputAmaxD: the cross-workgroup amax reduction emits s_atomic_dec, for which gfx1250 has no ISA opcode (HasSAtomic=false) -> segfault. - UseInitialStridesCD: the conversion/beta-clear kernel cannot pass initial stride params yet -> AssertionError in KernelWriterConversion. - General (pointer-array) batched (StridedBatched=False): loadBatchedAddress emits duplicate CFG labels (CFGBuilderPass abort/SIGABRT) and the general-batched conversion kernel is unimplemented. All three use ProblemType.get(..., default) so unit-test ProblemType dicts without these keys do not raise KeyError. Also dedup kernels before the parallel write/assemble pass: duplicate kernels share one .s/.o path, so assembling each unique kernel once avoids a FileNotFoundError race under parallel builds. Co-authored-by: Cursor <cursoragent@cursor.com>
…back The amax cross-workgroup reduction elected its final workgroup with a scalar s_atomic_dec, which has no ISA mapping on gfx1250 and made any OutputAmaxD solution fail to assemble (previously blocked by a hard reject in Solution.py). Route the synchronizer decrement through a new HasSAtomic-gated helper (generateAtomicDecU32): CDNA keeps the byte-identical scalar SAtomicDec, while gfx1250 uses the proven MBSK/GSU flat_atomic_dec_u32 fallback (lane-0 EXEC mask + XNACK drain + SCOPE_DEV). No host change: the synchronizer counter is the same shared buffer with zero-init and kernel self-reset. Also generalize amax_inter_wave_reduction's cross-wave mask from a hardcoded wave64 form (2-sgpr VCmp dst + SAndB64) to wave-width-aware (laneSGPRCount + SAndB32/SAndB64); this was latent-broken on wave32 and is exercised now that amax builds on gfx1250 (wave32). Remove the gfx1250 OutputAmaxD reject and add a gfx1250 amax build test derived from PR ROCm#11117. Verified: gfx1250 build-only emits flat_atomic_dec_u32 (scope:SCOPE_DEV, th:TH_ATOMIC_RETURN) with zero s_atomic_dec and correct wave32 masks (s_and_b32 vcc_lo); wave64 arches are byte-identical. Full tensilelite unit suite: 6144 passed, 0 failed. Co-authored-by: Cursor <cursoragent@cursor.com>
github-actions
Bot
added
project: hipblaslt
project: hipsparselt
ci:hipsparselt-fast
labels
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Codecov Report❌ Patch coverage is ❌ Your project check has failed because the head coverage (35.22%) 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 #11525 +/- ##
===========================================
- Coverage 69.27% 69.26% -0.01%
===========================================
Files 2811 2811
Lines 465757 465948 +191
Branches 68729 68769 +40
===========================================
+ Hits 322642 322718 +76
- Misses 119574 119666 +92
- Partials 23541 23564 +23
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
…deSize under-count The gfx1250 amax (OutputAmaxD) epilogue emits `v_max_f32 vD, vD, abs(vSrc)`. The abs source modifier forces the VOP3 (_e64) 8-byte encoding, but the stinkytofu static instruction-size model counted these compact VALU ops as 4 bytes. Across the fp8 epilogue's ~33 abs-max instructions this under-counted STINKY_TOTAL_INST_BYTES and tripped the gfx1250-only CheckASMCodeSize verifier (Tensile::FATAL, build exit 255). - stinkytofu: in getEffectiveBaseSizeInBytesImpl(), promote a compact VALU to 8 bytes when any abs/neg/clamp/omod modifier is active -- covering both the instruction-level VOP3Modifiers and the per-operand StinkyRegister isAbs/ isMinus form produced by rocisa's vgpr(..., isAbs=True). Add unit tests for both representations plus a no-modifier 4-byte baseline. - tensilelite: cleanly reject OutputAmaxD on gfx1250 unless the store path can accumulate amax. It is gathered only on the single-element ScaleD store route (GlobalWriteBatch gates gwvw==1); an even store width uses the packed VMulPKF32 path that never touches AmaxOut (silent amaxD=0). The store gwvw here is StoreVectorWidth (non-atomic buffer store), not GlobalWriteVectorWidth, so require StoreVectorWidth==1. An early dtype/UseScaleCD proxy guards before the store width is derived. - tests: add OCP fp8 amax / ScaleCD runnable coverage in amax_gfx1250.yaml. Co-authored-by: Cursor <cursoragent@cursor.com>
Contributor
Author
|
tox passed on v0 / v1 |
…t and bound its tail reads The gfx1250 amax (OutputAmaxD) reduction produced intermittently inflated amaxD values. Three independent defects in amax_output_result(), all in how the cross-workgroup workspace is handed off on gfx1250: - The per-workgroup partial was stored and read back with glc/slc. gfx1250 defaults VMEM to CU-local scope (DefaultScopeIsCULocal), so the reducer -- which runs on whichever CU won the flat_atomic_dec_u32 race -- could be served a stale partial out of its own L1. Both sides now carry scope:SCOPE_DEV, matching the GSU workspace handshake. - The atomic dec that advertises the partial was ordered only by s_waitcnt vlcnt(0), which does not track stores on gfx1250. Add vscnt=0 so the partial has landed before the counter is decremented. - The reducer walks the workspace in whole-wave (32 slot) strides, so when NumGroup is not a multiple of the wave size the last iteration reads past the slots this launch wrote. The code relied on Src+2 (num_records) clipping those reads to zero, but gfx1250 uses SrdUpperValue125X, whose dword3 layout carries no byte-range oob_select -- staticInit() leaves it 0 where gfx9/10/11/12XX set oob_select=3. The tail slots therefore returned a previous launch's partials (the workspace is reused and never zeroed), and any residue larger than the current amax skewed the result. Mask those lanes in software rather than relying on the SRD, gated on asmCaps["HasPartialOOB"] so every arch whose hardware does clip keeps byte-identical code. amax_gfx1250.yaml only exercised 128/127/256, which at MT16x16 are 8x8, 8x8 and 16x16 workgroups -- all multiples of 32, so the tail path was never reached and the test passed even with the masking removed. Add 176 (11x11=121) and 112 (7x7=49), ordered larger-first so the reused workspace still holds bigger partials when the smaller problem reads past its own slots. Also align the four groups on 8rn_amaxd.yaml's full bias/scale/activation epilogue and cover all four transposes. Verified on gfx1250 (gfx1250v0): - amax_gfx1250.yaml: 4/4 fail with the epilogue fix reverted (176x176 and 112x112 amaxD mismatch), 4/4 pass with it -- previously intermittent. - tox py3, -k 'gfx1250 and not streamk', CheckASMCodeSize=True: 2 failed, 433 passed. Both failures (mxf6_tdm_gfx1250, stinky_sia4) are present in the pre-change baseline; amax_gfx1250 moves from failed to passed. - unit + characterization: 6906 passed, 763 snapshots unchanged. - gfx942 unaffected: the 30 amax kernels generated from 8rn_amaxd.yaml are byte-identical before and after. Co-authored-by: Cursor <cursoragent@cursor.com>
1 task
jichangjichang
pushed a commit
that referenced
this pull request
Oct 13, 2026
JIRA ID : AIHPBLAS-4357, AIHPBLAS-4369, AIHPBLAS-4390, AIHPBLAS-4394 | Ticket | Summary | Covered by | |---|---|---| | AIHPBLAS-4357 | UnboundLocalError in `computeLoadSrd()` for non-batched GEMM with SupportUserArgs on gfx1250 | `Batched=False` reject | | AIHPBLAS-4369 | Kernel generation crash when StridedBatched is set to false | `StridedBatched=False` reject | | AIHPBLAS-4390 | UseInitialStridesCD=True triggers unimplemented AssertionError in KernelWriterConversion | `UseInitialStridesCD=True` reject | | AIHPBLAS-4394 | Segmentation fault with OutputAmaxD=true on gfx1250 | Main fix in [#11525](#11525); this PR fixes the amax instruction-size under-count that makes `CheckASMCodeSize` fail | ## Motivation Running the gfx1250 end-to-end TensileLite yamls from [#11117](#11117) shows three independent problems: 1. Some ProblemType configurations cannot be generated for gfx1250. Each one either aborts TensileCreateLibrary or produces a kernel that does not assemble, so the failure is a crash instead of a clean rejection. 2. OutputAmaxD kernels on gfx1250 fail `CheckASMCodeSize` (228 bytes short). StinkyTofu's instruction-size model counts `v_max_num_f32 vD, vD, |vSrc|` as 4 bytes, but the abs modifier forces the 8-byte VOP3 encoding. 3. Five benchmark yamls abort with `hipErrorNoBinaryForGpu` because the generated code object contains no kernels. ## Technical Details **1. Reject unsupported problem types on gfx1250** (`SolutionStructs/Solution.py`) The rejects apply to ISA (12, 5), which covers gfx1250, gfx1250-strict and gfx1250v0. Other architectures are unchanged. | Configuration | Why it fails | |---|---| | `StridedBatched=False` | `loadBatchedAddress` emits the `moduleWg` labels twice, which StinkyTofu's CFG builder aborts on. `KernelWriterConversion` also emits invalid HIP for Batch{C,D}. General batch stays available at runtime via ArgType==3 on StridedBatched SupportUserArgs kernels. (AIHPBLAS-4369) | | `Batched=False` | `computeLoadSrd` branches to `StridedBatchedGemmLoad` labels that only exist when there is a batch index. `KernelWriterConversion` still emits `batchIdx = id2`. (AIHPBLAS-4357) | | `UseInitialStridesCD=True` | The `KernelWriterConversion` and `KernelWriterBetaOnly` helpers hard-code `stride0 == 1` and assert on it. (AIHPBLAS-4390) | | `UseInitialStridesAB=True` | A and B each need `len(IndexAssignments)` strides (e.g. 3). rocisa's kernarg loader has no 3-dword form, and `DeviceUserArguments` holds only two strides per tensor. | Like the General Batched checks above them, these rejects have no early `return`. With `NoReject` set, derivation runs to completion instead of failing later with `KeyError 'UseDotInstruction'` ("Error processing permutation"). Optional keys are read with `.get()` so that partial ProblemTypes in unit tests still work. **2. Count abs/neg source modifiers as VOP3** (`shared/stinkytofu`, AIHPBLAS-4394) `getEffectiveBaseSizeInBytesImpl` now returns 8 bytes for a compact VALU instruction when either: - it carries `VOP3Modifiers` with any `neg_src*` or `abs_src*` set; or - a source register has `reg.isAbs` or `reg.isMinus` set. This is what the rocisa-to-stinky conversion produces for `vgpr(..., isAbs=True)`. `clamp` and `omod` are not counted because the emitter does not print them. The new `tensilelite/Tests/common/amaxd/gfx1250/amax_gfx1250.yaml` (taken from [#11525](#11525)) covers this path. It mirrors `common/amaxd/8rn_amaxd.yaml`, which skips gfx1250, so develop never exercised the under-count. On gfx1250 the yaml uses the `flat_atomic_dec_u32` synchronizer fallback, and every group has the full bias/scale epilogue with Activation enabled. **3. Assemble each unique kernel once** (`TensileCreateLibrary/Run.py`) Deduplication uses the full kernel key, but the `.s`/`.o` file name comes from `getRequiredParametersMin`. A duplicate kernel therefore writes the same path as the kernel it duplicates, and its source is empty. `ParallelMap2` assembled both in parallel; when the duplicate finished last, its empty `.s` replaced the real one. `writeSolutionsAndKernels` now drops duplicates before assembling. This is the same change as the `Run.py` hunk in [#11525](#11525). Known follow-up, not addressed here: MagicDivAlg, DebugStreamK, ScheduleGlobalRead/LocalWrite and StreamKAtomic change codegen but are not part of the kernel name, so variants that differ only in those parameters still map to the same file. ## Test Plan All testing ran on a gfx1250 B0 host (`HSA_DISABLE_GFX12_STRICT=0`) with ROCm `10.2.0a20261012`, the newest nightly at the time. This host cannot run gfx1250-strict (A0) binaries. Clean rebuild of this branch (`75ea1572641`, on develop `e2aef14a26f`), following the gfx1250-strict A0 install guide (sections 5, 6 and 8): - StinkyTofu `ctest`, including 10 new `InstructionSizeCostingTest` cases. - hipBLASLt and hipSPARSELt builds for gfx1250, gfx1250v0 and gfx1250-strict, with the library layout check. - TensileLite tox unit tests, `test_arch_steppings.py` and the packaging script tests. - Runtime smoke test, `hipblaslt-test` full run and `*quick*`. - FP4 regression: `hipblaslt-bench` over every FP4 solution, plus the mxf4 smoke tests. - TensileLite tox `py3` (`tensilelite/Tests/common -k "not streamk"`). - `hipsparselt-test` full run. The same plan also ran on the pre-rebase base (develop `3377e3006f4` plus these commits), with the same results apart from the new develop tests noted below. Earlier on the same branch, with the previous ROCm nightly: - `amax_gfx1250.yaml` on GPU with `CheckASMCodeSize` enabled. - The 94 gfx1250 e2e yamls from #11117, run before and after these changes. ## Test Result | Step | Result | |---|---| | StinkyTofu `ctest` | 1636/1636 passed | | hipBLASLt build | rc=0, 3 groups; `gfx1250` 412 / `gfx1250v0` 397 / `gfx1250-strict` 400 files; targets `gfx1250`, `gfx1250`, `gfx1250-strict`; layout OK | | hipSPARSELt build | rc=0; 131 files in each of the three subtrees | | TensileLite tox unit | 10451 passed, 0 failed (328 skipped, 2 xfailed, 1 xpassed) | | `test_arch_steppings.py` | 829 passed | | Packaging scripts | 18 passed | | Smoke | 1152 passed | | `hipblaslt-test` all | 67970 passed, 8 failed (see below), 167 skipped | | `hipblaslt-test` `*quick*` | 18043 passed | | FP4 bench | 345 of 346 solutions pass (see below) | | FP4 mxf4 smoke | 32 passed | | tox `py3` | 10072 passed, 0 failed, 1518 skipped, 3 xfailed | | `hipsparselt-test` | 165405 passed, 0 failed | Failures, none caused by this PR: - `hipblaslt-test`: the 8 failures are all new tests from [#13031](#13031), which arrived with the rebase. `matmul_fast_check_mx_fp8` and `matmul_fast_check_mx_fp8_large_k` run on `gpu_arch: '(950|1250)'` and cover every fp8 pair. Every case with a bf8 operand (bf8×bf8, f8×bf8, bf8×f8) fails with `NO solution found!`. gfx1250's logic has MX kernels for f8×f8, F4 and F8F4 only, with no bf8 MX logic, so develop fails the same way. The f8×f8 cases pass. - FP4 bench: the one failure is solution 1830 (`F4BS ... MT256x256x128_MI16x16x1`). It was added by [#13201](#13201) and returns `nan` or `inf` for k ≥ 384 on develop as well. It is only reached when every solution is listed or the solution index is forced. Other results: - `amax_gfx1250.yaml`: 34/34 passed. Without change 2 it fails `CheckASMCodeSize`, 228 bytes short. - The rejects turn the generator crashes for these configurations into clean "0 valid solutions" results. - Duplicate-kernel fix: `magic_div_alg`, `debug_flags`, `schedule_global_local` and `assert_params` go from a core dump to all passing (8, 7, 10 and 16 tests). `max_lds_occupancy` no longer core dumps; its last step finds 0 valid solutions because of a bug in the yaml (MaxOccupancy 4 needs 81920 B of LDS, above MaxLDS 65536). After this PR, #11117 needs to drop the `_xfail` suffix from the first four. - On the first pre-rebase `hipblaslt-test` attempt the GPU hung on `nightly_matmul_deepbench_bf16_rbf16_rbf16_rbf16_rf32_r_NN_576_64_28000_1_576_28000_1_576_576_3_t2` (gfxhub no-retry page fault, then MES `REMOVE_QUEUE` timeout) and the host had to be rebooted. That test has since passed on its own in both runs, and no GPU errors appeared in dmesg during either run, so the hang did not reproduce. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
mpanoop
pushed a commit
to mpanoop/rocm-libraries
that referenced
this pull request
Oct 14, 2026
JIRA ID : AIHPBLAS-4357, AIHPBLAS-4369, AIHPBLAS-4390, AIHPBLAS-4394 | Ticket | Summary | Covered by | |---|---|---| | AIHPBLAS-4357 | UnboundLocalError in `computeLoadSrd()` for non-batched GEMM with SupportUserArgs on gfx1250 | `Batched=False` reject | | AIHPBLAS-4369 | Kernel generation crash when StridedBatched is set to false | `StridedBatched=False` reject | | AIHPBLAS-4390 | UseInitialStridesCD=True triggers unimplemented AssertionError in KernelWriterConversion | `UseInitialStridesCD=True` reject | | AIHPBLAS-4394 | Segmentation fault with OutputAmaxD=true on gfx1250 | Main fix in [ROCm#11525](ROCm#11525); this PR fixes the amax instruction-size under-count that makes `CheckASMCodeSize` fail | ## Motivation Running the gfx1250 end-to-end TensileLite yamls from [ROCm#11117](ROCm#11117) shows three independent problems: 1. Some ProblemType configurations cannot be generated for gfx1250. Each one either aborts TensileCreateLibrary or produces a kernel that does not assemble, so the failure is a crash instead of a clean rejection. 2. OutputAmaxD kernels on gfx1250 fail `CheckASMCodeSize` (228 bytes short). StinkyTofu's instruction-size model counts `v_max_num_f32 vD, vD, |vSrc|` as 4 bytes, but the abs modifier forces the 8-byte VOP3 encoding. 3. Five benchmark yamls abort with `hipErrorNoBinaryForGpu` because the generated code object contains no kernels. ## Technical Details **1. Reject unsupported problem types on gfx1250** (`SolutionStructs/Solution.py`) The rejects apply to ISA (12, 5), which covers gfx1250, gfx1250-strict and gfx1250v0. Other architectures are unchanged. | Configuration | Why it fails | |---|---| | `StridedBatched=False` | `loadBatchedAddress` emits the `moduleWg` labels twice, which StinkyTofu's CFG builder aborts on. `KernelWriterConversion` also emits invalid HIP for Batch{C,D}. General batch stays available at runtime via ArgType==3 on StridedBatched SupportUserArgs kernels. (AIHPBLAS-4369) | | `Batched=False` | `computeLoadSrd` branches to `StridedBatchedGemmLoad` labels that only exist when there is a batch index. `KernelWriterConversion` still emits `batchIdx = id2`. (AIHPBLAS-4357) | | `UseInitialStridesCD=True` | The `KernelWriterConversion` and `KernelWriterBetaOnly` helpers hard-code `stride0 == 1` and assert on it. (AIHPBLAS-4390) | | `UseInitialStridesAB=True` | A and B each need `len(IndexAssignments)` strides (e.g. 3). rocisa's kernarg loader has no 3-dword form, and `DeviceUserArguments` holds only two strides per tensor. | Like the General Batched checks above them, these rejects have no early `return`. With `NoReject` set, derivation runs to completion instead of failing later with `KeyError 'UseDotInstruction'` ("Error processing permutation"). Optional keys are read with `.get()` so that partial ProblemTypes in unit tests still work. **2. Count abs/neg source modifiers as VOP3** (`shared/stinkytofu`, AIHPBLAS-4394) `getEffectiveBaseSizeInBytesImpl` now returns 8 bytes for a compact VALU instruction when either: - it carries `VOP3Modifiers` with any `neg_src*` or `abs_src*` set; or - a source register has `reg.isAbs` or `reg.isMinus` set. This is what the rocisa-to-stinky conversion produces for `vgpr(..., isAbs=True)`. `clamp` and `omod` are not counted because the emitter does not print them. The new `tensilelite/Tests/common/amaxd/gfx1250/amax_gfx1250.yaml` (taken from [ROCm#11525](ROCm#11525)) covers this path. It mirrors `common/amaxd/8rn_amaxd.yaml`, which skips gfx1250, so develop never exercised the under-count. On gfx1250 the yaml uses the `flat_atomic_dec_u32` synchronizer fallback, and every group has the full bias/scale epilogue with Activation enabled. **3. Assemble each unique kernel once** (`TensileCreateLibrary/Run.py`) Deduplication uses the full kernel key, but the `.s`/`.o` file name comes from `getRequiredParametersMin`. A duplicate kernel therefore writes the same path as the kernel it duplicates, and its source is empty. `ParallelMap2` assembled both in parallel; when the duplicate finished last, its empty `.s` replaced the real one. `writeSolutionsAndKernels` now drops duplicates before assembling. This is the same change as the `Run.py` hunk in [ROCm#11525](ROCm#11525). Known follow-up, not addressed here: MagicDivAlg, DebugStreamK, ScheduleGlobalRead/LocalWrite and StreamKAtomic change codegen but are not part of the kernel name, so variants that differ only in those parameters still map to the same file. ## Test Plan All testing ran on a gfx1250 B0 host (`HSA_DISABLE_GFX12_STRICT=0`) with ROCm `10.2.0a20261012`, the newest nightly at the time. This host cannot run gfx1250-strict (A0) binaries. Clean rebuild of this branch (`75ea1572641`, on develop `e2aef14a26f`), following the gfx1250-strict A0 install guide (sections 5, 6 and 8): - StinkyTofu `ctest`, including 10 new `InstructionSizeCostingTest` cases. - hipBLASLt and hipSPARSELt builds for gfx1250, gfx1250v0 and gfx1250-strict, with the library layout check. - TensileLite tox unit tests, `test_arch_steppings.py` and the packaging script tests. - Runtime smoke test, `hipblaslt-test` full run and `*quick*`. - FP4 regression: `hipblaslt-bench` over every FP4 solution, plus the mxf4 smoke tests. - TensileLite tox `py3` (`tensilelite/Tests/common -k "not streamk"`). - `hipsparselt-test` full run. The same plan also ran on the pre-rebase base (develop `3377e3006f4` plus these commits), with the same results apart from the new develop tests noted below. Earlier on the same branch, with the previous ROCm nightly: - `amax_gfx1250.yaml` on GPU with `CheckASMCodeSize` enabled. - The 94 gfx1250 e2e yamls from ROCm#11117, run before and after these changes. ## Test Result | Step | Result | |---|---| | StinkyTofu `ctest` | 1636/1636 passed | | hipBLASLt build | rc=0, 3 groups; `gfx1250` 412 / `gfx1250v0` 397 / `gfx1250-strict` 400 files; targets `gfx1250`, `gfx1250`, `gfx1250-strict`; layout OK | | hipSPARSELt build | rc=0; 131 files in each of the three subtrees | | TensileLite tox unit | 10451 passed, 0 failed (328 skipped, 2 xfailed, 1 xpassed) | | `test_arch_steppings.py` | 829 passed | | Packaging scripts | 18 passed | | Smoke | 1152 passed | | `hipblaslt-test` all | 67970 passed, 8 failed (see below), 167 skipped | | `hipblaslt-test` `*quick*` | 18043 passed | | FP4 bench | 345 of 346 solutions pass (see below) | | FP4 mxf4 smoke | 32 passed | | tox `py3` | 10072 passed, 0 failed, 1518 skipped, 3 xfailed | | `hipsparselt-test` | 165405 passed, 0 failed | Failures, none caused by this PR: - `hipblaslt-test`: the 8 failures are all new tests from [ROCm#13031](ROCm#13031), which arrived with the rebase. `matmul_fast_check_mx_fp8` and `matmul_fast_check_mx_fp8_large_k` run on `gpu_arch: '(950|1250)'` and cover every fp8 pair. Every case with a bf8 operand (bf8×bf8, f8×bf8, bf8×f8) fails with `NO solution found!`. gfx1250's logic has MX kernels for f8×f8, F4 and F8F4 only, with no bf8 MX logic, so develop fails the same way. The f8×f8 cases pass. - FP4 bench: the one failure is solution 1830 (`F4BS ... MT256x256x128_MI16x16x1`). It was added by [ROCm#13201](ROCm#13201) and returns `nan` or `inf` for k ≥ 384 on develop as well. It is only reached when every solution is listed or the solution index is forced. Other results: - `amax_gfx1250.yaml`: 34/34 passed. Without change 2 it fails `CheckASMCodeSize`, 228 bytes short. - The rejects turn the generator crashes for these configurations into clean "0 valid solutions" results. - Duplicate-kernel fix: `magic_div_alg`, `debug_flags`, `schedule_global_local` and `assert_params` go from a core dump to all passing (8, 7, 10 and 16 tests). `max_lds_occupancy` no longer core dumps; its last step finds 0 valid solutions because of a bug in the yaml (MaxOccupancy 4 needs 81920 B of LDS, above MaxLDS 65536). After this PR, ROCm#11117 needs to drop the `_xfail` suffix from the first four. - On the first pre-rebase `hipblaslt-test` attempt the GPU hung on `nightly_matmul_deepbench_bf16_rbf16_rbf16_rbf16_rf32_r_NN_576_64_28000_1_576_28000_1_576_576_3_t2` (gfxhub no-retry page fault, then MES `REMOVE_QUEUE` timeout) and the host had to be rebooted. That test has since passed on its own in both runs, and no GPU errors appeared in dmesg during either run, so the hang did not reproduce. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
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.
JIRA ID: AIHPBLAS-4394
Motivation
OutputAmaxD(the fused absolute-max reduction used for FP8 dynamicscaling) was unusable on gfx1250. The amax cross-workgroup reduction
elected its final workgroup with a scalar
s_atomic_dec, which has noISA mapping on gfx1250, so any
OutputAmaxDsolution failed toassemble. It was therefore hard-rejected at solution level, leaving
gfx1250 without amax support.
This PR enables
OutputAmaxDon gfx1250 with the minimum-risk approach:reuse the already-validated MBSK/GSU
flat_atomic_dec_u32fallbackinstead of introducing new rocisa opcodes.
Technical Details
HasSAtomic-gated helpergenerateAtomicDecU32(
KernelWriterAssembly.py): on CDNA it emits the byte-identical scalarSAtomicDec; on gfx1250 (HasSAtomic=False) it usesflat_atomic_dec_u32under a lane-0 EXEC mask, with the XNACK-replaydrain (
s_wait_xcnt) andscope:SCOPE_DEVthat gfx1250 requires. Thismirrors the proven fallback in
Components/GSU.py::_generateAtomicDec.amax_output_resultnow routes its synchronizer decrement through thishelper.
amax_inter_wave_reduction: the cross-wave mask washardcoded to wave64 (2-sgpr
v_cmpdst +s_and_b64). Generalized towave-width-aware (
self.states.laneSGPRCount+SAndB32/SAndB64).This path was latent-broken on wave32 and is now exercised because
gfx1250 amax runs wave32. Wave64 emission is byte-identical.
OutputAmaxDreject inSolution.py.inputs.Synchronizerbuffer used by MBSK/StreamK, with one-timezero-init and kernel self-reset. The
flat_atomic_dec_u32wrapsemantics (wrap value
NumGroup-1,TH_ATOMIC_RETURN) reset thecounter back to 0 identically to the scalar path.
Tensile/Tests/common/amaxd/gfx1250/amax_gfx1250.yaml,a gfx1250-only amax build test derived from the amax problem groups in
PR feat(tensilelite): add end-to-end test suite for gfx1250 #11117.
Cross-architecture safety: for every previously-working kernel on
gfx908/gfx90a/gfx942/gfx950 (wave64,
HasSAtomic=True) all three changesare byte-identical no-ops. On RDNA arches amax never assembled before
(raw
s_atomic_dec), so the change is an enablement/bugfix, not aregression; non-amax kernels are untouched.
Test Plan
--build-onlyof the newamax_gfx1250.yaml(FP16OutputAmaxD, and BF16UseScaleCD+OutputAmaxD); disassemble thegenerated code objects to confirm the emitted instructions.
(
misc_params,scale_cd) build on gfx1250.Test Result
shows
flat_atomic_dec_u32 v.., v[..:..], v.. th:TH_ATOMIC_RETURN scope:SCOPE_DEV,s_atomic_deccount = 0,s_wait_xcnt 0x0present,and correct wave32 masks (
s_and_b32 vcc_lo, single-sgprv_cmp,s_and_b64count = 0).6144 passed, 971 skipped, 18 xfailed, 0 failed.misc_paramsPASSES;scale_cdamax groupsPASS (its unrelated FP8
UseScaleAB: Truebool-vs-string YAML error ispre-existing and out of scope).
correctness fully verified in the functional-model environment).
Submission Checklist