ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-12 02:05:50 +00:00
Code Issues Packages Projects Releases Wiki Activity

clang-format

Browse Source
This commit is contained in:
Ville Pietilä
2026-01-14 09:29:04 -05:00
parent 096592eb99
commit 75d20e08f1
21 changed files with 432 additions and 378 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -64,8 +64,9 @@ concept InputTileThreadClusterDescriptor4D = requires(T t) {
};
template <typename T, size_t ThreadClusterRank>
concept InputTileThreadClusterDescriptor = (ThreadClusterRank == 3 && InputTileThreadClusterDescriptor3D<T>) ||
(ThreadClusterRank == 4 && InputTileThreadClusterDescriptor4D<T>);
concept InputTileThreadClusterDescriptor =
(ThreadClusterRank == 3 && InputTileThreadClusterDescriptor3D<T>) ||
(ThreadClusterRank == 4 && InputTileThreadClusterDescriptor4D<T>);
// Concept for thread cluster dimensions for GEMM output tensor.
template <typename T>
@@ -285,7 +286,6 @@ template <typename T>
concept TransposeTransferWellDefinedIfProvided =
!HasTransposeTransfer<T> || SpecifiesTransposeTransfer<T>;
/******************************************** */
/* Algorithm specialization concepts */
/******************************************** */
@@ -328,14 +328,12 @@ concept SpecifiesGenericInstance = !requires {
};
template <auto Algorithm>
concept SpecifiesXdl = requires {
requires Algorithm.warp_gemm.matrix_instruction == MatrixInstructionType::XDL;
};
concept SpecifiesXdl =
requires { requires Algorithm.warp_gemm.matrix_instruction == MatrixInstructionType::XDL; };
template <auto Algorithm>
concept SpecifiesWmma = requires {
requires Algorithm.warp_gemm.matrix_instruction == MatrixInstructionType::WMMA;
};
concept SpecifiesWmma =
requires { requires Algorithm.warp_gemm.matrix_instruction == MatrixInstructionType::WMMA; };
template <auto Algorithm>
concept SpecifiesValidWarpGemm = SpecifiesXdl<Algorithm> || SpecifiesWmma<Algorithm>;

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

@@ -14,31 +14,32 @@ concept SpecifiesTileTransferParameters =
// Base algorithm concepts
template <typename T, size_t ThreadClusterRank = 3>
concept ConvAlgorithm =
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> &&
SpecifiesTileTransferParameters<T, ThreadClusterRank> &&
SpecifiesWarpGemm<T> && SpecifiesGemmPipeline<T>;;
concept ConvAlgorithm = ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> &&
SpecifiesTileTransferParameters<T, ThreadClusterRank> &&
SpecifiesWarpGemm<T> && SpecifiesGemmPipeline<T>;
;
template <typename T>
concept FwdAlgorithm = ConvAlgorithm<T, 3> && SpecifiesFwdConvSpecialization<T>;
concept FwdAlgorithm = ConvAlgorithm<T, 3> && SpecifiesFwdConvSpecialization<T>;
template <typename T>
concept FwdAlgorithmV3 = FwdAlgorithm<T> && SpecifiesPipelineV3<T>;
template <typename T, size_t ThreadClusterRank = 3>
concept BwdAlgorithm = ConvAlgorithm<T, ThreadClusterRank> && SpecifiesBwdWeightConvSpecialization<T>;
concept BwdAlgorithm =
ConvAlgorithm<T, ThreadClusterRank> && SpecifiesBwdWeightConvSpecialization<T>;
template <typename T>
concept BwdAlgorithmV3 = BwdAlgorithm<T, 3> && SpecifiesPipelineV3<T>;
template <typename T>
concept DlAlgorithm =
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> &&
SpecifiesDlThreadConfig<T> && SpecifiesDlThreadCluster<T> && SpecifiesDlEpilogue<T>;
concept DlAlgorithm =
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesDlThreadConfig<T> &&
SpecifiesDlThreadCluster<T> && SpecifiesDlEpilogue<T>;
template <typename T>
concept FwdDlAlgorithmBase = DlAlgorithm<T> && SpecifiesFwdConvSpecialization<T> &&
SpecifiesDlFwdBlockTransfer<T> && SpecifiesGemmSpecialization<T>;
concept FwdDlAlgorithmBase = DlAlgorithm<T> && SpecifiesFwdConvSpecialization<T> &&
SpecifiesDlFwdBlockTransfer<T> && SpecifiesGemmSpecialization<T>;
template <auto Value>
concept FwdXdlAlgorithmBase = FwdAlgorithm<decltype(Value)> && SpecifiesXdl<Value>;
@@ -57,20 +58,23 @@ concept BwdWmmaV3AlgorithmBase = BwdAlgorithmV3<decltype(Value)> && SpecifiesWmm
// Reference algorithm concept
template <auto Value>
concept ReferenceAlgorithm = ConvAlgorithmDescriptor<decltype(Value)> && SpecifiesReferenceAlgorithm<decltype(Value)>;
concept ReferenceAlgorithm =
ConvAlgorithmDescriptor<decltype(Value)> && SpecifiesReferenceAlgorithm<decltype(Value)>;
// Tile-based algorithm concept
template <auto Value>
concept TileAlgorithm = ConvAlgorithmDescriptor<decltype(Value)> && SpecifiesTileThreadBlock<decltype(Value)> &&
SpecifiesTileTransfer<decltype(Value)> && SpecifiesTileConvSpecialization<decltype(Value)> &&
SpecifiesTileBlockGemm<decltype(Value)> && SpecifiesTileOptimizations<decltype(Value)>;
concept TileAlgorithm =
ConvAlgorithmDescriptor<decltype(Value)> && SpecifiesTileThreadBlock<decltype(Value)> &&
SpecifiesTileTransfer<decltype(Value)> && SpecifiesTileConvSpecialization<decltype(Value)> &&
SpecifiesTileBlockGemm<decltype(Value)> && SpecifiesTileOptimizations<decltype(Value)>;
// FWD XDL algorithm concepts
template <auto Value>
concept FwdXdlAlgorithm = FwdXdlAlgorithmBase<Value> && SpecifiesGenericInstance<decltype(Value)>;
template <auto Value>
concept LargeTensorAlgorithm = FwdXdlAlgorithmBase<Value> && SpecifiesLargeTensorSupport<decltype(Value)>;
concept LargeTensorAlgorithm =
FwdXdlAlgorithmBase<Value> && SpecifiesLargeTensorSupport<decltype(Value)>;
template <auto Value>
concept FwdXdlV3Algorithm = FwdAlgorithmV3<decltype(Value)> && SpecifiesXdl<Value>;
@@ -82,63 +86,68 @@ concept FwdWmmaAlgorithm = FwdAlgorithm<decltype(Value)> && SpecifiesWmma<Value>
// FWD DL algorithms
template <auto Value>
concept FwdDlAlgorithm = FwdDlAlgorithmBase<decltype(Value)>;
// BWD weight XDL algorithm concepts
template <auto Value>
concept BwdXdlAlgorithm =
BwdXdlAlgorithmBase<Value> && SpecifiesTransposeTransfer<decltype(Value)> &&
BwdXdlAlgorithmBase<Value> && SpecifiesTransposeTransfer<decltype(Value)> &&
SpecifiesGenericInstance<decltype(Value)>;
template <auto Value>
concept BwdMultiDXdlAlgorithm = BwdXdlAlgorithmBase<Value> && SpecifiesMultipleDSupport<decltype(Value)>;
concept BwdMultiDXdlAlgorithm =
BwdXdlAlgorithmBase<Value> && SpecifiesMultipleDSupport<decltype(Value)>;
template <auto Value>
concept BwdXdlV3Algorithm = BwdXdlV3AlgorithmBase<Value>;
template <auto Value>
concept BwdTwoStageXdlAlgorithm = BwdXdlV3AlgorithmBase<Value> && SpecifiesTransposeTransfer<decltype(Value)> &&
SpecifiesTwoStageSupport<decltype(Value)>;
concept BwdTwoStageXdlAlgorithm =
BwdXdlV3AlgorithmBase<Value> && SpecifiesTransposeTransfer<decltype(Value)> &&
SpecifiesTwoStageSupport<decltype(Value)>;
// BWD weight WMMA algorithm concepts
template <auto Value>
concept BwdWmmaAlgorithm =
BwdWmmaAlgorithmBase<Value> && SpecifiesNumPrefetchStages<decltype(Value)> &&
BwdWmmaAlgorithmBase<Value> && SpecifiesNumPrefetchStages<decltype(Value)> &&
SpecifiesGemmPipeline<decltype(Value)> && SpecifiesGenericInstance<decltype(Value)>;
template <auto Value>
concept BwdMultiDWmmaV3Algorithm = BwdWmmaV3AlgorithmBase<Value> && SpecifiesMultipleDSupport<decltype(Value)>;
concept BwdMultiDWmmaV3Algorithm =
BwdWmmaV3AlgorithmBase<Value> && SpecifiesMultipleDSupport<decltype(Value)>;
template <auto Value>
concept BwdWmmaV3Algorithm =
BwdWmmaV3AlgorithmBase<Value> && SpecifiesTransposeTransfer<decltype(Value)> &&
BwdWmmaV3AlgorithmBase<Value> && SpecifiesTransposeTransfer<decltype(Value)> &&
SpecifiesGenericInstance<decltype(Value)>;
template <auto Value>
concept BwdTwoStageWmmaV3Algorithm = BwdWmmaV3AlgorithmBase<Value> && SpecifiesTransposeTransfer<decltype(Value)> &&
SpecifiesTwoStageSupport<decltype(Value)>;
concept BwdTwoStageWmmaV3Algorithm =
BwdWmmaV3AlgorithmBase<Value> && SpecifiesTransposeTransfer<decltype(Value)> &&
SpecifiesTwoStageSupport<decltype(Value)>;
// BWD weight DL algorithms
template <auto Value>
concept BwdDlAlgorithm = DlAlgorithm<decltype(Value)> && SpecifiesBwdWeightConvSpecialization<decltype(Value)> &&
SpecifiesDlBwdBlockTransfer<decltype(Value)>;
concept BwdDlAlgorithm =
DlAlgorithm<decltype(Value)> && SpecifiesBwdWeightConvSpecialization<decltype(Value)> &&
SpecifiesDlBwdBlockTransfer<decltype(Value)>;
// Concepts for valid XDL/WMMA algorithms
template <auto Value>
concept SpecifiesValidFwdXdlAlgorithm =
FwdXdlAlgorithm<Value> || FwdXdlV3Algorithm<Value> || LargeTensorAlgorithm<Value>;
concept SpecifiesValidFwdXdlAlgorithm =
FwdXdlAlgorithm<Value> || FwdXdlV3Algorithm<Value> || LargeTensorAlgorithm<Value>;
template <auto Value>
concept SpecifiesValidFwdWmmaAlgorithm = FwdWmmaAlgorithm<Value>;
template <auto Value>
concept SpecifiesValidBwdXdlAlgorithm =
BwdXdlAlgorithm<Value> || BwdXdlV3Algorithm<Value> ||
BwdTwoStageXdlAlgorithm<Value> || BwdMultiDXdlAlgorithm<Value>;
concept SpecifiesValidBwdXdlAlgorithm =
BwdXdlAlgorithm<Value> || BwdXdlV3Algorithm<Value> || BwdTwoStageXdlAlgorithm<Value> ||
BwdMultiDXdlAlgorithm<Value>;
template <auto Value>
concept SpecifiesValidBwdWmmaAlgorithm =
BwdWmmaAlgorithm<Value> || BwdWmmaV3Algorithm<Value> ||
BwdTwoStageWmmaV3Algorithm<Value> || BwdMultiDWmmaV3Algorithm<Value>;
concept SpecifiesValidBwdWmmaAlgorithm =
BwdWmmaAlgorithm<Value> || BwdWmmaV3Algorithm<Value> || BwdTwoStageWmmaV3Algorithm<Value> ||
BwdMultiDWmmaV3Algorithm<Value>;
template <auto Value>
concept FwdWarpGemmOrDL = SpecifiesValidWarpGemm<Value> || FwdDlAlgorithm<Value>;

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

