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

Implementation descriptor

Browse Source
This commit is contained in:
Bartlomiej Kocot
2025-08-08 14:00:05 +00:00
parent 55a2a4cac7
commit 1b47cea817
3 changed files with 195 additions and 21 deletions
Download Patch File Download Diff File Expand all files Collapse all files

15
experimental/convolution_builder/convolution_example.cpp
View File

@@ -7,6 +7,7 @@
// Example of problem description for Forward Conv with default settings
struct GroupedConvFwdXdlImplicitGemm : public GroupedConvBaseXdlV1 {
static constexpr ConvolutionDirection ConvolutionDirection_ = ConvolutionDirection::Forward;
static constexpr ElementwiseOperation ElementwiseOperation_ = ElementwiseOperation::Bias;
};
// Example of problem description for Backward Weight Conv with default settings and Split K Two Stage
@@ -15,10 +16,20 @@ struct GroupedConvBwdWeightXdlImplicitGemmTwoStage : public GroupedConvBaseXdlV1
static constexpr SplitKSupport SplitKSupport_ = SplitKSupport::SupportedTwoStage;
};
struct ImplementationDescriptor : public NHWCImplementationBaseV1, public BF16ImplementationBaseV1 {
static constexpr ck::index_t BlockSize_ = 64;
static constexpr auto TileSizes_ = std::make_tuple(16, 16, 32);
static constexpr ck::index_t K1_ = 8;
static constexpr MFMAInstructionSize MFMAInstructionSize_ = MFMAInstructionSize::M16N16;
static constexpr auto XdlPerWave_ = std::make_tuple(16, 16);
static constexpr auto GlobalTransferVectorSize_ = std::make_tuple(1, 1, 1);
static constexpr auto LDSStoreVectorSize_ = std::make_tuple(4, 4);
};
int main () {
ConvolutionBuilder<GroupedConvFwdXdlImplicitGemm> builder_fwd;
ConvolutionBuilder<GroupedConvFwdXdlImplicitGemm, ImplementationDescriptor> builder_fwd;
std::cout << builder_fwd.GetInstanceName() << std::endl;
ConvolutionBuilder<GroupedConvBwdWeightXdlImplicitGemmTwoStage> builder_bwd_weight_two_stage;
ConvolutionBuilder<GroupedConvBwdWeightXdlImplicitGemmTwoStage, ImplementationDescriptor> builder_bwd_weight_two_stage;
std::cout << builder_bwd_weight_two_stage.GetInstanceName() << std::endl;
return 0;
}