[Test][gtest] Sl/test solver (#1877)

test/gtest/solver.hpp

@@ -0,0 +1,165 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2022 Advanced Micro Devices, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+#pragma once
+
+#include <gtest/gtest.h>
+#include "cpu_conv.hpp"
+#include "get_handle.hpp"
+#include "tensor_util.hpp"
+#include <fusionHost.hpp>
+#include <miopen/conv/data_invoke_params.hpp>
+
+template <typename T>
+miopenDataType_t GetDataType();
+
+template <>
+miopenDataType_t GetDataType<float>()
+{
+    return miopenFloat;
+}
+
+template <>
+miopenDataType_t GetDataType<half_float::half>()
+{
+    return miopenHalf;
+}
+
+struct ConvTestCase
+{
+    size_t N;
+    size_t C;
+    size_t H;
+    size_t W;
+    size_t k;
+    size_t y;
+    size_t x;
+    size_t pad_x;
+    size_t pad_y;
+    size_t stride_x;
+    size_t stride_y;
+    size_t dialtion_x;
+    size_t dilation_y;
+    miopenConvolutionMode_t conv_mode;
+    friend std::ostream& operator<<(std::ostream& os, const ConvTestCase& tc)
+    {
+        return os << "N: " << tc.N << " C:" << tc.C << " H:" << tc.H << " W:" << tc.W
+                  << " k: " << tc.k << " y:" << tc.y << " x:" << tc.x << " pad_y:" << tc.pad_y
+                  << " pad_x:" << tc.pad_x << " stride_y:" << tc.stride_y
+                  << " dilation_y:" << tc.dilation_y << " conv_mode:" << tc.conv_mode;
+    }
+
+    miopen::ConvolutionDescriptor GetConv()
+    {
+        return miopen::ConvolutionDescriptor{
+            {static_cast<int>(pad_y), static_cast<int>(pad_x)},
+            {static_cast<int>(stride_y), static_cast<int>(stride_x)},
+            {static_cast<int>(dilation_y), static_cast<int>(dilation_y)}};
+    }
+};
+
+std::vector<ConvTestCase> ConvTestConfigs()
+{ // n  c   h   w   k   y  x pad_x pad_y stri_x stri_y dia_x dia_y
+    return {{16, 128, 16, 16, 128, 3, 3, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {64, 128, 28, 28, 128, 3, 3, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {64, 256, 14, 14, 256, 3, 3, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {64, 512, 7, 7, 512, 3, 3, 1, 1, 1, 1, 1, 1, miopenConvolution},
+            {64, 1024, 14, 14, 1024, 3, 3, 1, 1, 1, 1, 1, 1, miopenConvolution}};
+}
+
+template <typename T = float>
+struct ConvFwdSolverTest
+    : public ::testing::TestWithParam<std::tuple<miopenConvFwdAlgorithm_t, ConvTestCase>>
+{
+protected:
+    void SetUp() override
+    {
+        test_skipped                = false;
+        std::tie(algo, conv_config) = GetParam();
+        input   = tensor<T>{conv_config.N, conv_config.C, conv_config.H, conv_config.W};
+        weights = tensor<T>{1, conv_config.k, conv_config.x, conv_config.y};
+        input.generate(tensor_elem_gen_integer{17});
+        weights.generate(tensor_elem_gen_integer{17});
+
+        conv_desc = conv_config.GetConv();
+
+        miopen::TensorDescriptor output_desc =
+            conv_desc.GetForwardOutputTensor(input.desc, weights.desc, GetDataType<T>());
+
+        output = tensor<T>{output_desc.GetLengths()};
+
+        std::fill(output.begin(), output.end(), std::numeric_limits<double>::quiet_NaN());
+
+        auto&& handle = get_handle();
+        in_dev        = handle.Write(input.data);
+        wei_dev       = handle.Write(weights.data);
+        out_dev       = handle.Write(output.data);
+    }
+    void TearDown() override
+    {
+        if(test_skipped)
+            return;
+
+        auto&& handle = get_handle();
+
+        miopen::TensorDescriptor output_desc =
+            conv_desc.GetForwardOutputTensor(input.desc, weights.desc, GetDataType<T>());
+        ref_out = tensor<T>{output_desc.GetLengths()};
+        cpu_convolution_forward(conv_desc.GetSpatialDimension(),
+                                input,
+                                weights,
+                                ref_out,
+                                conv_desc.GetConvPads(),
+                                conv_desc.GetConvStrides(),
+                                conv_desc.GetConvDilations(),
+                                conv_desc.GetGroupCount());
+
+        output.data = handle.Read<T>(out_dev, output.data.size());
+        EXPECT_FALSE(miopen::range_zero(ref_out)) << "Cpu data is all zeros";
+        EXPECT_FALSE(miopen::range_zero(output)) << "Gpu data is all zeros";
+        EXPECT_TRUE(miopen::range_distance(ref_out) == miopen::range_distance(output));
+
+        const double tolerance = 80;
+        double threshold       = std::numeric_limits<T>::epsilon() * tolerance;
+        auto error             = miopen::rms_range(ref_out, output);
+
+        EXPECT_FALSE(miopen::find_idx(ref_out, miopen::not_finite) >= 0)
+            << "Non finite number found in the CPU data";
+
+        EXPECT_TRUE(error < threshold)
+            << "Error beyond tolerance Error:" << error << ",  Threshold: " << threshold;
+    }
+    ConvTestCase conv_config;
+    miopen::ConvolutionDescriptor conv_desc;
+    tensor<T> input;
+    tensor<T> weights;
+    tensor<T> output;
+    tensor<T> ref_out;
+    miopen::Allocator::ManageDataPtr in_dev;
+    miopen::Allocator::ManageDataPtr wei_dev;
+    miopen::Allocator::ManageDataPtr out_dev;
+    miopenConvFwdAlgorithm_t algo = miopenConvolutionFwdAlgoDirect;
+    bool test_skipped             = false;
+};

test/gtest/solver_convasm3x3u.cpp

@@ -0,0 +1,90 @@
+/*******************************************************************************
+ *
+ * MIT License
+ *
+ * Copyright (c) 2022 Advanced Micro Devices, Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a copy
+ * of this software and associated documentation files (the "Software"), to deal
+ * in the Software without restriction, including without limitation the rights
+ * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ * copies of the Software, and to permit persons to whom the Software is
+ * furnished to do so, subject to the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included in all
+ * copies or substantial portions of the Software.
+ *
+ * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+ * SOFTWARE.
+ *
+ *******************************************************************************/
+#include "solver.hpp"
+
+struct ConvFwdSolverTestFloat : ConvFwdSolverTest<float>
+{
+};
+
+template <typename Solver>
+void SolverFwd(const miopen::TensorDescriptor& inputDesc,
+               ConstData_t input,
+               const miopen::TensorDescriptor& wDesc,
+               ConstData_t weight,
+               const miopen::TensorDescriptor& outputDesc,
+               Data_t output,
+               const miopen::ConvolutionDescriptor& convDesc,
+               const ConvTestCase& conv_config,
+               bool& test_skipped)
+{
+    auto&& handle = get_handle();
+
+    Solver solv{};
+
+    const auto tensors =
+        miopen::ConvFwdTensors{inputDesc, input, wDesc, weight, outputDesc, output};
+
+    auto ctx = miopen::ConvolutionContext{
+        inputDesc, wDesc, outputDesc, convDesc, miopen::conv::Direction::Forward};
+
+    ctx.SetStream(&handle);
+    ctx.DetectRocm();
+
+    if(!solv.IsApplicable(ctx))
+    {
+        test_skipped = true;
+        GTEST_SKIP() << solv.SolverDbId() << "ConvAsm3x3U Not Applicable for this problem"
+                     << conv_config;
+    }
+    const auto invoke_params = miopen::conv::DataInvokeParams{
+        tensors, nullptr, 0, convDesc.attribute.gfx90aFp16alt.GetFwd()};
+
+    ASSERT_TRUE(solv.IsApplicable(ctx));
+    auto sol = solv.GetSolution(ctx, solv.GetDefaultPerformanceConfig(ctx));
+    ASSERT_TRUE(sol.Succeeded());
+    ASSERT_TRUE(sol.invoker_factory);
+    const auto invoker = handle.PrepareInvoker(*sol.invoker_factory, sol.construction_params);
+    (invoker)(handle, invoke_params);
+    handle.Finish();
+}
+
+TEST_P(ConvFwdSolverTestFloat, ConvASM3x3UFwd)
+{
+    SolverFwd<miopen::solver::ConvAsm3x3U>(input.desc,
+                                           in_dev.get(),
+                                           weights.desc,
+                                           wei_dev.get(),
+                                           output.desc,
+                                           out_dev.get(),
+                                           conv_desc,
+                                           conv_config,
+                                           test_skipped);
+}
+
+INSTANTIATE_TEST_SUITE_P(ConvFwdTest,
+                         ConvFwdSolverTestFloat,
+                         testing::Combine(testing::Values(miopenConvolutionFwdAlgoDirect),
+                                          testing::ValuesIn(ConvTestConfigs())));

test/test.hpp

@@ -93,8 +93,6 @@ inline void expect_op(const TLeft& left,
 #define EXPECT_EQUAL(LEFT, RIGHT) EXPECT_OP(LEFT, ==, RIGHT)
 #define STATUS(...) EXPECT((__VA_ARGS__) == 0)
 
-#define FAIL(...) failed(__VA_ARGS__, __FILE__, __LINE__)
-
 template <class F>
 bool throws(F f)
 {