diff --git a/library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_bf16_instances.hpp b/library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_bf16_instances.hpp index 2bb3cd9cca..93fa3a1708 100644 --- a/library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_bf16_instances.hpp +++ b/library/include/ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_weight_bf16_instances.hpp @@ -22,14 +22,143 @@ using tile_grouped_conv_bwd_weight_bf16_instances = std::tuple< //#####################################| Dim| | | | Type| Type| Type| Elementwise| Elementwise| Elementwise| per| | | | | | | tile| tile| tile| size| size| size| //#####################################| Spatial| | | | | | | Operation| Operation| Operation| CU| | | | | | | size| size| size| A| B| C| //#####################################| | | | | | | | | | | - GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, GroupedConvolutionBackwardWeightInvoker, GroupedConvolutionBackwardWeightInvoker, GroupedConvolutionBackwardWeightInvoker, - GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, GroupedConvolutionBackwardWeightInvoker, GroupedConvolutionBackwardWeightInvoker, - GroupedConvolutionBackwardWeightInvoker + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + + //#####################################| Num| InLayout| WeiLayout| OutLayout| InData| WeiData| OutData| In| Wei| Out| K-block| M-tile| N-tile | K-tile | M-warp| N-warp| K-warp| M-warp| N-warp| K-warp| Vector| Vector| Vector| + //#####################################| Dim| | | | Type| Type| Type| Elementwise| Elementwise| Elementwise| per| | | | | | | tile| tile| tile| size| size| size| + //#####################################| Spatial| | | | | | | Operation| Operation| Operation| CU| | | | | | | size| size| size| A| B| C| + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker, + GroupedConvolutionBackwardWeightInvoker + // // clang-format on // clang-format on >; diff --git a/profiler/ck_tile/src/tile_profile_grouped_conv_bwd_weight.cpp b/profiler/ck_tile/src/tile_profile_grouped_conv_bwd_weight.cpp index faccc8fcce..d29df87ea6 100644 --- a/profiler/ck_tile/src/tile_profile_grouped_conv_bwd_weight.cpp +++ b/profiler/ck_tile/src/tile_profile_grouped_conv_bwd_weight.cpp @@ -119,16 +119,12 @@ int tile_profile_grouped_conv_bwd_weight(int argc, char* argv[]) using TF32 = ck::tf32_t; #endif - //using namespace ck::tensor_layout::convolution; - //using NWGC = ck_tile::tensor_layout::convolution::NWGC; using NHWGC = ck_tile::tensor_layout::convolution::NHWGC; using NDHWGC = ck_tile::tensor_layout::convolution::NDHWGC; - //using GKXC = ck_tile::tensor_layout::convolution::GKXC; using GKYXC = ck_tile::tensor_layout::convolution::GKYXC; using GKZYXC = ck_tile::tensor_layout::convolution::GKZYXC; - //using NWGK = ck_tile::tensor_layout::convolution::NWGK; using NHWGK = ck_tile::tensor_layout::convolution::NHWGK; using NDHWGK = ck_tile::tensor_layout::convolution::NDHWGK;