Commit Graph

866 Commits

Author SHA1 Message Date
Qianfeng Zhang
d190af2ef5 Tiny fix in trload with_softmax/no_softmax pipeline 2025-11-05 14:44:13 +00:00
Qianfeng Zhang
99993acca4 Improve both the with_softmax and no_softmax pipelines 2025-11-04 15:46:49 +00:00
Qianfeng Zhang
bc22b83b19 Add kUseTrLoad = false in non-trload pipeline 2025-11-03 12:40:16 +00:00
Qianfeng Zhang
e40ab20b9e Clarifying the using of CK_TILE_HOST and CK_TILE_HOST_DEVICE trying to save compiling time 2025-11-03 09:40:03 +00:00
Qianfeng Zhang
e31829384d Change in updating max_uih_seqlen in the example 2025-11-02 03:59:46 +00:00
Qianfeng Zhang
39cb8c33d1 Use supplement_array_by_last_element(num_targets, ) in example 2025-11-02 03:34:54 +00:00
Qianfeng Zhang
80e08b6efe Use supplement_array_by_last_element() in example to simplify the codes 2025-11-02 03:17:16 +00:00
Qianfeng Zhang
10133e5d51 Update to README.md 2025-11-02 03:16:48 +00:00
Qianfeng Zhang
8408ec0a02 Add scripts for testing the using of separate sequence lengths for k/v 2025-11-02 03:16:22 +00:00
Qianfeng Zhang
17e404be3b Support separate sequence lengths for q and kv 2025-11-02 03:14:53 +00:00
Qianfeng Zhang
eaf9650fed Use separate pipelines for using or not-using softmax situations 2025-10-30 10:01:52 +00:00
Qianfeng Zhang
207e6f10b8 Implementation of hstu attention pipeline using trload for v on mi350 2025-10-29 15:45:14 +00:00
Qianfeng Zhang
a464269bb6 Fix in the comments 2025-10-27 10:47:40 +00:00
Qianfeng Zhang
4eeb5cc917 Update to gemm_0's CBlockDistribution encoding so that it is compatible with gemm_1's ABlockDistribution encoding 2025-10-27 10:47:23 +00:00
Qianfeng Zhang
98a241a2eb Using separate tile settings for no-softmax and with-softmax hstu attention situations 2025-10-24 01:47:55 +00:00
Qianfeng Zhang
7c4012266a Update to benchmark scripts to consider for using softmax 2025-10-23 10:09:37 +00:00
Qianfeng Zhang
d1505786f8 Add support of softmax in hstu attention 2025-10-20 14:26:55 +00:00
Qianfeng Zhang
a874839dc2 Add template parameter to gemm_0 MakeCBlockTile() for the need of defining PcompBlockTileType 2025-10-20 14:26:29 +00:00
Qianfeng Zhang
1a8f2f21fb Move scaling by attn_scale to inside the main-loop 2025-10-20 14:22:18 +00:00
Qianfeng Zhang
bbda3f6f1c Let IsTokenPairInsideMask() return bool type 2025-10-20 14:21:26 +00:00
Qianfeng Zhang
fdb89d3e2f Add instances to consider for adding softmax support 2025-10-20 14:20:54 +00:00
Qianfeng Zhang
2072e53d1e Remove K0 from tile setting since it is not used 2025-10-14 07:15:26 +00:00
Qianfeng Zhang
22a7b31865 Change to pipeline so that it is easier to add support of using softmax 2025-10-12 06:09:55 +00:00
Qianfeng Zhang
d308b09fae Remove using IGLP method for instruction scheduling for kUseLocal true path 2025-10-12 06:09:25 +00:00
Qianfeng Zhang
6b40ce4074 Fix in GetQKBlockGemm() 2025-09-27 14:59:32 +00:00
Qianfeng Zhang
27b96b15c4 Simplify the warp_gemm definitions in GetQKBlockGemm and GetKVBlockGemm 2025-09-25 15:38:55 +00:00
Qianfeng Zhang
bd32cc0de0 Remove useless constant statement in the kernel 2025-09-19 07:24:29 +00:00
Qianfeng Zhang
db62a9f47e Remove un-necessary HSTU_CHECK() callings 2025-09-13 16:39:14 +00:00
Qianfeng Zhang
2427426640 Add HSTU_CHECK() and use it in example codes 2025-09-13 16:38:33 +00:00
Qianfeng Zhang
a5b7360862 Smalle update in reference hstu attention 2025-09-13 06:53:54 +00:00
Qianfeng Zhang
798fc3cc0b Detach HstuBlockMask from pipeline definition and construct the HstuBlockMask type in the kernel according to window_size 2025-09-12 09:11:47 +00:00
Qianfeng Zhang
7d10353fda Unify the license statements on all the source files 2025-09-11 10:27:00 +00:00
Qianfeng Zhang
1c030e8c3c Remove using MakeKargsImpl() to simplify the hstu kernel 2025-09-10 15:28:12 +00:00
Qianfeng Zhang
72eb4e95d8 Clarify the using of kSubQKHeaddim and kQKHeaddim so that less regular hdim (eg. 96, 160) can be efficiently supported 2025-09-09 12:55:01 +00:00
Qianfeng Zhang
f8dea2bc86 Use set_slice_tilie() to replace direct thread_buffer assignment 2025-09-09 12:54:32 +00:00
Qianfeng Zhang
4bf65d9fe5 Merge branch 'develop' into hstu_attention_mi350_fwd_bwd and change in using ck_tile::make_kernel 2025-09-01 07:35:35 +00:00
Aviral Goel
fcff0043ae chore(gemm): clang format to pass CI (#2758) 2025-08-29 00:38:46 -07:00
Vijay Krish
4208e28988 ck_tile kernel for gemm with groupwise quantized B tensor. (#2663)
* This change introduces new pipelines with Intrawave scheduler and block gemm primitives that loads the scale tensor to registers to perform dequantization post MFMA on C tensor in registers.

Scale tensor data, BQ is spliced across threads in registers and not stored in LDS.

Current support is for the following combinations, but it should be fairly straightforward to extend support to more formats.

fp8, fp8 -> f32
bf8, bf8 -> f32
fp8, i4 -> f32
bf8, i4 -> f32
Group size can go down to as low as K length of underlying WarpGemm primitive.

* Solve merge conflict

* [CK TILE] Update CHANGELOG.md

---------

Co-authored-by: Vijay Krishnamoorthy <vjkrish@fb.com>
Co-authored-by: ThomasNing <thomas.ning@amd.com>
Co-authored-by: Cong Ma <congma13@amd.com>
2025-08-28 23:43:02 -07:00
Cong Ma
428090f749 Support transposed C tile in Aquant (#2679)
The performance of Aquant has increased after enabling transposed C.

Do not need to exchange AQ elements among lanes after enabling
transposed C as one thread only holds data from one row.
2025-08-28 13:28:09 -07:00
asleepzzz
038ea82315 Revert "[CK_TILE] FMHA BWD Enable Tile 16x192 (#2741)" (#2757)
This reverts commit ead4447b20.
2025-08-28 22:50:42 +08:00
linqunAMD
4a49dac7c6 [Regression] Fix CK_TILE build error in grouped_convolution, copy_basic and fused_moegemm_kernel (#2728)
* fix copy basic build error

* fix other ck tile test build error
2025-08-28 20:30:30 +08:00
Yi DING
ead4447b20 [CK_TILE] FMHA BWD Enable Tile 16x192 (#2741)
* 16x192

* Use buffer_load_lds for lse/d

* Dispatch & cleanup

* Avoid zeroing dq & fix

* fix
2025-08-28 18:54:18 +08:00
Aviral Goel
f5f795c4d6 feat(HostTensor): Extend support for HostTensor class' >> operator to print more data types (#2691)
* feat(check_err): add a variable to adjust number of incorrect values to print

* feat(host_tensor): add printing capability for fp8 bf8 int8 int4

* fix(gemm_utils): update acceptable data type

* fix(host_tensor): print both 4 bit ints in pk_int4_t

* refactor(HostTensor): define pk_int4_t_to_int8x2_t and fix typo in vector_type.hpp

* feat(host_tensor): add print first n elements functions
2025-08-27 18:17:24 -07:00
Cong Ma
cd53e2e57e [CK TILE GEMM] Fix a merge conflict (#2753)
* Fixed a merge conflict in 245467f3
* Foramt the code
2025-08-27 11:08:09 -07:00
Cong Ma
245467f359 [CK TILE] Fix bugs in AQuant preshuffle (#2700)
* [CK TILE] Fix bugs in AQuant preshuffle

- Make Aquant works with block Mx64x256. `M` could be 16, 32, 64
- Make Aquant works with warp 16x16x32 and 32x32x16.

* [CK TILE] Rename Preshuffle to PreshuffleQuant

The new name, PreshuffleQuant, explicitly states the function's purpose:
to preshuffle the quantization matrix.

* [CK TILE Block Scale] Use GemmConfig to save tile properties

- Remove specialization of GemmQuantTypeConfig
- Pass GemmConfig around which contains tile properties. Stop using hard
  coded tile properties in `gemm_calc_aquant()`

* [CK TILE Block Scale] Rename GemmConfig used in block scale

    - Remove unused GemmConfig
    - Rename GemmConfig used in block scale

---------

Co-authored-by: ThomasNing <thomas.ning@amd.com>
2025-08-27 00:05:54 -07:00
Tianyuan Wu
e9605ed36d [CK_TILE] Fix the Wrong Output Generated by Gemm Examples on GFX11/12 (#2713)
* Introduce macro CK_TILE_USE_WMMA

Signed-off-by: Tianyuan Wu <Tianyuan.Wu@amd.com>

* Make CK_TILE_USE_WMMA global for all examples

Signed-off-by: Tianyuan Wu <Tianyuan.Wu@amd.com>

* Remove CK_TILE_USE_WMMA from config.hpp

Signed-off-by: Tianyuan Wu <Tianyuan.Wu@amd.com>

---------

Signed-off-by: Tianyuan Wu <Tianyuan.Wu@amd.com>
2025-08-25 12:55:35 -07:00
Yi DING
4cfa2c7158 [CK_TILE] FMHA BWD Fix Compilation with Bias (#2682)
* [CK_TILE] FMHA BWD Fix Compilation with Bias

* Fix appendkv kApplyRoPE
2025-08-22 10:01:10 +08:00
Qianfeng Zhang
d4f43f0653 Use xor transform to implement Q/K Lds descriptor for kKpack == 8 cases 2025-08-21 14:16:34 +00:00
Bartłomiej Kocot
4212bbc170 [CK Tile] Grouped convolution backward data (#2652)
* base working version for single groupped conv bwd data

* Fix 2d descriptor

* fix groups

* Add 3d support

* fixes

* fixes

* fixes

---------

Co-authored-by: Jakub Piasecki <jakpia21@gmail.com>
2025-08-20 05:29:57 -07:00
Qianfeng Zhang
19fc2a9051 Remove selectable VLayout for simplifying the codes since hdim is always fatest dimension 2025-08-20 09:06:01 +00:00