diff --git a/example/20_grouped_conv_bwd_weight/common.hpp b/example/20_grouped_conv_bwd_weight/common.hpp index 2386e35b48..a7cdacd0a3 100644 --- a/example/20_grouped_conv_bwd_weight/common.hpp +++ b/example/20_grouped_conv_bwd_weight/common.hpp @@ -41,15 +41,13 @@ struct CommonLayoutSetting using WeightLayout = WeightLay; using OutputLayout = OutputLay; }; -using ALayout = ck::tensor_layout::convolution::NHWGC; -using BLayout = ck::tensor_layout::convolution::GKYXC; -using ELayout = ck::tensor_layout::convolution::NHWGK; + namespace ctl = ck::tensor_layout::convolution; template struct CommonLayoutSettingSelector : CommonLayoutSetting>, ck::tuple_element_t>, ck::tuple_element_t>> { }; diff --git a/example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_xdl_fp16.cpp b/example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_xdl_fp16.cpp index b24125ed3a..d6f7ad44d2 100644 --- a/example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_xdl_fp16.cpp +++ b/example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_xdl_fp16.cpp @@ -34,7 +34,7 @@ static constexpr index_t SizeOfType = 2; static constexpr index_t Tile_Align_W = Tile_W + 1; static constexpr index_t ShareMemSize = Tile_H * Tile_Align_W * N_Pack * SizeOfType; static constexpr index_t ScratchSize = ShareMemSize / 64 / 4; -static constexpr index_t Num_Wave = 1; +static constexpr index_t Num_Wave = 4; #define MergeShareMem 1 template __device__ T warp_shuffle_up(const T& v_local, uint32_t lane_delta) @@ -605,9 +605,9 @@ struct DeviceGroupedConvBwdWeightNaive } } -using ALayout = ck::tensor_layout::convolution::NHWGC; +using ALayout = ck::tensor_layout::convolution::GNHWC; using BLayout = ck::tensor_layout::convolution::GKYXC; -using ELayout = ck::tensor_layout::convolution::NHWGK; +using ELayout = ck::tensor_layout::convolution::GNHWK; template using DeviceConvBwdWeightInstance =