Skip to content

Commit

Permalink
[TEST] forbid GPU tests without explicit datatypes (#3282)
Browse files Browse the repository at this point in the history 
* forbid GPU tests without explicit datatypes

* fix typo
  • Loading branch information
@CAHEK7
CAHEK7 authored Oct 1, 2024
1 parent c233b80 commit d68dbc1
Show file tree
Hide file tree
Showing 2 changed files with 157 additions and 117 deletions.
2 changes: 2 additions & 0 deletions test/gtest/check_names.py
View file
Original file line number Diff line number Diff line change
Expand Up @@ -94,6 +94,8 @@ def parse_tests(args):
mismatches[line] += " Hw"
if not datatype:
mismatches[line] += " Datatype"
if hw and hw.group() == "GPU" and datatype and ("NONE" in datatype.group()):
mismatches[line] += " Hw and Datatype combination (GPU+NONE)"

for l, k in mismatches.items():
logger.warning("Name: " + l + " Mismatch types:" + k)
Expand Down
272 changes: 155 additions & 117 deletions test/gtest/kernel_tuning_net.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -10,71 +10,83 @@ struct KernelTuningNetTestCase : AIModelTestCase
std::string arch;
};

std::vector<KernelTuningNetTestCase> GetConvAsm1x1UTestCases()
std::vector<KernelTuningNetTestCase> GetConvAsm1x1UTestCases_FP32()
{
return {{{{1, 512, 192, 288, {56, 56}, {1, 1}, {0, 0}, {1, 1}, {1, 1}},
miopen::conv::Direction::BackwardData,
miopenFloat,
miopenTensorNCHW},
"1,16,1,64,2,2,1,4",
"gfx908"},
{{{1, 256, 2048, 512, {7, 7}, {1, 1}, {0, 0}, {1, 1}, {1, 1}},
"gfx908"}};
}
std::vector<KernelTuningNetTestCase> GetConvAsm1x1UTestCases_FP16()
{
return {{{{1, 256, 2048, 512, {7, 7}, {1, 1}, {0, 0}, {1, 1}, {1, 1}},
miopen::conv::Direction::Forward,
miopenHalf,
miopenTensorNCHW},
"2,8,4,16,1,4,1,4",
"gfx908"}};
}

std::vector<KernelTuningNetTestCase> GetConvHipIgemmGroupFwdXdlopsTestCases()
std::vector<KernelTuningNetTestCase> GetConvHipIgemmGroupFwdXdlopsTestCases_FP32()
{
return {
{{{1, 128, 64, 128, {209, 209}, {3, 3}, {0, 0}, {2, 2}, {1, 1}},
miopen::conv::Direction::Forward,
miopenFloat,
miopenTensorNHWC},
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle<256, 128, 128, 16, Default, 32, 32, 2, 2, "
"4, 4, 4, 1, 1, 1>",
"gfx90a"},
{{{16, 256, 2016, 192, {7, 7}, {1, 1}, {0, 0}, {1, 1}, {1, 1}},
miopen::conv::Direction::Forward,
miopenHalf,
miopenTensorNHWC},
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle<64, 64, 64, 32, Filter1x1Stride1Pad0, 32, "
"32, 2, 2, 1, "
"1, 1, 1, 1>",
"gfx942"},
};
return {{{{1, 128, 64, 128, {209, 209}, {3, 3}, {0, 0}, {2, 2}, {1, 1}},
miopen::conv::Direction::Forward,
miopenFloat,
miopenTensorNHWC},
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle<256, 128, 128, 16, Default, 32, 32, 2, "
"2, 4, 4, 4, 1, 1, 1>",
"gfx90a"}};
}

std::vector<KernelTuningNetTestCase> GetConvHipIgemmGroupFwdXdlopsTestCases_FP16()
{
return {{{{16, 256, 2016, 192, {7, 7}, {1, 1}, {0, 0}, {1, 1}, {1, 1}},
miopen::conv::Direction::Forward,
miopenHalf,
miopenTensorNHWC},
"DeviceGroupedConvFwdMultipleABD_Xdl_CShuffle<64, 64, 64, 32, Filter1x1Stride1Pad0, "
"32, 32, 2, 2, 1, 1, 1, 1, 1>",
"gfx942"}};
}

std::vector<KernelTuningNetTestCase> GetConvHipIgemmGroupBwdXdlopsTestCases()
std::vector<KernelTuningNetTestCase> GetConvHipIgemmGroupBwdXdlopsTestCases_FP32()
{
return {{{{64, 96, 64, 64, {224, 224}, {3, 3}, {1, 1}, {1, 1}, {1, 1}},
miopen::conv::Direction::BackwardData,
miopenFloat,
miopenTensorNHWC},
"DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1<64, 64, 64, 32, 8, 8, Default, 32, "
"32, 2, 2, 1, 1, 1, 1>",
"gfx942"}};
}

std::vector<KernelTuningNetTestCase> GetConvHipIgemmGroupBwdXdlopsTestCases_FP16()
{
return {{{{32, 4, 256, 256, {59, 59}, {3, 3}, {1, 1}, {2, 2}, {1, 1}},
miopen::conv::Direction::BackwardData,
miopenHalf,
miopenTensorNHWC},
"DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1<128, 128, 32, 32, 8, 8, Default, "
"32, 32, 2, 1, 8, 8, 1, 1>",
"gfx90a"},
{{{64, 96, 64, 64, {224, 224}, {3, 3}, {1, 1}, {1, 1}, {1, 1}},
miopen::conv::Direction::BackwardData,
"gfx90a"}};
}

std::vector<KernelTuningNetTestCase> GetConvHipIgemmGroupWrwXdlopsTestCases_FP32()
{
return {{{{1, 512, 3, 64, {219, 219}, {11, 11}, {2, 2}, {4, 4}, {1, 1}},
miopen::conv::Direction::BackwardWeights,
miopenFloat,
miopenTensorNHWC},
"DeviceGroupedConvBwdDataMultipleD_Xdl_CShuffle_v1<64, 64, 64, 32, 8, 8, Default, 32, "
"32, 2, 2, 1, 1, 1, 1>",
"DeviceGroupedConvBwdWeight_Xdl_CShuffle<128, 128, 32, 4, Default, 4, 2, 1, 4, 4, 1, "
"1, 1, 1, 1>+128",
"gfx942"}};
}

std::vector<KernelTuningNetTestCase> GetConvHipIgemmGroupWrwXdlopsTestCases()
std::vector<KernelTuningNetTestCase> GetConvHipIgemmGroupWrwXdlopsTestCases_FP16()
{
return {
{{{1, 512, 3, 64, {219, 219}, {11, 11}, {2, 2}, {4, 4}, {1, 1}},
miopen::conv::Direction::BackwardWeights,
miopenFloat,
miopenTensorNHWC},
"DeviceGroupedConvBwdWeight_Xdl_CShuffle<128, 128, 32, 4, Default, 4, 2, 1, 4, 4, 1, 1, "
"1, 1, 1>+128",
"gfx942"},
{{{32, 1024, 480, 64, {14, 14}, {1, 1}, {0, 0}, {1, 1}, {1, 1}},
miopen::conv::Direction::BackwardWeights,
miopenHalf,
Expand All @@ -86,131 +98,157 @@ std::vector<KernelTuningNetTestCase> GetConvHipIgemmGroupWrwXdlopsTestCases()
miopen::conv::Direction::BackwardWeights,
miopenHalf,
miopenTensorNHWC},
"DeviceGroupedConvBwdWeight_Xdl_CShuffle<64, 64, 32, 4, Default, 8, 2, 1, 8, 4, 8, 2, "
"1, 1, 8>+1",
"DeviceGroupedConvBwdWeight_Xdl_CShuffle<64, 64, 32, 4, Default, 8, 2, 1, 8, 4, 8, 2, 1, "
"1, 8>+1",
"gfx90a"}};
}

struct KernelTuningNetTest : public ::testing::TestWithParam<KernelTuningNetTestCase>
template <typename Solver>
class KernelTuningNetTest : public ::testing::TestWithParam<KernelTuningNetTestCase>
{
protected:
void SetUp() override
void TestParameterPredictionModel()
{
#if MIOPEN_ENABLE_AI_KERNEL_TUNING
auto test_case = GetParam();
miopen::TensorDescriptor input_tensor_desc = miopen::TensorDescriptor(
auto test_case = GetParam();

auto&& handle = get_handle();
miopen::ExecutionContext ctx(&handle);

if(test_case.arch != ctx.GetStream().GetDeviceName())
GTEST_SKIP();

auto input_tensor_desc = miopen::TensorDescriptor(
test_case.data_type, test_case.layout, test_case.conv.GetInput());
miopen::TensorDescriptor weights_tensor_desc = miopen::TensorDescriptor(

auto weights_tensor_desc = miopen::TensorDescriptor(
test_case.data_type, test_case.layout, test_case.conv.GetWeights());
auto conv_desc = test_case.conv.GetConv();
miopen::TensorDescriptor output_desc = conv_desc.GetForwardOutputTensor(

auto conv_desc = test_case.conv.GetConv();

auto output_desc = conv_desc.GetForwardOutputTensor(
input_tensor_desc, weights_tensor_desc, test_case.data_type);
problem = (test_case.direction == miopen::conv::Direction::Forward)
? miopen::conv::ProblemDescription(input_tensor_desc,
weights_tensor_desc,
output_desc,
conv_desc,
test_case.direction)
: miopen::conv::ProblemDescription(output_desc,
weights_tensor_desc,
input_tensor_desc,
conv_desc,
test_case.direction);
expected = test_case.expected_config;
arch = test_case.arch;

auto problem = (test_case.direction == miopen::conv::Direction::Forward)
? miopen::conv::ProblemDescription(input_tensor_desc,
weights_tensor_desc,
output_desc,
conv_desc,
test_case.direction)
: miopen::conv::ProblemDescription(output_desc,
weights_tensor_desc,
input_tensor_desc,
conv_desc,
test_case.direction);

Solver perf_config;
ASSERT_TRUE(perf_config.IsModelApplicable(ctx, problem));

perf_config.HeuristicInit(ctx, problem);
ASSERT_EQ(perf_config.ToString(), test_case.expected_config);
#else
GTEST_SKIP();
#endif
}
miopen::conv::ProblemDescription problem;
std::string arch;
std::string expected;
};

template <typename T>
void TestParameterPredictionModel(miopen::conv::ProblemDescription problem,
std::string expected,
std::string arch)
using GPU_KernelTuningNetTestConvAsm1x1U_FP32 =
KernelTuningNetTest<miopen::solver::conv::PerformanceConfigConvAsm1x1U>;
using GPU_KernelTuningNetTestConvAsm1x1U_FP16 =
KernelTuningNetTest<miopen::solver::conv::PerformanceConfigConvAsm1x1U>;

TEST_P(GPU_KernelTuningNetTestConvAsm1x1U_FP32, ConvAsm1x1UParameterPredictionModel)
{
#if MIOPEN_ENABLE_AI_KERNEL_TUNING
auto&& handle = get_handle();
miopen::ExecutionContext ctx;
ctx.SetStream(&handle);
T perf_config;
if(arch != ctx.GetStream().GetDeviceName())
GTEST_SKIP();
if(!perf_config.IsModelApplicable(ctx, problem))
GTEST_SKIP();
perf_config.HeuristicInit(ctx, problem);
EXPECT_EQ(perf_config.ToString(), expected)
<< "Expected parameters: " << expected
<< "\nPredicted parameters: " << perf_config.ToString();
#else
std::ignore = problem;
std::ignore = expected;
std::ignore = arch;
GTEST_SKIP();
#endif
TestParameterPredictionModel();
}

struct GPU_KernelTuningNetTestConvAsm1x1U_NONE : KernelTuningNetTest
TEST_P(GPU_KernelTuningNetTestConvAsm1x1U_FP16, ConvAsm1x1UParameterPredictionModel)
{
};
TestParameterPredictionModel();
}

struct GPU_KernelTuningNetTestConvHipIgemmGroupFwdXdlops_NONE : KernelTuningNetTest
{
};
using GPU_KernelTuningNetTestConvHipIgemmGroupFwdXdlops_FP32 =
KernelTuningNetTest<miopen::solver::conv::PerformanceConfigHipImplicitGemmGroupFwdXdlops>;

struct GPU_KernelTuningNetTestConvHipIgemmGroupBwdXdlops_NONE : KernelTuningNetTest
{
};
using GPU_KernelTuningNetTestConvHipIgemmGroupFwdXdlops_FP16 =
KernelTuningNetTest<miopen::solver::conv::PerformanceConfigHipImplicitGemmGroupFwdXdlops>;

struct GPU_KernelTuningNetTestConvHipIgemmGroupWrwXdlops_NONE : KernelTuningNetTest
TEST_P(GPU_KernelTuningNetTestConvHipIgemmGroupFwdXdlops_FP32,
ConvHipIgemmGroupFwdXdlopsParameterPredictionModel)
{
};
TestParameterPredictionModel();
}

TEST_P(GPU_KernelTuningNetTestConvAsm1x1U_NONE, ConvAsm1x1UParameterPredictionModel)
TEST_P(GPU_KernelTuningNetTestConvHipIgemmGroupFwdXdlops_FP16,
ConvHipIgemmGroupFwdXdlopsParameterPredictionModel)
{
TestParameterPredictionModel<miopen::solver::conv::PerformanceConfigConvAsm1x1U>(
problem, expected, arch);
TestParameterPredictionModel();
}

TEST_P(GPU_KernelTuningNetTestConvHipIgemmGroupFwdXdlops_NONE,
ConvHipIgemmGroupFwdXdlopsParameterPredictionModel)
using GPU_KernelTuningNetTestConvHipIgemmGroupBwdXdlops_FP32 =
KernelTuningNetTest<miopen::solver::conv::PerformanceConfigHipImplicitGemmGroupBwdXdlops>;

using GPU_KernelTuningNetTestConvHipIgemmGroupBwdXdlops_FP16 =
KernelTuningNetTest<miopen::solver::conv::PerformanceConfigHipImplicitGemmGroupBwdXdlops>;

TEST_P(GPU_KernelTuningNetTestConvHipIgemmGroupBwdXdlops_FP32,
ConvHipIgemmGroupBwdXdlopsParameterPredictionModel)
{
TestParameterPredictionModel<
miopen::solver::conv::PerformanceConfigHipImplicitGemmGroupFwdXdlops>(
problem, expected, arch);
TestParameterPredictionModel();
}

TEST_P(GPU_KernelTuningNetTestConvHipIgemmGroupBwdXdlops_NONE,
TEST_P(GPU_KernelTuningNetTestConvHipIgemmGroupBwdXdlops_FP16,
ConvHipIgemmGroupBwdXdlopsParameterPredictionModel)
{
TestParameterPredictionModel<
miopen::solver::conv::PerformanceConfigHipImplicitGemmGroupBwdXdlops>(
problem, expected, arch);
TestParameterPredictionModel();
}

TEST_P(GPU_KernelTuningNetTestConvHipIgemmGroupWrwXdlops_NONE,
using GPU_KernelTuningNetTestConvHipIgemmGroupWrwXdlops_FP32 =
KernelTuningNetTest<miopen::solver::conv::PerformanceConfigHipImplicitGemmGroupWrwXdlops>;

using GPU_KernelTuningNetTestConvHipIgemmGroupWrwXdlops_FP16 =
KernelTuningNetTest<miopen::solver::conv::PerformanceConfigHipImplicitGemmGroupWrwXdlops>;

TEST_P(GPU_KernelTuningNetTestConvHipIgemmGroupWrwXdlops_FP32,
ConvHipIgemmGroupWrwXdlopsParameterPredictionModel)
{
TestParameterPredictionModel<
miopen::solver::conv::PerformanceConfigHipImplicitGemmGroupWrwXdlops>(
problem, expected, arch);
TestParameterPredictionModel();
}

TEST_P(GPU_KernelTuningNetTestConvHipIgemmGroupWrwXdlops_FP16,
ConvHipIgemmGroupWrwXdlopsParameterPredictionModel)
{
TestParameterPredictionModel();
}

INSTANTIATE_TEST_SUITE_P(Smoke,
GPU_KernelTuningNetTestConvAsm1x1U_FP32,
testing::ValuesIn(GetConvAsm1x1UTestCases_FP32()));

INSTANTIATE_TEST_SUITE_P(Smoke,
GPU_KernelTuningNetTestConvAsm1x1U_FP16,
testing::ValuesIn(GetConvAsm1x1UTestCases_FP16()));

INSTANTIATE_TEST_SUITE_P(Smoke,
GPU_KernelTuningNetTestConvHipIgemmGroupFwdXdlops_FP32,
testing::ValuesIn(GetConvHipIgemmGroupFwdXdlopsTestCases_FP32()));

INSTANTIATE_TEST_SUITE_P(Smoke,
GPU_KernelTuningNetTestConvHipIgemmGroupFwdXdlops_FP16,
testing::ValuesIn(GetConvHipIgemmGroupFwdXdlopsTestCases_FP16()));

INSTANTIATE_TEST_SUITE_P(Smoke,
GPU_KernelTuningNetTestConvAsm1x1U_NONE,
testing::ValuesIn(GetConvAsm1x1UTestCases()));
GPU_KernelTuningNetTestConvHipIgemmGroupBwdXdlops_FP32,
testing::ValuesIn(GetConvHipIgemmGroupBwdXdlopsTestCases_FP32()));

INSTANTIATE_TEST_SUITE_P(Smoke,
GPU_KernelTuningNetTestConvHipIgemmGroupFwdXdlops_NONE,
testing::ValuesIn(GetConvHipIgemmGroupFwdXdlopsTestCases()));
GPU_KernelTuningNetTestConvHipIgemmGroupBwdXdlops_FP16,
testing::ValuesIn(GetConvHipIgemmGroupBwdXdlopsTestCases_FP16()));

INSTANTIATE_TEST_SUITE_P(Smoke,
GPU_KernelTuningNetTestConvHipIgemmGroupBwdXdlops_NONE,
testing::ValuesIn(GetConvHipIgemmGroupBwdXdlopsTestCases()));
GPU_KernelTuningNetTestConvHipIgemmGroupWrwXdlops_FP32,
testing::ValuesIn(GetConvHipIgemmGroupWrwXdlopsTestCases_FP32()));

INSTANTIATE_TEST_SUITE_P(Smoke,
GPU_KernelTuningNetTestConvHipIgemmGroupWrwXdlops_NONE,
testing::ValuesIn(GetConvHipIgemmGroupWrwXdlopsTestCases()));
GPU_KernelTuningNetTestConvHipIgemmGroupWrwXdlops_FP16,
testing::ValuesIn(GetConvHipIgemmGroupWrwXdlopsTestCases_FP16()));

0 comments on commit d68dbc1

Please sign in to comment.