Repository navigation
feat(hipblaslt): support complex CGEMM/ZGEMM on gfx1250 - #10355
Conversation
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
Codecov Report❌ Patch coverage is Additional details and impacted files@@ Coverage Diff @@
## develop #10355 +/- ##
========================================
Coverage 69.18% 69.19%
========================================
Files 2762 2762
Lines 454772 454774 +2
Branches 67092 67096 +4
========================================
+ Hits 314625 314644 +19
+ Misses 117273 117257 -16
+ Partials 22874 22873 -1
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
b53191b to
22ba5bd
Compare
22ba5bd to
4ddb6be
Compare
…egen test fix(cgemm): use SIA0 for complex WMMA on gfx1250 SIA3 flat-schedules the mac items and recycles the de-interleave temps early, corrupting complex WMMA results. Reject SIA3 for the complex WMMA de-interleave path, regenerate gfx1250 CGEMM logic with SIA0, and move the gfx1250 CGEMM test to a dedicated gemm/gfx12 config (top-level cgemm.yaml keeps its MFMA-only SIA[1,3] fork for gfx942/950). test(hipblaslt-test): enable complex gtests on gfx1250 Add gfx1250 to the gpu_arch gate of the 9 complex matmul gtest cases. test(cgemm): drop LibraryLogic/LibraryClient from gfx1250 tox config test(complex): raise gfx1250 CGEMM/ZGEMM kernel coverage to 18 per group
test(complex): force gfx1250 target in codegen tests so CI non-gfx1250 hosts pass
4ddb6be to
8b17309
Compare
There was a problem hiding this comment.
Pull request overview
Adds hipBLASLt/Tensile support for complex GEMM on gfx1250, accounting for gfx1250 being WMMA-only (no f64 matrix ops) by using WMMA-based emulation for CGEMM and VALU f64 FMA for ZGEMM, plus associated validation and test coverage.
Changes:
- Extend Tensile assembly/codegen to correctly handle complex WMMA operand layout (de-interleave) and adjust matrix-instruction inst-type selection for complex.
- Add gfx1250-specific CGEMM/ZGEMM test configs and unit “codegen signature” tests to validate expected ISA patterns without requiring a GPU.
- Add gfx1250 library logic YAMLs and enable complex matmul gtests to run on gfx1250.
Reviewed changes
Copilot reviewed 29 out of 29 changed files in this pull request and generated 2 comments.
Show a summary per file
| File | Description |
|---|---|
| projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_zgemm_gfx1250_codegen.py | New unit test asserting gfx1250 ZGEMM codegen emits VALU v_fma_f64 and avoids f64 matrix ops. |
| projects/hipblaslt/tensilelite/Tensile/Tests/unit/test_cgemm_gfx1250_codegen.py | New unit test asserting gfx1250 CGEMM codegen emits WMMA f32 ops and expected de-interleave markers. |
| projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/zgemm_gfx1250.yaml | Adds gfx1250-specific ZGEMM benchmark/config coverage (VALU f64 path). |
| projects/hipblaslt/tensilelite/Tensile/Tests/common/gemm/gfx12/cgemm_gfx1250.yaml | Adds gfx1250-specific CGEMM benchmark/config coverage (WMMA/TDM path). |
| projects/hipblaslt/tensilelite/Tensile/SolutionStructs/Solution.py | Adds solution rejection rules for unsupported complex WMMA scheduling and f64-complex with TDM. |
| projects/hipblaslt/tensilelite/Tensile/KernelWriterAssembly.py | Updates inst-type selection for complex WMMA and implements complex WMMA de-interleave + negate handling in assembly writer. |
| projects/hipblaslt/tensilelite/Tensile/Components/TensorDataMover.py | Simplifies TDM data_size encoding based on element byte size (log2), with guard for unsupported sizes. |
| projects/hipblaslt/tensilelite/Tensile/Components/MAC_F64C.py | Modernizes complex-f64 VALU MAC component to use rocisa instruction objects and updated signature. |
| projects/hipblaslt/tensilelite/Tensile/Components/MAC_F32C.py | Modernizes complex-f32 VALU MAC component similarly (rocisa instruction objects + updated signature). |
| projects/hipblaslt/tensilelite/Tensile/Components/LraTileAssignment.py | Adds complex FP32 variant for transposed MFMA tile assignment consistent with non-LDSTr complex path. |
| projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx1250/GridBased/gfx1250_Cijk_Alik_C_Bjlk_Z_B_UserArgs.yaml | Adds gfx1250 ZGEMM logic (TT + conjugate variants) for user-args kernels. |
| projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx1250/GridBased/gfx1250_Cijk_Alik_C_Bjlk_C_B_UserArgs.yaml | Adds gfx1250 CGEMM WMMA/TDM logic for user-args kernels. |
| projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx1250/GridBased/gfx1250_Cijk_Alik_Bljk_Z_B_UserArgs.yaml | Adds additional gfx1250 ZGEMM logic variant. |
| projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx1250/GridBased/gfx1250_Cijk_Alik_Bjlk_C_C_B_UserArgs.yaml | Adds additional gfx1250 CGEMM logic variant (conjugate coverage). |
| projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx1250/GridBased/gfx1250_Cijk_Alik_Bjlk_C_B_UserArgs.yaml | Adds additional gfx1250 CGEMM logic variant. |
| projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx1250/GridBased/gfx1250_Cijk_Ailk_Bljk_C_B_UserArgs.yaml | Adds gfx1250 CGEMM logic for alternate index assignment/transpose variant. |
| projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx1250/GridBased/gfx1250_Cijk_Ailk_Bjlk_C_C_B_UserArgs.yaml | Adds gfx1250 CGEMM logic (conjugate coverage) for another transpose combination. |
| projects/hipblaslt/library/src/amd_detail/rocblaslt/src/Tensile/Logic/asm_full/gfx1250/GridBased/gfx1250_Cijk_Ailk_Bjlk_C_B_UserArgs.yaml | Adds gfx1250 CGEMM logic for another transpose combination. |
| projects/hipblaslt/clients/tests/data/matmul_gtest.yaml | Extends complex matmul gtest selection regex to include gfx1250. |
CGEMM (WMMA) now uses the TDM load path; TensorDataMover.dataSizeShift gains complex-single support (whole 8B element). ZGEMM's f64-complex is 16B, which exceeds TDM's 2-bit data_size field, so TDM is rejected and ZGEMM stays on the VALU path; both switch ScheduleIterAlg to 4. fix(complex): use numRegistersInB for B de-interleave; clarify SIA reject comment
8b17309 to
010175a
Compare
|
These CI failures are not related to this PR's changes. |
JIRA ID : https://amd-hub.atlassian.net/browse/AIHPBLAS-4258 Follow-up PR of #10355 ## Motivation gfx1250 CGEMM already supports WMMA+TDM codegen (#10355), and already add some CGEMM+TDM solutions. This PR adds more the CGEMM TDM solutions to logic yamls to improve coverage. <!-- Explain the purpose of this PR and the goals it aims to achieve. --> ## Technical Details Adds WMMA+TDM solutions to NN/NT/TT of the gfx1250 CGEMM logic yamls (`gfx1250_Cijk_*_C_B_UserArgs.yaml`), merged into the existing logic yamls. After this PR each transpose has >=3 TDM solutions (NT=3, NN=4, TT=3, TN=3). p.s. TN already had 3 TDM solutions from #10355 and is left unchanged in this PR. <!-- Explain the changes along with any relevant GitHub links. --> ## Test Plan Clean build for gfx1250, run `hipblaslt-test` including CGEMM (`--gtest_filter=*matmul*complex*`). <!-- Explain any relevant testing done to verify this PR. --> ## Test Result All passed <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: mengzcai <mengzcai@amd.com>
JIRA ID : N/A ## Motivation The gfx1250 complex CGEMM/ZGEMM codegen unit tests (`test_cgemm_gfx1250_codegen`, `test_zgemm_gfx1250_codegen`), added in #10355, fail with `AssertionError: no CGEMM/ZGEMM assembly generated` when run via the tox/CI path (`tensilelite tox common`), while passing standalone. Root cause: `--global-parameters` used argparse's default "store" action, so a second `--global-parameters` occurrence silently replaced the first. These tests pass `--global-parameters KeepBuildTmp=True`, but the `tensile_args` fixture then appends `--global-parameters CheckASMCodeSize=True`, clobbering `KeepBuildTmp`. With `KeepBuildTmp` gone, `removeTemporaries = not KeepBuildTmp` wiped `build_tmp` (which holds the emitted `.s`), so `rglob("*C_B*.s")` found nothing and the assertion failed. The codegen itself succeeded (`18/18 after KernelWriter`); only the artifacts were deleted before the assertion. ## Technical Details - `Tensile/Tensile.py`: switch the `--global-parameters` argument to `action="extend"` so repeated occurrences accumulate instead of the last one clobbering earlier ones. `default=[]` remains safe because argparse copies the list before extending, so the shared default is never mutated and parameters do not leak across successive `Tensile()` calls in a single pytest process. - No changes to the #10355 test files; the fix is at the CLI layer and also removes the same latent footgun for any other caller passing `--global-parameters` more than once (e.g. `test_keep_build_tmp`). - Added a regression test in `test_tensile_helpers_char.py` pinning the accumulate behavior. ## Test Plan - Reproduce the CI condition (append `--global-parameters CheckASMCodeSize=True` with `--prebuilt-client`) and run both codegen tests. - Run the full `Tensile/Tests/unit` suite with the same CI flags and `-n 4` to confirm no regression under warm-cache, full-session ordering. - Run the arg-parsing characterization tests to confirm defaults/behavior unchanged. ## Test Result - `test_cgemm_gfx1250_codegen` / `test_zgemm_gfx1250_codegen`: PASSED (previously FAILED via tox/CI). - Full `Tensile/Tests/unit` (CI flags, `-n 4`): 6410 passed, 206 skipped, 19 xfailed, 0 failed (~2:55); `test_keep_build_tmp[False]/[True]` also PASSED. - Arg-parsing tests: 41 passed, incl. the new accumulate regression test. ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. ## Risk level Low — single-line CLI parser change (`store` -> `extend`) plus one added test. No behavior change for single `--global-parameters` usage; the shared default is not mutated (argparse copies before extending), verified by the full unit suite. --- Fixes CI failure of the gfx1250 CGEMM/ZGEMM codegen tests introduced in #10355.
JIRA ID : https://amd-hub.atlassian.net/browse/AIHPBLAS-4258
Motivation
Support complex CGEMM/ZGEMM on gfx1250
Technical Details
gfx1250 is WMMA-only with no f64 matrix instruction.
CGEMM: uses the WMMA f32 path. WMMA K=4 needs planar operands, so the interleaved complex load is de-interleaved into temps before each WMMA. Restricted to ScheduleIterAlg=0 (SIA3 flat-schedules the mac items and recycles the temps early → wrong results).
ZGEMM: no f64 WMMA, so it runs the VALU v_fma_f64 path (MAC_F64C), like real DGEMM on this arch.
Added LraTileAssignment complex variant; fixed stale MAC_F32C/MAC_F64C (signature + modern rocisa API).
Added gfx1250 CGEMM/ZGEMM library logic, dedicated gemm/gfx12/{c,z}gemm_gfx1250.yaml tox configs, codegen unit tests, and enabled the complex hipblaslt-test gtests on gfx1250.
Test Plan
Test tox and hipblaslt-test on B0
Test Result
Test tox and hipblaslt-test passed on B0
Submission Checklist