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

Use tile and vector load template parameters.

Browse Source
This commit is contained in:
Ville Pietilä
2025-10-08 11:15:18 +00:00
parent 91465770b0
commit 36f75d061a
1 changed files with 17 additions and 12 deletions
Download Patch File Download Diff File Expand all files Collapse all files

29
example/ck_tile/20_grouped_convolution/grouped_convolution_forward_invoker.hpp
View File

@@ -8,6 +8,8 @@ struct GroupedConvolutionForwardInvoker
{
template <ck_tile::index_t NDimSpatial,
typename GemmWarpConfig,
typename GemmTileConfig,
typename GemmVectorLoads,
typename InDataType,
typename WeiDataType,
typename AccDataType,
@@ -15,29 +17,32 @@ struct GroupedConvolutionForwardInvoker
typename InLayout,
typename WeiLayout,
typename OutLayout,
typename DsDataType = ck_tile::tuple<>,
typename DsLayout = ck_tile::tuple<>,
typename CDEElementWise = ck_tile::element_wise::PassThrough>
ck_tile::index_t NumGroupMerge = 1,
typename DsDataType = ck_tile::tuple<>,
typename DsLayout = ck_tile::tuple<>,
typename CDEElementWise = ck_tile::element_wise::PassThrough>
static float grouped_conv_fwd(const ck_tile::GroupedConvFwdHostArgs& args,
const ck_tile::stream_config& s)
{
constexpr int kBlockPerCu = 1;
constexpr ck_tile::index_t M_Tile = 64;
constexpr ck_tile::index_t N_Tile = 64;
constexpr ck_tile::index_t K_Tile = 64;
constexpr int kBlockPerCu = 1;
constexpr ck_tile::index_t M_Warp = 2;
constexpr ck_tile::index_t N_Warp = 2;
constexpr ck_tile::index_t K_Warp = 1;
constexpr ck_tile::index_t M_Tile = GemmTileConfig::M_Tile;
constexpr ck_tile::index_t N_Tile = GemmTileConfig::N_Tile;
constexpr ck_tile::index_t K_Tile = GemmTileConfig::K_Tile;
constexpr ck_tile::index_t M_Warp = GemmWarpConfig::M_Warp;
constexpr ck_tile::index_t N_Warp = GemmWarpConfig::N_Warp;
constexpr ck_tile::index_t K_Warp = GemmWarpConfig::K_Warp;
constexpr ck_tile::index_t M_Warp_Tile = GemmWarpConfig::M_Warp_Tile;
constexpr ck_tile::index_t N_Warp_Tile = GemmWarpConfig::N_Warp_Tile;
constexpr ck_tile::index_t K_Warp_Tile = GemmWarpConfig::K_Warp_Tile;
constexpr ck_tile::index_t VectorSizeA = 8;
constexpr ck_tile::index_t VectorSizeB = 8;
constexpr ck_tile::index_t VectorSizeC = 8;
constexpr ck_tile::index_t VectorSizeA = GemmVectorLoads::VectorSizeA;
constexpr ck_tile::index_t VectorSizeB = GemmVectorLoads::VectorSizeB;
constexpr ck_tile::index_t VectorSizeC = GemmVectorLoads::VectorSizeC;
// Implicit GEMM Traits
using CodegenShape =