From bd40b582668752dd1d1c5bdbbfbca7c3816b7091 Mon Sep 17 00:00:00 2001 From: Kevin Abraham Date: Fri, 24 Oct 2025 19:54:05 +0000 Subject: [PATCH] implemented tests for instances from grouped_convolution_forward_convscale.hpp:100 --- ..._grouped_convolution_forward_convscale.cpp | 126 ++++++++++++++++-- .../builder/test/test_testing_utils.cpp | 2 +- 2 files changed, 119 insertions(+), 9 deletions(-) diff --git a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_convscale.cpp b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_convscale.cpp index 10835cde0b..997fadada4 100644 --- a/experimental/builder/test/test_ck_factory_grouped_convolution_forward_convscale.cpp +++ b/experimental/builder/test/test_ck_factory_grouped_convolution_forward_convscale.cpp @@ -1,6 +1,7 @@ // SPDX-License-Identifier: MIT // Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. +#include #include #include #include @@ -20,23 +21,33 @@ using OutLayout = ck::tensor_layout::convolution::NDHWGK; using ck::tensor_operation::device::DeviceGroupedConvFwdMultipleABD; 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 +template using DeviceOp = DeviceGroupedConvFwdMultipleABD; + Act, + AComputeType, + BComputeType>; } // namespace @@ -51,9 +62,88 @@ struct CkFactoryTestConvFwd : public testing::Test static auto get_expected_instances() { return InstanceSet(Case::expected); } }; +struct F8_ConvScale +{ + using DeviceOp = ::DeviceOp, + 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_comb1_ConvScale +{ + using DeviceOp = ::DeviceOp, + 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_comb2_ConvScale +{ + using DeviceOp = ::DeviceOp, + 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_comb3_ConvScale +{ + using DeviceOp = ::DeviceOp, + 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_ConvScaleRelu { - using DeviceOp = ::DeviceOp, ck::Tuple<>, ck::f8_t, ConvScaleRelu>; + using DeviceOp = ::DeviceOp, + 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 @@ -64,7 +154,14 @@ struct F8_ConvScaleRelu struct F8_CombConvScaleRelu { - using DeviceOp = ::DeviceOp, ck::Tuple<>, float, CombConvScaleRelu>; + using DeviceOp = ::DeviceOp, + ck::Tuple<>, + ck::f8_t, + ck::f8_t, + float, + CombConvScaleRelu, + ck::f8_t, + ck::f8_t>; constexpr static auto expected = { // clang-format off @@ -75,7 +172,14 @@ struct F8_CombConvScaleRelu struct F8_ConvScaleAdd { - using DeviceOp = ::DeviceOp, ck::Tuple, ck::f8_t, ConvScaleAdd>; + using DeviceOp = ::DeviceOp, + ck::Tuple, + ck::f8_t, + ck::f8_t, + ck::f8_t, + ConvScaleAdd, + ck::f8_t, + ck::f8_t>; constexpr static auto expected = { // clang-format off @@ -96,7 +200,13 @@ struct F8_ConvInvscale }; using TestTypes = - ::testing::Types; + ::testing::Types; TYPED_TEST_SUITE(CkFactoryTestConvFwd, TestTypes); diff --git a/experimental/builder/test/test_testing_utils.cpp b/experimental/builder/test/test_testing_utils.cpp index 79e9cbadda..d80eaed431 100644 --- a/experimental/builder/test/test_testing_utils.cpp +++ b/experimental/builder/test/test_testing_utils.cpp @@ -3,7 +3,7 @@ #include -#include "testing_utils.hpp" +#include "testing_utils.hpp" using ck_tile::test::InstanceMatcher; using ck_tile::test::InstanceSet;