Repository navigation
fix(tensilelite): fix TDM pointer-array batch handling - #11334
Conversation
Codecov Report❌ Patch coverage is ❌ Your project check has failed because the head coverage (38.52%) is below the target coverage (80.00%). You can increase the head coverage or adjust the target coverage. Additional details and impacted files@@ Coverage Diff @@
## develop #11334 +/- ##
===========================================
+ Coverage 69.64% 69.65% +0.02%
===========================================
Files 2826 2815 -11
Lines 469878 466858 -3020
Branches 69394 68907 -487
===========================================
- Hits 327211 325187 -2024
+ Misses 118978 118101 -877
+ Partials 23689 23570 -119
*This pull request uses carry forward flags. Click here to find out more.
🚀 New features to boost your workflow:
|
Mailbox ds_store/ds_load carry memTokenLdsBuffer0 so the scheduler places a publish fence between them and a release before the next LDS0 write. WG barriers stay untokened. SK5 extracts HybridMode from MagicShift bit 30, then clears that bit by XOR with (HybridMode<<30). The magic add + 5-bit shift is masked into a temp so the SKTiles overlay in MagicShiftItersPerTile is left intact. sk_hybrid_quick includes a TDMInst=3 HGEMM group (SK5, SIA 0/4).
JIRA ID : N/A
ArgType 3 supplies A and B as pointer arrays. TDM loads the selected batch
matrix, then applies tile offsets to that matrix base. A/B strided-batch WG2
offsets are suppressed after the pointer-array load.
Motivation
- Dereference the selected A and B batch pointers before applying TDM tile offsets on classic and subtile paths.
- Suppress strided-batch offsets after selecting a pointer-array matrix base, and only for A/B.
- Wait after the conventional A/B pointer-array SRD load before reusing that temp for batchOffset (gfx1250 XNACK replay).
- Route pointer tables into output-conversion C/D, and reject sparse, grouped, MBSK, and non-MultipleBuffer GSU>1.
- Initialize non-wave-separated TDM descriptors before applying their global offsets.
- Add standalone-client BatchMode plumbing for the pointer-array ABI.
- Add gfx1250 runtime regressions for GSU1/GSU2, static SK3/SK5, and subtile TDM, sweeping SIA 0 and 4.
Test Plan
- Characterization: r3 datamover, StreamK TDM SGPR-budget, ClientConfigIni.
- `tdm_general_batched_gfx1250.yaml`: non-StreamK pointer-array TDM, SIA 0 and 4.
- `sk_tdm_general_batched.yaml`: SK {3,4,5} x SIA {0,4} x hybrid {0,1}, BatchMode 1.
- `sk_tdm_subtile_general_batched.yaml`: subtile SK3, SIA 0 and 4.
- Rebuild gfx1250 device lib / hipBLASLt gtest; nightly general-batched BF16.
- tox on gfx1250
Test Results
- Characterization tests listed above passed.
- Non-StreamK pointer-array TDM passed GSU1/GSU2, SIA 0 and 4.
- SK3, SK4, and SK5 pointer-array TDM passed SIA 0 and 4.
- Subtile pointer-array TDM (SK3, SIA 0 and 4) passed.
- gfx1250 device-library and hipBLASLt gtest builds passed.
- The original nightly general-batched BF16 regression passed.
7ba08b2 to
b2f7586
Compare
76e3a0d to
f9a04e6
Compare
|
Thanks for your comments. |
✅ All Policy Checks Passed
📖 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. |
|
Fixes items added to |
|
@mengzcai Could you help to review this? |
|
This function should work for complex TDM as well (ROCM-30021: gfx1250 CGEMM via the pointer-array batched API). |
TheRock Submodule Bump ActivityNewest first
|
JIRA ID : ROCM-30021 ## Motivation Add a complex general-batch edge case covering the pointer-array (batch_mode=1) path over degenerate/boundary sizes and conjugate transposes, and toggle batch_mode [0, 1] on matmul_batch_medium_complex, to guard [ROCM-30021](https://amd-hub.atlassian.net/browse/ROCM-30021). <!-- Explain the purpose of this PR and the goals it aims to achieve. --> ## Technical Details Add batch_mode: [0, 1] to matmul_batch_medium_complex and one gtest case `matmul_batch_complex_general_batch_edge` reusing existing anchors (`one_matrix_size_range` + `complex_transA_transB_range`), covering batch_mode [0, 1]. No source changes. <!-- Explain the changes along with any relevant GitHub links. --> ## Test Plan `hipblaslt-test --gtest_filter=*matmul*complex*` on gfx1250. <!-- Explain any relevant testing done to verify this PR. --> ## Test Result 8684 complex tests pass, 0 failures (verified on develop branch that includes #11334). <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. [ROCM-30021]: https://amd-hub.atlassian.net/browse/ROCM-30021?atlOrigin=eyJpIjoiNWRkNTljNzYxNjVmNDY3MDlhMDU5Y2ZhYzA5YTRkZjUiLCJwIjoiZ2l0aHViLWNvbS1KU1cifQ Co-authored-by: mengzcai <mengzcai@amd.com>
…12074) #11653 Source commit: `a873b5eb2ab4cd7d58b4e2f3e3a42e64640e14fa` JIRA ID : ROCM-30021 ## Motivation Add a complex general-batch edge case covering the pointer-array (batch_mode=1) path over degenerate/boundary sizes and conjugate transposes, and toggle batch_mode [0, 1] on matmul_batch_medium_complex, to guard [ROCM-30021](https://amd-hub.atlassian.net/browse/ROCM-30021). <!-- Explain the purpose of this PR and the goals it aims to achieve. --> ## Technical Details Add batch_mode: [0, 1] to matmul_batch_medium_complex and one gtest case `matmul_batch_complex_general_batch_edge` reusing existing anchors (`one_matrix_size_range` + `complex_transA_transB_range`), covering batch_mode [0, 1]. No source changes. <!-- Explain the changes along with any relevant GitHub links. --> ## Test Plan `hipblaslt-test --gtest_filter=*matmul*complex*` on gfx1250. <!-- Explain any relevant testing done to verify this PR. --> ## Test Result 8684 complex tests pass, 0 failures (verified on develop branch that includes #11334). <!-- Briefly summarize test outcomes. --> ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. [ROCM-30021]: https://amd-hub.atlassian.net/browse/ROCM-30021?atlOrigin=eyJpIjoiNWRkNTljNzYxNjVmNDY3MDlhMDU5Y2ZhYzA5YTRkZjUiLCJwIjoiZ2l0aHViLWNvbS1KU1cifQ Co-authored-by: mengzcai <152131106+mengzcai@users.noreply.github.com> Co-authored-by: mengzcai <mengzcai@amd.com>
…#12072) #11334 Source commit: `6f0c83a90f7b00f873c1b51e82f54dd4e0e136bb` JIRA ID : AIHPBLAS-4515 ## Motivation ArgType 3 supplies A and B as pointer arrays. TDM loads the selected batch matrix, then applies tile offsets to that matrix base. A/B strided-batch WG2 offsets are suppressed after the pointer-array load. ## Technical Details - Dereference the selected A and B batch pointers before applying TDM tile offsets on classic and subtile paths. - Suppress strided-batch offsets after selecting a pointer-array matrix base, and only for A/B. - Wait after the conventional A/B pointer-array SRD load before reusing that temp for batchOffset (gfx1250 XNACK replay). - Route pointer tables into output-conversion C/D. - Initialize non-wave-separated TDM descriptors before applying their global offsets. - Add standalone-client BatchMode plumbing for the pointer-array ABI. - Add gfx1250 runtime regressions for GSU1/GSU2, static SK3/SK5, and subtile TDM, sweeping SIA 0 and 4. ## Test Plan - Characterization: r3 datamover, StreamK TDM SGPR-budget, ClientConfigIni. - `tdm_general_batched_gfx1250.yaml`: non-StreamK pointer-array TDM, SIA 0 and 4. - `sk_tdm_general_batched.yaml`: SK {3,4,5} x SIA {0,4} x hybrid {0,1}, BatchMode 1. - `sk_tdm_subtile_general_batched.yaml`: subtile SK3, SIA 0 and 4. - Rebuild gfx1250 device lib / hipBLASLt gtest; nightly general-batched BF16. - tox on gfx1250 ## Test Results - Characterization tests listed above passed. - Non-StreamK pointer-array TDM passed GSU1/GSU2, SIA 0 and 4. - SK3, SK4, and SK5 pointer-array TDM passed SIA 0 and 4. - Subtile pointer-array TDM (SK3, SIA 0 and 4) passed. - gfx1250 device-library and hipBLASLt gtest builds passed. - The original nightly general-batched BF16 regression passed. Co-authored-by: Nathan Henderson <nathan.henderson@amd.com>
JIRA ID : AIHPBLAS-4515, AIHPBLAS-4770
AIHPBLAS-4770
Motivation
ArgType 3 supplies A and B as pointer arrays. TDM loads the selected batch matrix, then applies tile offsets to that matrix base. A/B strided-batch WG2 offsets are suppressed after the pointer-array load.
Technical Details
Test Plan
tdm_general_batched_gfx1250.yaml: non-StreamK pointer-array TDM, SIA 0 and 4.sk_tdm_general_batched.yaml: SK {3,4,5} x SIA {0,4} x hybrid {0,1}, BatchMode 1.sk_tdm_subtile_general_batched.yaml: subtile SK3, SIA 0 and 4.Test Results