mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-17 03:19:48 +00:00
f9737ad5531d71ef88ff00a00d2eb578fe199e5a
12 Commits
| Author | SHA1 | Message | Date | |
|---|---|---|---|---|
Aviral Goel
|
b3c6098af2 |
ck_tile: add gtest unit tests for MX flatmm (gfx950) (#5082)
## Summary
- Add correctness unit tests for the MX-format flatmm kernel
(`example/ck_tile/18_flatmm/mxgemm`) under `test/ck_tile/flatmm/`
- Tests cover all five dtype combinations: FP4×FP4, FP8×FP8, FP6×FP6,
FP8×FP4, FP4×FP8
- Tests cover all four kernel dispatch paths (the `has_hot_loop` ×
`tail_num` product):
- `has_hot_loop=false, tail=ODD` (K=256, num_loop=1)
- `has_hot_loop=false, tail=EVEN` (K=512, num_loop=2)
- `has_hot_loop=true, tail=ODD` (K=768, num_loop=3)
- `has_hot_loop=true, tail=EVEN` (K=1024, num_loop=4)
- Remove unsupported `-split_k` CLI option from
`tile_example_mx_flatmm`; the pre-shuffled B layout is incompatible with
K-splitting and the option silently produced wrong results
## Changes
**New files (`test/ck_tile/flatmm/`):**
- `CMakeLists.txt` — builds 40 kernel instances as a shared OBJECT
library, links into 5 per-dtype test executables; forwards
`-DCK_TILE_USE_OCP_FP8` when `CK_USE_OCP_FP8` is ON
- `test_mx_flatmm_base.hpp` — base test fixture with
`run_test_with_validation(M, N, K, kbatch=1)`
- `test_mx_flatmm_fixtures.hpp` — concrete `TestMXFlatmm` typed test
class and type aliases
- `test_mx_flatmm_fp{4fp4,8fp8,6fp6,8fp4,4fp8}.cpp` — per-dtype
`TYPED_TEST_SUITE` files
**Modified files:**
- `example/ck_tile/18_flatmm/mxgemm/mx_flatmm_arch_traits.hpp` — moved
`preShuffleWeight` here (was in `mx_flatmm.cpp`) so it is includeable by
both the example and the tests
- `example/ck_tile/18_flatmm/mxgemm/mx_flatmm.cpp` / `run_mx_flatmm.inc`
— removed `-split_k` CLI arg, hardcoded `k_batch=1`, fixed `k_split`
formula, updated call sites after `preShuffleWeight` move
- `test/ck_tile/CMakeLists.txt` — added `add_subdirectory(flatmm)`
---------
Co-authored-by: Thomas Ning <Thomas.Ning@amd.com>
|
||
|
Andriy Roshchenko
|
44c42614c8 |
[CK TILE] Refactor MX FLATMM example (#4821)
Refactor the MX FLATMM example to support more pipelines across different architectures. This work facilitates the NPI team roadmap. |
||
|
ZheWang
|
418ee44844 |
Mx fp6 flatmm (#3601)
* add fp6 data-type and support sync/async dwordx3 load/store
* clang-format
* pre-commit
* 1st commit
* default mnk pass ut
* fix a distrubution
* fix
* fix bdram distr
* update
* pass ut
* improve perf
* update
* clean code
* resolve copilot comment
* reslove comment
* clang-format
---------
Co-authored-by: ZheWang <zhewan@amd.com>
[ROCm/composable_kernel commit:
|
||
|
Max Podkorytov
|
6cf89bbca9 |
[CK-Tile] move out memory operation from cshuffle epilogue class (#3359)
* initial poc
* factor out common parts in operator()
* cv4
* rest of the universal gemm pipelines
* fix test
* remove boilerplate from tile engine
* fix example
* fix example
* format
* fix tests build for gemm
* remove base pipeline codegen from gemm instance builder
* unify v3 logic with the rest of universal gemm pipelines
* fix build for multi abd test
* fix test gemm multi d
* fix build for weight preshuffle
* fix grouped gemm test
* fix grouped gemm multi d test
* fix grouped gemm preshuffle
* fix grouped gemm example except for quant
* fix gemm preshuffle
* fix splitk 2 stage example
* fix batched gemm example
* fix multid example
* fix multiabd example
* fix batched gemm test
* fixup
* fix examples build
* fix grouped gemm test build
* fix smoke builder
* hacky poc
* fix tile engine
* kill the lambda
* maybe fix test build
* more fixes
* clang-format
* save temp
* clang-format
* mostly fix examples
* clang-format
* remove dead code
* more cleanup
* fix fmha bwd build (default epilogue set/add appears to be broken)
* fix default epilogue tests but not correctness
* clang-format
* fix bquant
* clang-format
* cleanup dead code
* rearrange make windows for readability
* restore changes to IsSupportedArgument
* fix smoke-builder
* clang-format
* fixup rename class
* build fixes
* clang-format
* fix builder
* fixup
* remove set from builder tests
* fix test
* clang-format
* re-refactor the kernels
* clang-format
* fix header license
* remove memory operation from conv bwd test
* clang-format
* clang-format example,include
* clang-format test
* build fixes
* clang-format
* solve compilation error
* fix the CI
* solve compilation error
* clang format
* solve merge conflict
* solve merge conflict
* solve the gfx11 error
* solve test error
* moar build fixes
* remove AtomicAddRequiresKBatchGreaterThanOne test since the property is removed from the kernel scope
---------
Co-authored-by: Thomas Ning <Thomas.Ning@amd.com>
[ROCm/composable_kernel commit:
|
||
|
yadaish
|
fc3ffa0d75 |
[CK_TILE] support split-k a16w4 gemm1 (#3389)
* initial version to support moe gemm1 split-k
* add missing args
* fix build warning
* update reference
* for split-k disable bias and weight
* remove debug log
* fix format
* fix div by zero errors
* fix cmake config
* update
* resolve conflicts
* remove useless changes
* reformat
* fix
* remove useless changes
* fix ci
---------
Co-authored-by: lalala-sh <Jiaxing.Wen@amd.com>
Co-authored-by: root <root@smci355-ccs-aus-m01-25.cs-aus.dcgpu>
[ROCm/composable_kernel commit:
|
||
|
Yi DING
|
fb72fea980 |
[CK_TILE] Add FP8xF4 Flatmm (#3401)
* Refactor policy
* fix a bank conflict
* Enable mixed mx flatmm
* Update
[ROCm/composable_kernel commit:
|
||
|
Yi DING
|
9c7b0388a9 |
[CK_TILE] Generate random tensor values with multiple threads (#3324)
[ROCm/composable_kernel commit:
|
||
|
Yi DING
|
07158d16ad |
[CK_Tile] Flatmm MX Cleanup & Explicite Offset Calculation (#3286)
[ROCm/composable_kernel commit:
|
||
|
Aviral Goel
|
91ffc9dd1e |
chore(copyright): update copyright header for example directory (#3273)
* chore(copyright): update copyright header for codegen directory
* chore(copyright): update copyright header for example directory
[ROCm/composable_kernel commit:
|
||
|
Yi DING
|
e27e760d5a |
[CK_TILE] Add Flatmm MX FP8 (#3208)
* Use async for flatmm mxfp4
* Fix preshuffle
* Add flatmm mxfp8
* Thanks, Copilot
* Thanks Copilot again~
[ROCm/composable_kernel commit:
|
||
|
Yi DING
|
e2060bd1fb |
[CK_TILE] MX Flatmm Split kernel instances (#3207)
* [CK_TILE] MX Flatmm Split kernel instances
* Fix flatmm example compile
[ROCm/composable_kernel commit:
|
||
|
Yi DING
|
acec30dd09 |
[CK_TILE] Add mxfp4 flatmm (#3080)
* Squashed commit of the following: commit 3e1a851dad834776efbe4fe365ac82c4ed312010 Author: Ding, Yi <yi.ding@amd.com> Date: Thu Oct 23 06:10:54 2025 +0000 Fix & clean after rebase commit 1edf485092f44411da9a1796a4a6b72d5cdb67c6 Author: Ding, Yi <yi.ding@amd.com> Date: Wed Oct 22 10:46:13 2025 +0000 Squashed commit of the following: commit 5276b28a51dac7b5d2106fbae8e78de190ee0de1 Author: mtgu0705 <mtgu@amd.com> Date: Mon Sep 22 02:04:27 2025 -0500 fix bandwidth calculation commit d645bb20c6d879154c30ecd82bbff4d2a9206750 Author: mtgu0705 <mtgu@amd.com> Date: Mon Sep 22 00:58:59 2025 -0500 updates commit 0fa7e6b88aaf81a36034aa7607746de295de4263 Author: mtgu0705 <mtgu@amd.com> Date: Fri Sep 19 00:39:46 2025 -0500 fix a bug, set the A DS_read preload size to 4 for MXFP4 commit 50cafa824e2267f2b2f0dfeeb93e69a673630c61 Author: mtgu0705 <mtgu@amd.com> Date: Thu Sep 18 01:19:03 2025 -0500 fix a_wrap preload issue for large MPerBlock. commit e6333bbbc6ef540e24f92095040085f1ed59041e Author: mtgu0705 <mtgu@amd.com> Date: Wed Sep 17 21:34:03 2025 -0500 optimized the VGPR repack issue for MXFP4 commit e99e4932c401b9f6d1893dd5044c2827d6b3f145 Author: Gino Lu <gino.lu@amd.com> Date: Wed Sep 17 04:19:44 2025 -0500 fix time error commit 4586ce6da7fba0514f2e01a8124c76b7d494e124 Author: mtgu0705 <mtgu@amd.com> Date: Wed Sep 17 03:58:00 2025 -0500 updated, function passed. commit c4f25e7579573db5681b9160f6bdb1349f3566f1 Author: mtgu0705 <mtgu@amd.com> Date: Tue Sep 16 22:21:39 2025 -0500 fix, function partially passed commit a51b56eb6b00b99a4e8d2802dbf5b5b5277b54d8 Author: mtgu0705 <mtgu@amd.com> Date: Tue Sep 16 03:01:12 2025 -0500 fix, reference function passed, next check kernel function commit 5b02643ebab18960e8f9ba66c6bd2f91774f9cae Author: Gino Lu <gino.lu@amd.com> Date: Tue Sep 16 02:29:01 2025 -0500 let pack/unpack return pk_fp4_t commit 76d37c5d4b17530e95c6fced31bff66a35d54b8f Author: mtgu0705 <mtgu@amd.com> Date: Mon Sep 15 20:50:26 2025 -0500 fix commit e5be3e162b9a20e5355bd556d2b27afb6d8bf085 Author: Gino Lu <gino.lu@amd.com> Date: Mon Sep 15 05:51:06 2025 -0500 fix bug commit 39a024efe4aa773df589712b1290803bb5ab5d1d Author: mtgu0705 <mtgu@amd.com> Date: Mon Sep 15 04:02:05 2025 -0500 fix core dump issue, function is not correct. commit 16c49d268cfe065b5112b960b2d852b26552686a Author: mtgu0705 <mtgu@amd.com> Date: Mon Sep 15 03:03:02 2025 -0500 updates, build pass commit fe7a961852dee6eff3be3cf1e0d0fabec5cd42ee Author: mtgu0705 <mtgu@amd.com> Date: Mon Sep 15 00:05:18 2025 -0500 updates commit aaf9fe8022a72df59e04e4d5886dca3ba9c23400 Author: Gino Lu <gino.lu@amd.com> Date: Sun Sep 14 23:40:28 2025 -0500 fix bug commit a3da89290e1553b85fbf1171c07e93ac0f5584db Author: Gino Lu <gino.lu@amd.com> Date: Fri Sep 12 03:28:50 2025 -0500 fix interface commit c5ff747e72d877461ba61dc19a0fe15527d3161e Author: Gino Lu <gino.lu@amd.com> Date: Fri Sep 12 02:53:50 2025 -0500 add interface in warp_gemm_impl commit 0a48d369e601cc798589fc59e0784bdbfc0a22f9 Author: mtgu0705 <mtgu@amd.com> Date: Wed Sep 10 05:03:08 2025 -0500 updates some fixes. commit aaa2beca30ff5546d171a2028d1894fd4e131d4e Author: mtgu0705 <mtgu@amd.com> Date: Tue Sep 9 04:37:42 2025 -0500 fix after merge ginolu/add_wgmfma_dispatcher commit bf87449b09cba690922b2f3f78ba39bf1b1e472e Merge: 05ab58e3d 991d7fdbb Author: mtgu0705 <mtgu@amd.com> Date: Mon Sep 8 22:09:15 2025 -0500 Merge remote-tracking branch 'origin/ginolu/add_wgmfma_dispatcher' into mtgu/cktile_mxfp4_flatmm_dev commit 05ab58e3de2b708aceda63d704089c0fa89437ae Author: mtgu0705 <mtgu@amd.com> Date: Mon Sep 8 21:42:47 2025 -0500 update mx flatmm tail pipeline commit 991d7fdbb726d65091a91b5cc2800f798a6661fc Merge: ad046084a |