Skip to content

Commit

Permalink
[Tests] Replace test_conv_igemm_dynamic_xdlops_bwd with gtest (#2409)
Browse files Browse the repository at this point in the history
  • Loading branch information
@xinlipn
xinlipn authored Feb 15, 2024
1 parent ba2f070 commit b44f182
Show file tree
Hide file tree
Showing 6 changed files with 788 additions and 369 deletions.
310 changes: 0 additions & 310 deletions test/CMakeLists.txt
View file

Large diffs are not rendered by default.

112 changes: 61 additions & 51 deletions test/gtest/conv_igemm_dynamic_xdlops.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -31,9 +31,9 @@
#include "../conv2d.hpp"

MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL)
MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_GPU_XNACK_ENABLED)
MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLOAT_ARG)

namespace {
namespace conv_igemm_dynamic_xdlops {

auto GetTestCases()
{
Expand All @@ -44,7 +44,7 @@ auto GetTestCases()
"ConvAsmImplicitGemmGTCDynamicFwdXdlops;"
"ConvAsmImplicitGemmGTCDynamicWrwXdlops")}};

const std::string v = " --verbose";
const std::string cmd_v = " test_conv2d --verbose";
const std::string dis_bk_data = " --disable-backward-data";
const std::string dis_bk_wei = " --disable-backward-weights";
const std::string dis_fwd = " --disable-forward";
Expand All @@ -53,74 +53,82 @@ auto GetTestCases()
return std::vector{
// clang-format off
//bwd
std::pair{env_xdlops, 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},
std::pair{env_xdlops, 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},
std::pair{env_xdlops, 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},
std::pair{env_xdlops, v + " --input 64 224 17 17 --weights 224 224 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, v + " --input 128 128 35 35 --weights 256 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, v + " --input 128 128 64 64 --weights 256 128 3 3 --pads_strides_dilations 1 1 2 2 1 1" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, v + " --input 128 768 17 17 --weights 256 768 3 3 --pads_strides_dilations 1 1 1 1 2 2" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, v + " --input 3 256 28 28 --weights 80 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, v + " --input 2 256 12 18 --weights 256 256 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, v + " --input 400 256 7 7 --weights 1024 256 7 7 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, v + " --input 400 256 1 1 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, v + " --input 8 16 5 5 --weights 8 16 2 2 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, v + " --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, cmd_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},
std::pair{env_xdlops, cmd_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},
std::pair{env_xdlops, cmd_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},
std::pair{env_xdlops, cmd_v + " --input 64 224 17 17 --weights 224 224 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 128 128 35 35 --weights 256 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 128 128 64 64 --weights 256 128 3 3 --pads_strides_dilations 1 1 2 2 1 1" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 128 768 17 17 --weights 256 768 3 3 --pads_strides_dilations 1 1 1 1 2 2" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 3 256 28 28 --weights 80 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 2 256 12 18 --weights 256 256 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 400 256 7 7 --weights 1024 256 7 7 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 400 256 1 1 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 8 16 5 5 --weights 8 16 2 2 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_wei},
//fwd
//Be careful to add testings for (x=1, y=1, c % 8 != 0) due to WORKAROUND_SWDEV_306318
std::pair{env_xdlops, v + " --input 64 1024 14 14 --weights 1024 1024 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, v + " --input 64 256 56 56 --weights 512 256 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, v + " --input 64 2048 7 7 --weights 2048 2048 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, v + " --input 128 128 17 17 --weights 128 128 7 1 --pads_strides_dilations 3 0 1 1 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, 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},
std::pair{env_xdlops, v + " --input 128 192 17 17 --weights 320 192 3 3 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, v + " --input 128 256 35 35 --weights 64 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, v + " --input 128 48 35 35 --weights 64 48 5 5 --pads_strides_dilations 2 2 1 1 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, v + " --input 64 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, v + " --input 32 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, v + " --input 2 256 100 104 --weights 12 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, v + " --input 1 256 28 28 --weights 80 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 64 1024 14 14 --weights 1024 1024 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 64 256 56 56 --weights 512 256 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 64 2048 7 7 --weights 2048 2048 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 128 128 17 17 --weights 128 128 7 1 --pads_strides_dilations 3 0 1 1 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, cmd_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},
std::pair{env_xdlops, cmd_v + " --input 128 192 17 17 --weights 320 192 3 3 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 128 256 35 35 --weights 64 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 128 48 35 35 --weights 64 48 5 5 --pads_strides_dilations 2 2 1 1 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 64 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 32 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 2 256 100 104 --weights 12 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 1 256 28 28 --weights 80 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei},
//ho=wo=1 stride=2
std::pair{env_xdlops, v + " --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1 " + dis_bk_data + dis_bk_wei},
std::pair{env_xdlops, cmd_v + " --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1 " + dis_bk_data + dis_bk_wei},
//wrw
std::pair{env_xdlops, 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},
std::pair{env_xdlops, 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},
std::pair{env_xdlops, 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},
std::pair{env_xdlops, v + " --input 64 224 17 17 --weights 224 224 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, v + " --input 128 128 35 35 --weights 256 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, v + " --input 128 128 64 64 --weights 256 128 3 3 --pads_strides_dilations 1 1 2 2 1 1" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, v + " --input 128 768 17 17 --weights 256 768 3 3 --pads_strides_dilations 1 1 1 1 2 2" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, v + " --input 3 256 28 28 --weights 80 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, v + " --input 2 256 12 18 --weights 256 256 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, v + " --input 4 512 128 128 --weights 12 512 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, cmd_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},
std::pair{env_xdlops, cmd_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},
std::pair{env_xdlops, cmd_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},
std::pair{env_xdlops, cmd_v + " --input 64 224 17 17 --weights 224 224 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, cmd_v + " --input 128 128 35 35 --weights 256 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, cmd_v + " --input 128 128 64 64 --weights 256 128 3 3 --pads_strides_dilations 1 1 2 2 1 1" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, cmd_v + " --input 128 768 17 17 --weights 256 768 3 3 --pads_strides_dilations 1 1 1 1 2 2" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, cmd_v + " --input 3 256 28 28 --weights 80 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, cmd_v + " --input 2 256 12 18 --weights 256 256 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, cmd_v + " --input 4 512 128 128 --weights 12 512 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data},
//regression test for issue 540
std::pair{env_xdlops, v + " --input 4 32 79 141 --weights 64 32 5 10 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, v + " --input 400 256 7 7 --weights 1024 256 7 7 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, v + " --input 400 256 1 1 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, cmd_v + " --input 4 32 79 141 --weights 64 32 5 10 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, cmd_v + " --input 400 256 7 7 --weights 1024 256 7 7 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, cmd_v + " --input 400 256 1 1 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data},
//Regression test for SWDEV-295434 (FP16 only).
std::pair{env_xdlops, v + " --input 120 256 3 3 --weights 340 256 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_data},
std::pair{env_xdlops, cmd_v + " --input 120 256 3 3 --weights 340 256 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_data},
//ho=wo=1 stride=2
std::pair{env_xdlops, v + " --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1 " + dis_fwd + dis_bk_data}
std::pair{env_xdlops, cmd_v + " --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1 " + dis_fwd + dis_bk_data}
// clang-format on
};
}

using TestCase = decltype(GetTestCases())::value_type;

bool SkipTest()
static bool SkipTest(const std::string& float_arg)
{
return miopen::IsEnabled(ENV(MIOPEN_TEST_GPU_XNACK_ENABLED)) ||
miopen::IsDisabled(ENV(MIOPEN_TEST_ALL));
if(miopen::IsUnset(ENV(MIOPEN_TEST_ALL)))
return false;
if(miopen::IsEnabled(ENV(MIOPEN_TEST_ALL)))
if(miopen::GetStringEnv(ENV(MIOPEN_TEST_FLOAT_ARG)) == float_arg)
return false;
return true;
}

bool IsTestSupportedForDevice()
bool IsTestSupportedForDevice(const miopen::Handle& handle)
{
const auto target = handle.GetTargetProperties();
if(target.Xnack() && *target.Xnack())
return false;
using e_mask = enabled<Gpu::Default>;
using d_mask = disabled<Gpu::gfx900, Gpu::gfx906, Gpu::gfx90A>;
return ::IsTestSupportedForDevMask<d_mask, e_mask>();
}

} // namespace
} // namespace conv_igemm_dynamic_xdlops
using namespace conv_igemm_dynamic_xdlops;

class Conv2dDefaultFloat : public FloatTestCase<std::vector<TestCase>>
{
Expand All @@ -132,7 +140,8 @@ class Conv2dDefaultHalf : public HalfTestCase<std::vector<TestCase>>

TEST_P(Conv2dDefaultFloat, FloatTest_conv_igemm_dynamic_xdlops)
{
if(IsTestSupportedForDevice() && !SkipTest())
const auto& handle = get_handle();
if(IsTestSupportedForDevice(handle) && !SkipTest("--float"))
{
invoke_with_params<conv2d_driver, Conv2dDefaultFloat>(default_check);
}
Expand All @@ -144,7 +153,8 @@ TEST_P(Conv2dDefaultFloat, FloatTest_conv_igemm_dynamic_xdlops)

TEST_P(Conv2dDefaultHalf, HalfTest_conv_igemm_dynamic_xdlops)
{
if(IsTestSupportedForDevice() && !SkipTest())
const auto& handle = get_handle();
if(IsTestSupportedForDevice(handle) && !SkipTest("--half"))
{
invoke_with_params<conv2d_driver, Conv2dDefaultHalf>(default_check);
}
Expand Down
Loading

0 comments on commit b44f182

Please sign in to comment.