@@ -34,7 +34,7 @@ struct ConvBwdWeightMultiDWmmaV3Factory
static constexpr auto BWD_CONV_SPECIALIZATION =
internal::SetBwdWeightConvSpecialization<ALGORITHM>();
static constexpr auto BLOCK = internal::SetThreadBlockInfo<ALGORITHM>();
static constexpr auto BLOCK = internal::SetThreadBlockInfo<ALGORITHM>();
static constexpr auto WARP_GEMM = ALGORITHM.warp_gemm;
static constexpr auto A_BLOCK_TRANSFER =
internal::SetBwdConvBlockTransfer<ALGORITHM.transfer.a>();

4
experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_multi_d_xdl_factory.hpp
View File

@@ -34,8 +34,8 @@ struct ConvBwdWeightMultiDXdlFactory
static constexpr auto BWD_CONV_SPECIALIZATION =
internal::SetBwdWeightConvSpecialization<ALGORITHM>();
static constexpr auto BLOCK = internal::SetThreadBlockInfo<ALGORITHM>();
static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm;
static constexpr auto BLOCK = internal::SetThreadBlockInfo<ALGORITHM>();
static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm;
static constexpr auto A_BLOCK_TRANSFER =
internal::SetBwdConvBlockTransfer<ALGORITHM.transfer.a>();
static constexpr auto B_BLOCK_TRANSFER =

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

@@ -34,7 +34,7 @@ struct ConvBwdWeightTwoStageWmmaV3Factory
static constexpr auto BWD_CONV_SPECIALIZATION =
internal::SetBwdWeightConvSpecialization<ALGORITHM>();
static constexpr auto BLOCK = internal::SetThreadBlockInfo<ALGORITHM>();
static constexpr auto BLOCK = internal::SetThreadBlockInfo<ALGORITHM>();
static constexpr auto WARP_GEMM = ALGORITHM.warp_gemm;
static constexpr auto A_BLOCK_TRANSFER =
internal::SetBwdConvBlockTransfer<ALGORITHM.transfer.a>();

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

@@ -34,8 +34,8 @@ struct ConvBwdWeightTwoStageXdlFactory
static constexpr auto BWD_CONV_SPECIALIZATION =
internal::SetBwdWeightConvSpecialization<ALGORITHM>();
static constexpr auto BLOCK = internal::SetThreadBlockInfo<ALGORITHM>();
static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm;
static constexpr auto BLOCK = internal::SetThreadBlockInfo<ALGORITHM>();
static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm;
static constexpr auto A_BLOCK_TRANSFER =
internal::SetBwdConvBlockTransfer<ALGORITHM.transfer.a>();
static constexpr auto B_BLOCK_TRANSFER =

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

@@ -34,7 +34,7 @@ struct ConvBwdWeightWmmaFactory
static constexpr auto BWD_CONV_SPECIALIZATION =
internal::SetBwdWeightConvSpecialization<ALGORITHM>();
static constexpr auto BLOCK = internal::SetThreadBlockInfo<ALGORITHM>();
static constexpr auto BLOCK = internal::SetThreadBlockInfo<ALGORITHM>();
static constexpr auto WARP_GEMM = ALGORITHM.warp_gemm;
static constexpr auto GRIDWISE_GEMM_PIPELINE_VERSION =
internal::SetGridwiseGemmPipelineVersion<ALGORITHM>();

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

@@ -34,7 +34,7 @@ struct ConvBwdWeightWmmaV3Factory
static constexpr auto BWD_CONV_SPECIALIZATION =
internal::SetBwdWeightConvSpecialization<ALGORITHM>();
static constexpr auto BLOCK = internal::SetThreadBlockInfo<ALGORITHM>();
static constexpr auto BLOCK = internal::SetThreadBlockInfo<ALGORITHM>();
static constexpr auto WARP_GEMM = ALGORITHM.warp_gemm;
static constexpr auto A_BLOCK_TRANSFER =
internal::SetBwdConvBlockTransfer<ALGORITHM.transfer.a>();

5
experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_factory.hpp
View File

@@ -34,8 +34,8 @@ struct ConvBwdWeightXdlFactory
static constexpr auto BWD_CONV_SPECIALIZATION =
internal::SetBwdWeightConvSpecialization<ALGORITHM>();
static constexpr auto BLOCK = internal::SetThreadBlockInfo<ALGORITHM>();
static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm;
static constexpr auto BLOCK = internal::SetThreadBlockInfo<ALGORITHM>();
static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm;
static constexpr auto A_BLOCK_TRANSFER =
internal::SetBwdConvBlockTransfer<ALGORITHM.transfer.a>();
static constexpr auto B_BLOCK_TRANSFER =
@@ -57,7 +57,6 @@ struct ConvBwdWeightXdlFactory
"A nd B block transfer vector load size need to be the same");
static constexpr size_t GMEM_VECTOR_LOAD_SIZE = A_BLOCK_TRANSFER.global_memory_vector_load_size;
// The forward convolution kernel class instance.
using Instance = ck::tensor_operation::device::DeviceGroupedConvBwdWeight_Xdl_CShuffle<
SPATIAL_DIM,

4
experimental/builder/include/ck_tile/builder/factory/conv_bwd_weight_xdl_v3_factory.hpp
View File

@@ -34,8 +34,8 @@ struct ConvBwdWeightXdlV3Factory
static constexpr auto BWD_CONV_SPECIALIZATION =
internal::SetBwdWeightConvSpecialization<ALGORITHM>();
static constexpr auto BLOCK = internal::SetThreadBlockInfo<ALGORITHM>();
static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm;
static constexpr auto BLOCK = internal::SetThreadBlockInfo<ALGORITHM>();
static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm;
static constexpr auto A_BLOCK_TRANSFER =
internal::SetBwdConvBlockTransfer<ALGORITHM.transfer.a>();
static constexpr auto B_BLOCK_TRANSFER =

67
experimental/builder/include/ck_tile/builder/factory/conv_dispatcher.hpp
View File

