ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-16 08:44:55 +00:00
Code Issues Packages Projects Releases Wiki Activity

Refactor algorithm specialization and GEMM pipeline definitions.

Browse Source
This commit is contained in:
Ville Pietilä
2026-01-09 10:05:54 -05:00
parent f74e034ae9
commit 63fc27b0b1
14 changed files with 120 additions and 150 deletions
Download Patch File Download Diff File Expand all files Collapse all files

36
experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp
View File

@@ -282,16 +282,6 @@ concept SpecifiesNumGroupsToMerge = requires {
{ T::num_conv_groups_to_merge } -> SizeType;
};
template <typename T>
concept SpecifiesLoopScheduler = requires {
{ T::loop_scheduler } -> std::convertible_to<PipelineScheduler>;
};
template <typename T>
concept SpecifiesGenericInstance = !requires {
{ T::specialization };
};
template <typename T>
concept SpecifiesTransposeTransfer = requires {
{ T::max_transpose_transfer_src_scalar_per_vector } -> SizeType;
@@ -308,10 +298,6 @@ template <typename T>
concept TransposeTransferWellDefinedIfProvided =
!HasTransposeTransfer<T> || SpecifiesTransposeTransfer<T>;
template <typename T>
concept SpecifiesGemmBatchOptions = requires {
{ T::num_conv_groups_to_merge } -> SizeType;
};
/******************************************** */
/* Algorithm specialization concepts */
@@ -319,25 +305,39 @@ concept SpecifiesGemmBatchOptions = requires {
template <typename T>
concept SpecifiesLargeTensorSupport = requires {
{ T::specialization } -> std::convertible_to<ConvAlgorithmSpecialization>;
requires T::specialization == ConvAlgorithmSpecialization::LARGE_TENSOR;
requires !!(T::specialization & ConvAlgorithmSpecialization::LARGE_TENSOR);
};
template <typename T>
concept SpecifiesReferenceAlgorithm = requires {
{ T::specialization } -> std::convertible_to<ConvAlgorithmSpecialization>;
requires T::specialization == ConvAlgorithmSpecialization::REFERENCE;
requires !!(T::specialization & ConvAlgorithmSpecialization::REFERENCE);
};
template <typename T>
concept SpecifiesTwoStageSupport = requires {
{ T::specialization } -> std::convertible_to<ConvAlgorithmSpecialization>;
requires T::specialization == ConvAlgorithmSpecialization::TWO_STAGE;
requires !!(T::specialization & ConvAlgorithmSpecialization::TWO_STAGE);
};
template <typename T>
concept SpecifiesMultipleDSupport = requires {
{ T::specialization } -> std::convertible_to<ConvAlgorithmSpecialization>;
requires T::specialization == ConvAlgorithmSpecialization::MULTIPLE_D;
requires !!(T::specialization & ConvAlgorithmSpecialization::MULTIPLE_D);
};
template <typename T>
concept SpecifiesPipelineV3 = requires {
{ T::specialization } -> std::convertible_to<ConvAlgorithmSpecialization>;
requires !!(T::specialization & ConvAlgorithmSpecialization::PIPELINE_V3);
};
template <typename T>
concept SpecifiesGenericInstance = !requires {
{ T::specialization };
} || requires {
{ T::specialization } -> std::convertible_to<ConvAlgorithmSpecialization>;
requires !!(T::specialization == ConvAlgorithmSpecialization::NONE);
};
template <typename T>

25
experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp
View File

@@ -13,8 +13,7 @@ concept FwdXdlAlgorithmBase =
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesThreadDistribution3D<T> &&
SpecifiesLdsTransfer<T> && SpecifiesThreadClusterAccessOrder<T> &&
SpecifiesSourceAccessOrder<T> && SpecifiesWarpGemm<T> &&
SpecifiesFwdConvSpecialization<T> && SpecifiesGemmSpecialization<T> &&
SpecifiesNumPrefetchStages<T> && SpecifiesNumGroupsToMerge<T> && SpecifiesLoopScheduler<T> &&
SpecifiesFwdConvSpecialization<T> && SpecifiesGemmPipeline<T> &&
SpecifiesXdl<T>;
template <typename T>
@@ -29,7 +28,8 @@ concept BwdXdlV3AlgorithmBase =
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesThreadDistribution3D<T> &&
SpecifiesLdsTransfer<T> && SpecifiesThreadClusterAccessOrder<T> &&
SpecifiesSourceAccessOrder<T> && SpecifiesWarpGemm<T> &&
SpecifiesBwdWeightConvSpecialization<T> && SpecifiesGemmPipeline<T> && SpecifiesXdl<T>;
SpecifiesBwdWeightConvSpecialization<T> && SpecifiesGemmPipeline<T> && SpecifiesXdl<T> &&
SpecifiesPipelineV3<T>;
template <typename T>
concept BwdWmmaAlgorithmBase =
@@ -43,7 +43,8 @@ concept BwdWmmaV3AlgorithmBase =
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesThreadDistribution3D<T> &&
SpecifiesLdsTransfer<T> && SpecifiesThreadClusterAccessOrder<T> &&
SpecifiesSourceAccessOrder<T> && SpecifiesWarpGemm<T> &&
SpecifiesBwdWeightConvSpecialization<T> && SpecifiesGemmPipeline<T> && SpecifiesWmma<T>;
SpecifiesBwdWeightConvSpecialization<T> && SpecifiesGemmPipeline<T> && SpecifiesWmma<T> &&
SpecifiesPipelineV3<T>;
// Reference algorithm concept
template <typename T>
@@ -67,7 +68,8 @@ concept FwdXdlV3Algorithm =
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesThreadDistribution3D<T> &&
SpecifiesLdsTransfer<T> && SpecifiesThreadClusterAccessOrder<T> &&
SpecifiesSourceAccessOrder<T> && SpecifiesWarpGemm<T> &&
SpecifiesFwdConvSpecialization<T> && SpecifiesGemmSpecialization<T> && SpecifiesGemmPipeline<T> && SpecifiesXdl<T>;
SpecifiesFwdConvSpecialization<T> && SpecifiesGemmPipeline<T> && SpecifiesXdl<T> &&
SpecifiesPipelineV3<T>;
// FWD WMMA algorithm concepts
template <typename T>
@@ -75,8 +77,7 @@ concept FwdWmmaAlgorithm =
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesThreadDistribution3D<T> &&
SpecifiesLdsTransfer<T> && SpecifiesThreadClusterAccessOrder<T> &&
SpecifiesSourceAccessOrder<T> && SpecifiesWarpGemm<T> &&
SpecifiesFwdConvSpecialization<T> && SpecifiesGemmSpecialization<T> &&
SpecifiesNumPrefetchStages<T> && SpecifiesLoopScheduler<T> && SpecifiesGemmPipeline<T> && SpecifiesWmma<T>;
SpecifiesFwdConvSpecialization<T> && SpecifiesGemmPipeline<T> && SpecifiesWmma<T>;
// FWD DL algorithms
template <typename T>
@@ -94,17 +95,15 @@ template <typename T>
concept BwdMultiDXdlAlgorithm = BwdXdlAlgorithmBase<T> && SpecifiesMultipleDSupport<T>;
template <typename T>
concept BwdXdlV3Algorithm = BwdXdlV3AlgorithmBase<T> && SpecifiesGenericInstance<T>;
concept BwdXdlV3Algorithm = BwdXdlV3AlgorithmBase<T>;
template <typename T>
concept BwdTwoStageXdlAlgorithm = BwdXdlV3AlgorithmBase<T> && SpecifiesTransposeTransfer<T> &&
SpecifiesGemmBatchOptions<T> && SpecifiesTwoStageSupport<T>;
concept BwdTwoStageXdlAlgorithm = BwdXdlV3AlgorithmBase<T> && SpecifiesTransposeTransfer<T> && SpecifiesTwoStageSupport<T>;
// BWD weight WMMA algorithm concepts
template <typename T>
concept BwdWmmaAlgorithm =
BwdWmmaAlgorithmBase<T> && SpecifiesNumPrefetchStages<T> && SpecifiesLoopScheduler<T> &&
SpecifiesGemmPipeline<T> && SpecifiesGenericInstance<T>;
BwdWmmaAlgorithmBase<T> && SpecifiesNumPrefetchStages<T> && SpecifiesGemmPipeline<T> && SpecifiesGenericInstance<T>;
template <typename T>
concept BwdMultiDWmmaV3Algorithm = BwdWmmaV3AlgorithmBase<T> && SpecifiesMultipleDSupport<T>;
@@ -115,7 +114,7 @@ concept BwdWmmaV3Algorithm =
template <typename T>
concept BwdTwoStageWmmaV3Algorithm = BwdWmmaV3AlgorithmBase<T> && SpecifiesTransposeTransfer<T> &&
SpecifiesGemmBatchOptions<T> && SpecifiesTwoStageSupport<T>;
SpecifiesTwoStageSupport<T>;
// BWD weight DL algorithms
template <typename T>

2
experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp
View File

@@ -105,7 +105,7 @@ struct ConvBwdWeightTwoStageXdlFactory
C_BLOCK_TRANSFER.scalar_per_vector,
BLOCK_GEMM.scheduler,
BLOCK_GEMM.pipeline_version,
ALGORITHM.num_conv_groups_to_merge,
BLOCK_GEMM.num_conv_groups_to_merge,
typename Types::OutComputeType,
typename Types::InComputeType,
ALGORITHM.max_transpose_transfer_src_scalar_per_vector,

8
experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tuning_params.hpp
View File

@@ -31,6 +31,8 @@ ConvSpec(CONV_ENUM, GEMM_ENUM) -> ConvSpec<CONV_ENUM>;
struct BlockGemmSpec
{
size_t num_conv_groups_to_merge{1};
size_t num_gemm_k_prefetch_stages{1};
ck::BlockGemmPipelineVersion pipeline_version;
ck::BlockGemmPipelineScheduler scheduler;
};
@@ -63,7 +65,11 @@ consteval BlockGemmSpec SetBlockGemm()
default: throw "Unknown PipelineVersion";
}
return BlockGemmSpec{.pipeline_version = version, .scheduler = scheduler};
return BlockGemmSpec{
.num_conv_groups_to_merge = BG.num_conv_groups_to_merge,
.num_gemm_k_prefetch_stages = BG.num_gemm_k_prefetch_stages,
.pipeline_version = version,
.scheduler = scheduler};
}
template <ConvAlgorithmDescriptor auto ALGORITHM>

