mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-06-30 03:37:38 +00:00
tmp save
This commit is contained in:
Jakub Piasecki
@@ -32,9 +32,9 @@ struct GroupedConvolutionBackwardWeightInvoker
|
||||
GemmConfig::PermuteA,
|
||||
GemmConfig::PermuteB>;
|
||||
|
||||
constexpr ck_tile::index_t VectorSizeA = 4;
|
||||
constexpr ck_tile::index_t VectorSizeB = 8;
|
||||
constexpr ck_tile::index_t VectorSizeC = 8;
|
||||
constexpr ck_tile::index_t VectorSizeA = 1;
|
||||
constexpr ck_tile::index_t VectorSizeB = 2;
|
||||
constexpr ck_tile::index_t VectorSizeC = 2;
|
||||
|
||||
constexpr auto ConvSpec = ck_tile::ConvolutionSpecialization::Default;
|
||||
using TilePartitioner =
|
||||
|
||||
@@ -392,7 +392,7 @@ struct UniversalGemmBasePolicy
|
||||
}
|
||||
|
||||
template <typename Problem, bool IsWave32Host = false>
|
||||
CK_TILE_HOST_DEVICE static constexpr auto GetVectorSizeA()
|
||||
CK_TILE_HOST_DEVICE static constexpr index_t GetVectorSizeA()
|
||||
{
|
||||
using AsLayout = remove_cvref_t<typename Problem::AsLayoutTuple>;
|
||||
using AsDataType = remove_cvref_t<typename Problem::AsDataTypeTuple>;
|
||||
@@ -402,6 +402,10 @@ struct UniversalGemmBasePolicy
|
||||
using ALayout = remove_cvref_t<std::tuple_element_t<number<0>{}, AsLayout>>;
|
||||
using ADataType = remove_cvref_t<std::tuple_element_t<number<0>{}, AsDataType>>;
|
||||
|
||||
if constexpr(Problem::FixedVectorSize) {
|
||||
return Problem::VectorSizeA;
|
||||
}
|
||||
|
||||
if constexpr(std::is_same_v<ALayout, ck_tile::tensor_layout::gemm::RowMajor>)
|
||||
{
|
||||
return GetGlobalVectorLoadSize<Problem,
|
||||
@@ -421,7 +425,7 @@ struct UniversalGemmBasePolicy
|
||||
}
|
||||
|
||||
template <typename Problem, bool IsWave32Host = false>
|
||||
CK_TILE_HOST_DEVICE static constexpr auto GetVectorSizeB()
|
||||
CK_TILE_HOST_DEVICE static constexpr index_t GetVectorSizeB()
|
||||
{
|
||||
using BsLayout = remove_cvref_t<typename Problem::BsLayoutTuple>;
|
||||
using BsDataType = remove_cvref_t<typename Problem::BsDataTypeTuple>;
|
||||
@@ -431,6 +435,10 @@ struct UniversalGemmBasePolicy
|
||||
using BLayout = remove_cvref_t<std::tuple_element_t<number<0>{}, BsLayout>>;
|
||||
using BDataType = remove_cvref_t<std::tuple_element_t<number<0>{}, BsDataType>>;
|
||||
|
||||
if constexpr(Problem::FixedVectorSize) {
|
||||
return Problem::VectorSizeB;
|
||||
}
|
||||
|
||||
if constexpr(std::is_same_v<BLayout, ck_tile::tensor_layout::gemm::RowMajor>)
|
||||
{
|
||||
return GetGlobalVectorLoadSize<Problem,
|
||||
|
||||
@@ -887,8 +887,8 @@ struct GroupedConvolutionBackwardDataKernel
|
||||
const auto& gemm_pad_views = MakeGemmPadViews(gemm_tensor_views_tuple);
|
||||
auto gemm_tile_windows = MakeGemmTileWindows(gemm_pad_views, block_idx_m, block_idx_n);
|
||||
|
||||
const index_t num_loop = amd_wave_read_first_lane(
|
||||
TilePartitioner::GetLoopNum(gemm_tile_windows.at(I0).get_length(I1)));
|
||||
const index_t num_loop = amd_wave_read_first_lane(TilePartitioner::GetLoopNum(
|
||||
gemm_pad_views.at(I0).get_tensor_descriptor().get_length(I1)));
|
||||
|
||||
// Run GEMM cooperatively by whole workgroup.
|
||||
const auto& a_block_window = gemm_tile_windows.at(I0);
|
||||
|
||||
6
include/ck_tile/utility.hpp
Normal file
6
include/ck_tile/utility.hpp
Normal file
@@ -0,0 +1,6 @@
|
||||
// SPDX-License-Identifier: MIT
|
||||
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
#pragma once
|
||||
|
||||
#include "ck_tile/utility/json_dump.hpp"
|
||||
@@ -0,0 +1,162 @@
|
||||
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
#pragma once
|
||||
|
||||
#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_instance_factory.hpp"
|
||||
#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_data_invoker.hpp"
|
||||
|
||||
namespace ck_tile {
|
||||
namespace ops {
|
||||
|
||||
using BF16 = ck_tile::bfloat16_t;
|
||||
|
||||
using DeviceOp2DBF16 = GroupedConvolutionBackwardDataBaseInvoker<2,
|
||||
NHWGC,
|
||||
GKYXC,
|
||||
NHWGK,
|
||||
BF16,
|
||||
BF16,
|
||||
BF16,
|
||||
PassThrough,
|
||||
PassThrough,
|
||||
PassThrough,
|
||||
BF16,
|
||||
BF16>;
|
||||
|
||||
template <ck_tile::index_t NDimSpatial,
|
||||
typename ALayout,
|
||||
typename BLayout,
|
||||
typename ELayout>
|
||||
using tile_grouped_conv_bwd_data_bf16_instances = std::tuple<
|
||||
// clang-format off
|
||||
//#####################################| 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| Double| GEMM|
|
||||
//#####################################| Dim| | | | Type| Type| Type| Elementwise| Elementwise| Elementwise| per| | | | | | | tile| tile| tile| size| size| size| smem| pipeline|
|
||||
//#####################################| Spatial| | | | | | | Operation| Operation| Operation| CU| | | | | | | size| size| size| A| B| C| buffer| version|
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 64, 32, 1, 1, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 64, 32, 1, 1, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 64, 32, 1, 1, 1,16, 16, 32, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 2, 2, 1,32, 32, 16, 1, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 2, 2, 1,32, 32, 16, 1, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 2, 2, 1,32, 32, 16, 1, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 2, 1, 1,32, 32, 16, 8, 1,1, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 2, 1, 1,32, 32, 16, 8, 1,1, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 2, 1, 1,32, 32, 16, 8, 1,1, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 256, 32, 2, 2, 1,32, 32, 16, 8, 4,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 256, 32, 2, 2, 1,32, 32, 16, 8, 4,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 256, 32, 2, 2, 1,32, 32, 16, 8, 4,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 1, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 2, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 2, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 2, 1, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>, //
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 64, 4, 1, 1,32, 32, 16, 16, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 64, 4, 1, 1,32, 32, 16, 16, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 64, 4, 1, 1,32, 32, 16, 16, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 4, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 4, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 4, 1, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 16, 4, 1, 1,32, 32, 16, 4, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 16, 4, 1, 1,32, 32, 16, 4, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 16, 4, 1, 1,32, 32, 16, 4, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 64, 4, 1, 1,32, 32, 16, 16, 2,2, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 64, 4, 1, 1,32, 32, 16, 16, 2,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 64, 4, 1, 1,32, 32, 16, 16, 2,2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 4, 1, 1,32, 32, 16, 8, 2,2, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 4, 1, 1,32, 32, 16, 8, 2,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 4, 1, 1,32, 32, 16, 8, 2,2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 16, 4, 1, 1,32, 32, 16, 4, 2,2, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 16, 4, 1, 1,32, 32, 16, 4, 2,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 16, 4, 1, 1,32, 32, 16, 4, 2,2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 64, 4, 1, 1,32, 32, 16, 16, 1,1, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 64, 4, 1, 1,32, 32, 16, 16, 1,1, false, CK_TILE_PIPELINE_COMPUTE_V3>, //
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 64, 4, 1, 1,32, 32, 16, 16, 1,1, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 4, 1, 1,32, 32, 16, 8, 1,1, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 4, 1, 1,32, 32, 16, 8, 1,1, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 4, 1, 1,32, 32, 16, 8, 1,1, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 16, 4, 1, 1,32, 32, 16, 4, 1,1, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 16, 4, 1, 1,32, 32, 16, 4, 1,1, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 16, 4, 1, 1,32, 32, 16, 4, 1,1, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 64, 4, 1, 1,16, 16, 32, 16, 8,4, false, CK_TILE_PIPELINE_MEMORY>, // test
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 64, 4, 1, 1,16, 16, 32, 16, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 64, 4, 1, 1,16, 16, 32, 16, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 32, 4, 1, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 32, 4, 1, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 32, 4, 1, 1,16, 16, 32, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 16, 4, 1, 1,16, 16, 32, 4, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 16, 4, 1, 1,16, 16, 32, 4, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 16, 4, 1, 1,16, 16, 32, 4, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 64, 4, 1, 1,16, 16, 32, 16, 2,2, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 64, 4, 1, 1,16, 16, 32, 16, 2,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 64, 4, 1, 1,16, 16, 32, 16, 2,2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 32, 4, 1, 1,16, 16, 32, 8, 2,2, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 32, 4, 1, 1,16, 16, 32, 8, 2,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 32, 4, 1, 1,16, 16, 32, 8, 2,2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 16, 4, 1, 1,16, 16, 32, 4, 2,2, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 16, 4, 1, 1,16, 16, 32, 4, 2,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 16, 4, 1, 1,16, 16, 32, 4, 2,2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 64, 4, 1, 1,16, 16, 32, 16, 1,1, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 64, 4, 1, 1,16, 16, 32, 16, 1,1, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 64, 4, 1, 1,16, 16, 32, 16, 1,1, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 32, 4, 1, 1,16, 16, 32, 8, 1,1, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 32, 4, 1, 1,16, 16, 32, 8, 1,1, false, CK_TILE_PIPELINE_COMPUTE_V3>, //
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 32, 4, 1, 1,16, 16, 32, 8, 1,1, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 16, 4, 1, 1,16, 16, 32, 4, 1,1, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 16, 4, 1, 1,16, 16, 32, 4, 1,1, false, CK_TILE_PIPELINE_COMPUTE_V3> //
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 16, 4, 1, 1,16, 16, 32, 4, 1,1, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 64, 32, 1, 1, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
// GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 64, 32, 1, 1, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 64, 32, 1, 1, 1,16, 16, 32, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 64, 32, 1, 1, 1,16, 16, 32, 8, 1,1, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 64, 32, 1, 1, 1,16, 16, 32, 8, 1,1, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 64, 32, 1, 1, 1,16, 16, 32, 8, 1,1, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 64, 32, 1, 1, 1,16, 16, 32, 1, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 64, 32, 1, 1, 1,16, 16, 32, 1, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 64, 32, 1, 1, 1,16, 16, 32, 1, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 32, 1, 1, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 32, 1, 1, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 32, 1, 1, 1,16, 16, 32, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 32, 1, 1, 1,16, 16, 32, 8, 1,1, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 32, 1, 1, 1,16, 16, 32, 8, 1,1, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 32, 1, 1, 1,16, 16, 32, 8, 1,1, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 32, 1, 1, 1,16, 16, 32, 1, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 32, 1, 1, 1,16, 16, 32, 1, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>
|
||||
// // GroupedConvolutionBackwardDataInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 32, 1, 1, 1,16, 16, 32, 1, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>
|
||||
// clang-format on
|
||||
>;
|
||||
|
||||
} // namespace ops
|
||||
} // namespace ck_tile
|
||||
@@ -0,0 +1,131 @@
|
||||
// SPDX-License-Identifier: MIT
|
||||
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <cstdlib>
|
||||
#include <vector>
|
||||
#include <memory>
|
||||
#include <type_traits>
|
||||
|
||||
#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_instance_factory.hpp"
|
||||
#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_data_invoker.hpp"
|
||||
|
||||
namespace ck_tile {
|
||||
namespace ops {
|
||||
|
||||
using DeviceOp2DF16 = GroupedConvolutionBackwardDataBaseInvoker<2,
|
||||
NHWGC,
|
||||
GKYXC,
|
||||
NHWGK,
|
||||
ck_tile::half_t,
|
||||
ck_tile::half_t,
|
||||
ck_tile::half_t,
|
||||
ck_tile::element_wise::PassThrough,
|
||||
ck_tile::element_wise::PassThrough,
|
||||
ck_tile::element_wise::PassThrough,
|
||||
ck_tile::half_t,
|
||||
ck_tile::half_t>;
|
||||
|
||||
using DeviceOp2DBF16 = GroupedConvolutionBackwardDataBaseInvoker<2,
|
||||
NHWGC,
|
||||
GKYXC,
|
||||
NHWGK,
|
||||
ck_tile::bfloat16_t,
|
||||
ck_tile::bfloat16_t,
|
||||
ck_tile::bfloat16_t,
|
||||
ck_tile::element_wise::PassThrough,
|
||||
ck_tile::element_wise::PassThrough,
|
||||
ck_tile::element_wise::PassThrough,
|
||||
ck_tile::bfloat16_t,
|
||||
ck_tile::bfloat16_t>;
|
||||
|
||||
using DeviceOp2DF32 = GroupedConvolutionBackwardDataBaseInvoker<2,
|
||||
NHWGC,
|
||||
GKYXC,
|
||||
NHWGK,
|
||||
float,
|
||||
float,
|
||||
float,
|
||||
ck_tile::element_wise::PassThrough,
|
||||
ck_tile::element_wise::PassThrough,
|
||||
ck_tile::element_wise::PassThrough,
|
||||
float,
|
||||
float>;
|
||||
|
||||
// Forward declarations for instance factory functions
|
||||
// void add_grouped_conv2d_bwd_weight_f16_instances(std::vector<std::unique_ptr<DeviceOp2DF16>>& instances);
|
||||
void add_grouped_conv2d_bwd_data_bf16_instances(std::vector<std::unique_ptr<DeviceOp2DBF16>>& instances);
|
||||
// void add_grouped_conv2d_bwd_weight_bf16_instances_opt(std::vector<std::unique_ptr<DeviceOp2DBF16>>& instances);
|
||||
|
||||
template <ck_tile::index_t NumDimSpatial,
|
||||
typename InLayout,
|
||||
typename WeiLayout,
|
||||
typename OutLayout,
|
||||
typename InDataType,
|
||||
typename WeiDataType,
|
||||
typename OutDataType,
|
||||
typename ComputeTypeA,
|
||||
typename ComputeTypeB>
|
||||
struct DeviceOperationInstanceFactory<GroupedConvolutionBackwardDataBaseInvoker<
|
||||
NumDimSpatial,
|
||||
InLayout,
|
||||
WeiLayout,
|
||||
OutLayout,
|
||||
InDataType,
|
||||
WeiDataType,
|
||||
OutDataType,
|
||||
PassThrough,
|
||||
PassThrough,
|
||||
PassThrough,
|
||||
ComputeTypeA,
|
||||
ComputeTypeB>>
|
||||
{
|
||||
using DeviceOp = GroupedConvolutionBackwardDataBaseInvoker<NumDimSpatial,
|
||||
InLayout,
|
||||
WeiLayout,
|
||||
OutLayout,
|
||||
InDataType,
|
||||
WeiDataType,
|
||||
OutDataType,
|
||||
ck_tile::element_wise::PassThrough,
|
||||
ck_tile::element_wise::PassThrough,
|
||||
ck_tile::element_wise::PassThrough,
|
||||
ComputeTypeA,
|
||||
ComputeTypeB>;
|
||||
|
||||
static auto GetInstances()
|
||||
{
|
||||
std::vector<std::unique_ptr<DeviceOp>> op_ptrs;
|
||||
|
||||
if constexpr(NumDimSpatial == 2)
|
||||
{
|
||||
if constexpr(std::is_same_v<InLayout, NHWGC> && std::is_same_v<WeiLayout, GKYXC> &&
|
||||
std::is_same_v<OutLayout, NHWGK>)
|
||||
{
|
||||
if constexpr(std::is_same_v<InDataType, ck_tile::half_t> &&
|
||||
std::is_same_v<WeiDataType, ck_tile::half_t> &&
|
||||
std::is_same_v<OutDataType, ck_tile::half_t> &&
|
||||
std::is_same_v<ComputeTypeA, ck_tile::half_t> &&
|
||||
std::is_same_v<ComputeTypeB, ck_tile::half_t>)
|
||||
{
|
||||
// add_grouped_conv2d_bwd_weight_f16_instances(op_ptrs);
|
||||
}
|
||||
if constexpr(std::is_same_v<InDataType, ck_tile::bfloat16_t> &&
|
||||
std::is_same_v<WeiDataType, ck_tile::bfloat16_t> &&
|
||||
std::is_same_v<OutDataType, ck_tile::bfloat16_t> &&
|
||||
std::is_same_v<ComputeTypeA, ck_tile::bfloat16_t> &&
|
||||
std::is_same_v<ComputeTypeB, ck_tile::bfloat16_t>)
|
||||
{
|
||||
add_grouped_conv2d_bwd_data_bf16_instances(op_ptrs);
|
||||
// add_grouped_conv2d_bwd_weight_bf16_instances_opt(op_ptrs);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return op_ptrs;
|
||||
}
|
||||
};
|
||||
|
||||
} // namespace ops
|
||||
} // namespace ck_tile
|
||||
@@ -0,0 +1,274 @@
|
||||
// SPDX-License-Identifier: MIT
|
||||
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <cstdlib>
|
||||
#include <vector>
|
||||
#include <memory>
|
||||
|
||||
#include "ck_tile/core.hpp"
|
||||
#include "ck_tile/host/kernel_launch.hpp"
|
||||
#include "ck_tile/ops/epilogue.hpp"
|
||||
#include "ck_tile/ops/gemm.hpp"
|
||||
#include "ck_tile/ops/grouped_convolution.hpp"
|
||||
#include "ck_tile/library/tensor_operation_instance/gpu/gemm_configs.hpp"
|
||||
|
||||
namespace ck_tile {
|
||||
namespace ops {
|
||||
|
||||
template <ck_tile::index_t NDimSpatial,
|
||||
typename InLayout,
|
||||
typename WeiLayout,
|
||||
typename OutLayout,
|
||||
typename InDataType,
|
||||
typename WeiDataType,
|
||||
typename OutDataType,
|
||||
typename InElementwiseOperation,
|
||||
typename WeiElementwiseOperation,
|
||||
typename OutElementwiseOperation,
|
||||
typename ComputeTypeA = InDataType,
|
||||
typename ComputeTypeB = ComputeTypeA>
|
||||
struct GroupedConvolutionBackwardDataBaseInvoker
|
||||
{
|
||||
virtual bool IsSupportedArgument(const ck_tile::GroupedConvBwdDataHostArgs& args) const = 0;
|
||||
virtual float Run(const ck_tile::GroupedConvBwdDataHostArgs& args, bool time_kernel, int n_warmup, int n_repeat) const = 0;
|
||||
virtual std::string GetName(const ck_tile::GroupedConvBwdDataHostArgs& args) const = 0;
|
||||
GroupedConvolutionBackwardDataBaseInvoker() = default;
|
||||
GroupedConvolutionBackwardDataBaseInvoker(const GroupedConvolutionBackwardDataBaseInvoker&) = default;
|
||||
GroupedConvolutionBackwardDataBaseInvoker& operator=(const GroupedConvolutionBackwardDataBaseInvoker&) = default;
|
||||
GroupedConvolutionBackwardDataBaseInvoker(GroupedConvolutionBackwardDataBaseInvoker&&) = default;
|
||||
GroupedConvolutionBackwardDataBaseInvoker& operator=(GroupedConvolutionBackwardDataBaseInvoker&&) = default;
|
||||
virtual ~GroupedConvolutionBackwardDataBaseInvoker() = default;
|
||||
};
|
||||
|
||||
template <
|
||||
ck_tile::index_t NDimSpatial,
|
||||
typename InLayout,
|
||||
typename WeiLayout,
|
||||
typename OutLayout,
|
||||
typename InDataType,
|
||||
typename WeiDataType,
|
||||
typename OutDataType,
|
||||
typename InElementwiseOperation,
|
||||
typename WeiElementwiseOperation,
|
||||
typename OutElementwiseOperation,
|
||||
int kBlockPerCu,
|
||||
ck_tile::index_t M_Tile,
|
||||
ck_tile::index_t N_Tile,
|
||||
ck_tile::index_t K_Tile,
|
||||
ck_tile::index_t M_Warp,
|
||||
ck_tile::index_t N_Warp,
|
||||
ck_tile::index_t K_Warp,
|
||||
ck_tile::index_t M_Warp_Tile,
|
||||
ck_tile::index_t N_Warp_Tile,
|
||||
ck_tile::index_t K_Warp_Tile,
|
||||
ck_tile::index_t VectorSizeA,
|
||||
ck_tile::index_t VectorSizeB,
|
||||
ck_tile::index_t VectorSizeC,
|
||||
bool DoubleSmemBuffer,
|
||||
ck_tile::index_t PipelineVersion>
|
||||
struct GroupedConvolutionBackwardDataInvoker :
|
||||
public GroupedConvolutionBackwardDataBaseInvoker<NDimSpatial,
|
||||
InLayout,
|
||||
WeiLayout,
|
||||
OutLayout,
|
||||
InDataType,
|
||||
WeiDataType,
|
||||
OutDataType,
|
||||
InElementwiseOperation,
|
||||
WeiElementwiseOperation,
|
||||
OutElementwiseOperation>
|
||||
{
|
||||
using GemmShape = ck_tile::TileGemmShape<
|
||||
ck_tile::sequence<M_Tile, N_Tile, K_Tile>,
|
||||
ck_tile::sequence<M_Warp, N_Warp, K_Warp>,
|
||||
ck_tile::sequence<M_Warp_Tile, N_Warp_Tile, K_Warp_Tile>,
|
||||
GemmConfigBase::PermuteA,
|
||||
GemmConfigBase::PermuteB>;
|
||||
|
||||
static constexpr auto ConvSpec = ck_tile::ConvolutionSpecialization::Default;
|
||||
|
||||
using TilePartitioner =
|
||||
ck_tile::GemmSpatiallyLocalTilePartitioner<GemmShape,
|
||||
GemmConfigBase::TileParitionerGroupNum,
|
||||
GemmConfigBase::TileParitionerM01>;
|
||||
|
||||
using GroupedConvTraitsType = ck_tile::GroupedConvTraits<NDimSpatial,
|
||||
ConvSpec,
|
||||
InLayout,
|
||||
WeiLayout,
|
||||
ck_tile::tuple<>, // = DsLayout
|
||||
OutLayout,
|
||||
VectorSizeA,
|
||||
VectorSizeB,
|
||||
VectorSizeC>;
|
||||
|
||||
using GemmUniversalTraits = ck_tile::TileGemmUniversalTraits<
|
||||
GemmConfigBase::kPadM,
|
||||
GemmConfigBase::kPadN,
|
||||
GemmConfigBase::kPadK,
|
||||
DoubleSmemBuffer,
|
||||
typename GroupedConvTraitsType::GroupedConvImplicitGemmTraitsBwdData::AsLayout,
|
||||
typename GroupedConvTraitsType::GroupedConvImplicitGemmTraitsBwdData::BsLayout,
|
||||
typename GroupedConvTraitsType::GroupedConvImplicitGemmTraitsBwdData::CLayout,
|
||||
GemmConfigBase::TransposeC,
|
||||
GemmConfigBase::UseStructuredSparsity,
|
||||
false, // Persistent,
|
||||
GemmConfigBase::NumWaveGroups>;
|
||||
|
||||
using AccDataType = float;
|
||||
using GemmPipelineProblem = ck_tile::GemmPipelineProblem<
|
||||
OutDataType,
|
||||
WeiDataType,
|
||||
AccDataType,
|
||||
GemmShape,
|
||||
typename GroupedConvTraitsType::GroupedConvImplicitGemmTraitsBwdData,
|
||||
ck_tile::element_wise::PassThrough,
|
||||
ck_tile::element_wise::PassThrough,
|
||||
InDataType,
|
||||
true,
|
||||
VectorSizeA,
|
||||
VectorSizeB>;
|
||||
|
||||
using BaseGemmPipeline = typename PipelineTypeTraits<PipelineVersion>::template UniversalGemmPipeline<GemmPipelineProblem>;
|
||||
|
||||
template <bool HasHotLoop, ck_tile::TailNumber TailNumber, ck_tile::memory_operation_enum MemOp>
|
||||
auto CreateKernel() const
|
||||
{
|
||||
constexpr auto scheduler = GemmConfigBase::Scheduler;
|
||||
|
||||
using UniversalGemmProblem =
|
||||
ck_tile::UniversalGemmPipelineProblem<OutDataType,
|
||||
WeiDataType,
|
||||
AccDataType,
|
||||
GemmShape,
|
||||
GemmUniversalTraits,
|
||||
scheduler,
|
||||
HasHotLoop,
|
||||
TailNumber,
|
||||
ck_tile::element_wise::PassThrough,
|
||||
ck_tile::element_wise::PassThrough,
|
||||
InDataType,
|
||||
true,
|
||||
VectorSizeA,
|
||||
VectorSizeB>;
|
||||
|
||||
using GemmPipeline = typename PipelineTypeTraits<PipelineVersion>::template GemmPipeline<UniversalGemmProblem>;
|
||||
|
||||
using CDEElementWise = ck_tile::element_wise::PassThrough;
|
||||
|
||||
using ConvEpilogue = ck_tile::CShuffleEpilogue<ck_tile::CShuffleEpilogueProblem<
|
||||
OutDataType,
|
||||
WeiDataType,
|
||||
ck_tile::tuple<>, // = DsDataType
|
||||
AccDataType,
|
||||
InDataType,
|
||||
typename GroupedConvTraitsType::ImplicitGemmDsLayout,
|
||||
ck_tile::tensor_layout::gemm::RowMajor,
|
||||
CDEElementWise,
|
||||
TilePartitioner::MPerBlock,
|
||||
TilePartitioner::NPerBlock,
|
||||
M_Warp,
|
||||
N_Warp,
|
||||
M_Warp_Tile,
|
||||
N_Warp_Tile,
|
||||
K_Warp_Tile,
|
||||
GemmConfigBase::TransposeC,
|
||||
MemOp,
|
||||
1,
|
||||
true,
|
||||
GroupedConvTraitsType::VectorSizeC>>;
|
||||
|
||||
return ck_tile::GroupedConvolutionBackwardDataKernel<GroupedConvTraitsType,
|
||||
TilePartitioner,
|
||||
GemmPipeline,
|
||||
ConvEpilogue>{};
|
||||
}
|
||||
|
||||
bool IsSupportedArgument(const ck_tile::GroupedConvBwdDataHostArgs& args) const override
|
||||
{
|
||||
if (args.k_batch > 1)
|
||||
{
|
||||
using Kernel = decltype(CreateKernel<false, ck_tile::TailNumber::Empty, ck_tile::memory_operation_enum::atomic_add>());
|
||||
return Kernel::IsSupportedArgument(args);
|
||||
}
|
||||
using Kernel = decltype(CreateKernel<false, ck_tile::TailNumber::Empty, ck_tile::memory_operation_enum::set>());
|
||||
return Kernel::IsSupportedArgument(args);
|
||||
};
|
||||
|
||||
float Run(const ck_tile::GroupedConvBwdDataHostArgs& args, bool time_kernel, int n_warmup=5, int n_repeat=50) const override
|
||||
{
|
||||
const ck_tile::index_t gemm_k =
|
||||
args.N_ * std::accumulate(args.output_spatial_lengths_.begin(),
|
||||
args.output_spatial_lengths_.end(),
|
||||
1,
|
||||
std::multiplies<ck_tile::index_t>());
|
||||
|
||||
const ck_tile::index_t k_grain = args.k_batch * K_Tile;
|
||||
const ck_tile::index_t K_split = (gemm_k + k_grain - 1) / k_grain * K_Tile;
|
||||
const ck_tile::index_t num_loop = TilePartitioner::GetLoopNum(K_split);
|
||||
const bool has_hot_loop = BaseGemmPipeline::BlockHasHotloop(num_loop);
|
||||
const ck_tile::TailNumber tail_num = BaseGemmPipeline::GetBlockLoopTailNum(num_loop);
|
||||
float ave_time{0};
|
||||
|
||||
const auto Run = [&](const auto has_hot_loop_,
|
||||
const auto tail_number_,
|
||||
const auto memory_operation_) {
|
||||
constexpr bool has_hot_loop_v = has_hot_loop_.value;
|
||||
constexpr auto tail_number_v = tail_number_.value;
|
||||
constexpr auto memory_operation = memory_operation_.value;
|
||||
|
||||
auto kernel = CreateKernel<has_hot_loop_v, tail_number_v, memory_operation>();
|
||||
using Kernel = decltype(kernel);
|
||||
|
||||
auto kargs = Kernel::MakeKernelArgs(args);
|
||||
const dim3 grids = Kernel::GridSize(args);
|
||||
const dim3 blocks = Kernel::BlockSize();
|
||||
|
||||
ck_tile::stream_config s {nullptr, time_kernel, 1, n_warmup, n_repeat};
|
||||
|
||||
ave_time = ck_tile::launch_kernel(
|
||||
s, ck_tile::make_kernel<kBlockPerCu>(kernel, grids, blocks, 0, kargs));
|
||||
|
||||
return ave_time;
|
||||
};
|
||||
|
||||
const auto RunSplitk = [&](const auto has_hot_loop_, const auto tail_number_) {
|
||||
if(args.k_batch == 1)
|
||||
{
|
||||
Run(has_hot_loop_, tail_number_, MemoryOpSet{});
|
||||
}
|
||||
else
|
||||
{
|
||||
Run(has_hot_loop_, tail_number_, MemoryOpAtomicAdd{});
|
||||
}
|
||||
};
|
||||
|
||||
BaseGemmPipeline::TailHandler(RunSplitk, has_hot_loop, tail_num);
|
||||
return ave_time;
|
||||
};
|
||||
|
||||
std::string GetName(const ck_tile::GroupedConvBwdDataHostArgs& args) const override
|
||||
{
|
||||
std::stringstream min_occupancy;
|
||||
min_occupancy << "_blk_per_cu_" << kBlockPerCu;
|
||||
if (args.k_batch > 1)
|
||||
{
|
||||
using Kernel = decltype(CreateKernel<false, ck_tile::TailNumber::Empty, ck_tile::memory_operation_enum::atomic_add>());
|
||||
return Kernel::GetName() + min_occupancy.str();
|
||||
}
|
||||
using Kernel = decltype(CreateKernel<false, ck_tile::TailNumber::Empty, ck_tile::memory_operation_enum::set>());
|
||||
return Kernel::GetName() + min_occupancy.str();
|
||||
};
|
||||
|
||||
GroupedConvolutionBackwardDataInvoker() = default;
|
||||
GroupedConvolutionBackwardDataInvoker(const GroupedConvolutionBackwardDataInvoker&) = default;
|
||||
GroupedConvolutionBackwardDataInvoker& operator=(const GroupedConvolutionBackwardDataInvoker&) = default;
|
||||
GroupedConvolutionBackwardDataInvoker(GroupedConvolutionBackwardDataInvoker&&) = default;
|
||||
GroupedConvolutionBackwardDataInvoker& operator=(GroupedConvolutionBackwardDataInvoker&&) = default;
|
||||
~GroupedConvolutionBackwardDataInvoker() override = default;
|
||||
};
|
||||
|
||||
}
|
||||
}
|
||||
@@ -63,24 +63,49 @@ using tile_grouped_conv_bwd_weight_bf16_instances = std::tuple<
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 16, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 16, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 1, 1, 1, 32, 32, 16, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 32, 32, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 32, 32, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 8, 8, 8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 32, 32, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 32, 32, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 32, 32, 64, 1, 1, 1, 16, 16, 32, 8, 8, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 32, 32, 1, 1, 1,32, 32, 16, 2, 2,2, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 32, 32, 1, 1, 1,32, 32, 16, 2, 2,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 32, 32, 1, 1, 1,32, 32, 16, 2, 2,2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 32, 1, 1, 1,32, 32, 16, 4, 4,2, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 32, 1, 1, 1,32, 32, 16, 4, 4,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 32, 1, 1, 1,32, 32, 16, 4, 4,2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 128, 32, 1, 1, 1,32, 32, 16, 8, 8,2, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 128, 32, 1, 1, 1,32, 32, 16, 8, 8,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 128, 32, 1, 1, 1,32, 32, 16, 8, 8,2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 4, 4,2, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 4, 4,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 4, 4,2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,2, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 2, 1, 1,32, 32, 16, 8, 1,1, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 2, 1, 1,32, 32, 16, 8, 1,1, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 2, 1, 1,32, 32, 16, 8, 1,1, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 32, 1, 1, 1,32, 32, 16, 1, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 32, 1, 1, 1,32, 32, 16, 1, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 32, 1, 1, 1,32, 32, 16, 1, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>
|
||||
// clang-format on
|
||||
>;
|
||||
|
||||
|
||||
@@ -41,7 +41,43 @@ using tile_grouped_conv_bwd_weight_bf16_instances_opt = std::tuple<
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 128, 128, 16, 2, 2, 1, 32, 32, 16, 8, 8, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 128, 128, 16, 2, 2, 1, 32, 32, 16, 8, 8, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 128, 128, 16, 2, 2, 1, 32, 32, 16, 8, 8, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 128, 128, 16, 2, 2, 1, 32, 32, 16, 8, 8, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 128, 128, 16, 2, 2, 1, 32, 32, 16, 8, 8, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
|
||||
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 1, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 1, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 1, 1,32, 32, 16, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 2, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 2, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 2, 1,32, 32, 16, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 2, 1, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 2, 1, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 2, 1, 1,32, 32, 16, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionBackwardWeightInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>
|
||||
|
||||
|
||||
// clang-format on
|
||||
>;
|
||||
|
||||
|
||||
@@ -32,37 +32,89 @@ using tile_grouped_conv_fwd_bf16_instances = std::tuple<
|
||||
//##############################| 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| Double| GEMM|
|
||||
//##############################| Dim| | | | Type| Type| Type| Elementwise| Elementwise| Elementwise| per| | | | | | | tile| tile| tile| size| size| size| smem| pipeline|
|
||||
//##############################| Spatial| | | | | | | Operation| Operation| Operation| CU| | | | | | | size| size| size| A| B| C| buffer| version|
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 16, 2, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 8, 2, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 2, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 4, 4, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 8, 8, 8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 8, 2, 2, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 16, 2, 2, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 16, 16, 32, 2, 2, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 1, 1, 1,32, 32, 16, 1, 1,1, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 1, 1, 1,32, 32, 16, 1, 1,1, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 1, 1, 1,32, 32, 16, 1, 1,1, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,1, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,1, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,1, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 1, 1,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 1, 1,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 1, 1,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 1, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>
|
||||
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 16, 2, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 8, 2, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 2, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 4, 4, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 8, 8, 8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 8, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 16, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 16, 16, 32, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>
|
||||
// old
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 1, 1, 1,32, 32, 16, 1, 1,1, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,1, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 1, 1,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 2, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 1, 2,1, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 2, 1,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 4, 4,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
|
||||
// // GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 16, 2, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 8, 2, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 2, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 4, 4, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 8, 8, 8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 8, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 16, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 16, 16, 32, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>
|
||||
// clang-format on
|
||||
>;
|
||||
|
||||
|
||||
@@ -0,0 +1,147 @@
|
||||
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
#pragma once
|
||||
|
||||
#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_instance_factory.hpp"
|
||||
#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_fwd_invoker.hpp"
|
||||
|
||||
namespace ck_tile {
|
||||
namespace ops {
|
||||
|
||||
using BF16 = ck_tile::bfloat16_t;
|
||||
|
||||
using DeviceOpFwd2DBF16 = GroupedConvolutionForwardBaseInvoker<2,
|
||||
NHWGC,
|
||||
GKYXC,
|
||||
NHWGK,
|
||||
BF16,
|
||||
BF16,
|
||||
BF16,
|
||||
PassThrough,
|
||||
PassThrough,
|
||||
PassThrough,
|
||||
BF16,
|
||||
BF16>;
|
||||
|
||||
template <ck_tile::index_t NDimSpatial,
|
||||
typename ALayout,
|
||||
typename BLayout,
|
||||
typename ELayout>
|
||||
using tile_grouped_conv_fwd_bf16_instances_2 = std::tuple<
|
||||
// clang-format off
|
||||
//##############################| 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| Double| GEMM|
|
||||
//##############################| Dim| | | | Type| Type| Type| Elementwise| Elementwise| Elementwise| per| | | | | | | tile| tile| tile| size| size| size| smem| pipeline|
|
||||
//##############################| Spatial| | | | | | | Operation| Operation| Operation| CU| | | | | | | size| size| size| A| B| C| buffer| version|
|
||||
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 16, 64, 2, 1, 1,16, 16, 32, 8, 8,2, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 16, 64, 2, 1, 1,16, 16, 32, 8, 8,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 16, 64, 2, 1, 1,16, 16, 32, 8, 8,2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 16, 128, 1, 1, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 16, 128, 1, 1, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 16, 128, 1, 1, 1,16, 16, 32, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 16, 64, 1, 1, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 16, 64, 1, 1, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 16, 64, 1, 1, 1,16, 16, 32, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 32, 64, 1, 2, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 32, 64, 1, 2, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 32, 64, 1, 2, 1,16, 16, 32, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 32, 64, 4, 1, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 32, 64, 4, 1, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 32, 64, 4, 1, 1,32, 32, 16, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 16, 64, 4, 1, 1,16, 16, 32, 8, 8,2, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 16, 64, 4, 1, 1,16, 16, 32, 8, 8,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 16, 64, 4, 1, 1,16, 16, 32, 8, 8,2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 64, 2, 1, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 64, 2, 1, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 64, 2, 1, 1,32, 32, 16, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 16, 64, 2, 1, 1,16, 16, 32, 8, 8,2, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 16, 64, 2, 1, 1,16, 16, 32, 8, 8,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 16, 64, 2, 1, 1,16, 16, 32, 8, 8,2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 64, 2, 1, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 64, 2, 1, 1,32, 32, 16, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 64, 2, 1, 1,32, 32, 16, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 64, 2, 1, 1,16, 16, 32, 8, 8,2, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 64, 2, 1, 1,16, 16, 32, 8, 8,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 16, 64, 2, 1, 1,16, 16, 32, 8, 8,2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 16, 64, 2, 1, 1,16, 16, 32, 8, 8,2, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 16, 64, 2, 1, 1,16, 16, 32, 8, 8,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 16, 64, 2, 1, 1,16, 16, 32, 8, 8,2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 16, 128, 1, 1, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 16, 128, 1, 1, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 16, 128, 1, 1, 1,16, 16, 32, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 16, 64, 1, 1, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 16, 64, 1, 1, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 16, 64, 1, 1, 1,16, 16, 32, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 32, 64, 1, 2, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 32, 64, 1, 2, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 32, 64, 1, 2, 1,16, 16, 32, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 64, 64, 1, 2, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 64, 64, 1, 2, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 64, 64, 1, 2, 1,16, 16, 32, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 64, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 64, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 64, 1, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 128, 64, 1, 2, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 128, 64, 1, 2, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 128, 64, 1, 2, 1,16, 16, 32, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 128, 64, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 128, 64, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 128, 64, 1, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 256, 64, 1, 4, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 256, 64, 1, 4, 1,16, 16, 32, 8, 8,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 16, 256, 64, 1, 4, 1,16, 16, 32, 8, 8,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 256, 64, 1, 4, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 256, 64, 1, 4, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 256, 64, 1, 4, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>
|
||||
|
||||
|
||||
|
||||
// old
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 1, 1, 1,32, 32, 16, 1, 1,1, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,1, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 1, 1,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 2, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 1, 2,1, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 2, 1,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 4, 4,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
|
||||
// // GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 16, 2, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 8, 2, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 2, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 4, 4, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 8, 8, 8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 8, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 16, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 16, 16, 32, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>
|
||||
// clang-format on
|
||||
>;
|
||||
|
||||
} // namespace ops
|
||||
} // namespace ck_tile
|
||||
@@ -0,0 +1,151 @@
|
||||
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
#pragma once
|
||||
|
||||
#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_instance_factory.hpp"
|
||||
#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_fwd_invoker.hpp"
|
||||
|
||||
namespace ck_tile {
|
||||
namespace ops {
|
||||
|
||||
using BF16 = ck_tile::bfloat16_t;
|
||||
|
||||
using DeviceOpFwd2DBF16 = GroupedConvolutionForwardBaseInvoker<2,
|
||||
NHWGC,
|
||||
GKYXC,
|
||||
NHWGK,
|
||||
BF16,
|
||||
BF16,
|
||||
BF16,
|
||||
PassThrough,
|
||||
PassThrough,
|
||||
PassThrough,
|
||||
BF16,
|
||||
BF16>;
|
||||
|
||||
template <ck_tile::index_t NDimSpatial,
|
||||
typename ALayout,
|
||||
typename BLayout,
|
||||
typename ELayout>
|
||||
using tile_grouped_conv_fwd_bf16_instances_3 = std::tuple<
|
||||
// clang-format off
|
||||
//##############################| 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| Double| GEMM|
|
||||
//##############################| Dim| | | | Type| Type| Type| Elementwise| Elementwise| Elementwise| per| | | | | | | tile| tile| tile| size| size| size| smem| pipeline|
|
||||
//##############################| Spatial| | | | | | | Operation| Operation| Operation| CU| | | | | | | size| size| size| A| B| C| buffer| version|
|
||||
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 1, 2,1, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 1, 2,1, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 1, 2,1, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 2, 1,2, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 2, 1,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 2, 1,2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 4, 4,4, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 4, 4,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 4, 4,4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 64, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 64, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 64, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 64, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 64, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 64, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 64, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 64, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 64, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 64, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 64, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 64, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 64, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 64, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 64, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 256, 32, 2, 2, 1,16, 16, 32, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 256, 32, 2, 2, 1,16, 16, 32, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 256, 32, 2, 2, 1,16, 16, 32, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 224, 256, 64, 2, 2, 1,16, 16, 32, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 224, 256, 64, 2, 2, 1,16, 16, 32, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 224, 256, 64, 2, 2, 1,16, 16, 32, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 224, 64, 2, 2, 1,16, 16, 32, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 224, 64, 2, 2, 1,16, 16, 32, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 224, 64, 2, 2, 1,16, 16, 32, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 64, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 64, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 64, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 64, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 64, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 64, 2, 2, 1,32, 32, 16, 8, 8,8, true, CK_TILE_PIPELINE_COMPUTE_V4>
|
||||
|
||||
// old
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 1, 1, 1,32, 32, 16, 1, 1,1, false, CK_TILE_PIPELINE_MEMORY>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,1, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 1, 1,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 256, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 256, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 64, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 128, 32, 2, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 128, 32, 32, 2, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 128, 32, 1, 2, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 32, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 32, 64, 32, 1, 1, 1,32, 32, 16, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 1, 2,1, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 2, 1,2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 4, 4,4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16,BF16, BF16, PassThrough, PassThrough, PassThrough,1, 64, 64, 32, 2, 2, 1,16, 16, 32, 8, 8,8, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
|
||||
// // GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
// // GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, false, CK_TILE_PIPELINE_COMPUTE_V3>,
|
||||
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 16, 2, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 32, 32, 8, 2, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 2, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 4, 4, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 64, 64, 64, 2, 2, 1, 16, 16, 32, 8, 8, 8, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 8, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 32, 32, 16, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 64, 64, 64, 1, 1, 1, 16, 16, 32, 2, 2, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 8, 8, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 2, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 1, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>,
|
||||
// GroupedConvolutionForwardInvoker<NDimSpatial, ALayout, BLayout, ELayout, BF16, BF16, BF16, PassThrough, PassThrough, PassThrough, 2, 16, 16, 64, 1, 1, 1, 16, 16, 32, 4, 4, 4, true, CK_TILE_PIPELINE_COMPUTE_V4>
|
||||
// clang-format on
|
||||
>;
|
||||
|
||||
} // namespace ops
|
||||
} // namespace ck_tile
|
||||
@@ -41,6 +41,9 @@ using DeviceOpFwd2DF16 = GroupedConvolutionForwardBaseInvoker<2,
|
||||
F16>;
|
||||
|
||||
void add_grouped_conv2d_fwd_bf16_instances(std::vector<std::unique_ptr<DeviceOpFwd2DBF16>>& instances);
|
||||
void add_grouped_conv2d_fwd_bf16_instances_2(std::vector<std::unique_ptr<DeviceOpFwd2DBF16>>& instances);
|
||||
void add_grouped_conv2d_fwd_bf16_instances_3(std::vector<std::unique_ptr<DeviceOpFwd2DBF16>>& instances);
|
||||
|
||||
void add_grouped_conv2d_fwd_f16_instances(std::vector<std::unique_ptr<DeviceOpFwd2DF16>>& instances);
|
||||
|
||||
template <ck_tile::index_t NumDimSpatial,
|
||||
@@ -103,6 +106,8 @@ struct DeviceOperationInstanceFactory<GroupedConvolutionForwardBaseInvoker<
|
||||
std::is_same_v<ComputeTypeB, ck_tile::bfloat16_t>)
|
||||
{
|
||||
add_grouped_conv2d_fwd_bf16_instances(op_ptrs);
|
||||
add_grouped_conv2d_fwd_bf16_instances_2(op_ptrs);
|
||||
add_grouped_conv2d_fwd_bf16_instances_3(op_ptrs);
|
||||
}
|
||||
else
|
||||
{
|
||||
|
||||
@@ -3,6 +3,10 @@
|
||||
# Moving these large template instantiations to separate .cpp files enables parallel compilation
|
||||
|
||||
|
||||
add_library(ck_tile_grouped_conv_bwd_data_bf16_instances STATIC
|
||||
tile_grouped_conv_bwd_data_bf16_instances.cpp
|
||||
)
|
||||
|
||||
add_library(ck_tile_grouped_conv_bwd_weight_fp16_instances STATIC
|
||||
tile_grouped_conv_bwd_weight_fp16_instances.cpp
|
||||
)
|
||||
@@ -23,6 +27,19 @@ add_library(ck_tile_grouped_conv_fwd_bf16_instances STATIC
|
||||
tile_grouped_conv_fwd_bf16_instances.cpp
|
||||
)
|
||||
|
||||
add_library(ck_tile_grouped_conv_fwd_bf16_instances_2 STATIC
|
||||
tile_grouped_conv_fwd_bf16_instances_2.cpp
|
||||
)
|
||||
|
||||
add_library(ck_tile_grouped_conv_fwd_bf16_instances_3 STATIC
|
||||
tile_grouped_conv_fwd_bf16_instances_3.cpp
|
||||
)
|
||||
|
||||
target_include_directories(ck_tile_grouped_conv_bwd_data_bf16_instances PRIVATE
|
||||
${PROJECT_SOURCE_DIR}/include
|
||||
${PROJECT_SOURCE_DIR}/library/include
|
||||
)
|
||||
|
||||
target_include_directories(ck_tile_grouped_conv_bwd_weight_fp16_instances PRIVATE
|
||||
${PROJECT_SOURCE_DIR}/include
|
||||
${PROJECT_SOURCE_DIR}/library/include
|
||||
@@ -47,3 +64,13 @@ target_include_directories(ck_tile_grouped_conv_fwd_bf16_instances PRIVATE
|
||||
${PROJECT_SOURCE_DIR}/include
|
||||
${PROJECT_SOURCE_DIR}/library/include
|
||||
)
|
||||
|
||||
target_include_directories(ck_tile_grouped_conv_fwd_bf16_instances_2 PRIVATE
|
||||
${PROJECT_SOURCE_DIR}/include
|
||||
${PROJECT_SOURCE_DIR}/library/include
|
||||
)
|
||||
|
||||
target_include_directories(ck_tile_grouped_conv_fwd_bf16_instances_3 PRIVATE
|
||||
${PROJECT_SOURCE_DIR}/include
|
||||
${PROJECT_SOURCE_DIR}/library/include
|
||||
)
|
||||
|
||||
@@ -0,0 +1,20 @@
|
||||
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_instance_factory.hpp"
|
||||
#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_data_bf16_instances.hpp"
|
||||
|
||||
namespace ck_tile {
|
||||
namespace ops {
|
||||
|
||||
void add_grouped_conv2d_bwd_data_bf16_instances(std::vector<std::unique_ptr<DeviceOp2DBF16>>& instances)
|
||||
{
|
||||
add_device_operation_instances(instances,
|
||||
tile_grouped_conv_bwd_data_bf16_instances<
|
||||
2,
|
||||
NHWGC,
|
||||
GKYXC,
|
||||
NHWGK>{});
|
||||
}
|
||||
|
||||
} // namespace ops
|
||||
} // namespace ck_tile
|
||||
@@ -0,0 +1,20 @@
|
||||
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_instance_factory.hpp"
|
||||
#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_fwd_bf16_instances_2.hpp"
|
||||
|
||||
namespace ck_tile {
|
||||
namespace ops {
|
||||
|
||||
void add_grouped_conv2d_fwd_bf16_instances_2(std::vector<std::unique_ptr<DeviceOpFwd2DBF16>>& instances)
|
||||
{
|
||||
add_device_operation_instances(instances,
|
||||
tile_grouped_conv_fwd_bf16_instances_2<
|
||||
2,
|
||||
NHWGC,
|
||||
GKYXC,
|
||||
NHWGK>{});
|
||||
}
|
||||
|
||||
} // namespace ops
|
||||
} // namespace ck_tile
|
||||
@@ -0,0 +1,20 @@
|
||||
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_instance_factory.hpp"
|
||||
#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_fwd_bf16_instances_3.hpp"
|
||||
|
||||
namespace ck_tile {
|
||||
namespace ops {
|
||||
|
||||
void add_grouped_conv2d_fwd_bf16_instances_3(std::vector<std::unique_ptr<DeviceOpFwd2DBF16>>& instances)
|
||||
{
|
||||
add_device_operation_instances(instances,
|
||||
tile_grouped_conv_fwd_bf16_instances_3<
|
||||
2,
|
||||
NHWGC,
|
||||
GKYXC,
|
||||
NHWGK>{});
|
||||
}
|
||||
|
||||
} // namespace ops
|
||||
} // namespace ck_tile
|
||||
@@ -0,0 +1,299 @@
|
||||
// SPDX-License-Identifier: MIT
|
||||
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <algorithm>
|
||||
#include <iomanip>
|
||||
#include <iostream>
|
||||
#include <iterator>
|
||||
#include <typeinfo>
|
||||
#include <sstream>
|
||||
|
||||
#include "ck_tile/host.hpp"
|
||||
#include "ck_tile/host/convolution_parameter.hpp"
|
||||
#include "ck_tile/ops/elementwise/unary_element_wise_operation.hpp"
|
||||
#include "ck_tile/ops/grouped_convolution/utils/grouped_convolution_utils.hpp"
|
||||
#include "ck_tile/library/tensor_operation_instance/gpu/tile_grouped_conv_bwd_data_factory.hpp"
|
||||
#include "ck_tile/ops/grouped_convolution/kernel/grouped_convolution_backward_data_kernel.hpp"
|
||||
#include "ck_tile/host/reference/reference_grouped_conv_bwd_data.hpp"
|
||||
namespace ck_tile {
|
||||
namespace profiler {
|
||||
|
||||
template <typename InDataType, typename WeiDataType, typename AccDataType, typename OutDataType>
|
||||
auto calculate_rtol_atol(const ck_tile::index_t GemmK,
|
||||
const ck_tile::index_t kbatch,
|
||||
const float max_accumulated_value)
|
||||
{
|
||||
using ComputeType =
|
||||
std::conditional_t<sizeof(InDataType) < sizeof(WeiDataType), InDataType, WeiDataType>;
|
||||
// Calculate thresholds
|
||||
const auto rtol = ck_tile::get_relative_threshold<ComputeType, OutDataType, AccDataType>(
|
||||
ck_tile::integer_divide_ceil(GemmK, kbatch));
|
||||
const auto atol = ck_tile::get_absolute_threshold<ComputeType, OutDataType, AccDataType>(
|
||||
max_accumulated_value / kbatch, ck_tile::integer_divide_ceil(GemmK, kbatch));
|
||||
// Calculate error due to split_k accumulation
|
||||
const auto rtol_split_k =
|
||||
ck_tile::get_relative_threshold<OutDataType, OutDataType, OutDataType>(kbatch);
|
||||
const auto atol_split_k =
|
||||
ck_tile::get_absolute_threshold<OutDataType, OutDataType, OutDataType>(
|
||||
max_accumulated_value, kbatch);
|
||||
// Use higher threshold
|
||||
return ck_tile::make_tuple(std::max(rtol, rtol_split_k), std::max(atol, atol_split_k));
|
||||
}
|
||||
|
||||
template <ck_tile::index_t NDimSpatial,
|
||||
typename InLayout,
|
||||
typename WeiLayout,
|
||||
typename OutLayout,
|
||||
typename InDataType,
|
||||
typename WeiDataType,
|
||||
typename OutDataType,
|
||||
typename ComputeTypeA = InDataType,
|
||||
typename ComputeTypeB = ComputeTypeA>
|
||||
bool profile_grouped_conv_bwd_data_impl(int do_verification,
|
||||
int init_method,
|
||||
bool /*do_log*/,
|
||||
bool time_kernel,
|
||||
const ck_tile::conv::ConvParam& conv_param,
|
||||
const std::string& split_k,
|
||||
ck_tile::index_t instance_index = -1)
|
||||
{
|
||||
using AccDataType = float;
|
||||
using InElementOp = ck_tile::element_wise::PassThrough;
|
||||
using WeiElementOp = ck_tile::element_wise::PassThrough;
|
||||
using OutElementOp = ck_tile::element_wise::PassThrough;
|
||||
|
||||
const auto in_g_n_c_wis_desc =
|
||||
ck_tile::conv::make_input_host_tensor_descriptor_g_n_c_wis_packed<InLayout>(conv_param);
|
||||
const auto wei_g_k_c_xs_desc =
|
||||
ck_tile::conv::make_weight_host_tensor_descriptor_g_k_c_xs_packed<WeiLayout>(conv_param);
|
||||
const auto out_g_n_k_wos_desc =
|
||||
ck_tile::conv::make_output_host_tensor_descriptor_g_n_k_wos_packed<OutLayout>(conv_param);
|
||||
|
||||
ck_tile::HostTensor<InDataType> input(in_g_n_c_wis_desc);
|
||||
ck_tile::HostTensor<WeiDataType> weight(wei_g_k_c_xs_desc);
|
||||
ck_tile::HostTensor<OutDataType> output(out_g_n_k_wos_desc);
|
||||
|
||||
std::cout << "input: " << input.mDesc << std::endl;
|
||||
std::cout << "weight: " << weight.mDesc << std::endl;
|
||||
std::cout << "output: " << output.mDesc << std::endl;
|
||||
|
||||
switch(init_method)
|
||||
{
|
||||
case 0:
|
||||
ck_tile::FillUniformDistribution<WeiDataType>{-1.f, 1.f}(weight);
|
||||
ck_tile::FillUniformDistribution<OutDataType>{-1.f, 1.f}(output);
|
||||
break;
|
||||
case 1:
|
||||
ck_tile::FillMonotonicSeq<WeiDataType>{}(weight);
|
||||
ck_tile::FillMonotonicSeq<OutDataType>{}(output);
|
||||
break;
|
||||
case 2:
|
||||
ck_tile::FillUniformDistribution<WeiDataType>{1.f, 1.f}(weight);
|
||||
ck_tile::FillUniformDistribution<OutDataType>{1.f, 1.f}(output);
|
||||
break;
|
||||
default:
|
||||
input.SetZero();
|
||||
output.SetZero();
|
||||
}
|
||||
|
||||
using DeviceOp = ops::GroupedConvolutionBackwardDataBaseInvoker<
|
||||
NDimSpatial,
|
||||
InLayout,
|
||||
WeiLayout,
|
||||
OutLayout,
|
||||
InDataType,
|
||||
WeiDataType,
|
||||
OutDataType,
|
||||
InElementOp,
|
||||
WeiElementOp,
|
||||
OutElementOp,
|
||||
ComputeTypeA,
|
||||
ComputeTypeB>;
|
||||
|
||||
// get device op instances
|
||||
const auto ops = ck_tile::ops::DeviceOperationInstanceFactory<DeviceOp>::GetInstances();
|
||||
|
||||
std::cout << "found " << ops.size() << " instances" << std::endl;
|
||||
|
||||
std::string best_op_name;
|
||||
float best_avg_time = 0;
|
||||
float best_tflops = 0;
|
||||
float best_gb_per_sec = 0;
|
||||
std::string best_split_k("1");
|
||||
|
||||
std::vector<ck_tile::index_t> split_k_list = {1, 2, 4, 8, 16, 32, 64, 128};
|
||||
if(split_k != "all")
|
||||
{
|
||||
try
|
||||
{
|
||||
ck_tile::index_t split_k_value = std::stoi(split_k);
|
||||
split_k_list = {split_k_value};
|
||||
}
|
||||
catch(const std::exception& e)
|
||||
{
|
||||
std::cerr << e.what() << '\n';
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
}
|
||||
|
||||
// First, calculate the reference result if verification is needed.
|
||||
ck_tile::HostTensor<InDataType> input_host_ref(in_g_n_c_wis_desc);
|
||||
input_host_ref.SetZero();
|
||||
if (do_verification)
|
||||
{
|
||||
ck_tile::reference_grouped_conv_bwd_data<NDimSpatial, InDataType, WeiDataType, OutDataType>(
|
||||
input_host_ref,
|
||||
weight,
|
||||
output,
|
||||
conv_param.conv_filter_strides_,
|
||||
conv_param.conv_filter_dilations_,
|
||||
conv_param.input_left_pads_,
|
||||
conv_param.input_right_pads_);
|
||||
}
|
||||
|
||||
|
||||
index_t num_kernel = 0;
|
||||
bool all_pass = true;
|
||||
for(auto& op : ops)
|
||||
{
|
||||
for(std::size_t split_k_id = 0; split_k_id < split_k_list.size(); split_k_id++)
|
||||
{
|
||||
auto split_k_value = split_k_list[split_k_id];
|
||||
auto split_k_param_str = std::to_string(split_k_value);
|
||||
|
||||
ck_tile::DeviceMem input_dev_buf(input.get_element_space_size_in_bytes());
|
||||
ck_tile::DeviceMem weight_dev_buf(weight.get_element_space_size_in_bytes());
|
||||
ck_tile::DeviceMem output_dev_buf(output.get_element_space_size_in_bytes());
|
||||
|
||||
input_dev_buf.SetZero();
|
||||
weight_dev_buf.ToDevice(weight.data());
|
||||
output_dev_buf.ToDevice(output.data());
|
||||
|
||||
ck_tile::GroupedConvBwdDataHostArgs args(conv_param,
|
||||
input_dev_buf.GetDeviceBuffer(),
|
||||
weight_dev_buf.GetDeviceBuffer(),
|
||||
{},
|
||||
output_dev_buf.GetDeviceBuffer(),
|
||||
split_k_value);
|
||||
|
||||
// Split-K autodeduction is not supported.
|
||||
if(op->IsSupportedArgument(args) && split_k_value >= 1)
|
||||
{
|
||||
num_kernel++;
|
||||
if((instance_index != -1) && (instance_index + 1 != num_kernel))
|
||||
{
|
||||
// skip test if instance_index is specified
|
||||
continue;
|
||||
}
|
||||
|
||||
std::string op_name = op->GetName(args);
|
||||
std::cout << op_name << ", SplitK " << split_k_param_str << " is profiled..." << std::endl;
|
||||
|
||||
// Run verification first. If it doesn't pass, no need to do performance measurement.
|
||||
bool pass = false;
|
||||
if(do_verification)
|
||||
{
|
||||
constexpr int n_warmup = 0;
|
||||
constexpr int n_repeat = 1;
|
||||
|
||||
op->Run(args, false, n_warmup, n_repeat);
|
||||
input_dev_buf.FromDevice(input.data());
|
||||
|
||||
const ck_tile::index_t GemmK = weight.get_element_size() / (conv_param.G_ * conv_param.K_);
|
||||
const float max_accumulated_value =
|
||||
*std::max_element(input_host_ref.mData.begin(), input_host_ref.mData.end());
|
||||
const auto rtol_atol =
|
||||
calculate_rtol_atol<InDataType, WeiDataType, AccDataType, OutDataType>(
|
||||
GemmK, split_k_value, max_accumulated_value);
|
||||
|
||||
pass = ck_tile::check_err(input,
|
||||
input_host_ref,
|
||||
"Error: Incorrect results!",
|
||||
rtol_atol.at(ck_tile::number<0>{}),
|
||||
rtol_atol.at(ck_tile::number<1>{}));
|
||||
|
||||
std::cout << "Relative error threshold: " << rtol_atol.at(ck_tile::number<0>{})
|
||||
<< " Absolute error threshold: " << rtol_atol.at(ck_tile::number<1>{})
|
||||
<< std::endl;
|
||||
std::cout << "The CPU verification result is:" << (pass ? "correct" : "fail") << std::endl;
|
||||
|
||||
all_pass &= pass;
|
||||
}
|
||||
|
||||
bool is_valid = do_verification ? pass : true;
|
||||
|
||||
if (is_valid)
|
||||
{
|
||||
constexpr int n_warmup = 5;
|
||||
constexpr int n_repeat = 50;
|
||||
float avg_time = op->Run(args, time_kernel, n_warmup, n_repeat);
|
||||
|
||||
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 gb_per_sec = num_btype / 1.E6 / avg_time;
|
||||
|
||||
std::cout << "Perf: " << std::setw(10) << avg_time << " ms, " << tflops
|
||||
<< " TFlops, " << gb_per_sec << " GB/s, " << op_name << ", SplitK "
|
||||
<< split_k_param_str << std::endl;
|
||||
|
||||
if(tflops > best_tflops)
|
||||
{
|
||||
best_op_name = op_name;
|
||||
best_tflops = tflops;
|
||||
best_avg_time = avg_time;
|
||||
best_gb_per_sec = gb_per_sec;
|
||||
best_split_k = split_k_param_str;
|
||||
}
|
||||
}
|
||||
}
|
||||
else
|
||||
{
|
||||
std::cout << op->GetName(args) << ", SplitK " << split_k_param_str
|
||||
<< " does not support this problem." << std::endl;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
std::stringstream ss;
|
||||
ss << "\n********************************"
|
||||
<< "\nCK Tile best configuration parameters:"
|
||||
<< "\n********************************"
|
||||
<< "\nname: " << best_op_name
|
||||
<< "\navg_time: " << best_avg_time << "\ntflops: " << best_tflops
|
||||
<< "\nGB/s: " << best_gb_per_sec
|
||||
<< "\nSplitK: " << best_split_k
|
||||
<< std::endl;
|
||||
|
||||
std::cout << ss.str();
|
||||
|
||||
const char* log_file = std::getenv("CK_TILE_PROFILER_LOG_FILE");
|
||||
if(log_file != nullptr)
|
||||
{
|
||||
std::ofstream out(log_file, std::ios::app);
|
||||
if(out.is_open())
|
||||
{
|
||||
std::stringstream out_ss;
|
||||
out_ss << "CK Tile best configuration:" << std::endl
|
||||
<< "name: " << best_op_name << std::endl
|
||||
<< "avg_time: " << best_avg_time << std::endl
|
||||
<< "SplitK: " << best_split_k << std::endl
|
||||
<< "all_pass " << (all_pass ? "true" : "false") << std::endl;
|
||||
out << out_ss.str();
|
||||
out.close();
|
||||
}
|
||||
}
|
||||
|
||||
if(instance_index != -1)
|
||||
{
|
||||
std::cout << "grouped_conv_bwd_data_instance (" << instance_index << "/" << num_kernel
|
||||
<< "): Passed" << std::endl;
|
||||
}
|
||||
return all_pass;
|
||||
}
|
||||
|
||||
} // namespace profiler
|
||||
} // namespace ck_tile
|
||||
@@ -280,7 +280,8 @@ bool profile_grouped_conv_bwd_weight_impl(int do_verification,
|
||||
out_ss << "CK Tile best configuration:" << std::endl
|
||||
<< "name: " << best_op_name << std::endl
|
||||
<< "avg_time: " << best_avg_time << std::endl
|
||||
<< "SplitK: " << best_split_k << std::endl;
|
||||
<< "SplitK: " << best_split_k << std::endl
|
||||
<< "all_pass " << (all_pass ? "true" : "false") << std::endl;
|
||||
out << out_ss.str();
|
||||
out.close();
|
||||
}
|
||||
|
||||
@@ -91,8 +91,8 @@ bool profile_grouped_conv_fwd_impl(int do_verification,
|
||||
}
|
||||
else if(init_method == 2)
|
||||
{
|
||||
ck_tile::FillUniformDistribution<InDataType>{1.f, 1.f}(input);
|
||||
ck_tile::FillUniformDistribution<WeiDataType>{1.f, 1.f}(weight);
|
||||
ck_tile::FillUniformDistribution<InDataType>{0.f, 1.f}(input);
|
||||
ck_tile::FillUniformDistribution<WeiDataType>{0.f, 1.f}(weight);
|
||||
}
|
||||
else
|
||||
{
|
||||
@@ -133,7 +133,7 @@ bool profile_grouped_conv_fwd_impl(int do_verification,
|
||||
ck_tile::DeviceMem output_dev_buf(output.get_element_space_size_in_bytes());
|
||||
|
||||
input_dev_buf.ToDevice(input.data());
|
||||
weight_dev_buf.ToDevice(output.data());
|
||||
weight_dev_buf.ToDevice(weight.data());
|
||||
output_dev_buf.SetZero();
|
||||
|
||||
ck_tile::GroupedConvFwdHostArgs args(conv_param,
|
||||
@@ -233,6 +233,23 @@ bool profile_grouped_conv_fwd_impl(int do_verification,
|
||||
<< "\navg_time: " << best_avg_time << "\ntflops: " << best_tflops
|
||||
<< "\nGB/s: " << best_gb_per_sec << std::endl;
|
||||
|
||||
const char* log_file = std::getenv("CK_TILE_PROFILER_LOG_FILE");
|
||||
if(log_file != nullptr)
|
||||
{
|
||||
std::ofstream out(log_file, std::ios::app);
|
||||
if(out.is_open())
|
||||
{
|
||||
std::stringstream out_ss;
|
||||
out_ss << "CK Tile best configuration:" << std::endl
|
||||
<< "name: " << best_op_name << std::endl
|
||||
<< "avg_time: " << best_avg_time << std::endl
|
||||
<< "SplitK: " << 1 << std::endl
|
||||
<< "all_pass " << (all_pass ? "true" : "false") << std::endl;
|
||||
out << out_ss.str();
|
||||
out.close();
|
||||
}
|
||||
}
|
||||
|
||||
if(instance_index != -1)
|
||||
{
|
||||
std::cout << "grouped_conv_fwd_instance (" << instance_index << "/" << num_kernel
|
||||
|
||||
@@ -11,6 +11,7 @@ message(STATUS "CK_PROFILER_OP_FILTER: ${CK_PROFILER_OP_FILTER}")
|
||||
message(STATUS "CK_PROFILER_INSTANCE_FILTER: ${CK_PROFILER_INSTANCE_FILTER}")
|
||||
|
||||
if(SUPPORTED_GPU_TARGETS MATCHES "gfx9" OR SUPPORTED_GPU_TARGETS MATCHES "gfx1[12]")
|
||||
list(APPEND PROFILER_OPS tile_profile_grouped_conv_bwd_data.cpp)
|
||||
list(APPEND PROFILER_OPS tile_profile_grouped_conv_bwd_weight.cpp)
|
||||
list(APPEND PROFILER_OPS tile_profile_grouped_conv_fwd.cpp)
|
||||
endif()
|
||||
@@ -39,11 +40,15 @@ endif()
|
||||
|
||||
set(DEVICE_INSTANCES "")
|
||||
if(SUPPORTED_GPU_TARGETS MATCHES "gfx9" OR SUPPORTED_GPU_TARGETS MATCHES "gfx1[12]")
|
||||
list(APPEND DEVICE_INSTANCES ck_tile_grouped_conv_bwd_data_bf16_instances)
|
||||
list(APPEND DEVICE_INSTANCES ck_tile_grouped_conv_bwd_weight_fp16_instances)
|
||||
list(APPEND DEVICE_INSTANCES ck_tile_grouped_conv_bwd_weight_bf16_instances)
|
||||
list(APPEND DEVICE_INSTANCES ck_tile_grouped_conv_bwd_weight_bf16_instances_opt)
|
||||
list(APPEND DEVICE_INSTANCES ck_tile_grouped_conv_fwd_fp16_instances)
|
||||
list(APPEND DEVICE_INSTANCES ck_tile_grouped_conv_fwd_bf16_instances)
|
||||
list(APPEND DEVICE_INSTANCES ck_tile_grouped_conv_fwd_bf16_instances_2)
|
||||
list(APPEND DEVICE_INSTANCES ck_tile_grouped_conv_fwd_bf16_instances_3)
|
||||
|
||||
endif()
|
||||
|
||||
if(SUPPORTED_GPU_TARGETS MATCHES "gfx9" OR SUPPORTED_GPU_TARGETS MATCHES "gfx1[12]")
|
||||
|
||||
219
profiler/ck_tile/src/tile_profile_grouped_conv_bwd_data.cpp
Normal file
219
profiler/ck_tile/src/tile_profile_grouped_conv_bwd_data.cpp
Normal file
@@ -0,0 +1,219 @@
|
||||
// SPDX-License-Identifier: MIT
|
||||
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
#include <cstdlib>
|
||||
#include <initializer_list>
|
||||
#include <iostream>
|
||||
#include <numeric>
|
||||
|
||||
#include "tile_profile_grouped_conv_bwd_data_impl.hpp"
|
||||
#include "tile_profiler_operation_registry.hpp"
|
||||
|
||||
// CK Tile library dependencies
|
||||
#include "ck_tile/core/numeric/integral_constant.hpp"
|
||||
#include "ck_tile/ops/common/tensor_layout.hpp"
|
||||
|
||||
namespace {
|
||||
|
||||
enum struct ConvLayout
|
||||
{
|
||||
GNHWC_GKYXC_GNHWK, // 0
|
||||
NHWGC_GKYXC_NHWGK, // 1
|
||||
NGCHW_GKYXC_NGKHW, // 2
|
||||
NGCHW_GKCYX_NGKHW, // 3
|
||||
};
|
||||
|
||||
enum struct ConvDataType
|
||||
{
|
||||
F32_F32_F32, // 0
|
||||
F16_F16_F16, // 1
|
||||
BF16_BF16_BF16, // 2
|
||||
};
|
||||
|
||||
#define OP_NAME "grouped_conv_bwd_data"
|
||||
#define OP_DESC "Grouped Convolution Backward Data"
|
||||
|
||||
static void print_helper_msg()
|
||||
{
|
||||
std::string conv_param_parser_helper_msg;
|
||||
|
||||
conv_param_parser_helper_msg += "Following arguments (depending on number of spatial dims):\n"
|
||||
" Number of spatial dimensions (1=Conv1d, 2=Conv2d, 3=Conv3d)\n"
|
||||
" G, N, K, C, \n"
|
||||
" <filter spatial dimensions>, (ie Y, X for 2D)\n"
|
||||
" <input image spatial dimensions>, (ie Hi, Wi for 2D)\n"
|
||||
" <strides>, (ie Sy, Sx for 2D)\n"
|
||||
" <dilations>, (ie Dy, Dx for 2D)\n"
|
||||
" <left padding>, (ie LeftPy, LeftPx for 2D)\n"
|
||||
" <right padding>, (ie RightPy, RightPx for 2D)\n";
|
||||
|
||||
std::cout << "arg1: tensor operation (" OP_NAME ": " OP_DESC ")\n"
|
||||
<< "arg2: data type (0: Input fp32, Weight fp32, Output fp32\n"
|
||||
<< " 1: Input fp16, Weight fp16, Output fp16\n"
|
||||
<< " 2: Input bf16, Weight fp32, Output bf16\n"
|
||||
<< " 3: Input fp16, Weight fp16, Output fp16, Gemm bf8@fp8\n"
|
||||
<< " 4: Input int8, Weight int8, Output int8\n"
|
||||
<< " 5: Input bf16, Weight bf16, Output bf16\n"
|
||||
<< " 6: Input fp32, Weight fp32, Output fp32, Compute tf32)\n"
|
||||
<< "arg3: tensor layout (0: Input[G, N, C, Hi, Wi], Weight[G, K, C, Y, X], Output[G, "
|
||||
"N, K, Ho, Wo]\n"
|
||||
<< " 1: Input[G, N, Hi, Wi, C], Weight[G, K, Y, X, C], Output[G, "
|
||||
"N, Ho, Wo, K]\n"
|
||||
<< " 2: Input[N, Hi, Wi, G, C], Weight[G, K, Y, X, C], Output[N, "
|
||||
"Ho, Wo, G, K]\n"
|
||||
<< " 3: Input[N, G, C, Hi, Wi], Weight[G, K, Y, X, C], Output[N, "
|
||||
"G, K, Ho, Wo]\n"
|
||||
<< " 4: Input[N, G, C, Hi, Wi], Weight[G, K, C, Y, X], Output[N, "
|
||||
"G, K, Ho, Wo]\n"
|
||||
<< "arg4: verification (0: no, 1: yes)\n"
|
||||
<< "arg5: initialization (0: no init, 1: integer value, 2: decimal value)\n"
|
||||
<< "arg6: print tensor value (0: no; 1: yes)\n"
|
||||
<< "arg7: time kernel (0: no, 1: yes)\n"
|
||||
<< conv_param_parser_helper_msg
|
||||
<< " SplitK (-1 for internally computed split-K value, positive value to set k "
|
||||
"batches explicitly, or 'all' to test all internal split-K values)\n"
|
||||
<< std::endl;
|
||||
}
|
||||
|
||||
} // namespace
|
||||
|
||||
int tile_profile_grouped_conv_bwd_data(int argc, char* argv[])
|
||||
{
|
||||
// 8 for control, 1 for num_dim_spatial
|
||||
if(argc < 9)
|
||||
{
|
||||
print_helper_msg();
|
||||
return 1;
|
||||
}
|
||||
|
||||
const auto data_type = static_cast<ConvDataType>(std::stoi(argv[2]));
|
||||
const auto layout = static_cast<ConvLayout>(std::stoi(argv[3]));
|
||||
const bool do_verification = std::stoi(argv[4]);
|
||||
const int init_method = std::stoi(argv[5]);
|
||||
const bool do_log = std::stoi(argv[6]);
|
||||
const bool time_kernel = std::stoi(argv[7]);
|
||||
const int num_dim_spatial = std::stoi(argv[8]);
|
||||
|
||||
// 8 for control, 1 for num_dim_spatial, 4 for G/N/K/C, and 6 * num_dim_spatial, 1 for split-K
|
||||
if(argc != 8 + 1 + 4 + 6 * num_dim_spatial + 1)
|
||||
{
|
||||
print_helper_msg();
|
||||
return 1;
|
||||
}
|
||||
|
||||
const auto params = ck_tile::conv::parse_conv_param(num_dim_spatial, 9, argv);
|
||||
|
||||
const auto& split_k = std::string(argv[8 + 1 + 4 + 6 * num_dim_spatial]);
|
||||
|
||||
// using F32 = float;
|
||||
// using F16 = ck_tile::half_t;
|
||||
using BF16 = ck_tile::bfloat16_t;
|
||||
// using F8 = ck_tile::fp8_t;
|
||||
// using BF8 = ck_tile::bf8_t;
|
||||
|
||||
using NHWGC = ck_tile::tensor_layout::convolution::NHWGC;
|
||||
// using NDHWGC = ck_tile::tensor_layout::convolution::NDHWGC;
|
||||
|
||||
using GKYXC = ck_tile::tensor_layout::convolution::GKYXC;
|
||||
// using GKZYXC = ck_tile::tensor_layout::convolution::GKZYXC;
|
||||
|
||||
using NHWGK = ck_tile::tensor_layout::convolution::NHWGK;
|
||||
// using NDHWGK = ck_tile::tensor_layout::convolution::NDHWGK;
|
||||
|
||||
constexpr auto I2 = ck_tile::number<2>{};
|
||||
//constexpr auto I3 = ck_tile::number<3>{};
|
||||
|
||||
auto profile = [&](auto num_dim_spatial_tmp,
|
||||
auto in_layout,
|
||||
auto wei_layout,
|
||||
auto out_layout,
|
||||
auto in_type,
|
||||
auto wei_type,
|
||||
auto out_type,
|
||||
auto compute_type_a,
|
||||
auto compute_type_b) {
|
||||
constexpr ck_tile::index_t NDimSpatial = num_dim_spatial_tmp.value;
|
||||
|
||||
using InLayout = decltype(in_layout);
|
||||
using WeiLayout = decltype(wei_layout);
|
||||
using OutLayout = decltype(out_layout);
|
||||
|
||||
using InDataType = decltype(in_type);
|
||||
using WeiDataType = decltype(wei_type);
|
||||
using OutDataType = decltype(out_type);
|
||||
|
||||
using ComputeTypeA = decltype(compute_type_a);
|
||||
using ComputeTypeB = decltype(compute_type_b);
|
||||
|
||||
bool pass = ck_tile::profiler::profile_grouped_conv_bwd_data_impl<NDimSpatial,
|
||||
InLayout,
|
||||
WeiLayout,
|
||||
OutLayout,
|
||||
InDataType,
|
||||
WeiDataType,
|
||||
OutDataType,
|
||||
ComputeTypeA,
|
||||
ComputeTypeB>(
|
||||
do_verification, init_method, do_log, time_kernel, params, split_k);
|
||||
|
||||
return pass ? 0 : 1;
|
||||
};
|
||||
|
||||
if(num_dim_spatial == 2 && layout == ConvLayout::NHWGC_GKYXC_NHWGK)
|
||||
{
|
||||
// if(data_type == ConvDataType::F32_F32_F32)
|
||||
// {
|
||||
// return profile(I2, NHWGC{}, GKYXC{}, NHWGK{}, F32{}, F32{}, F32{}, F32{}, F32{});
|
||||
// }
|
||||
// if(data_type == ConvDataType::F16_F16_F16)
|
||||
// {
|
||||
// return profile(I2, NHWGC{}, GKYXC{}, NHWGK{}, F16{}, F16{}, F16{}, F16{}, F16{});
|
||||
// }
|
||||
// if(data_type == ConvDataType::BF16_F32_BF16)
|
||||
// {
|
||||
// // fp32 atomic add is used for weight tensor in bf16 kernel
|
||||
// return profile(I2, NHWGC{}, GKYXC{}, NHWGK{}, BF16{}, F32{}, BF16{}, BF16{}, BF16{});
|
||||
// }
|
||||
if(data_type == ConvDataType::BF16_BF16_BF16)
|
||||
{
|
||||
return profile(I2, NHWGC{}, GKYXC{}, NHWGK{}, BF16{}, BF16{}, BF16{}, BF16{}, BF16{});
|
||||
}
|
||||
}
|
||||
|
||||
// if(num_dim_spatial == 3 && layout == ConvLayout::NHWGC_GKYXC_NHWGK)
|
||||
// {
|
||||
// if(data_type == ConvDataType::F32_F32_F32)
|
||||
// {
|
||||
// return profile(I3, NDHWGC{}, GKZYXC{}, NDHWGK{}, F32{}, F32{}, F32{}, F32{}, F32{});
|
||||
// }
|
||||
// if(data_type == ConvDataType::F16_F16_F16)
|
||||
// {
|
||||
// return profile(I3, NDHWGC{}, GKZYXC{}, NDHWGK{}, F16{}, F16{}, F16{}, F16{}, F16{});
|
||||
// }
|
||||
// if(data_type == ConvDataType::BF16_F32_BF16)
|
||||
// {
|
||||
// // fp32 atomic add is used for weight tensor in bf16 kernel
|
||||
// return profile(I3, NDHWGC{}, GKZYXC{}, NDHWGK{}, BF16{}, F32{}, BF16{}, BF16{}, BF16{});
|
||||
// }
|
||||
// if(data_type == ConvDataType::BF16_BF16_BF16)
|
||||
// {
|
||||
// return profile(
|
||||
// I3, NDHWGC{}, GKZYXC{}, NDHWGK{}, BF16{}, BF16{}, BF16{}, BF16{}, BF16{});
|
||||
// }
|
||||
// if(data_type == ConvDataType::F16_F16_F16_BF8_F8)
|
||||
// {
|
||||
// return profile(I3, NDHWGC{}, GKZYXC{}, NDHWGK{}, F16{}, F16{}, F16{}, BF8{}, F8{});
|
||||
// }
|
||||
// else if(data_type == ConvDataType::I8_I8_I8)
|
||||
// {
|
||||
// return profile(
|
||||
// I3, NDHWGC{}, GKZYXC{}, NDHWGK{}, int8_t{}, int8_t{}, int8_t{}, int8_t{}, int8_t{});
|
||||
// }
|
||||
// }
|
||||
|
||||
std::cout << "this data_type & layout is not implemented" << std::endl;
|
||||
|
||||
return 1;
|
||||
}
|
||||
|
||||
REGISTER_PROFILER_OPERATION(OP_NAME, OP_DESC, tile_profile_grouped_conv_bwd_data);
|
||||
@@ -66,11 +66,10 @@ static void print_helper_msg()
|
||||
"G, K, Ho, Wo]\n"
|
||||
<< " 3: Input[N, G, C, Hi, Wi], Weight[G, K, C, Y, X], Output[N, "
|
||||
"G, K, Ho, Wo])\n"
|
||||
<< "arg4: indexing data type (0: 32-bit, 1: 64-bit)\n"
|
||||
<< "arg5: verification (0: no, 1: yes)\n"
|
||||
<< "arg6: initialization (0: no init, 1: integer value, 2: decimal value)\n"
|
||||
<< "arg7: print tensor value (0: no; 1: yes)\n"
|
||||
<< "arg8: time kernel (0: no, 1: yes)\n"
|
||||
<< "arg4: verification (0: no, 1: yes)\n"
|
||||
<< "arg5: initialization (0: no init, 1: integer value, 2: decimal value)\n"
|
||||
<< "arg6: print tensor value (0: no; 1: yes)\n"
|
||||
<< "arg7: time kernel (0: no, 1: yes)\n"
|
||||
<< conv_param_parser_helper_msg << std::endl;
|
||||
// clang-format on
|
||||
}
|
||||
@@ -80,7 +79,7 @@ static void print_helper_msg()
|
||||
int tile_profile_grouped_conv_fwd(int argc, char* argv[])
|
||||
{
|
||||
// 8 for control, 1 for num_dim_spatial
|
||||
if(argc < 10)
|
||||
if(argc < 9)
|
||||
{
|
||||
print_helper_msg();
|
||||
return 1;
|
||||
@@ -88,20 +87,21 @@ int tile_profile_grouped_conv_fwd(int argc, char* argv[])
|
||||
|
||||
const auto data_type = static_cast<ConvDataType>(std::stoi(argv[2]));
|
||||
const auto layout = static_cast<ConvLayout>(std::stoi(argv[3]));
|
||||
const bool do_verification = std::stoi(argv[5]);
|
||||
const int init_method = std::stoi(argv[6]);
|
||||
const bool do_log = std::stoi(argv[7]);
|
||||
const bool time_kernel = std::stoi(argv[8]);
|
||||
const int num_dim_spatial = std::stoi(argv[9]);
|
||||
const bool do_verification = std::stoi(argv[4]);
|
||||
const int init_method = std::stoi(argv[5]);
|
||||
const bool do_log = std::stoi(argv[6]);
|
||||
const bool time_kernel = std::stoi(argv[7]);
|
||||
const int num_dim_spatial = std::stoi(argv[8]);
|
||||
|
||||
// 9 for control, 1 for num_dim_spatial, 4 for G/N/K/C, and 6 * num_dim_spatial
|
||||
if(argc != 9 + 1 + 4 + 6 * num_dim_spatial)
|
||||
if(argc != 8 + 1 + 4 + 6 * num_dim_spatial + 1)
|
||||
{
|
||||
std::cout << argc << std::endl;
|
||||
print_helper_msg();
|
||||
return 1;
|
||||
}
|
||||
|
||||
const auto params = ck_tile::conv::parse_conv_param(num_dim_spatial, 10, argv);
|
||||
const auto params = ck_tile::conv::parse_conv_param(num_dim_spatial, 9, argv);
|
||||
constexpr ck_tile::index_t k_batch = 1;
|
||||
|
||||
using F32 = float;
|
||||
|
||||
@@ -36,7 +36,7 @@ bool profile_grouped_conv_bwd_data_impl(int do_verification,
|
||||
bool do_log,
|
||||
bool time_kernel,
|
||||
const ck::utils::conv::ConvParam& conv_param,
|
||||
ck::index_t split_k = 1,
|
||||
ck::index_t split_k = 0,
|
||||
index_t instance_index = -1)
|
||||
{
|
||||
using OutElementOp = ck::tensor_operation::element_wise::PassThrough;
|
||||
@@ -316,6 +316,23 @@ bool profile_grouped_conv_bwd_data_impl(int do_verification,
|
||||
<< "\navg_time: " << best_avg_time << "\ntflops: " << best_tflops
|
||||
<< "\nGB/s: " << best_gb_per_sec << ", SplitK " << best_split_k << std::endl;
|
||||
|
||||
const char* log_file = std::getenv("CK_PROFILER_LOG_FILE");
|
||||
if(log_file != nullptr)
|
||||
{
|
||||
std::ofstream out_stream(log_file, std::ios::app);
|
||||
if(out_stream.is_open())
|
||||
{
|
||||
std::stringstream out_ss;
|
||||
out_ss << "CK best configuration:" << std::endl
|
||||
<< "name: " << best_op_name << std::endl
|
||||
<< "avg_time: " << best_avg_time << std::endl
|
||||
<< "SplitK: " << best_split_k << std::endl;
|
||||
out_stream << out_ss.str();
|
||||
out_stream.close();
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
if(instance_index != -1)
|
||||
{
|
||||
std::cout << "grouped_conv_bwd_data_instance (" << instance_index << "/" << num_kernel
|
||||
|
||||
@@ -210,6 +210,8 @@ bool profile_grouped_conv_fwd_impl(int do_verification,
|
||||
{
|
||||
std::cout << op_ptr->GetTypeString() << " does not support this problem" << std::endl;
|
||||
}
|
||||
|
||||
|
||||
};
|
||||
|
||||
using DeviceOp = ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD<NDimSpatial,
|
||||
@@ -261,6 +263,23 @@ bool profile_grouped_conv_fwd_impl(int do_verification,
|
||||
std::cout << "Best configuration parameters:" << "\nname: " << best_op_name
|
||||
<< "\navg_time: " << best_avg_time << "\ntflops: " << best_tflops
|
||||
<< "\nGB/s: " << best_gb_per_sec << std::endl;
|
||||
|
||||
const char* log_file = std::getenv("CK_PROFILER_LOG_FILE");
|
||||
if(log_file != nullptr)
|
||||
{
|
||||
std::ofstream out(log_file, std::ios::app);
|
||||
if(out.is_open())
|
||||
{
|
||||
std::stringstream out_ss;
|
||||
out_ss << "CK best configuration:" << std::endl
|
||||
<< "name: " << best_op_name << std::endl
|
||||
<< "avg_time: " << best_avg_time << std::endl
|
||||
<< "SplitK: " << 1 << std::endl;
|
||||
out << out_ss.str();
|
||||
out.close();
|
||||
}
|
||||
}
|
||||
|
||||
if(instance_index != -1)
|
||||
{
|
||||
std::cout << "grouped_conv_fwd_instance (" << instance_index << "/" << num_kernel
|
||||
|
||||
@@ -79,7 +79,7 @@ int profile_grouped_conv_bwd_data(int argc, char* argv[])
|
||||
|
||||
const auto params = ck::utils::conv::parse_conv_param(num_dim_spatial, 9, argv);
|
||||
|
||||
ck::index_t split_k = std::stoi(argv[8 + 1 + 4 + 6 * num_dim_spatial]);
|
||||
//ck::index_t split_k = std::stoi(argv[8 + 1 + 4 + 6 * num_dim_spatial]);
|
||||
|
||||
using F32 = float;
|
||||
using F16 = ck::half_t;
|
||||
@@ -120,7 +120,7 @@ int profile_grouped_conv_bwd_data(int argc, char* argv[])
|
||||
WeiDataType,
|
||||
InDataType,
|
||||
ComputeDataType>(
|
||||
do_verification, init_method, do_log, time_kernel, params, split_k);
|
||||
do_verification, init_method, do_log, time_kernel, params);
|
||||
|
||||
return pass ? 0 : 1;
|
||||
};
|
||||
|
||||
@@ -61,11 +61,10 @@ static void print_helper_msg()
|
||||
"G, K, Ho, Wo]\n"
|
||||
<< " 3: Input[N, G, C, Hi, Wi], Weight[G, K, C, Y, X], Output[N, "
|
||||
"G, K, Ho, Wo])\n"
|
||||
<< "arg4: indexing data type (0: 32-bit, 1: 64-bit)\n"
|
||||
<< "arg5: verification (0: no, 1: yes)\n"
|
||||
<< "arg6: initialization (0: no init, 1: integer value, 2: decimal value)\n"
|
||||
<< "arg7: print tensor value (0: no; 1: yes)\n"
|
||||
<< "arg8: time kernel (0: no, 1: yes)\n"
|
||||
<< "arg4: verification (0: no, 1: yes)\n"
|
||||
<< "arg5: initialization (0: no init, 1: integer value, 2: decimal value)\n"
|
||||
<< "arg6: print tensor value (0: no; 1: yes)\n"
|
||||
<< "arg7: time kernel (0: no, 1: yes)\n"
|
||||
<< ck::utils::conv::get_conv_param_parser_helper_msg() << std::endl;
|
||||
// clang-format on
|
||||
}
|
||||
@@ -75,7 +74,7 @@ static void print_helper_msg()
|
||||
int profile_grouped_conv_fwd(int argc, char* argv[])
|
||||
{
|
||||
// 8 for control, 1 for num_dim_spatial
|
||||
if(argc < 10)
|
||||
if(argc < 9)
|
||||
{
|
||||
print_helper_msg();
|
||||
return 1;
|
||||
@@ -83,21 +82,21 @@ int profile_grouped_conv_fwd(int argc, char* argv[])
|
||||
|
||||
const auto data_type = static_cast<ConvDataType>(std::stoi(argv[2]));
|
||||
const auto layout = static_cast<ConvLayout>(std::stoi(argv[3]));
|
||||
const auto index_type = static_cast<IndexType>(std::stoi(argv[4]));
|
||||
const bool do_verification = std::stoi(argv[5]);
|
||||
const int init_method = std::stoi(argv[6]);
|
||||
const bool do_log = std::stoi(argv[7]);
|
||||
const bool time_kernel = std::stoi(argv[8]);
|
||||
const int num_dim_spatial = std::stoi(argv[9]);
|
||||
const auto index_type = IndexType::INDEX_T;
|
||||
const bool do_verification = std::stoi(argv[4]);
|
||||
const int init_method = std::stoi(argv[5]);
|
||||
const bool do_log = std::stoi(argv[6]);
|
||||
const bool time_kernel = std::stoi(argv[7]);
|
||||
const int num_dim_spatial = std::stoi(argv[8]);
|
||||
|
||||
// 9 for control, 1 for num_dim_spatial, 4 for G/N/K/C, and 6 * num_dim_spatial
|
||||
if(argc != 9 + 1 + 4 + 6 * num_dim_spatial)
|
||||
if(argc != 8 + 1 + 4 + 6 * num_dim_spatial + 1)
|
||||
{
|
||||
print_helper_msg();
|
||||
return 1;
|
||||
}
|
||||
|
||||
const auto params = ck::utils::conv::parse_conv_param(num_dim_spatial, 10, argv);
|
||||
const auto params = ck::utils::conv::parse_conv_param(num_dim_spatial, 9, argv);
|
||||
|
||||
using F32 = float;
|
||||
using F16 = ck::half_t;
|
||||
|
||||
@@ -9,6 +9,8 @@ import matplotlib.pyplot as plt
|
||||
plt.switch_backend('Agg')
|
||||
import numpy as np
|
||||
|
||||
import xlsxwriter
|
||||
|
||||
def parse_cli_args():
|
||||
"""Parse command line arguments"""
|
||||
parser = argparse.ArgumentParser(description="Run CK and CK Tile convolution profilers.")
|
||||
@@ -45,7 +47,7 @@ def run_ck_profiler_cmd(cmd_args, profiler_type, bin_path, results_file, log_to_
|
||||
subprocess.run(cmd)
|
||||
else:
|
||||
with open(os.devnull, 'w') as devnull:
|
||||
timeoutInSec = 15 * 60 # 15 minutes timeout
|
||||
timeoutInSec = 300 * 60 # 300 minutes timeout
|
||||
try:
|
||||
subprocess.run(cmd, stdout=devnull, stderr=devnull, timeout=timeoutInSec, env=env)
|
||||
except subprocess.TimeoutExpired:
|
||||
@@ -77,8 +79,8 @@ def run_analysis(results_file):
|
||||
while i < len(lines):
|
||||
line = lines[i].strip()
|
||||
|
||||
# Look for grouped_conv_bwd_weight command lines
|
||||
if line.startswith('grouped_conv_bwd_weight'):
|
||||
# Look for grouped_conv_* command lines
|
||||
if line.startswith('grouped_conv_'):
|
||||
current_case = {'command': line}
|
||||
i += 1
|
||||
|
||||
@@ -97,6 +99,9 @@ def run_analysis(results_file):
|
||||
if i < len(lines) and lines[i].strip().startswith('SplitK:'):
|
||||
current_case['ck_tile_splitk'] = lines[i].strip().replace('SplitK:', '').strip()
|
||||
i += 1
|
||||
if i < len(lines) and lines[i].strip().startswith('all_pass'):
|
||||
current_case['ck_tile_all_pass'] = lines[i].strip().replace('all_pass', '').strip()
|
||||
i += 1
|
||||
|
||||
# Parse CK results
|
||||
while i < len(lines) and not lines[i].strip().startswith('CK best configuration:'):
|
||||
@@ -129,25 +134,82 @@ def run_analysis(results_file):
|
||||
ck_times = []
|
||||
ck_tile_times = []
|
||||
case_labels = []
|
||||
|
||||
|
||||
workbook = xlsxwriter.Workbook('conv_perf.xlsx')
|
||||
worksheet = workbook.add_worksheet()
|
||||
|
||||
header_format = workbook.add_format()
|
||||
header_format.set_bold()
|
||||
|
||||
offset = 4
|
||||
|
||||
worksheet.write(offset, 0, "command", header_format)
|
||||
worksheet.set_column(0, 0, 66)
|
||||
worksheet.write(offset, 1, "CK Time", header_format)
|
||||
worksheet.set_column(1, 1, 11)
|
||||
worksheet.write(offset, 2, "CK Tile Time", header_format)
|
||||
worksheet.set_column(2, 2, 11)
|
||||
worksheet.write(offset, 3, "CK / CK Tile", header_format)
|
||||
worksheet.set_column(3, 3, 11)
|
||||
worksheet.write(offset, 4, "All pass", header_format)
|
||||
worksheet.set_column(4, 4, 11)
|
||||
worksheet.write(offset, 5, "CK best kernel", header_format)
|
||||
worksheet.set_column(5, 5, 25)
|
||||
worksheet.write(offset, 6, "CK tile best kernel", header_format)
|
||||
worksheet.set_column(6, 6, 25)
|
||||
|
||||
offset += 1
|
||||
|
||||
num_of_ck_tile_slower = 0
|
||||
|
||||
for i, case in enumerate(test_cases):
|
||||
worksheet.write(i + offset, 0, case['command'])
|
||||
worksheet.write(i + offset, 1, case['ck_time'])
|
||||
worksheet.write(i + offset, 2, case['ck_tile_time'])
|
||||
|
||||
format = workbook.add_format()
|
||||
ratio = case['ck_time'] / case['ck_tile_time']
|
||||
|
||||
if ratio < 1.0:
|
||||
format.set_bg_color('red')
|
||||
num_of_ck_tile_slower += 1
|
||||
else:
|
||||
format.set_bg_color('green')
|
||||
|
||||
all_pass = case['ck_tile_all_pass']
|
||||
|
||||
worksheet.write(i + offset, 3, ratio, format)
|
||||
|
||||
format2 = workbook.add_format()
|
||||
format2.set_bg_color('green' if all_pass == "true" else 'red')
|
||||
worksheet.write(i + offset, 4, all_pass, format2)
|
||||
worksheet.write(i + offset, 5, case['ck_name'])
|
||||
worksheet.write(i + offset, 6, case['ck_tile_name'])
|
||||
|
||||
ck_time = case['ck_time']
|
||||
ck_tile_time = case['ck_tile_time']
|
||||
|
||||
# Performance ratio: CK_time / CK_Tile_time * 100%
|
||||
# >100% means CK Tile is faster, <100% means CK is faster
|
||||
ratio = (ck_time / ck_tile_time) * 100
|
||||
performance_ratios.append(ratio)
|
||||
ck_times.append(ck_time)
|
||||
ck_tile_times.append(ck_tile_time)
|
||||
# ratio = (ck_time / ck_tile_time) * 100
|
||||
# performance_ratios.append(ratio)
|
||||
# ck_times.append(ck_time)
|
||||
# ck_tile_times.append(ck_tile_time)
|
||||
|
||||
# Create a short label for the test case
|
||||
cmd_parts = case['command'].split()
|
||||
if len(cmd_parts) >= 8:
|
||||
label = f"G{cmd_parts[8]}_N{cmd_parts[9]}_K{cmd_parts[10]}_C{cmd_parts[11]}"
|
||||
else:
|
||||
label = f"Case_{i+1}"
|
||||
case_labels.append(label)
|
||||
# # Create a short label for the test case
|
||||
# cmd_parts = case['command'].split()
|
||||
# if len(cmd_parts) >= 8:
|
||||
# label = f"G{cmd_parts[8]}_N{cmd_parts[9]}_K{cmd_parts[10]}_C{cmd_parts[11]}"
|
||||
# else:
|
||||
# label = f"Case_{i+1}"
|
||||
# case_labels.append(label)
|
||||
|
||||
worksheet.write(0, 0, f"all cases: {len(test_cases)}")
|
||||
worksheet.write(1, 0, f"ck tile slower: {num_of_ck_tile_slower}")
|
||||
worksheet.write(2, 0, f"ck tile slower: {(num_of_ck_tile_slower / len(test_cases) * 100):2.1f}%")
|
||||
|
||||
workbook.close()
|
||||
return
|
||||
|
||||
|
||||
max_cases_to_detailed_plot = 10
|
||||
|
||||
@@ -33,7 +33,9 @@ def parse_miopen_command(miopen_cmd):
|
||||
def determine_operation_type(params):
|
||||
"""Determine the operation type based on MIOpen parameters"""
|
||||
# TODO: Current data is for bwd weight.
|
||||
return "grouped_conv_bwd_weight"
|
||||
return "grouped_conv_bwd_data"
|
||||
#return "grouped_conv_bwd_weight"
|
||||
#return "grouped_conv_fwd"#"grouped_conv_bwd_weight"
|
||||
|
||||
def convert_miopen_to_ck_profiler(miopen_cmd):
|
||||
"""Convert MIOpen driver command to CK profiler command"""
|
||||
@@ -44,8 +46,8 @@ def convert_miopen_to_ck_profiler(miopen_cmd):
|
||||
# Determine operation type
|
||||
operation = determine_operation_type(params)
|
||||
|
||||
data_type = 5 # BF16
|
||||
layout = 2 # channels last
|
||||
data_type = 2 #2 for bwd data 2 FOR FWD 5 FOR BWD WEI # BF16
|
||||
layout = 1 #1 FIR BWD DATA 1 FOR FWD 2 FOR BWE WEI # channels last
|
||||
verification = 1 # with verification
|
||||
init_method = 2 # uniform data
|
||||
print_output = 0 # no print output
|
||||
@@ -172,11 +174,14 @@ def parse_arguments():
|
||||
|
||||
def main():
|
||||
args = parse_arguments()
|
||||
print(args.input_file)
|
||||
|
||||
# Check if input file exists
|
||||
if not os.path.exists(args.input_file):
|
||||
print(f"Error: Input file '{args.input_file}' does not exist")
|
||||
sys.exit(1)
|
||||
|
||||
print(args.input_file)
|
||||
|
||||
# Generate output filename if not provided
|
||||
output_file = args.output if args.output else generate_output_filename(args.input_file)
|
||||
|
||||
128
script/convert_old_ck_conv_bwd_data_to_ck_tile.py
Normal file
128
script/convert_old_ck_conv_bwd_data_to_ck_tile.py
Normal file
@@ -0,0 +1,128 @@
|
||||
import re
|
||||
|
||||
def extract_template_parameters(template_str):
|
||||
# Extract everything inside the outermost <>
|
||||
match = re.search(r"<(.*)>", template_str, re.DOTALL)
|
||||
if not match:
|
||||
return []
|
||||
|
||||
inside = match.group(1).strip()
|
||||
|
||||
params = []
|
||||
current = []
|
||||
depth = 0 # track nested < >
|
||||
|
||||
for char in inside:
|
||||
if char == '<':
|
||||
depth += 1
|
||||
current.append(char)
|
||||
elif char == '>':
|
||||
depth -= 1
|
||||
current.append(char)
|
||||
elif char == ',' and depth == 0:
|
||||
param = ''.join(current).strip()
|
||||
if param:
|
||||
params.append(param)
|
||||
current = []
|
||||
else:
|
||||
current.append(char)
|
||||
|
||||
# Append last parameter if any
|
||||
if current:
|
||||
params.append(''.join(current).strip())
|
||||
|
||||
return params
|
||||
|
||||
|
||||
input_path = "inputkernel.txt"
|
||||
output_path = "outputkernel_bwd_data.txt"
|
||||
|
||||
with open(input_path, 'r') as f:
|
||||
lines = f.readlines()
|
||||
|
||||
for line in lines:
|
||||
|
||||
# Example usage
|
||||
#input_str = " DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle<NDimSpatial,ALayout,BLayout, DsLayout,ELayout, F16, F16, F32, F16, DsDataTypes, F16, PassThrough, PassThrough, OutElementOp, ConvSpec, GemmMNKPadding, 1, 64, 32, 64, 32, 8, 8, 32, 32, 1, 2, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 16, 1, 4>, 8>"
|
||||
|
||||
params = extract_template_parameters(line)
|
||||
|
||||
NDimSpatial = params[0]
|
||||
ALayout = params[1]
|
||||
BLayout = params[2]
|
||||
DsLayout = params[3]
|
||||
ELayout = params[4]
|
||||
|
||||
ADataType = params[5]
|
||||
BDataType = params[6]
|
||||
AccDataType = params[7]
|
||||
CshuffleDataType= params[8]
|
||||
DsDataTypes = params[9]
|
||||
EDataType = params[10]
|
||||
|
||||
AElementwiseOp = params[11]
|
||||
BElementwiseOp = params[12]
|
||||
CElementwiseOp = "PassThrough"#params[13]
|
||||
ConvFwdSpec = params[14]
|
||||
|
||||
DoPadGemmM = params[15]
|
||||
DoPadGemmN = params[16]
|
||||
NumGemmK = params[17]
|
||||
|
||||
BlockSize = params[18]
|
||||
MPerBlock = params[19]
|
||||
NPerBlock = params[20]
|
||||
KPerBlock = params[21]
|
||||
AK1 = params[22]
|
||||
BK1 = params[23]
|
||||
MPerXDL = params[24]
|
||||
NPerXDL = params[25]
|
||||
MXdlPerWave = params[26]
|
||||
NXdlPerWave = params[27]
|
||||
ABlockTransferClusterLengths = params[28]
|
||||
ABlockTransferArrangeOrder = params[29]
|
||||
ABlockTransferSrcAccessOrder = params[30]
|
||||
ABlockTransferSrcVectorDim = params[31]
|
||||
ABlockTransferSrcScalarPerVector = params[32]
|
||||
ABlockTransferDstScalarPerVector_K1 = params[33]
|
||||
ABlockLdsAddExtraM = params[34]
|
||||
BBlockTransferClusterLengths = params[35]
|
||||
BBlockTransferArrangeOrder = params[36]
|
||||
BBlockTransferSrcVectorDim = params[37]
|
||||
BBlockTransferSrcAccessOrder = params[38]
|
||||
BBlockTransferSrcScalarPerVector = params[39]
|
||||
BBlockTransferDstScalarPerVector_K1 = params[40]
|
||||
BBlockLdsAddExtraM = params[41]
|
||||
CShuffleMXdlPerwave = params[42]
|
||||
CShuffleNXdlPerWavePerShuffle = params[43]
|
||||
CBlockTransferClusterLengths = params[44]
|
||||
CBlockTransferScalarPerVector = params[45]
|
||||
|
||||
|
||||
KBlockPerCu = 1
|
||||
MWarp = int(MPerBlock) // (int(MPerXDL) * int(MXdlPerWave))
|
||||
NWarp = int(NPerBlock) // (int(NPerXDL) * int(NXdlPerWave))
|
||||
KWarp = 1
|
||||
KPerXdl = 16 if MPerXDL == "32" else 32
|
||||
DoubleSMemBuffer = 'false'
|
||||
GemmPipelineVersion = "CK_TILE_PIPELINE_COMPUTE_V3"
|
||||
|
||||
print(MPerBlock, NPerBlock, KPerBlock)
|
||||
|
||||
pipelines = ["CK_TILE_PIPELINE_MEMORY", "CK_TILE_PIPELINE_COMPUTE_V3", "CK_TILE_PIPELINE_COMPUTE_V4"]
|
||||
|
||||
for pipeline in pipelines:
|
||||
DoubleSMemBuffer = 'false' if pipeline != 'CK_TILE_PIPELINE_COMPUTE_V4' else 'true'
|
||||
with open(output_path, 'a') as f:
|
||||
f.write(f'GroupedConvolutionBackwardDataInvoker<{NDimSpatial}, {ALayout}, {BLayout}, {ELayout}, {ADataType},'
|
||||
f'{BDataType}, {EDataType}, {AElementwiseOp}, {BElementwiseOp}, {CElementwiseOp},'
|
||||
f'{KBlockPerCu}, {MPerBlock}, {NPerBlock}, {KPerBlock}, {MWarp}, {NWarp}, {KWarp},'
|
||||
f'{MPerXDL}, {NPerXDL}, {KPerXdl}, {ABlockTransferSrcScalarPerVector}, {BBlockTransferSrcScalarPerVector},'
|
||||
f'{CBlockTransferScalarPerVector}, {DoubleSMemBuffer}, {pipeline}>,\n')
|
||||
|
||||
|
||||
# print(params[0])
|
||||
|
||||
# # Print each parameter as a separate variable
|
||||
# for i, p in enumerate(params, start=1):
|
||||
# print(f"param_{i} = '{p}'")
|
||||
120
script/convert_old_ck_conv_bwd_wei_to_ck_tile.py
Normal file
120
script/convert_old_ck_conv_bwd_wei_to_ck_tile.py
Normal file
@@ -0,0 +1,120 @@
|
||||
import re
|
||||
|
||||
def extract_template_parameters(template_str):
|
||||
# Extract everything inside the outermost <>
|
||||
match = re.search(r"<(.*)>", template_str, re.DOTALL)
|
||||
if not match:
|
||||
return []
|
||||
|
||||
inside = match.group(1).strip()
|
||||
|
||||
params = []
|
||||
current = []
|
||||
depth = 0 # track nested < >
|
||||
|
||||
for char in inside:
|
||||
if char == '<':
|
||||
depth += 1
|
||||
current.append(char)
|
||||
elif char == '>':
|
||||
depth -= 1
|
||||
current.append(char)
|
||||
elif char == ',' and depth == 0:
|
||||
param = ''.join(current).strip()
|
||||
if param:
|
||||
params.append(param)
|
||||
current = []
|
||||
else:
|
||||
current.append(char)
|
||||
|
||||
# Append last parameter if any
|
||||
if current:
|
||||
params.append(''.join(current).strip())
|
||||
|
||||
return params
|
||||
|
||||
|
||||
input_path = "inputkernel.txt"
|
||||
output_path = "outputkernel_bwd_wei.txt"
|
||||
|
||||
with open(input_path, 'r') as f:
|
||||
lines = f.readlines()
|
||||
|
||||
for line in lines:
|
||||
|
||||
# Example usage
|
||||
#input_str = " DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle<NDimSpatial,ALayout,BLayout, DsLayout,ELayout, F16, F16, F32, F16, DsDataTypes, F16, PassThrough, PassThrough, OutElementOp, ConvSpec, GemmMNKPadding, 1, 64, 32, 64, 32, 8, 8, 32, 32, 1, 2, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 16, 1, 4>, 8>"
|
||||
|
||||
params = extract_template_parameters(line)
|
||||
|
||||
NDimSpatial = params[0]
|
||||
ALayout = params[1]
|
||||
BLayout = params[2]
|
||||
ELayout = params[3]
|
||||
|
||||
ADataType = params[4]
|
||||
BDataType = params[5]
|
||||
EDataType = params[6]
|
||||
AccDataType = params[7]
|
||||
|
||||
AElementwiseOp = params[8]
|
||||
BElementwiseOp = params[9]
|
||||
CElementwiseOp = "PassThrough"#params[13]
|
||||
ConvFwdSpec = params[11]
|
||||
|
||||
BlockSize = params[12]
|
||||
MPerBlock = params[13]
|
||||
NPerBlock = params[14]
|
||||
KPerBlock = params[15]
|
||||
K1 = params[16]
|
||||
MPerXDL = params[17]
|
||||
NPerXDL = params[18]
|
||||
MXdlPerWave = params[19]
|
||||
NXdlPerWave = params[20]
|
||||
ABlockTransferClusterLengths = params[21]
|
||||
ABlockTransferArrangeOrder = params[22]
|
||||
ABlockTransferSrcAccessOrder = params[23]
|
||||
ABlockTransferSrcVectorDim = params[24]
|
||||
ABlockTransferSrcScalarPerVector = params[25]
|
||||
ABlockTransferDstScalarPerVector_K1 = params[26]
|
||||
ABlockLdsAddExtraM = params[27]
|
||||
BBlockTransferClusterLengths = params[28]
|
||||
BBlockTransferArrangeOrder = params[29]
|
||||
BBlockTransferSrcVectorDim = params[30]
|
||||
BBlockTransferSrcAccessOrder = params[31]
|
||||
BBlockTransferSrcScalarPerVector = params[32]
|
||||
BBlockTransferDstScalarPerVector_K1 = params[33]
|
||||
BBlockLdsAddExtraM = params[34]
|
||||
CShuffleMXdlPerwave = params[35]
|
||||
CShuffleNXdlPerWavePerShuffle = params[36]
|
||||
CBlockTransferClusterLengths = params[37]
|
||||
CBlockTransferScalarPerVector = params[38]
|
||||
|
||||
|
||||
KBlockPerCu = 1
|
||||
MWarp = int(MPerBlock) // (int(MPerXDL) * int(MXdlPerWave))
|
||||
NWarp = int(NPerBlock) // (int(NPerXDL) * int(NXdlPerWave))
|
||||
KWarp = 1
|
||||
KPerXdl = 16 if MPerXDL == "32" else 32
|
||||
DoubleSMemBuffer = 'false'
|
||||
GemmPipelineVersion = "CK_TILE_PIPELINE_COMPUTE_V3"
|
||||
|
||||
print(MPerBlock, NPerBlock, KPerBlock)
|
||||
|
||||
pipelines = ["CK_TILE_PIPELINE_MEMORY", "CK_TILE_PIPELINE_COMPUTE_V3", "CK_TILE_PIPELINE_COMPUTE_V4"]
|
||||
|
||||
for pipeline in pipelines:
|
||||
DoubleSMemBuffer = 'false' if pipeline != 'CK_TILE_PIPELINE_COMPUTE_V4' else 'true'
|
||||
with open(output_path, 'a') as f:
|
||||
f.write(f'GroupedConvolutionBackwardWeightInvoker<{NDimSpatial}, {ALayout}, {BLayout}, {ELayout}, {ADataType},'
|
||||
f'{BDataType}, {EDataType}, {AElementwiseOp}, {BElementwiseOp}, {CElementwiseOp},'
|
||||
f'{KBlockPerCu}, {MPerBlock}, {NPerBlock}, {KPerBlock}, {MWarp}, {NWarp}, {KWarp},'
|
||||
f'{MPerXDL}, {NPerXDL}, {KPerXdl}, {ABlockTransferSrcScalarPerVector}, {BBlockTransferSrcScalarPerVector},'
|
||||
f'{CBlockTransferScalarPerVector}, {DoubleSMemBuffer}, {pipeline}>,\n')
|
||||
|
||||
|
||||
# print(params[0])
|
||||
|
||||
# # Print each parameter as a separate variable
|
||||
# for i, p in enumerate(params, start=1):
|
||||
# print(f"param_{i} = '{p}'")
|
||||
121
script/convert_old_ck_conv_fwd_to_ck_tile.py
Normal file
121
script/convert_old_ck_conv_fwd_to_ck_tile.py
Normal file
@@ -0,0 +1,121 @@
|
||||
import re
|
||||
|
||||
def extract_template_parameters(template_str):
|
||||
# Extract everything inside the outermost <>
|
||||
match = re.search(r"<(.*)>", template_str, re.DOTALL)
|
||||
if not match:
|
||||
return []
|
||||
|
||||
inside = match.group(1).strip()
|
||||
|
||||
params = []
|
||||
current = []
|
||||
depth = 0 # track nested < >
|
||||
|
||||
for char in inside:
|
||||
if char == '<':
|
||||
depth += 1
|
||||
current.append(char)
|
||||
elif char == '>':
|
||||
depth -= 1
|
||||
current.append(char)
|
||||
elif char == ',' and depth == 0:
|
||||
param = ''.join(current).strip()
|
||||
if param:
|
||||
params.append(param)
|
||||
current = []
|
||||
else:
|
||||
current.append(char)
|
||||
|
||||
# Append last parameter if any
|
||||
if current:
|
||||
params.append(''.join(current).strip())
|
||||
|
||||
return params
|
||||
|
||||
|
||||
input_path = "inputkernel.txt"
|
||||
output_path = "outputkernel.txt"
|
||||
|
||||
with open(input_path, 'r') as f:
|
||||
lines = f.readlines()
|
||||
|
||||
for line in lines:
|
||||
|
||||
# Example usage
|
||||
#input_str = " DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle<NDimSpatial,ALayout,BLayout, DsLayout,ELayout, F16, F16, F32, F16, DsDataTypes, F16, PassThrough, PassThrough, OutElementOp, ConvSpec, GemmMNKPadding, 1, 64, 32, 64, 32, 8, 8, 32, 32, 1, 2, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, S<4, 16, 1>, S<1, 0, 2>, S<1, 0, 2>, 2, 8, 8, 1, 1, 1, S<1, 16, 1, 4>, 8>"
|
||||
|
||||
params = extract_template_parameters(line)
|
||||
|
||||
NDimSpatial = params[0]
|
||||
ALayout = params[1]
|
||||
BLayout = params[2]
|
||||
DsLayout = params[3]
|
||||
ELayout = params[4]
|
||||
ADataType = params[5]
|
||||
BDataType = params[6]
|
||||
AccDataType = params[7]
|
||||
CShuffleDataType= params[8]
|
||||
DsDataTypes = params[9]
|
||||
EDataType = params[10]
|
||||
AElementwiseOp = params[11]
|
||||
BElementwiseOp = params[12]
|
||||
CElementwiseOp = "PassThrough"#params[13]
|
||||
ConvFwdSpec = params[14]
|
||||
GemmSpec = params[15]
|
||||
NummGemmKPref = params[16]
|
||||
BlockSize = params[17]
|
||||
MPerBlock = params[18]
|
||||
NPerBlock = params[19]
|
||||
KPerBlock = params[20]
|
||||
AK1 = params[21]
|
||||
BK1 = params[22]
|
||||
MPerXDL = params[23]
|
||||
NPerXDL = params[24]
|
||||
MXdlPerWave = params[25]
|
||||
NXdlPerWave = params[26]
|
||||
ABlockTransferClusterLengths = params[27]
|
||||
ABlockTransferArrangeOrder = params[28]
|
||||
ABlockTransferSrcAccessOrder = params[29]
|
||||
ABlockTransferSrcVectorDim = params[30]
|
||||
ABlockTransferSrcScalarPerVector = params[31]
|
||||
ABlockTransferDstScalarPerVector_K1 = params[32]
|
||||
ABlockLdsAddExtraM = params[33]
|
||||
BBlockTransferClusterLengths = params[34]
|
||||
BBlockTransferArrangeOrder = params[35]
|
||||
BBlockTransferSrcVectorDim = params[36]
|
||||
BBlockTransferSrcAccessOrder = params[37]
|
||||
BBlockTransferSrcScalarPerVector = params[38]
|
||||
BBlockTransferDstScalarPerVector_K1 = params[39]
|
||||
BBlockLdsAddExtraM = params[40]
|
||||
CShuffleMXdlPerwave = params[41]
|
||||
CShuffleNXdlPerWavePerShuffle = params[42]
|
||||
CBlockTransferClusterLengths = params[43]
|
||||
CBlockTransferScalarPerVector = params[44]
|
||||
|
||||
|
||||
KBlockPerCu = 1
|
||||
MWarp = int(MPerBlock) // (int(MPerXDL) * int(MXdlPerWave))
|
||||
NWarp = int(NPerBlock) // (int(NPerXDL) * int(NXdlPerWave))
|
||||
KWarp = 1
|
||||
KPerXdl = 16 if MPerXDL == "32" else 32
|
||||
DoubleSMemBuffer = 'false'
|
||||
GemmPipelineVersion = "CK_TILE_PIPELINE_COMPUTE_V3"
|
||||
|
||||
pipelines = ["CK_TILE_PIPELINE_MEMORY", "CK_TILE_PIPELINE_COMPUTE_V3", "CK_TILE_PIPELINE_COMPUTE_V4"]
|
||||
|
||||
for pipeline in pipelines:
|
||||
DoubleSMemBuffer = 'false' if pipeline != 'CK_TILE_PIPELINE_COMPUTE_V4' else 'true'
|
||||
with open(output_path, 'a') as f:
|
||||
f.write(f'GroupedConvolutionForwardInvoker<{NDimSpatial}, {ALayout}, {BLayout}, {ELayout}, {ADataType},'
|
||||
f'{BDataType}, {EDataType}, {AElementwiseOp}, {BElementwiseOp}, {CElementwiseOp},'
|
||||
f'{KBlockPerCu}, {MPerBlock}, {NPerBlock}, {KPerBlock}, {MWarp}, {NWarp}, {KWarp},'
|
||||
f'{MPerXDL}, {NPerXDL}, {KPerXdl}, {ABlockTransferSrcScalarPerVector}, {BBlockTransferSrcScalarPerVector},'
|
||||
f'{CBlockTransferScalarPerVector}, {DoubleSMemBuffer}, {pipeline}>,\n')
|
||||
|
||||
|
||||
# print(params[0])
|
||||
|
||||
# # Print each parameter as a separate variable
|
||||
# for i, p in enumerate(params, start=1):
|
||||
# print(f"param_{i} = '{p}'")
|
||||
Reference in New Issue
Block a user