ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-04-19 22:39:03 +00:00
Code Issues Packages Projects Releases Wiki Activity

[CK_BUILDER] Factory tests (#3071)

Browse Source 
This pull requests adds some initial "factory tests" - these check that the instances which are used in MIOpen are actually present in CK. The main reason for this is documentation and sanity checking. Its likely that these tests get outdated fast, so we'll have to maintain them, but fortunately this is quite straight forward and shouldn't take a lot of time once they are in place.
This commit is contained in:
Robin Voetter
2025-10-28 18:27:42 +01:00
committed by GitHub
parent 155d63f4fe
commit 6f58d6e457
18 changed files with 7295 additions and 18 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -2,7 +2,7 @@ include(gtest)
# Helper function to create a gtest executable with common properties
function(add_ck_builder_test test_name)
add_executable(${test_name} ${ARGN})
add_executable(${test_name} ${ARGN} testing_utils.cpp)
target_compile_features(${test_name} PRIVATE cxx_std_20)
target_include_directories(${test_name} PRIVATE
"${PROJECT_SOURCE_DIR}/experimental/builder/include"
@@ -20,9 +20,28 @@ endfunction()
add_ck_builder_test(test_conv_builder
test_conv_builder.cpp
test_instance_traits.cpp
test_instance_traits_util.cpp
testing_utils.cpp)
test_instance_traits_util.cpp)
# Testing the virtual GetInstanceString methods requires kernel compilation.
add_ck_builder_test(test_get_instance_string
test_get_instance_string.cpp)
add_ck_builder_test(test_inline_diff test_inline_diff.cpp)
function(add_ck_factory_test test_name)
add_ck_builder_test(${test_name} ${ARGN})
target_link_libraries(${test_name} PRIVATE composablekernels::device_conv_operations)
endfunction()
add_ck_factory_test(test_testing_utils test_testing_utils.cpp)
add_ck_factory_test(test_ck_factory_grouped_convolution_forward test_ck_factory_grouped_convolution_forward.cpp)
add_ck_factory_test(test_ck_factory_grouped_convolution_forward_clamp test_ck_factory_grouped_convolution_forward_clamp.cpp)
add_ck_factory_test(test_ck_factory_grouped_convolution_forward_convscale test_ck_factory_grouped_convolution_forward_convscale.cpp)
add_ck_factory_test(test_ck_factory_grouped_convolution_forward_bilinear test_ck_factory_grouped_convolution_forward_bilinear.cpp)
add_ck_factory_test(test_ck_factory_grouped_convolution_forward_scale test_ck_factory_grouped_convolution_forward_scale.cpp)
add_ck_factory_test(test_ck_factory_grouped_convolution_forward_scaleadd_ab test_ck_factory_grouped_convolution_forward_scaleadd_ab.cpp)
add_ck_factory_test(test_ck_factory_grouped_convolution_forward_bias_clamp test_ck_factory_grouped_convolution_forward_bias_clamp.cpp)
add_ck_factory_test(test_ck_factory_grouped_convolution_forward_bias_bnorm_clamp test_ck_factory_grouped_convolution_forward_bias_bnorm_clamp.cpp)
add_ck_factory_test(test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu.cpp)
add_ck_factory_test(test_ck_factory_grouped_convolution_forward_dynamic_op test_ck_factory_grouped_convolution_forward_dynamic_op.cpp)

2345
experimental/builder/test/test_ck_factory_grouped_convolution_forward.cpp Normal file
View File

File diff suppressed because it is too large Load Diff

1272
experimental/builder/test/test_ck_factory_grouped_convolution_forward_bias_bnorm_clamp.cpp Normal file
View File

File diff suppressed because it is too large Load Diff

1269
experimental/builder/test/test_ck_factory_grouped_convolution_forward_bias_clamp.cpp Normal file
View File

File diff suppressed because it is too large Load Diff

118
experimental/builder/test/test_ck_factory_grouped_convolution_forward_bilinear.cpp Normal file
View File

@@ -0,0 +1,118 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_bilinear.hpp>
#include "ck/utility/data_type.hpp"
#include "testing_utils.hpp"
using ck_tile::test::InstanceSet;
using ck_tile::test::InstancesMatch;
namespace {
constexpr static auto NumDimSpatial = 3;
using InLayout = ck::tensor_layout::convolution::NDHWGC;
using WeiLayout = ck::tensor_layout::convolution::GKZYXC;
using OutLayout = ck::tensor_layout::convolution::NDHWGK;
using DsLayout = ck::Tuple<ck::tensor_layout::convolution::NDHWGK>;
using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD;
using ck::tensor_operation::element_wise::Bilinear;
using ck::tensor_operation::element_wise::PassThrough;
template <typename type, typename computeType = type>
using DeviceOp = DeviceGroupedConvFwdMultipleABD<NumDimSpatial,
InLayout,
WeiLayout,
DsLayout,
OutLayout,
type, // InDataType
type, // WeiDataType
ck::Tuple<type>,
type, // OutDataType
PassThrough,
PassThrough,
Bilinear,
computeType,
computeType>;
} // namespace
template <typename Case>
struct CkFactoryTestBilinearFwd : public testing::Test
{
static auto get_actual_instances()
{
return InstanceSet::from_factory<typename Case::DeviceOp>();
}
static auto get_expected_instances() { return InstanceSet(Case::expected); }
};
struct Bilinear_F32
{
using DeviceOp = ::DeviceOp<float>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct Bilinear_F32_TF32
{
using DeviceOp = ::DeviceOp<float, ck::tf32_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct Bilinear_F16
{
using DeviceOp = ::DeviceOp<ck::half_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct Bilinear_BF16
{
using DeviceOp = ::DeviceOp<ck::bhalf_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct Bilinear_INT8
{
using DeviceOp = ::DeviceOp<int8_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
using TestTypes =
::testing::Types<Bilinear_F32, Bilinear_F32_TF32, Bilinear_F16, Bilinear_BF16, Bilinear_INT8>;
TYPED_TEST_SUITE(CkFactoryTestBilinearFwd, TestTypes);
TYPED_TEST(CkFactoryTestBilinearFwd, TestInstances)
{
auto actual = TestFixture::get_actual_instances();
auto expected = TestFixture::get_expected_instances();
EXPECT_THAT(actual, InstancesMatch(expected));
}

1249
experimental/builder/test/test_ck_factory_grouped_convolution_forward_clamp.cpp Normal file
View File

File diff suppressed because it is too large Load Diff

246
experimental/builder/test/test_ck_factory_grouped_convolution_forward_convscale.cpp Normal file
View File

@@ -0,0 +1,246 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_convscale.hpp>
#include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_convscale_relu.hpp>
#include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_convscale_add.hpp>
#include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_convinvscale.hpp>
#include <ck/library/tensor_operation_instance/device_operation_instance_factory.hpp>
#include "testing_utils.hpp"
using ck_tile::test::InstanceSet;
using ck_tile::test::InstancesMatch;
namespace {
constexpr static auto NumDimSpatial = 3;
using InLayout = ck::tensor_layout::convolution::NDHWGC;
using WeiLayout = ck::tensor_layout::convolution::GKZYXC;
using OutLayout = ck::tensor_layout::convolution::NDHWGK;
using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD;
using ck::tensor_operation::device::instance::CombConvScale;
using ck::tensor_operation::device::instance::CombConvScaleRelu;
using ck::tensor_operation::element_wise::ConvInvscale;
using ck::tensor_operation::element_wise::ConvScale;
using ck::tensor_operation::element_wise::ConvScaleAdd;
using ck::tensor_operation::element_wise::ConvScaleRelu;
using ck::tensor_operation::element_wise::PassThrough;
template <typename DsLayout,
typename DsDataType,
typename InDataType,
typename WeiDataType,
typename OutDataType,
typename Act,
typename AComputeType,
typename BComputeType>
using DeviceOp = DeviceGroupedConvFwdMultipleABD<NumDimSpatial,
InLayout,
WeiLayout,
DsLayout,
OutLayout,
InDataType, // InDataType
WeiDataType, // WeiDataType
DsDataType,
OutDataType, // OutDataType
PassThrough,
PassThrough,
Act,
AComputeType,
BComputeType>;
} // namespace
template <typename Case>
struct CkFactoryTestConvFwd : public testing::Test
{
static auto get_actual_instances()
{
return InstanceSet::from_factory<typename Case::DeviceOp>();
}
static auto get_expected_instances() { return InstanceSet(Case::expected); }
};
struct F8_ConvScale
{
using DeviceOp = ::DeviceOp<ck::Tuple<>,
ck::Tuple<>,
ck::f8_t,
ck::f8_t,
ck::f8_t,
ConvScale,
ck::f8_t,
ck::f8_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct F8_BF8_comb1_ConvScale
{
using DeviceOp = ::DeviceOp<ck::Tuple<>,
ck::Tuple<>,
ck::bf8_t,
ck::bf8_t,
ck::f8_t,
ConvScale,
ck::bf8_t,
ck::bf8_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct F8_BF8_comb2_ConvScale
{
using DeviceOp = ::DeviceOp<ck::Tuple<>,
ck::Tuple<>,
ck::f8_t,
ck::bf8_t,
ck::f8_t,
ConvScale,
ck::f8_t,
ck::bf8_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct F8_BF8_comb3_ConvScale
{
using DeviceOp = ::DeviceOp<ck::Tuple<>,
ck::Tuple<>,
ck::bf8_t,
ck::f8_t,
ck::f8_t,
ConvScale,
ck::bf8_t,
ck::f8_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct F8_float_CombConvScale
{
using DeviceOp = ::DeviceOp<ck::Tuple<>,
ck::Tuple<>,
ck::f8_t,
ck::f8_t,
float,
CombConvScale,
ck::f8_t,
ck::f8_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct F8_ConvScaleRelu
{
using DeviceOp = ::DeviceOp<ck::Tuple<>,
ck::Tuple<>,
ck::f8_t,
ck::f8_t,
ck::f8_t,
ConvScaleRelu,
ck::f8_t,
ck::f8_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct F8_CombConvScaleRelu
{
using DeviceOp = ::DeviceOp<ck::Tuple<>,
ck::Tuple<>,
ck::f8_t,
ck::f8_t,
float,
CombConvScaleRelu,
ck::f8_t,
ck::f8_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct F8_ConvScaleAdd
{
using DeviceOp = ::DeviceOp<ck::Tuple<OutLayout>,
ck::Tuple<float>,
ck::f8_t,
ck::f8_t,
ck::f8_t,
ConvScaleAdd,
ck::f8_t,
ck::f8_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct F8_ConvInvscale
{
using DeviceOp = ::DeviceOp<ck::Tuple<>,
ck::Tuple<>,
ck::f8_t,
ck::f8_t,
ck::f8_t,
ConvInvscale,
ck::f8_t,
ck::f8_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
using TestTypes = ::testing::Types<F8_ConvScale,
F8_BF8_comb1_ConvScale,
F8_BF8_comb2_ConvScale,
F8_BF8_comb3_ConvScale,
F8_float_CombConvScale,
F8_ConvScaleRelu,
F8_CombConvScaleRelu,
F8_ConvScaleAdd,
F8_ConvInvscale>;
TYPED_TEST_SUITE(CkFactoryTestConvFwd, TestTypes);
TYPED_TEST(CkFactoryTestConvFwd, TestInstances)
{
auto actual = TestFixture::get_actual_instances();
auto expected = TestFixture::get_expected_instances();
EXPECT_THAT(actual, InstancesMatch(expected));
}

187
experimental/builder/test/test_ck_factory_grouped_convolution_forward_dynamic_op.cpp Normal file
View File

@@ -0,0 +1,187 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_dynamic_op.hpp>
#include "ck/utility/data_type.hpp"
#include "testing_utils.hpp"
using ck_tile::test::InstanceSet;
using ck_tile::test::InstancesMatch;
namespace {
using InLayout = ck::tensor_layout::convolution::NDHWGC;
using WeiLayout = ck::tensor_layout::convolution::GKZYXC;
using OutLayout = ck::tensor_layout::convolution::NDHWGK;
using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD;
using ck::tensor_operation::element_wise::DynamicUnaryOp;
using ck::tensor_operation::element_wise::PassThrough;
template <ck::index_t NumDimSpatial, typename T>
struct DeviceOpHelper;
template <typename T>
struct DeviceOpHelper<2, T>
{
using InLayout = ck::tensor_layout::convolution::NHWGC;
using WeiLayout = ck::tensor_layout::convolution::GKYXC;
using OutLayout = ck::tensor_layout::convolution::NHWGK;
using Type = DeviceGroupedConvFwdMultipleABD<2,
InLayout,
WeiLayout,
ck::Tuple<>, // DsLayout
OutLayout,
T, // InDataType
T, // WeiDataType
ck::Tuple<>, // DsDataType
T, // OutDataType
PassThrough,
PassThrough,
DynamicUnaryOp>;
};
template <typename T>
struct DeviceOpHelper<3, T>
{
using InLayout = ck::tensor_layout::convolution::NDHWGC;
using WeiLayout = ck::tensor_layout::convolution::GKZYXC;
using OutLayout = ck::tensor_layout::convolution::NDHWGK;
using Type = DeviceGroupedConvFwdMultipleABD<3,
InLayout,
WeiLayout,
ck::Tuple<>, // DsLayout
OutLayout,
T, // InDataType
T, // WeiDataType
ck::Tuple<>, // DsDataType
T, // OutDataType
PassThrough,
PassThrough,
DynamicUnaryOp>;
};
template <ck::index_t NumDimSpatial, typename T>
using DeviceOp = DeviceOpHelper<NumDimSpatial, T>::Type;
} // namespace
template <typename Case>
struct CkFactoryTestBilinearFwd : public testing::Test
{
static auto get_actual_instances()
{
return InstanceSet::from_factory<typename Case::DeviceOp>();
}
static auto get_expected_instances() { return InstanceSet(Case::expected); }
};
struct DyOp_F32_2
{
using DeviceOp = ::DeviceOp<2, float>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct DyOp_F32_3
{
using DeviceOp = ::DeviceOp<3, float>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct DyOp_F16_2
{
using DeviceOp = ::DeviceOp<2, ck::half_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct DyOp_F16_3
{
using DeviceOp = ::DeviceOp<3, ck::half_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct DyOp_BF16_2
{
using DeviceOp = ::DeviceOp<2, ck::bhalf_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct DyOp_BF16_3
{
using DeviceOp = ::DeviceOp<3, ck::bhalf_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct DyOp_INT8_2
{
using DeviceOp = ::DeviceOp<2, int8_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct DyOp_INT8_3
{
using DeviceOp = ::DeviceOp<3, int8_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
using TestTypes = ::testing::Types<DyOp_F32_2,
DyOp_F32_3,
DyOp_F16_2,
DyOp_F16_3,
DyOp_BF16_2,
DyOp_BF16_3,
DyOp_INT8_2,
DyOp_INT8_3>;
TYPED_TEST_SUITE(CkFactoryTestBilinearFwd, TestTypes);
TYPED_TEST(CkFactoryTestBilinearFwd, TestInstances)
{
auto actual = TestFixture::get_actual_instances();
auto expected = TestFixture::get_expected_instances();
EXPECT_THAT(actual, InstancesMatch(expected));
}

115
experimental/builder/test/test_ck_factory_grouped_convolution_forward_scale.cpp Normal file
View File

@@ -0,0 +1,115 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_scale.hpp>
#include "testing_utils.hpp"
using ck_tile::test::InstanceSet;
using ck_tile::test::InstancesMatch;
namespace {
constexpr static auto NumDimSpatial = 3;
using InLayout = ck::tensor_layout::convolution::NDHWGC;
using WeiLayout = ck::tensor_layout::convolution::GKZYXC;
using OutLayout = ck::tensor_layout::convolution::NDHWGK;
using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD;
using ck::tensor_operation::element_wise::PassThrough;
using ck::tensor_operation::element_wise::Scale;
template <typename T, typename ComputeType = T>
using DeviceOp = DeviceGroupedConvFwdMultipleABD<NumDimSpatial,
InLayout,
WeiLayout,
ck::Tuple<>, // DsLayout
OutLayout,
T, // InDataType
T, // WeiDataType
ck::Tuple<>, // DsDataType
T, // OutDataType
PassThrough,
PassThrough,
Scale,
ComputeType>;
} // namespace
template <typename Case>
struct CkFactoryTestConvFwd : public testing::Test
{
static auto get_actual_instances()
{
return InstanceSet::from_factory<typename Case::DeviceOp>();
}
static auto get_expected_instances() { return InstanceSet(Case::expected); }
};
struct F32
{
using DeviceOp = ::DeviceOp<float>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct F32_TF32
{
using DeviceOp = ::DeviceOp<float, ck::tf32_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct F16
{
using DeviceOp = ::DeviceOp<ck::half_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct BF16
{
using DeviceOp = ::DeviceOp<ck::bhalf_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct S8
{
using DeviceOp = ::DeviceOp<int8_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
using TestTypes = ::testing::Types<F32, F32_TF32, F16, BF16, S8>;
TYPED_TEST_SUITE(CkFactoryTestConvFwd, TestTypes);
TYPED_TEST(CkFactoryTestConvFwd, TestInstances)
{
auto actual = TestFixture::get_actual_instances();
auto expected = TestFixture::get_expected_instances();
EXPECT_THAT(actual, InstancesMatch(expected));
}

104
experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_ab.cpp Normal file
View File

@@ -0,0 +1,104 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_scaleadd_ab.hpp>
#include "testing_utils.hpp"
using ck_tile::test::InstanceSet;
using ck_tile::test::InstancesMatch;
namespace {
constexpr static auto NumDimSpatial = 3;
using InLayout = ck::tensor_layout::convolution::NDHWGC;
using WeiLayout = ck::tensor_layout::convolution::GKZYXC;
using OutLayout = ck::tensor_layout::convolution::NDHWGK;
using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD;
using ck::tensor_operation::element_wise::PassThrough;
using ck::tensor_operation::element_wise::ScaleAdd;
template <typename T>
using DeviceOp = DeviceGroupedConvFwdMultipleABD<NumDimSpatial,
InLayout,
WeiLayout,
ck::Tuple<>, // DsLayout
OutLayout,
ck::Tuple<T, T>, // InDataType
ck::Tuple<T, T>, // WeiDataType
ck::Tuple<>, // DsDataType
T, // OutDataType
ScaleAdd,
ScaleAdd,
PassThrough,
T>; // ComputeType
} // namespace
template <typename Case>
struct CkFactoryTestConvFwd : public testing::Test
{
static auto get_actual_instances()
{
return InstanceSet::from_factory<typename Case::DeviceOp>();
}
static auto get_expected_instances() { return InstanceSet(Case::expected); }
};
struct F32
{
using DeviceOp = ::DeviceOp<float>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct F16
{
using DeviceOp = ::DeviceOp<ck::half_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct BF16
{
using DeviceOp = ::DeviceOp<ck::bhalf_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct S8
{
using DeviceOp = ::DeviceOp<int8_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
using TestTypes = ::testing::Types<F32, F16, BF16, S8>;
TYPED_TEST_SUITE(CkFactoryTestConvFwd, TestTypes);
TYPED_TEST(CkFactoryTestConvFwd, TestInstances)
{
auto actual = TestFixture::get_actual_instances();
auto expected = TestFixture::get_expected_instances();
EXPECT_THAT(actual, InstancesMatch(expected));
}

105
experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu.cpp Normal file
View File

@@ -0,0 +1,105 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_scaleadd_scaleadd_relu.hpp>
#include "testing_utils.hpp"
using ck_tile::test::InstanceSet;
using ck_tile::test::InstancesMatch;
namespace {
constexpr static auto NumDimSpatial = 3;
using InLayout = ck::tensor_layout::convolution::NDHWGC;
using WeiLayout = ck::tensor_layout::convolution::GKZYXC;
using OutLayout = ck::tensor_layout::convolution::NDHWGK;
using ck::tensor_layout::convolution::G_K;
using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD;
using ck::tensor_operation::element_wise::PassThrough;
using ck::tensor_operation::element_wise::ScaleAddScaleAddRelu;
template <typename T, typename U = T>
using DeviceOp = DeviceGroupedConvFwdMultipleABD<NumDimSpatial,
InLayout,
WeiLayout,
ck::Tuple<OutLayout, G_K>, // DsLayout
OutLayout,
T, // InDataType
T, // WeiDataType
ck::Tuple<U, U>, // DsDataType
T, // OutDataType
PassThrough,
PassThrough,
ScaleAddScaleAddRelu,
T>; // ComputeType
} // namespace
template <typename Case>
struct CkFactoryTestConvFwd : public testing::Test
{
static auto get_actual_instances()
{
return InstanceSet::from_factory<typename Case::DeviceOp>();
}
static auto get_expected_instances() { return InstanceSet(Case::expected); }
};
struct F32
{
using DeviceOp = ::DeviceOp<float>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct F16
{
using DeviceOp = ::DeviceOp<ck::half_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct BF16
{
using DeviceOp = ::DeviceOp<ck::bhalf_t>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
struct S8
{
using DeviceOp = ::DeviceOp<int8_t, float>;
constexpr static auto expected = {
// clang-format off
""
// clang-format on
};
};
using TestTypes = ::testing::Types<F32, F16, BF16, S8>;
TYPED_TEST_SUITE(CkFactoryTestConvFwd, TestTypes);
TYPED_TEST(CkFactoryTestConvFwd, TestInstances)
{
auto actual = TestFixture::get_actual_instances();
auto expected = TestFixture::get_expected_instances();
EXPECT_THAT(actual, InstancesMatch(expected));
}

98
experimental/builder/test/test_testing_utils.cpp Normal file
View File

@@ -0,0 +1,98 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved.
#include <ck/library/tensor_operation_instance/gpu/grouped_convolution_forward.hpp>
#include "testing_utils.hpp"
using ck_tile::test::InstanceMatcher;
using ck_tile::test::InstanceSet;
using ck_tile::test::StringEqWithDiff;
TEST(InstanceSet, FromFactory)
{
using DeviceOp = ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD<
2, // NDimSpatial
ck::tensor_operation::device::instance::NHWGC, // InLayout
ck::tensor_operation::device::instance::GKYXC, // WeiLayout
ck::tensor_operation::device::instance::Empty_Tuple, // DsLayout
ck::tensor_operation::device::instance::NHWGK, // OutLayout
ck::half_t, // ADataType
ck::half_t, // BDataType
ck::Tuple<>, // DsDataType
ck::half_t, // EDataType
ck::tensor_operation::element_wise::PassThrough, // AElementwiseOperation
ck::tensor_operation::element_wise::PassThrough, // BElementwiseOperation
ck::tensor_operation::element_wise::PassThrough, // CDEElementwiseOperation
ck::half_t, // AComputeType
ck::half_t>; // BComputeType
const auto instances = InstanceSet::from_factory<DeviceOp>();
EXPECT_THAT(instances.instances.size(), testing::Gt(0));
const auto* el =
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle_V3<2,NHWGC,GKYXC,EmptyTuple,NHWGK,fp16,fp16,"
"fp32,fp16,EmptyTuple,fp16,PassThrough,PassThrough,PassThrough,OddC,MNKPadding,64,16,16,64,"
"8,8,16,16,1,1,Seq(8,8,1),Seq(1,0,2),Seq(1,0,2),2,8,8,0,Seq(8,8,1),Seq(1,0,2),Seq(1,0,2),2,"
"8,8,0,1,1,Seq(1,16,1,4),4,Intrawave,v2,fp16,fp16>";
EXPECT_THAT(instances.instances, testing::Contains(el));
}
TEST(InstanceMatcher, Basic)
{
auto a = InstanceSet{
"python",
"cobra",
"boa",
};
auto b = InstanceSet{
"cobra",
"boa",
"python",
};
auto c = InstanceSet{
"adder",
"boa",
"cobra",
};
auto d = InstanceSet{
"boa",
"python",
};
EXPECT_THAT(a, InstancesMatch(b));
EXPECT_THAT(c, Not(InstancesMatch(b)));
EXPECT_THAT(a, Not(InstancesMatch(d)));
EXPECT_THAT(d, Not(InstancesMatch(b)));
}
TEST(InstanceMatcher, ExplainMatchResult)
{
auto actual = InstanceSet{
"python",
"cobra",
"boa",
};
auto expected = InstanceSet{
"adder",
"boa",
"cobra",
"rattlesnake",
};
testing::StringMatchResultListener listener;
EXPECT_TRUE(!ExplainMatchResult(InstancesMatch(expected), actual, &listener));
EXPECT_THAT(listener.str(),
StringEqWithDiff("\n"
" Missing: 2\n"
"- adder\n"
"- rattlesnake\n"
"Unexpected: 1\n"
"- python\n"));
}

101
experimental/builder/test/testing_utils.cpp
View File

@@ -1,21 +1,18 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
#include <string>
#include <sstream>
#include <vector>
#include <algorithm>
#include <unistd.h>
#include "testing_utils.hpp"
#include <gtest/gtest.h>
#include <gmock/gmock.h>
#include "testing_utils.hpp"
#include <unistd.h>
#include <string>
#include <sstream>
#include <ostream>
#include <vector>
#include <algorithm>
namespace ck_tile::test {
namespace {
} // namespace
// Wagner-Fischer Algorithm for Computing Edit Distance and Inline Diff
//
// OUTPUT FORMAT: [expected|actual] for differences, plain text for matches
@@ -216,4 +213,88 @@ void StringEqWithDiffMatcher::DescribeNegationTo(std::ostream* os) const
return ::testing::MakeMatcher(new StringEqWithDiffMatcher(expected));
}
std::ostream& operator<<(std::ostream& os, const InstanceSet& set)
{
// These sets can grow very large, and so its not very nice or useful to print them
// in the event of a mismatch. Just print a brief description here, and use
// InstancesMatcher to print a more useful message.
return (os << "(set of " << set.instances.size() << " instances)");
}
InstanceMatcher::InstanceMatcher(const InstanceSet& expected) : expected_(expected) {}
::testing::Matcher<InstanceSet> InstancesMatch(const InstanceSet& expected)
{
return ::testing::MakeMatcher(new InstanceMatcher(expected));
}
bool InstanceMatcher::MatchAndExplain(InstanceSet actual,
::testing::MatchResultListener* listener) const
{
if(actual.instances == expected_.instances)
{
return true;
}
if(listener->IsInterested())
{
std::vector<std::string> instances;
std::set_difference(expected_.instances.begin(),
expected_.instances.end(),
actual.instances.begin(),
actual.instances.end(),
std::back_inserter(instances));
*listener << "\n";
if(instances.size() > 0)
{
*listener << " Missing: " << instances.size() << "\n";
for(const auto& instance : instances)
{
if(instance == "")
{
*listener << "- (empty string)\n";
}
else
{
*listener << "- " << instance << "\n";
}
}
}
instances.clear();
std::set_difference(actual.instances.begin(),
actual.instances.end(),
expected_.instances.begin(),
expected_.instances.end(),
std::back_inserter(instances));
if(instances.size() > 0)
{
*listener << "Unexpected: " << instances.size() << "\n";
for(const auto& instance : instances)
{
if(instance == "")
{
*listener << "- (empty string)\n";
}
else
{
*listener << "- " << instance << "\n";
}
}
}
}
return false;
}
void InstanceMatcher::DescribeTo(std::ostream* os) const { *os << expected_; }
void InstanceMatcher::DescribeNegationTo(std::ostream* os) const
{
*os << "is not equal to " << expected_;
}
} // namespace ck_tile::test

69
experimental/builder/test/testing_utils.hpp
View File

@@ -1,10 +1,15 @@
// SPDX-License-Identifier: MIT
// Copyright (c) 2018-2025, Advanced Micro Devices, Inc. All rights reserved.
#include <ck/library/tensor_operation_instance/device_operation_instance_factory.hpp>
#include <gtest/gtest.h>
#include <gmock/gmock.h>
#include <string>
#include <sstream>
#include <iosfwd>
#include <set>
#include <vector>
#include <array>
namespace ck_tile::test {
@@ -40,4 +45,68 @@ class StringEqWithDiffMatcher : public ::testing::MatcherInterface<std::string>
// Factory function for the StringEqWithDiff matcher
::testing::Matcher<std::string> StringEqWithDiff(const std::string& expected);
using ck::tensor_operation::device::instance::DeviceOperationInstanceFactory;
// This utility concept checks whether a type is a valid "Device Operation" -
// that is, there is a valid specialization of `DeviceOperationInstanceFactory`
// for it available.
template <typename DeviceOp>
concept HasCkFactory = requires {
{
DeviceOperationInstanceFactory<DeviceOp>::GetInstances()
} -> std::convertible_to<std::vector<std::unique_ptr<DeviceOp>>>;
};
// This structure represents a (unique) set of instances, either a statically
// defined one (for testing) or one obtained from DeviceOperationInstanceFactory.
// The idea is that we use this structure as a utility to compare a set of
// instances. Instances are stored in a set so that they can be lexicographically
// compared, this helps generating readable error messages which just contain
// the differenses between sets.
struct InstanceSet
{
explicit InstanceSet() {}
explicit InstanceSet(std::initializer_list<const char*> items)
: instances(items.begin(), items.end())
{
}
template <HasCkFactory DeviceOp>
static InstanceSet from_factory()
{
auto set = InstanceSet();
const auto ops = DeviceOperationInstanceFactory<DeviceOp>::GetInstances();
for(const auto& op : ops)
{
set.instances.insert(op->GetInstanceString());
}
return set;
}
std::set<std::string> instances;
};
std::ostream& operator<<(std::ostream& os, const InstanceSet& set);
// This is a custom Google Test matcher which can be used to compare two sets
// of instance names, with utility functions that print a helpful error
// message about the difference between the checked sets. Use `InstancesMatch`
// to obtain an instance of this type.
struct InstanceMatcher : public ::testing::MatcherInterface<InstanceSet>
{
explicit InstanceMatcher(const InstanceSet& expected);
bool MatchAndExplain(InstanceSet actual,
::testing::MatchResultListener* listener) const override;
void DescribeTo(std::ostream* os) const override;
void DescribeNegationTo(std::ostream* os) const override;
InstanceSet expected_;
};
::testing::Matcher<InstanceSet> InstancesMatch(const InstanceSet& expected);
} // namespace ck_tile::test

2
library/include/ck/library/tensor_operation_instance/gpu/grouped_convolution_forward.hpp
View File

@@ -300,7 +300,7 @@ struct DeviceOperationInstanceFactory<ck::tensor_operation::device::DeviceGroupe
#endif
}
// layout NGCHW/GKYXC/NGKHW
// layout NGCHW/GKCYX/NGKHW
if constexpr(NumDimSpatial == 2 && is_same_v<InLayout, NGCHW> &&
is_same_v<WeiLayout, GKCYX> && is_same_v<OutLayout, NGKHW>)
{

4
library/include/ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_dynamic_op.hpp
View File

@@ -7,7 +7,7 @@
#include <memory>
#include "ck/ck.hpp"
#include "ck/tensor_operation/gpu/device/device_grouped_conv_fwd_dynamic.hpp"
#include "ck/tensor_operation/gpu/device/device_grouped_conv_fwd_multiple_abd.hpp"
#include "ck/tensor_operation/gpu/device/tensor_layout.hpp"
#include "ck/tensor_operation/gpu/element/element_wise_operation.hpp"
@@ -161,7 +161,7 @@ template <ck::index_t NumDimSpatial,
typename DDataTypes,
typename OutDataType,
typename AComputeType,
typename BComputeType = AComputeType>
typename BComputeType>
struct DeviceOperationInstanceFactory<ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD<
NumDimSpatial,
InLayout,

2
library/include/ck/library/tensor_operation_instance/gpu/grouped_convolution_forward_scale.hpp
View File

@@ -69,7 +69,7 @@ void add_device_grouped_conv3d_fwd_xdl_scale_ndhwgc_gkzyxc_ndhwgk_f32_instances(
PassThrough,
Scale>>>& instances);
void add_device_grouped_conv3d_fwd_xdl_scale_ndhwgc_gkzyxc_ndhwgk_f32_instances(
void add_device_grouped_conv3d_fwd_xdl_scale_ndhwgc_gkzyxc_ndhwgk_f32_tf32_instances(
std::vector<std::unique_ptr<DeviceGroupedConvFwdMultipleABD<3,
NDHWGC,
GKZYXC,

2
script/remove_exec_bit.sh
View File

@@ -1,4 +1,4 @@
#!/bin/bash
#!/usr/bin/env bash
# Copyright © Advanced Micro Devices, Inc., or its affiliates.
# SPDX-License-Identifier: MIT