Po Yen Chen 95e722a3b3 [CK_TILE] Optimize fmha splitkv & splitkv combine kernels (#1577) ...

* Use smaller width for lse_accum dist tensor

* Update pipeline comment

* Fix wrong distribution for lse_accum

* Remove duplicate dim in lse_accum dist encoding

* Decide fmha splitkv combine kernel kBlockSize by kM0

* Remove assumption of MPerThread=1

* Add log<4> & log<8> specialization

* Enlarge occupancy array

* Fix vector size for small tile

* Add support for kMaxSplits=8

* Re-format gemm.hpp

* Use 16x16x16 warp gemm for fwd_splitkv

* Centralize policy code changes

* Leave fp8/bf8 tile settings unchanged

2024-10-21 10:52:11 +08:00

..

01_gemm

Fix default stride value (#1559)

2024-10-10 07:37:09 -07:00

02_gemm_bilinear

Fix issue with multiple targets and remove smfmac tests from unsupported test targets (#1372)

2024-07-03 23:34:38 -07:00

03_gemm_bias_relu

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

04_gemm_add_add_fastgelu

Merging the gfx12 code into public repo. (#1362)

2024-06-27 00:33:34 -07:00

09_convnd_fwd

re-enable convnd_fwd_xdl_fp64 testing (#1289)

2024-05-10 22:48:28 -07:00

10_convnd_fwd_multiple_d_multiple_reduce

Add Grouped Conv Fwd Large Tensor kernel (#1432)

2024-08-06 10:06:10 +02:00

12_reduce

Support large: 12d tensor size for reduction kenrel (#1465)

2024-08-13 16:15:47 +02:00

13_pool2d_fwd

Refactoring cmake files to build data types separately. (#932)

2023-09-20 22:15:56 -07:00

14_gemm_quantization

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

15_grouped_gemm

Switch to universal gemm in grouped gemm tile loop (#1335)

2024-06-18 09:01:49 -05:00

16_gemm_multi_d_multi_reduces

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

17_convnd_bwd_data

Add Grouped Conv Fwd Large Tensor kernel (#1432)

2024-08-06 10:06:10 +02:00

18_batched_gemm_reduce

Fixing most of the cppcheck errors. (#1142)

2024-01-24 13:47:48 -08:00

19_binary_elementwise

Refactor elementwise kernels (#1222)

2024-04-19 13:31:17 +02:00

20_grouped_conv_bwd_weight

[GEMM] gemm_universal related optimization (#1453)

2024-08-14 10:42:30 +08:00

21_gemm_layernorm

Refactor elementwise kernels (#1222)

2024-04-19 13:31:17 +02:00

22_cgemm

Clean DTYPES conditions in CMake (#974)

2023-10-18 11:14:14 -05:00

23_softmax

update copyright headers (#726)

2023-05-31 18:46:57 -05:00

24_batched_gemm

Clean DTYPES conditions in CMake (#974)

2023-10-18 11:14:14 -05:00

25_gemm_bias_e_permute

Refactoring cmake files to build data types separately. (#932)

2023-09-20 22:15:56 -07:00

26_contraction

Code clean-up (#1285)

2024-05-10 09:41:39 -07:00

27_layernorm2d_fwd

Split the static library into several files. (#1044)

2023-11-28 11:17:37 -08:00

28_grouped_gemm_bias_e_permute

Refactoring cmake files to build data types separately. (#932)

2023-09-20 22:15:56 -07:00

29_batched_gemm_bias_e_permute

Merging the gfx12 code into public repo. (#1362)

2024-06-27 00:33:34 -07:00

30_grouped_conv_fwd_multiple_d

Fix issue with multiple targets and remove smfmac tests from unsupported test targets (#1372)

2024-07-03 23:34:38 -07:00

31_batched_gemm_gemm

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

32_batched_gemm_scale_softmax_gemm

Fix issue with multiple targets and remove smfmac tests from unsupported test targets (#1372)

2024-07-03 23:34:38 -07:00

33_multiple_reduce

update copyright headers (#726)

2023-05-31 18:46:57 -05:00

34_batchnorm

Refactor elementwise kernels (#1222)

2024-04-19 13:31:17 +02:00

35_splitK_gemm

Universal gemm splitk using reduce (with multi-d) (#1341)

2024-07-19 22:01:22 +08:00

36_sparse_embedding

update copyright headers (#726)

2023-05-31 18:46:57 -05:00

37_batched_gemm_add_add_relu_gemm_add

Refactoring cmake files to build data types separately. (#932)

2023-09-20 22:15:56 -07:00

38_grouped_conv_bwd_data_multiple_d

Fix issue with multiple targets and remove smfmac tests from unsupported test targets (#1372)

2024-07-03 23:34:38 -07:00

39_permute

Clean DTYPES conditions in CMake (#974)

2023-10-18 11:14:14 -05:00

40_conv2d_fwd_quantization

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

41_grouped_conv_conv_fwd

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

42_groupnorm_fwd

Split the static library into several files. (#1044)

2023-11-28 11:17:37 -08:00

43_splitk_gemm_bias_e_permute

Refactoring cmake files to build data types separately. (#932)

2023-09-20 22:15:56 -07:00

44_elementwise_permute

Add transpose scale amax example (#1547)

2024-10-14 17:39:38 +02:00

45_elementwise_normalization

Layernorm and groupnorm support to save mean and inverse std in forward (#929)

2023-10-19 07:36:29 +08:00

46_gemm_add_multiply

Code clean-up (#1285)

2024-05-10 09:41:39 -07:00

47_gemm_bias_softmax_gemm_permute

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

48_pool3d_fwd

Fixing most of the cppcheck errors. (#1142)

2024-01-24 13:47:48 -08:00

49_maxpool2d_bwd

Refactoring cmake files to build data types separately. (#932)

2023-09-20 22:15:56 -07:00

50_put_element

Refactoring cmake files to build data types separately. (#932)

2023-09-20 22:15:56 -07:00

51_avgpool3d_bwd

Fixing most of the cppcheck errors. (#1142)

2024-01-24 13:47:48 -08:00

52_im2col_col2im

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

53_layernorm2d_bwd

layernorm and groupnorm backward data (#1083)

2023-12-19 04:23:11 +08:00

54_groupnorm_bwd

layernorm and groupnorm backward data (#1083)

2023-12-19 04:23:11 +08:00

59_grouped_gemm_multi_ABD

Fix example CMakeLists.txt (#1267)

2024-04-30 08:28:19 -07:00

60_gemm_multi_ABD

bf16A_Int8B with fastgelu/bias (#1264)

2024-04-26 07:26:30 -05:00

61_contraction_multi_ABD

add gemm_bias_add example (#1361)

2024-07-11 18:08:07 -07:00

62_convnd_activ

Adding Instances and Examples for FP8-based Scaled Convolution and AMAX Reduction. (#1473)

2024-08-21 15:22:41 -07:00

63_layernorm4d_fwd

Split the static library into several files. (#1044)

2023-11-28 11:17:37 -08:00

64_fpAintB_gemm

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07:00

65_gemm_multiply_multiply

[GEMM] gemm_universal related optimization (#1453)

2024-08-14 10:42:30 +08:00

66_complex_contraction_bilinear

Complex Contraction CK Bilinear Example (#1061)

2024-09-30 21:05:42 -06:00

ck_tile

[CK_TILE] Optimize fmha splitkv & splitkv combine kernels (#1577)

2024-10-21 10:52:11 +08:00

CMakeLists.txt

Fix build logic using GRU_ARCHS. (#1536)

2024-10-07 08:18:23 -07:00