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

Fixes problem descritpor

Browse Source
This commit is contained in:
Bartlomiej Kocot
2025-08-07 10:28:54 +00:00
parent 4d2ecedab2
commit 55a2a4cac7
5 changed files with 70 additions and 25 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
experimental/convolution_builder/CMakeLists.txt
View File

@@ -1,3 +1,7 @@
set(CMAKE_CXX_STANDARD 20)
add_executable(convolution_example convolution_example.cpp)
add_executable(convolution_example convolution_example.cpp)
set(EXAMPLE_COMPILE_OPTIONS)
list(APPEND EXAMPLE_COMPILE_OPTIONS -g -fverbose-asm --save-temps -Wno-gnu-line-marker)
target_compile_options(convolution_example PRIVATE ${EXAMPLE_COMPILE_OPTIONS})

54
experimental/convolution_builder/convolution_builder.hpp
View File

@@ -1,31 +1,55 @@
#pragma once
#include "convolution_solution_descriptor.hpp"
#include "convolution_problem_descriptor.hpp"
#include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_bwd_weight_two_stage_xdl_cshuffle.hpp"
#include "ck/tensor_operation/gpu/device/impl/device_grouped_conv_fwd_multiple_abd_xdl_cshuffle.hpp"
template<SolutionDescriptorV1 Solution>
struct ConvolutionBuilder {
template<typename Problem>
struct ConvolutionBuilder;
static constexpr auto GetSolution() {
if constexpr(Solution::GemmImplementationType_ == GemmImplementationType::XDL) {
if constexpr(Solution::ConvolutionDirection_ == ConvolutionDirection::Forward) {
return ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle< 2, ck::tensor_layout::convolution::NHWGC, ck::tensor_layout::convolution::GKYXC, ck::Tuple<>, ck::tensor_layout::convolution::NHWGK, float, float, float, float, ck::Tuple<>, float, ck::tensor_operation::element_wise::PassThrough, ck::tensor_operation::element_wise::PassThrough, ck::tensor_operation::element_wise::PassThrough, ck::tensor_operation::device::ConvolutionForwardSpecialization::Default, ck::tensor_operation::device::GemmSpecialization::MNKPadding, 1, 256, 128, 256, 32, 8, 8, 32, 32, 2, 4, ck::Sequence<4, 64, 1>, ck::Sequence<1, 0, 2>, ck::Sequence<1, 0, 2>, 2, 8, 8, 1, ck::Sequence<4, 64, 1>, ck::Sequence<1, 0, 2>, ck::Sequence<1, 0, 2>, 2, 8, 8, 1, 1, 1, ck::Sequence<1, 32, 1, 8>, 8>{};
} else if constexpr(Solution::ConvolutionDirection_ == ConvolutionDirection::BackwardData) {
return std::tuple<>{};
} else if constexpr(Solution::ConvolutionDirection_ == ConvolutionDirection::BackwardWeight) {
return ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle<2, ck::tensor_layout::convolution::NHWGC, ck::tensor_layout::convolution::GKYXC, ck::tensor_layout::convolution::NHWGK, float, float, float, float, ck::tensor_operation::element_wise::PassThrough, ck::tensor_operation::element_wise::PassThrough, ck::tensor_operation::element_wise::PassThrough, ck::tensor_operation::device::ConvolutionBackwardWeightSpecialization::Default, 64, 16, 16, 32, 8, 16, 16, 1, 1, ck::Sequence<4, 8, 1>, ck::Sequence<2, 0, 1>, ck::Sequence<1, 0, 2>, 1, 1, 4, false, ck::Sequence<4, 8, 1>, ck::Sequence<2, 0, 1>, ck::Sequence<1, 0, 2>, 1, 1, 4, false, 1, 1, ck::Sequence<1, 8, 1, 8>, 1>{};
template<ProblemDescriptorV1 Problem>
struct ConvolutionBuilder<Problem> {
enum class Kernel {
GroupedConvFwdMultipleABD_Xdl_CShuffle,
GroupedConvBwdWeightTwoStage_Xdl_CShuffle
};
static constexpr Kernel GetKernel() {
if constexpr(Problem::GemmImplementationType_ == GemmImplementationType::XDL) {
if constexpr(Problem::ConvolutionDirection_ == ConvolutionDirection::Forward) {
return Kernel::GroupedConvFwdMultipleABD_Xdl_CShuffle;
} else if constexpr(Problem::ConvolutionDirection_ == ConvolutionDirection::BackwardData) {
static_assert("Instance not found!");
} else if constexpr(Problem::ConvolutionDirection_ == ConvolutionDirection::BackwardWeight) {
return Kernel::GroupedConvBwdWeightTwoStage_Xdl_CShuffle;
} else {
return std::tuple<>{};
static_assert("Instance not found!");
}
} else {
return std::tuple<>{};
static_assert("Instance not found!");
}
}
static constexpr auto GetInstance() {
using FwdInstance = ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle< 2, ck::tensor_layout::convolution::NHWGC, ck::tensor_layout::convolution::GKYXC, ck::Tuple<>, ck::tensor_layout::convolution::NHWGK, float, float, float, float, ck::Tuple<>, float, ck::tensor_operation::element_wise::PassThrough, ck::tensor_operation::element_wise::PassThrough, ck::tensor_operation::element_wise::PassThrough, ck::tensor_operation::device::ConvolutionForwardSpecialization::Default, ck::tensor_operation::device::GemmSpecialization::MNKPadding, 1, 256, 128, 256, 32, 8, 8, 32, 32, 2, 4, ck::Sequence<4, 64, 1>, ck::Sequence<1, 0, 2>, ck::Sequence<1, 0, 2>, 2, 8, 8, 1, ck::Sequence<4, 64, 1>, ck::Sequence<1, 0, 2>, ck::Sequence<1, 0, 2>, 2, 8, 8, 1, 1, 1, ck::Sequence<1, 32, 1, 8>, 8>;
using BwdWeightInstance = ck::tensor_operation::device::DeviceGroupedConvBwdWeightTwoStage_Xdl_CShuffle<2, ck::tensor_layout::convolution::NHWGC, ck::tensor_layout::convolution::GKYXC, ck::tensor_layout::convolution::NHWGK, float, float, float, float, ck::tensor_operation::element_wise::PassThrough, ck::tensor_operation::element_wise::PassThrough, ck::tensor_operation::element_wise::PassThrough, ck::tensor_operation::device::ConvolutionBackwardWeightSpecialization::Default, 64, 16, 16, 32, 8, 16, 16, 1, 1, ck::Sequence<4, 8, 1>, ck::Sequence<2, 0, 1>, ck::Sequence<1, 0, 2>, 1, 1, 4, false, ck::Sequence<4, 8, 1>, ck::Sequence<2, 0, 1>, ck::Sequence<1, 0, 2>, 1, 1, 4, false, 1, 1, ck::Sequence<1, 8, 1, 8>, 1>;
using SelectedInstance = std::conditional_t<GetKernel() == Kernel::GroupedConvFwdMultipleABD_Xdl_CShuffle, FwdInstance, BwdWeightInstance>;
return SelectedInstance{};
}
std::string GetKernelName() const
std::string GetInstanceName() const
{
return GetSolution().GetTypeString();
if constexpr(GetKernel() == Kernel::GroupedConvFwdMultipleABD_Xdl_CShuffle) {
return "GroupedConvFwdMultipleABD_Xdl_CShuffle";
} else if constexpr (GetKernel() == Kernel::GroupedConvBwdWeightTwoStage_Xdl_CShuffle) {
return "GroupedConvBwdWeightTwoStage_Xdl_CShuffle";
} else {
return "Not found.";
}
// const auto instance = GetInstance();
// return instance.GetTypeString();
}
};

14
experimental/convolution_builder/convolution_example.cpp
View File

@@ -4,19 +4,21 @@
#include "convolution_builder.hpp"
// Example of solution description for Forward Conv with default settings
struct GroupedConvFwdXdlImplicitGemm : public GroupedConvBaseXdl {
// Example of problem description for Forward Conv with default settings
struct GroupedConvFwdXdlImplicitGemm : public GroupedConvBaseXdlV1 {
static constexpr ConvolutionDirection ConvolutionDirection_ = ConvolutionDirection::Forward;
};
// Example of solution description for Backward Weight Conv with default settings and Split K Two Stage
struct GroupedConvBwdWeightXdlImplicitGemm : public GroupedConvBaseXdl {
// Example of problem description for Backward Weight Conv with default settings and Split K Two Stage
struct GroupedConvBwdWeightXdlImplicitGemmTwoStage : public GroupedConvBaseXdlV1 {
static constexpr ConvolutionDirection ConvolutionDirection_ = ConvolutionDirection::BackwardWeight;
static constexpr SplitKSupport SplitKSupport_ = SplitKSupport::SupportedTwoStage;
};
int main () {
ConvolutionBuilder<GroupedConvFwdXdlImplicitGemm> builder;
std::cout << builder.GetKernelName() << std::endl;
ConvolutionBuilder<GroupedConvFwdXdlImplicitGemm> builder_fwd;
std::cout << builder_fwd.GetInstanceName() << std::endl;
ConvolutionBuilder<GroupedConvBwdWeightXdlImplicitGemmTwoStage> builder_bwd_weight_two_stage;
std::cout << builder_bwd_weight_two_stage.GetInstanceName() << std::endl;
return 0;
}

5
experimental/convolution_builder/convolution_implementation_descriptor.hpp Normal file
View File

@@ -0,0 +1,5 @@
#pragma once
#include <concepts>
template <typename T>
concept

16
experimental/convolution_builder/convolution_solution_descriptor.hpp → experimental/convolution_builder/convolution_problem_descriptor.hpp
View File

@@ -1,7 +1,11 @@
#pragma once
#include <iostream>
#include <concepts>
enum class ProblemDescriptorVersion
{
V1
};
enum class GemmImplementationType
{
XDL,
@@ -79,7 +83,8 @@ enum class ElementwiseOperation {
template <typename T>
concept SolutionDescriptorV1 = requires {
concept ProblemDescriptorV1 = requires {
{T::ProblemDescriptorVersion_} -> std::convertible_to<ProblemDescriptorVersion>;
{T::GemmImplementationType_} -> std::convertible_to<GemmImplementationType>;
{T::ConvolutionDirection_} -> std::convertible_to<ConvolutionDirection>;
{T::GemmPipelineVersion_} -> std::convertible_to<const GemmPipelineVersion>;
@@ -89,7 +94,7 @@ concept SolutionDescriptorV1 = requires {
{T::LargeTensorSupport_} -> std::convertible_to<const LargeTensorSupport>;
{T::ImplementationType_} -> std::convertible_to<const ImplementationType>;
{T::ElementwiseOperation_} -> std::convertible_to<const ElementwiseOperation>;
};
} && (T::ProblemDescriptorVersion_ == ProblemDescriptorVersion::V1);
struct GroupedConvBase {
static constexpr GemmPipelineVersion GemmPipelineVersion_ = GemmPipelineVersion::V1;
@@ -104,3 +109,8 @@ struct GroupedConvBase {
struct GroupedConvBaseXdl : public GroupedConvBase {
static constexpr GemmImplementationType GemmImplementationType_ = GemmImplementationType::XDL;
};
struct GroupedConvBaseXdlV1 : public GroupedConvBaseXdl {
static constexpr ProblemDescriptorVersion ProblemDescriptorVersion_ = ProblemDescriptorVersion::V1;
};