@@ -117,7 +117,7 @@ constexpr auto make_conv_instance()
// Forward direction (supports most algorithm variants)
else if constexpr(ConvDirectionIsForward<SIGNATURE>)
{
if constexpr (SpecifiesXdl<ALGORITHM>)
if constexpr(SpecifiesXdl<ALGORITHM>)
{
if constexpr(FwdXdlV3Algorithm<ALGORITHM>)
{
@@ -129,30 +129,31 @@ constexpr auto make_conv_instance()
}
else if constexpr(LargeTensorAlgorithm<ALGORITHM>)
{
return typename ConvFwdLargeTensorFactory<SIGNATURE, ALGORITHM, VERSION>::Instance{};
return
typename ConvFwdLargeTensorFactory<SIGNATURE, ALGORITHM, VERSION>::Instance{};
}
else
{
static_assert(
SpecifiesValidFwdXdlAlgorithm<ALGORITHM>,
"No suitable forward convolution XDL kernel factory found for the provided ALGORITHM. "
"No suitable forward convolution XDL kernel factory found for the provided "
"ALGORITHM. "
"The ALGORITHM must satisfy requirements for one of: XDL V3, generic XDL, "
"DL (NHWC layout), or Large Tensor variant.");
}
}
else if constexpr (SpecifiesWmma<ALGORITHM>)
else if constexpr(SpecifiesWmma<ALGORITHM>)
{
if constexpr(FwdWmmaAlgorithm<ALGORITHM>)
{
return typename ConvFwdWmmaFactory<SIGNATURE, ALGORITHM, VERSION>::Instance{};
}
else
else
{
static_assert(FwdWmmaAlgorithm<ALGORITHM>,
"Did not find matching WMMA factory.");
static_assert(FwdWmmaAlgorithm<ALGORITHM>, "Did not find matching WMMA factory.");
}
}
else if constexpr (FwdDlAlgorithm<ALGORITHM>)
else if constexpr(FwdDlAlgorithm<ALGORITHM>)
{
return typename ConvFwdDlFactory<SIGNATURE, ALGORITHM, VERSION>::Instance{};
}
@@ -173,22 +174,23 @@ constexpr auto make_conv_instance()
// Backward weight direction (will expand with more algorithms in the future)
else if constexpr(ConvDirectionIsBackwardWeight<SIGNATURE>)
{
if constexpr (SpecifiesXdl<ALGORITHM>)
if constexpr(SpecifiesXdl<ALGORITHM>)
{
// Start from more specialized and end with least specialized.
if constexpr(BwdTwoStageXdlAlgorithm<ALGORITHM>)
{
return
typename ConvBwdWeightTwoStageXdlFactory<SIGNATURE, ALGORITHM, VERSION>::Instance{};
return typename ConvBwdWeightTwoStageXdlFactory<SIGNATURE, ALGORITHM, VERSION>::
Instance{};
}
else if constexpr(BwdMultiDXdlAlgorithm<ALGORITHM>)
{
return
typename ConvBwdWeightMultiDXdlFactory<SIGNATURE, ALGORITHM, VERSION>::Instance{};
return typename ConvBwdWeightMultiDXdlFactory<SIGNATURE, ALGORITHM, VERSION>::
Instance{};
}
else if constexpr(BwdXdlV3Algorithm<ALGORITHM>)
{
return typename ConvBwdWeightXdlV3Factory<SIGNATURE, ALGORITHM, VERSION>::Instance{};
return
typename ConvBwdWeightXdlV3Factory<SIGNATURE, ALGORITHM, VERSION>::Instance{};
}
else if constexpr(BwdXdlAlgorithm<ALGORITHM>)
{
@@ -196,46 +198,51 @@ constexpr auto make_conv_instance()
}
else
{
static_assert(
SpecifiesValidBwdXdlAlgorithm<ALGORITHM>,
"No suitable backward weight convolution XDL kernel factory found for the provided ALGORITHM. "
"The ALGORITHM must satisfy requirements for one of: Two-Stage XDL, Multi-D XDL, DL, "
"generic XDL, or XDL V3 variant.");
static_assert(SpecifiesValidBwdXdlAlgorithm<ALGORITHM>,
"No suitable backward weight convolution XDL kernel factory found "
"for the provided ALGORITHM. "
"The ALGORITHM must satisfy requirements for one of: Two-Stage XDL, "
"Multi-D XDL, DL, "
"generic XDL, or XDL V3 variant.");
}
}
else if constexpr (SpecifiesWmma<ALGORITHM>)
else if constexpr(SpecifiesWmma<ALGORITHM>)
{
// Start from more specialized and end with least specialized.
if constexpr(BwdTwoStageWmmaV3Algorithm<ALGORITHM>)
{
return typename ConvBwdWeightTwoStageWmmaV3Factory<SIGNATURE, ALGORITHM, VERSION>::Instance{};
return typename ConvBwdWeightTwoStageWmmaV3Factory<SIGNATURE, ALGORITHM, VERSION>::
Instance{};
}
else if constexpr(BwdMultiDWmmaV3Algorithm<ALGORITHM>)
{
return typename ConvBwdWeightMultiDWmmaV3Factory<SIGNATURE, ALGORITHM, VERSION>::Instance{};
return typename ConvBwdWeightMultiDWmmaV3Factory<SIGNATURE, ALGORITHM, VERSION>::
Instance{};
}
else if constexpr(BwdWmmaV3Algorithm<ALGORITHM>)
{
return typename ConvBwdWeightWmmaV3Factory<SIGNATURE, ALGORITHM, VERSION>::Instance{};
return
typename ConvBwdWeightWmmaV3Factory<SIGNATURE, ALGORITHM, VERSION>::Instance{};
}
else if constexpr(BwdWmmaAlgorithm<ALGORITHM>)
{
return typename ConvBwdWeightWmmaFactory<SIGNATURE, ALGORITHM, VERSION>::Instance{};
}
else
else
{
static_assert(
SpecifiesValidBwdWmmaAlgorithm<ALGORITHM>,
"No suitable backward weight convolution WMMA kernel factory found for the provided ALGORITHM. "
"The ALGORITHM must satisfy requirements for one of: Two-Stage WMMA V3, Multi-D WMMA V3, "
"WMMA V3, or generic WMMA variant.");
static_assert(SpecifiesValidBwdWmmaAlgorithm<ALGORITHM>,
"No suitable backward weight convolution WMMA kernel factory found "
"for the provided ALGORITHM. "
"The ALGORITHM must satisfy requirements for one of: Two-Stage WMMA "
"V3, Multi-D WMMA V3, "
"WMMA V3, or generic WMMA variant.");
}
}
else if constexpr(BwdDlAlgorithm<ALGORITHM>)
{
return typename ConvBwdWeightDlFactory<SIGNATURE, ALGORITHM, VERSION>::Instance{};
}
else
else
{
static_assert(BwdWarpGemmOrDL<ALGORITHM>,
"Backward convolution: Algorithm must specify either DL, XDL or WMMA.");

4
experimental/builder/include/ck_tile/builder/factory/conv_fwd_v3_factory.hpp
View File

@@ -41,8 +41,8 @@ struct ConvFwdXdlV3Factory
static constexpr internal::ConvSpec SPECIALIZATION{.conv_spec = FWD_CONV_SPECIALIZATION,
.gemm_spec = GEMM_SPECIALIZATION};
static constexpr auto BLOCK = internal::SetThreadBlockInfo<ALGORITHM>();
static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm;
static constexpr auto BLOCK = internal::SetThreadBlockInfo<ALGORITHM>();
static constexpr auto XDL_PARAMS = ALGORITHM.warp_gemm;
static constexpr auto A_BLOCK_TRANSFER =
internal::SetFwdConvBlockTransfer<ALGORITHM.transfer.a>();
static constexpr auto B_BLOCK_TRANSFER =

76
experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_block_transfer.hpp
View File

@@ -17,11 +17,11 @@ struct BlockTransfer
ck::Array<size_t, ThreadClusterRank> thread_cluster_order{};
ck::Array<size_t, ThreadClusterRank> src_access_order{};
size_t global_memory_vector_load_size = 0;
size_t src_vector_dim = 0;
size_t src_scalar_per_vector = 0;
size_t lds_dst_scalar_per_vector = 0;
bool is_direct_load = false;
bool lds_padding = false;
size_t src_vector_dim = 0;
size_t src_scalar_per_vector = 0;
size_t lds_dst_scalar_per_vector = 0;
bool is_direct_load = false;
bool lds_padding = false;
};
template <auto TRANSFER>
@@ -33,15 +33,15 @@ constexpr BlockTransfer<> SetFwdConvBlockTransfer()
auto& lds_cfg = TRANSFER.lds_transfer;
return BlockTransfer<>{
.thread_cluster_dims = {block_xfer.k0, block_xfer.m_n, block_xfer.k1},
.thread_cluster_order = {block_order.order[0], block_order.order[1], block_order.order[2]},
.src_access_order = {src_order.order[0], src_order.order[1], src_order.order[2]},
.thread_cluster_dims = {block_xfer.k0, block_xfer.m_n, block_xfer.k1},
.thread_cluster_order = {block_order.order[0], block_order.order[1], block_order.order[2]},
.src_access_order = {src_order.order[0], src_order.order[1], src_order.order[2]},
.global_memory_vector_load_size = lds_cfg.global_memory_vector_load_size,
.src_vector_dim = lds_cfg.src_vector_dim,
.src_scalar_per_vector = lds_cfg.src_scalar_per_vector,
.lds_dst_scalar_per_vector = lds_cfg.lds_dst_scalar_per_vector,
.is_direct_load = lds_cfg.is_direct_load,
.lds_padding = lds_cfg.lds_padding,
.src_vector_dim = lds_cfg.src_vector_dim,
.src_scalar_per_vector = lds_cfg.src_scalar_per_vector,
.lds_dst_scalar_per_vector = lds_cfg.lds_dst_scalar_per_vector,
.is_direct_load = lds_cfg.is_direct_load,
.lds_padding = lds_cfg.lds_padding,
};
}
@@ -60,38 +60,38 @@ constexpr auto SetBwdConvBlockTransfer()
if constexpr(array_length == 3)
{
return BlockTransfer<3>{
.thread_cluster_dims = {block_xfer.k0, block_xfer.m_n, block_xfer.k1},
.thread_cluster_order = {block_order.order[0],
block_order.order[1],
block_order.order[2]},
.src_access_order = {src_order.order[0], src_order.order[1], src_order.order[2]},
.thread_cluster_dims = {block_xfer.k0, block_xfer.m_n, block_xfer.k1},
.thread_cluster_order = {block_order.order[0],
block_order.order[1],
block_order.order[2]},
.src_access_order = {src_order.order[0], src_order.order[1], src_order.order[2]},
.global_memory_vector_load_size = lds_cfg.global_memory_vector_load_size,
.src_vector_dim = lds_cfg.src_vector_dim,
.src_scalar_per_vector = lds_cfg.src_scalar_per_vector,
.lds_dst_scalar_per_vector = lds_cfg.lds_dst_scalar_per_vector,
.lds_padding = lds_cfg.lds_padding,
.src_vector_dim = lds_cfg.src_vector_dim,
.src_scalar_per_vector = lds_cfg.src_scalar_per_vector,
.lds_dst_scalar_per_vector = lds_cfg.lds_dst_scalar_per_vector,
.lds_padding = lds_cfg.lds_padding,
};
}
else if constexpr(array_length == 4)
{
return BlockTransfer<4>{
.thread_cluster_dims = {block_xfer.k_batch_size,
block_xfer.k0,
block_xfer.m_n,
block_xfer.k1},
.thread_cluster_order = {block_order.order[0],
block_order.order[1],
block_order.order[2],
block_order.order[3]},
.src_access_order = {src_order.order[0],
src_order.order[1],
src_order.order[2],
src_order.order[3]},
.thread_cluster_dims = {block_xfer.k_batch_size,
block_xfer.k0,
block_xfer.m_n,
block_xfer.k1},
.thread_cluster_order = {block_order.order[0],
block_order.order[1],
block_order.order[2],
block_order.order[3]},
.src_access_order = {src_order.order[0],
src_order.order[1],
src_order.order[2],
src_order.order[3]},
.global_memory_vector_load_size = lds_cfg.global_memory_vector_load_size,
.src_vector_dim = lds_cfg.src_vector_dim,
.src_scalar_per_vector = lds_cfg.src_scalar_per_vector,
.lds_dst_scalar_per_vector = lds_cfg.lds_dst_scalar_per_vector,
.lds_padding = lds_cfg.lds_padding,
.src_vector_dim = lds_cfg.src_vector_dim,
.src_scalar_per_vector = lds_cfg.src_scalar_per_vector,
.lds_dst_scalar_per_vector = lds_cfg.lds_dst_scalar_per_vector,
.lds_padding = lds_cfg.lds_padding,
};
}
else

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

@@ -65,11 +65,10 @@ consteval BlockGemmSpec SetBlockGemm()
default: throw "Unknown PipelineVersion";
}
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};
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>

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

