ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-15 03:30:11 +00:00
Code Issues Packages Projects Releases Wiki Activity

add new instance for tuning

Browse Source
This commit is contained in:
Qun Lin
2025-06-07 14:49:01 +08:00
parent 9b5f86cec6
commit 397c9c9865
4 changed files with 187 additions and 88 deletions
Download Patch File Download Diff File Expand all files Collapse all files

22
example/20_grouped_conv_bwd_weight/device_grouped_conv_bwd_weight_dl_v4.hpp
View File

@@ -1,20 +1,11 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#include "common.hpp"
#include "ck/utility/common_header.hpp"
// #include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_xdl_cshuffle.hpp"
#include "ck/utility/blkgemmpipe_scheduler.hpp"
#include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp"
using InDataType = F16;
using WeiDataType = F16;
using OutDataType = F16;
using AccDataType = F32;
using InElementOp = PassThrough;
using WeiElementOp = PassThrough;
using OutElementOp = PassThrough;
#include "ck/tensor_operation/gpu/device/device_grouped_conv_bwd_weight.hpp"
#include "ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp"
#include "ck/host_utility/kernel_launch.hpp"
#define ENABLE_PIPELINE_V2 1
@@ -104,6 +95,7 @@ template <index_t BlockSize,
typename InDataType,
typename WeiDataType,
typename OutDataType,
typename AccDataType,
typename BlockTileSize, // input, without padding
index_t FilterSize, //
typename FilterParam, // tuple<dilation, stride, padding>
@@ -186,7 +178,7 @@ struct GridwiseGroupedConv2DBwdWeightDlV4
static constexpr index_t SubTileIn_Stride = SubTileIn_Max_W;
static constexpr index_t SubTileOut_Stride = SubTileOut_W;
static constexpr index_t SubTileIn_Pack_W = GetAlignedPackW<SubTileIn_Max_W, InScalarPerVector>();
static constexpr index_t SubTileIn_Pack_W = (WSplit == 1) ? GetAlignedPackW<Tile_W, InScalarPerVector>() : GetAlignedPackW<SubTileIn_Max_W, InScalarPerVector>();
static constexpr index_t TileIn_Pack_Group = WaveSize / SubTileIn_Pack_W;
static constexpr index_t TileIn_Pack_H = math::integer_divide_ceil(Tile_H, TileIn_Pack_Group);
static constexpr index_t TileIn_Align_H =
@@ -1028,6 +1020,7 @@ template <index_t NDimSpatial,
typename InDataType,
typename WeiDataType,
typename OutDataType,
typename AccDataType,
typename BlockTileSize, // input, without include pading
index_t FilterSize, // seqence<x, y, [z]>
typename FilterParam, // tuple<dilation, stride, padding>
@@ -1072,6 +1065,7 @@ struct DeviceGroupedConvBwdWeightDlV4 : public DeviceGroupedConvBwdWeight<NDimSp
InDataType,
WeiDataType,
OutDataType,
AccDataType,
BlockTileSize,
FilterSize,
FilterParam,

102
example/20_grouped_conv_bwd_weight/grouped_conv_bwd_weight_dl_v4_fp16.cpp
View File

