ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-15 18:42:06 +00:00
Code Issues Packages Projects Releases Wiki Activity
1,863 Commits 745 Branches 38 Tags
29cf1610f14e06ca8a652d80fd83c40bba78218c
Commit Graph

205 Commits

Author SHA1 Message Date
Qianfeng Zhang
29cf1610f1 Enable RTN fp32 to bf16 conversion by adding compiler option in CMakeLists.txt 2025-05-20 07:52:01 +00:00
Qianfeng Zhang
fac03abbda Change do-while main-loop to while-do and remove early exiting check 2025-05-19 15:39:31 +00:00
Qianfeng Zhang
14ab6f154d Adjust the codes before the main-loop 2025-05-19 13:59:57 +00:00
Qianfeng Zhang
f411d676f2 Move k_tile loading and v_tile loading earlier in the loop 2025-05-19 13:59:24 +00:00
Qianfeng Zhang
902b1c645c Move k_tile loading in the loop earlier 2025-05-19 13:58:43 +00:00
Qianfeng Zhang
f582c21418 Replace s_acc and pcomp tile array by single tile object for simplification 2025-05-19 13:57:11 +00:00
Qianfeng Zhang
4e65469fe8 Add _builtin_amdgcn_sched_barrier(0) for instructing the compiler for better codes isolation 2025-05-18 16:20:47 +00:00
Qianfeng Zhang
e4e70f8b0a Set the block_per_cu to 3 for hdim-128 2025-05-18 15:58:57 +00:00
Qianfeng Zhang
ff3415d97d Prefetch b_warp_tensor for next nIter and move b_warp_windows construction into n-iteration in block_gemm_areg_bsmem_creg for gemm-1 2025-05-18 15:26:31 +00:00
Qianfeng Zhang
694295a9d3 Move b_warp_windows construction into k-iteration in block_gemm_areg_bsmem_creg for gemm-0 2025-05-18 15:25:43 +00:00
Qianfeng Zhang
afd7793e92 Prefetch K for next iteration from LDS in block_gemm_areg_bsmem_creg for gemm-0 2025-05-18 15:25:05 +00:00
Qianfeng Zhang
7c0ac51b4b Hack block_gemm_areg_bsmem_creg_v2 for gemm_1 2025-05-18 15:24:17 +00:00
Qianfeng Zhang
473fbc374b Rename the hacked block_gemm_areg_bsmem_creg_v2 2025-05-18 15:23:40 +00:00
Qianfeng Zhang
58e45ec53a Move the lambda for dividing by max_seqlen from kernel to pipeline 2025-05-18 07:58:52 +00:00
Qianfeng Zhang
0771390a28 Move the dividing by max_seqlen out of f_silu to be handle outside the main-loop 2025-05-18 03:41:31 +00:00
Qianfeng Zhang
586968785a Set example option -save_mask default to 0 2025-05-14 13:56:54 +00:00
Qianfeng Zhang
b0d3704390 Add scripts (test_ck_hstu_mask.sh and test_pytorch_hstu_mask.py) for checking mask 2025-05-14 02:00:22 +00:00
Qianfeng Zhang
5b0a2618fd Add -save_mask option to the example to output int8 mask tensor 2025-05-14 01:54:45 +00:00
Qianfeng Zhang
c3761c3bd6 Update the rules of hstu masking 2025-05-13 10:37:19 +00:00
Qianfeng Zhang
3a320bcdf6 Add test cases for better functional verification 2025-05-10 16:04:06 +00:00
Qianfeng Zhang
79cd1f0653 Fix sequence dim length for o_dram descriptor in the kernel 2025-05-10 16:02:52 +00:00
Qianfeng Zhang
1d1dd8f1eb Revert "Temporarily close the instance for hdim64 and hdim256 to save compiling time" 
This reverts commit 2972de4c88.
 2025-05-07 13:37:15 +00:00
Qianfeng Zhang
d32851e15c Simplification in the static iterations of block_gemm_areg_bsmem_creg_v2_hack 2025-05-07 10:09:23 +00:00
Qianfeng Zhang
632fd06a7a Use kK1=16 2025-05-07 09:51:16 +00:00
Qianfeng Zhang
079f7e3a03 Use type_convert rather than static_cast in f_silu 2025-05-07 07:11:54 +00:00
Qianfeng Zhang
72d55d1b40 Add max_seqlen as divider in siLu 2025-05-06 16:18:52 +00:00
Qianfeng Zhang
374e0626e6 Remove using cast_tile_pk_fp16_fp32 for better accuracy for fp16 hstu attention 2025-05-06 08:24:03 +00:00
Qianfeng Zhang
611f2ce1f9 Override and fix GetAlignmentK() 2025-05-03 16:23:54 +00:00
Qianfeng Zhang
da89540ee0 Use kN0=32 2025-04-30 05:42:43 +00:00
Qianfeng Zhang
2972de4c88 Temporarily close the instance for hdim64 and hdim256 to save compiling time 2025-04-30 02:20:41 +00:00
Qianfeng Zhang
d63dab90da Hack block_gemm_areg_bsmem_creg_v2 to let s_acc for gemm_0 not need be cleared first 2025-04-28 15:24:13 +00:00
Qianfeng Zhang
f1f4e249a6 Adjust the v_tile and k_tile loading location 2025-04-28 09:25:09 +00:00
Qianfeng Zhang
f53be61a74 Put two gemms call inside one n0loop unroll 2025-04-28 06:41:37 +00: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
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
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
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