31
experimental/builder/include/ck_tile/builder/types.hpp
View File

@@ -232,12 +232,35 @@ enum class PipelineScheduler
enum class ConvAlgorithmSpecialization
{
LARGE_TENSOR,
REFERENCE, // GPU reference implementation for validation,
TWO_STAGE,
MULTIPLE_D
NONE = 0,
LARGE_TENSOR = 1 << 0,
REFERENCE = 1 << 1, // GPU reference implementation for validation,
TWO_STAGE = 1 << 2,
MULTIPLE_D = 1 << 3,
PIPELINE_V3 = 1 << 4
};
constexpr ConvAlgorithmSpecialization operator|(ConvAlgorithmSpecialization lhs,
ConvAlgorithmSpecialization rhs)
{
using T = std::underlying_type_t<ConvAlgorithmSpecialization>;
return static_cast<ConvAlgorithmSpecialization>(static_cast<T>(lhs) | static_cast<T>(rhs));
}
constexpr ConvAlgorithmSpecialization operator&(ConvAlgorithmSpecialization lhs,
ConvAlgorithmSpecialization rhs)
{
using T = std::underlying_type_t<ConvAlgorithmSpecialization>;
return static_cast<ConvAlgorithmSpecialization>(static_cast<T>(lhs) & static_cast<T>(rhs));
}
// Enable direct boolean conversion for flag checks
constexpr bool operator!(ConvAlgorithmSpecialization spec)
{
using T = std::underlying_type_t<ConvAlgorithmSpecialization>;
return static_cast<T>(spec) == 0;
}
enum class MatrixInstructionType
{
XDL,

2
experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_two_stage_xdl_cshuffle.cpp
View File

@@ -19,7 +19,7 @@ constexpr auto SIGNATURE =
.weight = {.config = {.layout = ckb::TensorLayout::GKYXC}},
.output = {.config = {.layout = ckb::TensorLayout::GNHWK}}};
constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_CShuffle{}
constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_CShuffle_V3{}
.with_thread_block(cku::ThreadBlock_64_32x32x32)
.with_gemm_config(cku::BwdGemmParams_Xdl_1x1_per_wave)
.with_transfer(cku::BwdTransfer_4x8x1_4x16x1_v3)

