Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[TEST] forbid GPU tests without explicit datatypes #3282

Merged
merged 2 commits into from
Oct 1, 2024
Merged

[TEST] forbid GPU tests without explicit datatypes #3282

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
e360174
forbid GPU tests without explicit datatypes
CAHEK7 Sep 30, 2024
5cd1e1b
fix typo
CAHEK7 Oct 1, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 2 additions & 0 deletions test/gtest/check_names.py
View file
Open in desktop
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
Open in desktop
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));
BrianHarrisonAMD marked this conversation as resolved.
Show resolved Hide resolved

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()));
Loading