Haocong WANG 3049b5467c [GEMM] gemm_universal related optimization (#1453) ...

* replace buffer_atomic with global_atomic

* fixed global_atomic_add

* added bf16 atomic_add

* format

* clang-format-12

* clean

* clean

* add guards

* Update gtest.cmake

* enabled splitk_gemm_multi_d

* format

* add ckProfiler

* format

* fixed naming

* format

* clean

* clean

* add guards

* fix clang format

* format

* add kbatch printout

* clean

* Add rocm6.2 related gemm optimization

* Limit bf16 atomic usage

* remove redundant RCR gemm_universal instance

* Add RRR fp8 gemm universal instance

* Bug fix

* Add GPU_TARGET guard to FP8/BF8 target

* bug fix

* update cmake

* remove all fp8/bf8 example if arch not support

* Enable fp8 RRR support in ckProfiler

* limit greedy-reverse flag to gemm_universal in ckProfiler

---------

Co-authored-by: Jing Zhang <jizhan@fb.com>
Co-authored-by: Jing Zhang <jizhan@meta.com>
Co-authored-by: zjing14 <zhangjing14@gmail.com>
Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>
Co-authored-by: illsilin <Illia.Silin@amd.com>

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

..

batchnorm_multiblock

[HotFix] add config and version files to pass on build info (#856)

2023-08-23 11:36:17 -07:00

gemm_layernorm

Add fp8 @ bf8 gemm support and example (#933)

2023-10-02 16:39:03 -05:00

normalization

layernorm and groupnorm backward data (#1083)

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

block_to_ctile_map.hpp

Universal streamk with atomics (#1360)

2024-07-05 21:40:30 -07:00

gridwise_2d_multiple_reduction_multiblock.hpp

Fix for beta!=0 in reduce (#1440)

2024-08-06 09:10:39 -07:00

gridwise_2d_multiple_reduction_threadwise.hpp

update copyright headers (#726)

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

gridwise_2d_reduction_multiblock.hpp

Fix for beta!=0 in reduce (#1440)

2024-08-06 09:10:39 -07:00

gridwise_2d_reduction_threadwise_multi_d.hpp

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

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

gridwise_2d_reduction_threadwise.hpp

update copyright headers (#726)

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

gridwise_batched_gemm_gemm_xdl_cshuffle_v1.hpp

Add fp8 @ bf8 gemm support and example (#933)

2023-10-02 16:39:03 -05:00

gridwise_batched_gemm_multiple_d_gemm_multiple_d_xdl_cshuffle_v1.hpp

Navi3 rel (#1176)

2024-03-08 17:11:51 -08:00

gridwise_batched_gemm_multiple_d_softmax_gemm_xdl_cshuffle_v1.hpp

Add fp8 @ bf8 gemm support and example (#933)

2023-10-02 16:39:03 -05:00

gridwise_batched_gemm_softmax_gemm_wmma_cshuffle.hpp

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

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

gridwise_batched_gemm_softmax_gemm_xdl_cshuffle_v1.hpp

Add fp8 @ bf8 gemm support and example (#933)

2023-10-02 16:39:03 -05:00

gridwise_batchnorm_backward_blockwise_welford.hpp

update copyright headers (#726)

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

gridwise_batchnorm_forward_blockwise_welford.hpp

update copyright headers (#726)

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

gridwise_elementwise_1d_scale.hpp

Refactor elementwise kernels (#1222)

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

gridwise_elementwise_2d.hpp

Refactor elementwise kernels (#1222)

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

gridwise_elementwise_layernorm_welford_variance.hpp

enable compilation of INSTANCES_ONLY for Windows (#1082)

2023-12-20 14:34:53 -08:00

gridwise_fpAintB_gemm_wmma.hpp

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

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

gridwise_gemm_bias_add_reduce_xdl_cshuffle_v1.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_gemm_dl_multiple_d.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432)

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

gridwise_gemm_dl_v1r3.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432)

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

gridwise_gemm_dpp.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_gemm_multiple_abd_xdl_cshuffle.hpp

bf16A_Int8B with fastgelu/bias (#1264)

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

gridwise_gemm_multiple_d_multiple_r_xdl_cshuffle.hpp

Add fp8 @ bf8 gemm support and example (#933)

2023-10-02 16:39:03 -05:00

gridwise_gemm_multiple_d_wmma_cshuffle.hpp

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

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

gridwise_gemm_multiple_d_xdl_cshuffle_lds_direct_load.hpp

Implement direct loads split-K GEMM kernel (#1137)

2024-02-07 01:08:34 +01:00

gridwise_gemm_multiple_d_xdl_cshuffle.hpp

Grouped GEMM Multiple D tile loop. (#1247)

2024-04-25 15:12:53 -05:00

gridwise_gemm_multiple_d_xdl_splitk_cshuffle.hpp

improved zeroing (#1221)

2024-04-02 11:02:52 -05:00

gridwise_gemm_pipeline_selector.hpp

Grouped GEMM Multiple D tile loop. (#1247)

2024-04-25 15:12:53 -05:00

gridwise_gemm_pipeline_v1.hpp

Navi3 rel (#1176)

2024-03-08 17:11:51 -08:00

gridwise_gemm_pipeline_v2.hpp

Fixed Gemmv2r3 kpad (#938)

2023-09-26 18:40:00 -05:00

gridwise_gemm_pipeline_v3.hpp

initial stream-k implementation with example (#699)

2023-07-26 14:18:15 -05:00

gridwise_gemm_pipeline_v4_direct_load.hpp

Add support for double buffering in direct load GEMM kernel (#1052)

2023-12-03 23:08:47 +01:00

gridwise_gemm_reduce_xdl_cshuffle_v1.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_gemm_split_k_multiple_d_xdl_cshuffle_v2.hpp

Fp16/fp8 mixed-precision Gemm with multiply+add fusion (#865)

2023-08-28 16:27:32 -05:00

gridwise_gemm_split_k_multiple_d_xdl_cshuffle.hpp

Navi3 rel (#1176)

2024-03-08 17:11:51 -08:00

gridwise_gemm_waveletmodel.hpp

update copyright headers (#726)

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

gridwise_gemm_wmma.hpp

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

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

gridwise_gemm_xdl_cshuffle_bwd_weight_v3.hpp

Optimize grouped conv bwd weight for small M and N (#1303)

2024-05-22 21:01:01 +02:00

gridwise_gemm_xdl_cshuffle_streamk_v3.hpp

Universal streamk with atomics (#1360)

2024-07-05 21:40:30 -07:00

gridwise_gemm_xdl_cshuffle_v1.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_gemm_xdl_cshuffle_v2.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_gemm_xdl_cshuffle_v3_multi_abd.hpp

Optimize grouped conv bwd weight for small M and N (#1303)

2024-05-22 21:01:01 +02:00

gridwise_gemm_xdl_cshuffle_v3_multi_d_ab_scale.hpp

[GEMM] F8 GEMM, performance optimized. (#1384)

2024-07-19 22:06:52 +08:00

gridwise_gemm_xdl_cshuffle_v3_multi_d.hpp

[GEMM] gemm_universal related optimization (#1453)

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

gridwise_gemm_xdl_cshuffle_v3.hpp

[GEMM] gemm_universal related optimization (#1453)

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

gridwise_gemm_xdl_layernorm_cshuffle_v1.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_gemm_xdl_waveletmodel_cshuffle.hpp

Add fp8 @ bf8 gemm support and example (#933)

2023-10-02 16:39:03 -05:00

gridwise_gemm_xdlops_bwd_weight.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_gemm_xdlops_skip_b_lds_v1.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_gemm_xdlops_splitk_lds_direct_load.hpp

aggregate device macros in ck_tile config header (#1297)

2024-05-20 08:34:45 -07:00

gridwise_gemm_xdlops_streamk.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_gemm_xdlops_v2r3.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_gemm_xdlops_v2r4.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_gemm_xdlops_v2r4r2.hpp

replace the ENV macro with CK_ENV (#1296)

2024-05-17 10:42:51 -07:00

gridwise_gemm_xdlops_v3r1.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_gemm_xdlops_v3r2.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_gemm_xdlops_v3r3.hpp

Add support for more Navi2x and Navi3x models. (#1152)

2024-02-02 11:35:26 -08:00

gridwise_permute.hpp

update copyright headers (#726)

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

gridwise_put_element_1d.hpp

Average pool backward deviceOP and example (#797)

2023-08-10 12:04:35 +08:00

gridwise_set_buffer_value.hpp

update copyright headers (#726)

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

gridwise_set_multiple_buffer_value.hpp

update copyright headers (#726)

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

gridwise_softmax.hpp

update copyright headers (#726)

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

gridwise_sparse_embeddings_forward_layernorm.hpp

[HotFix] add config and version files to pass on build info (#856)

2023-08-23 11:36:17 -07:00

gridwise_tensor_rearrange.hpp

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

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