From e0bc9d36cf3bbe6021c9b4142eb7e278af02c6ba Mon Sep 17 00:00:00 2001 From: Jakub Piasecki Date: Fri, 23 May 2025 13:55:36 +0000 Subject: [PATCH] working version with out transpose to nchw --- ...d_preshuffle_multiple_abd_xdl_cshuffle.hpp | 196 +++++++++--------- .../grid/gridwise_gemm_pipeline_v1_nchw.hpp | 65 ++++-- .../transform_conv_fwd_to_gemm.hpp | 32 +-- ...ouped_conv_fwd_preshuffle_xdl_instance.hpp | 52 ++--- .../device_grouped_conv_fwd_xdl_instance.hpp | 4 +- 5 files changed, 196 insertions(+), 153 deletions(-) diff --git a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_preshuffle_multiple_abd_xdl_cshuffle.hpp b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_preshuffle_multiple_abd_xdl_cshuffle.hpp index 4fad2bc2e9..f629e42dee 100644 --- a/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_preshuffle_multiple_abd_xdl_cshuffle.hpp +++ b/include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_preshuffle_multiple_abd_xdl_cshuffle.hpp @@ -409,14 +409,14 @@ struct DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle template static auto MakeEGridDescriptor_M_N(const ConvToGemmFwdTransformer& conv_to_gemm_transformer) { - // namespace ctc = tensor_layout::convolution; - // using Layout = std::conditional_t< - // is_NGCHW_NGKHW(), - // ctc::NHWGK, - // std::conditional_t(), ctc::NDHWGK, ELay>>; + namespace ctc = tensor_layout::convolution; + using Layout = std::conditional_t< + is_NGCHW_NGKHW(), + ctc::NHWGK, + std::conditional_t(), ctc::NDHWGK, ELay>>; const auto out_gemmmraw_gemmnraw_desc = - conv_to_gemm_transformer.template MakeCDescriptor_M_N(); + conv_to_gemm_transformer.template MakeCDescriptor_M_N(); const auto out_gemmm_gemmn_desc = matrix_padder.PadCDescriptor_M_N(out_gemmmraw_gemmnraw_desc); @@ -548,12 +548,12 @@ struct DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle remove_cvref_t; using Block2TileMapElementwise = BlockToCTileMap_M00_N0_M01Adapt; - // using NGCHWTransposeDescType = - // remove_cvref_t({}, {}))>; - // using NHWGCTransposeDescType = - // remove_cvref_t({}, {}))>; + using NGCHWTransposeDescType = + remove_cvref_t({}, {}))>; + using NHWGCTransposeDescType = + remove_cvref_t({}, {}))>; // using GKCYXTransposeDescType = // remove_cvref_t; - // using GridwiseElementwiseOutputTranspose = - // GridwiseElementwise, - // Tuple, - // Tuple, - // Tuple, - // Block2TileMapElementwise, - // element_wise::PassThrough, - // ElementwiseBlocksize, - // NPerBlock, - // NPerBlock, - // NPerBlock / ClusterLengthNPerBlock, - // NPerBlock / ClusterLengthNPerBlock, - // Sequence<1, 0>, - // Sequence, - // Sequence, - // I0, - // I1>; + using GridwiseElementwiseOutputTranspose = + GridwiseElementwise, + Tuple, + Tuple, + Tuple, + Block2TileMapElementwise, + element_wise::PassThrough, + ElementwiseBlocksize, + NPerBlock, + NPerBlock, + NPerBlock / ClusterLengthNPerBlock, + NPerBlock / ClusterLengthNPerBlock, + Sequence<1, 0>, + Sequence, + Sequence, + I0, + I1>; // Argument struct Argument : public BaseArgument @@ -658,7 +658,8 @@ struct DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle ds_g_n_k_wos_lengths_{ds_g_n_k_wos_lengths}, ds_g_n_k_wos_strides_{ds_g_n_k_wos_strides}, e_g_n_k_wos_lengths_{e_g_n_k_wos_lengths}, - e_g_n_k_wos_strides_{e_g_n_k_wos_strides}, + e_g_n_k_wos_strides_{conv_ngchw_to_nhwgc_transformer.TransposeInOutStrides( + e_g_n_k_wos_lengths, e_g_n_k_wos_strides)}, conv_filter_strides_{conv_filter_strides}, conv_filter_dilations_{conv_filter_dilations}, input_left_pads_{input_left_pads}, @@ -837,38 +838,38 @@ struct DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle } } - // if constexpr(is_NGCHW_NGKHW() || - // is_NGCDHW_NGKDHW()) - // { - // // Use not modified base strides - // a_in_transpose_desc_ = - // conv_ngchw_to_nhwgc_transformer.template MakeNGCHWTransposeDesc( - // a_g_n_c_wis_lengths, a_g_n_c_wis_strides); - // a_out_transpose_desc_ = - // conv_ngchw_to_nhwgc_transformer.template MakeNHWGCTransposeDesc( - // a_g_n_c_wis_lengths, a_g_n_c_wis_strides); + if constexpr(is_NGCHW_NGKHW() || + is_NGCDHW_NGKDHW()) + { + // Use not modified base strides + // a_in_transpose_desc_ = + // conv_ngchw_to_nhwgc_transformer.template MakeNGCHWTransposeDesc( + // a_g_n_c_wis_lengths, a_g_n_c_wis_strides); + // a_out_transpose_desc_ = + // conv_ngchw_to_nhwgc_transformer.template MakeNHWGCTransposeDesc( + // a_g_n_c_wis_lengths, a_g_n_c_wis_strides); - // b_in_transpose_desc_ = - // conv_ngchw_to_nhwgc_transformer.template MakeGKCYXTransposeDesc( - // b_g_k_c_xs_lengths, b_g_k_c_xs_strides); - // b_out_transpose_desc_ = - // conv_ngchw_to_nhwgc_transformer.template MakeGKYXCTransposeDesc( - // b_g_k_c_xs_lengths, b_g_k_c_xs_strides); + // b_in_transpose_desc_ = + // conv_ngchw_to_nhwgc_transformer.template MakeGKCYXTransposeDesc( + // b_g_k_c_xs_lengths, b_g_k_c_xs_strides); + // b_out_transpose_desc_ = + // conv_ngchw_to_nhwgc_transformer.template MakeGKYXCTransposeDesc( + // b_g_k_c_xs_lengths, b_g_k_c_xs_strides); - // e_in_transpose_desc_ = - // conv_ngchw_to_nhwgc_transformer.template MakeNHWGCTransposeDesc( - // e_g_n_k_wos_lengths, e_g_n_k_wos_strides); - // e_out_transpose_desc_ = - // conv_ngchw_to_nhwgc_transformer.template MakeNGCHWTransposeDesc( - // e_g_n_k_wos_lengths, e_g_n_k_wos_strides); + e_in_transpose_desc_ = + conv_ngchw_to_nhwgc_transformer.template MakeNHWGCTransposeDesc( + e_g_n_k_wos_lengths, e_g_n_k_wos_strides); + e_out_transpose_desc_ = + conv_ngchw_to_nhwgc_transformer.template MakeNGCHWTransposeDesc( + e_g_n_k_wos_lengths, e_g_n_k_wos_strides); - // elementwise_block_2_ctile_map_transpose_a_ = Block2TileMapElementwise{ - // a_in_transpose_desc_.GetLength(I0), a_in_transpose_desc_.GetLength(I1)}; - // elementwise_block_2_ctile_map_transpose_b_ = Block2TileMapElementwise{ - // b_in_transpose_desc_.GetLength(I0), b_in_transpose_desc_.GetLength(I1)}; - // elementwise_block_2_ctile_map_transpose_e_ = Block2TileMapElementwise{ - // e_in_transpose_desc_.GetLength(I0), e_in_transpose_desc_.GetLength(I1)}; - // } + // elementwise_block_2_ctile_map_transpose_a_ = Block2TileMapElementwise{ + // a_in_transpose_desc_.GetLength(I0), a_in_transpose_desc_.GetLength(I1)}; + // elementwise_block_2_ctile_map_transpose_b_ = Block2TileMapElementwise{ + // b_in_transpose_desc_.GetLength(I0), b_in_transpose_desc_.GetLength(I1)}; + elementwise_block_2_ctile_map_transpose_e_ = Block2TileMapElementwise{ + e_in_transpose_desc_.GetLength(I0), e_in_transpose_desc_.GetLength(I1)}; + } } std::size_t GetWorkspaceATensorSizeBytes() const @@ -920,7 +921,8 @@ struct DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle std::size_t GetWorkspaceSizeBytes() const { - return 0; + return GetWorkspaceATensorSizeBytes() + GetWorkspaceBTensorSizeBytes() + + GetWorkspaceETensorSizeBytes(); } void Print() const @@ -980,8 +982,8 @@ struct DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle Block2TileMapElementwise elementwise_block_2_ctile_map_transpose_a_, elementwise_block_2_ctile_map_transpose_b_, elementwise_block_2_ctile_map_transpose_e_; - // NGCHWTransposeDescType a_in_transpose_desc_, e_out_transpose_desc_; - // NHWGCTransposeDescType a_out_transpose_desc_, e_in_transpose_desc_; + NGCHWTransposeDescType a_in_transpose_desc_, e_out_transpose_desc_; + NHWGCTransposeDescType a_out_transpose_desc_, e_in_transpose_desc_; // GKCYXTransposeDescType b_in_transpose_desc_; // GKYXCTransposeDescType b_out_transpose_desc_; @@ -1086,20 +1088,20 @@ struct DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle // p_a_grid = type_convert(arg.p_workspace_); // p_b_grid = type_convert(arg.p_workspace_) + // arg.GetWorkspaceATensorSizeBytes() / sizeof(BDataType); - // p_e_grid = type_convert(arg.p_workspace_) + - // (arg.GetWorkspaceATensorSizeBytes() + - // arg.GetWorkspaceBTensorSizeBytes()) / - // sizeof(EDataType); - } - else if constexpr(is_NGCHW_GKYXC_NGKHW() || - is_NGCDHW_GKZYXC_NGKDHW()) - { - p_a_grid = type_convert(arg.p_workspace_); p_e_grid = type_convert(arg.p_workspace_) + (arg.GetWorkspaceATensorSizeBytes() + arg.GetWorkspaceBTensorSizeBytes()) / sizeof(EDataType); } + else if constexpr(is_NGCHW_GKYXC_NGKHW() || + is_NGCDHW_GKZYXC_NGKDHW()) + { + // p_a_grid = type_convert(arg.p_workspace_); + // p_e_grid = type_convert(arg.p_workspace_) + + // (arg.GetWorkspaceATensorSizeBytes() + + // arg.GetWorkspaceBTensorSizeBytes()) / + // sizeof(EDataType); + } const auto kernel = kernel_grouped_conv_fwd_preshuffle_multiple_abd_xdl_cshuffle< GridwiseGemm, @@ -1218,36 +1220,36 @@ struct DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle if constexpr(is_NGCHW_NGKHW() || is_NGCDHW_NGKDHW()) { - // const index_t grid_size = - // arg.elementwise_block_2_ctile_map_transpose_e_.CalculateGridSize( - // arg.e_in_transpose_desc_); + const index_t grid_size = + arg.elementwise_block_2_ctile_map_transpose_e_.CalculateGridSize( + arg.e_in_transpose_desc_); - // const EDataType* p_e_in_grid = - // type_convert(arg.p_workspace_) + - // (arg.GetWorkspaceATensorSizeBytes() + arg.GetWorkspaceBTensorSizeBytes()) / - // sizeof(EDataType); + const EDataType* p_e_in_grid = + type_convert(arg.p_workspace_) + + (arg.GetWorkspaceATensorSizeBytes() + arg.GetWorkspaceBTensorSizeBytes()) / + sizeof(EDataType); - // EDataType* p_e_out_grid = arg.p_e_grid_; + EDataType* p_e_out_grid = arg.p_e_grid_; - // auto kernel_transpose = kernel_elementwise, - // ck::Tuple, - // ck::Tuple, - // ck::Tuple, - // Block2TileMapElementwise, - // element_wise::PassThrough>; + [[maybe_unused]] auto kernel_transpose = kernel_elementwise, + ck::Tuple, + ck::Tuple, + ck::Tuple, + Block2TileMapElementwise, + element_wise::PassThrough>; - // avg_time += launch_and_time_kernel(stream_config, - // kernel_transpose, - // dim3(grid_size), - // dim3(ElementwiseBlocksize), - // 0, - // make_tuple(arg.e_in_transpose_desc_), - // make_tuple(arg.e_out_transpose_desc_), - // make_tuple(p_e_in_grid), - // make_tuple(p_e_out_grid), - // arg.elementwise_block_2_ctile_map_transpose_e_, - // element_wise::PassThrough{}); + avg_time += launch_and_time_kernel(stream_config, + kernel_transpose, + dim3(grid_size), + dim3(ElementwiseBlocksize), + 0, + make_tuple(arg.e_in_transpose_desc_), + make_tuple(arg.e_out_transpose_desc_), + make_tuple(p_e_in_grid), + make_tuple(p_e_out_grid), + arg.elementwise_block_2_ctile_map_transpose_e_, + element_wise::PassThrough{}); } return avg_time; diff --git a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_pipeline_v1_nchw.hpp b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_pipeline_v1_nchw.hpp index 36fb73432b..1cf9314e6e 100644 --- a/include/ck/tensor_operation/gpu/grid/gridwise_gemm_pipeline_v1_nchw.hpp +++ b/include/ck/tensor_operation/gpu/grid/gridwise_gemm_pipeline_v1_nchw.hpp @@ -36,7 +36,7 @@ struct GridwiseGemmPipeline_v1_nchw<1, true, true> typename AGridBuffer, typename ABlock1Buffer, typename ABlock2Buffer, - typename ABlockTransferStep, + typename ABlockTransferSteps, typename BGridDesc, typename BBlockDesc, typename BBlockTransfer, @@ -54,7 +54,7 @@ struct GridwiseGemmPipeline_v1_nchw<1, true, true> [[maybe_unused]] const AGridBuffer& a_grid_buf, [[maybe_unused]] ABlock1Buffer& a_block1_buf, [[maybe_unused]] ABlock2Buffer& a_block2_buf, - [[maybe_unused]] const ABlockTransferStep& a_block_copy_step, + [[maybe_unused]] const ABlockTransferSteps& a_block_copy_step, [[maybe_unused]] const BGridDesc& b_grid_desc, [[maybe_unused]] const BBlockDesc& b_block_desc, [[maybe_unused]] BBlockTransfer& b_blockwise_copy, @@ -85,7 +85,8 @@ struct GridwiseGemmPipeline_v1_nchw<1, true, true> //block_sync_lds(); // if(threadIdx.x == 0) { - // for(int c=0; c<4; ++c) { + // printf("Next iter\n"); + // for(int c=0; c<5; ++c) { // for(int h=0; h<2;++h) { // for(int w=0; w<64;++w) { // printf("Ag[%d][%d][%d]=%f\n", c, h, w, static_cast(a_grid_buf[c * 64 * 2 + h * 64 + w])); @@ -93,9 +94,9 @@ struct GridwiseGemmPipeline_v1_nchw<1, true, true> // } // } - // // A[3][2][33]=0.000000 + // // // A[3][2][33]=0.000000 - // for(int c=0; c<8; ++c) { + // for(int c=0; c<4; ++c) { // for(int h=0; h<4;++h) { // for(int w=0; w<64;++w) { // printf("A[%d][%d][%d]=%f\n", c, h, w, static_cast(a_block1_buf[c * 64 * 4 + h * 64 + w])); @@ -104,14 +105,14 @@ struct GridwiseGemmPipeline_v1_nchw<1, true, true> // } // for(int m=0; m<128; ++m) { - // for(int k0=0; k0<8; ++k0) { + // for(int k0=0; k0<4; ++k0) { // for(int k1=0;k1<8;++k1) { // printf("A[%d][%d]=%f\n", m, k0*8+k1, static_cast(a_block2_buf[k0*8*128 + m*8 + k1])); - // int k = k0*8+k1; - // int x = k%3; - // int y = (k/3)%3; - // int c = k/(9); + // // int k = k0*8+k1; + // // int x = k%3; + // // int y = (k/3)%3; + // // int c = k/(9); // // A[m][k] = A[c, y+m/64,x+m%64] z oryginalnego obrazka @@ -142,6 +143,7 @@ struct GridwiseGemmPipeline_v1_nchw<1, true, true> do { // ask bartek how to structure this pipeline a_blockwise_global_to_lds1_copy.RunRead(a_grid_desc, a_grid_buf); // A Global -> VGPR + block_sync_lds(); b_blockwise_copy.RunRead(b_grid_desc, b_grid_buf); // B Global -> VGPR // @@ -153,11 +155,50 @@ struct GridwiseGemmPipeline_v1_nchw<1, true, true> a_blockwise_global_to_lds1_copy.MoveSrcSliceWindow(a_grid_desc, a_block_copy_step); b_blockwise_copy.MoveSrcSliceWindow(b_grid_desc, b_block_copy_step); - //block_sync_lds(); a_blockwise_lds1_to_lds2_copy.RunWrite(a_block2_desc_ak0_m_k1, a_block2_buf); b_blockwise_copy.RunWrite(b_block_desc, b_block_buf); - //block_sync_lds(); + // block_sync_lds(); // delete later + + // if(threadIdx.x == 0) { + // printf("Next iter\n"); + + // // // A[3][2][33]=0.000000 + + // for(int c=0; c<4; ++c) { + // for(int h=0; h<4;++h) { + // for(int w=0; w<64;++w) { + // printf("A[%d][%d][%d]=%f\n", c, h, w, static_cast(a_block1_buf[c * 64 * 4 + h * 64 + w])); + // } + // } + // } + + // for(int m=0; m<128; ++m) { + // for(int k0=0; k0<4; ++k0) { + // for(int k1=0;k1<8;++k1) { + // printf("A[%d][%d]=%f\n", m, k0*8+k1, static_cast(a_block2_buf[k0*8*128 + m*8 + k1])); + + // // int k = k0*8+k1; + // // int x = k%3; + // // int y = (k/3)%3; + // // int c = k/(9); + + // // A[m][k] = A[c, y+m/64,x+m%64] z oryginalnego obrazka + + // // if(x+m%64 > 0 && x+m%64 < 127) { + // // auto lds1 = static_cast(a_block1_buf[c * 64 * 4 + (m/64 + y) * 64 + m%64 + x - 1]); + // // auto lds2 = static_cast(a_block2_buf[k0*8*128 + m*8 + k1]); + // // printf("lds2[%d][%d]:%f lds1[%d][%d][%d]:%f %s\n", m, k0*8+k1, lds2, c, y + m/64, m%64 + x - 1, lds1, (lds1 > lds2)? "diff" : ""); + + // // } + + // // if(lds1 > lds2) { + // // printf("diff lds1[%d][%d][%d]:%f lds2[%d][%d]:%f\n", c, y + m/64, m%64 + x, lds1, m, k0*8+k1 - 1, lds2); + // // } + // } + // } + // } + // } ++i; } while(i < (num_loop - 1)); diff --git a/include/ck/tensor_operation/operator_transform/transform_conv_fwd_to_gemm.hpp b/include/ck/tensor_operation/operator_transform/transform_conv_fwd_to_gemm.hpp index eb15d8eeca..56fefeaa6b 100644 --- a/include/ck/tensor_operation/operator_transform/transform_conv_fwd_to_gemm.hpp +++ b/include/ck/tensor_operation/operator_transform/transform_conv_fwd_to_gemm.hpp @@ -1478,23 +1478,23 @@ struct TransformConvFwdToGemm } - template ), - bool>::type = false> - __host__ __device__ auto MakeCDescriptor_M_N() const - { - auto desc = make_naive_tensor_descriptor(make_tuple(N_, K_, Ho_, Wo_), - make_tuple(K_ * Wo_ * Ho_ , Ho_* Wo_ , Wo_, I1)); + // template ), + // bool>::type = false> + // __host__ __device__ auto MakeCDescriptor_M_N() const + // { + // auto desc = make_naive_tensor_descriptor(make_tuple(N_, K_, Ho_, Wo_), + // make_tuple(K_ * Wo_ * Ho_ , Ho_* Wo_ , Wo_, I1)); - return transform_tensor_descriptor( - desc, - make_tuple(make_merge_transform(make_tuple(N_, Ho_, Wo_)), - make_pass_through_transform(K_)), - make_tuple(Sequence<0, 2, 3>{}, Sequence<1>{}), - make_tuple(Sequence<0>{}, Sequence<1>{})); - } + // return transform_tensor_descriptor( + // desc, + // make_tuple(make_merge_transform(make_tuple(N_, Ho_, Wo_)), + // make_pass_through_transform(K_)), + // make_tuple(Sequence<0, 2, 3>{}, Sequence<1>{}), + // make_tuple(Sequence<0>{}, Sequence<1>{})); + // } template , S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, +// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, +// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, +// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, +// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, +// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, +// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, -// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, -// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 32, 1, 8>, 1>, - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, -// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 16, 1, 8>, 8>, - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 16, 1, 8>, 8>, - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, // best - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 64, 1, 4>, 8>, // best +// // DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, +// // DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, +// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, +// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, +// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 32, 1, 8>, 1>, +// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, +// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, +// // DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, +// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, +// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, +// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, +// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 16, 1, 8>, 8>, +// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, +// DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 16, 1, 8>, 8>, + DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 32, 1, 8>, 8> // best + // DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 64, 1, 4>, 8>, // best - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, // best - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 64, 1, 4>, 8>, // best + // DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, // best + // DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 64, 1, 4>, 8>, // best - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, - DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 4>, 8> + // DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, + // DeviceGroupedConvFwdPreshuffleMultipleABD_Xdl_CShuffle, S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 4>, 8> // clang-format on diff --git a/library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_fwd/device_grouped_conv_fwd_xdl_instance.hpp b/library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_fwd/device_grouped_conv_fwd_xdl_instance.hpp index a3b9228508..7402a3dce4 100644 --- a/library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_fwd/device_grouped_conv_fwd_xdl_instance.hpp +++ b/library/include/ck/library/tensor_operation_instance/gpu/grouped_conv_fwd/device_grouped_conv_fwd_xdl_instance.hpp @@ -127,7 +127,7 @@ using device_grouped_conv_fwd_xdl_f16_instances = std::tuple< //########################################| | | | | | | | | | | | Operation| Operation| Operation| | | Stage| | | | | | | | | Wave| Wave| Lengths_K0_M_K1| ArrangeOrder| | | PerVector| PerVector_K1| | Lengths_K0_N_K1| ArrangeOrder| | | PerVector| PerVector_K1| | PerShuffle| PerShuffle| _NBlock_NWaveNPerXdl| _NWaveNPerXdl| //########################################| | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | | // // generic instance - DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 1, 0, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 1, 0, 1, 1, S<1, 16, 1, 4>, 1> + // DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 1, 0, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 1, 0, 1, 1, S<1, 16, 1, 4>, 1> // DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, // DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, // DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 1, 8, 1, 1, 1, S<1, 16, 1, 4>, 1>, @@ -144,7 +144,7 @@ using device_grouped_conv_fwd_xdl_f16_instances = std::tuple< // DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 4>, 8>, // DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 16, 1, 8>, 8>, // DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 16, 1, 4>, 8>, - // DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, + // DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8> // DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 64, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 8>, 8>, // DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 32, 1, 4>, 8>, // DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 32, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 16, 1, 8>, 8>,