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

..

codegen_device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432)

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

device_avgpool3d_bwd_ndhwc_ndhwc.hpp

Average pool backward deviceOP and example (#797)

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

device_batched_contraction_multiple_d_wmma_cshuffle.hpp

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

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

device_batched_contraction_multiple_d_xdl_cshuffle.hpp

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

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

device_batched_gemm_e_permute_xdl.hpp

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

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

device_batched_gemm_gemm_xdl_cshuffle.hpp

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

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

device_batched_gemm_multi_d_xdl.hpp

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

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

device_batched_gemm_multiple_d_dl.hpp

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

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

device_batched_gemm_multiple_d_gemm_multiple_d_xdl_cshuffle.hpp

replace the ENV macro with CK_ENV (#1296)

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

device_batched_gemm_reduce_xdl_cshuffle.hpp

replace the ENV macro with CK_ENV (#1296)

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

device_batched_gemm_softmax_gemm_permute_wmma_cshuffle.hpp

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

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

device_batched_gemm_softmax_gemm_permute_xdl_cshuffle.hpp

replace the ENV macro with CK_ENV (#1296)

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

device_batched_gemm_softmax_gemm_xdl_cshuffle.hpp

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

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

device_batched_gemm_xdl.hpp

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

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

device_batchnorm_backward_impl.hpp

Fixed GroupedGemmFixedNK with hipGraph (#1065)

2023-11-30 15:09:27 -06:00

device_batchnorm_forward_impl_obsolete.hpp

Fixed GroupedGemmFixedNK with hipGraph (#1065)

2023-11-30 15:09:27 -06:00

device_batchnorm_forward_impl.hpp

Fixed GroupedGemmFixedNK with hipGraph (#1065)

2023-11-30 15:09:27 -06:00

device_cgemm_4gemm_xdl_cshuffle.hpp

Refactor elementwise kernels (#1222)

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

device_column_to_image_impl.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432)

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

device_contraction_multiple_abd_xdl_cshuffle.hpp

Fix continous dim selection in contraction (#1336)

2024-06-18 10:26:49 +02:00

device_contraction_multiple_d_xdl_cshuffle.hpp

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

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

device_contraction_utils.hpp

Fix continous dim selection in contraction (#1336)

2024-06-18 10:26:49 +02:00

device_conv2d_backward_weight_xdl_c_shuffle_nhwc_kyxc_nhwk.hpp

Use asynchronous version of hipMemset (#850)

2023-08-18 11:14:59 +08:00

device_conv2d_bwd_data_xdl_nhwc_kyxc_nhwk.hpp

replace the ENV macro with CK_ENV (#1296)

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

device_conv2d_fwd_xdl_c_shuffle_bias_activation_add_nhwc_kyxc_nhwk.hpp

replace the ENV macro with CK_ENV (#1296)

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

device_conv2d_fwd_xdl_c_shuffle_bias_activation_nhwc_kyxc_nhwk.hpp

replace the ENV macro with CK_ENV (#1296)

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

device_conv2d_fwd_xdl_c_shuffle_nhwc_kyxc_nhwk.hpp

replace the ENV macro with CK_ENV (#1296)

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

device_conv2d_fwd_xdl_nhwc_kyxc_nhwk.hpp

replace the ENV macro with CK_ENV (#1296)

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

device_conv3d_fwd_naive_ndhwc_kzyxc_ndhwk.hpp

update copyright headers (#726)

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

device_conv3d_fwd_xdl_ndhwc_kzyxc_ndhwk.hpp

replace the ENV macro with CK_ENV (#1296)

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

device_convnd_bwd_data_nwc_kxc_nwk_dl.hpp

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

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

device_convnd_bwd_data_nwc_kxc_nwk_xdl.hpp

replace the ENV macro with CK_ENV (#1296)

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

device_elementwise_dynamic_vector_dims_impl.hpp

Refactor elementwise kernels (#1222)

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

device_elementwise_normalization_impl.hpp

update copyright headers (#726)

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

device_elementwise_scale_impl.hpp

Refactor elementwise kernels (#1222)

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

device_fpAintB_gemm_wmma.hpp

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

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

device_gemm_bias_add_reduce_xdl_cshuffle.hpp

Disable XDL kernels on unsupported HW Add ck::is_xdl_supported (#768)

2023-07-26 07:19:55 -07:00

device_gemm_dl.hpp

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

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

device_gemm_dpp.hpp

Code clean-up (#1285)

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

device_gemm_multiple_abd_xdl_cshuffle.hpp

bf16A_Int8B with fastgelu/bias (#1264)

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

device_gemm_multiple_d_dl.hpp

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

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

device_gemm_multiple_d_layernorm_xdl_cshuffle.hpp

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

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

device_gemm_multiple_d_multiple_r_xdl_cshuffle.hpp

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

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

device_gemm_multiple_d_wmma_cshuffle.hpp

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

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

device_gemm_multiple_d_xdl_cshuffle_lds_direct_load.hpp

Add basic support for direct loads from global to LDS (#999)

2023-11-25 13:35:22 +01:00

device_gemm_multiple_d_xdl_cshuffle_v3_ab_scale.hpp

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

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

device_gemm_multiple_d_xdl_cshuffle_v3.hpp

[GEMM] gemm_universal related optimization (#1453)

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

device_gemm_multiple_d_xdl_cshuffle.hpp

Add host lib (#1134)

2024-03-05 17:08:43 -08:00

device_gemm_reduce_xdl_cshuffle.hpp

replace the ENV macro with CK_ENV (#1296)

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

device_gemm_wmma.hpp

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

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

device_gemm_xdl_cshuffle_lds_direct_load.hpp

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

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

device_gemm_xdl_cshuffle_streamk_v3.hpp

Universal streamk with atomics (#1360)

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

device_gemm_xdl_cshuffle_v2.hpp

[GEMM] Optimization for MI200/300. (#1135)

2024-01-19 07:02:22 -06:00

device_gemm_xdl_cshuffle_v3.hpp

[GEMM] gemm_universal related optimization (#1453)

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

device_gemm_xdl_cshuffle_v3r1.hpp

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

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

device_gemm_xdl_cshuffle.hpp

Add Gemm instances for performance improvement (#1018)

2023-11-07 09:09:58 -06:00

device_gemm_xdl_layernorm_cshuffle.hpp

replace the ENV macro with CK_ENV (#1296)

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

device_gemm_xdl_skip_b_lds.hpp

replace the ENV macro with CK_ENV (#1296)

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

device_gemm_xdl_splitk_c_shuffle_lds_direct_load.hpp

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

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

device_gemm_xdl_splitk_c_shuffle.hpp

Optimizing fp8_fp16 mixedprec gemm (#1150)

2024-02-12 09:45:42 -08:00

device_gemm_xdl_streamk.hpp

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

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

device_gemm_xdl_waveletmodel_cshuffle.hpp

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

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

device_gemm_xdl.hpp

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

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

device_grouped_contraction_multiple_d_xdl_cshuffle.hpp

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

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

device_grouped_conv_bwd_data_multiple_d_wmma_cshuffle.hpp

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

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

device_grouped_conv_bwd_data_multiple_d_xdl_cshuffle_v1.hpp

Add read_first_lane function for int64 (#1347)

2024-06-18 15:05:30 -05:00

device_grouped_conv_bwd_weight_dl.hpp

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

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

device_grouped_conv_bwd_weight_multiple_d_xdl_cshuffle.hpp

Add read_first_lane function for int64 (#1347)

2024-06-18 15:05:30 -05:00

device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp

Support access per groups and filter3x3 in grouped conv fwd (#1382)

2024-07-12 11:08:42 -07:00

device_grouped_conv_bwd_weight_wmma_cshuffle.hpp

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

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

device_grouped_conv_bwd_weight_xdl_cshuffle.hpp

Support large tensors in grouped conv fwd (#1332)

2024-06-14 09:53:03 -05:00

device_grouped_conv_fwd_dl_multiple_d_nhwc_kyxc_nhwk.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432)

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

device_grouped_conv_fwd_dl_nhwc_kyxc_nhwk.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432)

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

device_grouped_conv_fwd_multiple_abd_xdl_cshuffle_v3.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432)

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

device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp

Fix bug with n block id calculation in DeviceGroupedConvXdlCShuffle (#1457)

2024-08-10 13:12:05 +02:00

device_grouped_conv_fwd_multiple_d_multiple_r_xdl_cshuffle.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432)

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

device_grouped_conv_fwd_multiple_d_multiple_r.hpp

update copyright headers (#726)

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

device_grouped_conv_fwd_multiple_d_wmma_cshuffle.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432)

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

device_grouped_conv_fwd_multiple_d_xdl_cshuffle.hpp

Add instances for conv_scale with fp8@bf8->fp8 (#1220)

2024-04-03 09:08:08 -05:00

device_grouped_conv_fwd_multiple_d_xdl_large_tensor_cshuffle.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432)

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

device_grouped_conv_utils.hpp

Support access per groups and filter3x3 in grouped conv fwd (#1382)

2024-07-12 11:08:42 -07:00

device_grouped_gemm_multi_abd_xdl_fixed_nk.hpp

aggregate device macros in ck_tile config header (#1297)

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

device_grouped_gemm_multiple_d_dl.hpp

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

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

device_grouped_gemm_multiple_d_splitk_xdl_cshuffle_two_stage.hpp

replace the ENV macro with CK_ENV (#1296)

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

device_grouped_gemm_multiple_d_xdl_cshuffle_tile_loop.hpp

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

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

device_grouped_gemm_softmax_gemm_permute_xdl_cshuffle.hpp

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

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

device_grouped_gemm_xdl_fixed_nk.hpp

improved zeroing (#1221)

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

device_grouped_gemm_xdl_splitk_cshuffle.hpp

replace the ENV macro with CK_ENV (#1296)

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

device_grouped_gemm_xdl.hpp

replace the ENV macro with CK_ENV (#1296)

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

device_grouped_query_attention_forward_wmma.hpp

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

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

device_image_to_column_impl.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432)

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

device_max_pool_bwd_impl.hpp

Refactor elementwise kernels (#1222)

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

device_multi_query_attention_forward_wmma.hpp

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

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

device_multiple_reduce_multiblock.hpp

update copyright headers (#726)

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

device_multiple_reduce_threadwise.hpp

update copyright headers (#726)

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

device_normalization_bwd_data_impl.hpp

layernorm and groupnorm backward data (#1083)

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

device_normalization_bwd_gamma_beta_impl.hpp

layernorm and groupnorm backward data (#1083)

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

device_normalization_fwd_impl.hpp

layernorm and groupnorm backward data (#1083)

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

device_normalization_fwd_splitk_impl.hpp

layernorm and groupnorm backward data (#1083)

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

device_permute_impl.hpp

update copyright headers (#726)

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

device_pool2d_fwd_nhwc_nhwc.hpp

Refactor pool fwd (#815)

2023-08-15 02:25:28 +08:00

device_pool3d_fwd_ndhwc_ndhwc.hpp

Refactor pool fwd (#815)

2023-08-15 02:25:28 +08:00

device_put_element_impl.hpp

Maxpool bwd (#750)

2023-06-19 09:44:22 -05:00

device_reduce_common.hpp

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

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

device_reduce_multiblock.hpp

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

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

device_reduce_threadwise_multi_d.hpp

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

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

device_reduce_threadwise.hpp

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

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

device_softmax_impl.hpp

Revert "Grouped Gemm with looping over the tiles. (#788)" (#982)

2023-10-11 14:27:29 -05:00

device_sparse_embeddings_forward_layernorm.hpp

update copyright headers (#726)

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

device_splitk_contraction_multiple_d_xdl_cshuffle.hpp

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

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