diff --git a/experimental/builder/test/CMakeLists.txt b/experimental/builder/test/CMakeLists.txt index d1f439682d..16bba3373b 100644 --- a/experimental/builder/test/CMakeLists.txt +++ b/experimental/builder/test/CMakeLists.txt @@ -37,3 +37,4 @@ add_ck_factory_test(test_ck_factory_grouped_convolution_forward_convscale test_c 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_scaleadd_scaleadd_relu test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu.cpp) diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu.cpp new file mode 100644 index 0000000000..ba8726a643 --- /dev/null +++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_scaleadd_scaleadd_relu.cpp @@ -0,0 +1,105 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. + +#include + +#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 +using DeviceOp = DeviceGroupedConvFwdMultipleABD, // DsLayout + OutLayout, + T, // InDataType + T, // WeiDataType + ck::Tuple, // DsDataType + T, // OutDataType + PassThrough, + PassThrough, + ScaleAddScaleAddRelu, + T>; // ComputeType + +} // namespace + +template +struct CkFactoryTestConvFwd : public testing::Test +{ + static auto get_actual_instances() + { + return InstanceSet::from_factory(); + } + + static auto get_expected_instances() { return InstanceSet(Case::expected); } +}; + +struct F32 +{ + using DeviceOp = ::DeviceOp; + + constexpr static auto expected = { + // clang-format off + "" + // clang-format on + }; +}; + +struct F16 +{ + using DeviceOp = ::DeviceOp; + + constexpr static auto expected = { + // clang-format off + "" + // clang-format on + }; +}; + +struct BF16 +{ + using DeviceOp = ::DeviceOp; + + constexpr static auto expected = { + // clang-format off + "" + // clang-format on + }; +}; + +struct S8 +{ + using DeviceOp = ::DeviceOp; + + constexpr static auto expected = { + // clang-format off + "" + // clang-format on + }; +}; + +using TestTypes = ::testing::Types; + +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)); +}