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,926 Commits 749 Branches 38 Tags
3f9f2fa736dfcabff3f0558d5c1611b91efb7944
Commit Graph

635 Commits

Author SHA1 Message Date
Yi DING
8add2cf45d Fix fp8 convert & add option for basic example (#2129) 2025-04-27 16:26:05 -07:00
Qianfeng Zhang
1af27022ef Add IsFullTileInsideMask() to avoid pixel-by-pixel checking when kUseCausl=true but kUseLocal=false 2025-04-27 09:31:38 +00:00
Po Yen Chen
3d4d70d2fc Avoid using store_tile_raw() for fp32 tensors (#2072) 2025-04-26 23:07:41 -07:00
Qianfeng Zhang
054c397e05 Replace set_tile_if() by sweep_tile_span() to reduce branching 2025-04-27 05:00:09 +00:00
Qianfeng Zhang
95c93ba92e Update the GridSize() and GetTileIndex() in hstu kernel 2025-04-26 10:01:23 +00:00
Qianfeng Zhang
1b463e915d Add scripts for measuring jagged with/no causal cases 2025-04-25 15:59:51 +00:00
Qianfeng Zhang
9996270087 Tiny update in IsTokenPairInsideMask() 2025-04-25 15:36:58 +00:00
Qianfeng Zhang
27f7ab4f2c Use compiler builtin directly in f_silu for float type 2025-04-25 15:05:39 +00:00
Qianfeng Zhang
80677eb6e0 Code re-arrangement in pipeline 2025-04-25 14:41:37 +00:00
Qianfeng Zhang
4a49119d98 Update the seqlen_k_curr inside the first gemm loop 2025-04-25 13:59:48 +00:00
joyeamd
41541aff7a SWDEV-52596 for hdim=256, when use splitkv pipeline, two new pipelines need to be added (#2126) 2025-04-25 16:31:09 +08:00
Qianfeng Zhang
7818cce1c3 Rename the performance measurement scripts 2025-04-25 06:09:17 +00:00
Qianfeng Zhang
05910ebe0b Add support for WarpGem-16x16x32 in QK-BlockGemm (which enables using ds_write/read_b128 for K 2025-04-25 06:06:50 +00:00
Qianfeng Zhang
a41371f734 Update in K-Lds laying-out to consider for both WarpGemm-32x32x16 and WarpGemm-16x16x16 2025-04-24 15:02:57 +00:00
rocking
02ce6d39ea Only generate specific hdim (#2120) 2025-04-24 18:52:58 +08:00
Qianfeng Zhang
cea919aefb Use 16x16x16 WarpGemm 2025-04-24 08:14:09 +00:00
Qianfeng Zhang
7848d15d39 Using __builtin_amdgcn_rcpf in siLU function 2025-04-24 06:28:16 +00:00
Qianfeng Zhang
aec19176d4 Combine minus with scale_s 2025-04-24 05:47:24 +00:00
Qianfeng Zhang
ce4665262b Move silu calculation to gemm1 iteration and try to interleave gemm_1 and silu 2025-04-24 04:49:58 +00:00
Qianfeng Zhang
2d2e1941a8 Update in using masking for the case where kMasking is false and kPadSeqLenK is true 2025-04-23 10:47:27 +00:00
Qianfeng Zhang
8dcde8d10f Fix in generate_instances.py and re-generated the instances 2025-04-23 10:30:40 +00: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
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
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
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
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
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
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
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
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
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
Qianfeng Zhang
71697d9cb9 Add output of estimated TFLOPS 2025-04-09 14:50:18 +00:00
Qianfeng Zhang
1766e6d3be Update to the scripts and error thresholds 2025-04-09 10:34:37 +00:00