@@ -232,12 +232,12 @@ enum class PipelineScheduler
enum class ConvAlgorithmSpecialization
{
NONE = 0,
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
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,

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

@@ -19,12 +19,13 @@ constexpr auto SIGNATURE = ckt::ConvSignature{.spatial_dim = 3,
.weight = {.config = {.layout = GKZYXC}},
.output = {.config = {.layout = NGKDHW}}};
constexpr auto ALGORITHM = cku::ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle{}
.with_thread_block(cku::ThreadBlock_64_32x32x32)
.with_gemm_config(cku::GemmParams_Wmma_16x16_2x1_per_wave)
.with_transfer(cku::BwdTransfer_4x8x1_4x16x1_v3)
.with_bwd_specialization(ckb::ConvSpecialization::DEFAULT)
.with_gemm_pipeline(ckb::PipelineVersion::V1, ckb::PipelineScheduler::DEFAULT);
constexpr auto ALGORITHM =
cku::ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle{}
.with_thread_block(cku::ThreadBlock_64_32x32x32)
.with_gemm_config(cku::GemmParams_Wmma_16x16_2x1_per_wave)
.with_transfer(cku::BwdTransfer_4x8x1_4x16x1_v3)
.with_bwd_specialization(ckb::ConvSpecialization::DEFAULT)
.with_gemm_pipeline(ckb::PipelineVersion::V1, ckb::PipelineScheduler::DEFAULT);
using Builder = ckb::ConvBuilder<SIGNATURE, ALGORITHM>;
using Instance = Builder::Instance;

15
experimental/builder/test/conv/ck/test_ckb_conv_fwd_2d_bf16.cpp
View File

@@ -62,14 +62,13 @@ TEST(FwdConvInstances,
.weight = {.config = {.layout = GKYXC}},
.output = {.config = {.layout = NHWGK}}};
constexpr auto FwdConvAlgorithm =
ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_CShuffle_V3{}
.with_thread_block(ThreadBlock_256_256x256x32)
.with_gemm_config(FwdGemmParams_Xdl_4x4_per_wave)
.with_transfer(Transfer_4x64x1)
.with_fwd_specializations(ConvSpecialization::FILTER_3x3,
GemmSpecialization::MNKPadding)
.with_gemm_pipeline(BlockGemmDesc_v5_intrawave);
constexpr auto FwdConvAlgorithm = ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_CShuffle_V3{}
.with_thread_block(ThreadBlock_256_256x256x32)
.with_gemm_config(FwdGemmParams_Xdl_4x4_per_wave)
.with_transfer(Transfer_4x64x1)
.with_fwd_specializations(ConvSpecialization::FILTER_3x3,
GemmSpecialization::MNKPadding)
.with_gemm_pipeline(BlockGemmDesc_v5_intrawave);
using Builder = ConvBuilder<FwdConvSignature, FwdConvAlgorithm>;

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

@@ -31,10 +31,10 @@ static_assert(ckb::ThreadBlockDescriptor<ThreadBlock>);
struct WarpGemmParams
{
MatrixInstructionType matrix_instruction;
size_t gemm_m_per_instruction = 0;
size_t gemm_n_per_instruction = 0;
size_t gemm_m_iters_per_wave = 0;
size_t gemm_n_iters_per_wave = 0;
size_t gemm_m_per_instruction = 0;
size_t gemm_n_per_instruction = 0;
size_t gemm_m_iters_per_wave = 0;
size_t gemm_n_iters_per_wave = 0;
};
static_assert(ckb::WarpGemmDescriptor<WarpGemmParams>);
@@ -370,7 +370,7 @@ 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<GemmPipeline_, ConvAlgorithmTemplate>);
auto result = *this;
auto result = *this;
result.gemm_pipeline.num_conv_groups_to_merge = num_groups_to_merge;
return result;
}
@@ -378,7 +378,7 @@ struct ConvAlgorithmTemplate : Components...
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;
auto result = *this;
result.gemm_pipeline.num_gemm_k_prefetch_stages = num_prefetch_stages;
return result;
}
@@ -387,7 +387,7 @@ struct ConvAlgorithmTemplate : Components...
constexpr auto with_gemm_pipeline(const PL& pl) const
{
static_assert(std::is_base_of_v<GemmPipeline_, ConvAlgorithmTemplate>);
auto result = *this;
auto result = *this;
result.gemm_pipeline = pl;
return result;
}
@@ -395,7 +395,7 @@ struct ConvAlgorithmTemplate : Components...
constexpr auto with_gemm_pipeline(const PipelineVersion plv) const
{
static_assert(std::is_base_of_v<GemmPipeline_, ConvAlgorithmTemplate>);
auto result = *this;
auto result = *this;
result.gemm_pipeline.pipeline_version = plv;
return result;
}
@@ -403,7 +403,7 @@ struct ConvAlgorithmTemplate : Components...
constexpr auto with_gemm_pipeline(const PipelineScheduler sch) const
{
static_assert(std::is_base_of_v<GemmPipeline_, ConvAlgorithmTemplate>);
auto result = *this;
auto result = *this;
result.gemm_pipeline.scheduler = sch;
return result;
}
@@ -411,7 +411,7 @@ struct ConvAlgorithmTemplate : Components...
constexpr auto with_gemm_pipeline(const PipelineVersion plv, const PipelineScheduler sch) const
{
static_assert(std::is_base_of_v<GemmPipeline_, ConvAlgorithmTemplate>);
auto result = *this;
auto result = *this;
result.gemm_pipeline.pipeline_version = plv;
result.gemm_pipeline.scheduler = sch;
return result;
@@ -535,7 +535,7 @@ using ConvAlgorithm_Tile_GroupedConvolutionKernel = ConvAlgorithmTemplate<TileTh
TileOptimizations_>;
// Reference algorithm descriptor - for GPU reference validation
using ConvAlgorithm_Reference = ConvAlgorithmTemplate<AlgorithmSpecialization_<REFERENCE>>;
using ConvAlgorithm_Reference = ConvAlgorithmTemplate<AlgorithmSpecialization_<REFERENCE>>;
// Bwd weight algorithm types
using ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuffle =
@@ -543,7 +543,8 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuffle =
WarpGemm_,
InputOutputTileTransfer_<4>,
ConvSpecializationBwdWeight_,
GemmPipeline_, // Not needed, but we need this to meet the ConvAlgorithm concept.
GemmPipeline_, // Not needed, but we need this to meet the ConvAlgorithm
// concept.
TransposeParams_,
AlgorithmSpecialization_<>>;
@@ -593,7 +594,8 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle =
ConvAlgorithmTemplate<ThreadBlock_,
WarpGemm_,
InputOutputTileTransfer_<4>,
GemmPipeline_, // Not needed, but we need this to meet the ConvAlgorithm concept.
GemmPipeline_, // Not needed, but we need this to meet the ConvAlgorithm
// concept.
ConvSpecializationBwdWeight_,
AlgorithmSpecialization_<MULTIPLE_D>>;

53
experimental/builder/test/test_conv_description.cpp
View File

@@ -117,57 +117,60 @@ static_assert(!ckb::ConvSignatureDescriptor<ConvSignatureWithInvalidOptionalPara
struct DefaultAlgorithm
{
static constexpr ckb::ConvAlgorithmSpecialization specialization =
static constexpr ckb::ConvAlgorithmSpecialization specialization =
ckb::ConvAlgorithmSpecialization::PIPELINE_V3;
ckb::test::ThreadBlock thread_block{.block_size = 256,
.tile_size = {.m = 256, .n = 256, .k = 32}};
ckb::test::WarpGemmParams warp_gemm{
.matrix_instruction = ckb::MatrixInstructionType::XDL,
.gemm_m_per_instruction = 16, .gemm_n_per_instruction = 16, .gemm_m_iters_per_wave = 8, .gemm_n_iters_per_wave = 8};
ckb::test::WarpGemmParams warp_gemm{.matrix_instruction = ckb::MatrixInstructionType::XDL,
.gemm_m_per_instruction = 16,
.gemm_n_per_instruction = 16,
.gemm_m_iters_per_wave = 8,
.gemm_n_iters_per_wave = 8};
ckb::test::InputOutputTileTransfer<> transfer{
.a =
{
.thread_cluster = {.k0 = 1, .m_n = 128, .k1 = 2},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 2,
.src_scalar_per_vector = 2,
.lds_dst_scalar_per_vector = 2,
.is_direct_load = false,
.lds_padding = false},
.thread_cluster = {.k0 = 1, .m_n = 128, .k1 = 2},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 2,
.src_scalar_per_vector = 2,
.lds_dst_scalar_per_vector = 2,
.is_direct_load = false,
.lds_padding = false},
.thread_cluster_access_order = {.order = {0, 1, 2}},
.src_access_order = {.order = {0, 1, 2}},
},
.b =
{
.thread_cluster = {.k0 = 1, .m_n = 128, .k1 = 2},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 2,
.src_scalar_per_vector = 2,
.lds_dst_scalar_per_vector = 2,
.is_direct_load = false,
.lds_padding = false},
.thread_cluster = {.k0 = 1, .m_n = 128, .k1 = 2},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 2,
.src_scalar_per_vector = 2,
.lds_dst_scalar_per_vector = 2,
.is_direct_load = false,
.lds_padding = false},
.thread_cluster_access_order = {.order = {0, 1, 2}},
.src_access_order = {.order = {0, 1, 2}},
},
.c =
{
.thread_cluster =
{.gemm_m_block_size = 1, .gemm_m_per_block = 32, .gemm_n_block_size = 1, .gemm_n_per_block = 8},
.epilogue = {.m_xdl_per_wave_per_shuffle = 1,
.n_per_wave_per_shuffle = 1,
.scalar_per_vector = 2},
.thread_cluster = {.gemm_m_block_size = 1,
.gemm_m_per_block = 32,
.gemm_n_block_size = 1,
.gemm_n_per_block = 8},
.epilogue = {.m_xdl_per_wave_per_shuffle = 1,
.n_per_wave_per_shuffle = 1,
.scalar_per_vector = 2},
},
};
ckb::ConvSpecialization fwd_specialization = ckb::ConvSpecialization::DEFAULT;
ckb::GemmSpecialization gemm_specialization = ckb::GemmSpecialization::Default;
ckb::test::GemmPipeline gemm_pipeline{.pipeline_version = ckb::PipelineVersion::V4,
.scheduler =ckb::PipelineScheduler::INTRAWAVE};
.scheduler = ckb::PipelineScheduler::INTRAWAVE};
};
static_assert(ckb::ConvAlgorithmDescriptor<DefaultAlgorithm>);

