Split k f16 (#97)

* init for splitk f16 * a working prototype * debug * perf debug * update example * instances for mk kn * add instances for all layers * clean * clean * add tuning * format * add mn_padding into irregular tile * clean Co-authored-by: Chao Liu <[email protected]>
ROCm · Feb 25, 2022 · e221d11 · e221d11
1 parent bdedf64
commit e221d11
Show file tree

Hide file tree

Showing 11 changed files with 1,713 additions and 30 deletions.
diff --git a/composable_kernel/include/tensor_operation/gridwise_gemm_xdlops_v2r4r2.hpp b/composable_kernel/include/tensor_operation/gridwise_gemm_xdlops_v2r4r2.hpp
diff --git a/device_operation/CMakeLists.txt b/device_operation/CMakeLists.txt
@@ -31,7 +31,11 @@ set(DEVICE_GEMM_INSTANCE_SOURCE
    ${PROJECT_SOURCE_DIR}/device_operation/src/device_gemm_xdl_splitk_f32_f32_f32_mk_nk_mn_instance.cpp;
    ${PROJECT_SOURCE_DIR}/device_operation/src/device_gemm_xdl_splitk_f32_f32_f32_km_kn_mn_instance.cpp;
    ${PROJECT_SOURCE_DIR}/device_operation/src/device_gemm_xdl_splitk_f32_f32_f32_km_nk_mn_instance.cpp;
-)
+   ${PROJECT_SOURCE_DIR}/device_operation/src/device_gemm_xdl_splitk_f16_f16_f16_mk_kn_mn_instance.cpp;
+   ${PROJECT_SOURCE_DIR}/device_operation/src/device_gemm_xdl_splitk_f16_f16_f16_mk_nk_mn_instance.cpp;
+   ${PROJECT_SOURCE_DIR}/device_operation/src/device_gemm_xdl_splitk_f16_f16_f16_km_kn_mn_instance.cpp;
+   ${PROJECT_SOURCE_DIR}/device_operation/src/device_gemm_xdl_splitk_f16_f16_f16_km_nk_mn_instance.cpp;
+) 
 
 # device_gemm_bias_2d_instance
 set(DEVICE_GEMM_BIAS_2D_INSTANCE_SOURCE

diff --git a/device_operation/include/conv_utils.hpp b/device_operation/include/conv_utils.hpp
@@ -39,12 +39,12 @@ std::size_t GetFlops(ck::index_t N,
            std::accumulate(std::begin(output_spatial_lengths),
                            std::end(output_spatial_lengths),
                            static_cast<std::size_t>(1),
-                           std::multiplies<std::size_t>()) * 
+                           std::multiplies<std::size_t>()) *
            C *
            std::accumulate(std::begin(filter_spatial_lengths),
                            std::end(filter_spatial_lengths),
                            static_cast<std::size_t>(1),
-                           std::multiplies<std::size_t>());       
+                           std::multiplies<std::size_t>());
 }
 
 /**