// Copyright (c) Advanced Micro Devices, Inc., or its affiliates. // SPDX-License-Identifier: MIT #include #include #include "ck_tile/builder/factory/helpers/ck/conv_tensor_layout.hpp" #include "impl/conv_signature_types.hpp" namespace { namespace ckb = ck_tile::builder; using ck_tile::builder::DataType; using ck_tile::builder::ElementwiseOperation; using ck_tile::builder::TensorLayout; using ck_tile::builder::factory::internal::AuxiliaryTensorLayouts; using ck_tile::builder::factory::internal::ConvTensorLayouts; using ck_tile::builder::factory::internal::LayoutToCK; using ck_tile::builder::test::ConvolutionTensor; using ck_tile::builder::test::ConvSignature; using ck_tile::builder::test::TensorConfig; using ck_tile::builder::test::TensorOperation; namespace enums { using enum ck_tile::builder::ConvDirection; using enum ck_tile::builder::TensorLayout; using enum ck_tile::builder::DataType; } // namespace enums TEST(ConvTensorLayout, AssignsLayoutsFor1D_NWGC_GKXC_NWGK) { using namespace enums; static constexpr auto sig = ConvSignature<>{.spatial_dim = 1, .direction = FORWARD, .data_type = FP16, .accumulation_data_type = FP32, .input = {.config = {.layout = NWGC}}, .weight = {.config = {.layout = GKXC}}, .output = {.config = {.layout = NWGK}}}; using TensorLayouts = ConvTensorLayouts; EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } TEST(ConvTensorLayout, AssignsLayoutsFor1D_NGCW_GKXC_NGKW) { using namespace enums; static constexpr auto sig = ConvSignature<>{.spatial_dim = 1, .direction = FORWARD, .data_type = FP16, .accumulation_data_type = FP32, .input = {.config = {.layout = NGCW}}, .weight = {.config = {.layout = GKXC}}, .output = {.config = {.layout = NGKW}}}; using TensorLayouts = ConvTensorLayouts; EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } TEST(ConvTensorLayout, AssignsLayoutsFor1D_GNWC_GKXC_GNWK) { using namespace enums; static constexpr auto sig = ConvSignature<>{.spatial_dim = 1, .direction = FORWARD, .data_type = FP16, .accumulation_data_type = FP32, .input = {.config = {.layout = GNWC}}, .weight = {.config = {.layout = GKXC}}, .output = {.config = {.layout = GNWK}}}; using TensorLayouts = ConvTensorLayouts; EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } TEST(ConvTensorLayout, AssignsLayoutsFor1D_NGCW_GKCX_NGKW) { using namespace enums; static constexpr auto sig = ConvSignature<>{.spatial_dim = 1, .direction = FORWARD, .data_type = FP16, .accumulation_data_type = FP32, .input = {.config = {.layout = NGCW}}, .weight = {.config = {.layout = GKCX}}, .output = {.config = {.layout = NGKW}}}; using TensorLayouts = ConvTensorLayouts; EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } TEST(ConvTensorLayout, AssignsLayoutsFor2D_NGCHW_GKYXC_NGKHW) { using namespace enums; static constexpr auto sig = ConvSignature<>{.spatial_dim = 2, .direction = FORWARD, .data_type = FP16, .accumulation_data_type = FP32, .input = {.config = {.layout = NGCHW}}, .weight = {.config = {.layout = GKYXC}}, .output = {.config = {.layout = NGKHW}}}; using TensorLayouts = ConvTensorLayouts; EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } TEST(ConvTensorLayout, AssignsLayoutsFor2D_NHWGC_GKYXC_NHWGK) { using namespace enums; static constexpr auto sig = ConvSignature<>{.spatial_dim = 2, .direction = FORWARD, .data_type = DataType::FP16, .accumulation_data_type = DataType::FP32, .input = {.config = {.layout = NHWGC}}, .weight = {.config = {.layout = GKYXC}}, .output = {.config = {.layout = NHWGK}}}; using TensorLayouts = ConvTensorLayouts; EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } TEST(ConvTensorLayout, AssignsLayoutsFor2D_GNHWC_GKYXC_GNHWK) { using namespace enums; static constexpr auto sig = ConvSignature<>{.spatial_dim = 2, .direction = FORWARD, .data_type = DataType::FP16, .accumulation_data_type = DataType::FP32, .input = {.config = {.layout = GNHWC}}, .weight = {.config = {.layout = GKYXC}}, .output = {.config = {.layout = GNHWK}}}; using TensorLayouts = ConvTensorLayouts; EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } TEST(ConvTensorLayout, AssignsLayoutsFor2D_NGCHW_GKCYX_NGKHW) { using namespace enums; static constexpr auto sig = ConvSignature<>{.spatial_dim = 2, .direction = FORWARD, .data_type = DataType::FP16, .accumulation_data_type = DataType::FP32, .input = {.config = {.layout = NGCHW}}, .weight = {.config = {.layout = GKCYX}}, .output = {.config = {.layout = NGKHW}}}; using TensorLayouts = ConvTensorLayouts; EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } TEST(ConvTensorLayout, AssignsLayoutsFor3D_NGCDHW_GKCZYX_NGKDHW) { using namespace enums; static constexpr auto sig = ConvSignature<>{.spatial_dim = 3, .direction = FORWARD, .data_type = DataType::FP16, .accumulation_data_type = DataType::FP32, .input = {.config = {.layout = NGCDHW}}, .weight = {.config = {.layout = GKCZYX}}, .output = {.config = {.layout = NGKDHW}}}; using TensorLayouts = ConvTensorLayouts; EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } TEST(ConvTensorLayout, AssignsLayoutsFor3D_NDHWGC_GKZYXC_NDHWGK) { using namespace enums; static constexpr auto sig = ConvSignature<>{.spatial_dim = 3, .direction = FORWARD, .data_type = DataType::FP16, .accumulation_data_type = DataType::FP32, .input = {.config = {.layout = NDHWGC}}, .weight = {.config = {.layout = GKZYXC}}, .output = {.config = {.layout = NDHWGK}}}; using TensorLayouts = ConvTensorLayouts; EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } TEST(ConvTensorLayout, AssignsLayoutsFor3D_GNDHWC_GKZYXC_GNDHWK) { using namespace enums; static constexpr auto sig = ConvSignature<>{.spatial_dim = 3, .direction = FORWARD, .data_type = DataType::FP16, .accumulation_data_type = DataType::FP32, .input = {.config = {.layout = GNDHWC}}, .weight = {.config = {.layout = GKZYXC}}, .output = {.config = {.layout = GNDHWK}}}; using TensorLayouts = ConvTensorLayouts; EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v>)); } TEST(AuxiliaryTensorLayout, AssignsLayoutForG_K_strided) { using CKLayout = LayoutToCK::type; EXPECT_TRUE((std::is_same_v)); } TEST(AuxiliaryTensorLayout, AssignsLayoutForGC) { using CKLayout = LayoutToCK::type; EXPECT_TRUE((std::is_same_v)); } TEST(AuxiliaryTensorLayout, AssignsLayoutForG_C_strided) { using CKLayout = LayoutToCK::type; EXPECT_TRUE((std::is_same_v)); } TEST(AuxiliaryTensorLayout, EmptyAuxiliaryTensorLayoutIsEmptyTuple) { using ::ck_tile::builder::factory::internal::EmptyAuxiliaryTensorLayout; using EmptyLayout = EmptyAuxiliaryTensorLayout::type; EXPECT_TRUE((std::is_same_v>)); } struct MockAuxiliaryTensorConfig { TensorLayout layout; }; TEST(AuxiliaryTensorLayoutIntegration, SingleBiasTensorWithG_K_Layout) { using namespace enums; static constexpr std::array aux_configs = { MockAuxiliaryTensorConfig{.layout = G_K_strided}}; using AuxLayouts = AuxiliaryTensorLayouts; EXPECT_EQ(AuxLayouts::Size, 1); using ExpectedType = ck::Tuple; EXPECT_TRUE((std::is_same_v)); } TEST(AuxiliaryTensorLayoutIntegration, SingleBiasTensorWithGC_Layout) { using namespace enums; static constexpr std::array aux_configs = { MockAuxiliaryTensorConfig{.layout = TensorLayout::GC}}; using AuxLayouts = AuxiliaryTensorLayouts; EXPECT_EQ(AuxLayouts::Size, 1); using ExpectedType = ck::Tuple; EXPECT_TRUE((std::is_same_v)); } TEST(AuxiliaryTensorLayoutIntegration, SingleBiasTensorWithG_C_Layout) { using namespace enums; static constexpr std::array aux_configs = { MockAuxiliaryTensorConfig{.layout = G_C_strided}}; using AuxLayouts = AuxiliaryTensorLayouts; EXPECT_EQ(AuxLayouts::Size, 1); using ExpectedType = ck::Tuple; EXPECT_TRUE((std::is_same_v)); } TEST(AuxiliaryTensorLayoutIntegration, TwoAuxiliaryTensors) { using namespace enums; static constexpr std::array aux_configs = { MockAuxiliaryTensorConfig{.layout = TensorLayout::G_K_strided}, MockAuxiliaryTensorConfig{.layout = GC}}; using AuxLayouts = AuxiliaryTensorLayouts; EXPECT_EQ(AuxLayouts::Size, 2); using ExpectedType = ck::Tuple; EXPECT_TRUE((std::is_same_v)); } TEST(AuxiliaryTensorLayoutIntegration, ThreeAuxiliaryTensors) { using namespace enums; static constexpr std::array aux_configs = { MockAuxiliaryTensorConfig{.layout = G_K_strided}, MockAuxiliaryTensorConfig{.layout = GC}, MockAuxiliaryTensorConfig{.layout = G_C_strided}}; using AuxLayouts = AuxiliaryTensorLayouts; EXPECT_EQ(AuxLayouts::Size, 3); using ExpectedType = ck::Tuple; EXPECT_TRUE((std::is_same_v)); } TEST(AuxiliaryTensorLayoutIntegration, WorksWith1DConvolution) { using namespace enums; static constexpr std::array aux_configs = { MockAuxiliaryTensorConfig{.layout = G_K_strided}}; using AuxLayouts = AuxiliaryTensorLayouts; EXPECT_EQ(AuxLayouts::Size, 1); using ExpectedType = ck::Tuple; EXPECT_TRUE((std::is_same_v)); } TEST(AuxiliaryTensorLayoutIntegration, WorksWith3DConvolution) { using namespace enums; static constexpr std::array aux_configs = { MockAuxiliaryTensorConfig{.layout = GC}}; using AuxLayouts = AuxiliaryTensorLayouts; EXPECT_EQ(AuxLayouts::Size, 1); using ExpectedType = ck::Tuple; EXPECT_TRUE((std::is_same_v)); } TEST(ConvTensorLayoutsWithAuxiliary, Conv2DWithSingleBiasG_K) { using namespace enums; using OutputOp = TensorOperation; static constexpr auto sig = ConvSignature, ConvolutionTensor<>, ConvolutionTensor>{ .spatial_dim = 2, .direction = FORWARD, .data_type = DataType::FP16, .accumulation_data_type = DataType::FP32, .input = {.config = {.layout = NGCHW}}, .weight = {.config = {.layout = GKYXC}}, .output = {.config = {.layout = NGKHW}, .operation = OutputOp{.elementwise_operation = ElementwiseOperation::SCALE}}}; using TensorLayouts = ConvTensorLayouts; EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); using ExpectedDsLayout = ck::Tuple; EXPECT_TRUE((std::is_same_v)); } TEST(ConvTensorLayoutsWithAuxiliary, Conv2DWithSingleBiasGC) { using namespace enums; using OutputOp = TensorOperation; static constexpr auto sig = ConvSignature, ConvolutionTensor<>, ConvolutionTensor>{ .spatial_dim = 2, .direction = FORWARD, .data_type = DataType::BF16, .accumulation_data_type = DataType::FP32, .input = {.config = {.layout = NHWGC}}, .weight = {.config = {.layout = GKYXC}}, .output = {.config = {.layout = NHWGK}, .operation = OutputOp{.elementwise_operation = ElementwiseOperation::SCALE}}}; using TensorLayouts = ConvTensorLayouts; EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); using ExpectedDsLayout = ck::Tuple; EXPECT_TRUE((std::is_same_v)); } TEST(ConvTensorLayoutsWithAuxiliary, Conv2DWithTwoAuxiliaryTensors) { using namespace enums; using OutputOp = TensorOperation; static constexpr auto sig = ConvSignature, ConvolutionTensor<>, ConvolutionTensor>{ .spatial_dim = 2, .direction = FORWARD, .data_type = DataType::FP16, .accumulation_data_type = DataType::FP32, .input = {.config = {.layout = GNHWC}}, .weight = {.config = {.layout = GKYXC}}, .output = {.config = {.layout = GNHWK}, .operation = OutputOp{.elementwise_operation = ElementwiseOperation::SCALEADD_SCALEADD_RELU}}}; using TensorLayouts = ConvTensorLayouts; EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); using ExpectedDsLayout = ck::Tuple; EXPECT_TRUE((std::is_same_v)); } TEST(ConvTensorLayoutsWithAuxiliary, Conv1DWithBias) { using namespace enums; using OutputOp = TensorOperation; static constexpr auto sig = ConvSignature, ConvolutionTensor<>, ConvolutionTensor>{ .spatial_dim = 1, .direction = FORWARD, .data_type = DataType::FP32, .accumulation_data_type = DataType::FP32, .input = {.config = {.layout = NWGC}}, .weight = {.config = {.layout = GKXC}}, .output = {.config = {.layout = NWGK}, .operation = OutputOp{.elementwise_operation = ElementwiseOperation::SCALE}}}; using TensorLayouts = ConvTensorLayouts; EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); using ExpectedDsLayout = ck::Tuple; EXPECT_TRUE((std::is_same_v)); } TEST(ConvTensorLayoutsWithAuxiliary, Conv3DWithBias) { using namespace enums; using OutputOp = TensorOperation; static constexpr auto sig = ConvSignature, ConvolutionTensor<>, ConvolutionTensor>{ .spatial_dim = 3, .direction = FORWARD, .data_type = DataType::FP16, .accumulation_data_type = DataType::FP32, .input = {.config = {.layout = NDHWGC}}, .weight = {.config = {.layout = GKZYXC}}, .output = {.config = {.layout = NDHWGK}, .operation = OutputOp{.elementwise_operation = ElementwiseOperation::BIAS_BNORM_CLAMP}}}; using TensorLayouts = ConvTensorLayouts; EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); EXPECT_TRUE((std::is_same_v)); using ExpectedDsLayout = ck::Tuple; EXPECT_TRUE((std::is_same_v)); } } // namespace