From d2909a6f6b7607b4b7cbb23b62eb8ef63dc508ab Mon Sep 17 00:00:00 2001 From: xinlipn Date: Mon, 25 Sep 2023 23:30:28 -0700 Subject: [PATCH] [tests] convert test_conv_igemm_mlir_fwd to gTest (#2291) --- test/CMakeLists.txt | 42 ------- test/gtest/CMakeLists.txt | 2 +- test/gtest/conv_igemm_dynamic.cpp | 189 ++++++++++++++++++++++++++++++ 3 files changed, 190 insertions(+), 43 deletions(-) create mode 100644 test/gtest/conv_igemm_dynamic.cpp diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 9f6432ec6d..d2d8a83436 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1218,48 +1218,6 @@ set(DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_BWD_ENVS ${DYNAMIC_IMPLICITGEMM_COMMON} MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC) -if(${CODECOV_TEST}) - add_custom_test(test_conv_igemm_dynamic_small GFX908_DISABLED GFX90A_DISABLED SKIP_XNACK_ON - COMMAND ${DYNAMIC_IMPLICITGEMM_ENVS} $ --verbose --input 32 32 17 17 --weights 32 32 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-backward-data --disable-backward-weights --disable-validation - COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS} $ --verbose --input 64 64 28 28 --weights 32 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data --disable-validation - COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS} $ --verbose --input 64 64 28 28 --weights 16 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights --disable-validation - ) - set_tests_properties(test_conv_igemm_dynamic_small PROPERTIES COST 800) -else() - add_custom_test(test_conv_igemm_dynamic_small GFX908_DISABLED GFX90A_DISABLED SKIP_XNACK_ON - COMMAND ${DYNAMIC_IMPLICITGEMM_ENVS} $ --verbose --input 16 16 56 56 --weights 64 16 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND ${DYNAMIC_IMPLICITGEMM_ENVS} $ --verbose --input 16 64 34 34 --weights 64 64 3 3 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND ${DYNAMIC_IMPLICITGEMM_ENVS} $ --verbose --input 32 32 17 17 --weights 32 32 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND ${DYNAMIC_IMPLICITGEMM_1X1_ENVS} $ --verbose --input 16 384 8 8 --weights 64 384 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS} $ --verbose --input 64 64 28 28 --weights 32 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data - COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS} $ --verbose --input 16 128 36 36 --weights 32 128 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data - COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS} $ --verbose --input 64 64 28 28 --weights 16 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS} $ --verbose --input 16 128 36 36 --weights 32 128 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights - ) -endif() #if CODECOV_TEST - -add_custom_test(test_conv_igemm_dynamic SKIP_UNLESS_ALL GFX908_DISABLED GFX90A_DISABLED SKIP_XNACK_ON -COMMAND ${DYNAMIC_IMPLICITGEMM_ENVS} $ --verbose --input 64 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_ENVS} $ --verbose --input 64 256 34 34 --weights 256 256 3 3 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_ENVS} $ --verbose --input 128 128 35 35 --weights 128 128 3 3 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_ENVS} $ --verbose --input 64 1536 8 8 --weights 256 1536 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_ENVS} $ --verbose --input 128 48 7 7 --weights 128 48 5 5 --pads_strides_dilations 2 2 1 1 1 1 --disable-backward-data --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_ENVS} $ --verbose --input 128 128 17 17 --weights 128 128 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-backward-data --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_1X1_ENVS} $ --verbose --input 128 256 28 28 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_1X1_ENVS} $ --verbose --input 64 1536 8 8 --weights 256 1536 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_1X1_ENVS} $ --verbose --input 128 768 17 17 --weights 128 768 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS} $ --verbose --input 64 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data -COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS} $ --verbose --input 32 128 34 34 --weights 64 128 3 3 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data -COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS} $ --verbose --input 128 128 35 35 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-data -COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS} $ --verbose --input 128 256 56 56 --weights 64 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data -COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS} $ --verbose --input 64 512 28 28 --weights 256 512 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-data -COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS} $ --verbose --input 64 512 14 14 --weights 256 512 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data -COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS} $ --verbose --input 64 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS} $ --verbose --input 32 128 34 34 --weights 64 128 3 3 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS} $ --verbose --input 128 128 35 35 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS} $ --verbose --input 128 256 56 56 --weights 64 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights -) - # gfx90a is disabled due to WORKAROUND_ISSUE_1187 add_custom_test(test_conv_igemm_dynamic_xdlops_bwd SKIP_UNLESS_ALL HALF_ENABLED GFX90A_DISABLED GFX94X_ENABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 28 28 --weights 16 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights diff --git a/test/gtest/CMakeLists.txt b/test/gtest/CMakeLists.txt index 61966d1888..4b7f2aaf08 100644 --- a/test/gtest/CMakeLists.txt +++ b/test/gtest/CMakeLists.txt @@ -35,7 +35,7 @@ function(add_gtest TEST_NAME) target_link_libraries(test_${TEST_NAME} gtest_main MIOpen ${Boost_LIBRARIES} hip::host $) endif() # Enable CMake to discover the test binary - gtest_discover_tests(test_${TEST_NAME} PROPERTIES ENVIRONMENT "MIOPEN_USER_DB_PATH=${CMAKE_CURRENT_BINARY_DIR};MIOPEN_TEST_ALL=${MIOPEN_TEST_ALL};MIOPEN_TEST_MLIR=${MIOPEN_TEST_MLIR};MIOPEN_TEST_COMPOSABLEKERNEL=${MIOPEN_TEST_COMPOSABLEKERNEL}") + gtest_discover_tests(test_${TEST_NAME} PROPERTIES ENVIRONMENT "MIOPEN_USER_DB_PATH=${CMAKE_CURRENT_BINARY_DIR};MIOPEN_TEST_FLOAT_ARG=${MIOPEN_TEST_FLOAT_ARG};MIOPEN_TEST_ALL=${MIOPEN_TEST_ALL};MIOPEN_TEST_MLIR=${MIOPEN_TEST_MLIR};MIOPEN_TEST_COMPOSABLEKERNEL=${MIOPEN_TEST_COMPOSABLEKERNEL}") endif() endfunction() diff --git a/test/gtest/conv_igemm_dynamic.cpp b/test/gtest/conv_igemm_dynamic.cpp new file mode 100644 index 0000000000..25a4e179c5 --- /dev/null +++ b/test/gtest/conv_igemm_dynamic.cpp @@ -0,0 +1,189 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (c) 2023 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 + +#include +#include +#include +#include "../conv2d.hpp" +#include "get_handle.hpp" + +using TestCase = std::tuple, std::string>; + +MIOPEN_DECLARE_ENV_VAR(MIOPEN_TEST_GPU_XNACK_ENABLED) + +static bool SkipTest(void) { return miopen::IsEnabled(MIOPEN_TEST_GPU_XNACK_ENABLED{}); } + +void GetArgs(const TestCase& param, std::vector& tokens) +{ + auto env_vars = std::get<0>(param); + for(auto& elem : env_vars) + { + putenv(elem.data()); + } + + auto cmd = std::get<1>(param); + + std::stringstream ss(cmd); + std::istream_iterator begin(ss); + std::istream_iterator end; + while(begin != end) + tokens.push_back(*begin++); +} + +class Conv2dFloat : public testing::TestWithParam> +{ +}; + +void Run2dDriver(miopenDataType_t prec) +{ + + std::vector params; + switch(prec) + { + case miopenFloat: params = Conv2dFloat::GetParam(); break; + case miopenHalf: + case miopenInt8: + case miopenBFloat16: + case miopenInt8x4: + case miopenInt32: + case miopenDouble: + case miopenFloat8: + case miopenBFloat8: + FAIL() << "miopenHalf, miopenInt8, miopenBFloat16, miopenInt8x4, miopenInt32, " + "miopenDouble, miopenFloat8, miopenBFloat8 " + "data type not supported by conv_igemm_dynamic test"; + + default: params = Conv2dFloat::GetParam(); + } + + for(const auto& test_value : params) + { + std::vector tokens; + GetArgs(test_value, tokens); + std::vector ptrs; + + std::transform(tokens.begin(), + tokens.end(), + std::back_inserter(ptrs), + [](const std::string& str) { return str.data(); }); + + testing::internal::CaptureStderr(); + test_drive(ptrs.size(), ptrs.data()); + auto capture = testing::internal::GetCapturedStderr(); + std::cout << capture; + } +}; + +bool IsTestSupportedForDevice(const miopen::Handle& handle) +{ + std::string devName = handle.GetDeviceName(); + if(devName == "gfx900" || devName == "gfx906") + return true; + else + return false; +} + +TEST_P(Conv2dFloat, FloatTest) +{ + const auto& handle = get_handle(); + if(IsTestSupportedForDevice(handle) && !SkipTest()) + { + Run2dDriver(miopenFloat); + } + else + { + GTEST_SKIP(); + } +}; + +std::vector GetTestCases(const std::string& precision) +{ + + std::vector env = { + "MIOPEN_FIND_MODE=normal", + "MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicFwd"}; + std::vector env_1x1 = { + "MIOPEN_FIND_MODE=normal", + "MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicFwd_1x1"}; + std::vector env_wrw = { + "MIOPEN_FIND_MODE=normal", + "MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicWrw"}; + std::vector env_bwd = { + "MIOPEN_FIND_MODE=normal", + "MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicBwd"}; + + std::string v = " --verbose"; + std::string dis_bk_data = " --disable-backward-data"; + std::string dis_bk_wei = " --disable-backward-weights"; + std::string dis_fwd = " --disable-forward"; + std::string dis_vali = " --disable-validation"; + + const std::vector test_cases = { + // clang-format off +#if CODECOV_TEST + TestCase{env, precision + v + " --input 32 32 17 17 --weights 32 32 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_bk_data + dis_bk_wei + dis_vali}, + TestCase{env_wrw, precision + v + " --input 64 64 28 28 --weights 32 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data + dis_vali}, + TestCase{env_bwd, precision + v + " --input 64 64 28 28 --weights 16 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei + dis_vali}, +#else + TestCase{env, precision + v + " --input 16 16 56 56 --weights 64 16 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + TestCase{env, precision + v + " --input 16 64 34 34 --weights 64 64 3 3 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + TestCase{env, precision + v + " --input 32 32 17 17 --weights 32 32 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_bk_data + dis_bk_wei}, + TestCase{env_1x1, precision + v + " --input 16 384 8 8 --weights 64 384 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + TestCase{env_wrw, precision + v + " --input 64 64 28 28 --weights 32 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + TestCase{env_wrw, precision + v + " --input 16 128 36 36 --weights 32 128 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + TestCase{env_bwd, precision + v + " --input 64 64 28 28 --weights 16 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, + TestCase{env_bwd, precision + v + " --input 16 128 36 36 --weights 32 128 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, +#endif + +#if MIOPEN_TEST_ALL + //SKIP_UNLESS_ALL + TestCase{env, precision + v + " --input 64 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + TestCase{env, precision + v + " --input 64 256 34 34 --weights 256 256 3 3 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + TestCase{env, precision + v + " --input 128 128 35 35 --weights 128 128 3 3 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei}, + TestCase{env, precision + v + " --input 64 1536 8 8 --weights 256 1536 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + TestCase{env, precision + v + " --input 128 48 7 7 --weights 128 48 5 5 --pads_strides_dilations 2 2 1 1 1 1" + dis_bk_data + dis_bk_wei}, + TestCase{env, precision + v + " --input 128 128 17 17 --weights 128 128 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_bk_data + dis_bk_wei}, + TestCase{env_1x1, precision + v + " --input 128 256 28 28 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + TestCase{env_1x1, precision + v + " --input 64 1536 8 8 --weights 256 1536 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + TestCase{env_1x1, precision + v + " --input 128 768 17 17 --weights 128 768 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + TestCase{env_wrw, precision + v + " --input 64 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + TestCase{env_wrw, precision + v + " --input 32 128 34 34 --weights 64 128 3 3 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + TestCase{env_wrw, precision + v + " --input 128 128 35 35 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_data}, + TestCase{env_wrw, precision + v + " --input 128 256 56 56 --weights 64 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + TestCase{env_wrw, precision + v + " --input 64 512 28 28 --weights 256 512 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_data}, + TestCase{env_wrw, precision + v + " --input 64 512 14 14 --weights 256 512 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + TestCase{env_bwd, precision + v + " --input 64 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, + TestCase{env_bwd, precision + v + " --input 32 128 34 34 --weights 64 128 3 3 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, + TestCase{env_bwd, precision + v + " --input 128 128 35 35 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_wei}, + TestCase{env_bwd, precision + v + " --input 128 256 56 56 --weights 64 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei} +#endif + // clang-format on + }; + return test_cases; +} + +INSTANTIATE_TEST_SUITE_P(ConvIgemmDynamic, Conv2dFloat, testing::Values(GetTestCases("--float")));