From dbc623f45519f5ad10f5845ac877574ba0dece1f Mon Sep 17 00:00:00 2001 From: Bartlomiej Kocot Date: Mon, 29 Sep 2025 22:54:04 +0000 Subject: [PATCH] tmp --- ...uped_convolution_backward_data_invoker.hpp | 4 +- ...ed_convolution_backward_weight_invoker.hpp | 4 +- ...tion_backward_weight_two_stage_invoker.hpp | 4 +- ...ouped_convolution_backward_data_kernel.hpp | 39 +++++++++---------- .../utils/grouped_convolution_utils.hpp | 14 ++----- .../utils/transform_conv_bwd_data_to_gemm.hpp | 26 ++++++------- 6 files changed, 41 insertions(+), 50 deletions(-) diff --git a/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_data_invoker.hpp b/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_data_invoker.hpp index 1b3d45427d..890b9c63be 100644 --- a/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_data_invoker.hpp +++ b/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_data_invoker.hpp @@ -36,8 +36,8 @@ struct GroupedConvolutionBackwardDataInvoker constexpr ck_tile::index_t N_Warp_Tile = GemmWarpConfig::N_Warp_Tile; constexpr ck_tile::index_t K_Warp_Tile = GemmWarpConfig::K_Warp_Tile; - constexpr ck_tile::index_t VectorSizeA = 1; - constexpr ck_tile::index_t VectorSizeB = 1; + constexpr ck_tile::index_t VectorSizeA = 8; + constexpr ck_tile::index_t VectorSizeB = 8; constexpr ck_tile::index_t VectorSizeC = 8; // Implicit GEMM Traits diff --git a/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_weight_invoker.hpp b/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_weight_invoker.hpp index f46707d1d2..74969e4db1 100644 --- a/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_weight_invoker.hpp +++ b/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_weight_invoker.hpp @@ -35,8 +35,8 @@ struct GroupedConvolutionBackwardWeightInvoker constexpr ck_tile::index_t N_Warp_Tile = GemmWarpConfig::N_Warp_Tile; constexpr ck_tile::index_t K_Warp_Tile = GemmWarpConfig::K_Warp_Tile; - constexpr ck_tile::index_t VectorSizeA = 1; - constexpr ck_tile::index_t VectorSizeB = 1; + constexpr ck_tile::index_t VectorSizeA = 8; + constexpr ck_tile::index_t VectorSizeB = 8; constexpr ck_tile::index_t VectorSizeC = 8; // Implicit GEMM Traits diff --git a/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_weight_two_stage_invoker.hpp b/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_weight_two_stage_invoker.hpp index a8f321e7f9..58c73f48b7 100644 --- a/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_weight_two_stage_invoker.hpp +++ b/example/ck_tile/20_grouped_convolution/grouped_convolution_backward_weight_two_stage_invoker.hpp @@ -37,8 +37,8 @@ struct GroupedConvolutionBackwardWeightTwoStageInvoker constexpr ck_tile::index_t N_Warp_Tile = GemmWarpConfig::N_Warp_Tile; constexpr ck_tile::index_t K_Warp_Tile = GemmWarpConfig::K_Warp_Tile; - constexpr ck_tile::index_t VectorSizeA = 1; - constexpr ck_tile::index_t VectorSizeB = 1; + constexpr ck_tile::index_t VectorSizeA = 8; + constexpr ck_tile::index_t VectorSizeB = 8; constexpr ck_tile::index_t VectorSizeC = 1; // Implicit GEMM Traits 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 e68a510a0c..94bba53ae0 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 @@ -109,7 +109,7 @@ struct GroupedConvBwdDataKernelArgs GroupedConvTraitsType_::NDimSpatial>(1); a_grid_descs_m_k[gemm_count] = grid_descs.at(number<0>{}); - b_grid_descs_n_k[gemm_count] = grid_descs.at(number<1>{}); + b_grid_descs_k_n[gemm_count] = grid_descs.at(number<1>{}); c_grid_descs_m_n[gemm_count] = grid_descs.at(number<2>{}); const index_t grid_size_grp = @@ -225,7 +225,7 @@ struct GroupedConvBwdDataKernelArgs GroupedConvTraitsType_::NDimSpatial>(1); a_grid_descs_m_k[gemm_count] = grid_descs.at(number<0>{}); - b_grid_descs_n_k[gemm_count] = grid_descs.at(number<1>{}); + b_grid_descs_k_n[gemm_count] = grid_descs.at(number<1>{}); c_grid_descs_m_n[gemm_count] = grid_descs.at(number<2>{}); const index_t grid_size_grp = @@ -357,7 +357,7 @@ struct GroupedConvBwdDataKernelArgs GroupedConvTraitsType_::NDimSpatial>(1); a_grid_descs_m_k[gemm_count] = grid_descs.at(number<0>{}); - b_grid_descs_n_k[gemm_count] = grid_descs.at(number<1>{}); + b_grid_descs_k_n[gemm_count] = grid_descs.at(number<1>{}); c_grid_descs_m_n[gemm_count] = grid_descs.at(number<2>{}); const index_t grid_size_grp = @@ -416,7 +416,7 @@ struct GroupedConvBwdDataKernelArgs const void* wei_ptr; array a_grid_descs_m_k; - array b_grid_descs_n_k; + array b_grid_descs_k_n; array c_grid_descs_m_n; array block_starts; @@ -471,10 +471,6 @@ template struct GroupedConvolutionBackwardDataKernel { - // Todo: Enable Vector Load Size > 1 - static_assert(GroupedConvTraitsType_::VectorSizeA == 1 && - GroupedConvTraitsType_::VectorSizeB == 1); - static constexpr index_t NDimSpatial = GroupedConvTraitsType_::NDimSpatial_; static constexpr ConvolutionSpecialization ConvSpecialization = GroupedConvTraitsType_::ConvSpecialization; @@ -516,13 +512,10 @@ struct GroupedConvolutionBackwardDataKernel static_assert(GemmPipeline::kPadM && GemmPipeline::kPadN && GemmPipeline::kPadK, "Not supported!"); - static_assert(std::is_same_v, "Not supported!"); - static_assert(std::is_same_v, "Not supported!"); - // TODO: Change to and enable vector load - // static_assert(std::is_same_v, - // "Not supported A GEMM layout!"); - // static_assert(std::is_same_v, - // "Not supported B GEMM layout!"); + static_assert(std::is_same_v, + "Not supported A GEMM layout!"); + static_assert(std::is_same_v, + "Not supported B GEMM layout!"); static_assert(std::is_same_v, "Not supported C GEMM layout!"); @@ -703,7 +696,7 @@ struct GroupedConvolutionBackwardDataKernel const auto& b_tensor_view = [&]() { return make_tensor_view( b_ptr, - kargs.b_grid_descs_n_k[group_id]); // B: weight + kargs.b_grid_descs_k_n[group_id]); // B: weight }(); const auto& c_tensor_view = [&]() { @@ -742,8 +735,10 @@ struct GroupedConvolutionBackwardDataKernel const auto& b_pad_view = [&]() { const auto& b_tensor_view = views.at(I1); return pad_tensor_view(b_tensor_view, - make_tuple(number{}, - number{}), + make_tuple( + number{}, + number{} + ), sequence{}); }(); @@ -788,9 +783,11 @@ struct GroupedConvolutionBackwardDataKernel const auto& b_block_window = [&]() { return make_tile_window(b_pad_view, - make_tuple(number{}, - number{}), - {i_n, i_k}); + make_tuple( + number{}, + number{} + ), + {i_k, i_n}); }(); const auto ds_block_window = generate_tuple( diff --git a/include/ck_tile/ops/grouped_convolution/utils/grouped_convolution_utils.hpp b/include/ck_tile/ops/grouped_convolution/utils/grouped_convolution_utils.hpp index c745aee622..d3c43f2e67 100644 --- a/include/ck_tile/ops/grouped_convolution/utils/grouped_convolution_utils.hpp +++ b/include/ck_tile/ops/grouped_convolution/utils/grouped_convolution_utils.hpp @@ -81,21 +81,15 @@ struct GroupedConvTraits TileGemmTraits; using GroupedConvImplicitGemmTraitsBwdWeight = TileGemmTraits; static constexpr ck_tile::index_t VectorSizeA = VectorSizeA_; static constexpr ck_tile::index_t VectorSizeB = VectorSizeB_; diff --git a/include/ck_tile/ops/grouped_convolution/utils/transform_conv_bwd_data_to_gemm.hpp b/include/ck_tile/ops/grouped_convolution/utils/transform_conv_bwd_data_to_gemm.hpp index c68a0a1400..d238473207 100644 --- a/include/ck_tile/ops/grouped_convolution/utils/transform_conv_bwd_data_to_gemm.hpp +++ b/include/ck_tile/ops/grouped_convolution/utils/transform_conv_bwd_data_to_gemm.hpp @@ -502,7 +502,7 @@ struct TransformConvBwdDataToGemm // TODO Add support for NumGroupsToMerge > 1 return make_naive_tensor_descriptor(make_tuple(N_, Hi_, Wi_, C_), make_tuple(NStride, HiStride, WiStride, CStride), - number{}, + number{}, I1); } @@ -512,7 +512,7 @@ struct TransformConvBwdDataToGemm // GKYXC return make_naive_tensor_descriptor(make_tuple(K_, Y_, X_, C_), make_tuple(C_ * X_ * Y_, C_ * X_, C_, I1), - number{}, + number{}, I1); } @@ -547,7 +547,7 @@ struct TransformConvBwdDataToGemm return make_naive_tensor_descriptor( make_tuple(N_, Di_, Hi_, Wi_, C_), make_tuple(NStride, DiStride, HiStride, WiStride, CStride), - number{}, + number{}, I1); } @@ -558,7 +558,7 @@ struct TransformConvBwdDataToGemm return make_naive_tensor_descriptor( make_tuple(K_, Z_, Y_, X_, C_), make_tuple(C_ * X_ * Y_ * Z_, C_ * X_ * Y_, C_ * X_, C_, I1), - number{}, + number{}, I1); } // TODO: implement ck_tile::tensor_layout::convolution that describe packed/strided dimemsion as @@ -637,12 +637,12 @@ struct TransformConvBwdDataToGemm make_tuple(sequence<0>{}, sequence<1>{}, sequence<2>{}, sequence<3>{}), make_tuple(sequence<0>{}, sequence<1>{}, sequence<>{}, sequence<2>{})); - const auto wei_gemmn_gemmkraw_grid_desc = + const auto wei_gemmkraw_gemmn_grid_desc = transform_tensor_descriptor(wei_k_xdotslice_c_grid_desc, make_tuple(make_merge_transform(make_tuple(XDotSlice, K_)), make_pass_through_transform(C_)), make_tuple(sequence<1, 0>{}, sequence<2>{}), - make_tuple(sequence<1>{}, sequence<0>{})); + make_tuple(sequence<0>{}, sequence<1>{})); // c: input const auto in_n_wip_c_grid_desc = transform_tensor_descriptor( @@ -679,7 +679,7 @@ struct TransformConvBwdDataToGemm make_tuple(sequence<0>{}, sequence<1>{})); return make_tuple(out_gemmm_gemmkraw_grid_desc, - wei_gemmn_gemmkraw_grid_desc, + wei_gemmkraw_gemmn_grid_desc, in_gemmmraw_gemmnraw_grid_desc); } @@ -792,12 +792,12 @@ struct TransformConvBwdDataToGemm sequence<>{}, sequence<3>{})); - const auto wei_gemmn_gemmkraw_grid_desc = transform_tensor_descriptor( + const auto wei_gemmkraw_gemmn_grid_desc = transform_tensor_descriptor( wei_k_ydotslice_xdotslice_c_grid_desc, make_tuple(make_merge_transform(make_tuple(YDotSlice, XDotSlice, K_)), make_pass_through_transform(C_)), make_tuple(sequence<1, 2, 0>{}, sequence<3>{}), - make_tuple(sequence<1>{}, sequence<0>{})); + make_tuple(sequence<0>{}, sequence<1>{})); // c: input const auto in_n_hip_wip_c_grid_desc = transform_tensor_descriptor( @@ -849,7 +849,7 @@ struct TransformConvBwdDataToGemm make_tuple(sequence<0>{}, sequence<1>{})); return make_tuple(out_gemmm_gemmkraw_grid_desc, - wei_gemmn_gemmkraw_grid_desc, + wei_gemmkraw_gemmn_grid_desc, in_gemmmraw_gemmnraw_grid_desc); } @@ -994,12 +994,12 @@ struct TransformConvBwdDataToGemm sequence<>{}, sequence<4>{})); - const auto wei_gemmn_gemmkraw_grid_desc = transform_tensor_descriptor( + const auto wei_gemmkraw_gemmn_grid_desc = transform_tensor_descriptor( wei_k_ydotslice_xdotslice_c_grid_desc, make_tuple(make_merge_transform(make_tuple(ZDotSlice, YDotSlice, XDotSlice, K_)), make_pass_through_transform(C_)), make_tuple(sequence<1, 2, 3, 0>{}, sequence<4>{}), - make_tuple(sequence<1>{}, sequence<0>{})); + make_tuple(sequence<0>{}, sequence<1>{})); // c: input const auto in_n_hip_wip_c_grid_desc = transform_tensor_descriptor( @@ -1064,7 +1064,7 @@ struct TransformConvBwdDataToGemm make_tuple(sequence<0>{}, sequence<1>{})); return make_tuple(out_gemmm_gemmkraw_grid_desc, - wei_gemmn_gemmkraw_grid_desc, + wei_gemmkraw_gemmn_grid_desc, in_gemmmraw_gemmnraw_grid_desc); }