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

Fix test files for convolution builder.

Browse Source
This commit is contained in:
John Shumway
2025-08-28 02:01:46 +00:00
parent 897f966df6
commit cee90b800e
3 changed files with 45 additions and 94 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
experimental/builder/test/CMakeLists.txt
View File

@@ -1,13 +1,8 @@
set(CMAKE_CXX_STANDARD 20)
add_executable(gemm_example gemm_example.cpp)
target_include_directories(gemm_example PRIVATE
"${PROJECT_SOURCE_DIR}/include"
)
include(gtest)
add_executable(test_conv_builder builder.cpp)
add_executable(test_conv_builder test_conv_builder.cpp)
target_include_directories(test_conv_builder PRIVATE
"${PROJECT_SOURCE_DIR}/experimental/builder/include"
"${PROJECT_SOURCE_DIR}/include"

88
experimental/builder/test/builder.cpp
View File

@@ -1,88 +0,0 @@
#include <gtest/gtest.h>
#include <ck_tile/builder/conv_builder.hpp>
namespace {
namespace ckb = ck_tile::builder;
// Example of kernel 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 kernel description for Backward Weight Conv with default settings and Split K Two
// Stage
struct GroupedConvBwdWeightXdlImplicitGemmTwoStage : public GroupedConvBaseXdlV1
{
[[maybe_unused]] static constexpr ConvolutionDirection ConvolutionDirection_ =
ConvolutionDirection::BackwardWeight;
[[maybe_unused]] static constexpr SplitKSupport SplitKSupport_ =
SplitKSupport::SupportedTwoStage;
};
struct Implementation16x16 : ImplementationDefaultV1
{
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);
};
struct ProblemBF16NHWGC : public BF16ProblemBaseV1, public NHWGCProblemBaseV1
{
};
TEST(ConvBuilderTest, TestBuilderV0_0_0)
{
ConvolutionBuilder<GroupedConvFwdXdlImplicitGemm, ProblemBF16NHWGC, Implementation16x16>
builder_fwd;
EXPECT_EQ(builder_fwd.GetInstanceName(),
"GroupedConvFwdMultipleABD_Xdl_CShuffle<64, 16, 16, 32, Default, 8, 16x16, 16, 16, "
"1, 4, 1, 4, 1, Intrawave, v1, 1>");
// It would be nice if this worked, but it fails.
// [[maybe_unused]] auto instance = builder_fwd.GetInstance();
}
struct FwdConvSignature
{
static constexpr int SPATIAL_DIM = 2;
static constexpr auto DIRECTION = ckb::ConvDirection::Forward;
static constexpr auto LAYOUT = ckb::GroupConvLayout::NHWGC_GKYXC_NHWGK;
static constexpr auto DATA_TYPE = ckb::DataType::FP16;
};
TEST(ConvBuilderTest, TestSignature)
{
static_assert(ckb::ConvSignature<FwdConvSignature>);
SUCCEED();
}
struct FwdConvAlgorithm
{
//
};
TEST(ConvBuilderTest, TestAlgorithm)
{
static_assert(ckb::ConvAlgorithm<FwdConvAlgorithm>);
SUCCEED();
}
static constexpr char API_VERSION[] = "0.1.0";
using FwdConvBuilder = ckb::ConvBuilder<FwdConvSignature, FwdConvAlgorithm, API_VERSION>;
TEST(ConvBuilderTest, TestKernel)
{
EXPECT_EQ(
FwdConvBuilder::Instance::TypeString(),
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 256, 256, 32, Default, 32, 32, 4, 4, "
"8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v4>");
}
} // namespace

44
experimental/builder/test/test_conv_builder.cpp Normal file
View File

@@ -0,0 +1,44 @@
#include <gtest/gtest.h>
#include <ck_tile/builder/conv_builder.hpp>
namespace {
namespace ckb = ck_tile::builder;
struct FwdConvSignature
{
static constexpr int SPATIAL_DIM = 2;
static constexpr auto DIRECTION = ckb::ConvDirection::Forward;
static constexpr auto LAYOUT = ckb::GroupConvLayout::NHWGC_GKYXC_NHWGK;
static constexpr auto DATA_TYPE = ckb::DataType::FP16;
};
TEST(ConvBuilderTest, TestSignature)
{
static_assert(ckb::ConvSignature<FwdConvSignature>);
SUCCEED();
}
struct FwdConvAlgorithm
{
// TODO: Add algorithm info.
};
TEST(ConvBuilderTest, TestAlgorithm)
{
static_assert(ckb::ConvAlgorithm<FwdConvAlgorithm>);
SUCCEED();
}
static constexpr char API_VERSION[] = "0.1.0";
using FwdConvBuilder = ckb::ConvBuilder<FwdConvSignature, FwdConvAlgorithm, API_VERSION>;
TEST(ConvBuilderTest, TestInstance)
{
EXPECT_EQ(
FwdConvBuilder::Instance::TypeString(),
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<256, 256, 256, 32, Default, 32, 32, 4, 4, "
"8, 8, 8, 1, 1, BlkGemmPipelineScheduler: Intrawave, BlkGemmPipelineVersion: v4>");
}
} // namespace