@@ -1,10 +1,8 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
#include "common.hpp"
#include "ck/utility/blkgemmpipe_scheduler.hpp"
#include "device_grouped_conv_bwd_weight_dl_v4.hpp"
#include "common.hpp"
#define ENABLE_CONV_FACTORY 1
@@ -31,6 +29,7 @@ using DeviceConvBwdWeightInstance =
InDataType,
WeiDataType,
OutDataType,
AccDataType,
S<28, 28>,
5,
ck::Tuple<S<1,1>, S<1,1>, S<2,2>>,
@@ -44,23 +43,86 @@ using DeviceConvBwdWeightInstance =
2, // DstScalarPerVector
false>;
using DeviceConvBwdWeightFactory = std::tuple<
// NDimSpatial BlockSize InLayout WeiLayout OutLayout InDataType WeiDataType OutDataType BlockTileSize FilterSize FilterParam(dilation, stride, pad) NBatch NumWavePerTile InScalarPerVector OutScalarPerVector DstScalarPerVector RequirePadding
// ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<28, 28>, 5, ck::Tuple<S<1,1>, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 2, 2, 2, false>
// ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<56, 56>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 1, 1, 2, false, 1>
// , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<14, 14>, 5, ck::Tuple<S<1,1>, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 8, 1, 2, 2, 8, false>
// , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<7, 7>, 5, ck::Tuple<S<1,1>, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 16, 1, 1, 1, 8, false>
// , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 128, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<56, 56>, 5, ck::Tuple<S<1,1>, S<2,2>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 2, 2, 2, 2, 2, false>
// , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<14, 14>, 5, ck::Tuple<S<1,1>, S<2,2>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 8, 1, 2, 1, 8, false>
ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<112, 112>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 1, 4, 8, 8, 1, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<112, 112>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 4, 4, 4, 2, false, 2>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<112, 112>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 4, 2, 2, 2, false, 4>
// , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 128, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<56, 56>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 1, 2, 4, 4, 1, false>
// , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<28, 28>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 4, 4, 2, false>
// , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<14, 14>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 4, 1, 2, 2, 4, false>
// , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<7, 7>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 16, 1, 1, 1, 8, false>
// , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<112, 112>, 3, ck::Tuple<S<1,1>, S<2,2>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 1, 4, 8, 4, 1, false>
// , ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, S<28, 28>, 3, ck::Tuple<S<1,1>, S<2,2>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 4, 2, 2, false>
using DeviceConvBwdWeightFactory = std::tuple<
// NDimSpatial BlockSize InLayout WeiLayout OutLayout InDataType WeiDataType OutDataType AccDatType BlockTileSize FilterSize FilterParam(dilation, stride, pad) NBatch NumWavePerTile InScalarPerVector OutScalarPerVector DstScalarPerVector RequirePadding
ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<28, 28>, 5, ck::Tuple<S<1,1>, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 2, 2, 2, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<14, 14>, 5, ck::Tuple<S<1,1>, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 8, 1, 2, 2, 8, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<7, 7>, 5, ck::Tuple<S<1,1>, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 16, 1, 1, 1, 8, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 128, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<56, 56>, 5, ck::Tuple<S<1,1>, S<2,2>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 2, 2, 2, 2, 2, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<14, 14>, 5, ck::Tuple<S<1,1>, S<2,2>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 8, 1, 2, 1, 8, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<112, 112>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 1, 4, 8, 8, 1, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 128, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<56, 56>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 1, 2, 4, 4, 1, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<28, 28>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 4, 4, 2, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<14, 14>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 4, 1, 2, 2, 4, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<7, 7>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 16, 1, 1, 1, 8, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<112, 112>, 3, ck::Tuple<S<1,1>, S<2,2>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 1, 4, 8, 4, 1, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<28, 28>, 3, ck::Tuple<S<1,1>, S<2,2>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 4, 2, 2, false>
// 28 x 5 x 1
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<28, 28>, 5, ck::Tuple<S<1,1>, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 2, 2, 2, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<28, 28>, 5, ck::Tuple<S<1,1>, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 4, 4, 2, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<28, 28>, 5, ck::Tuple<S<1,1>, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 4, 4, 2, false>
// 14 x 5 x 1
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<14, 14>, 5, ck::Tuple<S<1,1>, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 4, 1, 2, 2, 4, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<14, 14>, 5, ck::Tuple<S<1,1>, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 2, 2, 2, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<14, 14>, 5, ck::Tuple<S<1,1>, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 4, 1, 2, 2, 4, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<14, 14>, 5, ck::Tuple<S<1,1>, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 2, 2, 2, false>
// 7 x 5 x 1
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<7, 7>, 5, ck::Tuple<S<1,1>, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 8, 1, 1, 1, 8, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<7, 7>, 5, ck::Tuple<S<1,1>, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 4, 1, 1, 1, 4, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<7, 7>, 5, ck::Tuple<S<1,1>, S<1,1>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 1, 1, 2, false>
// 56 x 5 x 2
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 128, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<56, 56>, 5, ck::Tuple<S<1,1>, S<2,2>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 2, 2, 4, 2, 2, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 128, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<56, 56>, 5, ck::Tuple<S<1,1>, S<2,2>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 2, 2, 4, 2, 2, false, 2>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 128, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<56, 56>, 5, ck::Tuple<S<1,1>, S<2,2>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 2, 2, 2, 2, 2, false, 2>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<56, 56>, 5, ck::Tuple<S<1,1>, S<2,2>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 2, 2, 2, false, 2>
// 14 x 5 x 2
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<14, 14>, 5, ck::Tuple<S<1,1>, S<2,2>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 4, 1, 2, 1, 4, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<14, 14>, 5, ck::Tuple<S<1,1>, S<2,2>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 8, 1, 2, 1, 8, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<14, 14>, 5, ck::Tuple<S<1,1>, S<2,2>, S<2,2>>, InElementOp, WeiElementOp, OutElementOp, 4, 1, 2, 1, 4, false>
// 112 x 3 x 1
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<112, 112>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 4, 4, 4, 2, false, 2>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<112, 112>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 4, 2, 2, 2, false, 4>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<112, 112>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 4, 2, 2, 2, false, 2>
// 56 x 3 x 1
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<56, 56>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 4, 4, 4, 2, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 128, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<56, 56>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 2, 2, 2, 2, false, 2>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 128, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<56, 56>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 2, 2, 2, 2, false, 4>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<56, 56>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 4, 2, 2, 2, false, 2>
// 28 x 3 x 1
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<28, 28>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 2, 2, 2, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 128, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<28, 28>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 4, 4, 2, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<28, 28>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 4, 4, 2, false>
// 14 x 3 x 1
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<14, 14>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 8, 1, 2, 2, 8, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<14, 14>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 2, 2, 2, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 128, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<14, 14>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 4, 1, 2, 2, 4, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<14, 14>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 4, 1, 2, 2, 4, false>
// 7 x 3 x 1
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<7, 7>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 8, 1, 1, 1, 8, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<7, 7>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 4, 1, 1, 1, 4, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<7, 7>, 3, ck::Tuple<S<1,1>, S<1,1>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 1, 1, 2, false>
// 112 x 3 x 2
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<112, 112>, 3, ck::Tuple<S<1,1>, S<2,2>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 4, 4, 2, 2, false, 2>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<112, 112>, 3, ck::Tuple<S<1,1>, S<2,2>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 4, 2, 2, 2, false, 2>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<112, 112>, 3, ck::Tuple<S<1,1>, S<2,2>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 4, 2, 2, 2, false, 4>
// 28 x 3 x 3
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<28, 28>, 3, ck::Tuple<S<1,1>, S<2,2>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 4, 2, 2, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 128, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<28, 28>, 3, ck::Tuple<S<1,1>, S<2,2>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 2, 1, 4, 2, 2, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 64, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<28, 28>, 3, ck::Tuple<S<1,1>, S<2,2>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 4, 1, 4, 2, 4, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 128, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<28, 28>, 3, ck::Tuple<S<1,1>, S<2,2>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 4, 1, 4, 2, 4, false>
, ck::tensor_operation::device::DeviceGroupedConvBwdWeightDlV4<2, 256, ALayout, BLayout, ELayout, InDataType, WeiDataType, OutDataType, AccDataType, S<28, 28>, 3, ck::Tuple<S<1,1>, S<2,2>, S<1,1>>, InElementOp, WeiElementOp, OutElementOp, 4, 1, 4, 2, 4, false>
>;
template <ck::index_t NDimSpatial>

