ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-14 18:17:44 +00:00
Code Issues Packages Projects Releases Wiki Activity
1,841 Commits 743 Branches 38 Tags
cfa85fa761fbd6af03316129fe5bebb4393a2fc3
Commit Graph

1841 Commits

This Branch
This Branch
All Branches
Author SHA1 Message Date
lalala-sh
cfa85fa761 fix moe i4 example bug (#2139) 
[ROCm/composable_kernel commit: 83394e40d2]
 2025-04-28 09:49:31 -07:00
Anton Gorenko
420de0f22d 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 cc666c6a19dabc2cce8141e7ae23bd460ceef331

* 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>

[ROCm/composable_kernel commit: edd92fc546]
 2025-04-28 10:14:21 +05:00
Yi DING
77c7fb1e6b Fix fp8 convert & add option for basic example (#2129) 
[ROCm/composable_kernel commit: 8add2cf45d]
 2025-04-27 16:26:05 -07:00
Po Yen Chen
88d4c6baab Avoid using store_tile_raw() for fp32 tensors (#2072) 
[ROCm/composable_kernel commit: 3d4d70d2fc]
 2025-04-26 23:07:41 -07:00
joyeamd
32053c0b27 SWDEV-52596 for hdim=256, when use splitkv pipeline, two new pipelines need to be added (#2126) 
[ROCm/composable_kernel commit: 41541aff7a]
 2025-04-25 16:31:09 +08:00
Khushbu Agarwal
5c17305a82 MFMA_32x32x16 for gfx950 (#2121) 
* Enable MFMA_32x32x16 for fp16/BF16 for gfx950

* clang formatted

[ROCm/composable_kernel commit: a2ed34a112]
 2025-04-24 10:20:22 -07:00
Illia Silin
08f859b460 make code compliant with std=c++20 (#2123) 
[ROCm/composable_kernel commit: 01cb8379cd]
 2025-04-24 10:14:52 -07:00
alexxu-amd
441a7d9847 Setup Doxygen API reference for Docs (#2115) 
* setup Doxygen settings

* add api_reference to requirements.txt

* add doxygen file header

* omit latex generation

* remove testing entry

* update Doxyfile

[ROCm/composable_kernel commit: ba97363acd]
 2025-04-24 09:35:06 -06:00
rocking
8100d99c5d Only generate specific hdim (#2120) 
[ROCm/composable_kernel commit: 02ce6d39ea]
 2025-04-24 18:52:58 +08:00
carlushuang
65199be10e [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>

[ROCm/composable_kernel commit: 5487289fc4]
 2025-04-23 12:40:18 -07:00
John Afaganis
fd63cc6e10 Update CODEOWNERS (#2119) 
[ROCm/composable_kernel commit: 854159fd00]
 2025-04-23 10:25:41 -07:00
lalala-sh
8f426b1216 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>

[ROCm/composable_kernel commit: 39ba03f25d]
 2025-04-23 10:35:34 +08:00
Khushbu Agarwal
c6b7f48326 Adding include directory in tile_engine (#2116) 
[ROCm/composable_kernel commit: 94662b02d0]
 2025-04-22 15:55:19 -07:00
Gino Lu
983dac1699 [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>

[ROCm/composable_kernel commit: 504f563f78]
 2025-04-22 15:52:36 -07:00
Rostyslav Geyyer
dbcaae42bf Temporarily disable MX FP4 device tests (#2112) 
[ROCm/composable_kernel commit: 416e851584]
 2025-04-22 16:08:48 -05:00
Thomas Ning
005e61ce63 GEMM Multiply Multiply Fix (#2102) 
* fix the type convert and increase the BF16 conversion + the profile comment

* fix the CI

[ROCm/composable_kernel commit: 0cca8fa28f]
 2025-04-22 01:13:22 -07:00
Thomas Ning
863ec4eb88 update code owner (#2113) 
[ROCm/composable_kernel commit: 4bef60aa57]
 2025-04-21 13:53:03 -07:00
Muhammed Emin Ozturk
eb8306b5f0 MI308 fix for streamk 1-Tile floating point exception (#2101) 
[ROCm/composable_kernel commit: b092c18da7]
 2025-04-21 11:44:07 -07:00
Thomas Ning
83197c769f 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>

[ROCm/composable_kernel commit: a738e43445]
 2025-04-21 10:21:35 -07:00
Illia Silin
b29fc0efce fix daily gfx942 build (#2106) 
[ROCm/composable_kernel commit: ce61759538]
 2025-04-21 08:48:22 -07:00
Khushbu Agarwal
74210a9dfc multi instance generation for CkTileEngine (#2080) 
* Add support for multi-instance verification, print detail for each instance, documentation fix

* clang formatted

* Added Readme file

* updated readme

* Addressing review comments

* clang formatted

* Updated ReadMe and GPU reference code

* simplified dispatch kernel code

* indentation

[ROCm/composable_kernel commit: 7cadf187e2]
 2025-04-21 08:39:45 -07:00
solin
dd2c3289c9 fix CI build fail 
[ROCm/composable_kernel commit: c318ec0778]
 2025-04-21 16:00:12 +08:00
lalala-sh
dc58110b06 enable do top k weights in moe stage1 gemm (#2094) 
* add switch for mul topk weights

* fix bf16/f16 bugs

* complete

[ROCm/composable_kernel commit: bcf5bb41be]
 2025-04-18 10:45:49 +08:00
Andriy Roshchenko
7972a39081 MX GEMM - Parameterized Test Template (#2088) 
* Tests for MX FP8 GEMM

* Improve documentation


[ROCm/composable_kernel commit: 213b203a3c]
 2025-04-16 19:56:00 -06:00
Andriy Roshchenko
a96e96ab4e 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


[ROCm/composable_kernel commit: da54464cce]
 2025-04-16 15:25:02 -06:00
Illia Silin
729c668e9d Upgrade default docker to Ubuntu24.04 (#2090) 
* upgrade docker to Ubuntu24.04

* add break-system-packages flag to pip install

* fix dockerfile

[ROCm/composable_kernel commit: 3bb62f16cd]
 2025-04-16 12:10:15 -07:00
aledudek
d02e727b51 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>

[ROCm/composable_kernel commit: 7c32652e03]
 2025-04-16 11:00:55 +02:00
BingYuan.Zhou
f6b51a1c7b [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>

[ROCm/composable_kernel commit: eaf1f0bf3b]
 2025-04-16 16:51:17 +08:00
felix
ad04db9d00 add preshuffle gemm fp16 (#2036) 
* add preshuffle gemm fp16

* clang format and test ok

* Update gemm_multiply_multiply_xdl_fp16_bpreshuffle.cpp

remove useless comments in example

* Update gemm_multiply_multiply_xdl_fp16_bpreshuffle.cpp

remove 2

---------

Co-authored-by: coderfeli <coderfeli@163.com>

[ROCm/composable_kernel commit: c5975529bb]
 2025-04-16 10:53:21 +08:00
joyeamd
b1afc03c6a fmha hdim256 vectorize improve (#2086) 
For hdim 256, will not have vectorized buffer load when seqlen % 256 != 0 and hdim % 256 = 0; this commit tries to solve this condition.

[ROCm/composable_kernel commit: 94d47b1680]
 2025-04-16 09:21:04 +08:00
Andriy Roshchenko
5e2bd20672 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

[ROCm/composable_kernel commit: 7106976a72]
 2025-04-15 17:17:07 -06:00
Illia Silin
1a8132e9f9 Upgrade default docker image to ROCm6.4 release. (#2082) 
* upgrade to rocm6.4

* fix gfx10 generic target syntax

* use gfx1101 target for unit tests

* use gfx1201 target for unit tests

* do not use generic targets until 6.4.1 release

* update target list and dockerfile.compiler

[ROCm/composable_kernel commit: d55c9cb313]
 2025-04-14 16:41:47 -07:00
Mingtao Gu
3cbda9a11c 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>

[ROCm/composable_kernel commit: 56378f810f]
 2025-04-14 16:58:57 +08:00
Thomas Ning
1d62f5edb9 Solve the Static Encoding Pattern compile error when the tile size is too small (#2079) 
[ROCm/composable_kernel commit: 269f4f6af5]
 2025-04-13 20:09:30 -07:00
Illia Silin
c91d046350 Fix build issues for multiple targets. (#2077) 
* build for multiple targets on gfx942

* add missing ignore statements

[ROCm/composable_kernel commit: 0d4f145078]
 2025-04-11 12:12:53 -07:00
Muhammed Emin Ozturk
342df94444 CkProfiler StreamK GemmUniversal Fix and Split Gemm_universal Test Redo PR #2044 (#2070) 
* fix and split gemm_universal test


* Update test_gemm_universal_streamk_ut_cases_fp8.inc

[ROCm/composable_kernel commit: 74fda2e796]
 2025-04-11 10:17:29 -07:00
jakpiase
addcd203eb [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

[ROCm/composable_kernel commit: 6c61f4d237]
 2025-04-11 12:18:26 +02:00
slippedJim
cca9cca699 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>

[ROCm/composable_kernel commit: 5f885d2b7a]
 2025-04-10 23:21:13 +08:00
Juan Manuel Martinez Caamaño
c4e0659e51 Replace inline assembly with builtins in FHMA (#2067) 
* Replace inline assembly with builtins in FHMA

---------

Co-authored-by: illsilin <Illia.Silin@amd.com>

[ROCm/composable_kernel commit: f14e648e7c]
 2025-04-10 09:48:37 +02:00
Illia Silin
9cc561987f enable gfx115x support (#2065) 
[ROCm/composable_kernel commit: 3e6d21adeb]
 2025-04-09 10:06:42 -07:00
MHYang-gh
4e5b335b5e Make buffer coherence configurable in tensor view (#2041) 
* Make buffer coherence configurable in tensor view

* Fix clang-format for tensor_view.hpp

[ROCm/composable_kernel commit: 03ce8729fd]
 2025-04-08 15:34:11 -07:00
valarLip
aa6eeb7f66 add passthrough for int32->float32 (#2062) 
[ROCm/composable_kernel commit: 2c563fecf7]
 2025-04-08 15:16:30 -07:00
Khushbu Agarwal
8a65c29ffb 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>

[ROCm/composable_kernel commit: 263ff689e0]
 2025-04-08 15:14:53 -07:00
spolifroni-amd
6491cd870a fixed broken github link (#2063) 
[ROCm/composable_kernel commit: 2c8132126c]
 2025-04-08 10:20:31 -07:00
dependabot[bot]
409f3dccc6 Bump rocm-docs-core from 1.18.1 to 1.18.2 in /docs/sphinx (#2047) 
Bumps [rocm-docs-core](https://github.com/ROCm/rocm-docs-core) from 1.18.1 to 1.18.2.
- [Release notes](https://github.com/ROCm/rocm-docs-core/releases)
- [Changelog](https://github.com/ROCm/rocm-docs-core/blob/develop/CHANGELOG.md)
- [Commits](https://github.com/ROCm/rocm-docs-core/compare/v1.18.1...v1.18.2)

---
updated-dependencies:
- dependency-name: rocm-docs-core
  dependency-version: 1.18.2
  dependency-type: direct:production
  update-type: version-update:semver-patch
...

Signed-off-by: dependabot[bot] <support@github.com>
Co-authored-by: dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>

[ROCm/composable_kernel commit: b12cd6580b]
 2025-04-08 09:06:38 -07:00
Max Podkorytov
1eec5cc073 simplify generate_tuple (#2043) 
[ROCm/composable_kernel commit: 6ce0797dad]
 2025-04-08 09:00:51 -07:00
aledudek
a42081842c [CK_TILE] Fix GEMM Memory Pipeline (#2034) 
* [CK_TILE] Fix GEMM Memory Pipeline

* Fix transpose tile

* Add comments

[ROCm/composable_kernel commit: 80aae6119b]
 2025-04-08 12:40:04 +02:00
Illia Silin
09ace4ab2f Fix a couple of CI issues. (#2050) 
* fix jenkins jobs

* fix perf log name for gfx908

* only run gemm perf tests on gfx908

[ROCm/composable_kernel commit: 72c0261ef1]
 2025-04-07 12:48:34 -07:00
Illia Silin
ad65c947c9 fix codegen issues (#2052) 
[ROCm/composable_kernel commit: 1793228422]
 2025-04-07 07:08:39 -07:00
Illia Silin
6346304629 Revert "CkProfiler StreamK GemmUniversal Fix and Split Gemm_universal Test (…" (#2054) 
This reverts commit cddc580bae98ff61b4d429c0ba2fe510e060a6de.

[ROCm/composable_kernel commit: 29f7266216]
 2025-04-07 06:49:36 -07:00