3
experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle.cpp
View File

@@ -24,8 +24,7 @@ constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CS
.with_gemm_config(cku::GemmParams_Wmma_16x16_2x1_per_wave)
.with_transfer(cku::BwdTransfer_4x8x1_4x16x1_v3)
.with_bwd_specialization(ckb::ConvSpecialization::DEFAULT)
.with_prefetch_config(1, ckb::PipelineScheduler::DEFAULT)
.with_gemm_pipeline(ckb::PipelineVersion::V1);
.with_gemm_pipeline(ckb::PipelineVersion::V1, ckb::PipelineScheduler::DEFAULT);
using Builder = ckb::ConvBuilder<SIGNATURE, ALGORITHM>;
using Instance = Builder::Instance;

1
experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_fp16.cpp
View File

@@ -31,7 +31,6 @@ TEST(FwdConvInstances,
.with_gemm_config(FwdGemmParams_Xdl_2x1_per_wave)
.with_transfer(Transfer_4x16x1)
.with_fwd_specializations(ConvSpecialization::DEFAULT, GemmSpecialization::MNKPadding)
.with_prefetch_config(1, PipelineScheduler::DEFAULT)
.with_num_conv_groups_to_merge(2);
using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;

4
experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_i8.cpp
View File

@@ -33,9 +33,9 @@ TEST(FwdConvInstances,
.with_gemm_config(GemmParams_Wmma_2x1_per_wave)
.with_transfer(Transfer_4x32x1)
.with_fwd_specializations(ConvSpecialization::DEFAULT, GemmSpecialization::MNKPadding)
.with_prefetch_config(1, PipelineScheduler::INTRAWAVE)
.with_num_conv_groups_to_merge(2)
.with_gemm_pipeline(PipelineVersion::V1);
.with_with_num_gemm_k_prefetch_stages(3)
.with_gemm_pipeline(PipelineVersion::V1, PipelineScheduler::INTRAWAVE);
using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;

1
experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_bf16_scaleadd_relu.cpp
View File

@@ -36,7 +36,6 @@ TEST(FwdConvInstances,
.with_gemm_config(FwdGemmParams_Xdl_2x1_per_wave)
.with_transfer(Transfer_4x16x1)
.with_fwd_specializations(ConvSpecialization::DEFAULT, GemmSpecialization::MNKPadding)
.with_prefetch_config(1, PipelineScheduler::DEFAULT)
.with_num_conv_groups_to_merge(1);
using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;

1
experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp8.cpp
View File

@@ -31,7 +31,6 @@ TEST(FwdConvInstances,
.with_gemm_config(FwdGemmParams_Xdl_4x2_per_wave)
.with_transfer(Transfer_4x64x1_fp8)
.with_fwd_specializations(ConvSpecialization::DEFAULT, GemmSpecialization::MNKPadding)
.with_prefetch_config(1, PipelineScheduler::DEFAULT)
.with_num_conv_groups_to_merge(1);
using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;

2
experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_large_tensor_fp16.cpp
View File

@@ -30,7 +30,6 @@ TEST(FwdConvInstances,
.with_gemm_config(FwdGemmParams_Xdl_2x1_per_wave)
.with_transfer(Transfer_4x16x1)
.with_fwd_specializations(ConvSpecialization::DEFAULT, GemmSpecialization::MNKPadding)
.with_prefetch_config(1, PipelineScheduler::DEFAULT)
.with_num_conv_groups_to_merge(1);
using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;
@@ -67,7 +66,6 @@ TEST(
.with_transfer(Transfer_4x16x1)
.with_fwd_specializations(ConvSpecialization::FILTER_1X1_PAD0,
GemmSpecialization::MNKPadding)
.with_prefetch_config(1, PipelineScheduler::DEFAULT)
.with_num_conv_groups_to_merge(1);
using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;

116
experimental/builder/test/impl/conv_algorithm_types.hpp
View File

@@ -40,6 +40,8 @@ static_assert(ckb::WarpGemmDescriptor<WarpGemmParams>);
struct GemmPipeline
{
size_t num_gemm_k_prefetch_stages{1};
size_t num_conv_groups_to_merge{1};
PipelineVersion pipeline_version;
PipelineScheduler scheduler{PipelineScheduler::DEFAULT};
};
@@ -195,23 +197,12 @@ struct ConvSpecializationBwdWeight_
ConvSpecialization bwd_weight_specialization;
};
struct Prefetch_
{
size_t num_gemm_k_prefetch_stages;
PipelineScheduler loop_scheduler;
};
struct TransposeParams_
{
size_t max_transpose_transfer_src_scalar_per_vector{1};
size_t max_transpose_transfer_dst_scalar_per_vector{1};
};
struct GemmBatchOptions_
{
size_t num_conv_groups_to_merge{1};
};
struct GemmPipeline_
{
GemmPipeline gemm_pipeline;
@@ -241,22 +232,10 @@ struct DlTransfer_
DlTransfer<Dim> transfer;
};
struct TwoStageSpecialization_
template <ConvAlgorithmSpecialization Specialization = ConvAlgorithmSpecialization::NONE>
struct AlgorithmSpecialization_
{
static constexpr ConvAlgorithmSpecialization specialization =
ConvAlgorithmSpecialization::TWO_STAGE;
};
struct MultipleDSpecialization_
{
static constexpr ConvAlgorithmSpecialization specialization =
ConvAlgorithmSpecialization::MULTIPLE_D;
};
struct LargeTensorSpecialization_
{
static constexpr ConvAlgorithmSpecialization specialization =
ConvAlgorithmSpecialization::LARGE_TENSOR;
static constexpr ConvAlgorithmSpecialization specialization = Specialization;
};
// Specify thread block dimensions for a GEMM (CK Tile).
@@ -378,15 +357,6 @@ struct ConvAlgorithmTemplate : Components...
return result;
}
constexpr auto with_prefetch_config(size_t k_prefetch_stages, PipelineScheduler scheduler) const
{
static_assert(std::is_base_of_v<Prefetch_, ConvAlgorithmTemplate>);
auto result = *this;
result.num_gemm_k_prefetch_stages = k_prefetch_stages;
result.loop_scheduler = scheduler;
return result;
}
constexpr auto with_transpose_params(size_t max_src_scalar_per_vector,
size_t max_dst_scalar_per_vector) const
{
@@ -399,9 +369,17 @@ struct ConvAlgorithmTemplate : Components...
constexpr auto with_num_conv_groups_to_merge(size_t num_groups_to_merge) const
{
static_assert(std::is_base_of_v<GemmBatchOptions_, ConvAlgorithmTemplate>);
static_assert(std::is_base_of_v<GemmPipeline_, ConvAlgorithmTemplate>);
auto result = *this;
result.num_conv_groups_to_merge = num_groups_to_merge;
result.gemm_pipeline.num_conv_groups_to_merge = num_groups_to_merge;
return result;
}
constexpr auto with_num_gemm_k_prefetch_stages(size_t num_prefetch_stages) const
{
static_assert(std::is_base_of_v<GemmPipeline_, ConvAlgorithmTemplate>);
auto result = *this;
result.gemm_pipeline.num_gemm_k_prefetch_stages = num_prefetch_stages;
return result;
}
@@ -422,6 +400,15 @@ struct ConvAlgorithmTemplate : Components...
return result;
}
constexpr auto with_gemm_pipeline(const PipelineVersion plv, const PipelineScheduler sch) const
{
static_assert(std::is_base_of_v<GemmPipeline_, ConvAlgorithmTemplate>);
auto result = *this;
result.gemm_pipeline.pipeline_version = plv;
result.gemm_pipeline.scheduler = sch;
return result;
}
template <typename TC>
constexpr auto with_dl_thread_config(const TC& tc) const
{
@@ -498,29 +485,24 @@ struct ConvAlgorithmTemplate : Components...
// Fwd algorithm types
using ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle =
using enum ckb::ConvAlgorithmSpecialization;
// Covers both XDL and WMMA variants for generic fwd convolution
using ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_CShuffle =
ConvAlgorithmTemplate<ThreadBlock_,
WarpGemm_,
InputOutputTileTransfer_<>,
ConvSpecializationFwd_,
Prefetch_,
GemmBatchOptions_>;
GemmPipeline_,
AlgorithmSpecialization_<>>;
using ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 =
ConvAlgorithmTemplate<ThreadBlock_,
WarpGemm_,
InputOutputTileTransfer_<>,
ConvSpecializationFwd_,
GemmPipeline_>;
using ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle =
ConvAlgorithmTemplate<ThreadBlock_,
WarpGemm_,
InputOutputTileTransfer_<>,
ConvSpecializationFwd_,
GemmPipeline_,
Prefetch_,
GemmBatchOptions_>;
AlgorithmSpecialization_<PIPELINE_V3>>;
using ConvAlgorithm_DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK =
ConvAlgorithmTemplate<ThreadBlock_,
@@ -534,9 +516,8 @@ using ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Xdl_CShuffle_Large_Tensor =
WarpGemm_,
InputOutputTileTransfer_<>,
ConvSpecializationFwd_,
Prefetch_,
GemmBatchOptions_,
LargeTensorSpecialization_>;
GemmPipeline_,
AlgorithmSpecialization_<LARGE_TENSOR | MULTIPLE_D>>;
// CK Tile algorithm
using ConvAlgorithm_Tile_GroupedConvolutionKernel = ConvAlgorithmTemplate<TileThreadBlock_,
@@ -546,13 +527,7 @@ using ConvAlgorithm_Tile_GroupedConvolutionKernel = ConvAlgorithmTemplate<TileTh
TileOptimizations_>;
// Reference algorithm descriptor - for GPU reference validation
// This is a simple algorithm that requires no complex configuration,
// just a specialization marker to identify it as a reference implementation.
struct ConvAlgorithm_Reference
{
static constexpr auto specialization = ckb::ConvAlgorithmSpecialization::REFERENCE;
// GPU reference uses simple algorithm, no tile configuration needed
};
using ConvAlgorithm_Reference = ConvAlgorithmTemplate<AlgorithmSpecialization_<REFERENCE>>;
// Bwd weight algorithm types
using ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuffle =
@@ -560,7 +535,8 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuffle =
WarpGemm_,
InputOutputTileTransfer_<4>,
ConvSpecializationBwdWeight_,
TransposeParams_>;
TransposeParams_,
AlgorithmSpecialization_<>>;
using ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle =
ConvAlgorithmTemplate<ThreadBlock_,
@@ -568,25 +544,25 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle =
InputOutputTileTransfer_<>,
ConvSpecializationBwdWeight_,
GemmPipeline_,
Prefetch_>;
AlgorithmSpecialization_<>>;
// Covers both XDL and WMMA variants
using ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_CShuffle =
using ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_CShuffle_V3 =
ConvAlgorithmTemplate<ThreadBlock_,
WarpGemm_,
InputOutputTileTransfer_<>,
ConvSpecializationBwdWeight_,
GemmPipeline_,
TransposeParams_,
GemmBatchOptions_,
TwoStageSpecialization_>;
AlgorithmSpecialization_<TWO_STAGE | PIPELINE_V3>>;
using ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuffle_V3 =
ConvAlgorithmTemplate<ThreadBlock_,
WarpGemm_,
InputOutputTileTransfer_<>,
ConvSpecializationBwdWeight_,
GemmPipeline_>;
GemmPipeline_,
AlgorithmSpecialization_<PIPELINE_V3>>;
using ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle_V3 =
ConvAlgorithmTemplate<ThreadBlock_,
@@ -594,7 +570,8 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle_V3 =
InputOutputTileTransfer_<>,
ConvSpecializationBwdWeight_,
GemmPipeline_,
TransposeParams_>;
TransposeParams_,
AlgorithmSpecialization_<PIPELINE_V3>>;
using ConvAlgorithm_DeviceGroupedConvBwdWeight_Dl =
ConvAlgorithmTemplate<ThreadBlock_,
@@ -608,7 +585,7 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle =
WarpGemm_,
InputOutputTileTransfer_<4>,
ConvSpecializationBwdWeight_,
MultipleDSpecialization_>;
AlgorithmSpecialization_<MULTIPLE_D>>;
using ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_Wmma_CShuffle_V3 =
ConvAlgorithmTemplate<ThreadBlock_,
@@ -617,8 +594,7 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_Wmma_CShuffle_V3 =
ConvSpecializationBwdWeight_,
GemmPipeline_,
TransposeParams_,
GemmBatchOptions_,
TwoStageSpecialization_>;
AlgorithmSpecialization_<TWO_STAGE | PIPELINE_V3>>;
using ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffle_V3 =
ConvAlgorithmTemplate<ThreadBlock_,
@@ -626,6 +602,6 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffle_V3 =
InputOutputTileTransfer_<>,
ConvSpecializationBwdWeight_,
GemmPipeline_,
MultipleDSpecialization_>;
AlgorithmSpecialization_<MULTIPLE_D | PIPELINE_V3>>;
} // namespace ck_tile::builder::test

38
experimental/builder/test/utils/conv_algorithm_type_utils.hpp
View File

@@ -98,7 +98,7 @@ template <>
inline std::string to_string<GemmPipeline>(GemmPipeline t)
{
std::ostringstream oss;
oss << to_string(t.scheduler) << "," << to_string(t.pipeline_version);
oss << t.num_gemm_k_prefetch_stages << "," << t.num_conv_groups_to_merge << "," << to_string(t.scheduler) << "," << to_string(t.pipeline_version);
return oss.str();
}
@@ -281,14 +281,6 @@ inline std::string to_string<ConvSpecializationBwdWeight_>(ConvSpecializationBwd
return oss.str();
}
template <>
inline std::string to_string<Prefetch_>(Prefetch_ t)
{
std::ostringstream oss;
oss << t.num_gemm_k_prefetch_stages << "," << to_string(t.loop_scheduler);
return oss.str();
}
template <>
inline std::string to_string<GemmPipeline_>(GemmPipeline_ t)
{
@@ -322,28 +314,8 @@ inline std::string to_string<DlTransfer_<5>>(DlTransfer_<5> t)
// Template specializations for algorithm types
template <>
inline std::string to_string<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle>(
ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle t)
{
std::ostringstream oss;
oss << to_string(static_cast<ThreadBlock_>(t)) << "," << to_string(static_cast<WarpGemm_>(t))
<< "," << to_string(static_cast<InputOutputTileTransfer_<>>(t));
return oss.str();
}
template <>
inline std::string to_string<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3>(
ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 t)
{
std::ostringstream oss;
oss << to_string(static_cast<ThreadBlock_>(t)) << "," << to_string(static_cast<WarpGemm_>(t))
<< "," << to_string(static_cast<InputOutputTileTransfer_<>>(t));
return oss.str();
}
template <>
inline std::string to_string<ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle>(
ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle t)
inline std::string to_string<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_CShuffle>(
ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_CShuffle t)
{
std::ostringstream oss;
oss << to_string(static_cast<ThreadBlock_>(t)) << "," << to_string(static_cast<WarpGemm_>(t))
@@ -425,8 +397,8 @@ inline std::string to_string<ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_W
// Covers both XDL and WMMA versions
template <>
inline std::string to_string<ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_CShuffle>(
ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_CShuffle t)
inline std::string to_string<ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_CShuffle_V3>(
ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_CShuffle_V3 t)
{
std::ostringstream oss;
oss << to_string(static_cast<ThreadBlock_>(t)) << "," << to_string(static_cast<WarpGemm_>(t))