121
example/20_grouped_conv_bwd_weight/run_grouped_conv_bwd_weight_example.inc
View File

@@ -98,58 +98,101 @@ bool run_grouped_conv_bwd_weight(const ExecutionConfig& config,
}
#if ENABLE_CONV_FACTORY
std::size_t best_tflops = 0;
std::size_t best_gb_per_sec = 0;
std::size_t best_avg_time = 0;
std::string best_kernel = "";
ck::index_t split_k_array[] = {1, 2, 4, 8, 16, 32};
ck::index_t split_k_count = 6;
ck::index_t instance_idx = 0;
ck::index_t best_split_k = 0;
if (split_k != -1)
{
split_k_count = 1;
split_k_array[0] = split_k;
}
bool found_kernel= false;
ck::static_for<0, std::tuple_size_v<DeviceConvBwdWeightFactory>, 1>{}([&](auto i) -> void {
const auto device_conv_bwd_weight_instance = std::get<i>(DeviceConvBwdWeightFactory{});
using DeviceConvBwdWeightInstance = ck::remove_cvref_t<decltype(device_conv_bwd_weight_instance)>;
auto conv = DeviceConvBwdWeightInstance{};
auto invoker = conv.MakeInvoker();
auto argument = conv.MakeArgument(static_cast<InDataType*>(in_device_buf.GetDeviceBuffer()),
static_cast<WeiDataType*>(wei_device_buf.GetDeviceBuffer()),
static_cast<OutDataType*>(out_device_buf.GetDeviceBuffer()),
input_lengths,
input_strides,
filter_lengths,
weights_strides,
output_lengths,
output_strides,
conv_filter_strides,
conv_filter_dilations,
input_left_pads,
input_right_pads,
InElementOp{},
WeiElementOp{},
OutElementOp{},
split_k);
DeviceMem gemm_workspace_dev(conv.GetWorkSpaceSize(&argument));
conv.SetWorkSpacePointer(&argument, gemm_workspace_dev.GetDeviceBuffer());
if(conv.IsSupportedArgument(argument))
for (ck::index_t j = 0; j < split_k_count; j++)
{
std::cout << "Run conv :" << conv.GetTypeString() << std::endl;
invoker.ShowInfo(argument);
float avg_time = invoker.Run(argument, StreamConfig{nullptr, config.time_kernel});
if(config.time_kernel)
ck::index_t cur_split_k = split_k_array[j];
auto invoker = conv.MakeInvoker();
auto argument = conv.MakeArgument(static_cast<InDataType*>(in_device_buf.GetDeviceBuffer()),
static_cast<WeiDataType*>(wei_device_buf.GetDeviceBuffer()),
static_cast<OutDataType*>(out_device_buf.GetDeviceBuffer()),
input_lengths,
input_strides,
filter_lengths,
weights_strides,
output_lengths,
output_strides,
conv_filter_strides,
conv_filter_dilations,
input_left_pads,
input_right_pads,
InElementOp{},
WeiElementOp{},
OutElementOp{},
cur_split_k);
DeviceMem gemm_workspace_dev(conv.GetWorkSpaceSize(&argument));
conv.SetWorkSpacePointer(&argument, gemm_workspace_dev.GetDeviceBuffer());
if(conv.IsSupportedArgument(argument))
{
std::size_t flop = conv_param.GetFlops();
std::size_t num_btype = conv_param.GetByte<InDataType, WeiDataType, OutDataType>();
found_kernel = true;
std::cout << "Run conv :" << conv.GetTypeString() << std::endl;
invoker.ShowInfo(argument);
float avg_time = invoker.Run(argument, StreamConfig{nullptr, config.time_kernel});
if(config.time_kernel)
{
std::size_t flop = conv_param.GetFlops();
std::size_t num_btype = conv_param.GetByte<InDataType, WeiDataType, OutDataType>();
float tflops = static_cast<float>(flop) / 1.E9 / avg_time;
float tflops = static_cast<float>(flop) / 1.E9 / avg_time;
float gb_per_sec = num_btype / 1.E6 / avg_time;
float gb_per_sec = num_btype / 1.E6 / avg_time;
std::cerr << "Perf: " << avg_time << " ms, " << tflops << " TFlops, " << gb_per_sec
<< " GB/s" << std::endl
<< "DeviceOp: " << conv.GetTypeString() << std::endl;
}
if(config.do_verification)
{
wei_device_buf.FromDevice(wei_device_result.mData.data());
std::cerr << "Perf: " << avg_time << " ms, " << tflops << " TFlops, " << gb_per_sec
<< " GB/s" << std::endl
<< "DeviceOp: " << conv.GetTypeString() << " Split_K: "<< cur_split_k << std::endl;
if (tflops > best_tflops)
{
best_tflops = tflops;
best_gb_per_sec = gb_per_sec;
best_avg_time = avg_time;
best_split_k = cur_split_k;
best_kernel = conv.GetTypeString();
instance_idx = i;
}
}
if(config.do_verification)
{
wei_device_buf.FromDevice(wei_device_result.mData.data());
ck::utils::check_err(wei_device_result.mData, wei_host_result.mData);
ck::utils::check_err(wei_device_result.mData, wei_host_result.mData);
}
}
}
});
if (found_kernel == false)
{
std::cerr << "wrong! device_conv with the specified compilation parameters does "
"not support this Conv problem"
<< std::endl;
return true;
}
if(config.time_kernel)
{
std::cerr << "Best Perf: " << best_avg_time << " ms, " << best_tflops << " TFlops, " << best_gb_per_sec
<< " GB/s" << std::endl
<< "Instance: " << instance_idx << " DeviceOp: " << best_kernel << " Split_K: "<< best_split_k << std::endl;
}
#else
auto conv = DeviceConvBwdWeightInstance<NDimSpatial>{};
auto invoker = conv.MakeInvoker();

