ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-09 00:27:01 +00:00
Code Issues Packages Projects Releases Wiki Activity
1,955 Commits 777 Branches 40 Tags
da5f962dfd120426e4c1fee0a1ae72dbe7bdb79e
Commit Graph

467 Commits

Author SHA1 Message Date
Anton Gorenko
2679c0aa2b Disable all tests and instances with f8 on gfx11 
Even though f8_f16 and f16_f8 don't require f8 WMMA instructions,
gfx11 still lacks hardware instructions for fast f8->f32 conversion.
 2025-05-23 15:17:36 +05:00
Anton Gorenko
9e70603bb5 Use PermuteB with f16_i4 in most instances (as xdl) 
Some instances use PermuteB = false for checking correctness.
See also the previous commit.
 2025-05-23 10:53:42 +05:00
Anton Gorenko
02bf56a6de Add missing add_device_gemm_wmma_universal_f8_f8_bf16 declarations 2025-05-23 10:53:42 +05:00
Anton Gorenko
b1f50b5ba7 Remove instances that are too slow (mostly because of register spilling) 2025-05-23 10:53:42 +05:00
Anton Gorenko
5082a9cd69 Add instances for Interwave and Intrawave v1 2025-05-23 10:53:41 +05:00
Anton Gorenko
17bc0fa4c5 Remove instances that are too slow (mostly because of register spilling) 2025-05-23 10:53:41 +05:00
Anton Gorenko
04d3fc7b07 Add instances with ABlockLdsExtraM and BBlockLdsExtraN = 0 2025-05-23 10:53:41 +05:00
Anton Gorenko
c94c3b4716 Add instances for Interwave and Intrawave v1 2025-05-23 10:53:41 +05:00
Anca Hamuraru
185ea0fdc4 Fix typo 2025-05-20 13:45:48 +00:00
Anca Hamuraru
0a5e6d48ec Fixed typos 2025-05-20 13:16:12 +00:00
Anca Hamuraru
a08ca63b36 Padding not supported for when BDataType is pk_i4_t. Added fix for correct check and removed padding instances. 2025-05-20 13:14:07 +00:00
Anca Hamuraru
7975c9cea6 Fixed typos for padded instances 2025-05-19 14:51:56 +00:00
Anca Hamuraru
a0a2bf27b1 Fixed typos for bf16xbf16xbf16 padding instances 2025-05-19 13:08:05 +00:00
Anca Hamuraru
bbda71f096 Added padding instances for f16xf8xf16, f16xi4xf16 2025-05-16 08:31:32 +00:00
Anca Hamuraru
e97a7f246d Added padding instances for f8xf16xf16 2025-05-15 11:31:50 +00:00
Anca Hamuraru
537d19c459 Added packed instances for bf16xi4xbf16 2025-05-15 10:36:50 +00:00
Anca Hamuraru
a3d99c696d Merge branch '32-add-the-remaining-combination-of-data-types-mixed-precision-for-gemm_universal-pipeline3-for' into '33-wip' 
Added mixed precision instances for gemm_universal

See merge request amd/ai/composable_kernel!11
 2025-05-15 06:53:55 +00:00
Anca Hamuraru
0482b83fe6 Added padding instances for bf16xbf16, f8xf8 2025-05-14 14:36:10 +00:00
Anca Hamuraru
e2855633ee Fixes for padding instances for f16xf16xf16 2025-05-14 10:50:34 +00:00
Anca Hamuraru
ab28ac0b81 Merge branch '33-wip' of projects.streamhpc.com:amd/ai/composable_kernel into 33-wip 2025-05-14 09:43:36 +00:00
Anca Hamuraru
647024db6f Ammending changes for adding support for padding instances for f16xf16xf16 2025-05-14 09:41:19 +00:00
Anca Hamuraru
8b52033aae Fixed cmake errors related to gemm_bilinear. Previously, if the above flags are set, cmake build fails: GPU_TARGETS="gfx1100;gfx1201" -D DTYPES="fp16;bf16;fp8" 
(cherry picked from commit 5bdc993dbf)

