Haocong WANG 65d6442b4c [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>

[ROCm/composable_kernel commit: 3049b5467c]

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

..

data_type_enum.hpp

update copyright headers (#726)

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

profile_avg_pool3d_bwd_impl.hpp

MaxPool & AvgPool bwd instances, test, ckProfiler, client example (#861)

2023-08-31 21:01:50 +08:00

profile_batched_gemm_add_relu_gemm_add_impl.hpp

update copyright headers (#726)

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

profile_batched_gemm_bias_softmax_gemm_permute_impl.hpp

update copyright headers (#726)

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

profile_batched_gemm_gemm_impl.hpp

update copyright headers (#726)

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

profile_batched_gemm_impl.hpp

Add DeviceBatchedGemmMultipleD_Dl (#732)

2023-06-12 08:37:15 -05:00

profile_batched_gemm_reduce_impl.hpp

update copyright headers (#726)

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

profile_batched_gemm_softmax_gemm_impl.hpp

update copyright headers (#726)

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

profile_batched_gemm_softmax_gemm_permute_impl.hpp

update copyright headers (#726)

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

profile_batchnorm_backward_impl.hpp

update copyright headers (#726)

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

profile_batchnorm_forward_impl.hpp

update copyright headers (#726)

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

profile_batchnorm_infer_impl.hpp

update copyright headers (#726)

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

profile_contraction_impl.hpp

Extend support for contraction 6D (#1207)

2024-04-09 23:46:21 +02:00

profile_contraction_utils.hpp

Extend support for contraction 6D (#1207)

2024-04-09 23:46:21 +02:00

profile_conv_bwd_data_impl.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432)

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

profile_conv_fwd_bias_relu_add_impl.hpp

update copyright headers (#726)

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

profile_conv_fwd_bias_relu_impl.hpp

update copyright headers (#726)

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

profile_conv_fwd_impl.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432)

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

profile_conv_tensor_rearrange_impl.hpp

Add support for groups in Img2Col/Col2Img (#1007)

2023-10-31 10:46:32 +01:00

profile_elementwise_layernorm_impl.hpp

Clip fp8 to +/-240 on all targets. (#1172)

2024-02-27 14:31:05 -06:00

profile_gemm_ab_scale_impl.hpp

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

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

profile_gemm_add_add_fastgelu_impl.hpp

update copyright headers (#726)

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

profile_gemm_add_fastgelu_impl.hpp

update copyright headers (#726)

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

profile_gemm_add_impl.hpp

Add support for mixed-precision f16bf16_int8 gemm (#1127)

2024-02-07 15:54:13 +01:00

profile_gemm_add_multiply_impl.hpp

update copyright headers (#726)

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

profile_gemm_add_relu_add_layernorm_impl.hpp

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

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

profile_gemm_add_relu_impl.hpp

Add support for mixed-precision f16bf16_int8 gemm (#1127)

2024-02-07 15:54:13 +01:00

profile_gemm_add_silu_impl.hpp

Add support for mixed-precision f16bf16_int8 gemm (#1127)

2024-02-07 15:54:13 +01:00

profile_gemm_bias_add_reduce_impl.hpp

update copyright headers (#726)

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

profile_gemm_bilinear_impl.hpp

update copyright headers (#726)

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

profile_gemm_fastgelu_impl.hpp

update copyright headers (#726)

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

profile_gemm_impl.hpp

fixed return (#1138)

2024-01-22 08:42:26 -08:00

profile_gemm_multiply_add_impl.hpp

fixed fp8 issues (#894)

2023-09-12 22:17:56 -05:00

profile_gemm_multiply_multiply_impl.hpp

[GEMM] gemm_universal related optimization (#1453)

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

profile_gemm_reduce_impl.hpp

update copyright headers (#726)

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

profile_gemm_splitk_impl.hpp

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

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

profile_gemm_streamk_impl.hpp

add generic instances (#858)

2023-08-23 09:18:10 -05:00

profile_gemm_universal_impl.hpp

[GEMM] gemm_universal related optimization (#1453)

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

profile_gemm_universal_reduce_impl.hpp

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

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

profile_gemm_universal_streamk_impl.hpp

Universal streamk with atomics (#1360)

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

profile_grouped_conv_bwd_data_impl.hpp

Support bf16/f32/f16 and NHWGC conv2d_bwd_data (#757)

2023-06-21 08:20:31 -05:00

profile_grouped_conv_bwd_weight_impl.hpp

Add two stage grouped conv bwd weight kernel (#1280)

2024-05-08 09:53:24 +02:00

profile_grouped_conv_fwd_impl.hpp

Add Grouped Conv Fwd Large Tensor kernel (#1432)

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

profile_grouped_conv_fwd_outelementop_impl.hpp

Add ckProfiler support for forward 3D convolutions with OUT element-wise operations. (#1354)

2024-07-08 10:55:54 -07:00

profile_grouped_gemm_fastgelu_impl.hpp

update copyright headers (#726)

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

profile_grouped_gemm_fixed_nk_impl.hpp

replace the ENV macro with CK_ENV (#1296)

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

profile_grouped_gemm_impl.hpp

replace the ENV macro with CK_ENV (#1296)

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

profile_grouped_gemm_multiply_tile_loop_impl.hpp

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

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

profile_grouped_gemm_tile_loop_impl.hpp

replace the ENV macro with CK_ENV (#1296)

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

profile_grouped_gemm_two_stage_impl.hpp

replace the ENV macro with CK_ENV (#1296)

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

profile_groupnorm_bwd_data_impl.hpp

layernorm and groupnorm backward data (#1083)

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

profile_groupnorm_bwd_gamma_beta_impl.hpp

layernorm & groupnorm bwd gamma beta (#1133)

2024-01-25 19:53:15 +08:00

profile_groupnorm_fwd_impl.hpp

Layernorm4d (#1022)

2023-11-09 08:34:51 +08:00

profile_layernorm_bwd_data_impl.hpp

layernorm and groupnorm backward data (#1083)

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

profile_layernorm_bwd_gamma_beta_impl.hpp

layernorm & groupnorm bwd gamma beta (#1133)

2024-01-25 19:53:15 +08:00

profile_layernorm_fwd_impl.hpp

Layernorm4d (#1022)

2023-11-09 08:34:51 +08:00

profile_max_pool3d_bwd_impl.hpp

MaxPool & AvgPool bwd instances, test, ckProfiler, client example (#861)

2023-08-31 21:01:50 +08:00

profile_permute_scale_impl.hpp

Introduce combined elementwise ops (#1217)

2024-04-02 17:23:49 -05:00

profile_pool3d_fwd_impl.hpp

Refactor pool fwd (#815)

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

profile_reduce_impl.hpp

update copyright headers (#726)

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

profile_softmax_impl.hpp

Padded Generic Kernel Instance (#730)

2023-06-16 23:43:11 -05:00

profile_transpose_impl.hpp

Refactor elementwise kernels (#1222)

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