ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-20 12:59:49 +00:00
Code Issues Packages Projects Releases Wiki Activity
2,374 Commits 749 Branches 38 Tags
7b68b6eebbd0cf4db40484e4ef41ba4e1aa9c5fc
Commit Graph

2374 Commits

This Branch
This Branch
All Branches
Author SHA1 Message Date
Khushbu Agarwal
94662b02d0 Adding include directory in tile_engine (#2116) 2025-04-22 15:55:19 -07: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
Qianfeng Zhang
022ed3fd8a Back to use exp() instead of exp2() since exp() in ck_tile using fast __builtin_amdgcn_exp2f() 2025-04-22 14:47:24 +00:00
Qianfeng Zhang
7316a44ff3 Update exp() in ck_tile/core/numeric/math.hpp to use __expf 2025-04-22 14:46:53 +00:00
Qianfeng Zhang
26db7e0b7c Use kN0=64 to save vgprs 2025-04-22 14:45:27 +00:00
Qianfeng Zhang
65ddb1a863 Fix the script name 2025-04-22 13:46:48 +00:00
Qianfeng Zhang
58ab5533a6 Fix in GetTileRangeAlongX 2025-04-22 13:46:23 +00:00
Qianfeng Zhang
677fd60d10 Add script compare_with_triton_2.sh for measuring the jagged cases of seqlen 1024/2048/4096/8192/16384/32768 2025-04-22 13:45:48 +00: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
Thomas Ning
4bef60aa57 update code owner (#2113) 2025-04-21 13:53:03 -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
Illia Silin
ce61759538 fix daily gfx942 build (#2106) 2025-04-21 08:48:22 -07:00
Khushbu Agarwal
7cadf187e2 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
 2025-04-21 08:39:45 -07:00
solin
c318ec0778 fix CI build fail 2025-04-21 16:00:12 +08:00
Qianfeng Zhang
2546e905ce Change gemm0 to iterate along kN0 so that BlockGemm can overlap with maksing and siLu 2025-04-20 13:23:15 +00:00
Qianfeng Zhang
ee259a8924 Fix the GetTileRangeAlongX() to align with the hstu masking definition when both causal=true and local=true 2025-04-18 15:37:49 +00:00
Qianfeng Zhang
efc786f6a3 Remove un-needed __builtin_amdgcn_sched_barrier(0) 2025-04-18 10:05:57 +00:00
Qianfeng Zhang
88e54a8989 Use shared ring Lds buffers for K/V to avoid over-lapping between first-K/last-V or last-K/first-V 2025-04-18 09:47:43 +00:00
Qianfeng Zhang
f12a47218f Tiny codes simplification in pipeline 2025-04-18 08:22:11 +00: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
Qianfeng Zhang
ca1ae84fc6 Remove one line of __builtin_amdgcn_sched_barrier(0) 2025-04-17 14:21:14 +00:00
Qianfeng Zhang
b0ae27046f Fix the integer overflow in total_flops calculation 2025-04-17 10:34:13 +00:00
Qianfeng Zhang
6086ead2f9 Add scripts for comparing with triton 2025-04-17 10:33:44 +00: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
Illia Silin
3bb62f16cd Upgrade default docker to Ubuntu24.04 (#2090) 
* upgrade docker to Ubuntu24.04

* add break-system-packages flag to pip install

* fix dockerfile
 2025-04-16 12:10:15 -07: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
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
Qianfeng Zhang
1351d9cd1b Use exp2() to calculate exp() for better performance 2025-04-16 06:54:06 +00:00
Qianfeng Zhang
226a254723 Remove the comparing of row/col to max_uih_len in masking 2025-04-16 04:35:42 +00:00
felix
c5975529bb 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>
 2025-04-16 10:53:21 +08:00
joyeamd
94d47b1680 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.
 2025-04-16 09:21:04 +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
Qianfeng Zhang
d1749b3aae Use kM0=128 kN0=64 to completely remove the vgprs spilling 2025-04-15 15:08:46 +00:00
Qianfeng Zhang
3cd1b13e46 Split HstuBlockMasking into HstuBlockMaskWithLocal and HstuBlockMaskNoLocal to save vgprs for non-local situations 2025-04-15 14:40:55 +00:00
Qianfeng Zhang
cad1356170 Use packed cast_tile for fp16 2025-04-15 14:29:30 +00:00
Qianfeng Zhang
fff13b6c76 Update to partially reduce the register spilling 2025-04-15 07:44:33 +00:00
Illia Silin
d55c9cb313 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
 2025-04-14 16:41:47 -07: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
Qianfeng Zhang
c2e6ab8516 Add IsFirstVLdsBufferOverlapLastKLdsBuffer() check to reduce call of s_barrier() 2025-04-13 11:00:22 +00:00
Qianfeng Zhang
238e78d82e Update the in pipeline codes 2025-04-13 09:43:58 +00:00
Qianfeng Zhang
53e567977e Fix in calculation of total_flops and update benchmark scripts 2025-04-13 08:50:00 +00: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
Muhammed Emin Ozturk
74fda2e796 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
 2025-04-11 10:17:29 -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