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

Fix GEMM pipeline concept usage.

Browse Source
This commit is contained in:
Ville Pietilä
2026-01-16 05:32:29 -05:00
parent 346b3fa5bd
commit e5be038a5e
3 changed files with 9 additions and 15 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
experimental/builder/include/ck_tile/builder/factory/README.md
View File

@@ -49,14 +49,13 @@ XDL and WMMA algorithms share a common base, while DL algorithms have their own
Both XDL and WMMA algorithms share the following foundational concepts:
```
ConvAlgorithm (Base Concept)
ConvWarpGemmAlgorithm (Base Concept)
│
│ Requirements:
│ • ConvAlgorithmDescriptor
│ • SpecifiesThreadBlock
│ • SpecifiesTileTransferParameters (ThreadClusters, LdsTransfer, AccessOrders)
│ • SpecifiesWarpGemm
│ • SpecifiesGemmPipeline
│
├─── FwdAlgorithm (Forward Convolution)
│ │
@@ -64,7 +63,7 @@ ConvAlgorithm (Base Concept)
│ │
│ └─── FwdAlgorithmV3
│ │
│ │ Additional: SpecifiesPipelineV3
│ │ Additional: SpecifiesPipelineV3 + SpecifiesGemmPipeline
│ │
│
└─── BwdAlgorithm (Backward Weight Convolution)
@@ -73,7 +72,7 @@ ConvAlgorithm (Base Concept)
│
└─── BwdAlgorithmV3
│
│ Additional: SpecifiesPipelineV3
│ Additional: SpecifiesPipelineV3 + SpecifiesGemmPipeline
│
```

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

@@ -14,23 +14,22 @@ concept SpecifiesTileTransferParameters =
// Base algorithm concepts
template <typename T, size_t ThreadClusterRank = 3>
concept ConvAlgorithm = ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> &&
concept ConvWarpGemmAlgorithm = ConvAlgorithmDescriptor<T> && SpecifiesThreadBlock<T> &&
SpecifiesTileTransferParameters<T, ThreadClusterRank> &&
SpecifiesWarpGemm<T> && SpecifiesGemmPipeline<T>;
;
SpecifiesWarpGemm<T>;
template <typename T>
concept FwdAlgorithm = ConvAlgorithm<T, 3> && SpecifiesFwdConvSpecialization<T>;
concept FwdAlgorithm = ConvWarpGemmAlgorithm<T, 3> && SpecifiesFwdConvSpecialization<T>;
template <typename T>
concept FwdAlgorithmV3 = FwdAlgorithm<T> && SpecifiesPipelineV3<T>;
concept FwdAlgorithmV3 = FwdAlgorithm<T> && SpecifiesPipelineV3<T> && SpecifiesGemmPipeline<T>;
template <typename T, size_t ThreadClusterRank = 3>
concept BwdAlgorithm =
ConvAlgorithm<T, ThreadClusterRank> && SpecifiesBwdWeightConvSpecialization<T>;
ConvWarpGemmAlgorithm<T, ThreadClusterRank> && SpecifiesBwdWeightConvSpecialization<T>;
template <typename T>
concept BwdAlgorithmV3 = BwdAlgorithm<T, 3> && SpecifiesPipelineV3<T>;
concept BwdAlgorithmV3 = BwdAlgorithm<T, 3> && SpecifiesPipelineV3<T> && SpecifiesGemmPipeline<T>;
template <typename T>
concept DlAlgorithm =

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

@@ -543,8 +543,6 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeight_Xdl_CShuffle =
WarpGemm_,
InputOutputTileTransfer_<4>,
ConvSpecializationBwdWeight_,
GemmPipeline_, // Not needed, but we need this to meet the ConvAlgorithm
// concept.
TransposeParams_,
AlgorithmSpecialization_<>>;
@@ -594,8 +592,6 @@ using ConvAlgorithm_DeviceGroupedConvBwdWeightMultipleD_Xdl_CShuffle =
ConvAlgorithmTemplate<ThreadBlock_,
WarpGemm_,
InputOutputTileTransfer_<4>,
GemmPipeline_, // Not needed, but we need this to meet the ConvAlgorithm
// concept.
ConvSpecializationBwdWeight_,
AlgorithmSpecialization_<MULTIPLE_D>>;