330
experimental/builder/test/utils/ckb_conv_test_configs.hpp
View File

@@ -55,13 +55,13 @@ constexpr DlTransfer<5> DlTransfer5D{.a = DlBlockTransfer_1x8x1x1x1,
constexpr InputOutputTileTransfer<> Transfer_4x64x1{
.a =
{
.thread_cluster = {.k0 = 4, .m_n = 64, .k1 = 1},
.thread_cluster = {.k0 = 4, .m_n = 64, .k1 = 1},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 2,
.src_scalar_per_vector = 2,
.lds_dst_scalar_per_vector = 8,
.is_direct_load = false,
.lds_padding = false},
.src_vector_dim = 2,
.src_scalar_per_vector = 2,
.lds_dst_scalar_per_vector = 8,
.is_direct_load = false,
.lds_padding = false},
.thread_cluster_access_order = {1, 0, 2},
.src_access_order = {1, 0, 2},
},
@@ -69,21 +69,23 @@ constexpr InputOutputTileTransfer<> Transfer_4x64x1{
{
.thread_cluster = {.k0 = 4, .m_n = 64, .k1 = 1},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 2,
.src_scalar_per_vector = 8,
.lds_dst_scalar_per_vector = 8,
.is_direct_load = false,
.lds_padding = false},
.src_vector_dim = 2,
.src_scalar_per_vector = 8,
.lds_dst_scalar_per_vector = 8,
.is_direct_load = false,
.lds_padding = false},
.thread_cluster_access_order = {1, 0, 2},
.src_access_order = {1, 0, 2},
},
.c =
{
.thread_cluster =
{.gemm_m_block_size = 1, .gemm_m_per_block = 32, .gemm_n_block_size = 1, .gemm_n_per_block = 8},
.epilogue = {.m_xdl_per_wave_per_shuffle = 1,
.n_per_wave_per_shuffle = 1,
.scalar_per_vector = 4},
.thread_cluster = {.gemm_m_block_size = 1,
.gemm_m_per_block = 32,
.gemm_n_block_size = 1,
.gemm_n_per_block = 8},
.epilogue = {.m_xdl_per_wave_per_shuffle = 1,
.n_per_wave_per_shuffle = 1,
.scalar_per_vector = 4},
},
};
@@ -92,11 +94,11 @@ constexpr InputOutputTileTransfer<4> BwdTransfer_4x64x1{
{
.thread_cluster = {.k0 = 4, .m_n = 64, .k1 = 1, .k_batch_size = 1},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 2,
.src_scalar_per_vector = 2,
.lds_dst_scalar_per_vector = 4,
.is_direct_load = false,
.lds_padding = true},
.src_vector_dim = 2,
.src_scalar_per_vector = 2,
.lds_dst_scalar_per_vector = 4,
.is_direct_load = false,
.lds_padding = true},
.thread_cluster_access_order = {0, 3, 1, 2},
.src_access_order = {0, 2, 1, 3},
},
@@ -104,21 +106,23 @@ constexpr InputOutputTileTransfer<4> BwdTransfer_4x64x1{
{
.thread_cluster = {.k0 = 4, .m_n = 64, .k1 = 1, .k_batch_size = 1},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 2,
.src_scalar_per_vector = 2,
.lds_dst_scalar_per_vector = 4,
.is_direct_load = false,
.lds_padding = true},
.src_vector_dim = 2,
.src_scalar_per_vector = 2,
.lds_dst_scalar_per_vector = 4,
.is_direct_load = false,
.lds_padding = true},
.thread_cluster_access_order = {0, 3, 1, 2},
.src_access_order = {0, 2, 1, 3},
},
.c =
{
.thread_cluster =
{.gemm_m_block_size = 1, .gemm_m_per_block = 32, .gemm_n_block_size = 1, .gemm_n_per_block = 8},
.epilogue = {.m_xdl_per_wave_per_shuffle = 1,
.n_per_wave_per_shuffle = 1,
.scalar_per_vector = 8},
.thread_cluster = {.gemm_m_block_size = 1,
.gemm_m_per_block = 32,
.gemm_n_block_size = 1,
.gemm_n_per_block = 8},
.epilogue = {.m_xdl_per_wave_per_shuffle = 1,
.n_per_wave_per_shuffle = 1,
.scalar_per_vector = 8},
},
};
@@ -127,11 +131,11 @@ constexpr InputOutputTileTransfer<> BwdTransfer_4x8x1_4x16x1_v3{
{
.thread_cluster = {.k0 = 4, .m_n = 8, .k1 = 1},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 1,
.src_scalar_per_vector = 2,
.lds_dst_scalar_per_vector = 2,
.is_direct_load = false,
.lds_padding = false},
.src_vector_dim = 1,
.src_scalar_per_vector = 2,
.lds_dst_scalar_per_vector = 2,
.is_direct_load = false,
.lds_padding = false},
.thread_cluster_access_order = {2, 0, 1},
.src_access_order = {1, 0, 2},
},
@@ -139,21 +143,23 @@ constexpr InputOutputTileTransfer<> BwdTransfer_4x8x1_4x16x1_v3{
{
.thread_cluster = {.k0 = 4, .m_n = 16, .k1 = 1},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 1,
.src_scalar_per_vector = 2,
.lds_dst_scalar_per_vector = 2,
.is_direct_load = false,
.lds_padding = false},
.src_vector_dim = 1,
.src_scalar_per_vector = 2,
.lds_dst_scalar_per_vector = 2,
.is_direct_load = false,
.lds_padding = false},
.thread_cluster_access_order = {2, 0, 1},
.src_access_order = {1, 0, 2},
},
.c =
{
.thread_cluster =
{.gemm_m_block_size = 1, .gemm_m_per_block = 8, .gemm_n_block_size = 1, .gemm_n_per_block = 8},
.epilogue = {.m_xdl_per_wave_per_shuffle = 1,
.n_per_wave_per_shuffle = 1,
.scalar_per_vector = 2},
.thread_cluster = {.gemm_m_block_size = 1,
.gemm_m_per_block = 8,
.gemm_n_block_size = 1,
.gemm_n_per_block = 8},
.epilogue = {.m_xdl_per_wave_per_shuffle = 1,
.n_per_wave_per_shuffle = 1,
.scalar_per_vector = 2},
},
};
@@ -162,11 +168,11 @@ constexpr InputOutputTileTransfer<> Transfer_4x64x1_fp8{
{
.thread_cluster = {.k0 = 4, .m_n = 64, .k1 = 1},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 2,
.src_scalar_per_vector = 8,
.lds_dst_scalar_per_vector = 8,
.is_direct_load = false,
.lds_padding = true},
.src_vector_dim = 2,
.src_scalar_per_vector = 8,
.lds_dst_scalar_per_vector = 8,
.is_direct_load = false,
.lds_padding = true},
.thread_cluster_access_order = {1, 0, 2},
.src_access_order = {1, 0, 2},
},
@@ -174,21 +180,23 @@ constexpr InputOutputTileTransfer<> Transfer_4x64x1_fp8{
{
.thread_cluster = {.k0 = 4, .m_n = 64, .k1 = 1},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 2,
.src_scalar_per_vector = 8,
.lds_dst_scalar_per_vector = 8,
.is_direct_load = false,
.lds_padding = true},
.src_vector_dim = 2,
.src_scalar_per_vector = 8,
.lds_dst_scalar_per_vector = 8,
.is_direct_load = false,
.lds_padding = true},
.thread_cluster_access_order = {1, 0, 2},
.src_access_order = {1, 0, 2},
},
.c =
{
.thread_cluster =
{.gemm_m_block_size = 1, .gemm_m_per_block = 32, .gemm_n_block_size = 1, .gemm_n_per_block = 8},
.epilogue = {.m_xdl_per_wave_per_shuffle = 1,
.n_per_wave_per_shuffle = 1,
.scalar_per_vector = 8},
.thread_cluster = {.gemm_m_block_size = 1,
.gemm_m_per_block = 32,
.gemm_n_block_size = 1,
.gemm_n_per_block = 8},
.epilogue = {.m_xdl_per_wave_per_shuffle = 1,
.n_per_wave_per_shuffle = 1,
.scalar_per_vector = 8},
},
};
@@ -196,34 +204,36 @@ constexpr InputOutputTileTransfer<> Transfer_4x16x1{
.a =
{
.thread_cluster = {.k0 = 4, .m_n = 16, .k1 = 1},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 2,
.src_scalar_per_vector = 8,
.lds_dst_scalar_per_vector = 8,
.is_direct_load = false,
.lds_padding = true},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 2,
.src_scalar_per_vector = 8,
.lds_dst_scalar_per_vector = 8,
.is_direct_load = false,
.lds_padding = true},
.thread_cluster_access_order = {1, 0, 2},
.src_access_order = {1, 0, 2},
},
.b =
{
.thread_cluster = {.k0 = 4, .m_n = 16, .k1 = 1},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 2,
.src_scalar_per_vector = 8,
.lds_dst_scalar_per_vector = 8,
.is_direct_load = false,
.lds_padding = true},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 2,
.src_scalar_per_vector = 8,
.lds_dst_scalar_per_vector = 8,
.is_direct_load = false,
.lds_padding = true},
.thread_cluster_access_order = {1, 0, 2},
.src_access_order = {1, 0, 2},
},
.c =
{
.thread_cluster =
{.gemm_m_block_size = 1, .gemm_m_per_block = 16, .gemm_n_block_size = 1, .gemm_n_per_block = 4},
.epilogue = {.m_xdl_per_wave_per_shuffle = 1,
.n_per_wave_per_shuffle = 1,
.scalar_per_vector = 8},
.thread_cluster = {.gemm_m_block_size = 1,
.gemm_m_per_block = 16,
.gemm_n_block_size = 1,
.gemm_n_per_block = 4},
.epilogue = {.m_xdl_per_wave_per_shuffle = 1,
.n_per_wave_per_shuffle = 1,
.scalar_per_vector = 8},
},
};
@@ -232,34 +242,36 @@ constexpr InputOutputTileTransfer<> Transfer_4x32x1{
.a =
{
.thread_cluster = {.k0 = 4, .m_n = 32, .k1 = 1},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 2,
.src_scalar_per_vector = 16,
.lds_dst_scalar_per_vector = 16,
.is_direct_load = false,
.lds_padding = true},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 2,
.src_scalar_per_vector = 16,
.lds_dst_scalar_per_vector = 16,
.is_direct_load = false,
.lds_padding = true},
.thread_cluster_access_order = {1, 0, 2},
.src_access_order = {1, 0, 2},
},
.b =
{
.thread_cluster = {.k0 = 4, .m_n = 32, .k1 = 1},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 2,
.src_scalar_per_vector = 16,
.lds_dst_scalar_per_vector = 16,
.is_direct_load = false,
.lds_padding = true},
.lds_transfer = {.global_memory_vector_load_size = 8,
.src_vector_dim = 2,
.src_scalar_per_vector = 16,
.lds_dst_scalar_per_vector = 16,
.is_direct_load = false,
.lds_padding = true},
.thread_cluster_access_order = {1, 0, 2},
.src_access_order = {1, 0, 2},
},
.c =
{
.thread_cluster =
{.gemm_m_block_size = 1, .gemm_m_per_block = 32, .gemm_n_block_size = 1, .gemm_n_per_block = 4},
.epilogue = {.m_xdl_per_wave_per_shuffle = 1,
.n_per_wave_per_shuffle = 1,
.scalar_per_vector = 8},
.thread_cluster = {.gemm_m_block_size = 1,
.gemm_m_per_block = 32,
.gemm_n_block_size = 1,
.gemm_n_per_block = 4},
.epilogue = {.m_xdl_per_wave_per_shuffle = 1,
.n_per_wave_per_shuffle = 1,
.scalar_per_vector = 8},
},
};
@@ -267,68 +279,94 @@ constexpr InputOutputTileTransfer<> Transfer_4x32x1_vector_load_16_generic{
.a =
{
.thread_cluster = {.k0 = 4, .m_n = 32, .k1 = 1},
.lds_transfer = {.global_memory_vector_load_size = 16,
.src_vector_dim = 2,
.src_scalar_per_vector = 1,
.lds_dst_scalar_per_vector = 16,
.is_direct_load = false,
.lds_padding = true},
.lds_transfer = {.global_memory_vector_load_size = 16,
.src_vector_dim = 2,
.src_scalar_per_vector = 1,
.lds_dst_scalar_per_vector = 16,
.is_direct_load = false,
.lds_padding = true},
.thread_cluster_access_order = {1, 0, 2},
.src_access_order = {1, 0, 2},
},
.b =
{
.thread_cluster = {.k0 = 4, .m_n = 32, .k1 = 1},
.lds_transfer = {.global_memory_vector_load_size = 16,
.src_vector_dim = 2,
.src_scalar_per_vector = 1,
.lds_dst_scalar_per_vector = 16,
.is_direct_load = false,
.lds_padding = true},
.lds_transfer = {.global_memory_vector_load_size = 16,
.src_vector_dim = 2,
.src_scalar_per_vector = 1,
.lds_dst_scalar_per_vector = 16,
.is_direct_load = false,
.lds_padding = true},
.thread_cluster_access_order = {1, 0, 2},
.src_access_order = {1, 0, 2},
},
.c =
{
.thread_cluster =
{.gemm_m_block_size = 1, .gemm_m_per_block = 32, .gemm_n_block_size = 1, .gemm_n_per_block = 4},
.epilogue = {.m_xdl_per_wave_per_shuffle = 1,
.n_per_wave_per_shuffle = 1,
.scalar_per_vector = 1},
.thread_cluster = {.gemm_m_block_size = 1,
.gemm_m_per_block = 32,
.gemm_n_block_size = 1,
.gemm_n_per_block = 4},
.epilogue = {.m_xdl_per_wave_per_shuffle = 1,
.n_per_wave_per_shuffle = 1,
.scalar_per_vector = 1},
},
};
constexpr WarpGemmParams BwdGemmParams_Xdl_4x4_per_wave{
.matrix_instruction = MatrixInstructionType::XDL,
.gemm_m_per_instruction = 32, .gemm_n_per_instruction = 32, .gemm_m_iters_per_wave = 4, .gemm_n_iters_per_wave = 4};
constexpr WarpGemmParams BwdGemmParams_Xdl_4x4_per_wave{.matrix_instruction =
MatrixInstructionType::XDL,
.gemm_m_per_instruction = 32,
.gemm_n_per_instruction = 32,
.gemm_m_iters_per_wave = 4,
.gemm_n_iters_per_wave = 4};
constexpr WarpGemmParams BwdGemmParams_Xdl_1x1_per_wave{
.matrix_instruction = MatrixInstructionType::XDL,
.gemm_m_per_instruction = 32, .gemm_n_per_instruction = 32, .gemm_m_iters_per_wave = 1, .gemm_n_iters_per_wave = 1};
constexpr WarpGemmParams BwdGemmParams_Xdl_1x1_per_wave{.matrix_instruction =
MatrixInstructionType::XDL,
.gemm_m_per_instruction = 32,
.gemm_n_per_instruction = 32,
.gemm_m_iters_per_wave = 1,
.gemm_n_iters_per_wave = 1};
constexpr WarpGemmParams FwdGemmParams_Xdl_4x4_per_wave{
.matrix_instruction = MatrixInstructionType::XDL,
.gemm_m_per_instruction = 32, .gemm_n_per_instruction = 32, .gemm_m_iters_per_wave = 4, .gemm_n_iters_per_wave = 4};
constexpr WarpGemmParams FwdGemmParams_Xdl_4x4_per_wave{.matrix_instruction =
MatrixInstructionType::XDL,
.gemm_m_per_instruction = 32,
.gemm_n_per_instruction = 32,
.gemm_m_iters_per_wave = 4,
.gemm_n_iters_per_wave = 4};
constexpr WarpGemmParams FwdGemmParams_Xdl_4x2_per_wave{
.matrix_instruction = MatrixInstructionType::XDL,
.gemm_m_per_instruction = 32, .gemm_n_per_instruction = 32, .gemm_m_iters_per_wave = 4, .gemm_n_iters_per_wave = 2};
constexpr WarpGemmParams FwdGemmParams_Xdl_4x2_per_wave{.matrix_instruction =
MatrixInstructionType::XDL,
.gemm_m_per_instruction = 32,
.gemm_n_per_instruction = 32,
.gemm_m_iters_per_wave = 4,
.gemm_n_iters_per_wave = 2};
constexpr WarpGemmParams FwdGemmParams_Xdl_2x2_per_wave{
.matrix_instruction = MatrixInstructionType::XDL,
.gemm_m_per_instruction = 32, .gemm_n_per_instruction = 32, .gemm_m_iters_per_wave = 2, .gemm_n_iters_per_wave = 2};
constexpr WarpGemmParams FwdGemmParams_Xdl_2x2_per_wave{.matrix_instruction =
MatrixInstructionType::XDL,
.gemm_m_per_instruction = 32,
.gemm_n_per_instruction = 32,
.gemm_m_iters_per_wave = 2,
.gemm_n_iters_per_wave = 2};
constexpr WarpGemmParams FwdGemmParams_Xdl_2x1_per_wave{
.matrix_instruction = MatrixInstructionType::XDL,
.gemm_m_per_instruction = 32, .gemm_n_per_instruction = 32, .gemm_m_iters_per_wave = 2, .gemm_n_iters_per_wave = 1};
constexpr WarpGemmParams FwdGemmParams_Xdl_2x1_per_wave{.matrix_instruction =
MatrixInstructionType::XDL,
.gemm_m_per_instruction = 32,
.gemm_n_per_instruction = 32,
.gemm_m_iters_per_wave = 2,
.gemm_n_iters_per_wave = 1};
constexpr WarpGemmParams GemmParams_Wmma_16x16_2x1_per_wave{
.matrix_instruction = MatrixInstructionType::WMMA,
.gemm_m_per_instruction = 16, .gemm_n_per_instruction = 16, .gemm_m_iters_per_wave = 2, .gemm_n_iters_per_wave = 1};
constexpr WarpGemmParams GemmParams_Wmma_16x16_2x1_per_wave{.matrix_instruction =
MatrixInstructionType::WMMA,
.gemm_m_per_instruction = 16,
.gemm_n_per_instruction = 16,
.gemm_m_iters_per_wave = 2,
.gemm_n_iters_per_wave = 1};
constexpr WarpGemmParams GemmParams_Wmma_16x16_2x2_per_wave{
.matrix_instruction = MatrixInstructionType::WMMA,
.gemm_m_per_instruction = 16, .gemm_n_per_instruction = 16, .gemm_m_iters_per_wave = 2, .gemm_n_iters_per_wave = 2};
constexpr WarpGemmParams GemmParams_Wmma_16x16_2x2_per_wave{.matrix_instruction =
MatrixInstructionType::WMMA,
.gemm_m_per_instruction = 16,
.gemm_n_per_instruction = 16,
.gemm_m_iters_per_wave = 2,
.gemm_n_iters_per_wave = 2};
constexpr ThreadBlock ThreadBlock_256_256x256x32{.block_size = 256,
.tile_size = {.m = 256, .n = 256, .k = 32}};
@@ -357,19 +395,19 @@ constexpr ThreadBlock ThreadBlock_128_128x128x32{.block_size = 128,
constexpr ThreadBlock ThreadBlock_128_64x64x64{.block_size = 128,
.tile_size = {.m = 64, .n = 64, .k = 64}};
constexpr GemmPipeline BlockGemmDesc_v1_intrawave = {
.pipeline_version = PipelineVersion::V1, .scheduler = PipelineScheduler::INTRAWAVE};
constexpr GemmPipeline BlockGemmDesc_v1_intrawave = {.pipeline_version = PipelineVersion::V1,
.scheduler = PipelineScheduler::INTRAWAVE};
constexpr GemmPipeline BlockGemmDesc_v2_intrawave = {
.pipeline_version = PipelineVersion::V2, .scheduler = PipelineScheduler::INTRAWAVE};
constexpr GemmPipeline BlockGemmDesc_v2_intrawave = {.pipeline_version = PipelineVersion::V2,
.scheduler = PipelineScheduler::INTRAWAVE};
constexpr GemmPipeline BlockGemmDesc_v3_intrawave = {
.pipeline_version = PipelineVersion::V3, .scheduler = PipelineScheduler::INTRAWAVE};
constexpr GemmPipeline BlockGemmDesc_v3_intrawave = {.pipeline_version = PipelineVersion::V3,
.scheduler = PipelineScheduler::INTRAWAVE};
constexpr GemmPipeline BlockGemmDesc_v4_intrawave = {
.pipeline_version = PipelineVersion::V4, .scheduler = PipelineScheduler::INTRAWAVE};
constexpr GemmPipeline BlockGemmDesc_v4_intrawave = {.pipeline_version = PipelineVersion::V4,
.scheduler = PipelineScheduler::INTRAWAVE};
constexpr GemmPipeline BlockGemmDesc_v5_intrawave = {
.pipeline_version = PipelineVersion::V5, .scheduler = PipelineScheduler::INTRAWAVE};
constexpr GemmPipeline BlockGemmDesc_v5_intrawave = {.pipeline_version = PipelineVersion::V5,
.scheduler = PipelineScheduler::INTRAWAVE};
} // namespace ck_tile::builder::test_utils

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

