From 122de35b8efb4dab0086c51ad7914a1df136a55d Mon Sep 17 00:00:00 2001 From: atamazov Date: Thu, 17 Nov 2022 17:39:37 +0300 Subject: [PATCH 01/66] ata-test-solver-02(01) [NFC] Typo in the error message in ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC. --- src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp b/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp index 0ae936c6ae..18a33d3be1 100644 --- a/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp +++ b/src/solver/conv_asm_implicit_gemm_gtc_wrw_nhwc.cpp @@ -1165,7 +1165,7 @@ ConvSolution ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC::GetSolution( if(workSpace == nullptr || workSpaceSize < required_workspace_size) MIOPEN_THROW("Not enough workspace has been provided for " - "ConvAsmImplicitGemmGTCDynamicWrwXdlops with fp16 and atomic " + "ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC with fp16 and atomic " "add."); auto trans_input_buf = trans_input_size == 0 From 36ad7055f365b23f7a7db5fe456891e7ae4e25bf Mon Sep 17 00:00:00 2001 From: atamazov Date: Thu, 17 Nov 2022 17:59:52 +0300 Subject: [PATCH 02/66] ata-test-solver-02(02) [tests] Added smoke test for solvers ConvAsmImplicitGemmGTCDynamic(Fwd|Bwd)Xdlops --- test/CMakeLists.txt | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index adb2abd880..e89510c935 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1864,9 +1864,13 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmV4R1Dynamic GFX908_DISABLED ${TEST_CONV_VERBOSE_W} --input 1 32 28 28 --weights 32 32 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicWrwXdlops GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED HALF_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlops GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED HALF_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicWrwXdlops $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_W} --input 2 256 12 18 --weights 256 256 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlops $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_B} --input 64 64 28 28 --weights 16 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlops $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 64 512 7 7 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) # Add here regression tests that should be run on Vega10/20 and GFX908 only with FP16. From fa5c77c2ab69d01999c43c64354446e0a09cbae9 Mon Sep 17 00:00:00 2001 From: atamazov Date: Fri, 18 Nov 2022 16:52:09 +0300 Subject: [PATCH 03/66] ata-test-solver-02(03) [tests] Added smoke test for ConvBinWinograd3x3U --- test/CMakeLists.txt | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index e89510c935..5615a4de0f 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1873,6 +1873,13 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlops GFX900_DIS ${TEST_CONV_VERBOSE_F} --input 64 512 7 7 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_conv_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinograd3x3U $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinograd3x3U $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_B} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + # Add here regression tests that should be run on Vega10/20 and GFX908 only with FP16. add_custom_test(test_regression_half_vega_gfx908 FLOAT_DISABLED HALF_ENABLED GFX90A_DISABLED # REGRESSION TEST for issue #894. From a5d96c17c61ddd0889c27abcd1c799adf1ac4b60 Mon Sep 17 00:00:00 2001 From: atamazov Date: Sat, 19 Nov 2022 20:59:29 +0300 Subject: [PATCH 04/66] ata-test-solver-02(04) [tests] Added smoke test for ConvBinWinogradRxS --- test/CMakeLists.txt | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 5615a4de0f..c755e5ffdd 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1880,6 +1880,24 @@ add_custom_test(smoke_conv_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK ${TEST_CONV_VERBOSE_B} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +# F16 is supported for 906 and 906 only, no WrW +add_custom_test(smoke_conv_solver_ConvBinWinogradRxS_f16 GFX900_DISABLED GFX90A_DISABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 1 40 20 20 --weights 20 40 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_B} --input 1 20 20 20 --weights 40 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + +# F32 is supported for 900, 906 and 908. +add_custom_test(smoke_conv_solver_ConvBinWinogradRxS_f32 GFX90A_DISABLED SKIP_XNACK_ON + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_B} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_W} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + # Add here regression tests that should be run on Vega10/20 and GFX908 only with FP16. add_custom_test(test_regression_half_vega_gfx908 FLOAT_DISABLED HALF_ENABLED GFX90A_DISABLED # REGRESSION TEST for issue #894. From 6d0721ef5c83974a4208afcb93dea955addc5ac9 Mon Sep 17 00:00:00 2001 From: atamazov Date: Sun, 20 Nov 2022 00:14:59 +0300 Subject: [PATCH 05/66] ata-test-solver-02(05) [tests] Added smoke test for ConvOclBwdWrW53 --- test/CMakeLists.txt | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index c755e5ffdd..e55ec4b833 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1848,11 +1848,13 @@ add_custom_test(smoke_conv_solver_ConvAsm_5x10_7x7 GFX90A_DISABLED SKIP_XNACK_ON ${TEST_CONV_VERBOSE_F} --input 1 3 224 224 --weights 64 3 7 7 --pads_strides_dilations 3 3 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvOclDirectFwd_11x11_Gen HALF_ENABLED BF16_ENABLED GFX103X_ENABLED +add_custom_test(smoke_conv_solver_ConvOcl_SET01 GFX103X_ENABLED HALF_ENABLED BF16_ENABLED COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclDirectFwd11x11 $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 1 44 44 --weights 1 1 11 11 --pads_strides_dilations 0 0 4 4 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclDirectFwdGen $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 1 6 6 --weights 1 1 3 3 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclBwdWrW53 $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_W} --input 16 1 7 7 --weights 1 1 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmV4R1Dynamic GFX908_DISABLED GFX90A_DISABLED SKIP_XNACK_ON From e6ff153eae9707f721621c6f118d7978a6dcad46 Mon Sep 17 00:00:00 2001 From: atamazov Date: Sun, 20 Nov 2022 01:42:37 +0300 Subject: [PATCH 06/66] ata-test-solver-02(06) [HOTFIX][tests] Fixed defect in #1844 which results in omission of all custom tests --- test/CMakeLists.txt | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index e55ec4b833..5dd8079fdb 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -244,7 +244,6 @@ message(STATUS "SKIP_ALL_EXCEPT_TESTS: ${SKIP_ALL_EXCEPT_TESTS}") set(XNACK_TESTS test_mdgraph) function(add_test_command NAME EXE) - # Restrict the use of SKIP_ALL_EXCEPT_TESTS list in the Int8, BF16 and MIOpenTensile tests if( (NOT (NAME IN_LIST SKIP_ALL_EXCEPT_TESTS) AND SKIP_ALL_EXCEPT_TESTS) OR (NAME IN_LIST SKIP_TESTS) ) @@ -567,7 +566,7 @@ function(add_custom_test NAME) if( (is_gfx900_check OR is_gfx906_check OR is_gfx908_check OR is_gfx103x_check OR is_gfx90a_check) AND is_full_check AND is_xnack_on_check - AND (is_miotensile_check AND is_mlir_check) + AND is_mlir_check AND (is_half_check OR is_bfloat16_check OR is_int8_check OR is_float_check) AND (is_ocl_check AND is_hip_check AND is_hip_nogpu_check) ) From a3266e7dd7a474c124a9e005ef3d2dd8454a77cf Mon Sep 17 00:00:00 2001 From: atamazov Date: Sun, 20 Nov 2022 19:36:00 +0300 Subject: [PATCH 07/66] ata-test-solver-02(07) [tests] Added smoke test for ConvOclBwdWrW1x1 --- test/CMakeLists.txt | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 5dd8079fdb..2f066baf45 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1856,6 +1856,12 @@ add_custom_test(smoke_conv_solver_ConvOcl_SET01 GFX103X_ENABLED HALF_ENABLED BF1 ${TEST_CONV_VERBOSE_W} --input 16 1 7 7 --weights 1 1 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +# GFX103X_DISABLED is due to WORKAROUND_SWDEV_266868 +add_custom_test(smoke_conv_solver_ConvOclBwdWrW1x1 GFX103X_DISABLED HALF_ENABLED BF16_ENABLED + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclBwdWrW1x1 $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_W} --input 1 16 14 14 --weights 16 16 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmV4R1Dynamic GFX908_DISABLED GFX90A_DISABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicFwd $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 16 16 16 16 --weights 16 16 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} From b8fca13874bd01b84395fa2495f81f4383dc87b3 Mon Sep 17 00:00:00 2001 From: atamazov Date: Sun, 20 Nov 2022 20:33:26 +0300 Subject: [PATCH 08/66] ata-test-solver-02(08) [tests] Added smoke test for fft --- test/CMakeLists.txt | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 2f066baf45..2efa80e5e9 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1837,6 +1837,13 @@ if(MIOPEN_TEST_FLOAT) add_custom_test(test_reduce_double SKIP_UNLESS_ALL GFX103X_ENABLED COMMAND $ --double --all --verbose) endif() +add_custom_test(smoke_conv_solver_fft GFX103X_ENABLED + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=fft $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=fft $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_B} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_conv_solver_ConvAsm_5x10_7x7 GFX90A_DISABLED SKIP_XNACK_ON # GFX90A_DISABLED is because of WORKAROUND_ISSUE_1146 COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm5x10u2v2f1 $ ${MIOPEN_TEST_FLOAT_ARG} From 8a3b199af4811975e8b9941578e8db54e8557133 Mon Sep 17 00:00:00 2001 From: atamazov Date: Sun, 20 Nov 2022 20:45:06 +0300 Subject: [PATCH 09/66] ata-test-solver-02(09) [tests] Added smoke test for ConvDirectNaiveConv --- test/CMakeLists.txt | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 2efa80e5e9..da00d4f0d4 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1844,6 +1844,18 @@ add_custom_test(smoke_conv_solver_fft GFX103X_ENABLED ${TEST_CONV_VERBOSE_B} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_conv_solver_ConvDirectNaiveConv_F GFX103X_ENABLED HALF_ENABLED BF16_ENABLED INT8_ENABLED + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvFwd $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + +add_custom_test(smoke_conv_solver_ConvDirectNaiveConv_BW GFX103X_ENABLED HALF_ENABLED BF16_ENABLED + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvBwd $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_B} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvWrw $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_W} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_conv_solver_ConvAsm_5x10_7x7 GFX90A_DISABLED SKIP_XNACK_ON # GFX90A_DISABLED is because of WORKAROUND_ISSUE_1146 COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm5x10u2v2f1 $ ${MIOPEN_TEST_FLOAT_ARG} From 69c0063f02eac6d0dff9850527b04be0bf91f05f Mon Sep 17 00:00:00 2001 From: atamazov Date: Mon, 21 Nov 2022 01:03:00 +0300 Subject: [PATCH 10/66] ata-test-solver-02(10) [tests] Print datatype used for testing --- test/CMakeLists.txt | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index da00d4f0d4..52f12f6168 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -173,6 +173,11 @@ else() set(MIOPENDRIVER_MODE_BN bnorm) endif() +message(STATUS "MIOPEN_TEST_FLOAT ${MIOPEN_TEST_FLOAT}") +message(STATUS "MIOPEN_TEST_HALF ${MIOPEN_TEST_HALF}") +message(STATUS "MIOPEN_TEST_BFLOAT16 ${MIOPEN_TEST_BFLOAT16}") +message(STATUS "MIOPEN_TEST_INT8 ${MIOPEN_TEST_INT8}") + set_var_to_condition(WORKAROUND_ISSUE_1187_DEFAULT MIOPEN_TEST_GFX90A AND MIOPEN_TEST_FLOAT) option( WORKAROUND_ISSUE_1187 "" ${WORKAROUND_ISSUE_1187_DEFAULT}) From cf8788ff9489d5189d9c27b9466d29f1bd35a49e Mon Sep 17 00:00:00 2001 From: atamazov Date: Mon, 21 Nov 2022 01:49:56 +0300 Subject: [PATCH 11/66] ata-test-solver-02(11) [tests] Added smoke test for ConvAsm1x1U (no tuning) --- test/CMakeLists.txt | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 52f12f6168..b1f52ac2e6 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1880,6 +1880,16 @@ add_custom_test(smoke_conv_solver_ConvOcl_SET01 GFX103X_ENABLED HALF_ENABLED BF1 ${TEST_CONV_VERBOSE_W} --input 16 1 7 7 --weights 1 1 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +# FP16 ALT attribute is disabled to enable the backward solver on MI200 for HALF. +add_custom_test(smoke_conv_solver_ConvAsm1x1U HALF_ENABLED SKIP_XNACK_ON + COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + # GFX103X_DISABLED is due to WORKAROUND_SWDEV_266868 add_custom_test(smoke_conv_solver_ConvOclBwdWrW1x1 GFX103X_DISABLED HALF_ENABLED BF16_ENABLED COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclBwdWrW1x1 $ ${MIOPEN_TEST_FLOAT_ARG} From 5b92535c6b6a854451a0b32cd809526992b9d07c Mon Sep 17 00:00:00 2001 From: atamazov Date: Mon, 21 Nov 2022 23:36:52 +0300 Subject: [PATCH 12/66] ata-test-solver-02(12) [NFC][tuning] Factor out some inline functions from header --- src/include/miopen/generic_search.hpp | 24 ------------------------ src/solver/conv_asm_1x1u_stride2.cpp | 8 ++++++++ 2 files changed, 8 insertions(+), 24 deletions(-) diff --git a/src/include/miopen/generic_search.hpp b/src/include/miopen/generic_search.hpp index 3e5939830d..a0b699e94c 100644 --- a/src/include/miopen/generic_search.hpp +++ b/src/include/miopen/generic_search.hpp @@ -211,30 +211,6 @@ class HeartBeat } }; -inline void InitRandomly(std::vector& vec, const double offset, const double factor) -{ - float* p = vec.data(); - for(unsigned long i = 0; i < vec.size(); ++i) - *p++ = static_cast( - (rand() * (1.0 / RAND_MAX) + offset) * // NOLINT (concurrency-mt-unsafe) - factor); -} - -inline void InitRandomly(std::vector& vec) -{ - float* p = vec.data(); - for(unsigned long i = 0; i < vec.size(); ++i) - *p++ = static_cast(rand() * (1.0 / RAND_MAX)); // NOLINT (concurrency-mt-unsafe) -} - -inline size_t divide_round_plus_inf(const size_t x, const unsigned y) -{ - assert(y > 0); - if(x % y != 0) - return x / y + 1; - return x / y; -} - /// Solver member function requirements: /// * GetDefaultPerformanceConfig shall be implemented. /// - Its return type shall be suitable for instantiation of the ComputedContainer. diff --git a/src/solver/conv_asm_1x1u_stride2.cpp b/src/solver/conv_asm_1x1u_stride2.cpp index 17bd4291b3..f8847bec4c 100644 --- a/src/solver/conv_asm_1x1u_stride2.cpp +++ b/src/solver/conv_asm_1x1u_stride2.cpp @@ -99,6 +99,14 @@ inline static bool IsLinear(const int v) return L <= v && v <= H; } +static inline size_t divide_round_plus_inf(const size_t x, const unsigned y) +{ + assert(x >= 0 && y > 0); + if(x % y != 0) + return x / y + 1; + return x / y; +} + enum class MemLayout : int { NCHW = 0, From 076fe70bafb07346c854b55d53568c332a2e4978 Mon Sep 17 00:00:00 2001 From: atamazov Date: Tue, 22 Nov 2022 00:13:20 +0300 Subject: [PATCH 13/66] ata-test-solver-02(13) [tuning] Add support for MIOPEN_DEBUG_TUNING_ITERATIONS_MAX --- src/CMakeLists.txt | 1 + src/generic_search.cpp | 44 +++++++++++++++++++++++++++ src/include/miopen/generic_search.hpp | 9 ++++++ 3 files changed, 54 insertions(+) create mode 100644 src/generic_search.cpp diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 4896209a01..13485d28ff 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -112,6 +112,7 @@ set( MIOpen_Source find_db.cpp fused_api.cpp fusion.cpp + generic_search.cpp handle_api.cpp invoker_cache.cpp kernel_build_params.cpp diff --git a/src/generic_search.cpp b/src/generic_search.cpp new file mode 100644 index 0000000000..47a9d09fa8 --- /dev/null +++ b/src/generic_search.cpp @@ -0,0 +1,44 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (c) 2022 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 + +namespace miopen { +namespace solver { + +MIOPEN_DECLARE_ENV_VAR(MIOPEN_DEBUG_TUNING_ITERATIONS_MAX) + +std::size_t GetTuningIterationsMax() +{ + return Value(MIOPEN_DEBUG_TUNING_ITERATIONS_MAX{}, std::numeric_limits::max()); +} + +} // namespace solver +} // namespace miopen diff --git a/src/include/miopen/generic_search.hpp b/src/include/miopen/generic_search.hpp index a0b699e94c..e46a695658 100644 --- a/src/include/miopen/generic_search.hpp +++ b/src/include/miopen/generic_search.hpp @@ -296,6 +296,8 @@ auto GenericSearch(const Solver s, return GenericSearch(s, ctx, invoke_ctx); } +std::size_t GetTuningIterationsMax(); + template auto GenericSearch(const Solver s, const Context& context_, const AnyInvokeParams& invoke_ctx_) -> decltype(s.GetDefaultPerformanceConfig(context_)) @@ -333,8 +335,11 @@ auto GenericSearch(const Solver s, const Context& context_, const AnyInvokeParam if(!miopen::IsCacheDisabled()) // Otherwise precompilation is useless. { std::vector kernels; + size_t n_current = 0; for(const auto& current_config : all_configs) { + if(n_current >= GetTuningIterationsMax()) + break; ConvSolution current_solution = s.GetSolution(context, current_config); for(auto&& kernel : current_solution.construction_params) { @@ -342,6 +347,7 @@ auto GenericSearch(const Solver s, const Context& context_, const AnyInvokeParam continue; kernels.push_back(kernel); } + ++n_current; } std::ignore = PrecompileKernels(profile_h, kernels); } @@ -351,6 +357,9 @@ auto GenericSearch(const Solver s, const Context& context_, const AnyInvokeParam size_t n_current = 0; for(const auto& current_config : all_configs) { + if(n_current >= GetTuningIterationsMax()) + break; + float elapsed_time = 0.0f; int ret = 0; MIOPEN_LOG_I2('#' << n_current << '/' << n_failed << '/' << n_runs_total << ' ' From a81434e1c15ae9d7137fbf9a99a0f424808fb7d5 Mon Sep 17 00:00:00 2001 From: atamazov Date: Tue, 22 Nov 2022 01:22:58 +0300 Subject: [PATCH 14/66] ata-test-solver-02(14) [tuning] Cleaner logging during tuning --- src/include/miopen/generic_search.hpp | 11 +++++++---- 1 file changed, 7 insertions(+), 4 deletions(-) diff --git a/src/include/miopen/generic_search.hpp b/src/include/miopen/generic_search.hpp index e46a695658..bc5bb8aa7c 100644 --- a/src/include/miopen/generic_search.hpp +++ b/src/include/miopen/generic_search.hpp @@ -38,6 +38,7 @@ #include #include +#include #include #include #include @@ -322,8 +323,10 @@ auto GenericSearch(const Solver s, const Context& context_, const AnyInvokeParam auto& profile_h = context.GetStream(); AutoEnableProfiling enableProfiling{profile_h}; - auto all_configs = GetAllConfigs(s, context); - const int n_runs_total = std::distance(all_configs.begin(), all_configs.end()); + auto all_configs = GetAllConfigs(s, context); + const std::size_t n_runs_total = + std::min(static_cast(std::distance(all_configs.begin(), all_configs.end())), + GetTuningIterationsMax()); bool is_passed = false; // left false only if all iterations failed. float best_time = std::numeric_limits::max(); @@ -338,7 +341,7 @@ auto GenericSearch(const Solver s, const Context& context_, const AnyInvokeParam size_t n_current = 0; for(const auto& current_config : all_configs) { - if(n_current >= GetTuningIterationsMax()) + if(n_current >= n_runs_total) break; ConvSolution current_solution = s.GetSolution(context, current_config); for(auto&& kernel : current_solution.construction_params) @@ -357,7 +360,7 @@ auto GenericSearch(const Solver s, const Context& context_, const AnyInvokeParam size_t n_current = 0; for(const auto& current_config : all_configs) { - if(n_current >= GetTuningIterationsMax()) + if(n_current >= n_runs_total) break; float elapsed_time = 0.0f; From 049514d941505f5d55231eceb656913cda9d6d8f Mon Sep 17 00:00:00 2001 From: atamazov Date: Tue, 22 Nov 2022 02:34:38 +0300 Subject: [PATCH 15/66] ata-test-solver-02(15) [tests] Added tuning to smoke test for ConvAsm1x1U --- test/CMakeLists.txt | 37 +++++++++++++++++++++++++++++++++++-- 1 file changed, 35 insertions(+), 2 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index b1f52ac2e6..26aaab748b 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -455,6 +455,7 @@ endfunction() # Testing mode: # SKIP_UNLESS_ALL - The test should be only run if MIOPEN_TEST_ALL=TRUE. Intended for long tests. # TEST_PERF_DB_RECORD_NOT_FOUND - Test should fail if output contains: "Perf Db: record not found". +# TEST_ANY_ERROR - In addition to the standard checks, the test should fail if output contains "Error" or "failed". # SKIP_XNACK_ON - Do not run the test if XNACK mode is enabled (xnack+) on the GPU. # SKIP_UNLESS_MLIR - The test should be only run if MIOPEN_TEST_MLIR=TRUE. # @@ -468,7 +469,7 @@ function(add_custom_test NAME) BF16_ENABLED BF16_DISABLED HALF_ENABLED HALF_DISABLED INT8_ENABLED INT8_DISABLED FLOAT_ENABLED FLOAT_DISABLED GFX900_ENABLED GFX900_DISABLED GFX906_ENABLED GFX906_DISABLED GFX908_ENABLED GFX908_DISABLED GFX103X_ENABLED GFX103X_DISABLED GFX90A_ENABLED GFX90A_DISABLED - SKIP_UNLESS_MLIR SKIP_UNLESS_ALL TEST_PERF_DB_RECORD_NOT_FOUND SKIP_XNACK_ON + SKIP_UNLESS_MLIR SKIP_UNLESS_ALL TEST_PERF_DB_RECORD_NOT_FOUND TEST_ANY_ERROR SKIP_XNACK_ON OCL_ENABLED OCL_DISABLED HIP_ENABLED HIP_DISABLED HIP_NOGPU_ENABLED HIP_NOGPU_DISABLED ) set(oneValueArgs) @@ -575,11 +576,18 @@ function(add_custom_test NAME) AND (is_half_check OR is_bfloat16_check OR is_int8_check OR is_float_check) AND (is_ocl_check AND is_hip_check AND is_hip_nogpu_check) ) + if(PARSE_TEST_PERF_DB_RECORD_NOT_FOUND AND PARSE_TEST_ANY_ERROR) + message(FATAL_ERROR " TEST_PERF_DB_RECORD_NOT_FOUND and TEST_ANY_ERROR should not be used together") + endif() + if(PARSE_TEST_PERF_DB_RECORD_NOT_FOUND) set_tests_properties(${NAME} PROPERTIES FAIL_REGULAR_EXPRESSION "(FAILED)|(Perf Db: record not found)") + elseif(PARSE_TEST_ANY_ERROR) + set_tests_properties(${NAME} PROPERTIES FAIL_REGULAR_EXPRESSION "(FAILED)|(Error)|(failed)") else() set_tests_properties(${NAME} PROPERTIES FAIL_REGULAR_EXPRESSION "FAILED") endif() + rate_added_test(${NAME}) else() set_tests_properties(${NAME} PROPERTIES DISABLED On) @@ -1880,14 +1888,39 @@ add_custom_test(smoke_conv_solver_ConvOcl_SET01 GFX103X_ENABLED HALF_ENABLED BF1 ${TEST_CONV_VERBOSE_W} --input 16 1 7 7 --weights 1 1 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +# NOTES ON WRITEING TESTS FOR TUNABLE SOLVERS +# +# At first, enforce tuning (SEARCH_DB_UPDATE). Use TEST_ANY_ERROR. +# It leads to test failure in case of any "Error" message output to the log, +# which happen if something is broken in the tuning machinery. +# +# Use MIOPEN_DEBUG_TUNING_ITERATIONS_MAX to save testing time. +# Note the the side effect of the limitation is that sub-optimal tuning +# results are written to the user-perf-db. +# +# That is why the second step, DB_CLEAN, is nececcary. It cleans the sub-optimal values from +# the user-perf-db. This prevents performance degradation on the user systems after the tests. + # FP16 ALT attribute is disabled to enable the backward solver on MI200 for HALF. -add_custom_test(smoke_conv_solver_ConvAsm1x1U HALF_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_conv_solver_ConvAsm1x1U HALF_ENABLED SKIP_XNACK_ON TEST_ANY_ERROR COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE + MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_ENFORCE=DB_CLEAN + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-validation ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE + MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_ENFORCE=DB_CLEAN + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-validation ${MIOPEN_TEST_FLAGS_ARGS} ) # GFX103X_DISABLED is due to WORKAROUND_SWDEV_266868 From ef476c0502f9d7ab19b36d8d4d5aa3916308a1a1 Mon Sep 17 00:00:00 2001 From: atamazov Date: Wed, 23 Nov 2022 19:37:08 +0300 Subject: [PATCH 16/66] ata-test-solver-02(16) [NFC] Removed excessive IsGfx90aFp16altRequired() check from ConvAsm1x1UV2 is it does not support FP16 --- src/solver/conv_asm_1x1u_stride2.cpp | 3 --- 1 file changed, 3 deletions(-) diff --git a/src/solver/conv_asm_1x1u_stride2.cpp b/src/solver/conv_asm_1x1u_stride2.cpp index f8847bec4c..09c430eb44 100644 --- a/src/solver/conv_asm_1x1u_stride2.cpp +++ b/src/solver/conv_asm_1x1u_stride2.cpp @@ -505,9 +505,6 @@ bool ConvAsm1x1UV2::IsApplicable(const ConvolutionContext& ctx, return false; } - if(name == "gfx90a" && problem.conv_problem.IsGfx90aFp16altRequired()) - return false; - const auto elements_in_dword = 4 / GetTypeSize(problem.in_data_type); // clang-format off const auto img_hw = problem.out_height * problem.out_width; From 06372303e4418733f78c1bfec5b2aa41a0368125 Mon Sep 17 00:00:00 2001 From: atamazov Date: Wed, 23 Nov 2022 21:11:24 +0300 Subject: [PATCH 17/66] ata-test-solver-02(17) [tests] Added smoke test for ConvAsm1x1UV2 (with tuning) --- test/CMakeLists.txt | 15 +++++++++++++++ 1 file changed, 15 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 26aaab748b..a1d8f77c57 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1923,6 +1923,21 @@ add_custom_test(smoke_conv_solver_ConvAsm1x1U HALF_ENABLED SKIP_XNACK_ON TEST_AN ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-validation ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_conv_solver_ConvAsm1x1UV2 SKIP_XNACK_ON TEST_ANY_ERROR + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1UV2 $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_ENFORCE=DB_CLEAN + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1UV2 $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-validation ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1UV2 $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_ENFORCE=DB_CLEAN + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1UV2 $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-validation ${MIOPEN_TEST_FLAGS_ARGS} +) + # GFX103X_DISABLED is due to WORKAROUND_SWDEV_266868 add_custom_test(smoke_conv_solver_ConvOclBwdWrW1x1 GFX103X_DISABLED HALF_ENABLED BF16_ENABLED COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclBwdWrW1x1 $ ${MIOPEN_TEST_FLOAT_ARG} From daa7a0f224a751c5917422acfe4dcbf20b00535e Mon Sep 17 00:00:00 2001 From: atamazov Date: Wed, 23 Nov 2022 22:11:11 +0300 Subject: [PATCH 18/66] ata-test-solver-02(18) [tests] Use MIOPEN_USER_DB_PATH instead of DB_CLEAN --- test/CMakeLists.txt | 50 +++++++++++++++------------------------------ 1 file changed, 16 insertions(+), 34 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index a1d8f77c57..96f2cde5f9 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1888,54 +1888,36 @@ add_custom_test(smoke_conv_solver_ConvOcl_SET01 GFX103X_ENABLED HALF_ENABLED BF1 ${TEST_CONV_VERBOSE_W} --input 16 1 7 7 --weights 1 1 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -# NOTES ON WRITEING TESTS FOR TUNABLE SOLVERS -# -# At first, enforce tuning (SEARCH_DB_UPDATE). Use TEST_ANY_ERROR. -# It leads to test failure in case of any "Error" message output to the log, -# which happen if something is broken in the tuning machinery. -# -# Use MIOPEN_DEBUG_TUNING_ITERATIONS_MAX to save testing time. -# Note the the side effect of the limitation is that sub-optimal tuning -# results are written to the user-perf-db. -# -# That is why the second step, DB_CLEAN, is nececcary. It cleans the sub-optimal values from -# the user-perf-db. This prevents performance degradation on the user systems after the tests. +# NOTES ON WRITING TESTS FOR TUNABLE SOLVERS +# * Enforce tuning (SEARCH_DB_UPDATE). +# * Use TEST_ANY_ERROR. This flag leads to test failure in case of any "Error" +# message output to the log, which happens if something is broken in the tuning machinery. +# * Use MIOPEN_DEBUG_TUNING_ITERATIONS_MAX to save testing time. +# Note the the side effect of the limitation is that sub-optimal tuning +# results are written to the user-perf-db. +# * That is why MIOPEN_USER_DB_PATH=. is used. The user-perf-db is written in +# "test" directory. This prevents performance degradation on the user systems +# after the tests complete. # FP16 ALT attribute is disabled to enable the backward solver on MI200 for HALF. add_custom_test(smoke_conv_solver_ConvAsm1x1U HALF_ENABLED SKIP_XNACK_ON TEST_ANY_ERROR - COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE - MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_USER_DB_PATH=. MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_ENFORCE=DB_CLEAN - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U $ ${MIOPEN_TEST_FLOAT_ARG} - ${TEST_CONV_VERBOSE_F} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-validation ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE - MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_USER_DB_PATH=. MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_ENFORCE=DB_CLEAN - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U $ ${MIOPEN_TEST_FLOAT_ARG} - ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-validation ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_conv_solver_ConvAsm1x1UV2 SKIP_XNACK_ON TEST_ANY_ERROR - COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_USER_DB_PATH=. MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1UV2 $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_ENFORCE=DB_CLEAN - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1UV2 $ ${MIOPEN_TEST_FLOAT_ARG} - ${TEST_CONV_VERBOSE_F} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-validation ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_USER_DB_PATH=. MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1UV2 $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_ENFORCE=DB_CLEAN - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1UV2 $ ${MIOPEN_TEST_FLOAT_ARG} - ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-validation ${MIOPEN_TEST_FLAGS_ARGS} ) # GFX103X_DISABLED is due to WORKAROUND_SWDEV_266868 From 094cbaf2c133a2a1100e308c1f39ce96971c2845 Mon Sep 17 00:00:00 2001 From: atamazov Date: Thu, 24 Nov 2022 01:25:42 +0300 Subject: [PATCH 19/66] ata-test-solver-02(19) [tests] Added smoke test for ConvWinograd3x3MultipassWrW (3-3 only) --- test/CMakeLists.txt | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 96f2cde5f9..f71504d9c0 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1969,6 +1969,12 @@ add_custom_test(smoke_conv_solver_ConvBinWinogradRxS_f32 GFX90A_DISABLED SKIP_XN ${TEST_CONV_VERBOSE_W} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +# GFX90A_DISABLED is due to WORKAROUND_ISSUE_1146 +add_custom_test(smoke_conv_solver_ConvWinograd3x3MultipassWrW GFX90A_DISABLED HALF_ENABLED BF16_ENABLED SKIP_XNACK_ON OCL_DISABLED + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER='ConvWinograd3x3MultipassWrW<3-3>' $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_W} --input 1 64 30 30 --weights 64 64 3 3 --pads_strides_dilations 1 1 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + # Add here regression tests that should be run on Vega10/20 and GFX908 only with FP16. add_custom_test(test_regression_half_vega_gfx908 FLOAT_DISABLED HALF_ENABLED GFX90A_DISABLED # REGRESSION TEST for issue #894. From f557eea372fe80c89c36725639a9e76acae27ce5 Mon Sep 17 00:00:00 2001 From: atamazov Date: Tue, 29 Nov 2022 01:49:20 +0300 Subject: [PATCH 20/66] ata-test-solver-02(20) [Fix][tests] Use MIOPEN_USER_DB_PATH to prevent potential performance degradation after tests. --- test/CMakeLists.txt | 25 ++++++++++++++++--------- test/gtest/CMakeLists.txt | 2 +- 2 files changed, 17 insertions(+), 10 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index f71504d9c0..f721517388 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -280,6 +280,7 @@ function(add_test_command NAME EXE) add_test(NAME ${NAME} COMMAND ${EXE} ${ARGN}) endif() endif() + set_tests_properties(${NAME} PROPERTIES ENVIRONMENT "MIOPEN_USER_DB_PATH=${CMAKE_CURRENT_BINARY_DIR}") endfunction() separate_arguments(MIOPEN_TEST_FLAGS_ARGS UNIX_COMMAND ${MIOPEN_TEST_FLAGS}) @@ -560,6 +561,17 @@ function(add_custom_test NAME) add_custom_target(${NAME} ${PARSE_UNPARSED_ARGUMENTS}) add_test(NAME ${NAME} COMMAND ${CMAKE_COMMAND} --build ${CMAKE_CURRENT_BINARY_DIR} --target ${NAME}) + # The tests often change the contents of the user databases, which may affect performance after testing. + # For example, MIOPEN_DEBUG_FIND_ONLY_SOLVER results in non-optimal records in user-find-db. + # Another example is tuning tests. These use MIOPEN_DEBUG_TUNING_ITERATIONS_MAX to save testing time, + # but the side effect of such savings is sub-optimal tuning values in user-perf-db. + # + # MIOPEN_USER_DB_PATH setting resolves this potential problem. The user databases are written in + # the non-default directory. This preserves the state of the actual user databases + # and prevents performance degradation after the tests complete. + # + set_tests_properties(${NAME} PROPERTIES ENVIRONMENT "MIOPEN_USER_DB_PATH=${CMAKE_CURRENT_BINARY_DIR}") + if(WORKAROUND_ISSUE_1148 AND (${NAME} MATCHES "test_conv_3d" OR ${NAME} MATCHES "test_conv_group" @@ -1893,29 +1905,24 @@ add_custom_test(smoke_conv_solver_ConvOcl_SET01 GFX103X_ENABLED HALF_ENABLED BF1 # * Use TEST_ANY_ERROR. This flag leads to test failure in case of any "Error" # message output to the log, which happens if something is broken in the tuning machinery. # * Use MIOPEN_DEBUG_TUNING_ITERATIONS_MAX to save testing time. -# Note the the side effect of the limitation is that sub-optimal tuning -# results are written to the user-perf-db. -# * That is why MIOPEN_USER_DB_PATH=. is used. The user-perf-db is written in -# "test" directory. This prevents performance degradation on the user systems -# after the tests complete. # FP16 ALT attribute is disabled to enable the backward solver on MI200 for HALF. add_custom_test(smoke_conv_solver_ConvAsm1x1U HALF_ENABLED SKIP_XNACK_ON TEST_ANY_ERROR - COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_USER_DB_PATH=. MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_USER_DB_PATH=. MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_conv_solver_ConvAsm1x1UV2 SKIP_XNACK_ON TEST_ANY_ERROR - COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_USER_DB_PATH=. MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1UV2 $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_USER_DB_PATH=. MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1UV2 $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) diff --git a/test/gtest/CMakeLists.txt b/test/gtest/CMakeLists.txt index ef6f5b4845..6107f2488d 100644 --- a/test/gtest/CMakeLists.txt +++ b/test/gtest/CMakeLists.txt @@ -29,7 +29,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}) + gtest_discover_tests(test_${TEST_NAME} PROPERTIES ENVIRONMENT "MIOPEN_USER_DB_PATH=${CMAKE_CURRENT_BINARY_DIR}") endif() endfunction() From f9ce39fc360f10a30b37ee19ad990ae6c90f37e3 Mon Sep 17 00:00:00 2001 From: atamazov Date: Fri, 9 Dec 2022 00:53:28 +0300 Subject: [PATCH 21/66] ata-test-solver-02(22) [tests] Added smoke test for ConvAsm3x3U (with tuning) --- test/CMakeLists.txt | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index f721517388..9787c8d306 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1927,6 +1927,15 @@ add_custom_test(smoke_conv_solver_ConvAsm1x1UV2 SKIP_XNACK_ON TEST_ANY_ERROR ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_conv_solver_ConvAsm3x3U SKIP_XNACK_ON TEST_ANY_ERROR + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm3x3U $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 1 4 10 10 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm3x3U $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_B} --input 1 4 10 10 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + # GFX103X_DISABLED is due to WORKAROUND_SWDEV_266868 add_custom_test(smoke_conv_solver_ConvOclBwdWrW1x1 GFX103X_DISABLED HALF_ENABLED BF16_ENABLED COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclBwdWrW1x1 $ ${MIOPEN_TEST_FLOAT_ARG} From c12b4aa52fd46c2c3b78aadb61aa45d03a62884d Mon Sep 17 00:00:00 2001 From: atamazov Date: Mon, 12 Dec 2022 21:05:13 +0300 Subject: [PATCH 22/66] ata-test-solver-03(01) [tests] Added smoke test for ConvAsmBwdWrW1x1 --- test/CMakeLists.txt | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 9787c8d306..5a674e2131 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1936,6 +1936,13 @@ add_custom_test(smoke_conv_solver_ConvAsm3x3U SKIP_XNACK_ON TEST_ANY_ERROR ${TEST_CONV_VERBOSE_B} --input 1 4 10 10 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_conv_solver_ConvAsmBwdWrW1x1 HALF_ENABLED BF16_ENABLED SKIP_XNACK_ON TEST_ANY_ERROR + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmBwdWrW1x1 $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_W} --input 1 4 5 5 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + # GFX103X_DISABLED is due to WORKAROUND_SWDEV_266868 add_custom_test(smoke_conv_solver_ConvOclBwdWrW1x1 GFX103X_DISABLED HALF_ENABLED BF16_ENABLED COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclBwdWrW1x1 $ ${MIOPEN_TEST_FLOAT_ARG} From 831d6cf775c5bec5813871411b80a8ef3fc7b00a Mon Sep 17 00:00:00 2001 From: atamazov Date: Mon, 12 Dec 2022 22:39:52 +0300 Subject: [PATCH 23/66] ata-test-solver-03(02) [tests] Added smoke test for ConvAsmBwdWrW3x3 --- test/CMakeLists.txt | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 5a674e2131..e09c6b6de1 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1943,6 +1943,20 @@ add_custom_test(smoke_conv_solver_ConvAsmBwdWrW1x1 HALF_ENABLED BF16_ENABLED SKI ${TEST_CONV_VERBOSE_W} --input 1 4 5 5 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +# GFX90A_DISABLED for FP32 because of WORKAROUND_SWDEV_330460 +add_custom_test(smoke_conv_solver_ConvAsmBwdWrW3x3_fp32 GFX90A_DISABLED SKIP_XNACK_ON TEST_ANY_ERROR + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmBwdWrW3x3 $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_W} --input 2 4 3 3 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + +add_custom_test(smoke_conv_solver_ConvAsmBwdWrW3x3_fp16 FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON TEST_ANY_ERROR + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmBwdWrW3x3 $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_W} --input 2 4 3 3 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + # GFX103X_DISABLED is due to WORKAROUND_SWDEV_266868 add_custom_test(smoke_conv_solver_ConvOclBwdWrW1x1 GFX103X_DISABLED HALF_ENABLED BF16_ENABLED COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclBwdWrW1x1 $ ${MIOPEN_TEST_FLOAT_ARG} From 3d8642700977321e71f6ffae8e24e287960a4b88 Mon Sep 17 00:00:00 2001 From: atamazov Date: Mon, 12 Dec 2022 23:44:01 +0300 Subject: [PATCH 24/66] ata-test-solver-03(03) [tests] Added smoke test for ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC --- test/CMakeLists.txt | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index e09c6b6de1..90debe02b6 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1981,6 +1981,20 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlops GFX900_DIS ${TEST_CONV_VERBOSE_F} --input 64 512 7 7 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC_fp32_fp16 GFX900_DISABLED GFX906_DISABLED HALF_ENABLED SKIP_XNACK_ON + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 + --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} +) + +add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC_bf16 GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED FLOAT_DISABLED BF16_ENABLED SKIP_XNACK_ON + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 + --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_conv_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinograd3x3U $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} From 30911f2ca7c475213246c2b6926c4a9a5679920c Mon Sep 17 00:00:00 2001 From: atamazov Date: Mon, 12 Dec 2022 23:55:16 +0300 Subject: [PATCH 25/66] ata-test-solver-03(04) [tests] Added smoke test for ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC --- test/CMakeLists.txt | 12 ++++++++++-- 1 file changed, 10 insertions(+), 2 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 90debe02b6..c7229fce00 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1981,18 +1981,26 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlops GFX900_DIS ${TEST_CONV_VERBOSE_F} --input 64 512 7 7 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC_fp32_fp16 GFX900_DISABLED GFX906_DISABLED HALF_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamic_FB_XdlopsNHWC_fp32_fp16 GFX900_DISABLED GFX906_DISABLED HALF_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_B} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 + --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC_bf16 GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED FLOAT_DISABLED BF16_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamic_FB_XdlopsNHWC_bf16 GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED FLOAT_DISABLED BF16_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_B} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 + --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_conv_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON From 7058fbfe0f3f93f56f4800e86a3b68dc05a7b3c0 Mon Sep 17 00:00:00 2001 From: atamazov Date: Tue, 13 Dec 2022 00:02:22 +0300 Subject: [PATCH 26/66] ata-test-solver-03(05) [tests] Added smoke test for ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC --- test/CMakeLists.txt | 12 ++++++++++-- 1 file changed, 10 insertions(+), 2 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index c7229fce00..7e281f3384 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1981,7 +1981,7 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlops GFX900_DIS ${TEST_CONV_VERBOSE_F} --input 64 512 7 7 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamic_FB_XdlopsNHWC_fp32_fp16 GFX900_DISABLED GFX906_DISABLED HALF_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_fp32_fp16 GFX900_DISABLED GFX906_DISABLED HALF_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 @@ -1990,9 +1990,13 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamic_FB_XdlopsNHWC_fp MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_B} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_W} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 + --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamic_FB_XdlopsNHWC_bf16 GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED FLOAT_DISABLED BF16_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16 GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED FLOAT_DISABLED BF16_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 @@ -2001,6 +2005,10 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamic_FB_XdlopsNHWC_bf MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_B} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_W} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 + --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_conv_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON From cda245c8464d6cb224c45a5e1f5646eef8080bef Mon Sep 17 00:00:00 2001 From: atamazov Date: Tue, 13 Dec 2022 01:57:03 +0300 Subject: [PATCH 27/66] ata-test-solver-03(06) [tests] Added smoke test for ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC --- test/CMakeLists.txt | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 7e281f3384..6e03e10c76 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2011,6 +2011,13 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16 G --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED GFX90A_DISABLED GFX103X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 256 3 3 128 --pads_strides_dilations 0 0 1 1 1 1 + --in_layout NCHW --fil_layout CHWN --out_layout NCHW --tensor_vect 1 --vector_length 4 ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_conv_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinograd3x3U $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} From eaf9db04d0f4e044b7c94cf50bc6ec157e515c23 Mon Sep 17 00:00:00 2001 From: atamazov Date: Tue, 13 Dec 2022 02:14:34 +0300 Subject: [PATCH 28/66] ata-test-solver-03(07) [tests] Added missing TEST_ANY_ERROR to tuning tests --- test/CMakeLists.txt | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 6e03e10c76..ec577b4d41 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1981,7 +1981,8 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlops GFX900_DIS ${TEST_CONV_VERBOSE_F} --input 64 512 7 7 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_fp32_fp16 GFX900_DISABLED GFX906_DISABLED HALF_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_fp32_fp16 GFX900_DISABLED GFX906_DISABLED + HALF_ENABLED SKIP_XNACK_ON TEST_ANY_ERROR COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 @@ -1996,7 +1997,8 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_fp32_f --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16 GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED FLOAT_DISABLED BF16_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16 GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED + FLOAT_DISABLED BF16_ENABLED SKIP_XNACK_ON TEST_ANY_ERROR COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 @@ -2011,7 +2013,8 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16 G --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED GFX90A_DISABLED GFX103X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED GFX90A_DISABLED GFX103X_ENABLED + FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON TEST_ANY_ERROR COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 256 3 3 128 --pads_strides_dilations 0 0 1 1 1 1 From c3daed0c5d0491238a3d1b09dd585259346a23b5 Mon Sep 17 00:00:00 2001 From: atamazov Date: Tue, 13 Dec 2022 02:17:08 +0300 Subject: [PATCH 29/66] ata-test-solver-03(08) [tests] TEST_ANY_ERROR -> TEST_TUNING --- test/CMakeLists.txt | 30 +++++++++++++++--------------- 1 file changed, 15 insertions(+), 15 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index ec577b4d41..f7e974003c 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -456,7 +456,7 @@ endfunction() # Testing mode: # SKIP_UNLESS_ALL - The test should be only run if MIOPEN_TEST_ALL=TRUE. Intended for long tests. # TEST_PERF_DB_RECORD_NOT_FOUND - Test should fail if output contains: "Perf Db: record not found". -# TEST_ANY_ERROR - In addition to the standard checks, the test should fail if output contains "Error" or "failed". +# TEST_TUNING - In addition to the standard checks, the test should fail if output contains "Error" or "failed". # SKIP_XNACK_ON - Do not run the test if XNACK mode is enabled (xnack+) on the GPU. # SKIP_UNLESS_MLIR - The test should be only run if MIOPEN_TEST_MLIR=TRUE. # @@ -470,7 +470,7 @@ function(add_custom_test NAME) BF16_ENABLED BF16_DISABLED HALF_ENABLED HALF_DISABLED INT8_ENABLED INT8_DISABLED FLOAT_ENABLED FLOAT_DISABLED GFX900_ENABLED GFX900_DISABLED GFX906_ENABLED GFX906_DISABLED GFX908_ENABLED GFX908_DISABLED GFX103X_ENABLED GFX103X_DISABLED GFX90A_ENABLED GFX90A_DISABLED - SKIP_UNLESS_MLIR SKIP_UNLESS_ALL TEST_PERF_DB_RECORD_NOT_FOUND TEST_ANY_ERROR SKIP_XNACK_ON + SKIP_UNLESS_MLIR SKIP_UNLESS_ALL TEST_PERF_DB_RECORD_NOT_FOUND TEST_TUNING SKIP_XNACK_ON OCL_ENABLED OCL_DISABLED HIP_ENABLED HIP_DISABLED HIP_NOGPU_ENABLED HIP_NOGPU_DISABLED ) set(oneValueArgs) @@ -588,13 +588,13 @@ function(add_custom_test NAME) AND (is_half_check OR is_bfloat16_check OR is_int8_check OR is_float_check) AND (is_ocl_check AND is_hip_check AND is_hip_nogpu_check) ) - if(PARSE_TEST_PERF_DB_RECORD_NOT_FOUND AND PARSE_TEST_ANY_ERROR) - message(FATAL_ERROR " TEST_PERF_DB_RECORD_NOT_FOUND and TEST_ANY_ERROR should not be used together") + if(PARSE_TEST_PERF_DB_RECORD_NOT_FOUND AND PARSE_TEST_TUNING) + message(FATAL_ERROR " TEST_PERF_DB_RECORD_NOT_FOUND and TEST_TUNING should not be used together") endif() if(PARSE_TEST_PERF_DB_RECORD_NOT_FOUND) set_tests_properties(${NAME} PROPERTIES FAIL_REGULAR_EXPRESSION "(FAILED)|(Perf Db: record not found)") - elseif(PARSE_TEST_ANY_ERROR) + elseif(PARSE_TEST_TUNING) set_tests_properties(${NAME} PROPERTIES FAIL_REGULAR_EXPRESSION "(FAILED)|(Error)|(failed)") else() set_tests_properties(${NAME} PROPERTIES FAIL_REGULAR_EXPRESSION "FAILED") @@ -1902,12 +1902,12 @@ add_custom_test(smoke_conv_solver_ConvOcl_SET01 GFX103X_ENABLED HALF_ENABLED BF1 # NOTES ON WRITING TESTS FOR TUNABLE SOLVERS # * Enforce tuning (SEARCH_DB_UPDATE). -# * Use TEST_ANY_ERROR. This flag leads to test failure in case of any "Error" +# * Use TEST_TUNING. This flag leads to test failure in case of any "Error" # message output to the log, which happens if something is broken in the tuning machinery. # * Use MIOPEN_DEBUG_TUNING_ITERATIONS_MAX to save testing time. # FP16 ALT attribute is disabled to enable the backward solver on MI200 for HALF. -add_custom_test(smoke_conv_solver_ConvAsm1x1U HALF_ENABLED SKIP_XNACK_ON TEST_ANY_ERROR +add_custom_test(smoke_conv_solver_ConvAsm1x1U HALF_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U $ ${MIOPEN_TEST_FLOAT_ARG} @@ -1918,7 +1918,7 @@ add_custom_test(smoke_conv_solver_ConvAsm1x1U HALF_ENABLED SKIP_XNACK_ON TEST_AN ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsm1x1UV2 SKIP_XNACK_ON TEST_ANY_ERROR +add_custom_test(smoke_conv_solver_ConvAsm1x1UV2 SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1UV2 $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} @@ -1927,7 +1927,7 @@ add_custom_test(smoke_conv_solver_ConvAsm1x1UV2 SKIP_XNACK_ON TEST_ANY_ERROR ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsm3x3U SKIP_XNACK_ON TEST_ANY_ERROR +add_custom_test(smoke_conv_solver_ConvAsm3x3U SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm3x3U $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 4 10 10 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} @@ -1936,7 +1936,7 @@ add_custom_test(smoke_conv_solver_ConvAsm3x3U SKIP_XNACK_ON TEST_ANY_ERROR ${TEST_CONV_VERBOSE_B} --input 1 4 10 10 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsmBwdWrW1x1 HALF_ENABLED BF16_ENABLED SKIP_XNACK_ON TEST_ANY_ERROR +add_custom_test(smoke_conv_solver_ConvAsmBwdWrW1x1 HALF_ENABLED BF16_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmBwdWrW1x1 $ ${MIOPEN_TEST_FLOAT_ARG} @@ -1944,13 +1944,13 @@ add_custom_test(smoke_conv_solver_ConvAsmBwdWrW1x1 HALF_ENABLED BF16_ENABLED SKI ) # GFX90A_DISABLED for FP32 because of WORKAROUND_SWDEV_330460 -add_custom_test(smoke_conv_solver_ConvAsmBwdWrW3x3_fp32 GFX90A_DISABLED SKIP_XNACK_ON TEST_ANY_ERROR +add_custom_test(smoke_conv_solver_ConvAsmBwdWrW3x3_fp32 GFX90A_DISABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmBwdWrW3x3 $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_W} --input 2 4 3 3 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsmBwdWrW3x3_fp16 FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON TEST_ANY_ERROR +add_custom_test(smoke_conv_solver_ConvAsmBwdWrW3x3_fp16 FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmBwdWrW3x3 $ ${MIOPEN_TEST_FLOAT_ARG} @@ -1982,7 +1982,7 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlops GFX900_DIS ) add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_fp32_fp16 GFX900_DISABLED GFX906_DISABLED - HALF_ENABLED SKIP_XNACK_ON TEST_ANY_ERROR + HALF_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 @@ -1998,7 +1998,7 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_fp32_f ) add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16 GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED - FLOAT_DISABLED BF16_ENABLED SKIP_XNACK_ON TEST_ANY_ERROR + FLOAT_DISABLED BF16_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 @@ -2014,7 +2014,7 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16 G ) add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED GFX90A_DISABLED GFX103X_ENABLED - FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON TEST_ANY_ERROR + FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 256 3 3 128 --pads_strides_dilations 0 0 1 1 1 1 From 76e6402aaed17bc8161955bbd475c851bacd79ed Mon Sep 17 00:00:00 2001 From: atamazov Date: Tue, 13 Dec 2022 20:11:39 +0300 Subject: [PATCH 30/66] ata-test-solver-03(09) [tests] Added smoke test for ConvCkIgemmFwdV6r1DlopsNchw --- test/CMakeLists.txt | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index f7e974003c..94dff24bec 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2021,6 +2021,14 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC GFX --in_layout NCHW --fil_layout CHWN --out_layout NCHW --tensor_vect 1 --vector_length 4 ${MIOPEN_TEST_FLAGS_ARGS} ) +# MIOPEN_DEBUG_TUNING_ITERATIONS_MAX is set to 2 because kernels are very slow to build. +add_custom_test(smoke_conv_solver_ConvCkIgemmFwdV6r1DlopsNchw GFX103X_ENABLED HALF_ENABLED TEST_TUNING + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=2 + MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvCkIgemmFwdV6r1DlopsNchw $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 128 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_conv_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinograd3x3U $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} From 476f12a0ca1801d3a842e81a0219259c1a81b559 Mon Sep 17 00:00:00 2001 From: atamazov Date: Tue, 13 Dec 2022 22:18:24 +0300 Subject: [PATCH 31/66] ata-test-solver-03(10) [tests] Added smoke test for ConvHipImplicitGemmV4R1Fwd (with tuning) --- test/CMakeLists.txt | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 94dff24bec..b379c47586 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2029,6 +2029,13 @@ add_custom_test(smoke_conv_solver_ConvCkIgemmFwdV6r1DlopsNchw GFX103X_ENABLED HA ${TEST_CONV_VERBOSE_F} --input 128 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_conv_solver_ConvHipImplicitGemmV4R1Fwd GFX103X_ENABLED HALF_ENABLED BF16_ENABLED TEST_TUNING + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R1Fwd $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 256 32 27 27 --weights 128 32 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_conv_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinograd3x3U $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} From 3631c7c474b6acea4d2bd917a4f618b532e7cb2a Mon Sep 17 00:00:00 2001 From: atamazov Date: Tue, 13 Dec 2022 22:31:16 +0300 Subject: [PATCH 32/66] ata-test-solver-03(11) [tests] Added smoke test for ConvHipImplicitGemmV4R4Fwd (with tuning) --- test/CMakeLists.txt | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index b379c47586..d27defa5d3 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2036,6 +2036,12 @@ add_custom_test(smoke_conv_solver_ConvHipImplicitGemmV4R1Fwd GFX103X_ENABLED HAL ${TEST_CONV_VERBOSE_F} --input 256 32 27 27 --weights 128 32 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_conv_solver_ConvHipImplicitGemmV4R4Fwd GFX103X_ENABLED TEST_TUNING + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R4Fwd $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 2 16 28 28 --weights 32 16 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_conv_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinograd3x3U $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} From 020957a352a9fe800a0a1536ec7f291648fb943e Mon Sep 17 00:00:00 2001 From: atamazov Date: Tue, 13 Dec 2022 22:44:52 +0300 Subject: [PATCH 33/66] ata-test-solver-03(12) [tests] Added smoke test for ConvHipImplicitGemmV4R4WrW (with tuning) --- test/CMakeLists.txt | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index d27defa5d3..95bb836c75 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2036,10 +2036,13 @@ add_custom_test(smoke_conv_solver_ConvHipImplicitGemmV4R1Fwd GFX103X_ENABLED HAL ${TEST_CONV_VERBOSE_F} --input 256 32 27 27 --weights 128 32 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvHipImplicitGemmV4R4Fwd GFX103X_ENABLED TEST_TUNING +add_custom_test(smoke_conv_solver_ConvHipImplicitGemmV4R4 GFX103X_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R4Fwd $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 2 16 28 28 --weights 32 16 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R4WrW $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_W} --input 8 128 14 14 --weights 32 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_conv_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON From ea55432f96b24c253b37e34bad400bd2e01a81c2 Mon Sep 17 00:00:00 2001 From: atamazov Date: Wed, 14 Dec 2022 00:02:20 +0300 Subject: [PATCH 34/66] ata-test-solver-03(13) [tests] Added smoke test for ConvHipImplicitGemmV4R1WrW (with tuning) --- test/CMakeLists.txt | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 95bb836c75..59c1af4818 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2029,11 +2029,15 @@ add_custom_test(smoke_conv_solver_ConvCkIgemmFwdV6r1DlopsNchw GFX103X_ENABLED HA ${TEST_CONV_VERBOSE_F} --input 128 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvHipImplicitGemmV4R1Fwd GFX103X_ENABLED HALF_ENABLED BF16_ENABLED TEST_TUNING +add_custom_test(smoke_conv_solver_ConvHipImplicitGemmV4R1 GFX103X_ENABLED HALF_ENABLED BF16_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R1Fwd $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 256 32 27 27 --weights 128 32 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R1WrW $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_W} --input 64 64 55 55 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_conv_solver_ConvHipImplicitGemmV4R4 GFX103X_ENABLED TEST_TUNING From bda93ce6186e27da9d34b3f3ab54e38e583dc822 Mon Sep 17 00:00:00 2001 From: atamazov Date: Wed, 14 Dec 2022 00:19:29 +0300 Subject: [PATCH 35/66] ata-test-solver-03(14) [tests] Added smoke test for ConvHipImplicitGemmBwdDataV4R1 (with tuning) --- test/CMakeLists.txt | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 59c1af4818..589ed20eca 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2040,6 +2040,15 @@ add_custom_test(smoke_conv_solver_ConvHipImplicitGemmV4R1 GFX103X_ENABLED HALF_E ${TEST_CONV_VERBOSE_W} --input 64 64 55 55 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +# MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1=1 is necessary due to WORKAROUND_SWDEV_229277_227616_229195, +# which disables ConvHipImplicitGemmBwdDataV4R1, but we still want to check that the solver is not broken. +add_custom_test(smoke_conv_solver_ConvHipImplicitGemmBwdDataV4R1 GFX103X_ENABLED TEST_TUNING + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1=1 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV4R1 $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_B} --input 16 64 16 16 --weights 64 64 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_conv_solver_ConvHipImplicitGemmV4R4 GFX103X_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R4Fwd $ ${MIOPEN_TEST_FLOAT_ARG} From 7d2c55c60135ba0875083f4ac8ae1ef6597a9861 Mon Sep 17 00:00:00 2001 From: atamazov Date: Wed, 14 Dec 2022 00:46:14 +0300 Subject: [PATCH 36/66] ata-test-solver-03(15) [NFC] ConvHipImplicitGemmBwdDataV1R1: Improve W/A for issue 309. Do not observe FP16_ALT attribute (it is irrelevant as FP16 is not supported) --- src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp | 9 +++------ 1 file changed, 3 insertions(+), 6 deletions(-) diff --git a/src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp b/src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp index f3daccd977..e1d85ed9f2 100644 --- a/src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp +++ b/src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp @@ -644,19 +644,16 @@ bool ConvHipImplicitGemmBwdDataV1R1::IsApplicable(const ConvolutionContext& ctx, if(!problem.Is2d() && !(problem.Is3d() && problem.IsFp32())) return false; - // TBD Renable fp16 once the root cause of - // fp16 failures, observed in CI, is resolved. if(!(problem.IsFp32() || problem.IsBfp16())) return false; if(problem.group_counts != 1) return false; + #if WORKAROUND_ISSUE_309 if(problem.IsBfp16()) - return false; + if(!miopen::IsEnabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1{})) + return false; #endif - if(ctx.GetStream().GetDeviceName() == "gfx90a" && - problem.conv_problem.IsGfx90aFp16altRequired()) - return false; const auto k = ProblemInterpreter::GetOutputChannelK(problem); if(k % GetEPackLength(ctx, problem, false) != 0) From 9b713e96e91330905cc50c862676a542db875eb5 Mon Sep 17 00:00:00 2001 From: atamazov Date: Wed, 14 Dec 2022 00:53:07 +0300 Subject: [PATCH 37/66] ata-test-solver-03(16) [tests] Added smoke test for ConvHipImplicitGemmBwdDataV1R1 (with tuning) --- test/CMakeLists.txt | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 589ed20eca..577f7e1d79 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2029,6 +2029,12 @@ add_custom_test(smoke_conv_solver_ConvCkIgemmFwdV6r1DlopsNchw GFX103X_ENABLED HA ${TEST_CONV_VERBOSE_F} --input 128 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_conv_solver_ConvHipImplicitGemmBwdDataV1R1 GFX103X_ENABLED TEST_TUNING + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV1R1 $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_B} --input 32 128 32 32 --weights 12 128 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_conv_solver_ConvHipImplicitGemmV4R1 GFX103X_ENABLED HALF_ENABLED BF16_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 From f80866630b334fa778498dd2239f7920fdb21398 Mon Sep 17 00:00:00 2001 From: atamazov Date: Wed, 14 Dec 2022 01:22:37 +0300 Subject: [PATCH 38/66] ata-test-solver-03(17) [tests] Added smoke test for ConvHipImplicitGemmBwdDataV4R1Xdlops (with tuning) --- test/CMakeLists.txt | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 577f7e1d79..fb6f0644c4 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2064,6 +2064,18 @@ add_custom_test(smoke_conv_solver_ConvHipImplicitGemmV4R4 GFX103X_ENABLED TEST_T ${TEST_CONV_VERBOSE_W} --input 8 128 14 14 --weights 32 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +# WORKAROUND_ISSUE_1206 disables this solver for FP32 due to precision issues. +# WORKAROUND_SWDEV_329642 disables this solver on MI200 for BF16. +# However we still want to check that these cases are not broken and therefore use +# MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1_XDLOPS=1 to enable the solver. +add_custom_test(smoke_conv_solver_ConvHipImplicitGemmBwdDataV4R1Xdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1_XDLOPS=1 + MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV4R1Xdlops $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_B} --input 1 160 28 28 --weights 128 160 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_conv_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinograd3x3U $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} From af93f6afe835ae871a1d4369d13f04a761f2b5be Mon Sep 17 00:00:00 2001 From: atamazov Date: Wed, 14 Dec 2022 20:03:33 +0300 Subject: [PATCH 39/66] ata-test-solver-03(18) [tests] Added smoke test for ConvHipImplicitGemmForwardV4R4Xdlops (with tuning) --- test/CMakeLists.txt | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index fb6f0644c4..29a6070126 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2076,6 +2076,13 @@ add_custom_test(smoke_conv_solver_ConvHipImplicitGemmBwdDataV4R1Xdlops GFX900_DI ${TEST_CONV_VERBOSE_B} --input 1 160 28 28 --weights 128 160 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_conv_solver_ConvHipImplicitGemmForwardV4R4Xdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmForwardV4R4Xdlops $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 128 48 13 13 --weights 192 48 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_conv_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinograd3x3U $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} From 7adeb51ea355b4fbb503f276a766d901f07186de Mon Sep 17 00:00:00 2001 From: atamazov Date: Wed, 14 Dec 2022 20:14:59 +0300 Subject: [PATCH 40/66] ata-test-solver-03(19) [tests] Added smoke test for ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm (with tuning) --- test/CMakeLists.txt | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 29a6070126..68c8fc3bbc 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2083,6 +2083,13 @@ add_custom_test(smoke_conv_solver_ConvHipImplicitGemmForwardV4R4Xdlops GFX900_DI ${TEST_CONV_VERBOSE_F} --input 128 48 13 13 --weights 192 48 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_conv_solver_ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 16 1 7 7 --weights 1 1 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_conv_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinograd3x3U $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} From 6a0a0462ff6ebd51834da346989a997e7bb1a328 Mon Sep 17 00:00:00 2001 From: atamazov Date: Wed, 14 Dec 2022 20:20:55 +0300 Subject: [PATCH 41/66] ata-test-solver-03(20) [tests][NFC] smoke_conv_solver_Conv* -> smoke_solver_Conv* --- test/CMakeLists.txt | 58 ++++++++++++++++++++++----------------------- 1 file changed, 29 insertions(+), 29 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 68c8fc3bbc..f179a0c999 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1862,26 +1862,26 @@ if(MIOPEN_TEST_FLOAT) add_custom_test(test_reduce_double SKIP_UNLESS_ALL GFX103X_ENABLED COMMAND $ --double --all --verbose) endif() -add_custom_test(smoke_conv_solver_fft GFX103X_ENABLED +add_custom_test(smoke_solver_ConvFFT GFX103X_ENABLED COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=fft $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=fft $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_B} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvDirectNaiveConv_F GFX103X_ENABLED HALF_ENABLED BF16_ENABLED INT8_ENABLED +add_custom_test(smoke_solver_ConvDirectNaiveConv_F GFX103X_ENABLED HALF_ENABLED BF16_ENABLED INT8_ENABLED COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvFwd $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvDirectNaiveConv_BW GFX103X_ENABLED HALF_ENABLED BF16_ENABLED +add_custom_test(smoke_solver_ConvDirectNaiveConv_BW GFX103X_ENABLED HALF_ENABLED BF16_ENABLED COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvBwd $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_B} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvWrw $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_W} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsm_5x10_7x7 GFX90A_DISABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvAsm_5x10_7x7 GFX90A_DISABLED SKIP_XNACK_ON # GFX90A_DISABLED is because of WORKAROUND_ISSUE_1146 COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm5x10u2v2f1 $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 1 5 10 --weights 16 1 5 10 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} @@ -1891,7 +1891,7 @@ add_custom_test(smoke_conv_solver_ConvAsm_5x10_7x7 GFX90A_DISABLED SKIP_XNACK_ON ${TEST_CONV_VERBOSE_F} --input 1 3 224 224 --weights 64 3 7 7 --pads_strides_dilations 3 3 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvOcl_SET01 GFX103X_ENABLED HALF_ENABLED BF16_ENABLED +add_custom_test(smoke_solver_ConvOcl_SET01 GFX103X_ENABLED HALF_ENABLED BF16_ENABLED COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclDirectFwd11x11 $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 1 44 44 --weights 1 1 11 11 --pads_strides_dilations 0 0 4 4 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclDirectFwdGen $ ${MIOPEN_TEST_FLOAT_ARG} @@ -1907,7 +1907,7 @@ add_custom_test(smoke_conv_solver_ConvOcl_SET01 GFX103X_ENABLED HALF_ENABLED BF1 # * Use MIOPEN_DEBUG_TUNING_ITERATIONS_MAX to save testing time. # FP16 ALT attribute is disabled to enable the backward solver on MI200 for HALF. -add_custom_test(smoke_conv_solver_ConvAsm1x1U HALF_ENABLED SKIP_XNACK_ON TEST_TUNING +add_custom_test(smoke_solver_ConvAsm1x1U HALF_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U $ ${MIOPEN_TEST_FLOAT_ARG} @@ -1918,7 +1918,7 @@ add_custom_test(smoke_conv_solver_ConvAsm1x1U HALF_ENABLED SKIP_XNACK_ON TEST_TU ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsm1x1UV2 SKIP_XNACK_ON TEST_TUNING +add_custom_test(smoke_solver_ConvAsm1x1UV2 SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1UV2 $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} @@ -1927,7 +1927,7 @@ add_custom_test(smoke_conv_solver_ConvAsm1x1UV2 SKIP_XNACK_ON TEST_TUNING ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsm3x3U SKIP_XNACK_ON TEST_TUNING +add_custom_test(smoke_solver_ConvAsm3x3U SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm3x3U $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 4 10 10 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} @@ -1936,7 +1936,7 @@ add_custom_test(smoke_conv_solver_ConvAsm3x3U SKIP_XNACK_ON TEST_TUNING ${TEST_CONV_VERBOSE_B} --input 1 4 10 10 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsmBwdWrW1x1 HALF_ENABLED BF16_ENABLED SKIP_XNACK_ON TEST_TUNING +add_custom_test(smoke_solver_ConvAsmBwdWrW1x1 HALF_ENABLED BF16_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmBwdWrW1x1 $ ${MIOPEN_TEST_FLOAT_ARG} @@ -1944,13 +1944,13 @@ add_custom_test(smoke_conv_solver_ConvAsmBwdWrW1x1 HALF_ENABLED BF16_ENABLED SKI ) # GFX90A_DISABLED for FP32 because of WORKAROUND_SWDEV_330460 -add_custom_test(smoke_conv_solver_ConvAsmBwdWrW3x3_fp32 GFX90A_DISABLED SKIP_XNACK_ON TEST_TUNING +add_custom_test(smoke_solver_ConvAsmBwdWrW3x3_fp32 GFX90A_DISABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmBwdWrW3x3 $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_W} --input 2 4 3 3 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsmBwdWrW3x3_fp16 FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON TEST_TUNING +add_custom_test(smoke_solver_ConvAsmBwdWrW3x3_fp16 FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmBwdWrW3x3 $ ${MIOPEN_TEST_FLOAT_ARG} @@ -1958,12 +1958,12 @@ add_custom_test(smoke_conv_solver_ConvAsmBwdWrW3x3_fp16 FLOAT_DISABLED HALF_ENAB ) # GFX103X_DISABLED is due to WORKAROUND_SWDEV_266868 -add_custom_test(smoke_conv_solver_ConvOclBwdWrW1x1 GFX103X_DISABLED HALF_ENABLED BF16_ENABLED +add_custom_test(smoke_solver_ConvOclBwdWrW1x1 GFX103X_DISABLED HALF_ENABLED BF16_ENABLED COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclBwdWrW1x1 $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_W} --input 1 16 14 14 --weights 16 16 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmV4R1Dynamic GFX908_DISABLED GFX90A_DISABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvAsmImplicitGemmV4R1Dynamic GFX908_DISABLED GFX90A_DISABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicFwd $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 16 16 16 16 --weights 16 16 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicBwd $ ${MIOPEN_TEST_FLOAT_ARG} @@ -1972,7 +1972,7 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmV4R1Dynamic GFX908_DISABLED ${TEST_CONV_VERBOSE_W} --input 1 32 28 28 --weights 32 32 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlops GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED HALF_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED HALF_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicWrwXdlops $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_W} --input 2 256 12 18 --weights 256 256 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlops $ ${MIOPEN_TEST_FLOAT_ARG} @@ -1981,7 +1981,7 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlops GFX900_DIS ${TEST_CONV_VERBOSE_F} --input 64 512 7 7 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_fp32_fp16 GFX900_DISABLED GFX906_DISABLED +add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_fp32_fp16 GFX900_DISABLED GFX906_DISABLED HALF_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} @@ -1997,7 +1997,7 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_fp32_f --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16 GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED +add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16 GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED FLOAT_DISABLED BF16_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} @@ -2013,7 +2013,7 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16 G --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED GFX90A_DISABLED GFX103X_ENABLED +add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED GFX90A_DISABLED GFX103X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC $ ${MIOPEN_TEST_FLOAT_ARG} @@ -2022,20 +2022,20 @@ add_custom_test(smoke_conv_solver_ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC GFX ) # MIOPEN_DEBUG_TUNING_ITERATIONS_MAX is set to 2 because kernels are very slow to build. -add_custom_test(smoke_conv_solver_ConvCkIgemmFwdV6r1DlopsNchw GFX103X_ENABLED HALF_ENABLED TEST_TUNING +add_custom_test(smoke_solver_ConvCkIgemmFwdV6r1DlopsNchw GFX103X_ENABLED HALF_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=2 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvCkIgemmFwdV6r1DlopsNchw $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 128 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvHipImplicitGemmBwdDataV1R1 GFX103X_ENABLED TEST_TUNING +add_custom_test(smoke_solver_ConvHipImplicitGemmBwdDataV1R1 GFX103X_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV1R1 $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_B} --input 32 128 32 32 --weights 12 128 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvHipImplicitGemmV4R1 GFX103X_ENABLED HALF_ENABLED BF16_ENABLED TEST_TUNING +add_custom_test(smoke_solver_ConvHipImplicitGemmV4R1 GFX103X_ENABLED HALF_ENABLED BF16_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R1Fwd $ ${MIOPEN_TEST_FLOAT_ARG} @@ -2048,14 +2048,14 @@ add_custom_test(smoke_conv_solver_ConvHipImplicitGemmV4R1 GFX103X_ENABLED HALF_E # MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1=1 is necessary due to WORKAROUND_SWDEV_229277_227616_229195, # which disables ConvHipImplicitGemmBwdDataV4R1, but we still want to check that the solver is not broken. -add_custom_test(smoke_conv_solver_ConvHipImplicitGemmBwdDataV4R1 GFX103X_ENABLED TEST_TUNING +add_custom_test(smoke_solver_ConvHipImplicitGemmBwdDataV4R1 GFX103X_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1=1 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV4R1 $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_B} --input 16 64 16 16 --weights 64 64 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvHipImplicitGemmV4R4 GFX103X_ENABLED TEST_TUNING +add_custom_test(smoke_solver_ConvHipImplicitGemmV4R4 GFX103X_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R4Fwd $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 2 16 28 28 --weights 32 16 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} @@ -2068,7 +2068,7 @@ add_custom_test(smoke_conv_solver_ConvHipImplicitGemmV4R4 GFX103X_ENABLED TEST_T # WORKAROUND_SWDEV_329642 disables this solver on MI200 for BF16. # However we still want to check that these cases are not broken and therefore use # MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1_XDLOPS=1 to enable the solver. -add_custom_test(smoke_conv_solver_ConvHipImplicitGemmBwdDataV4R1Xdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING +add_custom_test(smoke_solver_ConvHipImplicitGemmBwdDataV4R1Xdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1_XDLOPS=1 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 @@ -2076,21 +2076,21 @@ add_custom_test(smoke_conv_solver_ConvHipImplicitGemmBwdDataV4R1Xdlops GFX900_DI ${TEST_CONV_VERBOSE_B} --input 1 160 28 28 --weights 128 160 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvHipImplicitGemmForwardV4R4Xdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING +add_custom_test(smoke_solver_ConvHipImplicitGemmForwardV4R4Xdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmForwardV4R4Xdlops $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 128 48 13 13 --weights 192 48 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING +add_custom_test(smoke_solver_ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 16 1 7 7 --weights 1 1 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinograd3x3U $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinograd3x3U $ ${MIOPEN_TEST_FLOAT_ARG} @@ -2098,7 +2098,7 @@ add_custom_test(smoke_conv_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK ) # F16 is supported for 906 and 906 only, no WrW -add_custom_test(smoke_conv_solver_ConvBinWinogradRxS_f16 GFX900_DISABLED GFX90A_DISABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvBinWinogradRxS_f16 GFX900_DISABLED GFX90A_DISABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 40 20 20 --weights 20 40 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${MIOPEN_TEST_FLOAT_ARG} @@ -2106,7 +2106,7 @@ add_custom_test(smoke_conv_solver_ConvBinWinogradRxS_f16 GFX900_DISABLED GFX90A_ ) # F32 is supported for 900, 906 and 908. -add_custom_test(smoke_conv_solver_ConvBinWinogradRxS_f32 GFX90A_DISABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvBinWinogradRxS_f32 GFX90A_DISABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${MIOPEN_TEST_FLOAT_ARG} @@ -2116,7 +2116,7 @@ add_custom_test(smoke_conv_solver_ConvBinWinogradRxS_f32 GFX90A_DISABLED SKIP_XN ) # GFX90A_DISABLED is due to WORKAROUND_ISSUE_1146 -add_custom_test(smoke_conv_solver_ConvWinograd3x3MultipassWrW GFX90A_DISABLED HALF_ENABLED BF16_ENABLED SKIP_XNACK_ON OCL_DISABLED +add_custom_test(smoke_solver_ConvWinograd3x3MultipassWrW GFX90A_DISABLED HALF_ENABLED BF16_ENABLED SKIP_XNACK_ON OCL_DISABLED COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER='ConvWinograd3x3MultipassWrW<3-3>' $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_W} --input 1 64 30 30 --weights 64 64 3 3 --pads_strides_dilations 1 1 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) From 450d8c15effc916a6de1ab97cd4a41c31bafa710 Mon Sep 17 00:00:00 2001 From: atamazov Date: Wed, 14 Dec 2022 20:38:49 +0300 Subject: [PATCH 42/66] ata-test-solver-03(21) [tests] Added smoke test for ConvHipImplicitGemmBwdDataV1R1Xdlops (with tuning) --- test/CMakeLists.txt | 11 +++++++++++ 1 file changed, 11 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index f179a0c999..df5e833a49 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2064,6 +2064,17 @@ add_custom_test(smoke_solver_ConvHipImplicitGemmV4R4 GFX103X_ENABLED TEST_TUNING ${TEST_CONV_VERBOSE_W} --input 8 128 14 14 --weights 32 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +# WORKAROUND_SWDEV_251757 disables this solver due to precision issues. +# However we still want to check that solver is not broken and therefore use +# MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1_XDLOPS=1 to enable it. +add_custom_test(smoke_solver_ConvHipImplicitGemmBwdDataV1R1Xdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1_XDLOPS=1 + MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV1R1Xdlops $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_B} --input 32 128 32 32 --weights 12 128 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + # WORKAROUND_ISSUE_1206 disables this solver for FP32 due to precision issues. # WORKAROUND_SWDEV_329642 disables this solver on MI200 for BF16. # However we still want to check that these cases are not broken and therefore use From a853d8b4e1e06c5b41e6b2cbd456233c4d68111c Mon Sep 17 00:00:00 2001 From: atamazov Date: Wed, 14 Dec 2022 21:34:27 +0300 Subject: [PATCH 43/66] ata-test-solver-03(22) [tests] Added SKIP_UNLESS_COMPOSABLEKERNEL option for custom tests --- test/CMakeLists.txt | 13 +++++++++---- 1 file changed, 9 insertions(+), 4 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index df5e833a49..09571acdee 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -43,6 +43,7 @@ option( MIOPEN_TEST_CONV Off) option( MIOPEN_TEST_DEEPBENCH Off) option( MIOPEN_TEST_DRIVER_ITER_MODE Off) option( MIOPEN_TEST_MLIR "Test for MLIR compilation backend" ${MIOPEN_USE_MLIR} ) +option( MIOPEN_TEST_COMPOSABLEKERNEL "Test with composable_kernel library" ${MIOPEN_USE_COMPOSABLEKERNEL} ) set_var_to_condition(MIOPEN_TEST_WITH_MIOPENDRIVER_DEFAULT MIOPEN_BUILD_DRIVER) option( MIOPEN_TEST_WITH_MIOPENDRIVER "Use MIOpenDriver in tests" ${MIOPEN_TEST_WITH_MIOPENDRIVER_DEFAULT}) @@ -459,6 +460,7 @@ endfunction() # TEST_TUNING - In addition to the standard checks, the test should fail if output contains "Error" or "failed". # SKIP_XNACK_ON - Do not run the test if XNACK mode is enabled (xnack+) on the GPU. # SKIP_UNLESS_MLIR - The test should be only run if MIOPEN_TEST_MLIR=TRUE. +# SKIP_UNLESS_COMPOSABLEKERNEL - The test should be only run if MIOPEN_TEST_COMPOSABLEKERNEL=TRUE. # # Backend: OCL HIP HIP_NOGPU # The option can be enabled or disabled by using '_ENABLED' and '_DISABLED' suffix. @@ -470,7 +472,7 @@ function(add_custom_test NAME) BF16_ENABLED BF16_DISABLED HALF_ENABLED HALF_DISABLED INT8_ENABLED INT8_DISABLED FLOAT_ENABLED FLOAT_DISABLED GFX900_ENABLED GFX900_DISABLED GFX906_ENABLED GFX906_DISABLED GFX908_ENABLED GFX908_DISABLED GFX103X_ENABLED GFX103X_DISABLED GFX90A_ENABLED GFX90A_DISABLED - SKIP_UNLESS_MLIR SKIP_UNLESS_ALL TEST_PERF_DB_RECORD_NOT_FOUND TEST_TUNING SKIP_XNACK_ON + SKIP_UNLESS_MLIR SKIP_UNLESS_COMPOSABLEKERNEL SKIP_UNLESS_ALL TEST_PERF_DB_RECORD_NOT_FOUND TEST_TUNING SKIP_XNACK_ON OCL_ENABLED OCL_DISABLED HIP_ENABLED HIP_DISABLED HIP_NOGPU_ENABLED HIP_NOGPU_DISABLED ) set(oneValueArgs) @@ -504,6 +506,10 @@ function(add_custom_test NAME) bool_not_f(${PARSE_SKIP_UNLESS_MLIR} is_mlir_check) bool_or_f(${is_mlir_check} ${MIOPEN_TEST_MLIR} is_mlir_check) + set(is_composablekernel_check) + bool_not_f(${PARSE_SKIP_UNLESS_COMPOSABLEKERNEL} is_composablekernel_check) + bool_or_f(${is_composablekernel_check} ${MIOPEN_TEST_COMPOSABLEKERNEL} is_composablekernel_check) + set(is_ocl_check) set(OCL_TEST_DEFAULT TRUE) option_support_check(${PARSE_OCL_ENABLED} ${PARSE_OCL_DISABLED} ${OCL_TEST_DEFAULT} is_ocl_check) @@ -585,6 +591,7 @@ function(add_custom_test NAME) AND is_full_check AND is_xnack_on_check AND is_mlir_check + AND is_composablekernel_check AND (is_half_check OR is_bfloat16_check OR is_int8_check OR is_float_check) AND (is_ocl_check AND is_hip_check AND is_hip_nogpu_check) ) @@ -765,14 +772,12 @@ add_custom_test(test_conv_igemm_mlir_xdlops_fwd SKIP_UNLESS_ALL HALF_ENABLED INT COMMAND ${IMPLICITGEMM_MLIR_ENV_F_XDLOPS} $ ${TEST_CONV_VERBOSE_F} --input 256 256 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --group-count 4 ) -if(MIOPEN_USE_COMPOSABLEKERNEL) -add_custom_test(test_conv_hip_igemm_xdlops SKIP_UNLESS_ALL OCL_DISABLED HALF_DISABLED FLOAT_DISABLED INT8_ENABLED GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED +add_custom_test(test_conv_hip_igemm_xdlops SKIP_UNLESS_ALL OCL_DISABLED HALF_DISABLED FLOAT_DISABLED INT8_ENABLED GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED SKIP_UNLESS_COMPOSABLEKERNEL COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 256 128 28 28 --weights 128 128 3 3 --output_type int8 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1 COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 128 512 7 7 --weights 512 512 3 3 --output_type int8 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1 COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 128 64 56 56 --weights 64 64 1 1 --output_type int8 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1 COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 256 256 56 56 --weights 256 64 1 1 --output_type int8 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1 ) -endif() add_custom_test(test_conv_igemm_mlir_xdlops_bwd_wrw SKIP_UNLESS_ALL HALF_ENABLED SKIP_UNLESS_MLIR GFX900_DISABLED GFX906_DISABLED COMMAND ${IMPLICITGEMM_MLIR_ENV_B_XDLOPS} $ ${TEST_CONV_VERBOSE_B} --input 256 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1 From 2153fa620c3af6212e43c2f2214ed44890a1b673 Mon Sep 17 00:00:00 2001 From: atamazov Date: Thu, 15 Dec 2022 00:39:33 +0300 Subject: [PATCH 44/66] ata-test-solver-03(23) [tests] Added smoke test for ConvHipImplicitGemmFwdXdlops (with tuning) --- test/CMakeLists.txt | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 09571acdee..587bf6c52f 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2106,6 +2106,15 @@ add_custom_test(smoke_solver_ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm GF ${TEST_CONV_VERBOSE_F} --input 16 1 7 7 --weights 1 1 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +# Note: "--output_type" is used only with "--int8", ignored otherwise. +add_custom_test(smoke_solver_ConvHipImplicitGemmFwdXdlops OCL_DISABLED GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED HALF_ENABLED INT8_ENABLED + SKIP_UNLESS_COMPOSABLEKERNEL TEST_TUNING + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmFwdXdlops $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 + --output_type int8 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinograd3x3U $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} From a8eca64d0f05d960964e2da95b0e6f2d99605075 Mon Sep 17 00:00:00 2001 From: atamazov Date: Thu, 15 Dec 2022 00:53:42 +0300 Subject: [PATCH 45/66] ata-test-solver-03(24) [CMake][Fix] Fix issue in set_var_to_condition, see https://github.com/ROCmSoftwarePlatform/MIOpen/pull/1253#pullrequestreview-1218296540 --- CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index df5754d9eb..097e777c7d 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -29,7 +29,7 @@ macro(set_var_to_condition var) if(${ARGN}) set(${var} TRUE) else() - set(${name} FALSE) + set(${var} FALSE) endif() endmacro() From c12d0d5eb39a31046926368930ed2640a683927b Mon Sep 17 00:00:00 2001 From: atamazov Date: Thu, 15 Dec 2022 01:06:20 +0300 Subject: [PATCH 46/66] ata-test-solver-03(25) [tests] Added support for the "--output_type" option of test_conv. --- test/CMakeLists.txt | 19 ++++++++++++------- 1 file changed, 12 insertions(+), 7 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 587bf6c52f..d84619b32a 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -155,6 +155,9 @@ endif() set(MIOPEN_TEST_FLOAT_ARG) set(MIOPEN_TEST_FLOAT FALSE) +set(MIOPEN_TEST_CONV_INT8_OUT_TYPE_INT8 "") +set(MIOPEN_TEST_CONV_INT8_OUT_TYPE_INT32 "") +set(MIOPEN_TEST_CONV_INT8_OUT_TYPE_FLOAT "") if(MIOPEN_TEST_HALF) set(MIOPEN_TEST_FLOAT_ARG --half) set(MIOPENDRIVER_MODE_CONV convfp16) @@ -163,6 +166,9 @@ elseif(MIOPEN_TEST_INT8) set(MIOPEN_TEST_FLOAT_ARG --int8) set(MIOPENDRIVER_MODE_CONV convint8) set(MIOPENDRIVER_MODE_BN NOT_SUPPORTED) + set(MIOPEN_TEST_CONV_INT8_OUT_TYPE_INT8 "--output_type int8") + set(MIOPEN_TEST_CONV_INT8_OUT_TYPE_INT32 "--output_type int32") + set(MIOPEN_TEST_CONV_INT8_OUT_TYPE_FLOAT "--output_type float") elseif(MIOPEN_TEST_BFLOAT16) set(MIOPEN_TEST_FLOAT_ARG --bfloat16) set(MIOPENDRIVER_MODE_CONV convbfp16) @@ -773,10 +779,10 @@ add_custom_test(test_conv_igemm_mlir_xdlops_fwd SKIP_UNLESS_ALL HALF_ENABLED INT ) add_custom_test(test_conv_hip_igemm_xdlops SKIP_UNLESS_ALL OCL_DISABLED HALF_DISABLED FLOAT_DISABLED INT8_ENABLED GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED SKIP_UNLESS_COMPOSABLEKERNEL - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 256 128 28 28 --weights 128 128 3 3 --output_type int8 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 128 512 7 7 --weights 512 512 3 3 --output_type int8 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 128 64 56 56 --weights 64 64 1 1 --output_type int8 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1 - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 256 256 56 56 --weights 256 64 1 1 --output_type int8 --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1 + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 256 128 28 28 --weights 128 128 3 3 ${MIOPEN_TEST_CONV_INT8_OUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1 + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 128 512 7 7 --weights 512 512 3 3 ${MIOPEN_TEST_CONV_INT8_OUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1 + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 128 64 56 56 --weights 64 64 1 1 ${MIOPEN_TEST_CONV_INT8_OUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1 + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 256 256 56 56 --weights 256 64 1 1 ${MIOPEN_TEST_CONV_INT8_OUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1 ) add_custom_test(test_conv_igemm_mlir_xdlops_bwd_wrw SKIP_UNLESS_ALL HALF_ENABLED SKIP_UNLESS_MLIR GFX900_DISABLED GFX906_DISABLED @@ -2106,13 +2112,12 @@ add_custom_test(smoke_solver_ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm GF ${TEST_CONV_VERBOSE_F} --input 16 1 7 7 --weights 1 1 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -# Note: "--output_type" is used only with "--int8", ignored otherwise. add_custom_test(smoke_solver_ConvHipImplicitGemmFwdXdlops OCL_DISABLED GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED HALF_ENABLED INT8_ENABLED SKIP_UNLESS_COMPOSABLEKERNEL TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmFwdXdlops $ ${MIOPEN_TEST_FLOAT_ARG} - ${TEST_CONV_VERBOSE_F} --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 - --output_type int8 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} + ${TEST_CONV_VERBOSE_F} --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_CONV_INT8_OUT_TYPE_INT8} + --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON From a136e045def455b4bcf0057d1ce24be8c89f6559 Mon Sep 17 00:00:00 2001 From: atamazov Date: Thu, 15 Dec 2022 02:16:57 +0300 Subject: [PATCH 47/66] ata-test-solver-03(26) [tests][NFC] Rename variable for clarity --- test/CMakeLists.txt | 22 +++++++++++----------- 1 file changed, 11 insertions(+), 11 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index d84619b32a..9e30b8d014 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -155,9 +155,9 @@ endif() set(MIOPEN_TEST_FLOAT_ARG) set(MIOPEN_TEST_FLOAT FALSE) -set(MIOPEN_TEST_CONV_INT8_OUT_TYPE_INT8 "") -set(MIOPEN_TEST_CONV_INT8_OUT_TYPE_INT32 "") -set(MIOPEN_TEST_CONV_INT8_OUT_TYPE_FLOAT "") +set(MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT8 "") +set(MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT32 "") +set(MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_FLOAT "") if(MIOPEN_TEST_HALF) set(MIOPEN_TEST_FLOAT_ARG --half) set(MIOPENDRIVER_MODE_CONV convfp16) @@ -166,9 +166,9 @@ elseif(MIOPEN_TEST_INT8) set(MIOPEN_TEST_FLOAT_ARG --int8) set(MIOPENDRIVER_MODE_CONV convint8) set(MIOPENDRIVER_MODE_BN NOT_SUPPORTED) - set(MIOPEN_TEST_CONV_INT8_OUT_TYPE_INT8 "--output_type int8") - set(MIOPEN_TEST_CONV_INT8_OUT_TYPE_INT32 "--output_type int32") - set(MIOPEN_TEST_CONV_INT8_OUT_TYPE_FLOAT "--output_type float") + set(MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT8 "--output_type int8") + set(MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT32 "--output_type int32") + set(MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_FLOAT "--output_type float") elseif(MIOPEN_TEST_BFLOAT16) set(MIOPEN_TEST_FLOAT_ARG --bfloat16) set(MIOPENDRIVER_MODE_CONV convbfp16) @@ -779,10 +779,10 @@ add_custom_test(test_conv_igemm_mlir_xdlops_fwd SKIP_UNLESS_ALL HALF_ENABLED INT ) add_custom_test(test_conv_hip_igemm_xdlops SKIP_UNLESS_ALL OCL_DISABLED HALF_DISABLED FLOAT_DISABLED INT8_ENABLED GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED SKIP_UNLESS_COMPOSABLEKERNEL - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 256 128 28 28 --weights 128 128 3 3 ${MIOPEN_TEST_CONV_INT8_OUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 128 512 7 7 --weights 512 512 3 3 ${MIOPEN_TEST_CONV_INT8_OUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 128 64 56 56 --weights 64 64 1 1 ${MIOPEN_TEST_CONV_INT8_OUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1 - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 256 256 56 56 --weights 256 64 1 1 ${MIOPEN_TEST_CONV_INT8_OUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1 + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 256 128 28 28 --weights 128 128 3 3 ${MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1 + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 128 512 7 7 --weights 512 512 3 3 ${MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 1 1 1 1 1 1 + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 128 64 56 56 --weights 64 64 1 1 ${MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1 + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --disable-backward-data --disable-backward-weights --verbose --input 256 256 56 56 --weights 256 64 1 1 ${MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC --pads_strides_dilations 0 0 1 1 1 1 ) add_custom_test(test_conv_igemm_mlir_xdlops_bwd_wrw SKIP_UNLESS_ALL HALF_ENABLED SKIP_UNLESS_MLIR GFX900_DISABLED GFX906_DISABLED @@ -2116,7 +2116,7 @@ add_custom_test(smoke_solver_ConvHipImplicitGemmFwdXdlops OCL_DISABLED GFX900_DI SKIP_UNLESS_COMPOSABLEKERNEL TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmFwdXdlops $ ${MIOPEN_TEST_FLOAT_ARG} - ${TEST_CONV_VERBOSE_F} --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_CONV_INT8_OUT_TYPE_INT8} + ${TEST_CONV_VERBOSE_F} --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} ) From e6ddb11f429fe0956b8dfb64e73e9df8f1c192e4 Mon Sep 17 00:00:00 2001 From: atamazov Date: Thu, 15 Dec 2022 20:16:50 +0300 Subject: [PATCH 48/66] ata-test-solver-03(27) [tests] Added smoke test for ConvHipImplicitGemmForwardV4R5Xdlops (with tuning) --- test/CMakeLists.txt | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 9e30b8d014..2a9d52c89b 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2112,6 +2112,13 @@ add_custom_test(smoke_solver_ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm GF ${TEST_CONV_VERBOSE_F} --input 16 1 7 7 --weights 1 1 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_solver_ConvHipImplicitGemmForwardV4R5Xdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmForwardV4R5Xdlops $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_F} --input 128 16 54 54 --weights 64 16 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_solver_ConvHipImplicitGemmFwdXdlops OCL_DISABLED GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED HALF_ENABLED INT8_ENABLED SKIP_UNLESS_COMPOSABLEKERNEL TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 From 6badec2663a6748c2fcfc8ec0975c740f11f3762 Mon Sep 17 00:00:00 2001 From: atamazov Date: Thu, 15 Dec 2022 20:55:07 +0300 Subject: [PATCH 49/66] ata-test-solver-03(28) [tests] Added smoke test for ConvHipImplicitGemmWrwV4R4Xdlops (with tuning) --- test/CMakeLists.txt | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 2a9d52c89b..bff4d41740 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2105,6 +2105,13 @@ add_custom_test(smoke_solver_ConvHipImplicitGemmForwardV4R4Xdlops GFX900_DISABLE ${TEST_CONV_VERBOSE_F} --input 128 48 13 13 --weights 192 48 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_solver_ConvHipImplicitGemmWrwV4R4Xdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmWrwV4R4Xdlops $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_W} --input 1 192 28 28 --weights 16 192 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_solver_ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 From b80aa6529c96c431dd21ed5f8fce61e0840ec6d7 Mon Sep 17 00:00:00 2001 From: atamazov Date: Fri, 16 Dec 2022 00:48:38 +0300 Subject: [PATCH 50/66] ata-test-solver-03(29) [tests] Added smoke test for ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm (with tuning) --- test/CMakeLists.txt | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index bff4d41740..86ade15084 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2119,6 +2119,13 @@ add_custom_test(smoke_solver_ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm GF ${TEST_CONV_VERBOSE_F} --input 16 1 7 7 --weights 1 1 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_solver_ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm $ ${MIOPEN_TEST_FLOAT_ARG} + ${TEST_CONV_VERBOSE_W} --input 256 2 5 5 --weights 1 2 3 3 --pads_strides_dilations 1 1 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_solver_ConvHipImplicitGemmForwardV4R5Xdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 From 9a185e6493b8ba7c4692d932cfb629c1a330d254 Mon Sep 17 00:00:00 2001 From: atamazov Date: Sat, 17 Dec 2022 01:16:47 +0300 Subject: [PATCH 51/66] ata-test-solver-03(30) [tests] Remove unnecessary MIOPEN_TEST_FLOAT_ARG from test_conv commands where TEST_CONV_VERBOSE_F/B/W is already used. --- test/CMakeLists.txt | 114 ++++++++++++++++++++++---------------------- 1 file changed, 57 insertions(+), 57 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 86ade15084..21aa03a316 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1874,40 +1874,40 @@ if(MIOPEN_TEST_FLOAT) endif() add_custom_test(smoke_solver_ConvFFT GFX103X_ENABLED - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=fft $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=fft $ ${TEST_CONV_VERBOSE_F} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=fft $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=fft $ ${TEST_CONV_VERBOSE_B} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvDirectNaiveConv_F GFX103X_ENABLED HALF_ENABLED BF16_ENABLED INT8_ENABLED - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvFwd $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvFwd $ ${TEST_CONV_VERBOSE_F} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvDirectNaiveConv_BW GFX103X_ENABLED HALF_ENABLED BF16_ENABLED - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvBwd $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvBwd $ ${TEST_CONV_VERBOSE_B} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvWrw $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvWrw $ ${TEST_CONV_VERBOSE_W} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvAsm_5x10_7x7 GFX90A_DISABLED SKIP_XNACK_ON # GFX90A_DISABLED is because of WORKAROUND_ISSUE_1146 - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm5x10u2v2f1 $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm5x10u2v2f1 $ ${TEST_CONV_VERBOSE_F} --input 1 1 5 10 --weights 16 1 5 10 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm5x10u2v2b1 $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm5x10u2v2b1 $ ${TEST_CONV_VERBOSE_B} --input 1 1 16 160 --weights 16 16 5 10 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm7x7c3h224w224k64u2v2p3q3f1 $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm7x7c3h224w224k64u2v2p3q3f1 $ ${TEST_CONV_VERBOSE_F} --input 1 3 224 224 --weights 64 3 7 7 --pads_strides_dilations 3 3 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvOcl_SET01 GFX103X_ENABLED HALF_ENABLED BF16_ENABLED - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclDirectFwd11x11 $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclDirectFwd11x11 $ ${TEST_CONV_VERBOSE_F} --input 1 1 44 44 --weights 1 1 11 11 --pads_strides_dilations 0 0 4 4 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclDirectFwdGen $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclDirectFwdGen $ ${TEST_CONV_VERBOSE_F} --input 1 1 6 6 --weights 1 1 3 3 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclBwdWrW53 $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclBwdWrW53 $ ${TEST_CONV_VERBOSE_W} --input 16 1 7 7 --weights 1 1 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) @@ -1921,89 +1921,89 @@ add_custom_test(smoke_solver_ConvOcl_SET01 GFX103X_ENABLED HALF_ENABLED BF16_ENA add_custom_test(smoke_solver_ConvAsm1x1U HALF_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U $ ${TEST_CONV_VERBOSE_F} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U $ ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvAsm1x1UV2 SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1UV2 $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1UV2 $ ${TEST_CONV_VERBOSE_F} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1UV2 $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1UV2 $ ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvAsm3x3U SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm3x3U $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm3x3U $ ${TEST_CONV_VERBOSE_F} --input 1 4 10 10 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm3x3U $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm3x3U $ ${TEST_CONV_VERBOSE_B} --input 1 4 10 10 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvAsmBwdWrW1x1 HALF_ENABLED BF16_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmBwdWrW1x1 $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmBwdWrW1x1 $ ${TEST_CONV_VERBOSE_W} --input 1 4 5 5 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) # GFX90A_DISABLED for FP32 because of WORKAROUND_SWDEV_330460 add_custom_test(smoke_solver_ConvAsmBwdWrW3x3_fp32 GFX90A_DISABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmBwdWrW3x3 $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmBwdWrW3x3 $ ${TEST_CONV_VERBOSE_W} --input 2 4 3 3 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvAsmBwdWrW3x3_fp16 FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmBwdWrW3x3 $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmBwdWrW3x3 $ ${TEST_CONV_VERBOSE_W} --input 2 4 3 3 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) # GFX103X_DISABLED is due to WORKAROUND_SWDEV_266868 add_custom_test(smoke_solver_ConvOclBwdWrW1x1 GFX103X_DISABLED HALF_ENABLED BF16_ENABLED - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclBwdWrW1x1 $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclBwdWrW1x1 $ ${TEST_CONV_VERBOSE_W} --input 1 16 14 14 --weights 16 16 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvAsmImplicitGemmV4R1Dynamic GFX908_DISABLED GFX90A_DISABLED SKIP_XNACK_ON - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicFwd $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicFwd $ ${TEST_CONV_VERBOSE_F} --input 16 16 16 16 --weights 16 16 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicBwd $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicBwd $ ${TEST_CONV_VERBOSE_B} --input 64 64 14 14 --weights 16 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicWrw $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicWrw $ ${TEST_CONV_VERBOSE_W} --input 1 32 28 28 --weights 32 32 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED HALF_ENABLED SKIP_XNACK_ON - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicWrwXdlops $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicWrwXdlops $ ${TEST_CONV_VERBOSE_W} --input 2 256 12 18 --weights 256 256 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlops $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlops $ ${TEST_CONV_VERBOSE_B} --input 64 64 28 28 --weights 16 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlops $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlops $ ${TEST_CONV_VERBOSE_F} --input 64 512 7 7 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_fp32_fp16 GFX900_DISABLED GFX906_DISABLED HALF_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC $ ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC $ ${TEST_CONV_VERBOSE_B} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC $ ${TEST_CONV_VERBOSE_W} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} ) @@ -2011,15 +2011,15 @@ add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_fp32_fp16 G add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16 GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED FLOAT_DISABLED BF16_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC $ ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC $ ${TEST_CONV_VERBOSE_B} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC $ ${TEST_CONV_VERBOSE_W} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} ) @@ -2027,7 +2027,7 @@ add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16 GFX900 add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED GFX90A_DISABLED GFX103X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC $ ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 256 3 3 128 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NCHW --fil_layout CHWN --out_layout NCHW --tensor_vect 1 --vector_length 4 ${MIOPEN_TEST_FLAGS_ARGS} ) @@ -2036,24 +2036,24 @@ add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC GFX900_D add_custom_test(smoke_solver_ConvCkIgemmFwdV6r1DlopsNchw GFX103X_ENABLED HALF_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=2 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvCkIgemmFwdV6r1DlopsNchw $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvCkIgemmFwdV6r1DlopsNchw $ ${TEST_CONV_VERBOSE_F} --input 128 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvHipImplicitGemmBwdDataV1R1 GFX103X_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV1R1 $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV1R1 $ ${TEST_CONV_VERBOSE_B} --input 32 128 32 32 --weights 12 128 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvHipImplicitGemmV4R1 GFX103X_ENABLED HALF_ENABLED BF16_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R1Fwd $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R1Fwd $ ${TEST_CONV_VERBOSE_F} --input 256 32 27 27 --weights 128 32 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R1WrW $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R1WrW $ ${TEST_CONV_VERBOSE_W} --input 64 64 55 55 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) @@ -2062,16 +2062,16 @@ add_custom_test(smoke_solver_ConvHipImplicitGemmV4R1 GFX103X_ENABLED HALF_ENABLE add_custom_test(smoke_solver_ConvHipImplicitGemmBwdDataV4R1 GFX103X_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1=1 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV4R1 $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV4R1 $ ${TEST_CONV_VERBOSE_B} --input 16 64 16 16 --weights 64 64 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvHipImplicitGemmV4R4 GFX103X_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R4Fwd $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R4Fwd $ ${TEST_CONV_VERBOSE_F} --input 2 16 28 28 --weights 32 16 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R4WrW $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R4WrW $ ${TEST_CONV_VERBOSE_W} --input 8 128 14 14 --weights 32 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) @@ -2082,7 +2082,7 @@ add_custom_test(smoke_solver_ConvHipImplicitGemmBwdDataV1R1Xdlops GFX900_DISABLE COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1_XDLOPS=1 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV1R1Xdlops $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV1R1Xdlops $ ${TEST_CONV_VERBOSE_B} --input 32 128 32 32 --weights 12 128 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) @@ -2094,81 +2094,81 @@ add_custom_test(smoke_solver_ConvHipImplicitGemmBwdDataV4R1Xdlops GFX900_DISABLE COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1_XDLOPS=1 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV4R1Xdlops $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV4R1Xdlops $ ${TEST_CONV_VERBOSE_B} --input 1 160 28 28 --weights 128 160 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvHipImplicitGemmForwardV4R4Xdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmForwardV4R4Xdlops $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmForwardV4R4Xdlops $ ${TEST_CONV_VERBOSE_F} --input 128 48 13 13 --weights 192 48 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvHipImplicitGemmWrwV4R4Xdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmWrwV4R4Xdlops $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmWrwV4R4Xdlops $ ${TEST_CONV_VERBOSE_W} --input 1 192 28 28 --weights 16 192 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm $ ${TEST_CONV_VERBOSE_F} --input 16 1 7 7 --weights 1 1 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm $ ${TEST_CONV_VERBOSE_W} --input 256 2 5 5 --weights 1 2 3 3 --pads_strides_dilations 1 1 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvHipImplicitGemmForwardV4R5Xdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmForwardV4R5Xdlops $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmForwardV4R5Xdlops $ ${TEST_CONV_VERBOSE_F} --input 128 16 54 54 --weights 64 16 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvHipImplicitGemmFwdXdlops OCL_DISABLED GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED HALF_ENABLED INT8_ENABLED SKIP_UNLESS_COMPOSABLEKERNEL TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmFwdXdlops $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmFwdXdlops $ ${TEST_CONV_VERBOSE_F} --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinograd3x3U $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinograd3x3U $ ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinograd3x3U $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinograd3x3U $ ${TEST_CONV_VERBOSE_B} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) # F16 is supported for 906 and 906 only, no WrW add_custom_test(smoke_solver_ConvBinWinogradRxS_f16 GFX900_DISABLED GFX90A_DISABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${TEST_CONV_VERBOSE_F} --input 1 40 20 20 --weights 20 40 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${TEST_CONV_VERBOSE_B} --input 1 20 20 20 --weights 40 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) # F32 is supported for 900, 906 and 908. add_custom_test(smoke_solver_ConvBinWinogradRxS_f32 GFX90A_DISABLED SKIP_XNACK_ON - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${TEST_CONV_VERBOSE_B} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${TEST_CONV_VERBOSE_W} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) # GFX90A_DISABLED is due to WORKAROUND_ISSUE_1146 add_custom_test(smoke_solver_ConvWinograd3x3MultipassWrW GFX90A_DISABLED HALF_ENABLED BF16_ENABLED SKIP_XNACK_ON OCL_DISABLED - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER='ConvWinograd3x3MultipassWrW<3-3>' $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER='ConvWinograd3x3MultipassWrW<3-3>' $ ${TEST_CONV_VERBOSE_W} --input 1 64 30 30 --weights 64 64 3 3 --pads_strides_dilations 1 1 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) From b3a49179ccf3d0e64e6b7656665d0288c7e91494 Mon Sep 17 00:00:00 2001 From: atamazov Date: Sat, 17 Dec 2022 02:05:23 +0300 Subject: [PATCH 52/66] ata-test-solver-03(31) [tests] Added smoke test for ConvMlirIgemmFwd/Bwd/WrW (with tuning). Removed test_conv_igemm_mlir_small. --- test/CMakeLists.txt | 22 ++++++++++++---------- 1 file changed, 12 insertions(+), 10 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 21aa03a316..7984dfcc86 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -712,16 +712,6 @@ set(TEST_CONV_VERBOSE_F ${MIOPEN_TEST_FLOAT_ARG} --verbose --disable-backward-da set(TEST_CONV_VERBOSE_B ${MIOPEN_TEST_FLOAT_ARG} --verbose --disable-forward --disable-backward-weights) set(TEST_CONV_VERBOSE_W ${MIOPEN_TEST_FLOAT_ARG} --verbose --disable-forward --disable-backward-data) -# Note: OpenCL Debug + Codecov test stage taking longer time than a smoke test should therefore disabling that scenario -string(TOUPPER ${CMAKE_BUILD_TYPE} CMAKE_BUILD_TYPE) -if ((NOT CMAKE_BUILD_TYPE MATCHES "DEBUG") OR (NOT ${CODECOV_TEST})) - add_custom_test(test_conv_igemm_mlir_small HALF_ENABLED SKIP_UNLESS_MLIR GFX900_DISABLED GFX908_DISABLED GFX90A_DISABLED - COMMAND ${IMPLICITGEMM_MLIR_ENV_F} $ ${TEST_CONV_VERBOSE_F} --input 64 128 14 14 --weights 128 128 1 1 --pads_strides_dilations 0 0 2 2 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND ${IMPLICITGEMM_MLIR_ENV_B} $ ${TEST_CONV_VERBOSE_B} --input 64 256 28 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --group-count 4 - COMMAND ${IMPLICITGEMM_MLIR_ENV_W} $ ${TEST_CONV_VERBOSE_W} --input 64 64 28 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 - ) -endif() - add_custom_test(test_pooling2d_asymmetric SKIP_UNLESS_ALL HALF_ENABLED GFX103X_ENABLED COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --all --dataset 1 --limit 0 ${MIOPEN_TEST_FLAGS_ARGS} ) @@ -2172,6 +2162,18 @@ add_custom_test(smoke_solver_ConvWinograd3x3MultipassWrW GFX90A_DISABLED HALF_EN ${TEST_CONV_VERBOSE_W} --input 1 64 30 30 --weights 64 64 3 3 --pads_strides_dilations 1 1 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_solver_ConvMlirIgemm GFX900_DISABLED GFX908_DISABLED GFX90A_DISABLED HALF_ENABLED SKIP_UNLESS_MLIR TEST_TUNING + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + ${IMPLICITGEMM_MLIR_ENV_F} $ ${TEST_CONV_VERBOSE_F} + --input 64 128 14 14 --weights 128 128 1 1 --pads_strides_dilations 0 0 2 2 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + ${IMPLICITGEMM_MLIR_ENV_B} $ ${TEST_CONV_VERBOSE_B} + --input 64 256 28 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --group-count 4 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + ${IMPLICITGEMM_MLIR_ENV_W} $ ${TEST_CONV_VERBOSE_W} + --input 64 64 28 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + # Add here regression tests that should be run on Vega10/20 and GFX908 only with FP16. add_custom_test(test_regression_half_vega_gfx908 FLOAT_DISABLED HALF_ENABLED GFX90A_DISABLED # REGRESSION TEST for issue #894. From 826b9e2de7617f0eaf019c6b6bd078e18f94be1d Mon Sep 17 00:00:00 2001 From: atamazov Date: Sat, 17 Dec 2022 19:29:55 +0300 Subject: [PATCH 53/66] ata-test-solver-03(32) [tests] Added smoke test for ConvMlirIgemmFwd/Bwd/WrWXdlops(with tuning). Removed test_conv_igemm_mlir_xdlops_small. --- test/CMakeLists.txt | 18 ++++++++++++------ 1 file changed, 12 insertions(+), 6 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 7984dfcc86..ceb890f26b 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -751,12 +751,6 @@ set(IMPLICITGEMM_MLIR_ENV_F_XDLOPS ${IMPLICITGEMM_MLIR_ENV_BASE} MIOPEN_DEBUG_FI set(IMPLICITGEMM_MLIR_ENV_B_XDLOPS ${IMPLICITGEMM_MLIR_ENV_BASE} MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvMlirIgemmBwdXdlops) set(IMPLICITGEMM_MLIR_ENV_W_XDLOPS ${IMPLICITGEMM_MLIR_ENV_BASE} MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvMlirIgemmWrWXdlops) -add_custom_test(test_conv_igemm_mlir_xdlops_small HALF_ENABLED SKIP_UNLESS_MLIR GFX900_DISABLED GFX906_DISABLED - COMMAND ${IMPLICITGEMM_MLIR_ENV_F_XDLOPS} $ ${TEST_CONV_VERBOSE_F} --input 64 128 14 14 --weights 128 128 1 1 --pads_strides_dilations 0 0 2 2 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND ${IMPLICITGEMM_MLIR_ENV_B_XDLOPS} $ ${TEST_CONV_VERBOSE_B} --input 64 256 28 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --group-count 4 - COMMAND ${IMPLICITGEMM_MLIR_ENV_W_XDLOPS} $ ${TEST_CONV_VERBOSE_W} --input 64 64 28 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 -) - add_custom_test(test_conv_igemm_mlir_xdlops_fwd SKIP_UNLESS_ALL HALF_ENABLED INT8_ENABLED SKIP_UNLESS_MLIR GFX900_DISABLED GFX906_DISABLED COMMAND ${IMPLICITGEMM_MLIR_ENV_F_XDLOPS} $ ${TEST_CONV_VERBOSE_F} --input 256 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1 COMMAND ${IMPLICITGEMM_MLIR_ENV_F_XDLOPS} $ ${TEST_CONV_VERBOSE_F} --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 @@ -2174,6 +2168,18 @@ add_custom_test(smoke_solver_ConvMlirIgemm GFX900_DISABLED GFX908_DISABLED GFX90 --input 64 64 28 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_solver_ConvMlirIgemmXdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED SKIP_UNLESS_MLIR TEST_TUNING + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + ${IMPLICITGEMM_MLIR_ENV_F_XDLOPS} $ ${TEST_CONV_VERBOSE_F} + --input 64 128 14 14 --weights 128 128 1 1 --pads_strides_dilations 0 0 2 2 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + ${IMPLICITGEMM_MLIR_ENV_B_XDLOPS} $ ${TEST_CONV_VERBOSE_B} + --input 64 256 28 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --group-count 4 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + ${IMPLICITGEMM_MLIR_ENV_W_XDLOPS} $ ${TEST_CONV_VERBOSE_W} + --input 64 64 28 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + # Add here regression tests that should be run on Vega10/20 and GFX908 only with FP16. add_custom_test(test_regression_half_vega_gfx908 FLOAT_DISABLED HALF_ENABLED GFX90A_DISABLED # REGRESSION TEST for issue #894. From 3e45f91b8fb6961b519737378f7db902ed406998 Mon Sep 17 00:00:00 2001 From: atamazov Date: Sun, 18 Dec 2022 02:04:14 +0300 Subject: [PATCH 54/66] ata-test-solver-03(33) [tests] Added smoke test for ConvBinWinogradRxSf2x3 (with tuning) --- test/CMakeLists.txt | 11 +++++++++++ 1 file changed, 11 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index ceb890f26b..c0f333c080 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2156,6 +2156,17 @@ add_custom_test(smoke_solver_ConvWinograd3x3MultipassWrW GFX90A_DISABLED HALF_EN ${TEST_CONV_VERBOSE_W} --input 1 64 30 30 --weights 64 64 3 3 --pads_strides_dilations 1 1 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3 GFX900_DISABLED GFX103X_ENABLED HALF_ENABLED SKIP_XNACK_ON + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=`ConvBinWinogradRxSf2x3` $ + ${TEST_CONV_VERBOSE_F} --input 1 40 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 2 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=`ConvBinWinogradRxSf2x3` $ + ${TEST_CONV_VERBOSE_B} --input 1 40 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 2 ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_solver_ConvMlirIgemm GFX900_DISABLED GFX908_DISABLED GFX90A_DISABLED HALF_ENABLED SKIP_UNLESS_MLIR TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 ${IMPLICITGEMM_MLIR_ENV_F} $ ${TEST_CONV_VERBOSE_F} From ad70061637a628216835924568a45ae2cc27cee0 Mon Sep 17 00:00:00 2001 From: atamazov Date: Sun, 18 Dec 2022 02:11:00 +0300 Subject: [PATCH 55/66] ata-test-solver-03(34) [tests] Rename couple of tests for clarity. Small fixes. --- test/CMakeLists.txt | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index c0f333c080..b6764bc984 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2133,7 +2133,7 @@ add_custom_test(smoke_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON ) # F16 is supported for 906 and 906 only, no WrW -add_custom_test(smoke_solver_ConvBinWinogradRxS_f16 GFX900_DISABLED GFX90A_DISABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvBinWinogradRxS_fp16 GFX900_DISABLED GFX90A_DISABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${TEST_CONV_VERBOSE_F} --input 1 40 20 20 --weights 20 40 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ @@ -2141,7 +2141,7 @@ add_custom_test(smoke_solver_ConvBinWinogradRxS_f16 GFX900_DISABLED GFX90A_DISAB ) # F32 is supported for 900, 906 and 908. -add_custom_test(smoke_solver_ConvBinWinogradRxS_f32 GFX90A_DISABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvBinWinogradRxS_fp32 GFX90A_DISABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS $ @@ -2159,11 +2159,11 @@ add_custom_test(smoke_solver_ConvWinograd3x3MultipassWrW GFX90A_DISABLED HALF_EN add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3 GFX900_DISABLED GFX103X_ENABLED HALF_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=`ConvBinWinogradRxSf2x3` $ + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3 $ ${TEST_CONV_VERBOSE_F} --input 1 40 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 2 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=`ConvBinWinogradRxSf2x3` $ + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3 $ ${TEST_CONV_VERBOSE_B} --input 1 40 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 2 ${MIOPEN_TEST_FLAGS_ARGS} ) From 8834f72733bb0b646460513384d764060cf7e590 Mon Sep 17 00:00:00 2001 From: atamazov Date: Sun, 18 Dec 2022 02:15:42 +0300 Subject: [PATCH 56/66] ata-test-solver-03(35) [tests] Added smoke test for ConvBinWinogradRxSf2x3, WrW (with tuning) --- test/CMakeLists.txt | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index b6764bc984..8fdf32864a 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2156,7 +2156,7 @@ add_custom_test(smoke_solver_ConvWinograd3x3MultipassWrW GFX90A_DISABLED HALF_EN ${TEST_CONV_VERBOSE_W} --input 1 64 30 30 --weights 64 64 3 3 --pads_strides_dilations 1 1 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3 GFX900_DISABLED GFX103X_ENABLED HALF_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3 GFX900_DISABLED GFX103X_ENABLED HALF_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3 $ @@ -2165,6 +2165,10 @@ add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3 GFX900_DISABLED GFX103X_ENAB MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3 $ ${TEST_CONV_VERBOSE_B} --input 1 40 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 2 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3 $ + ${TEST_CONV_VERBOSE_W} --input 1 40 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 2 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvMlirIgemm GFX900_DISABLED GFX908_DISABLED GFX90A_DISABLED HALF_ENABLED SKIP_UNLESS_MLIR TEST_TUNING From 0665cc869bfdd533426ef43a71050fe85b073c6f Mon Sep 17 00:00:00 2001 From: atamazov Date: Sun, 18 Dec 2022 02:37:45 +0300 Subject: [PATCH 57/66] ata-test-solver-03(36) [tests] Added smoke test for ConvBinWinogradRxSf2x3g1 --- test/CMakeLists.txt | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 8fdf32864a..ec1134e23c 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2171,6 +2171,12 @@ add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3 GFX900_DISABLED GFX103X_ENAB ${TEST_CONV_VERBOSE_W} --input 1 40 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 2 ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3g1 GFX900_DISABLED GFX103X_ENABLED HALF_ENABLED SKIP_XNACK_ON + COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ + --input 1 40 20 20 --weights 20 40 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_solver_ConvMlirIgemm GFX900_DISABLED GFX908_DISABLED GFX90A_DISABLED HALF_ENABLED SKIP_UNLESS_MLIR TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 ${IMPLICITGEMM_MLIR_ENV_F} $ ${TEST_CONV_VERBOSE_F} From 6cdc8f36359dc0b90ae449544eb8e2e5ac2867a1 Mon Sep 17 00:00:00 2001 From: atamazov Date: Mon, 19 Dec 2022 01:48:23 +0300 Subject: [PATCH 58/66] ata-test-solver-03(37) [tests] Added smoke test for ConvBinWinogradRxSf3x2 (with tuning) --- test/CMakeLists.txt | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index ec1134e23c..aee1b45195 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2177,6 +2177,13 @@ add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3g1 GFX900_DISABLED GFX103X_EN --input 1 40 20 20 --weights 20 40 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +add_custom_test(smoke_solver_ConvBinWinogradRxSf3x2 GFX900_DISABLED GFX103X_ENABLED HALF_ENABLED SKIP_XNACK_ON + COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ + --input 1 40 20 20 --weights 20 40 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +) + add_custom_test(smoke_solver_ConvMlirIgemm GFX900_DISABLED GFX908_DISABLED GFX90A_DISABLED HALF_ENABLED SKIP_UNLESS_MLIR TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 ${IMPLICITGEMM_MLIR_ENV_F} $ ${TEST_CONV_VERBOSE_F} From bf0172f46a4b3484a505f413b6485371949c4a77 Mon Sep 17 00:00:00 2001 From: atamazov Date: Mon, 19 Dec 2022 02:01:52 +0300 Subject: [PATCH 59/66] ata-test-solver-03(38) [tests] Rework test for ConvBinWinogradRxSf2x3 --- test/CMakeLists.txt | 10 +--------- 1 file changed, 1 insertion(+), 9 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index aee1b45195..b70c6333c9 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2160,15 +2160,7 @@ add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3 GFX900_DISABLED GFX103X_ENAB COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3 $ - ${TEST_CONV_VERBOSE_F} --input 1 40 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 2 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 - MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3 $ - ${TEST_CONV_VERBOSE_B} --input 1 40 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 2 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 - MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3 $ - ${TEST_CONV_VERBOSE_W} --input 1 40 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 2 ${MIOPEN_TEST_FLAGS_ARGS} + --input 1 40 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 2 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3g1 GFX900_DISABLED GFX103X_ENABLED HALF_ENABLED SKIP_XNACK_ON From 3d4399e9840e87f2c30b01de40db75355fcd84c0 Mon Sep 17 00:00:00 2001 From: atamazov Date: Sat, 24 Dec 2022 01:44:06 +0300 Subject: [PATCH 60/66] ata-test-solver-03(39) Fix bug in commit ata-test-solver-03(25) --- test/CMakeLists.txt | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index b70c6333c9..86c1920b3e 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -155,9 +155,9 @@ endif() set(MIOPEN_TEST_FLOAT_ARG) set(MIOPEN_TEST_FLOAT FALSE) -set(MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT8 "") -set(MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT32 "") -set(MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_FLOAT "") +set(MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT8) +set(MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT32) +set(MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_FLOAT) if(MIOPEN_TEST_HALF) set(MIOPEN_TEST_FLOAT_ARG --half) set(MIOPENDRIVER_MODE_CONV convfp16) @@ -166,9 +166,9 @@ elseif(MIOPEN_TEST_INT8) set(MIOPEN_TEST_FLOAT_ARG --int8) set(MIOPENDRIVER_MODE_CONV convint8) set(MIOPENDRIVER_MODE_BN NOT_SUPPORTED) - set(MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT8 "--output_type int8") - set(MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT32 "--output_type int32") - set(MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_FLOAT "--output_type float") + set(MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT8 --output_type int8) + set(MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT32 --output_type int32) + set(MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_FLOAT --output_type float) elseif(MIOPEN_TEST_BFLOAT16) set(MIOPEN_TEST_FLOAT_ARG --bfloat16) set(MIOPENDRIVER_MODE_CONV convbfp16) From 8def845b69664fa33df427d2039a5021cc512467 Mon Sep 17 00:00:00 2001 From: atamazov Date: Sat, 24 Dec 2022 02:07:41 +0300 Subject: [PATCH 61/66] ata-test-solver-03(40) Formatting --- src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp b/src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp index e1d85ed9f2..522f881206 100644 --- a/src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp +++ b/src/solver/conv_hip_implicit_gemm_bwd_v1r1.cpp @@ -651,8 +651,8 @@ bool ConvHipImplicitGemmBwdDataV1R1::IsApplicable(const ConvolutionContext& ctx, #if WORKAROUND_ISSUE_309 if(problem.IsBfp16()) - if(!miopen::IsEnabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1{})) - return false; + if(!miopen::IsEnabled(MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1{})) + return false; #endif const auto k = ProblemInterpreter::GetOutputChannelK(problem); From c6da4382357d1127be9f99c18045068d470d4026 Mon Sep 17 00:00:00 2001 From: atamazov Date: Fri, 13 Jan 2023 17:12:56 +0300 Subject: [PATCH 62/66] ata-test-solver-02(24) Added GFX110X_ENABLED where necessary --- test/CMakeLists.txt | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index c092c59310..4f9d253719 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1886,19 +1886,19 @@ if(MIOPEN_TEST_FLOAT) add_custom_test(test_reduce_double SKIP_UNLESS_ALL GFX103X_ENABLED GFX110X_ENABLED COMMAND $ --double --all --verbose) endif() -add_custom_test(smoke_conv_solver_fft GFX103X_ENABLED +add_custom_test(smoke_conv_solver_fft GFX103X_ENABLED GFX110X_ENABLED COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=fft $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=fft $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_B} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvDirectNaiveConv_F GFX103X_ENABLED HALF_ENABLED BF16_ENABLED INT8_ENABLED +add_custom_test(smoke_conv_solver_ConvDirectNaiveConv_F GFX103X_ENABLED GFX110X_ENABLED HALF_ENABLED BF16_ENABLED INT8_ENABLED COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvFwd $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvDirectNaiveConv_BW GFX103X_ENABLED HALF_ENABLED BF16_ENABLED +add_custom_test(smoke_conv_solver_ConvDirectNaiveConv_BW GFX103X_ENABLED GFX110X_ENABLED HALF_ENABLED BF16_ENABLED COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvBwd $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_B} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvWrw $ ${MIOPEN_TEST_FLOAT_ARG} From cee52553f81989776cff3b8cae647601b381417c Mon Sep 17 00:00:00 2001 From: atamazov Date: Mon, 16 Jan 2023 17:40:43 +0300 Subject: [PATCH 63/66] ata-test-solver-02(25) Fixed cppcheck issue --- src/solver/conv_asm_1x1u_stride2.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/src/solver/conv_asm_1x1u_stride2.cpp b/src/solver/conv_asm_1x1u_stride2.cpp index 09c430eb44..1db17e1a34 100644 --- a/src/solver/conv_asm_1x1u_stride2.cpp +++ b/src/solver/conv_asm_1x1u_stride2.cpp @@ -101,7 +101,6 @@ inline static bool IsLinear(const int v) static inline size_t divide_round_plus_inf(const size_t x, const unsigned y) { - assert(x >= 0 && y > 0); if(x % y != 0) return x / y + 1; return x / y; From 9b30926a6e81880fdb62fdca26ba99f49451d506 Mon Sep 17 00:00:00 2001 From: atamazov Date: Mon, 20 Feb 2023 00:38:45 +0300 Subject: [PATCH 64/66] ata-test-solver-03(49) [tests][NFC] more smoke_conv_solver_Conv* -> smoke_solver_Conv* --- test/CMakeLists.txt | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 1a5a68f4b2..a5a0fc8cce 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2188,7 +2188,7 @@ add_custom_test(smoke_solver_ConvBinWinogradRxS_fp32 GFX90A_DISABLED SKIP_XNACK_ ) # FP16 ALT attribute is disabled to enable the backward solver on MI200 for HALF. -add_custom_test(smoke_conv_solver_ConvBinWinogradRxSf2x3g1_f16 GFX103X_ENABLED GFX110X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3g1_f16 GFX103X_ENABLED GFX110X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 40 20 20 --weights 20 40 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} @@ -2200,7 +2200,7 @@ add_custom_test(smoke_conv_solver_ConvBinWinogradRxSf2x3g1_f16 GFX103X_ENABLED G ${TEST_CONV_VERBOSE_W} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvBinWinogradRxSf2x3g1_f32 GFX103X_ENABLED GFX110X_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3g1_f32 GFX103X_ENABLED GFX110X_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ ${MIOPEN_TEST_FLOAT_ARG} @@ -2210,7 +2210,7 @@ add_custom_test(smoke_conv_solver_ConvBinWinogradRxSf2x3g1_f32 GFX103X_ENABLED G ) # FP16 ALT attribute is disabled to enable the backward solver on MI200 for HALF. -add_custom_test(smoke_conv_solver_ConvBinWinogradRxSf3x2_f16 GFX103X_ENABLED GFX110X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvBinWinogradRxSf3x2_f16 GFX103X_ENABLED GFX110X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 40 20 20 --weights 20 40 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} @@ -2222,7 +2222,7 @@ add_custom_test(smoke_conv_solver_ConvBinWinogradRxSf3x2_f16 GFX103X_ENABLED GFX ${TEST_CONV_VERBOSE_W} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_conv_solver_ConvBinWinogradRxSf3x2_f32 GFX103X_ENABLED GFX110X_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvBinWinogradRxSf3x2_f32 GFX103X_ENABLED GFX110X_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ ${MIOPEN_TEST_FLOAT_ARG} ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ ${MIOPEN_TEST_FLOAT_ARG} From 775b2adbacde26bfe98ab88136b3ce9e4a3c608b Mon Sep 17 00:00:00 2001 From: atamazov Date: Mon, 20 Feb 2023 00:41:59 +0300 Subject: [PATCH 65/66] ata-test-solver-03(50) [tests] More removals of unnecessary MIOPEN_TEST_FLOAT_ARG where TEST_CONV_VERBOSE_F/B/W is already used. --- test/CMakeLists.txt | 24 ++++++++++++------------ 1 file changed, 12 insertions(+), 12 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index a5a0fc8cce..24352b5aad 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2190,44 +2190,44 @@ add_custom_test(smoke_solver_ConvBinWinogradRxS_fp32 GFX90A_DISABLED SKIP_XNACK_ # FP16 ALT attribute is disabled to enable the backward solver on MI200 for HALF. add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3g1_f16 GFX103X_ENABLED GFX110X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ ${TEST_CONV_VERBOSE_F} --input 1 40 20 20 --weights 20 40 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ ${TEST_CONV_VERBOSE_B} --input 1 20 20 20 --weights 40 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ ${TEST_CONV_VERBOSE_W} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3g1_f32 GFX103X_ENABLED GFX110X_ENABLED SKIP_XNACK_ON - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ ${TEST_CONV_VERBOSE_B} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ ${TEST_CONV_VERBOSE_W} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) # FP16 ALT attribute is disabled to enable the backward solver on MI200 for HALF. add_custom_test(smoke_solver_ConvBinWinogradRxSf3x2_f16 GFX103X_ENABLED GFX110X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ ${TEST_CONV_VERBOSE_F} --input 1 40 20 20 --weights 20 40 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ ${TEST_CONV_VERBOSE_B} --input 1 20 20 20 --weights 40 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 - MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ ${MIOPEN_TEST_FLOAT_ARG} + MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ ${TEST_CONV_VERBOSE_W} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) add_custom_test(smoke_solver_ConvBinWinogradRxSf3x2_f32 GFX103X_ENABLED GFX110X_ENABLED SKIP_XNACK_ON - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ ${TEST_CONV_VERBOSE_B} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ ${MIOPEN_TEST_FLOAT_ARG} + COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ ${TEST_CONV_VERBOSE_W} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) From eab21d1396f4432a652ee7035684a9de8cef4faf Mon Sep 17 00:00:00 2001 From: atamazov Date: Fri, 3 Mar 2023 20:04:29 +0300 Subject: [PATCH 66/66] ata-test-solver-03(52) Explicitly enable ConvHipImplicitGemmV4R1Fwd to override WORKAROUND_iGemm_936 from Jenkinsfile. --- test/CMakeLists.txt | 3 +++ 1 file changed, 3 insertions(+) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 028511cd2f..3cec8e6594 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -2067,8 +2067,11 @@ add_custom_test(smoke_solver_ConvHipImplicitGemmBwdDataV1R1 GFX103X_ENABLED TEST ${TEST_CONV_VERBOSE_B} --input 32 128 32 32 --weights 12 128 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +# MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1=1 is necessary due to WORKAROUND_iGemm_936 in Jenkinsfile, +# which disables ConvHipImplicitGemmV4R1Fwd, but we still want to check that the solver is not broken. add_custom_test(smoke_solver_ConvHipImplicitGemmV4R1 GFX103X_ENABLED HALF_ENABLED BF16_ENABLED TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 + MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1=1 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R1Fwd $ ${TEST_CONV_VERBOSE_F} --input 256 32 27 27 --weights 128 32 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS}