ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-14 19:18:35 +00:00
Code Issues Packages Projects Releases Wiki Activity

change layout to gnhwc

Browse Source
This commit is contained in:
Lin, Qun
2025-05-29 19:28:38 -05:00
parent 321dc5dd39
commit 26b8a1d972
2 changed files with 6 additions and 8 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
example/20_grouped_conv_bwd_weight/common.hpp
View File

@@ -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 <ck::index_t NDimSpatial>
struct CommonLayoutSettingSelector
: CommonLayoutSetting<ck::tuple_element_t<NDimSpatial - 1,
ck::Tuple<ck::tensor_layout::convolution::GNWC,
ck::tensor_layout::convolution::NHWGC,
ck::tensor_layout::convolution::GNHWC,
ck::tensor_layout::convolution::GNDHWC>>,
ck::tuple_element_t<NDimSpatial - 1,
ck::Tuple<ck::tensor_layout::convolution::GKXC,
@@ -57,7 +55,7 @@ struct CommonLayoutSettingSelector
ck::tensor_layout::convolution::GKZYXC>>,
ck::tuple_element_t<NDimSpatial - 1,
ck::Tuple<ck::tensor_layout::convolution::GNWK,
ck::tensor_layout::convolution::NHWGK,
ck::tensor_layout::convolution::GNHWK,
ck::tensor_layout::convolution::GNDHWK>>>
{
};

6
example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_xdl_fp16.cpp
View File

@@ -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 <typename T>
__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 <ck::index_t NDimSpatial>
using DeviceConvBwdWeightInstance =