Co-authored-by: Anca Hamuraru <anca@streamhpc.com>
 2025-05-14 08:45:34 +00:00
Anca Hamuraru
dd47f39a0a Added padding instances for f16xf16xf16 2025-05-14 07:54:46 +00:00
Bartłomiej Kocot
6fddb5708c Add grouped conv fwd bias relu instances (#2179) 
* Add grouped conv fwd bias relu instances

* fixes

* fix
 2025-05-09 22:52:34 +02:00
Andriy Roshchenko
cb27e7c77f Ensure MX GEMM Instances can be Cross-Compiled for Multiple Architectures (#2171) 
* Re-enable MX GEMM instances

* Fix compilation error when building MX GEMM for multiple architectures
 2025-05-08 13:26:03 -06:00
Anca Hamuraru
744262dd88 Pre-commit hooks fixes 2025-05-08 13:48:18 +00:00
Bartłomiej Kocot
397b9080a2 Move 16x16 grouped conv fwd instances from comp header (#2165) 
* Move 16x16 grouped conv fwd instances from comp header

* Improvements
 2025-05-07 08:04:31 -07:00
Rostyslav Geyyer
8a0d659f92 Add FP4 MX MFMA tests (#2151) 
* Add conversion tests

* Fix ctor

* Fix nan logic

* Fix conversion logic

* Permute packed f4_t values

* Fix conversion to float, repack vector elements

* Fix device tests

* Permute elements in a vector

* Add a repro test

* Add a conversion for a repro test

* Update test vectors

* Update conversion

* Fix the test

* Update test vector generator

* Fix vector sr conversion

* Permute conversion args

* Update conversion

* Test

* Fix packing

* Simplify conversion function

* Pack conversion in a loop

* Pack conversion in a loop

* Pack another conversion in a loop

* Pack one more conversion in a loop

* Pack the last conversion in a loop

* Clean up

* Add ops

* Add tests

* Add missing utils

* Update reference mx gemm

* Add f4x2 init mode

* Update host tensor utils

* Update chunk size for f4x2

* Add non scaled ops

* Add a type utility

* Update non scaled reference kernel

* Add non scaled tests

* Debug mfma arguments

* Add more debug info

* Update chunk size

* Update data layout

* Add more debugging

* Fix B stride

* Fix reference gemm

* Fix build

* One more reference fix

* Add more debug info

* Disable some tests

* Enable tests

* Add fp4 dimensions

* Update reference kernels

* Temp edits

* Remove leftovers

* Fix conflicts

* Clean up

* More clean up

* Revert "More clean up"

This reverts commit d8d35a0846.

* Add layouts to tests

---------

Co-authored-by: Andriy Roshchenko <107577548+andriy-ca@users.noreply.github.com>
 2025-05-06 09:24:00 -05:00
Andriy Roshchenko
79beaacdd1 Restrict MX GEMM instantiation to GFX950 arch (#2157) 2025-05-05 08:18:22 -07:00
Anca Hamuraru
55f160278b Added support for F16xI4xF16 to gemm_wmma_universal 2025-05-05 10:36:59 +00:00
Anca Hamuraru
c5d99d2295 Added support for BF16xI4xBF16 to gemm_wmma_universal 
(cherry picked from commit c6a4a69d2d43d59bae8bdabfae80d648646f217e)

Co-authored-by: Anca Hamuraru <anca@streamhpc.com>
 2025-05-05 09:55:29 +00:00
Anca Hamuraru
636fbd5517 Added support for F16xF8xF16 to gemm_wmma_universal 2025-05-05 08:47:30 +00:00
Andriy Roshchenko
79b0bfeb41 MX GEMM - Add FP8 GEMM Tests for Different Layouts (#2152) 
* Add gemm_mx_fp8_bf8 example with row-major B

* Add more overloads of MX MFMA instructions

* Add MK_KN (RRR) tests

* Add KM_NK (CCR) tests

* Add more problem sizes to Large tests

* Add test_gemm_mx to the list of regression tests
 2025-05-01 11:55:48 -06:00
Anca Hamuraru
977d8a6649 Added support for F8xF16xF16 to gemm_wmma_universal 
(cherry picked from commit 15c851de6daa513a12c2e3af299bab0176175fb5)

Co-authored-by: Anca Hamuraru <anca@streamhpc.com>
 2025-05-01 12:53:15 +00:00
Anca Hamuraru
03c0446ed2 Adding support for RRR, F8xF16xF16 gemm_universal_wmma - wip 
(cherry picked from commit f8c06322df0abcbd5945a56cdf5bffe56480f9f0)

Co-authored-by: Anca Hamuraru <anca@streamhpc.com>
 2025-05-01 12:52:18 +00:00
Anca Hamuraru
5bdc993dbf Fixed cmake errors related to gemm_bilinear. Previously, if the above flags are set, cmake build fails: GPU_TARGETS="gfx1100;gfx1201" -D DTYPES="fp16;bf16;fp8" 2025-04-30 13:02:57 +00:00
Bartłomiej Kocot
23de234dbe Add grouped conv fwd 16x16 mfma instruction instances (#2140) 
* Add grouped conv fwd 16x16 mfma instruction instances

* fix

* remove oddc

* fix

* fix
 2025-04-30 09:49:37 +02:00
Bartłomiej Kocot
4094ad158a Integrate universal gemm with conv bwd data and add SplitK (#1315) 
* Integrate universal gemm with conv bwd data

* Fix multi d kernel

* Add splitK support

* instances refactor

* instances refactor

* refactor

* fixeS

* fixes

* 16x16 instnaces

* Fixes

* Fix

* Fix

* Fix

* Fix

* Fix

* Fixes

* fix

* fix
 2025-04-28 23:54:49 +02: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
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
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
aledudek
7c32652e03 Add grouped conv fwd 3d GKCYX instances for f32, f16, bf16 (#2069) 
* Part1

* Add grouped conv fwd 3d GKCYX instances for f32, f16, bf16

* Add missing coma

* Add missing cpp instance files

* Fix 3d layout

* Add missing closing bracket

* Add missing comp x2 and part2 instances

* Fix typo in instance name

* fix

* Fix

---------

Co-authored-by: Bartlomiej Kocot <barkocot@amd.com>
 2025-04-16 11:00:55 +02:00
Illia Silin
3e6d21adeb enable gfx115x support (#2065) 2025-04-09 10:06:42 -07:00
Khushbu Agarwal
263ff689e0 New instances for gemm_multiply_multiply_weightpreshuffle operator (#2061) 
* Add new instances for weight_preshuffle for f8->bf16

* Add new instances for weight_preshuffle for f8->f16

* clang formatted

---------

Co-authored-by: Khushbu Agarwal <khuagar@amd.com>
Co-authored-by: Thomas Ning <Thomas.Ning@amd.com>
 2025-04-08 15:14:53 -07:00
Bartłomiej Kocot
2ccf914888 Add support for GKCYX grouped conv weight (#2023) 
* Grouped conv bwd weight GKCYX support

* fix and changelog

* fix

* fix

* fixes

* comments

* fix
 2025-04-02 23:59:49 +02:00
Bartłomiej Kocot
8c0ab61ece Grouped conv backward data GKCYX support (#2029) 
* Grouped conv backward data GKCYX support

* profiler

* Converter

* split instances
 2025-04-01 13:24:38 -07:00
Bartłomiej Kocot
6355ee7ca5 Improve compilation time for grouped conv fwd (#2039) 
* Improve compilation time for grouped conv fwd

* Fix
 2025-04-01 07:11:42 -07:00
Muhammed Emin Ozturk
dd4c12b155 f8/bf16 GEMM Stream-K (#1879) 2025-03-31 20:30:17 -06:00
jefyang1
16b15e336a Fix gemm universal and grouped_conv_fwd test failures on gfx950 (#2031) 2025-03-31 09:20:52 -07:00