30
script/test_example_conv_bwd.sh
View File

@@ -3,22 +3,22 @@ EXAMPLE="../build/bin/example_grouped_conv_bwd_weight_dl_v4_fp16"
set -x
# G N K C Y X H W Sy Sx Dy Dx Pad
$EXAMPLE 1 2 1 1 2 480 128 1 1 5 5 28 28 1 1 1 1 2 2 2 2
$EXAMPLE 1 2 1 1 2 960 128 1 1 5 5 14 14 1 1 1 1 2 2 2 2
$EXAMPLE 1 2 1 1 2 1344 128 1 1 5 5 14 14 1 1 1 1 2 2 2 2
$EXAMPLE 1 2 1 1 2 2304 128 1 1 5 5 7 7 1 1 1 1 2 2 2 2
$EXAMPLE 1 2 1 1 2 288 128 1 1 5 5 56 56 2 2 1 1 2 2 2 2
$EXAMPLE 1 2 1 1 2 1344 128 1 1 5 5 14 14 2 2 1 1 2 2 2 2
# G N K C Y X H W Sy Sx Dy Dx Pad
$EXAMPLE 1 2 1 -1 2 480 128 1 1 5 5 28 28 1 1 1 1 2 2 2 2
$EXAMPLE 1 2 1 -1 2 960 128 1 1 5 5 14 14 1 1 1 1 2 2 2 2
$EXAMPLE 1 2 1 -1 2 1344 128 1 1 5 5 14 14 1 1 1 1 2 2 2 2
$EXAMPLE 1 2 1 -1 2 2304 128 1 1 5 5 7 7 1 1 1 1 2 2 2 2
$EXAMPLE 1 2 1 -1 2 288 128 1 1 5 5 56 56 2 2 1 1 2 2 2 2
$EXAMPLE 1 2 1 -1 2 1344 128 1 1 5 5 14 14 2 2 1 1 2 2 2 2
$EXAMPLE 1 2 1 1 2 288 128 1 1 3 3 56 56 1 1 1 1 1 1 1 1
$EXAMPLE 1 2 1 1 2 64 128 1 1 3 3 112 112 1 1 1 1 1 1 1 1
$EXAMPLE 1 2 1 1 2 32 128 1 1 3 3 112 112 1 1 1 1 1 1 1 1
$EXAMPLE 1 2 1 1 2 960 128 1 1 3 3 14 14 1 1 1 1 1 1 1 1
$EXAMPLE 1 2 1 1 2 2304 128 1 1 3 3 7 7 1 1 1 1 1 1 1 1
$EXAMPLE 1 2 1 1 2 3840 128 1 1 3 3 7 7 1 1 1 1 1 1 1 1
$EXAMPLE 1 2 1 1 2 480 128 1 1 3 3 28 28 2 2 1 1 1 1 1 1
$EXAMPLE 1 2 1 1 2 192 128 1 1 3 3 112 112 2 2 1 1 1 1 1 1
$EXAMPLE 1 2 1 -1 2 288 128 1 1 3 3 56 56 1 1 1 1 1 1 1 1
$EXAMPLE 1 2 1 -1 2 64 128 1 1 3 3 112 112 1 1 1 1 1 1 1 1
$EXAMPLE 1 2 1 -1 2 32 128 1 1 3 3 112 112 1 1 1 1 1 1 1 1
$EXAMPLE 1 2 1 -1 2 960 128 1 1 3 3 14 14 1 1 1 1 1 1 1 1
$EXAMPLE 1 2 1 -1 2 2304 128 1 1 3 3 7 7 1 1 1 1 1 1 1 1
$EXAMPLE 1 2 1 -1 2 3840 128 1 1 3 3 7 7 1 1 1 1 1 1 1 1
$EXAMPLE 1 2 1 -1 2 480 128 1 1 3 3 28 28 2 2 1 1 1 1 1 1
$EXAMPLE 1 2 1 -1 2 192 128 1 1 3 3 112 112 2 2 1 1 1 1 1 1
set +x