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
|
4ae9acd712
|
Revert "Update exp() in ck_tile/core/numeric/math.hpp to use __expf"
This reverts commit 7316a44ff3.
|
2025-04-25 14:42:33 +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 |
|
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 |
|
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 |
|
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 |
|
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 |
|
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 |
|
Qianfeng Zhang
|
dd2cd2cbcb
|
Tune the input initialization to avoid over-flow in silu
|
2025-04-09 10:03:32 +00:00 |
|
Qianfeng Zhang
|
86c0e45987
|
Add benchmark_hstu_attention.sh
|
2025-04-09 08:28:05 +00:00 |
|
Qianfeng Zhang
|
9cb2dca958
|
Add several verification test cases
|
2025-04-08 16:38:35 +00:00 |
|
Qianfeng Zhang
|
561d490990
|
Fix in kernel and forward dispatch for jagged mode
|
2025-04-08 16:37:52 +00:00 |
|