ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-06-29 11:16:59 +00:00
Code Issues Packages Projects Releases Wiki Activity
1,865 Commits 776 Branches 40 Tags
mi355_transpose_load_dev
Commit Graph

776 Commits

Author SHA1 Message Date
joye
b406b584d7 fix clang format check 2025-04-28 08:56:00 +08:00
joye
292f89c860 fix a transpose index issue 2025-04-27 17:39:52 +08:00
joye
bd091d9a88 fix a transpose index issue 2025-04-27 17:33:32 +08:00
joye
de9407ed93 fix 16x16 related dimension transpose 2025-04-27 02:44:49 -05:00
joye
57e8e34705 exchange the iteration order 2025-04-26 21:37:59 -05:00
joye
25ab38f913 fix transpose related codes 2025-04-27 09:55:29 +08:00
joye
a823e00dc0 delete unused variables 2025-04-25 14:47:49 +08:00
joye
a6564da629 update output tensor distribution 2025-04-25 14:45:45 +08:00
joye
119a8e0e16 hack for transpose 16x16 2025-04-25 00:56:59 -05:00
joye
e2f3c95d24 miss output tile distribution mapping 2025-04-24 21:55:03 -05:00
joye
6beb585dad update tile transpose 2025-04-24 20:36:09 -05:00
joye
efa7243ee5 transpose load enable 2025-04-24 19:19:44 -05:00
joye
df9769afba can pass; but no logic 2025-04-24 19:05:47 -05:00
joye
90a4501869 add some fix 2025-04-24 18:29:38 +08:00
joye
34040f43b6 update some codes 2025-04-24 17:53:19 +08:00
joye
8d75983536 fix a distribution issue 2025-04-24 02:24:31 -05:00
joye
6893165818 add some fixes 2025-04-24 01:20:35 -05:00
joye
afb1cec9c4 update transpose load logic 2025-04-24 11:14:07 +08:00
joye
3918a35870 Merge branch 'mi355_transpose_load_dev' of https://github.com/ROCm/composable_kernel into mi355_transpose_load_dev 2025-04-24 08:28:14 +08:00
joye
aaab4aacc5 Merge branch 'develop' of https://github.com/ROCm/composable_kernel into mi355_transpose_load_dev 2025-04-24 08:27:53 +08:00
joye
c862437f05 fix some issues 2025-04-23 19:27:29 -05:00
carlushuang
5487289fc4 [CK_TILE] support gfx950 matrix core in 01_fmha fwd (#2110) 
* gfx950 01_fmha fwd

* fix comment

---------

Co-authored-by: Thomas Ning <Thomas.Ning@amd.com>
 2025-04-23 12:40:18 -07:00
joye
acce2df3bf fix some compile errors 2025-04-23 05:27:52 -05:00
joye
e14a16359f add transpose load; no real logic 2025-04-23 16:55:13 +08:00
lalala-sh
39ba03f25d Moe gemm activation (#2026) 
* fix useless code and remove usless oob

* clang format

* fix coredump in e2e test

* fix2

* fix clang format

* fix output oob

* impl int64 but result not correct

* int64 index ok now

* input output all ok

* fix uint32

* revert v1 test

* use uint32

* mork to support 13w tokens

* moe sorting fix moebuf

* fix merge

* update moe api fix aiter build

* fix buid

* fuse silu

* silu ok

* acale ok

* add silu

* change code

* gemm2 ok

* gufusion compatible ok, fix warnings

* gu fusion for m32 m64 ok

* support bf16 cshuffle

* i4 gemm2 ok

* i4 gemm2 ok and i4 gemm1 build

* 16x16 run ok

* change flops; change cshuffle dtype

* fuse gelu silu act in moe gemm1

* fp8 with act ready

* int4 act ready

* remove useless changes

* remove useless code change

* fix clang format

* add the arch limit of int4 moe gemm

* fuse moe activation

* fix fp8 16x16

* fix no quant case

* fix bugs

* fix fp8 gufusion bug

* remove useless comments

* refine activation code & complete moe example

* fix int8 bugs

* merge tkw1

---------

Co-authored-by: coderfeli <coderfeli@163.com>
Co-authored-by: feli <felix.li@amd.com>
Co-authored-by: illsilin <Illia.Silin@amd.com>
Co-authored-by: root <root@hjbog-srdc-51.amd.com>
Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
 2025-04-23 10:35:34 +08:00
Gino Lu
504f563f78 [CK-Tile] warp-gemm support for using V_MFMA_F32_16x16x32_BF16 (#2073) 
* draft v_mfma_f32_16x16x32_bf16

* fix error config and add debug code.

* Solve the CShuffle Problem

* draft v_mfma_f32_16x16x32_bf16

* fix error config and add debug code.

* Solve the CShuffle Problem

* fix error while testing new command

* Finished the feature of new mfma 16*16*32

* Addressed the comment

---------

Co-authored-by: ThomasNing <thomas.ning@amd.com>
 2025-04-22 15:52:36 -07:00
Rostyslav Geyyer
416e851584 Temporarily disable MX FP4 device tests (#2112) 2025-04-22 16:08:48 -05:00
Thomas Ning
0cca8fa28f GEMM Multiply Multiply Fix (#2102) 
* fix the type convert and increase the BF16 conversion + the profile comment

* fix the CI
 2025-04-22 01:13:22 -07:00
Muhammed Emin Ozturk
b092c18da7 MI308 fix for streamk 1-Tile floating point exception (#2101) 2025-04-21 11:44:07 -07:00
Thomas Ning
a738e43445 MFMA 16x16x32fp8 (#2103) 
* add mfma_16x16x32_fp8

* clang format code

* Finished the fix for gemm basic

* clang foramt

* rebuild CI

* recover gemm.hpp

* add MFMA 16*16*32bf8

---------

Co-authored-by: solin <bingzhou@amd.com>
 2025-04-21 10:21:35 -07:00
solin
c318ec0778 fix CI build fail 2025-04-21 16:00:12 +08:00
lalala-sh
bcf5bb41be enable do top k weights in moe stage1 gemm (#2094) 
* add switch for mul topk weights

* fix bf16/f16 bugs

* complete
 2025-04-18 10:45:49 +08:00
Andriy Roshchenko
213b203a3c MX GEMM - Parameterized Test Template (#2088) 
* Tests for MX FP8 GEMM

* Improve documentation
 2025-04-16 19:56:00 -06:00
Andriy Roshchenko
da54464cce MX GEMM - Add MX BF8 example (#2071) 
* Add MX GEMM example for MX BF8

* Verified MX FP8 with 16x16x128 scale builtin

* Verify MX BF8 GEMM with BF16 output
 2025-04-16 15:25:02 -06:00
BingYuan.Zhou
eaf1f0bf3b [flatmm] implement basic fp16 flatmm (#2089) 
* [flatmm] implement basic fp16 flatmm

* fix CI build fail

---------

Co-authored-by: root <root@hjbog-srdc-50.amd.com>
Co-authored-by: solin <bingzhou@amd.com>
 2025-04-16 16:51:17 +08:00
Andriy Roshchenko
7106976a72 MX GEMM - New GEMM pipeline for MX data types (#2059) 
* Allow selection of mfma_scale instructions

* Read B tensor from LDS to VGPR in chunks of 16 in MFMA order

* Add constexpr and synchronize return type for `get_exponent_value`

* Pass scales by reference and add comments to `mfma_scale_f32_32x32x64`

* Add support for microscaling instructions in `XdlopsGemm`

* Fix `mfma_scale_f32_16x16x128f8f6f4` wrapper

* Remove software implementation of MX GEMM

* Make interface of `intrin_mfma_scale_f32_16x16x128f8f6f4<16, 16>` consistent with the other scale instruction

* Update README

* Updated CHANGELOG

* Remove unused static methods
 2025-04-15 17:17:07 -06:00
Mingtao Gu
56378f810f CK pk_i4_t test failures fix (SWDEV-518629) (#2075) 
* fix pk_i4_v3 tests failures in Unbuntu env.

* fix pk_i4_t tests failure on Unbuntu issues.

* some fixed.

---------

Co-authored-by: mtgu0705 <mtgu@amd.com>
 2025-04-14 16:58:57 +08:00
Thomas Ning
269f4f6af5 Solve the Static Encoding Pattern compile error when the tile size is too small (#2079) 2025-04-13 20:09:30 -07:00
Illia Silin
0d4f145078 Fix build issues for multiple targets. (#2077) 
* build for multiple targets on gfx942

* add missing ignore statements
 2025-04-11 12:12:53 -07:00
jakpiase
6c61f4d237 [CK_TILE] Add 2:4 structured sparsity support for fp16 gemm (#1957) 
* add structured sparsity fp16 support for gemm

* added reviewer suggestions

* update changelog

* update changelog

* add reviewers suggestions

* Minor fix

* clang fix

* fix doxygen
 2025-04-11 12:18:26 +02:00
slippedJim
5f885d2b7a add fmha fwd splitkv receipt for aiter c++ api (#2068) 
* add s_randval for c++ api

* Fix bug of bias in splitkv

---------

Co-authored-by: rocking <ChunYu.Lai@amd.com>
 2025-04-10 23:21:13 +08:00
Juan Manuel Martinez Caamaño
f14e648e7c Replace inline assembly with builtins in FHMA (#2067) 
* Replace inline assembly with builtins in FHMA

---------

Co-authored-by: illsilin <Illia.Silin@amd.com>
 2025-04-10 09:48:37 +02:00
Illia Silin
3e6d21adeb enable gfx115x support (#2065) 2025-04-09 10:06:42 -07:00
MHYang-gh
03ce8729fd Make buffer coherence configurable in tensor view (#2041) 
* Make buffer coherence configurable in tensor view

* Fix clang-format for tensor_view.hpp
 2025-04-08 15:34:11 -07:00
valarLip
2c563fecf7 add passthrough for int32->float32 (#2062) 2025-04-08 15:16:30 -07:00
Max Podkorytov
6ce0797dad simplify generate_tuple (#2043) 2025-04-08 09:00:51 -07:00
aledudek
80aae6119b [CK_TILE] Fix GEMM Memory Pipeline (#2034) 
* [CK_TILE] Fix GEMM Memory Pipeline

* Fix transpose tile

* Add comments
 2025-04-08 12:40:04 +02:00
Illia Silin
1793228422 fix codegen issues (#2052) 2025-04-07 07:08:39 -07:00
Illia Silin
572cd820ce Split env.hpp header from the ck.hpp header. (#2049) 
* split env.hpp out of main headers

* fix namespace logic
 2025-04-03 15:30:21 -07:00
Rostyslav Geyyer
265af71a71 Add FP16/BF16<->FP8/BF8 conversions (#2035) 
* Move conversion functions and add missing conversions

* Add tests

* Add missing conversions

* Add missing conversions

* Add bf8 tests

* Update clipping for vectors

* Add missing conversions

* Add bf16 fp8 tests

* Add bf16 bf8 tests

* Fix device conversion

* Fix conversions

* Fix vector use

* Minor fix

* Add a workaround flag

* Add a workaround flag for bf16 conversion

* Add another workaround

* Add a workaround for fp16 to bf8 conversion

* Update type alias

* Add docstrings and missing wrappers

* Fix if defined macros

* Fix more if defined macros

* Add comments

* Remove __host__ specifier

* Add a gfx950 guard

* Update function naming
 2025-04-03 12:42:03 -05:00