diff --git a/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp b/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp index d036ab7ec6..946b485bdb 100644 --- a/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp +++ b/experimental/builder/include/ck_tile/builder/conv_algorithm_concepts.hpp @@ -282,16 +282,6 @@ concept SpecifiesNumGroupsToMerge = requires { { T::num_conv_groups_to_merge } -> SizeType; }; -template -concept SpecifiesLoopScheduler = requires { - { T::loop_scheduler } -> std::convertible_to; -}; - -template -concept SpecifiesGenericInstance = !requires { - { T::specialization }; -}; - template concept SpecifiesTransposeTransfer = requires { { T::max_transpose_transfer_src_scalar_per_vector } -> SizeType; @@ -308,10 +298,6 @@ template concept TransposeTransferWellDefinedIfProvided = !HasTransposeTransfer || SpecifiesTransposeTransfer; -template -concept SpecifiesGemmBatchOptions = requires { - { T::num_conv_groups_to_merge } -> SizeType; -}; /******************************************** */ /* Algorithm specialization concepts */ @@ -319,25 +305,39 @@ concept SpecifiesGemmBatchOptions = requires { template concept SpecifiesLargeTensorSupport = requires { { T::specialization } -> std::convertible_to; - requires T::specialization == ConvAlgorithmSpecialization::LARGE_TENSOR; + requires !!(T::specialization & ConvAlgorithmSpecialization::LARGE_TENSOR); }; template concept SpecifiesReferenceAlgorithm = requires { { T::specialization } -> std::convertible_to; - requires T::specialization == ConvAlgorithmSpecialization::REFERENCE; + requires !!(T::specialization & ConvAlgorithmSpecialization::REFERENCE); }; template concept SpecifiesTwoStageSupport = requires { { T::specialization } -> std::convertible_to; - requires T::specialization == ConvAlgorithmSpecialization::TWO_STAGE; + requires !!(T::specialization & ConvAlgorithmSpecialization::TWO_STAGE); }; template concept SpecifiesMultipleDSupport = requires { { T::specialization } -> std::convertible_to; - requires T::specialization == ConvAlgorithmSpecialization::MULTIPLE_D; + requires !!(T::specialization & ConvAlgorithmSpecialization::MULTIPLE_D); +}; + +template +concept SpecifiesPipelineV3 = requires { + { T::specialization } -> std::convertible_to; + requires !!(T::specialization & ConvAlgorithmSpecialization::PIPELINE_V3); +}; + +template +concept SpecifiesGenericInstance = !requires { + { T::specialization }; +} || requires { + { T::specialization } -> std::convertible_to; + requires !!(T::specialization == ConvAlgorithmSpecialization::NONE); }; template diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp index ff7e54546d..81bafa158d 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_algorithms.hpp @@ -13,8 +13,7 @@ concept FwdXdlAlgorithmBase = ConvAlgorithmDescriptor && SpecifiesThreadBlock && SpecifiesThreadDistribution3D && SpecifiesLdsTransfer && SpecifiesThreadClusterAccessOrder && SpecifiesSourceAccessOrder && SpecifiesWarpGemm && - SpecifiesFwdConvSpecialization && SpecifiesGemmSpecialization && - SpecifiesNumPrefetchStages && SpecifiesNumGroupsToMerge && SpecifiesLoopScheduler && + SpecifiesFwdConvSpecialization && SpecifiesGemmPipeline && SpecifiesXdl; template @@ -29,7 +28,8 @@ concept BwdXdlV3AlgorithmBase = ConvAlgorithmDescriptor && SpecifiesThreadBlock && SpecifiesThreadDistribution3D && SpecifiesLdsTransfer && SpecifiesThreadClusterAccessOrder && SpecifiesSourceAccessOrder && SpecifiesWarpGemm && - SpecifiesBwdWeightConvSpecialization && SpecifiesGemmPipeline && SpecifiesXdl; + SpecifiesBwdWeightConvSpecialization && SpecifiesGemmPipeline && SpecifiesXdl && + SpecifiesPipelineV3; template concept BwdWmmaAlgorithmBase = @@ -43,7 +43,8 @@ concept BwdWmmaV3AlgorithmBase = ConvAlgorithmDescriptor && SpecifiesThreadBlock && SpecifiesThreadDistribution3D && SpecifiesLdsTransfer && SpecifiesThreadClusterAccessOrder && SpecifiesSourceAccessOrder && SpecifiesWarpGemm && - SpecifiesBwdWeightConvSpecialization && SpecifiesGemmPipeline && SpecifiesWmma; + SpecifiesBwdWeightConvSpecialization && SpecifiesGemmPipeline && SpecifiesWmma && + SpecifiesPipelineV3; // Reference algorithm concept template @@ -67,7 +68,8 @@ concept FwdXdlV3Algorithm = ConvAlgorithmDescriptor && SpecifiesThreadBlock && SpecifiesThreadDistribution3D && SpecifiesLdsTransfer && SpecifiesThreadClusterAccessOrder && SpecifiesSourceAccessOrder && SpecifiesWarpGemm && - SpecifiesFwdConvSpecialization && SpecifiesGemmSpecialization && SpecifiesGemmPipeline && SpecifiesXdl; + SpecifiesFwdConvSpecialization && SpecifiesGemmPipeline && SpecifiesXdl && + SpecifiesPipelineV3; // FWD WMMA algorithm concepts template @@ -75,8 +77,7 @@ concept FwdWmmaAlgorithm = ConvAlgorithmDescriptor && SpecifiesThreadBlock && SpecifiesThreadDistribution3D && SpecifiesLdsTransfer && SpecifiesThreadClusterAccessOrder && SpecifiesSourceAccessOrder && SpecifiesWarpGemm && - SpecifiesFwdConvSpecialization && SpecifiesGemmSpecialization && - SpecifiesNumPrefetchStages && SpecifiesLoopScheduler && SpecifiesGemmPipeline && SpecifiesWmma; + SpecifiesFwdConvSpecialization && SpecifiesGemmPipeline && SpecifiesWmma; // FWD DL algorithms template @@ -94,17 +95,15 @@ template concept BwdMultiDXdlAlgorithm = BwdXdlAlgorithmBase && SpecifiesMultipleDSupport; template -concept BwdXdlV3Algorithm = BwdXdlV3AlgorithmBase && SpecifiesGenericInstance; +concept BwdXdlV3Algorithm = BwdXdlV3AlgorithmBase; template -concept BwdTwoStageXdlAlgorithm = BwdXdlV3AlgorithmBase && SpecifiesTransposeTransfer && - SpecifiesGemmBatchOptions && SpecifiesTwoStageSupport; +concept BwdTwoStageXdlAlgorithm = BwdXdlV3AlgorithmBase && SpecifiesTransposeTransfer && SpecifiesTwoStageSupport; // BWD weight WMMA algorithm concepts template concept BwdWmmaAlgorithm = - BwdWmmaAlgorithmBase && SpecifiesNumPrefetchStages && SpecifiesLoopScheduler && - SpecifiesGemmPipeline && SpecifiesGenericInstance; + BwdWmmaAlgorithmBase && SpecifiesNumPrefetchStages && SpecifiesGemmPipeline && SpecifiesGenericInstance; template concept BwdMultiDWmmaV3Algorithm = BwdWmmaV3AlgorithmBase && SpecifiesMultipleDSupport; @@ -115,7 +114,7 @@ concept BwdWmmaV3Algorithm = template concept BwdTwoStageWmmaV3Algorithm = BwdWmmaV3AlgorithmBase && SpecifiesTransposeTransfer && - SpecifiesGemmBatchOptions && SpecifiesTwoStageSupport; + SpecifiesTwoStageSupport; // BWD weight DL algorithms template diff --git a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp index 5eea36313f..764b061796 100644 --- a/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_two_stage_xdl_factory.hpp @@ -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, diff --git a/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tuning_params.hpp b/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tuning_params.hpp index 29cf3f8513..b64b7d336e 100644 --- a/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tuning_params.hpp +++ b/experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tuning_params.hpp @@ -31,6 +31,8 @@ ConvSpec(CONV_ENUM, GEMM_ENUM) -> ConvSpec; 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 diff --git a/experimental/builder/include/ck_tile/builder/types.hpp b/experimental/builder/include/ck_tile/builder/types.hpp index 91952c21a9..af9cc44115 100644 --- a/experimental/builder/include/ck_tile/builder/types.hpp +++ b/experimental/builder/include/ck_tile/builder/types.hpp @@ -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; + return static_cast(static_cast(lhs) | static_cast(rhs)); +} + +constexpr ConvAlgorithmSpecialization operator&(ConvAlgorithmSpecialization lhs, + ConvAlgorithmSpecialization rhs) +{ + using T = std::underlying_type_t; + return static_cast(static_cast(lhs) & static_cast(rhs)); +} + +// Enable direct boolean conversion for flag checks +constexpr bool operator!(ConvAlgorithmSpecialization spec) +{ + using T = std::underlying_type_t; + return static_cast(spec) == 0; +} + enum class MatrixInstructionType { XDL, diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_two_stage_xdl_cshuffle.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_two_stage_xdl_cshuffle.cpp index bda064c291..57b85fa8c7 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_two_stage_xdl_cshuffle.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_two_stage_xdl_cshuffle.cpp @@ -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) diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle.cpp index a559d3ee47..562599da93 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_bwd_weight_wmma_cshuffle.cpp @@ -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; using Instance = Builder::Instance; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_fp16.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_fp16.cpp index d3ace110c4..abeca068e7 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_fp16.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_fp16.cpp @@ -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; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_i8.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_i8.cpp index 8ee12a46ba..7ef33469a0 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_i8.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_1d_i8.cpp @@ -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; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_bf16_scaleadd_relu.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_bf16_scaleadd_relu.cpp index 23edef5436..395640dcf1 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_bf16_scaleadd_relu.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_bf16_scaleadd_relu.cpp @@ -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; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp8.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp8.cpp index b117e693fe..85a87fc3a2 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp8.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_fp8.cpp @@ -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; diff --git a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_large_tensor_fp16.cpp b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_large_tensor_fp16.cpp index 97bc0a00e5..a4dd6171d8 100644 --- a/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_large_tensor_fp16.cpp +++ b/experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_large_tensor_fp16.cpp @@ -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; @@ -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; diff --git a/experimental/builder/test/impl/conv_algorithm_types.hpp b/experimental/builder/test/impl/conv_algorithm_types.hpp index 2e0c79ec5c..8c1c3f898c 100644 --- a/experimental/builder/test/impl/conv_algorithm_types.hpp +++ b/experimental/builder/test/impl/conv_algorithm_types.hpp @@ -40,6 +40,8 @@ static_assert(ckb::WarpGemmDescriptor); 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 transfer; }; -struct TwoStageSpecialization_ +template +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); - 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); + static_assert(std::is_base_of_v); 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); + 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); + auto result = *this; + result.gemm_pipeline.pipeline_version = plv; + result.gemm_pipeline.scheduler = sch; + return result; + } + template 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, ConvSpecializationFwd_, - Prefetch_, - GemmBatchOptions_>; + GemmPipeline_, + AlgorithmSpecialization_<>>; using ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 = - ConvAlgorithmTemplate, - ConvSpecializationFwd_, - GemmPipeline_>; - -using ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle = ConvAlgorithmTemplate, ConvSpecializationFwd_, GemmPipeline_, - Prefetch_, - GemmBatchOptions_>; + AlgorithmSpecialization_>; using ConvAlgorithm_DeviceGroupedConvFwdDlMultipleD_NHWC_KYXC_NHWK = ConvAlgorithmTemplate, ConvSpecializationFwd_, - Prefetch_, - GemmBatchOptions_, - LargeTensorSpecialization_>; + GemmPipeline_, + AlgorithmSpecialization_>; // CK Tile algorithm using ConvAlgorithm_Tile_GroupedConvolutionKernel = ConvAlgorithmTemplate; // 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>; // 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, ConvSpecializationBwdWeight_, GemmPipeline_, - Prefetch_>; + AlgorithmSpecialization_<>>; // Covers both XDL and WMMA variants -using ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_CShuffle = +using ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_CShuffle_V3 = ConvAlgorithmTemplate, ConvSpecializationBwdWeight_, GemmPipeline_, TransposeParams_, - GemmBatchOptions_, - TwoStageSpecialization_>; + AlgorithmSpecialization_>; using ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuffle_V3 = ConvAlgorithmTemplate, ConvSpecializationBwdWeight_, - GemmPipeline_>; + GemmPipeline_, + AlgorithmSpecialization_>; using ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle_V3 = ConvAlgorithmTemplate, ConvSpecializationBwdWeight_, GemmPipeline_, - TransposeParams_>; + TransposeParams_, + AlgorithmSpecialization_>; using ConvAlgorithm_DeviceGroupedConvBwdWeight_Dl = ConvAlgorithmTemplate, ConvSpecializationBwdWeight_, - MultipleDSpecialization_>; + AlgorithmSpecialization_>; using ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_Wmma_CShuffle_V3 = ConvAlgorithmTemplate; + AlgorithmSpecialization_>; using ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Wmma_CShuffle_V3 = ConvAlgorithmTemplate, ConvSpecializationBwdWeight_, GemmPipeline_, - MultipleDSpecialization_>; + AlgorithmSpecialization_>; } // namespace ck_tile::builder::test diff --git a/experimental/builder/test/utils/conv_algorithm_type_utils.hpp b/experimental/builder/test/utils/conv_algorithm_type_utils.hpp index bad29a65c0..7b1bfc9e55 100644 --- a/experimental/builder/test/utils/conv_algorithm_type_utils.hpp +++ b/experimental/builder/test/utils/conv_algorithm_type_utils.hpp @@ -98,7 +98,7 @@ template <> inline std::string to_string(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(ConvSpecializationBwd return oss.str(); } -template <> -inline std::string to_string(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_ t) { @@ -322,28 +314,8 @@ inline std::string to_string>(DlTransfer_<5> t) // Template specializations for algorithm types template <> -inline std::string to_string( - ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle t) -{ - std::ostringstream oss; - oss << to_string(static_cast(t)) << "," << to_string(static_cast(t)) - << "," << to_string(static_cast>(t)); - return oss.str(); -} - -template <> -inline std::string to_string( - ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3 t) -{ - std::ostringstream oss; - oss << to_string(static_cast(t)) << "," << to_string(static_cast(t)) - << "," << to_string(static_cast>(t)); - return oss.str(); -} - -template <> -inline std::string to_string( - ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Wmma_CShuffle t) +inline std::string to_string( + ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_CShuffle t) { std::ostringstream oss; oss << to_string(static_cast(t)) << "," << to_string(static_cast(t)) @@ -425,8 +397,8 @@ inline std::string to_string -inline std::string to_string( - ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_CShuffle t) +inline std::string to_string( + ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_CShuffle_V3 t) { std::ostringstream oss; oss << to_string(static_cast(t)) << "," << to_string(static_cast(t))