Repository navigation
fix(rocke): address #13205 review: develop merge, sinks gates, case study - #5
Merged
AviralGoelAMD merged 129 commits intoOct 14, 2026
Conversation
…12752) ## Motivation JIRA ID : AIFFT-517 The bitwise reproducibility tests store a fingerprint (hash) of each FFT result and compare it on later runs. The hash had three problems: - It used `std::hash`, which can give different results on different compilers or runs, so stored fingerprints could stop matching even when the output was identical. - Each value was multiplied by a weight and rounded before hashing, so a change in the last bit could be lost. - Each element's position was computed as `row + column`, so a transposed output produced the same hash. ## Technical Details - `shared/fft_hash.h`: each element is hashed from its exact bits together with its real position (batch, then row-major index), using a fixed hash function. Results are still added together, so the thread count doesn't affect the hash. - `projects/hipfft/shared/fft_hash.h`: kept identical to the rocFFT copy. - `clients/tests/buffer_hash_test.cpp`: new tests check that a transpose, a 1-bit change and a batch swap all change the hash, and that padding doesn't. Fixed expected hashes cover half, single and double precision. - `testing.md`: updated to describe the new hash. SHA-256 (`shared/sha256.h`) isn't used because its result depends on the order of the data, so it would change with the number of threads. Existing stored fingerprints must be regenerated. ## Test Plan On MI300X (gfx942), ROCm 7.2.4: - Ran the new tests against the old hash and the new one. - Ran the bitwise repro test twice in separate processes with `--repro-db`, for 5 FFT cases. ## Test Result - The old hash fails the transpose, 1-bit and padding tests. The new hash passes all hash tests. - Repro-db: the first run stored all 5 fingerprints, and the second run matched all 5. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.
… predicate and KD-tree search" (#13479) Reverts #13293 Fix issue :#13480 test run passed: https://github.com/ROCm/TheRock/actions/runs/38203742146/job/114696224272
JIRA ID : AIRADSW-483
## Motivation
On gfx1201 the TN GEMM types with 1-byte inputs F8F8S, B8B8S and I8II
use GridBased logic files of about 780 solutions
each. A 19-solution Origami (Prediction) catalog is faster on average
and about 40x smaller in solutions.
## Technical Details
- Removes
`asm_full/gfx1201/GridBased/gfx1201_Cijk_Alik_Bljk_{F8F8S,B8B8S}_BH_Bias_SHB_HA_S_SAB_SCD_SAV_UserArgs.yaml`
and `..._I8II_BH_SHB_HA_S_SCD_SAV_UserArgs.yaml`, and adds Prediction
catalogs with the same names in
`asm_full/gfx1201/Origami/`,
F8F8S against stock GridBased:
| shape set | geomean speedup |
|---|---|
| 400 random TN shapes, blind | 1.153 |
| 200 real TN workload shapes, blind | 1.231 |
🤖 Generated with [Claude Code](https://claude.com/claude-code)
Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…13387) JIRA ID: AIRADSW-483 ## Motivation On gfx1100 (Radeon RX 7900 XTX, "navi31"), hipBLASLt chooses the GEMM kernel for fp16 and bf16 NN matrix multiplies from GridBased libraries. A GridBased library is a lookup table that maps tuned problem sizes to kernels; sizes it was not tuned for fall back to the nearest entry. The NN tables are thin (471 rows each, against about 9 700 for TN), so most NN sizes take a fallback kernel. This PR replaces the four NN tables with Origami Prediction libraries, the approach navi32 adopted in #11305. A Prediction library holds a short list of kernels, and the Origami analytical model ranks them for each problem at run time. On fresh problem sizes held out of tuning, the new libraries are faster in geomean (the mean of per-shape speedups) for every type: - fp16 NN: 34% faster off the stock table, 17% faster on the stock table's own rows. - bf16 NN: 41% faster. - Aux (UseE) NN: 13–17% faster. - On a blind 194-shape benchmark set that was never used for tuning, fp16 NN wins 19 of 20 categories: geomean 1.334, +1.54M GFLOPS net. The exceptions are listed under Known limits. The TN libraries are not touched by this PR. ## Technical Details Directory: `projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/navi31/`. Each new file in `Origami/` replaces the `GridBased/` file of the same name. Both files describe the same problem type, so they cannot be in one library together. There are no code changes; the last commit adds one CHANGELOG line. | file (`navi31_Cijk_…_UserArgs.yaml`) | types | before | after | |---|---|---|---| | `Ailk_Bljk_HHS_BH_Bias_HAS_SAV` | fp16 in/out, fp32 compute | 70 kernels + 471-row table | 42 kernels | | `Ailk_Bljk_BBS_BH_Bias_HAS_SAV` | bf16 | 73 kernels + 471-row table | 42 (port of HHS-NN) | | `Ailk_Bljk_HHS_BH_Bias_AuxH_HAS_SAV` | fp16, also writes E | 70 kernels + 471-row table | 39 (port, StreamK→0) | | `Ailk_Bljk_BBS_BH_Bias_AuxB_HAS_SAV` | bf16, also writes E | 73 kernels + 471-row table | 39 (port, StreamK→0) | **Ports.** - The bf16 and Aux files are the fp16 catalog with the whole problem type copied from the stock file of the target type. - Aux means UseE: the GEMM also writes the pre-activation output E. With UseE, StreamK=3 kernels fail code generation on this tree, so those members were converted to StreamK=0. One converted member duplicated an existing one and was removed (40 → 39). - The Aux files are based on the 40-kernel HHS-NN catalog with the member-13 rollback. They do not include the narrow-N update, whose new StreamK=3 kernel cannot be used with UseE. 🤖 Generated with [Claude Code](https://claude.com/claude-code) --------- Co-authored-by: Claude Opus 5.5 (1M context) <noreply@anthropic.com>
…950 (#13374) ## Motivation On gfx942 and gfx950, each XCD has its own L2 cache, and the L2 caches are not coherent with each other. Agent-scope acquire and release fences therefore lower to a full L2 invalidate (`buffer_inv sc1`) and a full L2 write-back (`buffer_wbl2`). The level-scheduled triangular solves and incomplete factorizations issue these fences for every row: the producer before it sets the row's done flag, and the consumer after it has waited for the flag. As a result, these routines spend most of their time on cache maintenance on these architectures. The analysis stage has a second issue. Every row updates the maximum row length with its own device-scope atomic on a single address, and on these architectures these atomics serialize across the XCDs. ## Technical details **Solve and factorization kernels.** On gfx942, gfx950 and gfx9-4-generic (`has_noncoherent_xcd_l2`), row results and done flags are exchanged through device-scope relaxed atomic loads and stores, which carry `sc1` and are coherent across the XCD L2s. They replace plain accesses ordered by agent-scope fences. - Producer: waits for its result stores to complete (`__builtin_amdgcn_s_waitcnt(0)` behind a workgroup release fence) before it stores the done flag. No `buffer_wbl2`. - Consumer: after the spin loop on the flag, a workgroup acquire fence keeps the result loads after the flag loads. No `buffer_inv`. The new helpers are in `rocsparse_common.hpp`: - `acquire_done_rows` - `load_done_row_data` - `store_done_row_data` - `store_done_row_result` (falls back to `nontemporal_store` on other architectures) - `release_done_row_data` - `set_done_flag` - `wait_store_completion` (gfx9 only; used by `release_done_row_data` and `set_done_flag` when `has_noncoherent_xcd_l2` is true) On all other architectures the helpers compile to the previous code: plain loads and stores, nontemporal stores, `__threadfence()`, and the agent-scope acquire fence and release store. **Small-block BSR factorizations.** In these kernels one wavefront factorizes a block row, so only the blocks that other block rows read need device-scope accesses. The blocks that only the same wavefront reads back use plain loads and stores, because the memory model needs no cache maintenance for coherence between the lanes of one wavefront. - bsrilu0 (block sizes up to 8): dependent rows read only the diagonal and upper blocks, so only those are stored with device scope. The block row's own loads and its lower-block stores are plain. - bsric0 (block sizes up to 16): dependent rows read every block a row stores, so all stores stay device-scope. The re-reads of the row's own blocks are plain loads. The store wait after each block is replaced with a wavefront-scope fence, because `set_done_flag` waits for all stores before it sets the flag. Affected kernels: - solves: csrsv (also used for CSC and COO), ellsv, csrsm, and bsrsv for block sizes up to 32 - factorizations: csrilu0, csric0, and csrildlt0 (hash and binsearch variants), bsrilu0 for block sizes up to 8, and bsric0 for block sizes up to 16 with at most 128 blocks per block row bsrsm, bsrsv for block sizes above 32, and the larger-block bsrilu0 and bsric0 kernels keep develop's fence-based synchronization. These kernels reread each dependency's result many times, and every device-scope read bypasses the local L2, so the fence pair per dependency is cheaper there. **Analysis (gfx942, gfx950 and gfx9-4-generic only).** Each row updates a per-workgroup maximum row length with a shared-memory atomic. The last row of the workgroup to finish publishes it with a single device-scope atomic. This applies to the analysis of the CSR, CSC, COO, and BSR solves and factorizations, including the analysis stage of `rocsparse_spilu0`, `rocsparse_spic0`, and `rocsparse_spildlt0`. Other architectures keep one device-scope atomic per row. JIRA ID : ROCM-32435 --------- Co-authored-by: Cursor <cursoragent@cursor.com>
…nd it (#12638) ## Motivation A user hit an intermittent hang on MI350 spawning 128 worker `wthreads` from a coordinator `wthread`, well below the reported scheduler capacity of 256. It wasn't the scheduler. The coordinator kept its worker handles in a local container and overflowed its device stack. A `wthread` callable's stack budget is `hipLimitStackSize`, 1024 bytes by default, and `sizeof(hip::wthread)` is 16 bytes on device so it could hold about 63 handles regardless of `hardware_concurrency()`. This wasn't documented anywhere, and the failure gives no usable diagnostic, so reading it as a scheduler bug was reasonable. This PR documents the budget so users can size callables deliberately. Related: ROCM-31414 Closes #12473 ## Technical Details Docs only. `conceptual/hipThreads-execution-model.rst` — new "Stack space for callables" section: why the budget applies (the scheduler dispatches callables indirectly, so the runtime falls back to `hipLimitStackSize` instead of a compiler-computed frame), what counts against it, the two ways to design around it (device heap, or raise the limit), and how to recognize an overflow. `reference/limitations.rst` — one line on the existing private-stacks bullet cross-referencing the above, so it stays discoverable. ## Test Result - Scheduler kernel compiles with `.uses_dynamic_stack: true` / `.private_segment_fixed_size: 72`; faulting dispatch shows `private_seg_size=1024`. - Defaults measured: `hipLimitStackSize 1024 B`, `hipLimitMallocHeapSize 8 MB`, `sizeof(hip::wthread) 16`. - Ceiling lands exactly at 1024/16: 63 workers pass, 64 faults. - Both documented workarounds run - An ordinary HIP kernel with a 4 KB local gets a 4112-byte frame and runs fine, confirming the limit doesn't bind for normal HIP code. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. --------- Co-authored-by: Daniel McIntosh <Daniel.Mcintosh@amd.com>
## Motivation First PR of a stack adding a generic sparse matrix sort, `rocsparse_spsort`. It sorts an input matrix A into an output matrix B, or in place when the same descriptor is passed for both. COO, CSR and CSC are supported, including uniform strided batched matrices (every batch has the same sparsity pattern, and only the values differ). This PR adds only the public API and its documentation, so the interface can be reviewed on its own: - `rocsparse_spsort_buffer_size` and `rocsparse_spsort` in `internal/generic/rocsparse_spsort.h`. - The `rocsparse_spsort_descr` descriptor with `rocsparse_spsort_descr_create`, `rocsparse_spsort_descr_destroy` and `rocsparse_spsort_set_input`. - The `rocsparse_spsort_alg`, `rocsparse_spsort_input` and `rocsparse_spsort_stage` enums. - Entries in the reference docs and in the hipGraph support table. The functions are declared but not implemented until the following PRs in the stack. The doc example snippets are added together with the examples. ## Test plan - Compiled a C and a C++ translation unit including `rocsparse.h` and using the new types. - Built the library with this PR. JIRA ID: AISPARSE-717 ## Stack This PR is part of a stack. Review and merge in order: 1. #12553 ⬅ this PR 2. #12555 3. #12556 4. #12557 5. #12558 6. #12559 7. #12560 --------- Co-authored-by: Cursor <cursoragent@cursor.com>
## Motivation Implements the SpSort descriptor, `rocsparse_spsort_buffer_size` and `rocsparse_spsort` for COO matrices: - Sorts by row (`rocsparse_direction_row`) or by column (`rocsparse_direction_column`). - Out of place (A into B) or in place (the same descriptor for A and B). - Uniform strided batched matrices, where A and B may have different batch strides. Every batch has the same sparsity pattern, so the indices are sorted once and the resulting permutation is applied to the values of every batch. Adds the internal `rocsparse::coosort_buffer_size` and `rocsparse::coosort` in the existing `rocsparse_coosort.cpp`. They take the algorithm, the direction, the source and target sparse matrix descriptors and the buffer size, and gather the values of each batch with the shared permutation. The index type dispatch happens right before the kernel launch through the existing `gcoosort` routines, so the coosort templates are unchanged. The SpSort descriptor is defined in `rocsparse_spsort.hpp`, with `rocsparse_spsort_descr.cpp` (descriptor and `rocsparse_spsort_set_input`), `rocsparse_spsort_buffer_size.cpp` and `rocsparse_spsort.cpp`. The extra gather buffer for the values is only allocated when sorting in place. CSR and CSC return `rocsparse_status_not_implemented` until the next PRs. ## Test plan - Built the library and ran `example_rocsparse_spsort_coo`. - Ran the existing `coosort` tests (quick and pre_checkin). - Ran the `spsort_coo` tests added later in the stack against this library (quick and pre_checkin). JIRA ID: AISPARSE-717 ## Stack This PR is part of a stack. Review and merge in order: 1. #12553 2. #12555 ⬅ this PR 3. #12556 4. #12557 5. #12558 6. #12559 7. #12560 --------- Co-authored-by: Cursor <cursoragent@cursor.com>
## Motivation Adds CSR support to `rocsparse_spsort`. It sorts the column indices, and the corresponding values, within each row. The direction must be `rocsparse_direction_row`, and any other direction returns `rocsparse_status_invalid_value`. Out of place, in place and strided batched matrices are supported. Batches must be uniform: every batch has the same sparsity pattern, so the column indices are sorted once and the resulting permutation is applied to the values of every batch. Batch rules: - The columns/values batch stride must be at least nnz. - The offsets batch stride must be 0 (shared row pointer) or at least m + 1. - B may only share its row pointer if A does. Changes: - Adds the internal `rocsparse::csrsort_buffer_size` and `rocsparse::csrsort` in the existing `rocsparse_csrsort.cpp`, which take the source and target sparse matrix descriptors like the COO ones. They wrap `rocsparse::csxsort_buffer_size` and `rocsparse::csxsort`, which CSC reuses in the next PR. - Turns the existing csrsort implementation in the same file into file local templates on the row pointer and column index types, dispatched on the index types right before the call. The `rocsparse_csrsort` C API now calls those templates. - Adds the mixed 32-bit key and 64-bit value segmented radix sort instantiations needed for 64-bit row pointers with 32-bit column indices. ## Test plan - Built the library and ran `example_rocsparse_spsort`. - Ran the existing `csrsort` and `cscsort` tests (quick and pre_checkin). - Ran the `spsort_csr` tests added later in the stack against this library (quick and pre_checkin). JIRA ID: AISPARSE-717 ## Stack This PR is part of a stack. Review and merge in order: 1. #12553 2. #12555 3. #12556 ⬅ this PR 4. #12557 5. #12558 6. #12559 7. #12560 --------- Co-authored-by: Cursor <cursoragent@cursor.com>
## Motivation Adds CSC support to `rocsparse_spsort`. It sorts the row indices, and the corresponding values, within each column. Out of place, in place and strided batched matrices are supported, with the same batch rules as CSR, using n + 1 for the column pointer. Adds the internal `rocsparse::cscsort_buffer_size` and `rocsparse::cscsort` in the existing `rocsparse_cscsort.cpp`. They take the source and target sparse matrix descriptors and reuse `rocsparse::csxsort_buffer_size` and `rocsparse::csxsort` from the CSR PR, which sort a CSC matrix as the CSR layout of its transpose. ## Test plan - Built the library. - Ran the existing `csrsort` and `cscsort` tests. - Ran the `spsort_csc` tests added later in the stack against this library (quick and pre_checkin). JIRA ID: AISPARSE-717 ## Stack This PR is part of a stack. Review and merge in order: 1. #12553 2. #12555 3. #12556 4. #12557 ⬅ this PR 5. #12558 6. #12559 7. #12560 --------- Co-authored-by: Cursor <cursoragent@cursor.com>
…12558) ## Motivation Adds tests and a benchmark for `rocsparse_spsort` with COO matrices: - Checks out of place results (A unchanged) and in place results against a host reference, for both directions. - Covers strided batched matrices, with A padded between batches to check that its stride is honoured. - Checks the SpSort argument handling: descriptor inputs, stage ordering, matching A and B, and batch counts and strides. - Adds the `spsort_coo` benchmark. The test times are about 0.6 s for quick, 6 s for pre_checkin and 67 s for nightly. ## Test plan - `rocsparse-test --gtest_filter='*spsort*'`: 382 tests passed. - `rocsparse-bench -f spsort_coo` JIRA ID: AISPARSE-717 ## Stack This PR is part of a stack. Review and merge in order: 1. #12553 2. #12555 3. #12556 4. #12557 5. #12558 ⬅ this PR 6. #12559 7. #12560 --------- Co-authored-by: Cursor <cursoragent@cursor.com>
…12559) ## Motivation Adds tests and a benchmark for `rocsparse_spsort` with CSR matrices: - Checks out of place and in place results against a host reference that sorts each row. - Covers all supported row pointer and column index type combinations (i32/i32, i64/i32, i64/i64). - Covers strided batched matrices, with A's offsets and columns/values padded between batches. - Checks the CSR specific argument handling: the direction and the batch stride rules. - Adds the `spsort_csr` benchmark. The test times are about 0.5 s for quick, 5 s for pre_checkin and 35 s for nightly. ## Test plan - `rocsparse-test --gtest_filter='*spsort*:-nightly*spsort_coo*'`: 486 tests passed, including 179 nightly `spsort_csr` tests. - `rocsparse-bench -f spsort_csr` JIRA ID: AISPARSE-717 ## Stack This PR is part of a stack. Review and merge in order: 1. #12553 2. #12555 3. #12556 4. #12557 5. #12558 6. #12559 ⬅ this PR 7. #12560 --------- Co-authored-by: Cursor <cursoragent@cursor.com>
…12560) ## Motivation Adds tests and a benchmark for `rocsparse_spsort` with CSC matrices: - Checks out of place and in place results against a host reference that sorts each column. - Covers all supported column pointer and row index type combinations. - Covers strided batched matrices. - Checks the CSC specific argument handling: the direction and the batch stride rules. - Adds the `spsort_csc` benchmark. The test times are about 0.5 s for quick, 5 s for pre_checkin and 33 s for nightly. ## Test plan - `rocsparse-test --gtest_filter='quick*spsort*:pre_checkin*spsort*:nightly*spsort_csc*'`: 651 tests passed. - `rocsparse-bench -f spsort_csc` JIRA ID: AISPARSE-717 ## Stack This PR is part of a stack. Review and merge in order: 1. #12553 2. #12555 3. #12556 4. #12557 5. #12558 6. #12559 7. #12560 ⬅ this PR --------- Co-authored-by: Cursor <cursoragent@cursor.com>
…ed odist (#13017) ## Motivation `rocfft_plan_create` fails ("Can't find valid buffer assignment with current buffers.") for 1D complex transforms of lengths involving large prime factors that use the multi-kernel fused Bluestein algorithm, when the output batch distance is larger than the transform length (e.g. length 8209, batch 2, odist 20000). JIRA ID: AIFFT-551 (issue found while working on that) ## Technical Details Buffer assignment described the output of fused Bluestein nodes as a single transform length, while pairing it with *all* of the node's output strides. `compute_ptrdiff` then mispaired lengths and strides (and could underflow), so `CheckAssignmentValid` rejected every valid candidate assignment. - Add `OutputFootprint(node)`: the output lengths of a node with their matching output strides. The final fused Bluestein stage (`BFT_INV_CHIRP_MUL` block-RC) only stores the first `transform_length` elements of its padded dims 0-1, so its footprint is that length plus the remaining dims. - `CheckAssignmentValid`: size the output with `OutputFootprint`, so lengths and strides are paired correctly. - `dataFits`: use `OutputFootprint` when matching against the user buffer's length/stride. The dimension decomposition check still uses the output lengths as before. For transposes/SBCC writing to a non-unit-stride user buffer, the length/stride matching now uses correctly paired dimensions. - `MatchingLengthStride`: return `false` instead of reading out of bounds on inconsistent inputs. Stack note: with this fix, plans for inner-batched layouts (idist = odist = 1) using fused Bluestein are created but can fault on the GPU. That is fixed by the next PR in the stack (`users/regan-amd/fix_fused_bluestein_unit_distance`). ## Test Plan - New adhoc accuracy token: `complex_forward_len_8209_single_op_batch_2_istride_1_CI_ostride_1_CI_idist_8209_odist_20000_ioffset_0_0_ooffset_0_0` - Regular test suite. ## Test Result - New token fails before this change (plan creation failure) and passes after it. - Regular test suite passes on gfx1201. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
…ce (#13020) ## Motivation JIRA ID: AIFFT-13, AIFFT-551 (issue found while working on that) 1D complex transforms of lengths involving large prime factors that use the multi-kernel fused Bluestein algorithm fail for inner-batched layouts (unit batch distance, i.e. idist = odist = 1 and istride = ostride = batch). On develop, such plans either fail in `rocfft_plan_create` or fault on the GPU at execution. With the previous PR in this stack, plan creation succeeds, but execution still faults. ## Technical Details The L1D_CC "dist 1" special case (`CC1DNode::AssignParams_internal`) reshapes its children's kernel launch geometry, swapping the batch and sub-FFT dimensions so that the batch dimension is the adjacent one. Fused Bluestein stages have already set their Bluestein strides and dists against the original geometry, so after the reshape their kernels index past the Bluestein buffer (GPU memory access fault in the `sbcc_*_chirp_mul` kernels). The reshape is now skipped for fused Bluestein nodes (`BT_MULTI_KERNEL_FUSED`), keeping their launch geometry consistent with their Bluestein strides. This keeps the fused path (no fallback to non-fused Bluestein). ## Test Plan - New `fused_bluestein_non_unit_stride` accuracy suite (9 tokens): - inner-batched (idist = odist = 1, istride = ostride = batch) fused Bluestein layouts: forward/inverse, single/double, in-place/out-of-place, including the AIFFT-13 reproducers (lengths 4182, 4167, 4128 with batches 2456, 1039, 8092); - general non-unit stride and non-unit dist fused Bluestein layouts. - Regression: `prime_*`, `adhoc*`, `fused_bluestein*` and all `*_idist_1_*` suites in full, plus a 5% sample of all other single-device suites. ## Test Result - Before this stack: 6 of the new tokens fail in `rocfft_plan_create`; `complex_inverse_len_8191_single_op_batch_8_..._idist_1_odist_1` faults on the GPU (`sbcc_inv_chirp_mul`). - After this change: all 9 new tests pass. - Regression (gfx1201): full suite pass, no failures. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
Updated all pinned TheRock workflow and source refs to `015ba17` due to submodule bump Baseline run ID: [37869200457](https://github.com/ROCm/TheRock/actions/runs/37869200457) Co-authored-by: therockbot <therockbot@amd.com>
JIRA ID : ROCM-32153 Targets `develop`. It removes the quarantine that #13005 (AIHPBLAS-5015, now merged) added for these cases. ## Motivation On gfx950, two f16 GEMM solutions return success with wrong values for integer inputs at 4,099 x 2,053 x 129: a TN custom-schedule MT256x160x64 direct-to-LDS kernel, and an NN MT32x192x64 kernel. Neither is a tail-loop or edge problem. The TN kernel fails whenever its custom main loop runs, and the NN kernel fails at every K from 64 to 256. The TN kernel matters most, because the heuristic picks it by default for f16 TN 1,536 x 6,144 x 3,072 on MI355X, and it gave wrong answers in 6 of 30 runs there. I didn't find a small, clearly correct codegen fix for either kernel, so this PR removes them from the gfx950 logic. The custom schedule itself is tracked in ROCM-32343, which also covers its bf16 copies. ## Technical Details - Three solution entries come out of the gfx950 logic files (under `library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx950/`): - the TN kernel in `gfx950/Equality/..._Alik_Bljk_HHS...yaml` (#135) and its copy in `gfx950_id75a3/Equality/` (#151); - the NN kernel in `gfx950/Equality/..._Ailk_Bljk_HHS...yaml` (#334). - Each removal renumbers the solutions after it and drops the one exact-size entry the solution was tuned for, so that size falls back to the remaining selection. Every other exact-size entry still points at the same solution, and the build's logic validation passes. No tests or tuning files reference these kernels. - `known_bugs.yaml`: the two f16 `M: 4099` entries come out. ## Why remove these kernels instead of keeping them I know these kernels are fast. I looked hard at keeping them, and I don't think either one can stay as it is. **TN (custom-schedule MT256x160x64).** It fails whenever its custom main loop runs (K > 128), not just at odd sizes. That includes the one size it was tuned for: at f16 TN 1,536 x 6,144 x 3,072 the heuristic picks it by default on MI355X, and it was wrong in 6 of 30 runs. The kernel the heuristic picks instead (MT192x192x64) takes 58.7 µs against 57.1 µs, so about 2.8% slower, and it was right every time. On MI350X the difference is smaller. There, 68.6 µs against 68.1 µs is about 1% slower (the old kernel happened to be right in 10 of 10 runs there, but it fails the 4,099 x 2,053 x 129 cases on MI350X too). I think a few percent on one shape is a fair price for not returning wrong answers. **NN (MT32x192x64, StreamK 3).** I hoped to keep this one and just restrict it to the size it was tuned for (24 x 168 x 96, batch 6,144). The data doesn't support that. At exactly that size it was wrong in 30 of 30 runs on MI355X, with integer inputs (max error 3,344 to 3,480 ULP, where a correct kernel gives 0). It's wrong on MI350X there too. It's correct for up to 256 batches and wrong from 304 batches up, so my guess is it breaks once there are more tiles than CUs. On MI350X it was the heuristic's default at that size, so users were getting wrong answers from it. The replacement default (MT32x192x32) takes 84.6 µs against 74.8 µs, so it's about 13% slower, but it's correct. On MI355X the default at that size is a different kernel, and it doesn't change. One thing that made the NN problem easy to miss: `hipblaslt-bench -v` doesn't flag it. The bench computes the allclose result and then discards it, and the atol and rtol columns it prints are tolerances, not a verdict (ROCM-32137). Here the norm error was about 108, but the bench still exited 0. `--ulp` shows the problem right away. ROCM-32343 tracks fixing the custom schedule, so the TN kernels can come back once it's right. If someone fixes the NN kernel's Stream-K path, it can come back the same way. ## Test Plan On gfx950 (MI355X and MI350X): the f16 every-solution cases at 4,099 x 2,053 x 129 with the same test binary against this branch's library and the stack tip's, repeated; the ticket's `hipblaslt-bench -v` repros; and the heuristic default and timing at the sizes the removed solutions were tuned for. ## Test Result ROCm 10.2.0 (TheRock 10.2.0a20260928): - MI355X: with the stack tip's library, 4 of 16 f16 every-solution tests fail (NN and TN at 4,099); with this branch, 16 of 16 pass, and 96 of 96 over 6 repeats. - MI350X: the same, 4 of 16 failing before and 16 of 16 passing after. - The ticket's bench repros fail on the stack tip with CPU-reference errors of 0.073 and 0.045 (TN, both copies) and 0.019 (NN). - The heuristic default at 4,099 x 2,053 x 129 is unchanged and correct (MT224x160x32 for TN, MT256x144x32 for NN). - f16 TN 1,536 x 6,144 x 3,072 on MI355X now uses MT192x192x64: 58.7 µs with 0 of 10 runs wrong, against 57.1 µs for the removed kernel, which was sometimes wrong. That's about 2.8% slower. - Math CI precheckin passed against develop. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (2/5): logic-file data only, no code. Three wrong solutions stop being offered on gfx950; one tuned size gets a correct kernel that is about 3% slower on MI355X. Co-authored-by: Cursor <cursoragent@cursor.com>
## Motivation Adds the `BUILD_WITH_SPSORT` build option, which gates `rocsparse_spsort` so that it can be enabled or disabled at build time, in the same way as `BUILD_WITH_DIAGONAL_SOLVE`. The option is `ON` by default. When the option is `OFF`, `ROCSPARSE_WITH_SPSORT` is not defined in `rocsparse-config.h` and the following are left out: - The `rocsparse_spsort` and `rocsparse_spsort_buffer_size` declarations, the `rocsparse_spsort_descr_create`, `rocsparse_spsort_descr_destroy` and `rocsparse_spsort_set_input` declarations, and the `rocsparse_spsort_descr`, `rocsparse_spsort_alg`, `rocsparse_spsort_input` and `rocsparse_spsort_stage` types. - The implementation of those routines. Their sources are still listed in the library, but compile to nothing. - The `spsort_coo`, `spsort_csr` and `spsort_csc` tests and benchmarks. - The `rocsparse_spsort` documentation examples, which build as a stub that prints that rocSPARSE was built without `rocsparse_spsort`. `ROCSPARSE_WITH_SPSORT` is also added to the Doxygen `PREDEFINED` list so the documentation is still generated. The legacy `coosort`, `csrsort` and `cscsort` routines, and the fixes made to them earlier in the stack, are not gated. This PR only adds the guards. It does not change the code of the PRs below it, so their CI results are unaffected. ## Test plan - Build with `-DBUILD_WITH_SPSORT=OFF`: the library, `rocsparse-test`, `rocsparse-bench` and the samples build, `librocsparse.so` exports no `rocsparse_spsort*` symbols, and the spsort examples print the stub message. `rocsparse-test --gtest_filter='*coosort*:*csrsort*:*cscsort*'`: 556 tests passed. - Build with `-DBUILD_WITH_SPSORT=ON`: `rocsparse-test --gtest_filter='*spsort*'`: 1482 tests passed. JIRA ID: AISPARSE-717 ## Stack This PR is part of a stack. Review and merge in order: 1. #12553 2. #12555 3. #12556 4. #12557 5. #12558 6. #12559 7. #12560 8. #13392 ⬅ this PR --------- Co-authored-by: Cursor <cursoragent@cursor.com>
JIRA ID : AIHPBLAS-4195 ## Motivation HostNumerics is a new CPU-only component under `shared/hostnumerics/` that will hold shared data-type definitions, test-data generation, and CPU reference math for hipBLASLt, TensileLite, and rocRoller. Its standalone automated GitHub checks currently build and test it with GCC only. Those checks do not catch invalid memory access, undefined behavior such as signed integer overflow, or bugs on code paths the tests never run. ## Technical Details - Renames the two existing jobs to start with "GCC, " and sets `CXX=g++` so the compiler is explicit. - Adds "Clang, ASan and UBSan", a Debug build with Clang 18 using AddressSanitizer (detects out-of-bounds and use-after-free memory access) and UndefinedBehaviorSanitizer (detects undefined behavior such as signed overflow). `-fno-sanitize-recover=all` makes any finding fail the test. - Adds "Clang-tidy", which runs clang-tidy 18, a static analyzer that inspects code without running it, over the C++ sources and HostNumerics headers. Any warning fails the job. - Adds `shared/hostnumerics/.clang-tidy`, enabling the Clang static analyzer checks plus 13 `bugprone-*` checks, including checks for narrowing, integer division, misplaced widening, signed-char misuse, and loop-variable width. - The clang-tidy job builds with Python on and analyzes `python/src/bindings.cpp` as well as the C++ tests and HostNumerics headers. Its header filter excludes nanobind headers. The sanitizer job keeps Python off; sanitizing the Python bindings remains follow-up work. - Static analysis stays in the standalone workflow because the shared clang-tidy workflow currently installs TheRock and ROCm on a ROCm worker; HostNumerics needs only a sparse component checkout and an Ubuntu CPU runner. - No product code or build files change. ## Test Plan Each job was shown to fail on a deliberate defect, using four throwaway PRs based on this one: 1. #13498 binds Python's `storage_bits` to the wrong field. 2. #13499 turns tests on by default for embedded consumers. 3. #13500 adds a heap buffer overflow and a signed integer overflow in separate tests. 4. #13501 adds a null pointer dereference on a path the runtime tests never execute. ## Test Result All four HostNumerics jobs passed on the original PR head ([run](https://github.com/ROCm/rocm-libraries/actions/runs/38009876405)). The review updates were validated locally by replaying all four jobs with GCC 13.3 and Clang 18.1.3. A temporary null dereference in `bindings.cpp` and deliberate defects for each added numeric check failed analysis as expected. All four HostNumerics jobs also pass on the updated head ([run](https://github.com/ROCm/rocm-libraries/actions/runs/38301400794)); its clang-tidy log confirms that `python/src/bindings.cpp` is analyzed. In each demonstration PR, only the targeted job failed and the other three passed: - #13498: [GCC, Python ON, tests ON](https://github.com/ROCm/rocm-libraries/actions/runs/38268525203/job/114876732319) failed its metadata test. - #13499: [GCC, Python OFF, tests OFF](https://github.com/ROCm/rocm-libraries/actions/runs/38268533527/job/114876759448) failed its embedded-defaults check. - #13500: [Clang, ASan and UBSan](https://github.com/ROCm/rocm-libraries/actions/runs/38268545089/job/114876797348) reported `heap-buffer-overflow` and `signed integer overflow`. - #13501: [Clang-tidy](https://github.com/ROCm/rocm-libraries/actions/runs/38268554101/job/114876843557) reported `clang-analyzer-core.NullDereference`. TheRock CI was skipped with `ci:skip` because only the HostNumerics workflow changes. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (2/5): changes only the HostNumerics CI workflow and adds a clang-tidy config used by that workflow. No product code changes, and these jobs are not required checks, so a false failure would not block other PRs.
JIRA ID : AIHPBLAS-4195 ## Motivation HostNumerics is a new CPU-only component under `shared/hostnumerics/`. It will hold the data-type definitions, test-data generation, and CPU reference math that hipBLASLt, TensileLite, and rocRoller each implement separately today. #13405 added `ScalarType`, an enum naming every number format the component supports. Templated C++ code also needs to find the `ScalarType` for a C++ type such as `float` or `std::complex<double>`, and there is currently no way to do that. ## Technical Details - Adds `nativeScalarType<T>`, a compile-time constant giving the `ScalarType` for a C++ type. For example, `nativeScalarType<const float&>` is `ScalarType::Float32`; `const`, `volatile`, and references are ignored. - Adds the `NativeScalar<T>` concept, a compile-time yes/no check, so callers can test for a mapping without causing a compile error. - Maps `bool`, the fixed-width integers `int8_t` through `uint64_t`, `float`, `double`, `std::complex<float>`, and `std::complex<double>`. Any other type, including `long double`, pointers, arrays, and strings, is a compile error. - The lookup table is in `detail::NativeScalarType` and is not meant to be extended by users. - Header-only change to `scalar_type.hpp`, which now includes `<complex>` and `<type_traits>`. No build files change, and no code outside `shared/hostnumerics/` uses the component yet. ## Test Plan `shared/hostnumerics/tests/scalar_type_test.cpp` adds 24 `static_assert` checks: each of the 13 mappings, qualifier handling on 3 types, 2 positive `NativeScalar` checks, and rejection of 6 unsupported types. Because these are compile-time checks, a wrong mapping fails the build of `hostnumerics-scalar-type-test`. ## Test Result The HostNumerics standalone CI job "Python ON, tests ON, Ninja Multi-Config" on Ubuntu 24.04 built and ran the test and passed. TheRock CI was skipped with the `ci:skip` label because no other ROCm component builds HostNumerics yet. No Windows build has been run. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low (2/5): adds new header-only declarations and changes no existing behavior. Nothing outside `shared/hostnumerics/` includes this header yet, so no product, GPU, or binary-interface behavior can change.
…13389) ## Summary - Register `rocsparse-unit-test` and `rocsparse-unit-test-device` on the `standard`, `comprehensive`, and `full` ctest labels, including the installed `bin/rocsparse/CTestTestfile.cmake` that TheRock runs. `quick` does not include them. - Each binary is registered twice. TheRock on gfx1151 requires the `ex_gpu_gfx1151` label, and on every other GPU it excludes `ex_gpu`, so one entry cannot satisfy both. - Include the gfx8/gfx9 wfreduce lane check from #13388. Without it, `rocsparse-unit-test-device` fails 14 cases on gfx94X and gfx950, which are TheRock Linux test targets. This supersedes #13388. ## Test plan - [ ] TheRock `standard` on gfx94X and gfx950 runs both unit-test binaries, and the wfreduce cases pass - [ ] gfx1151 selects the `ex_gpu_gfx1151` entries rather than the plain ones - [ ] `ctest -L quick` does not run either unit-test binary --------- Co-authored-by: Kim Liegeois <kliegeois@users.noreply.github.com> Co-authored-by: James Sandham <33790278+jsandham@users.noreply.github.com>
) Resolves: #13517 ## Motivation On ROCm 7.0.2, `hipblaslt-test --gtest_filter='*fast_check_placement*:*fast_check_mx_fp8_placement*-*known_bug*'` fails 4 of 30 tests with `unknown error` at the first `hipGetLastError()` check in `testing_matmul_with_bias` (`CHECK_DEVICE_ALLOCATION` after allocating A). The failing tests vary with allocation order (e.g. `PLACE_c`, nightly `PLACE_b`, mx-fp8 `PLACE_scale_b`), pass when run alone, and the error is not caused by the failing test: it is left behind by the previous test's `PlacedRegion` teardown, where `hipMemAddressFree` fails with `hsa_amd_vmem_address_free ... status:4097` (`HSA_STATUS_ERROR_INVALID_ARGUMENT`). ## Technical Details Root cause is a reference leak in the HIP runtime before ROCm 7.1: 1. `testing_matmul_with_bias` destroys its stream without synchronizing it. The last command on that stream is usually fast_check's poison refill of the placed buffer (`PlacedRegion::fill_poison`). 2. `HostQueue::terminate()` takes a reference to the stream's last command, and when a CPU submission batch is still pending it queues a marker and overwrites that pointer without releasing it. Fixed upstream in ROCm/rocm-systems@33085dd232 (SWDEV-533220, "Release marker when HostQueue is destroyed", ROCm 7.1.0+). 3. The leaked kernel command keeps the placed buffer window's `amd::Memory` alive, so `hipMemUnmap` does not destroy it, and the window keeps the VA reservation's object alive after `hipMemAddressFree`. 4. The next `hipMemAddressReserve` that returns the same base cannot replace HIP's stale `VirtualMemObjMap_` entry (`insert` does not overwrite). `hipMemAddressFree` ignores its size argument and frees using the stale object's size, which ROCr rejects when the sizes differ. The error stays as the thread's last error and is reported by the next test. This was confirmed with gdb reference-count tracing in the test process, and with a standalone HIP reproducer: map a window, launch a kernel on it on a stream, destroy the stream without syncing, unmap/free, then reserve a different size at the same base. The fix synchronizes the stream before `hipStreamDestroy`, so no batch is pending when the stream is destroyed and nothing leaks. This is harmless on runtimes that already have the upstream fix. ## Test Plan - Build `hipblaslt-test` against ROCm 7.0.2 (gfx950) and run `--gtest_filter='*fast_check_placement*:*fast_check_mx_fp8_placement*-*known_bug*'` with `AMD_LOG_LEVEL=1`, with and without this change. - Run the same filter on ROCm 7.2.1 (gfx950) to check for regressions. ## Test Result - ROCm 7.0.2, gfx950, without the change: 26/30 passed, 4 failed, 4 failed `hsa_amd_vmem_address_free` calls. - ROCm 7.0.2, gfx950, with the change: 30/30 passed in two consecutive runs, no failed `hsa_amd_vmem_address_free` calls. - ROCm 7.2.1, gfx950, with the change: the 26 non-MX placement tests pass and no frees fail. The 4 `pre_checkin_matmul_fast_check_mx_fp8_placement` tests fail with incorrect D values, identically with and without this change; that is a separate issue, unrelated to this fix. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. ## Risk level Low. Test-client-only change: one `hipStreamSynchronize` before destroying a stream that has already been used synchronously, at the end of the test function. Co-authored-by: Cursor <cursoragent@cursor.com>
JIRA ID: AICK-2310
# Add scheduled TDM FMHA for FP16 and BF16
## Summary
- Add the `qr_tdm_sched` forward FMHA pipeline for FP16/BF16 on gfx125:
Q/K–V dimensions 64/64, 128/128, and 192/128. The implementation uses
geometry-derived QK/PV scheduling, packed LDS, and scheduled
WMMA/LDS-read placement.
- Generate and try `qr_tdm_sched` first for its supported row-major V
cases with `max_seqlen_q >= 128`, `sink_ptr == nullptr`, and no bias,
dropout, soft cap, or quantization. Dense and supported causal masks are
included; LSE is optional.
- Select N64 for D64 and N128 for D128. D192 selects N128 when the
average K length is at least 512, K-length pointers are absent, the N128
workgroup count fits the CU count, and the mask is dense or unrestricted
square causal; other supported D192 cases select N64.
- Fall back to the existing `qr_tdm` candidate when a scheduled trait is
ineligible, including short Q or a non-null virtual-sink pointer. Shapes
or features outside `qr_tdm` continue through the existing `qr`
fallback.
- Separate the no-sink and finite-sink kernel specializations, and add
codegen, dispatch, layout, schedule, sink, and numerical regression
tests under `test/ck_tile/fmha`.
## BF16 performance
All performance numbers in this PR description are taken from the
unmodified `tile_example_fmha_fwd` output. The two binaries use the same
example main, inputs, GPU timer, and command-line options; one links
`qr_tdm_sched` from this PR and the other links upstream `qr_tdm`.
Kernel names were printed and checked on every run.
Measured on gfx1250 **B0 stepping** (256 CUs). SCLK was neither locked
nor logged by the benchmark, so no fixed SCLK value is claimed. The
historical 110 workloads per dimension use BSHD layout, `Sq = Sk` in
{256, 2048, 8192, 16384, 32768}, dense or causal masking, LSE, no sink,
and the same scale and input seed for both kernels. The example uses
GPU-event timing, five warmups, each workload’s historical repeat count,
and two adjacent alternating rounds. For each kernel, the table reports
the median of the two values printed by the example. Speedup divides the
two printed TFLOP/s medians; above 1× favors `qr_tdm_sched`. The example
prints latency to 0.001 ms and throughput to 0.01 TFLOP/s.
| Q/K–V dimension | Workloads faster with qr_tdm_sched (Introduced by
this PR) | Median speedup vs qr_tdm | Speedup range |
| --- | ---: | ---: | ---: |
| 64/64 | 110/110 | 1.317× | 1.054–1.819× |
| 128/128 | 110/110 | 1.319× | 1.145–1.882× |
| 192/128 | 110/110 | 1.179× | 1.035–1.882× |
The full per-workload results follow. Each kernel cell shows **execution
time in µs / effective TFLOP/s**. `B/H/Hkv` denotes batch size, query
heads, and key/value heads. Causal TFLOP/s uses triangular attention
work.
<details><summary>BF16 Q/K=64, V=64: all 110 workloads</summary>
| B/H/Hkv | Sq = Sk | Mask | qr_tdm_sched (Introduced by this PR)<br>µs
/ TFLOP/s | qr_tdm<br>µs / TFLOP/s | qr_tdm_sched / qr_tdm |
| --- | ---: | --- | ---: | ---: | ---: |
| 1/8/8 | 256 | dense | 6.5 / 20.60 | 8 / 16.43 | 1.254× |
| 1/8/8 | 256 | causal | 9 / 7.87 | 9.5 / 7.07 | 1.112× |
| 1/8/8 | 2048 | dense | 20 / 428.70 | 32 / 268.56 | 1.596× |
| 1/8/8 | 2048 | causal | 22 / 198.80 | 29 / 148.99 | 1.334× |
| 1/8/8 | 8192 | dense | 105 / 1306.69 | 143 / 961.35 | 1.359× |
| 1/8/8 | 8192 | causal | 76.5 / 899.52 | 104 / 658.12 | 1.367× |
| 1/8/8 | 16384 | dense | 352.5 / 1558.99 | 516 / 1065.49 | 1.463× |
| 1/8/8 | 16384 | causal | 205.5 / 1336.91 | 231 / 1191.24 | 1.122× |
| 1/8/8 | 32768 | dense | 1316 / 1670.73 | 1783.5 / 1233.18 | 1.355× |
| 1/8/8 | 32768 | causal | 683.5 / 1608.76 | 831 / 1323.27 | 1.216× |
| 1/64/8 | 256 | dense | 6 / 167.20 | 8 / 131.09 | 1.275× |
| 1/64/8 | 256 | causal | 9 / 61.97 | 10 / 55.78 | 1.111× |
| 1/64/8 | 2048 | dense | 54 / 1272.37 | 77 / 896.91 | 1.419× |
| 1/64/8 | 2048 | causal | 35 / 972.00 | 44 / 784.50 | 1.239× |
| 1/64/8 | 8192 | dense | 670.5 / 1640.03 | 889.5 / 1235.90 | 1.327× |
| 1/64/8 | 8192 | causal | 366.5 / 1499.62 | 443 / 1241.95 | 1.207× |
| 1/64/8 | 16384 | dense | 2584.5 / 1701.43 | 3442.5 / 1277.50 | 1.332×
|
| 1/64/8 | 16384 | causal | 1362 / 1614.49 | 1637.5 / 1343.09 | 1.202× |
| 1/64/8 | 32768 | dense | 10242.5 / 1717.59 | 13515.5 / 1301.62 |
1.320× |
| 1/64/8 | 32768 | causal | 5295.5 / 1661.06 | 6376.5 / 1379.45 | 1.204×
|
| 1/64/64 | 256 | dense | 6 / 166.81 | 8 / 129.20 | 1.291× |
| 1/64/64 | 256 | causal | 9 / 61.73 | 10 / 55.41 | 1.114× |
| 1/64/64 | 2048 | dense | 54.5 / 1260.94 | 77 / 895.85 | 1.408× |
| 1/64/64 | 2048 | causal | 36 / 961.11 | 44.5 / 776.22 | 1.238× |
| 1/64/64 | 8192 | dense | 700.5 / 1569.05 | 906 / 1214.18 | 1.292× |
| 1/64/64 | 8192 | causal | 375.5 / 1464.12 | 450 / 1222.10 | 1.198× |
| 1/64/64 | 16384 | dense | 2859.5 / 1538.04 | 3619.5 / 1214.97 | 1.266×
|
| 1/64/64 | 16384 | causal | 1461 / 1505.70 | 1733 / 1269.13 | 1.186× |
| 1/64/64 | 32768 | dense | 11411.5 / 1541.58 | 14538 / 1210.10 | 1.274×
|
| 1/64/64 | 32768 | causal | 5834.5 / 1507.61 | 6931.5 / 1269.02 |
1.188× |
| 1/12/12 | 256 | dense | 8 / 26.42 | 9 / 21.56 | 1.225× |
| 1/12/12 | 256 | causal | 10 / 10.31 | 11 / 9.39 | 1.098× |
| 1/12/12 | 2048 | dense | 23 / 565.14 | 35 / 368.93 | 1.532× |
| 1/12/12 | 2048 | causal | 23 / 279.37 | 30.5 / 211.40 | 1.321× |
| 1/12/12 | 8192 | dense | 147.5 / 1399.60 | 181.5 / 1137.03 | 1.231× |
| 1/12/12 | 8192 | causal | 92 / 1122.19 | 115 / 898.60 | 1.249× |
| 1/12/12 | 16384 | dense | 516.5 / 1596.95 | 670.5 / 1229.97 | 1.298× |
| 1/12/12 | 16384 | causal | 280 / 1471.83 | 349 / 1181.07 | 1.246× |
| 1/12/12 | 32768 | dense | 1977 / 1668.20 | 2623.5 / 1257.21 | 1.327× |
| 1/12/12 | 32768 | causal | 1021 / 1614.88 | 1219 / 1353.24 | 1.193× |
| 1/32/32 | 256 | dense | 6 / 85.71 | 8 / 67.95 | 1.261× |
| 1/32/32 | 256 | causal | 8 / 32.12 | 9.5 / 28.52 | 1.126× |
| 1/32/32 | 2048 | dense | 30 / 1143.37 | 41 / 831.19 | 1.376× |
| 1/32/32 | 2048 | causal | 25 / 696.09 | 32 / 541.97 | 1.284× |
| 1/32/32 | 8192 | dense | 349.5 / 1574.21 | 460 / 1194.55 | 1.318× |
| 1/32/32 | 8192 | causal | 192 / 1430.57 | 235 / 1171.82 | 1.221× |
| 1/32/32 | 16384 | dense | 1334 / 1648.75 | 1763 / 1247.51 | 1.322× |
| 1/32/32 | 16384 | causal | 696 / 1580.20 | 832.5 / 1321.26 | 1.196× |
| 1/32/32 | 32768 | dense | 5428.5 / 1620.33 | 7043.5 / 1248.83 | 1.297×
|
| 1/32/32 | 32768 | causal | 2760.5 / 1593.35 | 3289.5 / 1336.97 |
1.192× |
| 32/8/8 | 256 | dense | 13 / 341.98 | 15 / 295.50 | 1.157× |
| 32/8/8 | 256 | causal | 15 / 146.99 | 15 / 139.49 | 1.054× |
| 32/8/8 | 2048 | dense | 213 / 1290.18 | 268.5 / 1022.58 | 1.262× |
| 32/8/8 | 2048 | causal | 148.5 / 927.29 | 175.5 / 782.73 | 1.185× |
| 32/8/8 | 8192 | dense | 2637 / 1668.03 | 3477 / 1264.95 | 1.319× |
| 32/8/8 | 8192 | causal | 1474 / 1492.24 | 1785.5 / 1231.93 | 1.211× |
| 32/8/8 | 16384 | dense | 10304 / 1707.32 | 13556 / 1297.74 | 1.316× |
| 32/8/8 | 16384 | causal | 5525 / 1592.06 | 6652.5 / 1322.32 | 1.204× |
| 32/8/8 | 32768 | dense | 40970 / 1717.57 | 53682 / 1310.85 | 1.310× |
| 32/8/8 | 32768 | causal | 21349 / 1648.11 | 25695 / 1369.36 | 1.204× |
| 1/4/4 | 256 | dense | 6 / 11.62 | 8 / 8.68 | 1.339× |
| 1/4/4 | 256 | causal | 8 / 4.21 | 9 / 3.75 | 1.120× |
| 1/4/4 | 2048 | dense | 18 / 233.25 | 30 / 141.31 | 1.651× |
| 1/4/4 | 2048 | causal | 21 / 104.25 | 28 / 76.84 | 1.357× |
| 1/4/4 | 8192 | dense | 68.5 / 1005.61 | 114 / 604.32 | 1.664× |
| 1/4/4 | 8192 | causal | 67.5 / 508.64 | 97 / 353.75 | 1.438× |
| 1/4/4 | 16384 | dense | 197 / 1394.20 | 272.5 / 1008.87 | 1.382× |
| 1/4/4 | 16384 | causal | 144 / 952.23 | 200.5 / 685.69 | 1.389× |
| 1/4/4 | 32768 | dense | 681.5 / 1613.03 | 1019.5 / 1078.21 | 1.496× |
| 1/4/4 | 32768 | causal | 390 / 1411.11 | 442.5 / 1242.26 | 1.136× |
| 1/3/3 | 256 | dense | 6 / 8.69 | 8 / 6.55 | 1.327× |
| 1/3/3 | 256 | causal | 8 / 3.17 | 9 / 2.82 | 1.122× |
| 1/3/3 | 2048 | dense | 18 / 174.82 | 30 / 106.75 | 1.638× |
| 1/3/3 | 2048 | causal | 21 / 77.13 | 28 / 57.34 | 1.345× |
| 1/3/3 | 8192 | dense | 64 / 802.98 | 112 / 461.57 | 1.740× |
| 1/3/3 | 8192 | causal | 66.5 / 388.10 | 96 / 267.89 | 1.449× |
| 1/3/3 | 16384 | dense | 178.5 / 1155.61 | 263.5 / 783.27 | 1.475× |
| 1/3/3 | 16384 | causal | 137.5 / 749.85 | 194 / 531.33 | 1.411× |
| 1/3/3 | 32768 | dense | 528.5 / 1560.32 | 667 / 1236.47 | 1.262× |
| 1/3/3 | 32768 | causal | 316 / 1305.88 | 412 / 1000.44 | 1.305× |
| 1/2/2 | 256 | dense | 6 / 5.72 | 8 / 4.40 | 1.301× |
| 1/2/2 | 256 | causal | 8 / 2.21 | 9 / 1.97 | 1.124× |
| 1/2/2 | 2048 | dense | 17.5 / 123.09 | 29 / 73.66 | 1.671× |
| 1/2/2 | 2048 | causal | 20 / 54.11 | 27 / 39.39 | 1.374× |
| 1/2/2 | 8192 | dense | 61 / 561.33 | 109 / 316.07 | 1.776× |
| 1/2/2 | 8192 | causal | 65 / 264.22 | 95 / 181.44 | 1.456× |
| 1/2/2 | 16384 | dense | 127 / 1079.66 | 220 / 625.33 | 1.727× |
| 1/2/2 | 16384 | causal | 128 / 537.47 | 186.5 / 368.42 | 1.459× |
| 1/2/2 | 32768 | dense | 375.5 / 1463.87 | 533.5 / 1030.62 | 1.420× |
| 1/2/2 | 32768 | causal | 275 / 999.47 | 386 / 712.56 | 1.403× |
| 1/1/1 | 256 | dense | 5 / 3.51 | 6 / 2.59 | 1.355× |
| 1/1/1 | 256 | causal | 7 / 1.19 | 8 / 1.06 | 1.128× |
| 1/1/1 | 2048 | dense | 16 / 66.02 | 28 / 37.91 | 1.741× |
| 1/1/1 | 2048 | causal | 19 / 27.58 | 27 / 20.00 | 1.379× |
| 1/1/1 | 8192 | dense | 58 / 295.06 | 106 / 162.21 | 1.819× |
| 1/1/1 | 8192 | causal | 63 / 136.24 | 92.5 / 92.93 | 1.466× |
| 1/1/1 | 16384 | dense | 117 / 585.65 | 211 / 325.26 | 1.801× |
| 1/1/1 | 16384 | causal | 124 / 278.17 | 182 / 188.73 | 1.474× |
| 1/1/1 | 32768 | dense | 244 / 1128.26 | 427.5 / 642.97 | 1.755× |
| 1/1/1 | 32768 | causal | 247.5 / 556.17 | 365 / 376.88 | 1.476× |
| 32/1/1 | 256 | dense | 5 / 109.78 | 7 / 82.30 | 1.334× |
| 32/1/1 | 256 | causal | 7 / 37.70 | 8 / 33.35 | 1.130× |
| 32/1/1 | 2048 | dense | 27.5 / 1238.43 | 39 / 876.60 | 1.413× |
| 32/1/1 | 2048 | causal | 28 / 609.23 | 36 / 474.27 | 1.285× |
| 32/1/1 | 8192 | dense | 349 / 1575.91 | 460 / 1194.72 | 1.319× |
| 32/1/1 | 8192 | causal | 233 / 1180.41 | 299 / 920.61 | 1.282× |
| 32/1/1 | 16384 | dense | 1308.5 / 1680.83 | 1748.5 / 1258.01 | 1.336×
|
| 32/1/1 | 16384 | causal | 771 / 1426.37 | 953.5 / 1153.07 | 1.237× |
| 32/1/1 | 32768 | dense | 5088.5 / 1728.62 | 6778 / 1297.69 | 1.332× |
| 32/1/1 | 32768 | causal | 2802.5 / 1569.31 | 3492.5 / 1259.32 | 1.246×
|
</details>
<details><summary>BF16 Q/K=128, V=128: all 110 workloads</summary>
| B/H/Hkv | Sq = Sk | Mask | qr_tdm_sched (Introduced by this PR)<br>µs
/ TFLOP/s | qr_tdm<br>µs / TFLOP/s | qr_tdm_sched / qr_tdm |
| --- | ---: | --- | ---: | ---: | ---: |
| 1/8/8 | 256 | dense | 5 / 49.45 | 7 / 40.56 | 1.219× |
| 1/8/8 | 256 | causal | 7 / 19.76 | 10 / 14.05 | 1.406× |
| 1/8/8 | 2048 | dense | 19.5 / 883.62 | 26 / 652.13 | 1.355× |
| 1/8/8 | 2048 | causal | 21 / 411.26 | 36 / 238.12 | 1.727× |
| 1/8/8 | 8192 | dense | 139.5 / 1969.99 | 162.5 / 1694.40 | 1.163× |
| 1/8/8 | 8192 | causal | 91.5 / 1500.30 | 141 / 974.51 | 1.540× |
| 1/8/8 | 16384 | dense | 510 / 2154.63 | 589 / 1865.78 | 1.155× |
| 1/8/8 | 16384 | causal | 281.5 / 1951.12 | 383.5 / 1434.64 | 1.360× |
| 1/8/8 | 32768 | dense | 1956 / 2248.86 | 2269.5 / 1938.01 | 1.160× |
| 1/8/8 | 32768 | causal | 1020.5 / 2154.98 | 1377.5 / 1596.22 | 1.350×
|
| 1/64/8 | 256 | dense | 6 / 372.57 | 7 / 306.49 | 1.216× |
| 1/64/8 | 256 | causal | 7 / 153.30 | 10 / 108.67 | 1.411× |
| 1/64/8 | 2048 | dense | 78.5 / 1752.28 | 91 / 1513.26 | 1.158× |
| 1/64/8 | 2048 | causal | 50 / 1369.23 | 66 / 1035.78 | 1.322× |
| 1/64/8 | 8192 | dense | 1009 / 2180.05 | 1164 / 1888.82 | 1.154× |
| 1/64/8 | 8192 | causal | 560 / 1964.26 | 736 / 1493.53 | 1.315× |
| 1/64/8 | 16384 | dense | 3924 / 2241.60 | 4520.5 / 1945.76 | 1.152× |
| 1/64/8 | 16384 | causal | 2078 / 2116.28 | 2770 / 1587.78 | 1.333× |
| 1/64/8 | 32768 | dense | 16023 / 2195.86 | 18352 / 1917.18 | 1.145× |
| 1/64/8 | 32768 | causal | 8186.5 / 2148.93 | 10847 / 1621.95 | 1.325×
|
| 1/64/64 | 256 | dense | 6 / 368.88 | 7 / 306.31 | 1.204× |
| 1/64/64 | 256 | causal | 7 / 153.55 | 10 / 108.72 | 1.412× |
| 1/64/64 | 2048 | dense | 80.5 / 1708.78 | 92.5 / 1486.65 | 1.149× |
| 1/64/64 | 2048 | causal | 51 / 1351.80 | 67.5 / 1016.80 | 1.329× |
| 1/64/64 | 8192 | dense | 1096.5 / 2005.53 | 1267.5 / 1734.80 | 1.156×
|
| 1/64/64 | 8192 | causal | 584 / 1883.05 | 751 / 1464.62 | 1.286× |
| 1/64/64 | 16384 | dense | 4300 / 2045.51 | 5004.5 / 1757.77 | 1.164× |
| 1/64/64 | 16384 | causal | 2245.5 / 1958.67 | 2935 / 1498.59 | 1.307×
|
| 1/64/64 | 32768 | dense | 17102 / 2057.36 | 19826.5 / 1774.61 | 1.159×
|
| 1/64/64 | 32768 | causal | 8833 / 1991.76 | 11570.5 / 1520.44 | 1.310×
|
| 1/12/12 | 256 | dense | 6 / 68.67 | 7 / 57.87 | 1.187× |
| 1/12/12 | 256 | causal | 7 / 28.33 | 10 / 20.48 | 1.383× |
| 1/12/12 | 2048 | dense | 21 / 1222.07 | 28 / 923.80 | 1.323× |
| 1/12/12 | 2048 | causal | 22 / 592.16 | 37 / 348.87 | 1.697× |
| 1/12/12 | 8192 | dense | 210.5 / 1955.64 | 247.5 / 1664.91 | 1.175× |
| 1/12/12 | 8192 | causal | 125.5 / 1637.74 | 157.5 / 1311.42 | 1.249× |
| 1/12/12 | 16384 | dense | 775.5 / 2125.59 | 894.5 / 1844.38 | 1.152× |
| 1/12/12 | 16384 | causal | 420 / 1964.32 | 551.5 / 1495.67 | 1.313× |
| 1/12/12 | 32768 | dense | 3062 / 2154.51 | 3586 / 1839.65 | 1.171× |
| 1/12/12 | 32768 | causal | 1585.5 / 2079.95 | 2073 / 1591.16 | 1.307×
|
| 1/32/32 | 256 | dense | 6 / 192.74 | 7 / 157.75 | 1.222× |
| 1/32/32 | 256 | causal | 7 / 78.35 | 10 / 55.27 | 1.418× |
| 1/32/32 | 2048 | dense | 39.5 / 1733.26 | 46.5 / 1474.77 | 1.175× |
| 1/32/32 | 2048 | causal | 28 / 1225.74 | 41 / 838.59 | 1.462× |
| 1/32/32 | 8192 | dense | 520 / 2115.86 | 597 / 1842.40 | 1.148× |
| 1/32/32 | 8192 | causal | 287.5 / 1912.56 | 380 / 1445.95 | 1.323× |
| 1/32/32 | 16384 | dense | 2052 / 2143.53 | 2386 / 1843.25 | 1.163× |
| 1/32/32 | 16384 | causal | 1068.5 / 2058.40 | 1401.5 / 1569.14 |
1.312× |
| 1/32/32 | 32768 | dense | 8160 / 2155.86 | 9495 / 1852.79 | 1.164× |
| 1/32/32 | 32768 | causal | 4212.5 / 2088.30 | 5575 / 1577.80 | 1.324×
|
| 32/8/8 | 256 | dense | 9 / 948.69 | 11 / 807.17 | 1.175× |
| 32/8/8 | 256 | causal | 10 / 425.35 | 13 / 329.72 | 1.290× |
| 32/8/8 | 2048 | dense | 301.5 / 1822.08 | 354.5 / 1550.15 | 1.175× |
| 32/8/8 | 2048 | causal | 205 / 1339.68 | 264 / 1040.47 | 1.288× |
| 32/8/8 | 8192 | dense | 4004.5 / 2196.34 | 4650 / 1891.64 | 1.161× |
| 32/8/8 | 8192 | causal | 2218 / 1983.39 | 2948 / 1492.08 | 1.329× |
| 32/8/8 | 16384 | dense | 15650 / 2248.20 | 18080 / 1946.01 | 1.155× |
| 32/8/8 | 16384 | causal | 8324 / 2113.67 | 11086.5 / 1586.88 | 1.332×
|
| 32/8/8 | 32768 | dense | 62354 / 2257.07 | 71921.5 / 1956.83 | 1.153×
|
| 32/8/8 | 32768 | causal | 32432.5 / 2169.79 | 43190 / 1629.34 | 1.332×
|
| 1/4/4 | 256 | dense | 6 / 23.94 | 7 / 19.89 | 1.203× |
| 1/4/4 | 256 | causal | 7 / 9.84 | 10 / 6.97 | 1.410× |
| 1/4/4 | 2048 | dense | 18.5 / 464.11 | 25 / 338.91 | 1.369× |
| 1/4/4 | 2048 | causal | 20 / 211.88 | 36 / 120.81 | 1.754× |
| 1/4/4 | 8192 | dense | 80 / 1713.53 | 102 / 1345.16 | 1.274× |
| 1/4/4 | 8192 | causal | 71 / 961.78 | 129 / 531.40 | 1.810× |
| 1/4/4 | 16384 | dense | 268.5 / 2047.22 | 310.5 / 1771.69 | 1.156× |
| 1/4/4 | 16384 | causal | 172 / 1599.62 | 271.5 / 1013.66 | 1.578× |
| 1/4/4 | 32768 | dense | 996.5 / 2206.86 | 1157 / 1900.56 | 1.161× |
| 1/4/4 | 32768 | causal | 539.5 / 2037.87 | 746.5 / 1472.63 | 1.384× |
| 1/3/3 | 256 | dense | 6 / 17.51 | 7 / 14.69 | 1.192× |
| 1/3/3 | 256 | causal | 7 / 7.22 | 10 / 5.21 | 1.385× |
| 1/3/3 | 2048 | dense | 18 / 349.55 | 25 / 257.00 | 1.360× |
| 1/3/3 | 2048 | causal | 20 / 159.11 | 35 / 91.00 | 1.748× |
| 1/3/3 | 8192 | dense | 72 / 1433.12 | 96.5 / 1068.22 | 1.342× |
| 1/3/3 | 8192 | causal | 70 / 737.01 | 127.5 / 404.37 | 1.823× |
| 1/3/3 | 16384 | dense | 235 / 1756.28 | 277 / 1488.26 | 1.180× |
| 1/3/3 | 16384 | causal | 155.5 / 1327.24 | 262.5 / 785.17 | 1.690× |
| 1/3/3 | 32768 | dense | 761 / 2167.19 | 902.5 / 1826.72 | 1.186× |
| 1/3/3 | 32768 | causal | 423 / 1947.91 | 569.5 / 1448.76 | 1.345× |
| 1/2/2 | 256 | dense | 5 / 12.27 | 7 / 10.12 | 1.212× |
| 1/2/2 | 256 | causal | 7 / 4.96 | 10 / 3.50 | 1.414× |
| 1/2/2 | 2048 | dense | 18 / 238.93 | 25 / 174.05 | 1.373× |
| 1/2/2 | 2048 | causal | 20 / 108.42 | 35 / 61.23 | 1.771× |
| 1/2/2 | 8192 | dense | 66 / 1041.96 | 92 / 748.55 | 1.392× |
| 1/2/2 | 8192 | causal | 67.5 / 509.00 | 125.5 / 273.87 | 1.859× |
| 1/2/2 | 16384 | dense | 151.5 / 1813.10 | 193.5 / 1419.58 | 1.277× |
| 1/2/2 | 16384 | causal | 136 / 1010.78 | 250 / 549.94 | 1.838× |
| 1/2/2 | 32768 | dense | 512 / 2148.24 | 595 / 1848.09 | 1.162× |
| 1/2/2 | 32768 | causal | 326.5 / 1685.10 | 527.5 / 1041.99 | 1.617× |
| 1/1/1 | 256 | dense | 5 / 6.31 | 7 / 5.12 | 1.231× |
| 1/1/1 | 256 | causal | 7 / 2.51 | 10 / 1.77 | 1.421× |
| 1/1/1 | 2048 | dense | 17 / 123.37 | 24 / 89.28 | 1.382× |
| 1/1/1 | 2048 | causal | 20 / 54.98 | 35 / 30.91 | 1.778× |
| 1/1/1 | 8192 | dense | 62 / 553.37 | 88 / 389.71 | 1.420× |
| 1/1/1 | 8192 | causal | 65 / 262.91 | 123 / 139.94 | 1.879× |
| 1/1/1 | 16384 | dense | 126 / 1086.95 | 178 / 773.91 | 1.404× |
| 1/1/1 | 16384 | causal | 129 / 531.38 | 243 / 282.38 | 1.882× |
| 1/1/1 | 32768 | dense | 287.5 / 1911.71 | 372.5 / 1475.77 | 1.295× |
| 1/1/1 | 32768 | causal | 264.5 / 1038.78 | 487.5 / 563.97 | 1.842× |
| 32/1/1 | 256 | dense | 6 / 193.55 | 7 / 159.84 | 1.211× |
| 32/1/1 | 256 | causal | 7 / 78.92 | 10 / 55.59 | 1.420× |
| 32/1/1 | 2048 | dense | 39 / 1755.37 | 47 / 1456.70 | 1.205× |
| 32/1/1 | 2048 | causal | 35 / 980.40 | 51 / 675.58 | 1.451× |
| 32/1/1 | 8192 | dense | 524 / 2098.01 | 611 / 1798.82 | 1.166× |
| 32/1/1 | 8192 | causal | 334.5 / 1643.07 | 470 / 1170.80 | 1.403× |
| 32/1/1 | 16384 | dense | 1975 / 2226.80 | 2305.5 / 1907.85 | 1.167× |
| 32/1/1 | 16384 | causal | 1132 / 1942.68 | 1580 / 1391.70 | 1.396× |
| 32/1/1 | 32768 | dense | 7753.5 / 2268.94 | 9014.5 / 1951.53 | 1.163×
|
| 32/1/1 | 32768 | causal | 4220.5 / 2084.35 | 5883.5 / 1495.04 | 1.394×
|
</details>
<details><summary>BF16 Q/K=192, V=128: all 110 workloads</summary>
| B/H/Hkv | Sq = Sk | Mask | qr_tdm_sched (Introduced by this PR)<br>µs
/ TFLOP/s | qr_tdm<br>µs / TFLOP/s | qr_tdm_sched / qr_tdm |
| --- | ---: | --- | ---: | ---: | ---: |
| 1/8/8 | 256 | dense | 7 / 48.02 | 8 / 43.73 | 1.098× |
| 1/8/8 | 256 | causal | 9 / 19.61 | 11 / 16.02 | 1.224× |
| 1/8/8 | 2048 | dense | 23 / 925.25 | 35 / 617.25 | 1.499× |
| 1/8/8 | 2048 | causal | 26 / 416.77 | 44 / 242.61 | 1.718× |
| 1/8/8 | 8192 | dense | 182 / 1888.17 | 201.5 / 1707.81 | 1.106× |
| 1/8/8 | 8192 | causal | 133 / 1292.51 | 177 / 971.64 | 1.330× |
| 1/8/8 | 16384 | dense | 665.5 / 2064.98 | 745 / 1845.91 | 1.119× |
| 1/8/8 | 16384 | causal | 392 / 1752.87 | 452.5 / 1518.08 | 1.155× |
| 1/8/8 | 32768 | dense | 2587 / 2125.11 | 2912 / 1888.20 | 1.125× |
| 1/8/8 | 32768 | causal | 1408 / 1952.06 | 1712 / 1605.57 | 1.216× |
| 1/64/8 | 256 | dense | 7 / 369.48 | 8 / 333.00 | 1.110× |
| 1/64/8 | 256 | causal | 9 / 151.57 | 11 / 124.22 | 1.220× |
| 1/64/8 | 2048 | dense | 102 / 1678.51 | 112.5 / 1521.56 | 1.103× |
| 1/64/8 | 2048 | causal | 69 / 1250.12 | 77 / 1121.22 | 1.115× |
| 1/64/8 | 8192 | dense | 1322 / 2078.82 | 1481 / 1856.32 | 1.120× |
| 1/64/8 | 8192 | causal | 757 / 1816.07 | 891.5 / 1541.56 | 1.178× |
| 1/64/8 | 16384 | dense | 5171 / 2126.24 | 5815.5 / 1890.62 | 1.125× |
| 1/64/8 | 16384 | causal | 2859 / 1923.03 | 3381 / 1626.04 | 1.183× |
| 1/64/8 | 32768 | dense | 20964 / 2097.89 | 23605.5 / 1863.16 | 1.126×
|
| 1/64/8 | 32768 | causal | 11295.5 / 1946.87 | 13384.5 / 1642.97 |
1.185× |
| 1/64/64 | 256 | dense | 7 / 369.46 | 8 / 333.52 | 1.108× |
| 1/64/64 | 256 | causal | 9 / 152.22 | 11 / 124.11 | 1.226× |
| 1/64/64 | 2048 | dense | 105.5 / 1627.51 | 114.5 / 1499.87 | 1.085× |
| 1/64/64 | 2048 | causal | 70.5 / 1212.61 | 78.5 / 1094.97 | 1.107× |
| 1/64/64 | 8192 | dense | 1458 / 1885.32 | 1605.5 / 1712.16 | 1.101× |
| 1/64/64 | 8192 | causal | 798 / 1721.78 | 918 / 1496.87 | 1.150× |
| 1/64/64 | 16384 | dense | 5716.5 / 1923.35 | 6349 / 1731.78 | 1.111× |
| 1/64/64 | 16384 | causal | 3107.5 / 1769.16 | 3614 / 1521.28 | 1.163×
|
| 1/64/64 | 32768 | dense | 22649 / 1941.85 | 25245 / 1742.13 | 1.115× |
| 1/64/64 | 32768 | causal | 12231.5 / 1797.93 | 14312.5 / 1536.50 |
1.170× |
| 1/12/12 | 256 | dense | 7 / 68.65 | 8 / 63.35 | 1.084× |
| 1/12/12 | 256 | causal | 9 / 28.48 | 11 / 23.48 | 1.213× |
| 1/12/12 | 2048 | dense | 25.5 / 1273.39 | 37 / 877.98 | 1.450× |
| 1/12/12 | 2048 | causal | 27 / 594.66 | 46 / 353.45 | 1.682× |
| 1/12/12 | 8192 | dense | 275 / 1872.67 | 318 / 1619.62 | 1.156× |
| 1/12/12 | 8192 | causal | 161 / 1602.28 | 195 / 1324.34 | 1.210× |
| 1/12/12 | 16384 | dense | 1020 / 2021.97 | 1133.5 / 1819.31 | 1.111× |
| 1/12/12 | 16384 | causal | 567 / 1817.85 | 676 / 1525.85 | 1.191× |
| 1/12/12 | 32768 | dense | 4026.5 / 2047.83 | 4507.5 / 1829.59 | 1.119×
|
| 1/12/12 | 32768 | causal | 2155.5 / 1913.13 | 2552.5 / 1615.36 |
1.184× |
| 1/32/32 | 256 | dense | 7 / 187.38 | 8 / 168.97 | 1.109× |
| 1/32/32 | 256 | causal | 9 / 77.41 | 11 / 63.06 | 1.228× |
| 1/32/32 | 2048 | dense | 53 / 1621.17 | 58 / 1485.64 | 1.091× |
| 1/32/32 | 2048 | causal | 40 / 1071.55 | 51 / 844.96 | 1.268× |
| 1/32/32 | 8192 | dense | 687.5 / 1999.30 | 762 / 1803.20 | 1.109× |
| 1/32/32 | 8192 | causal | 391.5 / 1755.41 | 466.5 / 1473.47 | 1.191× |
| 1/32/32 | 16384 | dense | 2731 / 2013.10 | 3041 / 1807.96 | 1.113× |
| 1/32/32 | 16384 | causal | 1467.5 / 1873.78 | 1720.5 / 1597.70 |
1.173× |
| 1/32/32 | 32768 | dense | 10799.5 / 2036.24 | 12074 / 1821.28 | 1.118×
|
| 1/32/32 | 32768 | causal | 5818.5 / 1889.83 | 6846.5 / 1606.04 |
1.177× |
| 32/8/8 | 256 | dense | 12 / 910.68 | 12 / 880.00 | 1.035× |
| 32/8/8 | 256 | causal | 13 / 412.21 | 15 / 369.06 | 1.117× |
| 32/8/8 | 2048 | dense | 395.5 / 1737.20 | 434.5 / 1580.08 | 1.099× |
| 32/8/8 | 2048 | causal | 273 / 1260.08 | 313 / 1098.72 | 1.147× |
| 32/8/8 | 8192 | dense | 5265 / 2088.28 | 5900.5 / 1863.57 | 1.121× |
| 32/8/8 | 8192 | causal | 3038.5 / 1809.44 | 3584 / 1534.17 | 1.179× |
| 32/8/8 | 16384 | dense | 20551.5 / 2139.96 | 23140 / 1900.63 | 1.126×
|
| 32/8/8 | 16384 | causal | 11441.5 / 1922.14 | 13574 / 1620.11 | 1.186×
|
| 32/8/8 | 32768 | dense | 82542.5 / 2131.29 | 92859.5 / 1894.49 |
1.125× |
| 32/8/8 | 32768 | causal | 44795 / 1963.68 | 53240 / 1652.22 | 1.189× |
| 1/4/4 | 256 | dense | 7 / 23.82 | 8 / 21.95 | 1.085× |
| 1/4/4 | 256 | causal | 9 / 9.81 | 10.5 / 8.02 | 1.224× |
| 1/4/4 | 2048 | dense | 22 / 487.02 | 34 / 320.13 | 1.521× |
| 1/4/4 | 2048 | causal | 25 / 213.91 | 44 / 123.07 | 1.738× |
| 1/4/4 | 8192 | dense | 96 / 1791.19 | 136 / 1266.14 | 1.415× |
| 1/4/4 | 8192 | causal | 90 / 958.32 | 163 / 525.59 | 1.823× |
| 1/4/4 | 16384 | dense | 346 / 1984.35 | 386.5 / 1779.49 | 1.115× |
| 1/4/4 | 16384 | causal | 254.5 / 1350.47 | 342.5 / 1003.10 | 1.346× |
| 1/4/4 | 32768 | dense | 1310 / 2098.09 | 1470 / 1869.88 | 1.122× |
| 1/4/4 | 32768 | causal | 754 / 1823.20 | 878 / 1564.96 | 1.165× |
| 1/3/3 | 256 | dense | 7 / 17.62 | 8 / 16.11 | 1.094× |
| 1/3/3 | 256 | causal | 9 / 7.25 | 11 / 5.95 | 1.219× |
| 1/3/3 | 2048 | dense | 22 / 366.83 | 33 / 240.81 | 1.523× |
| 1/3/3 | 2048 | causal | 25 / 160.38 | 43 / 92.67 | 1.731× |
| 1/3/3 | 8192 | dense | 86 / 1496.30 | 130 / 988.91 | 1.513× |
| 1/3/3 | 8192 | causal | 87 / 739.11 | 161 / 400.90 | 1.844× |
| 1/3/3 | 16384 | dense | 307.5 / 1674.18 | 347.5 / 1482.52 | 1.129× |
| 1/3/3 | 16384 | causal | 238 / 1082.90 | 331 / 778.76 | 1.391× |
| 1/3/3 | 32768 | dense | 1007 / 2047.31 | 1181 / 1745.87 | 1.173× |
| 1/3/3 | 32768 | causal | 567.5 / 1816.11 | 715 / 1441.41 | 1.260× |
| 1/2/2 | 256 | dense | 7 / 12.18 | 8 / 11.12 | 1.094× |
| 1/2/2 | 256 | causal | 9 / 4.94 | 10 / 4.03 | 1.225× |
| 1/2/2 | 2048 | dense | 21 / 252.14 | 33 / 163.94 | 1.538× |
| 1/2/2 | 2048 | causal | 25 / 108.95 | 43 / 62.28 | 1.749× |
| 1/2/2 | 8192 | dense | 78 / 1099.51 | 125 / 684.67 | 1.606× |
| 1/2/2 | 8192 | causal | 85 / 504.75 | 159 / 269.97 | 1.870× |
| 1/2/2 | 16384 | dense | 179.5 / 1914.47 | 260 / 1321.57 | 1.449× |
| 1/2/2 | 16384 | causal | 171.5 / 1001.16 | 317.5 / 541.13 | 1.850× |
| 1/2/2 | 32768 | dense | 667.5 / 2059.55 | 744.5 / 1846.57 | 1.115× |
| 1/2/2 | 32768 | causal | 492.5 / 1395.38 | 670.5 / 1024.41 | 1.362× |
| 1/1/1 | 256 | dense | 7 / 6.20 | 7.5 / 5.61 | 1.105× |
| 1/1/1 | 256 | causal | 8 / 2.50 | 10 / 2.03 | 1.232× |
| 1/1/1 | 2048 | dense | 21 / 129.12 | 32 / 83.85 | 1.540× |
| 1/1/1 | 2048 | causal | 24 / 55.22 | 43 / 31.50 | 1.753× |
| 1/1/1 | 8192 | dense | 73 / 584.67 | 121 / 355.15 | 1.646× |
| 1/1/1 | 8192 | causal | 83 / 259.93 | 155 / 138.75 | 1.873× |
| 1/1/1 | 16384 | dense | 149.5 / 1148.66 | 243.5 / 705.19 | 1.629× |
| 1/1/1 | 16384 | causal | 163 / 526.39 | 307.5 / 279.66 | 1.882× |
| 1/1/1 | 32768 | dense | 341.5 / 2012.12 | 507.5 / 1354.29 | 1.486× |
| 1/1/1 | 32768 | causal | 331.5 / 1036.67 | 622.5 / 551.94 | 1.878× |
| 32/1/1 | 256 | dense | 7 / 191.09 | 8 / 175.12 | 1.091× |
| 32/1/1 | 256 | causal | 9 / 78.16 | 11 / 63.94 | 1.222× |
| 32/1/1 | 2048 | dense | 53 / 1614.35 | 58.5 / 1474.45 | 1.095× |
| 32/1/1 | 2048 | causal | 50.5 / 845.21 | 58.5 / 734.68 | 1.150× |
| 32/1/1 | 8192 | dense | 681 / 2018.51 | 760.5 / 1808.03 | 1.116× |
| 32/1/1 | 8192 | causal | 459 / 1496.40 | 565.5 / 1215.26 | 1.231× |
| 32/1/1 | 16384 | dense | 2588.5 / 2124.03 | 2905 / 1892.45 | 1.122× |
| 32/1/1 | 16384 | causal | 1559.5 / 1762.61 | 1917 / 1434.18 | 1.229× |
| 32/1/1 | 32768 | dense | 10155 / 2165.40 | 11427 / 1924.38 | 1.125× |
| 32/1/1 | 32768 | causal | 5901.5 / 1863.09 | 7211 / 1524.86 | 1.222× |
</details>
## Clang 24 performance disclosure
A separate matched Clang 24 study found slower `qr_tdm_sched`
performance in some D64/D128 causal configurations. That study used a
different timing harness, so its quantitative results are excluded from
this description to keep every reported performance number sourced from
`tile_example_fmha_fwd`. A matched Clang 24 example run is still needed
before assigning a case count or slowdown magnitude to this benchmark.
The dispatch intentionally keeps `qr_tdm_sched` priority for its
supported configurations; the Clang 23 results above should not be
extrapolated to Clang 24.
## Toolchain and build environment
| Component | Clang 23 example comparison | Clang 24 matched validation
/ separate diagnostic |
| --- | --- | --- |
| Device compiler | AMD Clang `23.0.0git`; LLVM
`2a67d831a09cc538194bd595279ee91258570c10`; `clang++` SHA256
`b552c9fa9e9158688cd4944a4aa899dfedc7426b700097ef9338f4476dc6f6f6` | AMD
Clang `24.0.0git`; LLVM `4f43f4746ede4cc49ab4129f473dbd54e0d9db4f`;
`clang++` SHA256
`1cfe274fcb8f92e5ccb423173e207dda6036e4a9efcfcdf29193049849d33979` |
| `hipcc -v` / HIP | Container `hipcc` reports Clang 23 and HIP 7.15.0;
`hipconfig --version` reports `7.15.0-0000000` | The pinned Clang 24
build invokes `clang++` directly, with `--hip-path=/opt/rocm`; container
HIP headers/runtime are 7.15.0. Container `hipcc` was not used for this
build. |
| LLVM linker and device libraries | AMD LLD 23.0.0;
`/opt/rocm/amdgcn/bitcode` | AMD LLD 24.0.0; pinned Clang 24
device-library bundle |
| Target and common build options | `gfx1250`, `-O3`, `-DNDEBUG`,
`CK_TILE_FMHA_FWD_FAST_EXP2=1`; scheduled objects use `-mllvm
-amdgpu-expert-scheduling-mode` | Same target and relevant
optimization/math settings; scheduled objects use expert scheduling |
| Source | This PR at `4688449fa9b` vs upstream `qr_tdm` at
`9e3337709130` (includes PR #12347) | Matched final-candidate source and
upstream `qr_tdm`; numerical results omitted here because they were not
produced by `tile_example_fmha_fwd` |
## Validation
- Current-candidate FP16/BF16 generated-dispatch correctness passed
924/924 configurations across D64/D128/D192; sequences through 4096
checked all rows, while longer sequences used deterministic sampled rows
plus whole-output finite and LSE scans.
- The generated-dispatch GTests passed 480/480; finite virtual-sink
fallback passed 96/96 and selected `qr_tdm`; direct layout, softmax,
fragment, and schedule GTests passed 79/79. Python codegen/dispatcher
tests passed 40/40.
- All 72 generated scheduled objects compiled under the pinned Clang 24
toolchain. Local benchmark scripts and codegen filters are excluded from
the PR.
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.
…fx942 (#13292) JIRA ID: https://amd-hub.atlassian.net/browse/AIROCBLAS-1268 JIRA ID: https://amd-hub.atlassian.net/browse/AIROCBLAS-1267 JIRA ID: https://amd-hub.atlassian.net/browse/AIHPBLAS-3019 ## Motivation rocBLAS currently defaults to the Tensile Solution Library for FP16, BF16, and FP64 GEMM on gfx942. As part of the GEMM backend unification initiative, the default backend is flipped to hipBLASLt for these three input types, following the performance parity threshold of within 5% of Tensile or acceptable performance. This extends the two preceding flips — DGEMM on gfx950 (#12424) and DGEMM on the 228-CU MI300A (#12724) — as follows: * FP64 now defaults to hipBLASLt on every gfx942 device (MI300X and MI300A alike), rather than MI300A (228 CUs) only. * FP16 and BF16 now default to hipBLASLt on the 304-CU MI300X only. MI300A (228 CUs) continues to default to Tensile for FP16/BF16 in this PR, but it will be changed by #13045. * ## Technical Details `useHipBLASLt()` (`library/src/tensile_host.cpp`) previously routed only `double` to hipBLASLt on gfx942, and only when `multiProcessorCount == 228` (MI300A). The gfx942 branch for MI300x now applies FP16, BF16, and FP64. No change is needed in `isDefaultHipBLASLtArch()` (`handle.hpp`): `942` was already added there by #12724, and that function only gates eager hipBLASLt handle construction, not per-datatype routing. `docs/conceptual/rocblas-design-notes.rst` is also updated. It previously stated that hipBLASLt is the default "for problems on the gfx12 architecture", which was already stale after #12424 and #12724 flipped FP64 on gfx950 and MI300A. It now carries a table of which architectures default to hipBLASLt and for which data types. ## Test Plan 1. Confirm backend routing per datatype on gfx942 with `ROCBLAS_LAYER=8`. 2. Run the GEMM gtest suite on MI300X for regressions. ## Test Result ### Backend routing (MI300X, gfx942, `ROCBLAS_LAYER=8`) | Input type | Backend selected | |---|---| | `f16_r` | `rocblas_gemm_hipblaslt_backend` | | `bf16_r` (via `gemm_ex`, compute `f32_r`) | `rocblas_gemm_hipblaslt_backend` | | `f64_r` | `rocblas_gemm_hipblaslt_backend` | | `f32_r` | `rocblas_gemm_tensile_backend` (unchanged) | ### gtest `rocblas-test --gtest_filter='*hipblaslt*:*non_hipblaslt*-*known_bug*'`: | PASS | FAIL | SKIP | total | TEST_EXIT | |---|---|---|---|---| | 2,208 | 0 | 0 | 2,208 | 0 | `./rocblas-test --gtest_filter=*gemm*f64_r*-*known_bug*`: | PASS | FAIL | SKIP | total | TEST_EXIT | |---|---|---|---|---| | 163,672 | 0 | 44 | 163,716 | 0 | ` ./rocblas-test --gtest_filter=*gemm*f16_r*:*gemm*bf16_r*-*known_bug*` | PASS | FAIL | SKIP | total | TEST_EXIT | |---|---|---|---|---| | 497,739 | 19 | 192 | 497,950 | 0 | Those failed tests are all stress_gemm with ILP tests, which requires appropriate ILP library link, so after I rebuilt the library, I reran those failed tests as below, and all the tests PASSED. `./rocblas-test --gtest_filter=_/gemm.*stress*ILP64*f16_r*:_/gemm_ex.*stress*ILP64*f16_r*:_/gemm_ex.*stress*ILP64*bf16_r*` | PASS | FAIL | SKIP | total | TEST_EXIT | |---|---|---|---|---| | 23 | 0 | 0 | 23 | 0 | ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
JIRA ID : AIHPBLAS-5103
## Motivation
Add some optimization to enhance the post loop performance.
The gfx1250 narrow-float (half/bf16/f8/bf8) store epilogue had two
avoidable
stalls:
1. The alpha multiply ran as a bulk pass over *all* batch elements
before the
per-element saturation/pack/store loop, so element 0's store could not
issue
until every element's alpha was emitted — delaying the first
buffer_store.
With GroupLoadStore, the up-front alpha of the next batch also overwrote
the
previous batch's store source VGPRs before that store had read them,
which ATT
showed as a ~500-cycle `s_wait_alu depctr_vm_vsrc(0)` at the start of
every
batch after the first.
2. The fp8 (E4M3) pk8 saturation clamped each element with
`v_cmp_class + v_med3 + v_cndmask`, which was expensive in the NonEdge
path.
## Technical Details
Two commits:
**perf(tensilelite): speed up gfx1250 fp8 epilogue saturation and
grouped stores**
- Saturation: add `v_maximumminimum_f32` (rocisa `VMaximumMinimumF32`,
StinkyTofu
opcode/map/logical-IR/unit test, adaptor shim, `HasVMaximumMinimumF32`
cap).
`PackData_FLOAT8` pk8 now emits `v_pk_fma_f32 x*0+x` (inf→NaN) + a
`v_maximumminimum_f32` clamp, falling back to the old sequence without
the cap.
Matches `__HIP_SATFINITE` over all 2^32 inputs except the NaN sign of
+inf;
`s_set_vgpr_msb` in the NonEdge path drops 1621→213.
- GroupLoadStore: emit one `s_wait_alu(va_vdst=0)` before the grouped
stores
instead of one per element when `ExpertSchedulingMode > 0`.
- `incrementToNextRow`: skip the `s_lshl_b32` and use the stride SGPR
directly
for 1-byte BPE and a single row.
**feat(tensilelite): add InterleaveAlpha=2 to fuse alpha into the
per-element pack**
- `InterleaveAlpha` gains value `2`: each element's alpha multiply is
deferred to
just before its own conversion+pack (in the store loop) instead of
scaling all
elements up front. Alpha still lands on the raw accumulator (before
beta*C,
bias, activation, pack), so it is numerically general across dest types
and
routes through the up-front alpha path like `0` — the win is register
lifetime,
letting the first store issue after only element 0's work.
- `ValidParameters`: `InterleaveAlpha` now `[0, 1, 2]`.
- `GlobalWriteBatch`: value `2` takes the up-front alpha path
(comparisons made
explicit — `!= 1` / `== 1` — since `2` is truthy);
`_interleaveAlphaWithPack()`
fires on `InterleaveAlpha == 2` for any narrow-float output.
- `Solution`: matrix-instruction kernels reject only `InterleaveAlpha ==
1`
(`2` is allowed); the ExpertSchedulingMode support gate treats `2` like
`0`.
- `RequiredParameters`: add `InterleaveAlpha` to the kernel-name set.
Value `2`
produces different assembly than `0`, so without it in the name the `0`
and `2`
kernels collide and one is dropped as a duplicate (`Naming._getName`),
leaving a
dangling `.s` and "named symbol not found" at launch.
- Adds
`Tests/common/gemm/gfx12/interleave_alpha2_group_store_gfx1250.yaml`,
which sweeps `NumElementsPerBatchStore [0,16] × GroupLoadStore
[False,True] ×
InterleaveAlpha [0,2]` for bf16 and fp8.
## Test Plan
- tensilelite tox
- hipblaslt-test
## Test Result
pass
## Submission Checklist
- [ ] Look over the contributing guidelines at
https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests.
## Risk level
Low–Medium. Scoped to the gfx1250 narrow-float store epilogue. The
saturation and
grouped-store changes fall back to the prior sequence when the new cap /
`ExpertSchedulingMode` is absent. `InterleaveAlpha=2` is opt-in (default
`0`
behavior is unchanged), but it does add `InterleaveAlpha` to the kernel
name,
which changes generated kernel names — relevant to any externally cached
library
logic keyed by name.
---------
Co-authored-by: Serge Lu <Serge.Lu@amd.com>
## Motivation Fix dgrad conv regressions after switch aot ## Technical Details Depthwise 7x7 stride-1 dgrad — the AOT row-streaming kernel was prefetching whole KH rows, blowing up to 377–512 VGPRs and spilling to scratch. Ported the forward-depthwise tricks (ceil(16/n_cols) row prefetch, i32 offset bounds instead of lane masks, empty-asm on accumulators, FMA loop over ceil(Ho/KH) groups + flush-only tail). Brought it down to ~180 VGPRs and ~2x faster; also added block_w=7 to the dgrad block-width table. Dense g=1 dgrad (direct MFMA path) — added channel caps (_DENSE_MFMA_DGRAD_CHANNELS) so these cases probe with groups=1 instead of failing to compile with block_groups>1; direct now beats igemm (0.117 vs 0.227 in one measured case). Patchify dgrad (stride == f ## Test Plan CI ## Test Result Pending ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/TheRock/blob/main/GOVERNANCE.md#pull-requests. JIRA ID: AICK-2527
## Motivation
Wgrad GEMMs usually have a filter-sized output (few output tiles) and a
very long reduction (N*Ho*Wo). Split-K cuts every tile the same number
of ways, so on these shapes it leaves CUs idle. This PR ports CK Tile's
stream-K to rocKE wgrad. Stream-K balances MAC iterations instead of
tiles:
- Tiles that fill the device run data-parallel.
- Only the remainder is spread across a pool of CTAs.
Stream-K is opt-in. The grouped-convolution dispatcher and its benchmark
are unchanged (they are due for a refactor), so no production path picks
stream-K on its own.
## Technical Details
- **Core IR:** three new leaf ops in both engines: `memref.fence`,
`memref.global_flag_store` and `memref.global_flag_wait_eq`.
- Each lowers to an LLVM fence, or to an atomic store/load with
release/acquire ordering at a chosen sync scope.
- The Python and C++ verifiers check arity, a global i32 flag pointer,
i32 index/value operands, and the allowed scope/ordering values.
- All four lowerers reject a missing or unknown scope/ordering with the
same message.
- **Partitioner:** `StreamKIterPartition` (Python) and
`rocke_streamk_iter_partition_t` (C++) mirror CK's DP + SK split.
- `iters_per_sk_cta` is clamped to at least 1, so an all-DP partition is
safe.
- `flags_bytes(n_waves)` sizes the per-(CTA, wave) flag table.
- The device-side emitters read the partition from runtime values.
- **Wgrad kernel:**
- Launch modes: `streamk = dp_sk | persistent`.
- Reductions: `linear` and `tree` are deterministic and hand off through
per-(CTA, wave) flags. `atomic` (fp32 dW only) and `workspace` (f32
scratch + Stage-2 reduce) are the other two.
- Conv groups are folded into GEMM-M per group, so the fold is exact for
any wg_M. The C++ engine now builds grouped wgrad on the stream-K path.
- **AOT:** the partition is passed as an `sk_*` kernarg block, appended
after `ks_count` so existing kernarg offsets don't change.
- Python side: `conv_abi` and `ConvArgs.streamk_partition`.
- C++ side: `rocke_conv_arg_names_ex`.
- The CTA pool is a launch parameter, so one binary serves any problem
shape and any pool. `streamk_ctas = -1` resolves to the arch default
pool, in both engines.
- **Sweep / benchmark (opt-in):** `benchmark_implicit_gemm_conv.py`
(also reachable through `benchmark_conv_compare.py`) gains `--streamk
{off,dp_sk,persistent,all}` (default `off`), `--streamk-reductions` and
`--streamk-ctas`.
- Stream-K wgrad candidates are added to the JIT sweep, the
`--compile-all` AOT cache (`kernel_sweep.py`), and `--run-from-cache`.
- With stream-K off, the cache identity and kernel labels are unchanged.
- The stream-K launcher allocates the reduction's workspace and zeroes
what each reduction accumulates into. It shares the Stage-2 reduce with
the two-stage path.
- **Constraints:** gfx942 and gfx950 only, MFMA wave64. Both engines
reject the following with identical messages: split-K, two-stage, async
DMA, unroll, cshuffle, chiplet swizzle and group merge.
## Test Plan
- Byte identity across both engines at llvm20, llvm22 and llvm23.
- Differential parity: new wgrad configs 21–30 and negative configs
104–108.
- Representative-IR golden: all 8 `{dp_sk, persistent} × {linear, tree,
atomic, workspace}` combos (dense gfx950 plus grouped gfx942).
- CPU tests:
- partitioner invariants
- validator agreement across both engines, including multi-cause rejects
and the arch gate
- admits-implies-builds matrix
- host plan and default-pool agreement between Python and C++
- stream-K ABI and launch values
- sweep launcher, including all-DP partitions
- dual-engine lowering
- C++ ctests: `rocke_streamk_iter_partition`,
`rocke_wgrad_streamk_host`, `rocke_future_intrinsic_lowering` and
`rocke_ir_serialize_roundtrip`.
- GPU tests:
- numerics for both launch modes × all four reductions on dense, strided
and grouped shapes, including a partial last M tile
- one binary launched on several shapes (AOT)
- 20-launch bit-exact repeat test for linear and tree, with partials and
dW poisoned before each launch
- flag ordering pinned in IR (acquire spin / release publish) and ISA
(cache invalidate / writeback)
- sweep launcher on all-DP partitions (linear and workspace)
## Test Result
- Byte identity GREEN for all families at llvm20, llvm22 and llvm23.
31/31 wgrad configs pass differential parity.
- Representative-IR golden OK at all three flavors after merging
develop.
- CPU tests, C++-binding tests and ctests pass.
- GPU tests pass on gfx950. gfx942 is covered by the IR, golden and CPU
gates only.
JIRA ID : AICK-1943
---------
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
…grad on gfx950 (#12951) ## Summary Depthwise backward-weight convolution has had merged-group (`Gm`) support in the kernel since #12359, but nothing ever asked for it. `group_merge` defaults to 1 and `dispatch_conv_grouped` never set it, so the only ways to reach a merged kernel were to hand-construct a spec in a test or pass `--group-merge N` on the benchmark CLI. Every production depthwise wgrad ran unmerged. This PR adds the selection policy — `_wgrad_merge_degree()` — plus the plumbing that lets a non-default degree survive the trip from spec selection to launch. No kernel capability is added or changed. Why it matters: in wgrad the per-group GEMM is `kpg x (Z*Y*X*cpg) x (N*Ho*Wo)`, so depthwise (`cpg == kpg == 1`) makes the free axis one element wide and every B load scalar. Merging makes that run `Gm` elements long — `Gm` *is* the load vector width. ## What changed `library/dispatch/grouped_convolution.py`: - **`_wgrad_merge_degree()`** — the policy. Picks the largest degree the kernel's own gate admits at the tile that already ships. - **`ConvGroupedSpec.group_merge`** field, forwarded by `to_wgrad_spec()`, and emitted by `kernel_name()` as `gm{N}` — merging changes the GEMM, so merged and unmerged kernels must not collide on the compile-cache key. - **`_wgrad_grid()`** — z extent becomes `(groups / Gm) * split_k`; without this we would launch `Gm x` too many CTAs. - **`_resolve_wgrad_split_k()`** — split-K is chosen from GEMM dims that merging scales, and merged + `split_k > 1` is routed to two-stage. The workspace sizing is deliberately left **unmerged**, mirroring `wgrad_two_stage_workspace_nbytes`: merging changes how many CTAs write `dW`, not how much `dW` there is. Scaling it by `gm` would overstate the allocation quadratically and start refusing two-stage on problems that fit. ### There is no fitted constant, and that is the result The forward selector (#12793) is a fitted closed form with eight constants. Backward-weight ended up with none. Forward puts `Gm` on GemmN+GemmK, trading vector width against a working set growing in `Y*X*Gm` — a genuine interior optimum. Wgrad puts `Gm` on GemmM+GemmN and the gate caps it at `Z*Y*X*cpg*Gm <= tile_n`; with the tile fixed, the admissible set is two or three degrees wide and its top is nearly always best. Three richer policies were fitted and rejected; `score_policy.py` re-derives the rejection from the raw sweeps. ## Measured evidence Relative ratios only, per `platform/AGENTS.md` §Compliance. No absolute latency, throughput, or achieved-FLOP figure appears here or in any commit message. - **Hardware target:** gfx950. - **Workload / axes:** backward-weight (`-F 4`) grouped convolution, depthwise (`cpg == kpg == 1`), fp16 and bf16, NHWC. 197 distinct shapes — ~half depthwise stages from 35 published architectures (mobilenet v1–v4, efficientnet b0–b6/v2, convnext/convnextv2, replknet/unireplknet, shufflenetv2, ghostnet, mnasnet, xception, yolo11/12, …), ~half a stratified random grid including a deliberate ~12% unmergeable minority. - **Baseline:** the configuration that shipped at this PR's merge-base (`2489f0b9153`) — same tile, `Gm = 1`. Ratios are against that, not against an oracle, because that is what the change is worth. - **Method:** `measure_degrees.py`, 2 warmup + 10 timed iterations per config, tile pinned, degree axis fully open, split-K open, epilogue pinned to two-stage so the comparison isolates the degree. Both arms measured in the same process on the same silicon. At these measured configurations, against the previously shipped point: | | ratio | |---|---| | geometric mean over 197 shapes | **1.98x** | | median | 1.25x | | p90 | 7.12x | | max | 36.7x | | min | 0.97x | 62/197 shapes select `Gm = 1` and are bit-for-bit the previous behaviour. 90/197 land within ±2%. One shape regresses beyond 2% (0.971x, `C128 K128 G128 5x5`). Against the best degree available at the shipped tile, the policy captures 0.998 of the oracle, and picks the exact best on 97 of the 135 shapes that have more than one admissible degree; of the other 38 the worst is 0.954x and none falls below 0.95x. **Honest losses:** the three rejected ladder policies all scored a higher geometric mean (up to 2.30x) by spending a wider `tile_n` — and all carried 20–45 regressions out of 197, worst 0.462x. Buying the mean with a long tail of halvings was not a trade worth shipping. Details and the rejected fits are in the case study. ## Correctness - **Numeric gate in the measured run:** `--verify` was forced for every config; 13,244 ranked configurations across three `tile_n` rungs, **13,244 PASS, 0 FAIL**. Checked against an independent torch reference at relative error `< 5e-2` for 16-bit compute (`benchmark_implicit_gemm_conv.py:288`). No unverified kernel was allowed to set an oracle or a ratio. - **On-GPU suite:** `library/tests/test_conv_wgrad_correctness.py` — 66 passed, 346 subtests. - **Selection suite:** 512 passed, 2279 subtests across `library/tests/dispatch/`, `test_conv_wgrad_group_merge_gate.py`, `test_library_layering.py`. `TestWgradMergeDegree::test_every_choice_is_admissible` is the load-bearing gate: at `split_k == -1` the kernel gate's split-K clause is vacuous, so a degree the gate refuses does not merely merge badly — `support()` rejects and depthwise wgrad stops dispatching entirely. That test replays every selected degree through the real `wgrad_group_merge_available`, so a clause that moves in the kernel and not in the selector fails loudly. The second commit pins a separate invariant: merged + `split_k > 1` must be two-stage. A merged tile computes a `Gm x Gm` block of group *pairs* and wants only the diagonal; the packed-atomic epilogue cannot mask off the rest, so taking it would accumulate an off-diagonal pair into a live `dW` element — a wrong gradient with nothing raised. ## Reproducing All of it, from `examples/gfx950/conv_wgrad/merge_degree_model/`: ```bash python3 gen_shapes.py --count 200 --model-cap 100 -o shapes_wgrad.txt python3 measure_degrees.py shapes_wgrad.txt -o deg_tn64.csv python3 score_policy.py --shipped deg_tn64.csv The corpus is checked in so results are reproducible; the sweep CSVs are not — they are raw per-kernel timings carrying absolute numbers. score_policy.py degrades cleanly when they are absent. Narrative and the rejected fits: wgrad_merged_groups_case_study.md. JIRA-ID : AICK-2486 --------- Co-authored-by: Claude Opus 5 <noreply@anthropic.com>
github-actions
Bot
added
project: hipblaslt-provider
rocke: platform
project: hipdnn-integration-tests
project: hipcub
project: hipdnn
project: hipfft
project: hipsolver
project: hipsparselt
ci:hipsparselt-fast
project: miopen
project: miopen-provider
project: rocblas
project: rocfft
project: rocprim
project: rocrand
project: rocsparse
project: rocthrust
project: rpp
shared: origami
shared: stinkytofu
shared: rocroller
shared: tensile
project: rocwmma
project: hiptensor
github actions
ck: convolution
ck: attention-moe
labels
AviralGoelAMD
merged commit Oct 14, 2026
e3ccb2e
into
Parthiv911:users/Parthiv911/rocke/gfx942-dense-cfvst-kdb-vpf
21 of 42 checks passed
Author
|
Superseded: these commits were pushed directly to |
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: AICK-1900
Why this change
Review fixes for ROCm#13205, stacked on its head (
6e35a8a6d3a). Merging this PR intousers/Parthiv911/rocke/gfx942-dense-cfvst-kdb-vpfbrings the branch up to date withdevelop(including ROCm#12620, gfx942 dense sinks) and addresses four of @aosewski's review comments. The changes are scoped to those comments.What changed
kernels/gfx942/attention_dense.py(lse_store)m_iplus develop's sinks-awarerow_has_mass. The pipeline'sloop.resultsbinding is renamedloop_res, so a staleres[0]read fails at build time.kernels/gfx942/attention_dense.py(supports_attention_dense)use_sinkson the bf16 CFVST andlds_num_buffers=2paths. ROCm#12620 removed the blanket sinks rejection, and sinks x CFVST/KDB-VPF has no golden or numeric coverage.tests/test_attention_dense_gfx942.pytests/golden/*.json,tests/test_attention_dense_gfx942_golden.pybuilders/.../gfx942_dense_cfvst_kdb_vprefetch_case_study.md_use_cfvstand module docstringsRegister usage (r4269756187)
gfx942, B1 S4096 Hq32 Hkv8 D128 causal. Values come from the code-object notes (
.vgpr_count/.vgpr_spill_count/.private_segment_fixed_size). Thewaves_per_eu=2cap is 256 VGPRs.llvm20: VGPR / spills / scratch Bllvm23: VGPR / spills / scratch Buse_cfvst=Truelds_num_buffers=2emit_lse=Truepersistent=True,num_persistent=304llvm23the KDB/VPF pipeline fits; only the persistent grid spills.llvm20, 7 atllvm23).llvm22was not measured.How we validated
Tested tree: this branch.
test_attention_dense_gfx942_numeric.py -k "cfvst_pipeline or kdbvpf or lse or sink": 165 passedllvm23), gfx942 304-CUres[0]without the rename:lse_kdbvpf_d128_bf16_causaldrifts in all 3 flavors.res[0]with the rename:UnboundLocalError. PR-siderow_has_key:lse_sink_persist_empty_window_d128_bf16drifts in all 3 flavors. New rejection rows: 3 fail with the gate removed.pre-commit run --from-ref origin/develop --to-ref HEADNotes
persistent/exp2_fast, hand-built or dispatch-resolved). The opening paragraph marks both as not yet recorded (see also r4269756253)._COHORTtightening, ablation provenance, dispatch reachability, docstring wording, and thelds_num_bufferssplit) are not addressed here.