Dan Yao 26840b623a CK Tile FA Training kernels (#1286) ...

* FA fwd dropout

* FA bwd

* epilogue reuse

* CMakeLists update

* [CK_TILE] support alibi (#1269)

* add alibi support

* fix code

* update code based on comment

* Support more hdim

* fix fp8 bias

* support seqlen_k=0 case

* remove unused printf

* fix format

---------

Co-authored-by: rocking <ChunYu.Lai@amd.com>

* now fwd/bwd can build

* bwd alibi

* add bwd validation stream_config

* update generated filenames

* update bwd kernel launch

* CK_TILE_HOST_DEVICE in philox

* Transpose -> transpose

* format

* format

* format

* Generate the instance for FA required

* format

* fix error in WarpGemm

---------

Co-authored-by: danyao12 <danyao12>
Co-authored-by: carlushuang <carlus.huang@amd.com>
Co-authored-by: rocking <ChunYu.Lai@amd.com>
Co-authored-by: Po Yen Chen <PoYen.Chen@amd.com>
Co-authored-by: Jing Zhang <jizhan@amd.com>

[ROCm/composable_kernel commit: 2cab8d39e3]

2024-06-04 13:12:45 -05:00

..

01_gemm

Code clean-up (#1285)

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

02_gemm_bilinear

Code clean-up (#1285)

2024-05-10 09:41:39 -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

Code clean-up (#1285)

2024-05-10 09:41:39 -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

Split the instances by architecture. (#1223)

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

12_reduce

Change 1d,2d,... to 1D,2D,... (#997)

2023-10-19 16:53:18 +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

Code clean-up (#1285)

2024-05-10 09:41:39 -07: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

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07: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

Add grouped conv bwd weight multi d kernel (#1237)

2024-04-18 23:35:04 +02: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

Split the instances by architecture. (#1223)

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

30_grouped_conv_fwd_multiple_d

Code clean-up (#1285)

2024-05-10 09:41:39 -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

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -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

Split the instances by architecture. (#1223)

2024-04-02 09:42:17 -07: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

Add an example (#1227)

2024-04-09 13:57:32 -05: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

Refactor elementwise kernels (#1222)

2024-04-19 13:31:17 +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

Added Multi_ABD support into Gemm and GroupedGemmFixedNK (#978)

2024-04-15 21:09:45 -05:00

62_convnd_activ

Add grouped conv bwd weight multi d kernel (#1237)

2024-04-18 23:35:04 +02: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

Post-merge fix of PR 1300 (#1313)

2024-05-31 22:46:41 -07:00

ck_tile

CK Tile FA Training kernels (#1286)

2024-06-04 13:12:45 -05:00

CMakeLists.txt

Select appropriate GPU targets for instances, tests, and examples. (#1304)

2024-05-22 11:45:27 -07:00