ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-14 10:09:41 +00:00
Code Issues Packages Projects Releases Wiki Activity
1,859 Commits 742 Branches 38 Tags
9d5b89ce8a3fa80ff9c73e88fe2cd109fa7ce21f
Commit Graph

1859 Commits

This Branch
This Branch
All Branches
Author SHA1 Message Date
Aviral Goel
9d5b89ce8a [Tile Engine] Improved README.md (#2134) 
* improved tile_engine readme

* changed ck tile explanation and json

* further improved readme

* fixed typo

[ROCm/composable_kernel commit: 1aea51d34e]
 2025-04-29 17:37:07 -07:00
Max Podkorytov
55ebbab430 try building ck4inductor and testing it inside a virtual environment (#2142) 
use system virtualenv

use python-full ubuntu package in docker image

---------

Co-authored-by: illsilin <Illia.Silin@amd.com>
Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>

[ROCm/composable_kernel commit: 6601931949]
 2025-04-29 17:22:38 -07:00
Illia Silin
b4edc2ef70 Run CI jobs as user jenkins (#2141) 
* run CI as jenkins

* remove user jenkins from docker image

* move inductor installation to a writeable path

* add a switch for inductor tests

[ROCm/composable_kernel commit: 8fcb4dff1a]
 2025-04-29 07:35:10 -07:00
Aviral Goel
97c866807c Add Matrix A and Matrix B Swizzle for LDS in Computev4 policy (#2136) 
* fixed computev4 policy bug for lds swizzle

* added swizzle for input matrix B

* Improved ComputeV4 policy and pipeline by swizzling A and B

* consolidated LDS descriptor functions in parent struct

[ROCm/composable_kernel commit: 65f182d617]
 2025-04-28 18:20:47 -07:00
Khushbu Agarwal
7795e976da Support for MFMA_16x16x128 for fp8/bf8 (#2125) 
* Adding 16x16x128 support for gfx950

* Support for fp8 and bf8

* fix input arguments for MFMA scale instruction

* clang-formatted

* Fixes for lwpck-3145 (#2138)

* Fix lds tile & cmake dep & default epilogue

* Fallback BTypeToUse to ADataType in WOQ cases

* reverting instance json file

* reverting instance json file

---------

Co-authored-by: Yi DING <yi.ding@amd.com>

[ROCm/composable_kernel commit: d107f3c3a5]
 2025-04-28 18:19:50 -07:00
Khushbu Agarwal
a75ab12f3a [TileEngine] Support for sparsity in codegen (#2128) 
* Added sparsity flag in codegen

* remove comments

* clan formatted

* added sparsity as runtime argument

* updated README

* updated stream config variable

* fix typo for tail_num in hot loop

[ROCm/composable_kernel commit: 768c99eca9]
 2025-04-28 18:19:23 -07:00
Bartłomiej Kocot
05f9b2dde3 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

[ROCm/composable_kernel commit: 4094ad158a]
 2025-04-28 23:54:49 +02:00
Daniel Su
02ef8bcfb1 Check max-ilp-scheduling compiler option for moe_gemm examples (#2127) 
[ROCm/composable_kernel commit: d9786f3363]
 2025-04-28 13:40:22 -07:00
jakpiase
d48f3b585b Add ck tile examples to package (#1880) 
* add ck tile examples to package

* Update jenkinsfile

* fix for jenkinsfile

* fix for building ck tile code on non gfx9

* compile ck tile examples only for gfx94

* include ck tile examples in all target

* fix for basic gemm UseStructuredSparsity

* Update CMakeLists.txt

* Update gemm_pipeline_problem.hpp

* add targets to rocm install

---------

Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>

[ROCm/composable_kernel commit: 434d19f696]
 2025-04-28 09:53:19 -07:00
lalala-sh
ffc41f64fd fix moe i4 example bug (#2139) 
[ROCm/composable_kernel commit: 83394e40d2]
 2025-04-28 09:49:31 -07:00
Anton Gorenko
db016cf6da 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 5585c3121e

* 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
790e7bbd4b 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
7b61b0e4fa Avoid using store_tile_raw() for fp32 tensors (#2072) 
[ROCm/composable_kernel commit: 3d4d70d2fc]
 2025-04-26 23:07:41 -07:00
joyeamd
c27b313602 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
9b67733bd1 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
93caa99277 make code compliant with std=c++20 (#2123) 
[ROCm/composable_kernel commit: 01cb8379cd]
 2025-04-24 10:14:52 -07:00
alexxu-amd
ade20b5ac2 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
cb729623b5 Only generate specific hdim (#2120) 
[ROCm/composable_kernel commit: 02ce6d39ea]
 2025-04-24 18:52:58 +08:00
carlushuang
0e2d4899d4 [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
c644088d7b Update CODEOWNERS (#2119) 
[ROCm/composable_kernel commit: 854159fd00]
 2025-04-23 10:25:41 -07:00
lalala-sh
7748793e09 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
03cdc5602a Adding include directory in tile_engine (#2116) 
[ROCm/composable_kernel commit: 94662b02d0]
 2025-04-22 15:55:19 -07:00
Gino Lu
d20a94b709 [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
9a8af1c82a Temporarily disable MX FP4 device tests (#2112) 
[ROCm/composable_kernel commit: 416e851584]
 2025-04-22 16:08:48 -05:00
Thomas Ning
79b01b0824 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
7e42b9a814 update code owner (#2113) 
[ROCm/composable_kernel commit: 4bef60aa57]
 2025-04-21 13:53:03 -07:00
Muhammed Emin Ozturk
16fd74e8d3 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
e3f4d5d28f 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
ea9d075348 fix daily gfx942 build (#2106) 
[ROCm/composable_kernel commit: ce61759538]
 2025-04-21 08:48:22 -07:00
Khushbu Agarwal
790dfe9bcd 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
4c9c2aa6e0 fix CI build fail 
[ROCm/composable_kernel commit: c318ec0778]
 2025-04-21 16:00:12 +08:00
lalala-sh
2d0b5aba13 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
6d0890b6f4 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
348760d56e 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
bf57a9a5f3 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
c26c2b1fdc 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
4ec293cb4b [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
4c44fa374e 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
19a9980cd5 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
d9c9f17c3d 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
49e4ad3278 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
e8db9f0220 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
1b61d3a0ed 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
90612d0e37 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
a7a4456753 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
d76ebf9795 [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
959225947a 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
7a42b06988 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
7546e4bafe enable gfx115x support (#2065) 
[ROCm/composable_kernel commit: 3e6d21adeb]
 2025-04-09 10:06:42 -07:00
MHYang-gh
62ce5b906b 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