ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-07-15 11:34:54 +00:00
Code Issues Packages Projects Releases Wiki Activity

clang-format

Browse Source
This commit is contained in:
Ville Pietilä
2026-01-13 04:20:08 -05:00
parent a8f1d44078
commit bf57fbf488
3 changed files with 25 additions and 26 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -9,7 +9,7 @@ namespace ck_tile::builder::factory {
// Base algorithm concepts
template <typename T, size_t ThreadClusterRank = 3>
concept TileTransferParameters =
concept TileTransferParameters =
SpecifiesBlockTransfer<T, ThreadClusterRank> && SpecifiesLdsTransfer<T> &&
SpecifiesThreadClusterAccessOrder<T> && SpecifiesSourceAccessOrder<T>;
@@ -19,36 +19,34 @@ concept SpecifiesTileTransferParameters3D = TileTransferParameters<T, 3>;
template <typename T>
concept SpecifiesTileTransferParameters4D = TileTransferParameters<T, 4>;
template <typename T>
concept FwdXdlAlgorithmBase =
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesTileTransferParameters3D<T> &&
SpecifiesGridwiseFwdXdlGemm<T> && SpecifiesFwdConvSpecialization<T> && SpecifiesGemmSpecialization<T> &&
SpecifiesNumPrefetchStages<T> && SpecifiesNumGroupsToMerge<T> && SpecifiesLoopScheduler<T>;
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesTileTransferParameters3D<T> &&
SpecifiesGridwiseFwdXdlGemm<T> && SpecifiesFwdConvSpecialization<T> &&
SpecifiesGemmSpecialization<T> && SpecifiesNumPrefetchStages<T> &&
SpecifiesNumGroupsToMerge<T> && SpecifiesLoopScheduler<T>;
template <typename T>
concept BwdXdlAlgorithmBase =
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> &&
SpecifiesTileTransferParameters4D<T> && SpecifiesGridwiseBwdXdlGemm<T> &&
SpecifiesBwdWeightConvSpecialization<T>;
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesTileTransferParameters4D<T> &&
SpecifiesGridwiseBwdXdlGemm<T> && SpecifiesBwdWeightConvSpecialization<T>;
template <typename T>
concept BwdXdlV3AlgorithmBase =
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> &&
SpecifiesTileTransferParameters3D<T> && SpecifiesGridwiseBwdXdlGemm<T> &&
SpecifiesBwdWeightConvSpecialization<T> && SpecifiesBlockGemm<T>;
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesTileTransferParameters3D<T> &&
SpecifiesGridwiseBwdXdlGemm<T> && SpecifiesBwdWeightConvSpecialization<T> &&
SpecifiesBlockGemm<T>;
template <typename T>
concept BwdWmmaAlgorithmBase =
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> &&
SpecifiesTileTransferParameters3D<T> && SpecifiesGridwiseWmmaGemm<T> &&
SpecifiesBwdWeightConvSpecialization<T>;
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesTileTransferParameters3D<T> &&
SpecifiesGridwiseWmmaGemm<T> && SpecifiesBwdWeightConvSpecialization<T>;
template <typename T>
concept BwdWmmaV3AlgorithmBase =
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> &&
SpecifiesTileTransferParameters3D<T> && SpecifiesGridwiseWmmaGemm<T> &&
SpecifiesBwdWeightConvSpecialization<T> && SpecifiesBlockGemm<T>;
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesTileTransferParameters3D<T> &&
SpecifiesGridwiseWmmaGemm<T> && SpecifiesBwdWeightConvSpecialization<T> &&
SpecifiesBlockGemm<T>;
// Reference algorithm concept
template <typename T>
@@ -69,17 +67,17 @@ concept LargeTensorAlgorithm = FwdXdlAlgorithmBase<T> && SpecifiesLargeTensorSup
template <typename T>
concept FwdXdlV3Algorithm =
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> &&
SpecifiesTileTransferParameters3D<T> && SpecifiesGridwiseFwdXdlGemm<T> &&
SpecifiesFwdConvSpecialization<T> && SpecifiesGemmSpecialization<T> && SpecifiesBlockGemm<T>;
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesTileTransferParameters3D<T> &&
SpecifiesGridwiseFwdXdlGemm<T> && SpecifiesFwdConvSpecialization<T> &&
SpecifiesGemmSpecialization<T> && SpecifiesBlockGemm<T>;
// FWD WMMA algorithm concepts
template <typename T>
concept FwdWmmaAlgorithm =
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> &&
SpecifiesTileTransferParameters3D<T> && SpecifiesGridwiseWmmaGemm<T> &&
SpecifiesFwdConvSpecialization<T> && SpecifiesGemmSpecialization<T> &&
SpecifiesNumPrefetchStages<T> && SpecifiesLoopScheduler<T> && SpecifiesGridwiseGemmPipeline<T>;
ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> && SpecifiesTileTransferParameters3D<T> &&
SpecifiesGridwiseWmmaGemm<T> && SpecifiesFwdConvSpecialization<T> &&
SpecifiesGemmSpecialization<T> && SpecifiesNumPrefetchStages<T> && SpecifiesLoopScheduler<T> &&
SpecifiesGridwiseGemmPipeline<T>;
// FWD DL algorithms
template <typename T>

2
experimental/builder/include/ck_tile/builder/factory/helpers/ck/conv_tensor_layout.hpp
View File

@@ -224,7 +224,7 @@ struct ConvTensorLayouts
{
using InLayout = decltype(TensorLayoutToCK<Signature.input.config.layout>());
using WeiLayout = decltype(TensorLayoutToCK<Signature.weight.config.layout>());
using OutLayout = decltype(TensorLayoutToCK<Signature.output.config.layout>());
using OutLayout = decltype(TensorLayoutToCK<Signature.output.config.layout>());
using DsLayout = decltype(GetAuxiliaryTensorLayouts<Signature, SPATIAL_DIM>())::type;
};

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

@@ -133,7 +133,8 @@ inline std::string to_string(BlockTransfer<ThreadClusterRank> t)
}
else
{
static_assert(ThreadClusterRank == 3 || ThreadClusterRank == 4, "Unsupported ThreadClusterRank");
static_assert(ThreadClusterRank == 3 || ThreadClusterRank == 4,
"Unsupported ThreadClusterRank");
}
}