@@ -98,7 +98,8 @@ template <>
inline std::string to_string<GemmPipeline>(GemmPipeline t)
{
std::ostringstream oss;
oss << t.num_gemm_k_prefetch_stages << "," << t.num_conv_groups_to_merge << "," << 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();
}
@@ -123,17 +124,17 @@ inline std::string to_string(InputThreadCluster<ThreadClusterRank> t)
template <>
inline std::string to_string<OutputThreadCluster>(OutputThreadCluster t)
{
return array_to_seq(
std::array<size_t, 4>{t.gemm_m_block_size, t.gemm_m_per_block, t.gemm_n_block_size, t.gemm_n_per_block});
return array_to_seq(std::array<size_t, 4>{
t.gemm_m_block_size, t.gemm_m_per_block, t.gemm_n_block_size, t.gemm_n_per_block});
}
template <>
inline std::string to_string<LdsTransfer>(LdsTransfer t)
{
std::ostringstream oss;
oss << t.global_memory_vector_load_size << "," << t.src_vector_dim << "," << t.src_scalar_per_vector << "," << t.lds_dst_scalar_per_vector
<< "," << (t.lds_padding ? "true" : "false") << ","
<< (t.is_direct_load ? "true" : "false");
oss << t.global_memory_vector_load_size << "," << t.src_vector_dim << ","
<< t.src_scalar_per_vector << "," << t.lds_dst_scalar_per_vector << ","
<< (t.lds_padding ? "true" : "false") << "," << (t.is_direct_load ? "true" : "false");
return oss.str();
}
@@ -259,7 +260,6 @@ inline std::string to_string<WarpGemm_>(WarpGemm_ t)
return to_string(t.warp_gemm);
}
template <size_t ThreadClusterRank = 3>
inline std::string to_string(InputOutputTileTransfer_<ThreadClusterRank> t)
{
@@ -319,21 +319,21 @@ inline std::string to_string<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_CShuf
ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_CShuffle t)
{
std::ostringstream oss;
if (t.warp_gemm.matrix_instruction == MatrixInstructionType::WMMA)
if(t.warp_gemm.matrix_instruction == MatrixInstructionType::WMMA)
{
oss << to_string(static_cast<ThreadBlock_>(t)) << ","
<< t.transfer.a.lds_transfer.global_memory_vector_load_size << ","
<< to_string(static_cast<WarpGemm_>(t))
<< "," << to_string(static_cast<InputOutputTileTransfer_<>>(t));
<< to_string(static_cast<WarpGemm_>(t)) << ","
<< to_string(static_cast<InputOutputTileTransfer_<>>(t));
}
else
else
{
oss << to_string(static_cast<ThreadBlock_>(t)) << ","
<< t.transfer.a.lds_transfer.global_memory_vector_load_size << ","
<< t.transfer.b.lds_transfer.global_memory_vector_load_size << ","
<< to_string(static_cast<WarpGemm_>(t))
<< "," << to_string(static_cast<InputOutputTileTransfer_<>>(t));
}
<< to_string(static_cast<WarpGemm_>(t)) << ","
<< to_string(static_cast<InputOutputTileTransfer_<>>(t));
}
return oss.str();
}
@@ -342,11 +342,11 @@ inline std::string to_string<ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_CShuf
ConvAlgorithm_DeviceGroupedConvFwdMultipleABD_CShuffle_V3 t)
{
std::ostringstream oss;
oss << to_string(static_cast<ThreadBlock_>(t)) << ","
oss << to_string(static_cast<ThreadBlock_>(t)) << ","
<< t.transfer.a.lds_transfer.global_memory_vector_load_size << ","
<< t.transfer.b.lds_transfer.global_memory_vector_load_size << ","
<< to_string(static_cast<WarpGemm_>(t))
<< "," << to_string(static_cast<InputOutputTileTransfer_<>>(t));
<< to_string(static_cast<WarpGemm_>(t)) << ","
<< to_string(static_cast<InputOutputTileTransfer_<>>(t));
return oss.str();
}
@@ -367,11 +367,11 @@ inline std::string to_string<ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Xdl_CSh
ConvAlgorithm_DeviceGroupedConvFwdMultipleD_Xdl_CShuffle_Large_Tensor t)
{
std::ostringstream oss;
oss << to_string(static_cast<ThreadBlock_>(t)) << ","
oss << to_string(static_cast<ThreadBlock_>(t)) << ","
<< t.transfer.a.lds_transfer.global_memory_vector_load_size << ","
<< t.transfer.b.lds_transfer.global_memory_vector_load_size << ","
<< to_string(static_cast<WarpGemm_>(t))
<< "," << to_string(static_cast<InputOutputTileTransfer_<>>(t));
<< to_string(static_cast<WarpGemm_>(t)) << ","
<< to_string(static_cast<InputOutputTileTransfer_<>>(t));
return oss.str();
}
@@ -380,10 +380,10 @@ inline std::string to_string<ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuff
ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuffle t)
{
std::ostringstream oss;
oss << to_string(static_cast<ThreadBlock_>(t)) << ","
oss << to_string(static_cast<ThreadBlock_>(t)) << ","
<< t.transfer.a.lds_transfer.global_memory_vector_load_size << ","
<< to_string(static_cast<WarpGemm_>(t))
<< "," << to_string(static_cast<InputOutputTileTransfer_<4>>(t));
<< to_string(static_cast<WarpGemm_>(t)) << ","
<< to_string(static_cast<InputOutputTileTransfer_<4>>(t));
return oss.str();
}
@@ -392,10 +392,10 @@ inline std::string to_string<ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuff
ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuffle_V3 t)
{
std::ostringstream oss;
oss << to_string(static_cast<ThreadBlock_>(t)) << ","
oss << to_string(static_cast<ThreadBlock_>(t)) << ","
<< t.transfer.a.lds_transfer.global_memory_vector_load_size << ","
<< to_string(static_cast<WarpGemm_>(t))
<< "," << to_string(static_cast<InputOutputTileTransfer_<>>(t));
<< to_string(static_cast<WarpGemm_>(t)) << ","
<< to_string(static_cast<InputOutputTileTransfer_<>>(t));
return oss.str();
}
@@ -404,10 +404,10 @@ inline std::string to_string<ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuf
ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle t)
{
std::ostringstream oss;
oss << to_string(static_cast<ThreadBlock_>(t)) << ","
oss << to_string(static_cast<ThreadBlock_>(t)) << ","
<< t.transfer.a.lds_transfer.global_memory_vector_load_size << ","
<< to_string(static_cast<WarpGemm_>(t))
<< "," << to_string(static_cast<InputOutputTileTransfer_<>>(t));
<< to_string(static_cast<WarpGemm_>(t)) << ","
<< to_string(static_cast<InputOutputTileTransfer_<>>(t));
return oss.str();
}
@@ -416,10 +416,10 @@ inline std::string to_string<ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuf
ConvAlgorithm_DeviceGroupedConvBwdWeight_Wmma_CShuffle_V3 t)
{
std::ostringstream oss;
oss << to_string(static_cast<ThreadBlock_>(t)) << ","
oss << to_string(static_cast<ThreadBlock_>(t)) << ","
<< t.transfer.a.lds_transfer.global_memory_vector_load_size << ","
<< to_string(static_cast<WarpGemm_>(t))
<< "," << to_string(static_cast<InputOutputTileTransfer_<>>(t));
<< to_string(static_cast<WarpGemm_>(t)) << ","
<< to_string(static_cast<InputOutputTileTransfer_<>>(t));
return oss.str();
}
@@ -430,8 +430,8 @@ inline std::string to_string<ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_W
std::ostringstream oss;
oss << to_string(static_cast<ThreadBlock_>(t)) << ","
<< t.transfer.a.lds_transfer.global_memory_vector_load_size << ","
<< to_string(static_cast<WarpGemm_>(t))
<< "," << to_string(static_cast<InputOutputTileTransfer_<>>(t));
<< to_string(static_cast<WarpGemm_>(t)) << ","
<< to_string(static_cast<InputOutputTileTransfer_<>>(t));
return oss.str();
}
@@ -441,14 +441,13 @@ inline std::string to_string<ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_C
ConvAlgorithm_DeviceGroupedConvBwdWeight_TwoStage_CShuffle_V3 t)
{
std::ostringstream oss;
oss << to_string(static_cast<ThreadBlock_>(t))
<< "," << t.transfer.a.lds_transfer.global_memory_vector_load_size
<< "," << to_string(static_cast<WarpGemm_>(t))
<< "," << to_string(static_cast<InputOutputTileTransfer_<>>(t));
oss << to_string(static_cast<ThreadBlock_>(t)) << ","
<< t.transfer.a.lds_transfer.global_memory_vector_load_size << ","
<< to_string(static_cast<WarpGemm_>(t)) << ","
<< to_string(static_cast<InputOutputTileTransfer_<>>(t));
return oss.str();
}
template <>
inline std::string to_string<ConvAlgorithm_DeviceGroupedConvBwdWeight_Dl>(
ConvAlgorithm_DeviceGroupedConvBwdWeight_Dl t)
@@ -466,10 +465,10 @@ inline std::string to_string<ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_X
ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle t)
{
std::ostringstream oss;
oss << to_string(static_cast<ThreadBlock_>(t)) << ","
oss << to_string(static_cast<ThreadBlock_>(t)) << ","
<< t.transfer.a.lds_transfer.global_memory_vector_load_size << ","
<< to_string(static_cast<WarpGemm_>(t))
<< "," << to_string(static_cast<InputOutputTileTransfer_<4>>(t));
<< to_string(static_cast<WarpGemm_>(t)) << ","
<< to_string(static_cast<InputOutputTileTransfer_<4>>(t));
return oss.str();
}