ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-17 03:19:48 +00:00
Code Issues Packages Projects Releases Wiki Activity
Files
ad0a8e4cd27e8b66781d18ca6a7b1190e0611597
composable_kernel/library/src/tensor_operation_instance/gpu/gemm
History
Bartlomiej Wroblewski ad0a8e4cd2 Optimize fp16 direct load GEMM instances (#1086) 
This PR optimizes fp16 instances of direct load GEMM kernel introduced in #999 and #1052.

Measured the performance of new instances on CDNA2 GPU and compared it against the performance of the best non-direct-load GEMM instances. Used 76 different GEMM problems.
On average, this change improves the performance of the tested problems by 47%. For cases known as latency-bound, the speedup is around 126%.
2023-12-18 11:09:10 +01:00
..
device_gemm_xdl_f16_f16_f16
Update tuning parameter & compilation options of DeviceGemmXdl<> instance (layout=TT) (#819)
2023-08-02 10:32:22 -05:00
CMakeLists.txt
disabling some fp8 gemm instances to reduce build time (#1084)
2023-12-11 17:49:27 -08:00
device_gemm_dl_f16_f16_f16_km_kn_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_dl_f16_f16_f16_km_kn_mn_irregular_instance.cpp
Fix DL GEMM instances with too large vector size (#901)
2023-09-18 14:08:23 +02:00
device_gemm_dl_f16_f16_f16_km_nk_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_dl_f16_f16_f16_km_nk_mn_irregular_instance.cpp
Fix DL GEMM instances with too large vector size (#901)
2023-09-18 14:08:23 +02:00
device_gemm_dl_f16_f16_f16_mk_kn_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_dl_f16_f16_f16_mk_kn_mn_irregular_instance.cpp
Fix DL GEMM instances with too large vector size (#901)
2023-09-18 14:08:23 +02:00
device_gemm_dl_f16_f16_f16_mk_nk_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_dl_f16_f16_f16_mk_nk_mn_irregular_instance.cpp
Add instances for fp16/int8 Gemm kernels (Navi21) (#717)
2023-05-30 07:07:17 -05:00
device_gemm_dl_f32_f32_f32_km_kn_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_dl_f32_f32_f32_km_nk_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_dl_f32_f32_f32_mk_kn_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_dl_f32_f32_f32_mk_nk_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_dl_i8_i8_i8_km_kn_mn_instance.cpp
[HotFix] add config and version files to pass on build info (#856)
2023-08-23 11:36:17 -07:00
device_gemm_dl_i8_i8_i8_km_kn_mn_irregular_instance.cpp
Fix DL GEMM instances with too large vector size (#901)
2023-09-18 14:08:23 +02:00
device_gemm_dl_i8_i8_i8_km_nk_mn_instance.cpp
[HotFix] add config and version files to pass on build info (#856)
2023-08-23 11:36:17 -07:00
device_gemm_dl_i8_i8_i8_km_nk_mn_irregular_instance.cpp
Fix DL GEMM instances with too large vector size (#901)
2023-09-18 14:08:23 +02:00
device_gemm_dl_i8_i8_i8_mk_kn_mn_instance.cpp
[HotFix] add config and version files to pass on build info (#856)
2023-08-23 11:36:17 -07:00
device_gemm_dl_i8_i8_i8_mk_kn_mn_irregular_instance.cpp
Fix DL GEMM instances with too large vector size (#901)
2023-09-18 14:08:23 +02:00
device_gemm_dl_i8_i8_i8_mk_nk_mn_instance.cpp
[HotFix] add config and version files to pass on build info (#856)
2023-08-23 11:36:17 -07:00
device_gemm_dl_i8_i8_i8_mk_nk_mn_irregular_instance.cpp
[HotFix] add config and version files to pass on build info (#856)
2023-08-23 11:36:17 -07:00
device_gemm_dpp_f16_f16_f16_km_kn_mn_instance.cpp
Add new instances and support for small cases in DPP8 GEMM (#896)
2023-09-12 10:05:23 -05:00
device_gemm_dpp_f16_f16_f16_km_kn_mn_irregular_instance.cpp
Add new instances and support for small cases in DPP8 GEMM (#896)
2023-09-12 10:05:23 -05:00
device_gemm_dpp_f16_f16_f16_km_nk_mn_instance.cpp
Add new instances and support for small cases in DPP8 GEMM (#896)
2023-09-12 10:05:23 -05:00
device_gemm_dpp_f16_f16_f16_km_nk_mn_irregular_instance.cpp
Add new instances and support for small cases in DPP8 GEMM (#896)
2023-09-12 10:05:23 -05:00
device_gemm_dpp_f16_f16_f16_mk_kn_mn_instance.cpp
Add new instances and support for small cases in DPP8 GEMM (#896)
2023-09-12 10:05:23 -05:00
device_gemm_dpp_f16_f16_f16_mk_kn_mn_irregular_instance.cpp
Add new instances and support for small cases in DPP8 GEMM (#896)
2023-09-12 10:05:23 -05:00
device_gemm_dpp_f16_f16_f16_mk_nk_mn_instance.cpp
Add new instances and support for small cases in DPP8 GEMM (#896)
2023-09-12 10:05:23 -05:00
device_gemm_dpp_f16_f16_f16_mk_nk_mn_irregular_instance.cpp
Add new instances and support for small cases in DPP8 GEMM (#896)
2023-09-12 10:05:23 -05:00
device_gemm_xdl_c_shuffle_2_stage_f16_f16_f16_mk_nk_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_xdl_c_shuffle_bf16_bf16_bf16_km_kn_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_xdl_c_shuffle_bf16_bf16_bf16_km_nk_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_xdl_c_shuffle_bf16_bf16_bf16_mk_kn_mn_instance.cpp
add more instances for bfp16 gemm (#1036)
2023-11-11 07:09:32 -08:00
device_gemm_xdl_c_shuffle_bf16_bf16_bf16_mk_nk_mn_instance.cpp
add more instances for bfp16 gemm (#1036)
2023-11-11 07:09:32 -08:00
device_gemm_xdl_c_shuffle_f16_f8_f16_mk_kn_mn_instance.cpp
fix clang format (#1095)
2023-12-08 14:32:37 -08:00
device_gemm_xdl_c_shuffle_f16_f8_f16_mk_nk_mn_instance.cpp
fix clang format (#1095)
2023-12-08 14:32:37 -08:00
device_gemm_xdl_c_shuffle_f16_f16_f16_km_kn_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_xdl_c_shuffle_f16_f16_f16_km_nk_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_xdl_c_shuffle_f16_f16_f16_mk_kn_mn_instance.cpp
Improve 4k gemm perf (#1047)
2023-11-17 07:06:24 -06:00
device_gemm_xdl_c_shuffle_f16_f16_f16_mk_nk_mn_instance.cpp
Add Gemm instances for performance improvement (#1018)
2023-11-07 09:09:58 -06:00
device_gemm_xdl_c_shuffle_f32_f32_f32_km_kn_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_xdl_c_shuffle_f32_f32_f32_km_nk_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_xdl_c_shuffle_f32_f32_f32_mk_kn_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_xdl_c_shuffle_f32_f32_f32_mk_nk_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_xdl_c_shuffle_fp8_fp8_fp8_km_kn_mn_instance.cpp
Add fp8 gemm instances (#920)
2023-09-26 14:59:33 -05:00
device_gemm_xdl_c_shuffle_fp8_fp8_fp8_km_nk_mn_instance.cpp
Add fp8 gemm instances (#920)
2023-09-26 14:59:33 -05:00
device_gemm_xdl_c_shuffle_fp8_fp8_fp8_mk_kn_mn_default_instance.cpp
disabling some fp8 gemm instances to reduce build time (#1084)
2023-12-11 17:49:27 -08:00
device_gemm_xdl_c_shuffle_fp8_fp8_fp8_mk_kn_mn_padded_instance.cpp
disabling some fp8 gemm instances to reduce build time (#1084)
2023-12-11 17:49:27 -08:00
device_gemm_xdl_c_shuffle_fp8_fp8_fp8_mk_nk_mn_instance.cpp
Enabled padding for regular gemm (#1004)
2023-10-23 16:52:53 -05:00
device_gemm_xdl_c_shuffle_i8_i8_i8_km_kn_mn_instance.cpp
[HotFix] add config and version files to pass on build info (#856)
2023-08-23 11:36:17 -07:00
device_gemm_xdl_c_shuffle_i8_i8_i8_km_nk_mn_instance.cpp
[HotFix] add config and version files to pass on build info (#856)
2023-08-23 11:36:17 -07:00
device_gemm_xdl_c_shuffle_i8_i8_i8_mk_kn_mn_instance.cpp
[HotFix] add config and version files to pass on build info (#856)
2023-08-23 11:36:17 -07:00
device_gemm_xdl_c_shuffle_i8_i8_i8_mk_nk_mn_instance.cpp
[HotFix] add config and version files to pass on build info (#856)
2023-08-23 11:36:17 -07:00
device_gemm_xdl_c_shuffle_lds_direct_load_f16_f16_f16_mk_nk_mn_instance.cpp
Optimize fp16 direct load GEMM instances (#1086)
2023-12-18 11:09:10 +01:00
device_gemm_xdl_c_shuffle_lds_direct_load_f32_f32_f32_km_kn_mn_instance.cpp
Add support for double buffering in direct load GEMM kernel (#1052)
2023-12-03 23:08:47 +01:00
device_gemm_xdl_c_shuffle_lds_direct_load_f32_f32_f32_km_nk_mn_instance.cpp
Add support for double buffering in direct load GEMM kernel (#1052)
2023-12-03 23:08:47 +01:00
device_gemm_xdl_c_shuffle_lds_direct_load_f32_f32_f32_mk_kn_mn_instance.cpp
Add support for double buffering in direct load GEMM kernel (#1052)
2023-12-03 23:08:47 +01:00
device_gemm_xdl_c_shuffle_lds_direct_load_f32_f32_f32_mk_nk_mn_instance.cpp
Add support for double buffering in direct load GEMM kernel (#1052)
2023-12-03 23:08:47 +01:00
device_gemm_xdl_f32_f32_f32_km_kn_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_xdl_f32_f32_f32_km_nk_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_xdl_f32_f32_f32_mk_kn_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_xdl_f32_f32_f32_mk_nk_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_xdl_f64_f64_f64_km_kn_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_xdl_f64_f64_f64_km_nk_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_xdl_f64_f64_f64_mk_kn_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00
device_gemm_xdl_f64_f64_f64_mk_nk_mn_instance.cpp
update copyright headers (#726)
2023-05-31 18:46:57 -05:00