ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-17 00:58:44 +00:00
Code Issues Packages Projects Releases Wiki Activity
1,956 Commits 779 Branches 40 Tags
6fb2b54ff4e827ee1fa42e0db73625bf4f22ddc3
Commit Graph

118 Commits

Author SHA1 Message Date
OscarXu
ec8d00d58d mx_moe_fp4 ready for aiter with clang-format. 2025-05-16 04:09:26 -05:00
OscarXu
f03a96ffcf align file and function name 2025-05-16 03:54:28 -05:00
OscarXu
39ff3fbf05 v3 function pass 2025-05-16 03:42:48 -05:00
OscarXu
f70f778e27 v1 compile pass. Function not ready 2025-05-15 08:01:56 -05:00
OscarXu
7c76e19e55 Merge remote-tracking branch 'origin/fp4_gu_moe' into fp4_gu_moe_gemm1 2025-05-15 16:37:08 +08:00
mtgu0705
b65bc1ba4a added mx moe block v3 support, function passed 2025-05-15 03:26:37 -05:00
OscarXu
bcb4f7e98c Add gemm1 v1 to selector 2025-05-15 16:22:45 +08:00
OscarXu
c0babbca62 Merge remote-tracking branch 'origin/fp4_gu_moe' into fp4_gu_moe_gemm1 2025-05-15 16:12:19 +08:00
OscarXu
17922821ec Add gemm1 v1 2025-05-15 16:11:43 +08:00
mtgu0705
4e2ec31e4d rename moe block selector and pipeline 2025-05-15 02:12:50 -05:00
Your Name
e060999d29 v3 2025-05-15 13:19:07 +08:00
mtgu0705
102151ebcf temp save 2025-05-14 08:13:47 -05:00
mtgu0705
2700b217be 16x16x128 input size blockscale function passed 2025-05-14 03:20:59 -05:00
mtgu0705
1bbb50b212 mfma using asm, device result correct, host result need to check 2025-05-13 20:57:34 -05:00
mtgu0705
5ba86c210b updated and build passed 2025-05-13 14:49:37 +08:00
mtgu0705
cc43f88f08 add code for mxfp4 gemm, blockscale not supported yet 2025-05-12 20:56:50 +08:00
Your Name
dff709c79f Revert "update" 
This reverts commit 960b2bce1c.
 2025-05-12 16:24:12 +08:00
Your Name
58f848cc07 Merge branch 'wjx/align_v3_pipeline' into fp4_gu_moe 2025-05-12 16:21:39 +08:00
aska-0096
22f8f5e5d8 fix bug for a lds read 2025-05-12 04:00:21 +00:00
mtgu0705
0bddd63d9c fix bugs, build passed 2025-05-11 17:49:43 +08:00
mtgu0705
d3f007c775 hotfix 2025-05-11 16:58:36 +08:00
mtgu0705
7d272f7f63 fixed some bugs 2025-05-10 22:01:20 +08:00
mtgu0705
70648240f9 added fp4_bpreshuffle example, build failures 2025-05-10 21:34:32 +08:00
Your Name
5421e71155 Merge branch 'wip-f4' into mt1 2025-05-10 14:56:21 +08:00
mtgu0705
bb49d2e623 fix the bug, functional test passed 2025-05-10 14:33:44 +08:00
mtgu0705
493a8639b3 fix a bug 2025-05-10 13:07:43 +08:00
aska-0096
087b20dc1d clang format 2025-05-09 16:15:10 +00:00
aska-0096
0987b0af44 remove unnecessary hacky 2025-05-09 16:07:22 +00:00
aska-0096
86bff9c46e Fix pipe v3 correctness issue 2025-05-09 15:54:43 +00:00
aska-0096
7bde4b8d34 Add pipeline v3. Have some runtime issue and register spill 2025-05-09 09:47:22 +00:00
mtgu0705
f2a474e2e9 fix update 2025-05-09 11:04:39 +08:00
mtgu0705
11f386108e some fixes 2025-05-08 23:38:27 +08:00
aska-0096
b2efb06315 Spilt the fp4 target. Fix the known bugs. 128x128x128 sanity checked; remove prints 2025-05-08 15:07:33 +00:00
mtgu0705
7c49f9dd31 add mx fp8 b_preshuffle support, function not yet tested. 2025-05-08 22:41:54 +08:00
lalala-sh
6c459e8c38 Merge branch 'develop' into wjx/align_v3_pipeline 2025-05-08 18:08:31 +08:00
lalala-sh
960b2bce1c update 2025-05-08 09:48:23 +00:00
aska-0096
32e6c8aad3 16x16x128 correct; 64x64x128 failed 2025-05-08 09:42:53 +00:00
Ding, Yi
6519ba3aa9 wip3 2025-05-08 07:44:22 +00:00
Ding, Yi
8d51a4ae96 wip2 2025-05-07 09:40:47 +00:00
Ding, Yi
d12b750bb0 wip 2025-05-07 08:17:07 +00:00
lalala-sh
0ab978584d fix bugs 2025-05-06 07:36:59 +00:00
lalala-sh
9c06c3817a [fix] align v3 gufusion pipeline 2025-04-30 02:27:39 +00:00
lalala-sh
b8427b812e align v3 gufusion pipeline 2025-04-30 01:35:49 +00:00
Anton Gorenko
edd92fc546 DeviceGemm_Wmma_CShuffleV3 with BlockGemmPipelineVersion::v3 (#2096) 
* Prepare files for DeviceGemm_Wmma_CShuffleV3

* Implement main part of CShuffleV3 with block pipeline v3 for WMMA

* Remove unused functions and template params for A/B descriptors

* Support both gfx11 and gfx12

* Enable SplitK for gfx12 and disable for gfx11

* Added RowColRow layout for DeviceGemmV2 fp16

* Added more instances for Row, Col, Row data layout

* Added instances for DeviceGemm_Wmma_CShuffleV3, Col, Row, Row data layout

* Added instances for DeviceGemm_Wmma_CShuffleV3, Col, Col, Row data layout

* Added more instances for DeviceGemm_Wmma_CShuffleV3, Row, Row, Row data layout

* Fix formatting

* Add documentation

Based on e5ad48a784

* Enable gemm_universal profiling for gfx11/12

* Add WMMA intrinsics for F8/BF8

* Support F8/BF8 DeviceGemm_Wmma_CShuffleV3, add basic instances

* Add BF16 instances and tests

* Fix test_gemm_universal_wmma_fp8 by adding CK_USE_WMMA_FP8

---------

Co-authored-by: Anca Hamuraru <anca@streamhpc.com>
 2025-04-28 10:14:21 +05:00
aska-0096
bc9c819aa4 generalized bpreshuffle pipeline optimization 2025-04-27 11:50:30 +00:00
aska-0096
49338edb1b tempsave 2025-04-27 08:05:20 +00:00
aska-0096
946a2119cd temp save 2025-04-25 05:12:47 +00:00
coderfeli
cd7955ce8b Merge branch 'swdev_528812' into dev/moe_opt 2025-04-25 02:51:33 +00:00
coderfeli
ddb5f36eeb add missing file 2025-04-24 11:10:31 +00:00
coderfeli
ceaa5a984b gu fusion v3 2025-04-24 11:05:08 +00:00