From fd12e33f27fbcd02ec4a7c87941537c8b7baa862 Mon Sep 17 00:00:00 2001 From: Illia Silin <98187287+illsilin@users.noreply.github.com> Date: Wed, 15 Oct 2025 07:43:11 -0700 Subject: [PATCH] re-enable clang-format by default (#3030) * re-enable clang-format by default * fix clang format [ROCm/composable_kernel commit: 3348f01e6fc65a7afcea3ea4167cc70e902e854a] --- Jenkinsfile | 1 - .../20_grouped_convolution/gemm_configs.hpp | 12 +-- .../38_block_scale_gemm/gemm_quant_basic.cpp | 5 +- ...ouped_convolution_backward_data_kernel.hpp | 70 ++++++++-------- ...ped_convolution_backward_weight_kernel.hpp | 70 ++++++++-------- .../grouped_convolution_forward_kernel.hpp | 84 +++++++++---------- .../profile_grouped_conv_bwd_data_impl.hpp | 8 +- 7 files changed, 126 insertions(+), 124 deletions(-) diff --git a/Jenkinsfile b/Jenkinsfile index 11a9d9eb74..d934e3b2b4 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -1190,7 +1190,6 @@ pipeline { when { beforeAgent true expression { env.SHOULD_RUN_CI.toBoolean() } - expression { params.RUN_CPPCHECK.toBoolean() } } parallel{ stage('Clang Format and Cppcheck') { diff --git a/example/ck_tile/20_grouped_convolution/gemm_configs.hpp b/example/ck_tile/20_grouped_convolution/gemm_configs.hpp index 37a63cd65c..77e1c3af1a 100644 --- a/example/ck_tile/20_grouped_convolution/gemm_configs.hpp +++ b/example/ck_tile/20_grouped_convolution/gemm_configs.hpp @@ -226,20 +226,20 @@ struct ConvTypeConfig; template <> struct ConvTypeConfig { - using InDataType = ck_tile::half_t; - using WeiDataType = ck_tile::half_t; + using InDataType = ck_tile::half_t; + using WeiDataType = ck_tile::half_t; using AccDataType = float; - using OutDataType = ck_tile::half_t; + using OutDataType = ck_tile::half_t; // ToDo: Add more bias config to support different categories of GEMM. }; template <> struct ConvTypeConfig { - using InDataType = ck_tile::bf16_t; - using WeiDataType = ck_tile::bf16_t; + using InDataType = ck_tile::bf16_t; + using WeiDataType = ck_tile::bf16_t; using AccDataType = float; - using OutDataType = ck_tile::bf16_t; + using OutDataType = ck_tile::bf16_t; }; template diff --git a/example/ck_tile/38_block_scale_gemm/gemm_quant_basic.cpp b/example/ck_tile/38_block_scale_gemm/gemm_quant_basic.cpp index c9cc56d033..0752dfdde4 100644 --- a/example/ck_tile/38_block_scale_gemm/gemm_quant_basic.cpp +++ b/example/ck_tile/38_block_scale_gemm/gemm_quant_basic.cpp @@ -451,4 +451,7 @@ int run_gemm_example(int argc, char* argv[]) } } -int main(int argc, char* argv[]) { return !run_gemm_example(argc, argv); } +int main(int argc, char* argv[]) +{ + return !run_gemm_example(argc, argv); +} diff --git a/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_data_kernel.hpp b/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_data_kernel.hpp index 071ea2dccc..15c56f9261 100644 --- a/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_data_kernel.hpp +++ b/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_data_kernel.hpp @@ -44,13 +44,13 @@ struct GroupedConvBwdDataKernelArgs CK_TILE_HOST GroupedConvBwdDataKernelArgs(const GroupedConvBwdDataHostArgs& args) { in_g_n_c_wis_lengths = {static_cast(args.G_), - static_cast(args.N_), - static_cast(args.C_), - static_cast(args.input_spatial_lengths_[0])}; + static_cast(args.N_), + static_cast(args.C_), + static_cast(args.input_spatial_lengths_[0])}; wei_g_k_c_xs_lengths = {static_cast(args.G_), - static_cast(args.K_), - static_cast(args.C_), - static_cast(args.filter_spatial_lengths_[0])}; + static_cast(args.K_), + static_cast(args.C_), + static_cast(args.filter_spatial_lengths_[0])}; out_g_n_k_wos_lengths = {static_cast(args.G_), static_cast(args.N_), static_cast(args.K_), @@ -145,15 +145,15 @@ struct GroupedConvBwdDataKernelArgs CK_TILE_HOST GroupedConvBwdDataKernelArgs(const GroupedConvBwdDataHostArgs& args) { in_g_n_c_wis_lengths = {static_cast(args.G_), - static_cast(args.N_), - static_cast(args.C_), - static_cast(args.input_spatial_lengths_[0]), - static_cast(args.input_spatial_lengths_[1])}; + static_cast(args.N_), + static_cast(args.C_), + static_cast(args.input_spatial_lengths_[0]), + static_cast(args.input_spatial_lengths_[1])}; wei_g_k_c_xs_lengths = {static_cast(args.G_), - static_cast(args.K_), - static_cast(args.C_), - static_cast(args.filter_spatial_lengths_[0]), - static_cast(args.filter_spatial_lengths_[1])}; + static_cast(args.K_), + static_cast(args.C_), + static_cast(args.filter_spatial_lengths_[0]), + static_cast(args.filter_spatial_lengths_[1])}; out_g_n_k_wos_lengths = {static_cast(args.G_), static_cast(args.N_), static_cast(args.K_), @@ -161,13 +161,13 @@ struct GroupedConvBwdDataKernelArgs static_cast(args.output_spatial_lengths_[1])}; conv_filter_strides = {static_cast(args.conv_filter_strides_[0]), - static_cast(args.conv_filter_strides_[1])}; + static_cast(args.conv_filter_strides_[1])}; conv_filter_dilations = {static_cast(args.conv_filter_dilations_[0]), static_cast(args.conv_filter_dilations_[1])}; input_left_pads = {static_cast(args.input_left_pads_[0]), - static_cast(args.input_left_pads_[1])}; + static_cast(args.input_left_pads_[1])}; input_right_pads = {static_cast(args.input_right_pads_[0]), - static_cast(args.input_right_pads_[1])}; + static_cast(args.input_right_pads_[1])}; k_batch = args.k_batch; @@ -262,17 +262,17 @@ struct GroupedConvBwdDataKernelArgs CK_TILE_HOST GroupedConvBwdDataKernelArgs(const GroupedConvBwdDataHostArgs& args) { in_g_n_c_wis_lengths = {static_cast(args.G_), - static_cast(args.N_), - static_cast(args.C_), - static_cast(args.input_spatial_lengths_[0]), - static_cast(args.input_spatial_lengths_[1]), - static_cast(args.input_spatial_lengths_[2])}; + static_cast(args.N_), + static_cast(args.C_), + static_cast(args.input_spatial_lengths_[0]), + static_cast(args.input_spatial_lengths_[1]), + static_cast(args.input_spatial_lengths_[2])}; wei_g_k_c_xs_lengths = {static_cast(args.G_), - static_cast(args.K_), - static_cast(args.C_), - static_cast(args.filter_spatial_lengths_[0]), - static_cast(args.filter_spatial_lengths_[1]), - static_cast(args.filter_spatial_lengths_[2])}; + static_cast(args.K_), + static_cast(args.C_), + static_cast(args.filter_spatial_lengths_[0]), + static_cast(args.filter_spatial_lengths_[1]), + static_cast(args.filter_spatial_lengths_[2])}; out_g_n_k_wos_lengths = {static_cast(args.G_), static_cast(args.N_), static_cast(args.K_), @@ -281,17 +281,17 @@ struct GroupedConvBwdDataKernelArgs static_cast(args.output_spatial_lengths_[2])}; conv_filter_strides = {static_cast(args.conv_filter_strides_[0]), - static_cast(args.conv_filter_strides_[1]), - static_cast(args.conv_filter_strides_[2])}; + static_cast(args.conv_filter_strides_[1]), + static_cast(args.conv_filter_strides_[2])}; conv_filter_dilations = {static_cast(args.conv_filter_dilations_[0]), static_cast(args.conv_filter_dilations_[1]), static_cast(args.conv_filter_dilations_[2])}; input_left_pads = {static_cast(args.input_left_pads_[0]), - static_cast(args.input_left_pads_[1]), - static_cast(args.input_left_pads_[2])}; + static_cast(args.input_left_pads_[1]), + static_cast(args.input_left_pads_[2])}; input_right_pads = {static_cast(args.input_right_pads_[0]), - static_cast(args.input_right_pads_[1]), - static_cast(args.input_right_pads_[2])}; + static_cast(args.input_right_pads_[1]), + static_cast(args.input_right_pads_[2])}; k_batch = args.k_batch; @@ -387,8 +387,8 @@ struct GroupedConvBwdDataKernelArgs static constexpr index_t MaxGroupedGemmGroupsNum = 128; - using ABCGridDescs = remove_cvref_t; + using ABCGridDescs = remove_cvref_t< + decltype(ConvToGemmTransformer{}.MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N(1))>; using AGridDescMK = remove_cvref_t{}])>; using BGridDescNK = remove_cvref_t{}])>; diff --git a/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp b/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp index 14a04615dd..83ecb34a79 100644 --- a/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp +++ b/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_weight_kernel.hpp @@ -40,13 +40,13 @@ struct GroupedConvBwdWeightKernelArgs CK_TILE_HOST GroupedConvBwdWeightKernelArgs(const GroupedConvBwdWeightHostArgs& args) { in_g_n_c_wis_lengths = {static_cast(args.G_), - static_cast(args.N_), - static_cast(args.C_), - static_cast(args.input_spatial_lengths_[0])}; + static_cast(args.N_), + static_cast(args.C_), + static_cast(args.input_spatial_lengths_[0])}; wei_g_k_c_xs_lengths = {static_cast(args.G_), - static_cast(args.K_), - static_cast(args.C_), - static_cast(args.filter_spatial_lengths_[0])}; + static_cast(args.K_), + static_cast(args.C_), + static_cast(args.filter_spatial_lengths_[0])}; out_g_n_k_wos_lengths = {static_cast(args.G_), static_cast(args.N_), static_cast(args.K_), @@ -109,15 +109,15 @@ struct GroupedConvBwdWeightKernelArgs CK_TILE_HOST GroupedConvBwdWeightKernelArgs(const GroupedConvBwdWeightHostArgs& args) { in_g_n_c_wis_lengths = {static_cast(args.G_), - static_cast(args.N_), - static_cast(args.C_), - static_cast(args.input_spatial_lengths_[0]), - static_cast(args.input_spatial_lengths_[1])}; + static_cast(args.N_), + static_cast(args.C_), + static_cast(args.input_spatial_lengths_[0]), + static_cast(args.input_spatial_lengths_[1])}; wei_g_k_c_xs_lengths = {static_cast(args.G_), - static_cast(args.K_), - static_cast(args.C_), - static_cast(args.filter_spatial_lengths_[0]), - static_cast(args.filter_spatial_lengths_[1])}; + static_cast(args.K_), + static_cast(args.C_), + static_cast(args.filter_spatial_lengths_[0]), + static_cast(args.filter_spatial_lengths_[1])}; out_g_n_k_wos_lengths = {static_cast(args.G_), static_cast(args.N_), static_cast(args.K_), @@ -125,13 +125,13 @@ struct GroupedConvBwdWeightKernelArgs static_cast(args.output_spatial_lengths_[1])}; conv_filter_strides = {static_cast(args.conv_filter_strides_[0]), - static_cast(args.conv_filter_strides_[1])}; + static_cast(args.conv_filter_strides_[1])}; conv_filter_dilations = {static_cast(args.conv_filter_dilations_[0]), static_cast(args.conv_filter_dilations_[1])}; input_left_pads = {static_cast(args.input_left_pads_[0]), - static_cast(args.input_left_pads_[1])}; + static_cast(args.input_left_pads_[1])}; input_right_pads = {static_cast(args.input_right_pads_[0]), - static_cast(args.input_right_pads_[1])}; + static_cast(args.input_right_pads_[1])}; k_batch = args.k_batch; @@ -185,17 +185,17 @@ struct GroupedConvBwdWeightKernelArgs CK_TILE_HOST GroupedConvBwdWeightKernelArgs(const GroupedConvBwdWeightHostArgs& args) { in_g_n_c_wis_lengths = {static_cast(args.G_), - static_cast(args.N_), - static_cast(args.C_), - static_cast(args.input_spatial_lengths_[0]), - static_cast(args.input_spatial_lengths_[1]), - static_cast(args.input_spatial_lengths_[2])}; + static_cast(args.N_), + static_cast(args.C_), + static_cast(args.input_spatial_lengths_[0]), + static_cast(args.input_spatial_lengths_[1]), + static_cast(args.input_spatial_lengths_[2])}; wei_g_k_c_xs_lengths = {static_cast(args.G_), - static_cast(args.K_), - static_cast(args.C_), - static_cast(args.filter_spatial_lengths_[0]), - static_cast(args.filter_spatial_lengths_[1]), - static_cast(args.filter_spatial_lengths_[2])}; + static_cast(args.K_), + static_cast(args.C_), + static_cast(args.filter_spatial_lengths_[0]), + static_cast(args.filter_spatial_lengths_[1]), + static_cast(args.filter_spatial_lengths_[2])}; out_g_n_k_wos_lengths = {static_cast(args.G_), static_cast(args.N_), static_cast(args.K_), @@ -204,17 +204,17 @@ struct GroupedConvBwdWeightKernelArgs static_cast(args.output_spatial_lengths_[2])}; conv_filter_strides = {static_cast(args.conv_filter_strides_[0]), - static_cast(args.conv_filter_strides_[1]), - static_cast(args.conv_filter_strides_[2])}; + static_cast(args.conv_filter_strides_[1]), + static_cast(args.conv_filter_strides_[2])}; conv_filter_dilations = {static_cast(args.conv_filter_dilations_[0]), static_cast(args.conv_filter_dilations_[1]), static_cast(args.conv_filter_dilations_[2])}; input_left_pads = {static_cast(args.input_left_pads_[0]), - static_cast(args.input_left_pads_[1]), - static_cast(args.input_left_pads_[2])}; + static_cast(args.input_left_pads_[1]), + static_cast(args.input_left_pads_[2])}; input_right_pads = {static_cast(args.input_right_pads_[0]), - static_cast(args.input_right_pads_[1]), - static_cast(args.input_right_pads_[2])}; + static_cast(args.input_right_pads_[1]), + static_cast(args.input_right_pads_[2])}; k_batch = args.k_batch; @@ -257,8 +257,8 @@ struct GroupedConvBwdWeightKernelArgs GemmBatch = args.G_; } - using ABCGridDescs = remove_cvref_t; + using ABCGridDescs = remove_cvref_t< + decltype(ConvToGemmTransformer{}.MakeABCGridDescriptor_A_K0_M_K1_B_K0_N_K1_C_M_N())>; using AGridDescKM = remove_cvref_t{}])>; using BGridDescKN = remove_cvref_t{}])>; diff --git a/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_forward_kernel.hpp b/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_forward_kernel.hpp index 7d7f8b1cf2..0363782d33 100644 --- a/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_forward_kernel.hpp +++ b/include/ck_tile/ops/grouped_convolution/kernel/grouped_convolution_forward_kernel.hpp @@ -41,13 +41,13 @@ struct GroupedConvFwdKernelArgs CK_TILE_HOST GroupedConvFwdKernelArgs(const GroupedConvFwdHostArgs& args) { in_g_n_c_wis_lengths = {static_cast(args.G_), - static_cast(args.N_), - static_cast(args.C_), - static_cast(args.input_spatial_lengths_[0])}; + static_cast(args.N_), + static_cast(args.C_), + static_cast(args.input_spatial_lengths_[0])}; wei_g_k_c_xs_lengths = {static_cast(args.G_), - static_cast(args.K_), - static_cast(args.C_), - static_cast(args.filter_spatial_lengths_[0])}; + static_cast(args.K_), + static_cast(args.C_), + static_cast(args.filter_spatial_lengths_[0])}; out_g_n_k_wos_lengths = {static_cast(args.G_), static_cast(args.N_), static_cast(args.K_), @@ -124,15 +124,15 @@ struct GroupedConvFwdKernelArgs CK_TILE_HOST GroupedConvFwdKernelArgs(const GroupedConvFwdHostArgs& args) { in_g_n_c_wis_lengths = {static_cast(args.G_), - static_cast(args.N_), - static_cast(args.C_), - static_cast(args.input_spatial_lengths_[0]), - static_cast(args.input_spatial_lengths_[1])}; + static_cast(args.N_), + static_cast(args.C_), + static_cast(args.input_spatial_lengths_[0]), + static_cast(args.input_spatial_lengths_[1])}; wei_g_k_c_xs_lengths = {static_cast(args.G_), - static_cast(args.K_), - static_cast(args.C_), - static_cast(args.filter_spatial_lengths_[0]), - static_cast(args.filter_spatial_lengths_[1])}; + static_cast(args.K_), + static_cast(args.C_), + static_cast(args.filter_spatial_lengths_[0]), + static_cast(args.filter_spatial_lengths_[1])}; out_g_n_k_wos_lengths = {static_cast(args.G_), static_cast(args.N_), static_cast(args.K_), @@ -140,13 +140,13 @@ struct GroupedConvFwdKernelArgs static_cast(args.output_spatial_lengths_[1])}; conv_filter_strides = {static_cast(args.conv_filter_strides_[0]), - static_cast(args.conv_filter_strides_[1])}; + static_cast(args.conv_filter_strides_[1])}; conv_filter_dilations = {static_cast(args.conv_filter_dilations_[0]), static_cast(args.conv_filter_dilations_[1])}; input_left_pads = {static_cast(args.input_left_pads_[0]), - static_cast(args.input_left_pads_[1])}; + static_cast(args.input_left_pads_[1])}; input_right_pads = {static_cast(args.input_right_pads_[0]), - static_cast(args.input_right_pads_[1])}; + static_cast(args.input_right_pads_[1])}; k_batch = args.k_batch; @@ -216,17 +216,17 @@ struct GroupedConvFwdKernelArgs CK_TILE_HOST GroupedConvFwdKernelArgs(const GroupedConvFwdHostArgs& args) { in_g_n_c_wis_lengths = {static_cast(args.G_), - static_cast(args.N_), - static_cast(args.C_), - static_cast(args.input_spatial_lengths_[0]), - static_cast(args.input_spatial_lengths_[1]), - static_cast(args.input_spatial_lengths_[2])}; + static_cast(args.N_), + static_cast(args.C_), + static_cast(args.input_spatial_lengths_[0]), + static_cast(args.input_spatial_lengths_[1]), + static_cast(args.input_spatial_lengths_[2])}; wei_g_k_c_xs_lengths = {static_cast(args.G_), - static_cast(args.K_), - static_cast(args.C_), - static_cast(args.filter_spatial_lengths_[0]), - static_cast(args.filter_spatial_lengths_[1]), - static_cast(args.filter_spatial_lengths_[2])}; + static_cast(args.K_), + static_cast(args.C_), + static_cast(args.filter_spatial_lengths_[0]), + static_cast(args.filter_spatial_lengths_[1]), + static_cast(args.filter_spatial_lengths_[2])}; out_g_n_k_wos_lengths = {static_cast(args.G_), static_cast(args.N_), static_cast(args.K_), @@ -235,17 +235,17 @@ struct GroupedConvFwdKernelArgs static_cast(args.output_spatial_lengths_[2])}; conv_filter_strides = {static_cast(args.conv_filter_strides_[0]), - static_cast(args.conv_filter_strides_[1]), - static_cast(args.conv_filter_strides_[2])}; + static_cast(args.conv_filter_strides_[1]), + static_cast(args.conv_filter_strides_[2])}; conv_filter_dilations = {static_cast(args.conv_filter_dilations_[0]), static_cast(args.conv_filter_dilations_[1]), static_cast(args.conv_filter_dilations_[2])}; input_left_pads = {static_cast(args.input_left_pads_[0]), - static_cast(args.input_left_pads_[1]), - static_cast(args.input_left_pads_[2])}; + static_cast(args.input_left_pads_[1]), + static_cast(args.input_left_pads_[2])}; input_right_pads = {static_cast(args.input_right_pads_[0]), - static_cast(args.input_right_pads_[1]), - static_cast(args.input_right_pads_[2])}; + static_cast(args.input_right_pads_[1]), + static_cast(args.input_right_pads_[2])}; k_batch = args.k_batch; @@ -306,15 +306,15 @@ struct GroupedConvFwdKernelArgs args.output_spatial_lengths_[2]; } - using AGridDescMK = remove_cvref_t())>; - using BGridDescNK = remove_cvref_t())>; - using CGridDescMN = remove_cvref_t())>; + using AGridDescMK = remove_cvref_t< + decltype(ConvToGemmFwdTransformer{} + .template MakeADescriptor_M_K())>; + using BGridDescNK = remove_cvref_t< + decltype(ConvToGemmFwdTransformer{} + .template MakeBDescriptor_N_K())>; + using CGridDescMN = remove_cvref_t< + decltype(ConvToGemmFwdTransformer{} + .template MakeCDescriptor_M_N())>; static constexpr index_t NonSpatialDims = 3; array in_g_n_c_wis_lengths; diff --git a/profiler/include/profiler/profile_grouped_conv_bwd_data_impl.hpp b/profiler/include/profiler/profile_grouped_conv_bwd_data_impl.hpp index 0b73fe7adf..2369b2eac8 100644 --- a/profiler/include/profiler/profile_grouped_conv_bwd_data_impl.hpp +++ b/profiler/include/profiler/profile_grouped_conv_bwd_data_impl.hpp @@ -177,12 +177,12 @@ bool profile_grouped_conv_bwd_data_impl(int do_verification, in_device_buf.FromDevice(in_device.mData.data()); using ComputeType_ = std::conditional_t; + OutDataType, + WeiDataType>; using ComputeType = std::conditional_t; + ComputeType_, + ComputeDataType>; using AccDataType = std::conditional_t, int32_t, float>; const index_t num_accums = conv_param.K_;