diff --git a/test/gtest/check_names.py b/test/gtest/check_names.py index 86ccd490e8..6104c952a0 100755 --- a/test/gtest/check_names.py +++ b/test/gtest/check_names.py @@ -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) diff --git a/test/gtest/kernel_tuning_net.cpp b/test/gtest/kernel_tuning_net.cpp index 047b63cd40..3886461795 100644 --- a/test/gtest/kernel_tuning_net.cpp +++ b/test/gtest/kernel_tuning_net.cpp @@ -10,15 +10,18 @@ struct KernelTuningNetTestCase : AIModelTestCase std::string arch; }; -std::vector GetConvAsm1x1UTestCases() +std::vector 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 GetConvAsm1x1UTestCases_FP16() +{ + return {{{{1, 256, 2048, 512, {7, 7}, {1, 1}, {0, 0}, {1, 1}, {1, 1}}, miopen::conv::Direction::Forward, miopenHalf, miopenTensorNCHW}, @@ -26,28 +29,40 @@ std::vector GetConvAsm1x1UTestCases() "gfx908"}}; } -std::vector GetConvHipIgemmGroupFwdXdlopsTestCases() +std::vector 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 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 GetConvHipIgemmGroupBwdXdlopsTestCases() +std::vector 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 GetConvHipIgemmGroupBwdXdlopsTestCases_FP16() { return {{{{32, 4, 256, 256, {59, 59}, {3, 3}, {1, 1}, {2, 2}, {1, 1}}, miopen::conv::Direction::BackwardData, @@ -55,26 +70,23 @@ std::vector GetConvHipIgemmGroupBwdXdlopsTestCases() 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 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 GetConvHipIgemmGroupWrwXdlopsTestCases() +std::vector 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, @@ -86,131 +98,157 @@ std::vector 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 +template +class KernelTuningNetTest : public ::testing::TestWithParam { 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 -void TestParameterPredictionModel(miopen::conv::ProblemDescription problem, - std::string expected, - std::string arch) +using GPU_KernelTuningNetTestConvAsm1x1U_FP32 = + KernelTuningNetTest; +using GPU_KernelTuningNetTestConvAsm1x1U_FP16 = + KernelTuningNetTest; + +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; -struct GPU_KernelTuningNetTestConvHipIgemmGroupBwdXdlops_NONE : KernelTuningNetTest -{ -}; +using GPU_KernelTuningNetTestConvHipIgemmGroupFwdXdlops_FP16 = + KernelTuningNetTest; -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( - problem, expected, arch); + TestParameterPredictionModel(); } -TEST_P(GPU_KernelTuningNetTestConvHipIgemmGroupFwdXdlops_NONE, - ConvHipIgemmGroupFwdXdlopsParameterPredictionModel) +using GPU_KernelTuningNetTestConvHipIgemmGroupBwdXdlops_FP32 = + KernelTuningNetTest; + +using GPU_KernelTuningNetTestConvHipIgemmGroupBwdXdlops_FP16 = + KernelTuningNetTest; + +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; + +using GPU_KernelTuningNetTestConvHipIgemmGroupWrwXdlops_FP16 = + KernelTuningNetTest; + +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()));