diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 0436208268..27032593c4 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -1026,24 +1026,6 @@ set(DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_BWD_ENVS ${DYNAMIC_IMPLICITGEMM_COMMON} MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC) -add_custom_test(test_conv_igemm_dynamic_xdlops_bwd SKIP_UNLESS_ALL HALF_ENABLED GFX90A_DISABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 28 28 --weights 16 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 128 36 36 --weights 32 128 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 224 17 17 --weights 224 224 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 128 128 35 35 --weights 256 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 128 128 64 64 --weights 256 128 3 3 --pads_strides_dilations 1 1 2 2 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 128 768 17 17 --weights 256 768 3 3 --pads_strides_dilations 1 1 1 1 2 2 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 3 256 28 28 --weights 80 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 256 12 18 --weights 256 256 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 400 256 7 7 --weights 1024 256 7 7 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 400 256 1 1 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 8 16 5 5 --weights 8 16 2 2 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights -# ho=wo=1 stride=2 - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-weights -) - # TODO: disabled for WORKAROUND_ISSUE_1979 #add_custom_test(test_conv_igemm_dynamic_xdlops_bwd_group SKIP_UNLESS_ALL HALF_ENABLED FLOAT_DISABLED GFX90A_DISABLED GFX94X_ENABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON # ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} @@ -1052,299 +1034,7 @@ add_custom_test(test_conv_igemm_dynamic_xdlops_bwd SKIP_UNLESS_ALL HALF_ENABLED # COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 8 128 56 56 --weights 128 16 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 8 --disable-forward --disable-backward-weights #) -add_custom_test(test_conv_igemm_dynamic_xdlops_bwd_float SKIP_UNLESS_ALL HALF_DISABLED FLOAT_ENABLED GFX90A_DISABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 512 128 128 --weights 12 512 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights -) - # Be careful to add testings for (x=1, y=1, c % 8 != 0) due to WORKAROUND_SWDEV_306318 -add_custom_test(test_conv_igemm_dynamic_xdlops_fwd SKIP_UNLESS_ALL HALF_ENABLED GFX90A_DISABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_FWD_GTC_DYNAMIC_XDLOPS_ENVS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 1024 14 14 --weights 1024 1024 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 256 56 56 --weights 512 256 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 2048 7 7 --weights 2048 2048 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 128 128 17 17 --weights 128 128 7 1 --pads_strides_dilations 3 0 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 128 128 17 17 --weights 128 128 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 128 192 17 17 --weights 320 192 3 3 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 128 256 35 35 --weights 64 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 128 48 35 35 --weights 64 48 5 5 --pads_strides_dilations 2 2 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 256 100 104 --weights 12 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 256 28 28 --weights 80 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights -## ho=wo=1 stride=2 - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights -) - -add_custom_test(test_conv_igemm_dynamic_xdlops_fwd_half SKIP_UNLESS_ALL HALF_ENABLED FLOAT_DISABLED GFX90A_DISABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_FWD_GTC_DYNAMIC_XDLOPS_ENVS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 3 224 224 --weights 64 3 7 7 --pads_strides_dilations 3 3 2 2 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 3 230 230 --weights 64 3 7 7 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights -) - -add_custom_test(test_conv_igemm_dynamic_xdlops_wrw SKIP_UNLESS_ALL GFX90A_DISABLED GFX900_DISABLED GFX906_DISABLED HALF_ENABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_WRW_ENVS_XDLOPS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 28 28 --weights 32 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 128 36 36 --weights 32 128 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 224 17 17 --weights 224 224 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 128 128 35 35 --weights 256 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 128 128 64 64 --weights 256 128 3 3 --pads_strides_dilations 1 1 2 2 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 128 768 17 17 --weights 256 768 3 3 --pads_strides_dilations 1 1 1 1 2 2 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 3 256 28 28 --weights 80 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 256 12 18 --weights 256 256 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 512 128 128 --weights 12 512 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data -#regression test for issue 540 - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 32 79 141 --weights 64 32 5 10 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-data - - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 400 256 7 7 --weights 1024 256 7 7 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 400 256 1 1 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data -# Regression test for SWDEV-295434 (FP16 only). - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 120 256 3 3 --weights 340 256 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-data -# ho=wo=1 stride=2 - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-data -) - -add_custom_test(test_conv_igemm_dynamic_xdlops_wrw_half SKIP_UNLESS_ALL GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED HALF_ENABLED FLOAT_DISABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_WRW_ENVS_XDLOPS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 3 32 32 --weights 1 3 11 11 --pads_strides_dilations 1 1 2 2 2 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 3 224 224 --weights 1 3 3 3 --pads_strides_dilations 0 0 1 1 2 2 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 1 8 8 --weights 1 1 2 2 --pads_strides_dilations 0 0 1 1 2 2 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 128 56 56 --weights 1 128 5 5 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-data -) - -add_custom_test(test_conv_igemm_dynamic_xdlops_nhwc_fwd SKIP_UNLESS_ALL HALF_ENABLED GFX900_DISABLED GFX906_DISABLED GFX94X_ENABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_FWD_ENVS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 59 57 --weights 12 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 3 78 78 --weights 64 3 7 7 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 192 17 17 --weights 224 192 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 3 17 17 --weights 64 3 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 19 19 --weights 510 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC -# tensor larger than 4GB - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2048 1 512 1024 --weights 1 1 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC -# ho=wo=1 stride=2 - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC -) - -add_custom_test(test_conv_igemm_dynamic_xdlops_nhwc_fwd_nchw SKIP_UNLESS_ALL HALF_ENABLED GFX900_DISABLED GFX906_DISABLED GFX94X_ENABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_FWD_ENVS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 59 57 --weights 12 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 3 78 78 --weights 64 3 7 7 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 192 17 17 --weights 224 192 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 3 17 17 --weights 64 3 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 19 19 --weights 510 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-backward-data --disable-backward-weights -# TODO: disabled for WORKAROUND_ISSUE_1979 - #COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 3 224 224 --weights 63 1 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 3 --disable-backward-data --disable-backward-weights -) - -add_custom_test(test_conv_igemm_dynamic_xdlops_nhwc_bwd SKIP_UNLESS_ALL HALF_ENABLED GFX900_DISABLED GFX906_DISABLED GFX94X_ENABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_BWD_ENVS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 32 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 16 25 25 --weights 64 16 3 3 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 15 256 1 1 --weights 340 256 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 15 128 10 10 --weights 340 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 19 19 --weights 510 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC -# tensor larger than 4GB - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2048 1 512 1024 --weights 1 1 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC -# ho=wo=1 stride=2 - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC -) - -add_custom_test(test_conv_igemm_dynamic_xdlops_nhwc_bwd_nchw SKIP_UNLESS_ALL HALF_ENABLED GFX900_DISABLED GFX906_DISABLED GFX94X_ENABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_BWD_ENVS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 32 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 16 25 25 --weights 64 16 3 3 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 15 256 1 1 --weights 340 256 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 15 128 10 10 --weights 340 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 19 19 --weights 510 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-weights -) - -add_custom_test(test_conv_igemm_dynamic_xdlops_nhwc_fwd_bf16 SKIP_UNLESS_ALL BF16_ENABLED FLOAT_DISABLED HALF_DISABLED GFX908_DISABLED GFX94X_ENABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_FWD_ENVS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 59 57 --weights 12 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 3 78 78 --weights 64 3 7 7 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 192 17 17 --weights 224 192 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 3 17 17 --weights 64 3 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC -) - -add_custom_test(test_conv_igemm_dynamic_xdlops_nhwc_bwd_bf16 SKIP_UNLESS_ALL BF16_ENABLED FLOAT_DISABLED HALF_DISABLED GFX908_DISABLED GFX94X_ENABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_BWD_ENVS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 32 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 16 25 25 --weights 64 16 3 3 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 15 256 1 1 --weights 340 256 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 15 128 10 10 --weights 340 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-weights --in_layout NHWC --fil_layout NHWC --out_layout NHWC -) - -set(DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_WRW_ENVS - ${DYNAMIC_IMPLICITGEMM_COMMON} - MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC) - -set(ARGS_NHWC_WRW - --disable-forward - --disable-backward-data - --in_layout NHWC - --fil_layout NHWC - --out_layout NHWC) - -add_custom_test(test_conv_igemm_dynamic_xdlops_nhwc_wrw SKIP_UNLESS_ALL HALF_ENABLED GFX900_DISABLED GFX906_DISABLED GFX94X_ENABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_WRW_ENVS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 32 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 16 25 25 --weights 64 16 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NHWC_WRW} - - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 32 79 141 --weights 64 32 5 10 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NHWC_WRW} - - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 400 256 7 7 --weights 1024 256 7 7 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 400 256 1 1 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NHWC_WRW} - - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 3 32 32 --weights 1 3 11 11 --pads_strides_dilations 1 1 2 2 2 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 3 224 224 --weights 1 3 3 3 --pads_strides_dilations 0 0 1 1 2 2 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 1 8 8 --weights 1 1 2 2 --pads_strides_dilations 0 0 1 1 2 2 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 128 56 56 --weights 1 128 5 5 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 19 19 --weights 510 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${ARGS_NHWC_WRW} -# ho=wo=1 stride=2 - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NHWC_WRW} -) - -add_custom_test(test_conv_igemm_dynamic_xdlops_nhwc_wrw_nchw SKIP_UNLESS_ALL HALF_ENABLED GFX900_DISABLED GFX906_DISABLED GFX94X_ENABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_WRW_ENVS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 32 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 16 25 25 --weights 64 16 3 3 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data - - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 32 79 141 --weights 64 32 5 10 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-data - - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 400 256 7 7 --weights 1024 256 7 7 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 400 256 1 1 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data - - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 3 32 32 --weights 1 3 11 11 --pads_strides_dilations 1 1 2 2 2 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 3 224 224 --weights 1 3 3 3 --pads_strides_dilations 0 0 1 1 2 2 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 1 8 8 --weights 1 1 2 2 --pads_strides_dilations 0 0 1 1 2 2 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 128 56 56 --weights 1 128 5 5 --pads_strides_dilations 0 0 2 2 1 1 --disable-forward --disable-backward-data - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 19 19 --weights 510 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-forward --disable-backward-data -) - -add_custom_test(test_conv_igemm_dynamic_xdlops_nhwc_wrw_bf16 SKIP_UNLESS_ALL BF16_ENABLED FLOAT_DISABLED HALF_DISABLED GFX908_DISABLED GFX94X_ENABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_WRW_ENVS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 32 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 16 25 25 --weights 64 16 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NHWC_WRW} - - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 32 79 141 --weights 64 32 5 10 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NHWC_WRW} - - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 400 256 7 7 --weights 1024 256 7 7 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 400 256 1 1 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NHWC_WRW} - - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 3 32 32 --weights 1 3 11 11 --pads_strides_dilations 1 1 2 2 2 1 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 3 224 224 --weights 1 3 3 3 --pads_strides_dilations 0 0 1 1 2 2 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 1 8 8 --weights 1 1 2 2 --pads_strides_dilations 0 0 1 1 2 2 ${ARGS_NHWC_WRW} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 128 56 56 --weights 1 128 5 5 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NHWC_WRW} -) set(DYNAMIC_IMPLICITGEMM_DLOPS_NCHWC_FWD_ENVS ${DYNAMIC_IMPLICITGEMM_COMMON} diff --git a/test/gtest/conv_igemm_dynamic_xdlops.cpp b/test/gtest/conv_igemm_dynamic_xdlops.cpp index 98a2e0c823..2c97f5bd9b 100644 --- a/test/gtest/conv_igemm_dynamic_xdlops.cpp +++ b/test/gtest/conv_igemm_dynamic_xdlops.cpp @@ -31,9 +31,9 @@ #include "../conv2d.hpp" MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL) -MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_GPU_XNACK_ENABLED) +MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLOAT_ARG) -namespace { +namespace conv_igemm_dynamic_xdlops { auto GetTestCases() { @@ -44,7 +44,7 @@ auto GetTestCases() "ConvAsmImplicitGemmGTCDynamicFwdXdlops;" "ConvAsmImplicitGemmGTCDynamicWrwXdlops")}}; - const std::string v = " --verbose"; + const std::string cmd_v = " test_conv2d --verbose"; const std::string dis_bk_data = " --disable-backward-data"; const std::string dis_bk_wei = " --disable-backward-weights"; const std::string dis_fwd = " --disable-forward"; @@ -53,74 +53,82 @@ auto GetTestCases() return std::vector{ // clang-format off //bwd - std::pair{env_xdlops, v + " --input 64 64 28 28 --weights 16 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 16 128 36 36 --weights 32 128 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 64 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 64 224 17 17 --weights 224 224 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_fwd + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 128 128 35 35 --weights 256 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 128 128 64 64 --weights 256 128 3 3 --pads_strides_dilations 1 1 2 2 1 1" + dis_fwd + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 128 768 17 17 --weights 256 768 3 3 --pads_strides_dilations 1 1 1 1 2 2" + dis_fwd + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 3 256 28 28 --weights 80 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 2 256 12 18 --weights 256 256 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 400 256 7 7 --weights 1024 256 7 7 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 400 256 1 1 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 8 16 5 5 --weights 8 16 2 2 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 64 64 28 28 --weights 16 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 16 128 36 36 --weights 32 128 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 64 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 64 224 17 17 --weights 224 224 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_fwd + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 128 128 35 35 --weights 256 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 128 128 64 64 --weights 256 128 3 3 --pads_strides_dilations 1 1 2 2 1 1" + dis_fwd + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 128 768 17 17 --weights 256 768 3 3 --pads_strides_dilations 1 1 1 1 2 2" + dis_fwd + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 3 256 28 28 --weights 80 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 2 256 12 18 --weights 256 256 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 400 256 7 7 --weights 1024 256 7 7 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 400 256 1 1 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 8 16 5 5 --weights 8 16 2 2 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_wei}, //fwd //Be careful to add testings for (x=1, y=1, c % 8 != 0) due to WORKAROUND_SWDEV_306318 - std::pair{env_xdlops, v + " --input 64 1024 14 14 --weights 1024 1024 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 64 256 56 56 --weights 512 256 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 64 2048 7 7 --weights 2048 2048 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 128 128 17 17 --weights 128 128 7 1 --pads_strides_dilations 3 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 128 128 17 17 --weights 128 128 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_bk_data + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 128 192 17 17 --weights 320 192 3 3 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 128 256 35 35 --weights 64 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 128 48 35 35 --weights 64 48 5 5 --pads_strides_dilations 2 2 1 1 1 1" + dis_bk_data + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 64 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_bk_data + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 32 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 2 256 100 104 --weights 12 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, - std::pair{env_xdlops, v + " --input 1 256 28 28 --weights 80 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 64 1024 14 14 --weights 1024 1024 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 64 256 56 56 --weights 512 256 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 64 2048 7 7 --weights 2048 2048 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 128 128 17 17 --weights 128 128 7 1 --pads_strides_dilations 3 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 128 128 17 17 --weights 128 128 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 128 192 17 17 --weights 320 192 3 3 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 128 256 35 35 --weights 64 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 128 48 35 35 --weights 64 48 5 5 --pads_strides_dilations 2 2 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 64 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 32 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 2 256 100 104 --weights 12 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 1 256 28 28 --weights 80 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, //ho=wo=1 stride=2 - std::pair{env_xdlops, v + " --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1 " + dis_bk_data + dis_bk_wei}, + std::pair{env_xdlops, cmd_v + " --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1 " + dis_bk_data + dis_bk_wei}, //wrw - std::pair{env_xdlops, v + " --input 64 64 28 28 --weights 32 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, - std::pair{env_xdlops, v + " --input 16 128 36 36 --weights 32 128 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, - std::pair{env_xdlops, v + " --input 64 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, - std::pair{env_xdlops, v + " --input 64 224 17 17 --weights 224 224 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_fwd + dis_bk_data}, - std::pair{env_xdlops, v + " --input 128 128 35 35 --weights 256 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_data}, - std::pair{env_xdlops, v + " --input 128 128 64 64 --weights 256 128 3 3 --pads_strides_dilations 1 1 2 2 1 1" + dis_fwd + dis_bk_data}, - std::pair{env_xdlops, v + " --input 128 768 17 17 --weights 256 768 3 3 --pads_strides_dilations 1 1 1 1 2 2" + dis_fwd + dis_bk_data}, - std::pair{env_xdlops, v + " --input 3 256 28 28 --weights 80 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, - std::pair{env_xdlops, v + " --input 2 256 12 18 --weights 256 256 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_data}, - std::pair{env_xdlops, v + " --input 4 512 128 128 --weights 12 512 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + std::pair{env_xdlops, cmd_v + " --input 64 64 28 28 --weights 32 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + std::pair{env_xdlops, cmd_v + " --input 16 128 36 36 --weights 32 128 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + std::pair{env_xdlops, cmd_v + " --input 64 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + std::pair{env_xdlops, cmd_v + " --input 64 224 17 17 --weights 224 224 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_fwd + dis_bk_data}, + std::pair{env_xdlops, cmd_v + " --input 128 128 35 35 --weights 256 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_data}, + std::pair{env_xdlops, cmd_v + " --input 128 128 64 64 --weights 256 128 3 3 --pads_strides_dilations 1 1 2 2 1 1" + dis_fwd + dis_bk_data}, + std::pair{env_xdlops, cmd_v + " --input 128 768 17 17 --weights 256 768 3 3 --pads_strides_dilations 1 1 1 1 2 2" + dis_fwd + dis_bk_data}, + std::pair{env_xdlops, cmd_v + " --input 3 256 28 28 --weights 80 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + std::pair{env_xdlops, cmd_v + " --input 2 256 12 18 --weights 256 256 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_data}, + std::pair{env_xdlops, cmd_v + " --input 4 512 128 128 --weights 12 512 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, //regression test for issue 540 - std::pair{env_xdlops, v + " --input 4 32 79 141 --weights 64 32 5 10 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_data}, - std::pair{env_xdlops, v + " --input 400 256 7 7 --weights 1024 256 7 7 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, - std::pair{env_xdlops, v + " --input 400 256 1 1 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + std::pair{env_xdlops, cmd_v + " --input 4 32 79 141 --weights 64 32 5 10 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_data}, + std::pair{env_xdlops, cmd_v + " --input 400 256 7 7 --weights 1024 256 7 7 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + std::pair{env_xdlops, cmd_v + " --input 400 256 1 1 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, //Regression test for SWDEV-295434 (FP16 only). - std::pair{env_xdlops, v + " --input 120 256 3 3 --weights 340 256 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_data}, + std::pair{env_xdlops, cmd_v + " --input 120 256 3 3 --weights 340 256 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_data}, //ho=wo=1 stride=2 - std::pair{env_xdlops, v + " --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1 " + dis_fwd + dis_bk_data} + std::pair{env_xdlops, cmd_v + " --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1 " + dis_fwd + dis_bk_data} // clang-format on }; } using TestCase = decltype(GetTestCases())::value_type; -bool SkipTest() +static bool SkipTest(const std::string& float_arg) { - return miopen::IsEnabled(ENV(MIOPEN_TEST_GPU_XNACK_ENABLED)) || - miopen::IsDisabled(ENV(MIOPEN_TEST_ALL)); + if(miopen::IsUnset(ENV(MIOPEN_TEST_ALL))) + return false; + if(miopen::IsEnabled(ENV(MIOPEN_TEST_ALL))) + if(miopen::GetStringEnv(ENV(MIOPEN_TEST_FLOAT_ARG)) == float_arg) + return false; + return true; } -bool IsTestSupportedForDevice() +bool IsTestSupportedForDevice(const miopen::Handle& handle) { + const auto target = handle.GetTargetProperties(); + if(target.Xnack() && *target.Xnack()) + return false; using e_mask = enabled; using d_mask = disabled; return ::IsTestSupportedForDevMask(); } -} // namespace +} // namespace conv_igemm_dynamic_xdlops +using namespace conv_igemm_dynamic_xdlops; class Conv2dDefaultFloat : public FloatTestCase> { @@ -132,7 +140,8 @@ class Conv2dDefaultHalf : public HalfTestCase> TEST_P(Conv2dDefaultFloat, FloatTest_conv_igemm_dynamic_xdlops) { - if(IsTestSupportedForDevice() && !SkipTest()) + const auto& handle = get_handle(); + if(IsTestSupportedForDevice(handle) && !SkipTest("--float")) { invoke_with_params(default_check); } @@ -144,7 +153,8 @@ TEST_P(Conv2dDefaultFloat, FloatTest_conv_igemm_dynamic_xdlops) TEST_P(Conv2dDefaultHalf, HalfTest_conv_igemm_dynamic_xdlops) { - if(IsTestSupportedForDevice() && !SkipTest()) + const auto& handle = get_handle(); + if(IsTestSupportedForDevice(handle) && !SkipTest("--half")) { invoke_with_params(default_check); } diff --git a/test/gtest/conv_igemm_dynamic_xdlops_half.cpp b/test/gtest/conv_igemm_dynamic_xdlops_half.cpp new file mode 100644 index 0000000000..d4e9886531 --- /dev/null +++ b/test/gtest/conv_igemm_dynamic_xdlops_half.cpp @@ -0,0 +1,167 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (c) 2023 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 "../conv2d.hpp" +#include "get_handle.hpp" + +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL) +MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_FIND_MODE) +MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_DEBUG_FIND_ONLY_SOLVER) +MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLOAT_ARG) + +namespace conv_igemm_dynamic_xdlops_half { + +static bool SkipTest(const std::string& float_arg) +{ + if(miopen::IsUnset(ENV(MIOPEN_TEST_ALL))) + return false; + if(miopen::IsEnabled(ENV(MIOPEN_TEST_ALL))) + if(miopen::GetStringEnv(ENV(MIOPEN_TEST_FLOAT_ARG)) == float_arg) + return false; + return true; +} + +void SetupEnvVar(void) +{ + miopen::UpdateEnvVar(ENV(MIOPEN_FIND_MODE), std::string("normal")); + miopen::UpdateEnvVar( + ENV(MIOPEN_DEBUG_FIND_ONLY_SOLVER), + std::string( + "ConvAsmImplicitGemmGTCDynamicFwdXdlops;ConvAsmImplicitGemmGTCDynamicWrwXdlops")); +} + +void GetArgs(const std::string& param, std::vector& tokens) +{ + std::stringstream ss(param); + std::istream_iterator begin(ss); + std::istream_iterator end; + while(begin != end) + tokens.push_back(*begin++); +} + +class Conv2dHalf_conv_igemm_dynamic_xdlops_half + : public testing::TestWithParam> +{ +}; + +void Run2dDriver(miopenDataType_t prec) +{ + + std::vector params; + switch(prec) + { + case miopenHalf: params = Conv2dHalf_conv_igemm_dynamic_xdlops_half::GetParam(); break; + case miopenFloat: + case miopenInt8: + case miopenBFloat16: + case miopenInt32: + case miopenDouble: + case miopenFloat8: + case miopenBFloat8: + FAIL() << "miopenFloat, miopenInt8, miopenBFloat16, miopenInt32, " + "miopenDouble, miopenFloat8, miopenBFloat8 " + "data type not supported by conv_igemm_dynamic_xdlops_half test"; + + default: params = Conv2dHalf_conv_igemm_dynamic_xdlops_half::GetParam(); + } + + SetupEnvVar(); + + for(const auto& test_value : params) + { + std::vector tokens; + GetArgs(test_value, tokens); + std::vector ptrs; + + std::transform(tokens.begin(), + tokens.end(), + std::back_inserter(ptrs), + [](const std::string& str) { return str.data(); }); + + testing::internal::CaptureStderr(); + test_drive(ptrs.size(), ptrs.data()); + auto capture = testing::internal::GetCapturedStderr(); + std::cout << capture; + } +}; + +bool IsTestSupportedForDevice(const miopen::Handle& handle) +{ + const auto target = handle.GetTargetProperties(); + std::string devName = handle.GetDeviceName(); + if(target.Xnack() && *target.Xnack()) + return false; + + if(devName == "gfx908") + return true; + else + return false; +} + +std::vector GetTestCases(const std::string& precision) +{ + const std::string flags = "test_conv2d " + precision + " --verbose "; + const std::string dis_bk_data = " --disable-backward-data"; + const std::string dis_bk_wei = " --disable-backward-weights"; + const std::string dis_fwd = " --disable-forward"; + + const std::vector test_cases = { + // clang-format off + //fwd + {flags + " --input 64 3 224 224 --weights 64 3 7 7 --pads_strides_dilations 3 3 2 2 1 1" + dis_bk_data + dis_bk_wei}, + {flags + " --input 64 3 230 230 --weights 64 3 7 7 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei}, + + //wrw + {flags + " --input 1 3 32 32 --weights 1 3 11 11 --pads_strides_dilations 1 1 2 2 2 1" + dis_fwd + dis_bk_data}, + {flags + " --input 1 3 224 224 --weights 1 3 3 3 --pads_strides_dilations 0 0 1 1 2 2" + dis_fwd + dis_bk_data}, + {flags + " --input 1 1 8 8 --weights 1 1 2 2 --pads_strides_dilations 0 0 1 1 2 2" + dis_fwd + dis_bk_data}, + {flags + " --input 1 128 56 56 --weights 1 128 5 5 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_data} + // clang-format on + }; + return test_cases; +} + +} // namespace conv_igemm_dynamic_xdlops_half +using namespace conv_igemm_dynamic_xdlops_half; + +TEST_P(Conv2dHalf_conv_igemm_dynamic_xdlops_half, HalfTest_conv_igemm_dynamic_xdlops_half) +{ + const auto& handle = get_handle(); + if(IsTestSupportedForDevice(handle) && !SkipTest("--half")) + { + Run2dDriver(miopenHalf); + } + else + { + GTEST_SKIP(); + } +}; + +INSTANTIATE_TEST_SUITE_P(ConvIgemmDynamicXdlopsFwdWrw, + Conv2dHalf_conv_igemm_dynamic_xdlops_half, + testing::Values(GetTestCases("--half"))); diff --git a/test/gtest/conv_igemm_dynamic_xdlops_nhwc_bf16.cpp b/test/gtest/conv_igemm_dynamic_xdlops_nhwc_bf16.cpp new file mode 100644 index 0000000000..f8e6777725 --- /dev/null +++ b/test/gtest/conv_igemm_dynamic_xdlops_nhwc_bf16.cpp @@ -0,0 +1,223 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (c) 2023 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 "../conv2d.hpp" +#include "get_handle.hpp" + +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL) +MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_FIND_MODE) +MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_DEBUG_FIND_ONLY_SOLVER) +MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLOAT_ARG) + +namespace conv_igemm_dynamic_xdlops_nhwc_bf16 { + +static bool SkipTest(const std::string& float_arg) +{ + if(miopen::IsUnset(ENV(MIOPEN_TEST_ALL))) + return false; + if(miopen::IsEnabled(ENV(MIOPEN_TEST_ALL))) + if(miopen::GetStringEnv(ENV(MIOPEN_TEST_FLOAT_ARG)) == float_arg) + return false; + return true; +} + +void SetupEnvVar(void) +{ + miopen::UpdateEnvVar(ENV(MIOPEN_FIND_MODE), std::string("normal")); + miopen::UpdateEnvVar( + ENV(MIOPEN_DEBUG_FIND_ONLY_SOLVER), + std::string( + "ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC;ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC;" + "ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC")); +} + +void GetArgs(const std::string& param, std::vector& tokens) +{ + std::stringstream ss(param); + std::istream_iterator begin(ss); + std::istream_iterator end; + while(begin != end) + tokens.push_back(*begin++); +} + +class Conv2dBf16 : public testing::TestWithParam> +{ +}; + +void Run2dDriver(miopenDataType_t prec) +{ + + std::vector params; + switch(prec) + { + case miopenBFloat16: params = Conv2dBf16::GetParam(); break; + case miopenFloat: + case miopenHalf: + case miopenInt8: + case miopenInt32: + case miopenDouble: + case miopenFloat8: + case miopenBFloat8: + FAIL() << "miopenFloat, miopenHalf, miopenInt8, miopenInt32, " + "miopenDouble, miopenFloat8, miopenBFloat8 " + "data type not supported by conv_igemm_dynamic_xdlops_nhwc_bf16 test"; + + default: params = Conv2dBf16::GetParam(); + } + + SetupEnvVar(); + + for(const auto& test_value : params) + { + std::vector tokens; + GetArgs(test_value, tokens); + std::vector ptrs; + + std::transform(tokens.begin(), + tokens.end(), + std::back_inserter(ptrs), + [](const std::string& str) { return str.data(); }); + + testing::internal::CaptureStderr(); + test_drive(ptrs.size(), ptrs.data()); + auto capture = testing::internal::GetCapturedStderr(); + std::cout << capture; + } +}; + +bool IsTestSupportedForDevice(const miopen::Handle& handle) +{ + const auto target = handle.GetTargetProperties(); + std::string devName = handle.GetDeviceName(); + if(target.Xnack() && *target.Xnack()) + return false; + + if(devName == "gfx90a" || miopen::StartsWith(devName, "gfx94")) + return true; + else + return false; +} + +std::vector GetTestCases(const std::string& precision) +{ + const std::string flags = "test_conv2d " + precision + " --verbose "; + const std::string dis_bk_data = " --disable-backward-data"; + const std::string dis_bk_wei = " --disable-backward-weights"; + const std::string dis_fwd = " --disable-forward"; + const std::string dis_vali = " --disable-validation"; + const std::string in_nhwc = " --in_layout NHWC"; + const std::string fil_nhwc = " --fil_layout NHWC"; + const std::string out_nhwc = " --out_layout NHWC"; + const std::string args_nhwc_fwd = dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc; + const std::string args_nhwc_bwd = dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc; + const std::string args_nhwc_wrw = dis_fwd + dis_bk_data + in_nhwc + fil_nhwc + out_nhwc; + + const std::vector test_cases = { + // clang-format off + //fwd + {flags + " --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_fwd}, + {flags + " --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_fwd}, + {flags + " --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_fwd}, + {flags + " --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_fwd}, + {flags + " --input 2 64 59 57 --weights 12 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_fwd}, + {flags + " --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1" + args_nhwc_fwd}, + {flags + " --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1" + args_nhwc_fwd}, + {flags + " --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1" + args_nhwc_fwd}, + {flags + " --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1" + args_nhwc_fwd}, + {flags + " --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1" + args_nhwc_fwd}, + {flags + " --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1" + args_nhwc_fwd}, + {flags + " --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1" + args_nhwc_fwd}, + {flags + " --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1" + args_nhwc_fwd}, + {flags + " --input 64 3 78 78 --weights 64 3 7 7 --pads_strides_dilations 0 0 2 2 1 1" + args_nhwc_fwd}, + {flags + " --input 16 192 17 17 --weights 224 192 1 7 --pads_strides_dilations 0 3 1 1 1 1" + args_nhwc_fwd}, + {flags + " --input 16 3 17 17 --weights 64 3 1 1 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_fwd}, + + //nhwc_bwd + {flags + " --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_bwd}, + {flags + " --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_bwd}, + {flags + " --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_bwd}, + {flags + " --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_bwd}, + {flags + " --input 2 64 32 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_bwd}, + {flags + " --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1" + args_nhwc_bwd}, + {flags + " --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1" + args_nhwc_bwd}, + {flags + " --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1" + args_nhwc_bwd}, + {flags + " --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1" + args_nhwc_bwd}, + {flags + " --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1" + args_nhwc_bwd}, + {flags + " --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1" + args_nhwc_bwd}, + {flags + " --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1" + args_nhwc_bwd}, + {flags + " --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1" + args_nhwc_bwd}, + {flags + " --input 16 16 25 25 --weights 64 16 3 3 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_bwd}, + {flags + " --input 15 256 1 1 --weights 340 256 3 3 --pads_strides_dilations 1 1 1 1 1 1" + args_nhwc_bwd}, + {flags + " --input 15 128 10 10 --weights 340 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + args_nhwc_bwd}, + + //nhwc_wrw + {flags + " --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 " + args_nhwc_wrw}, + {flags + " --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1 " + args_nhwc_wrw}, + {flags + " --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 " + args_nhwc_wrw}, + {flags + " --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 " + args_nhwc_wrw}, + {flags + " --input 2 64 32 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 " + args_nhwc_wrw}, + {flags + " --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1 " + args_nhwc_wrw}, + {flags + " --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1 " + args_nhwc_wrw}, + {flags + " --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1 " + args_nhwc_wrw}, + {flags + " --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1 " + args_nhwc_wrw}, + {flags + " --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1 " + args_nhwc_wrw}, + {flags + " --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1 " + args_nhwc_wrw}, + {flags + " --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1 " + args_nhwc_wrw}, + {flags + " --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 " + args_nhwc_wrw}, + {flags + " --input 16 16 25 25 --weights 64 16 3 3 --pads_strides_dilations 0 0 1 1 1 1 " + args_nhwc_wrw}, + {flags + " --input 4 32 79 141 --weights 64 32 5 10 --pads_strides_dilations 0 0 2 2 1 1 " + args_nhwc_wrw}, + {flags + " --input 400 256 7 7 --weights 1024 256 7 7 --pads_strides_dilations 0 0 1 1 1 1 " + args_nhwc_wrw}, + {flags + " --input 400 256 1 1 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 " + args_nhwc_wrw}, + {flags + " --input 1 3 32 32 --weights 1 3 11 11 --pads_strides_dilations 1 1 2 2 2 1 " + args_nhwc_wrw}, + {flags + " --input 1 3 224 224 --weights 1 3 3 3 --pads_strides_dilations 0 0 1 1 2 2 " + args_nhwc_wrw}, + {flags + " --input 1 1 8 8 --weights 1 1 2 2 --pads_strides_dilations 0 0 1 1 2 2 " + args_nhwc_wrw}, + {flags + " --input 1 128 56 56 --weights 1 128 5 5 --pads_strides_dilations 0 0 2 2 1 1 " + args_nhwc_wrw} + // clang-format on + }; + return test_cases; +} + +} // namespace conv_igemm_dynamic_xdlops_nhwc_bf16 +using namespace conv_igemm_dynamic_xdlops_nhwc_bf16; + +TEST_P(Conv2dBf16, Bf16Test_conv_igemm_dynamic_xdlops_nhwc_bf16) +{ + const auto& handle = get_handle(); + if(IsTestSupportedForDevice(handle) && !SkipTest("--bfloat16")) + { + Run2dDriver(miopenBFloat16); + } + else + { + GTEST_SKIP(); + } +}; + +INSTANTIATE_TEST_SUITE_P(ConvIgemmDynamicXdlopsNhwc, + Conv2dBf16, + testing::Values(GetTestCases("--bfloat16"))); diff --git a/test/gtest/conv_igemm_dynamic_xdlops_nhwc_nchw.cpp b/test/gtest/conv_igemm_dynamic_xdlops_nhwc_nchw.cpp new file mode 100644 index 0000000000..657e14fe19 --- /dev/null +++ b/test/gtest/conv_igemm_dynamic_xdlops_nhwc_nchw.cpp @@ -0,0 +1,321 @@ +/******************************************************************************* + * + * MIT License + * + * Copyright (c) 2023 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 "../conv2d.hpp" +#include "get_handle.hpp" + +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL) +MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_FIND_MODE) +MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_DEBUG_FIND_ONLY_SOLVER) +MIOPEN_DECLARE_ENV_VAR_STR(MIOPEN_TEST_FLOAT_ARG) + +namespace conv_igemm_dynamic_xdlops_nhwc_nchw { + +static bool SkipTest(const std::string& float_arg) +{ + if(miopen::IsUnset(ENV(MIOPEN_TEST_ALL))) + return false; + if(miopen::IsEnabled(ENV(MIOPEN_TEST_ALL))) + if(miopen::GetStringEnv(ENV(MIOPEN_TEST_FLOAT_ARG)) == float_arg) + return false; + return true; +} + +void SetupEnvVar(void) +{ + miopen::UpdateEnvVar(ENV(MIOPEN_FIND_MODE), std::string("normal")); + miopen::UpdateEnvVar( + ENV(MIOPEN_DEBUG_FIND_ONLY_SOLVER), + std::string( + "ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC;ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC;" + "ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC")); +} + +void GetArgs(const std::string& param, std::vector& tokens) +{ + std::stringstream ss(param); + std::istream_iterator begin(ss); + std::istream_iterator end; + while(begin != end) + tokens.push_back(*begin++); +} + +class Conv2dFloat_conv_igemm_dynamic_xdlops_nhwc_nchw + : public testing::TestWithParam> +{ +}; + +class Conv2dHalf_conv_igemm_dynamic_xdlops_nhwc_nchw + : public testing::TestWithParam> +{ +}; + +void Run2dDriver(miopenDataType_t prec) +{ + + std::vector params; + switch(prec) + { + case miopenFloat: params = Conv2dFloat_conv_igemm_dynamic_xdlops_nhwc_nchw::GetParam(); break; + case miopenHalf: params = Conv2dHalf_conv_igemm_dynamic_xdlops_nhwc_nchw::GetParam(); break; + case miopenInt8: + case miopenBFloat16: + case miopenInt32: + case miopenDouble: + case miopenFloat8: + case miopenBFloat8: + default: + FAIL() << "miopenInt8, miopenBFloat16, miopenInt32, " + "miopenDouble, miopenFloat8, miopenBFloat8 " + "data type not supported by conv_igemm_dynamic_xdlops_nhwc_nchw test"; + } + + SetupEnvVar(); + + for(const auto& test_value : params) + { + std::vector tokens; + GetArgs(test_value, tokens); + std::vector ptrs; + + std::transform(tokens.begin(), + tokens.end(), + std::back_inserter(ptrs), + [](const std::string& str) { return str.data(); }); + + testing::internal::CaptureStderr(); + test_drive(ptrs.size(), ptrs.data()); + auto capture = testing::internal::GetCapturedStderr(); + std::cout << capture; + } +}; + +bool IsTestSupportedForDevice(const miopen::Handle& handle) +{ + const auto target = handle.GetTargetProperties(); + std::string devName = handle.GetDeviceName(); + if(target.Xnack() && *target.Xnack()) + return false; + if(devName == "gfx908" || devName == "gfx90a" || miopen::StartsWith(devName, "gfx94")) + return true; + else + return false; +} + +std::vector GetTestCases(const std::string& precision) +{ + const std::string flags = "test_conv2d " + precision + " --verbose "; + const std::string dis_bk_data = " --disable-backward-data"; + const std::string dis_bk_wei = " --disable-backward-weights"; + const std::string dis_fwd = " --disable-forward"; + const std::string dis_vali = " --disable-validation"; + const std::string in_nhwc = " --in_layout NHWC"; + const std::string fil_nhwc = " --fil_layout NHWC"; + const std::string out_nhwc = " --out_layout NHWC"; + const std::string args_nhwc_wrw = dis_fwd + dis_bk_data + in_nhwc + fil_nhwc + out_nhwc; + + const std::vector test_cases = { + // clang-format off + //nhwc_fwd + {flags + " --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 2 64 59 57 --weights 12 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 64 3 78 78 --weights 64 3 7 7 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 16 192 17 17 --weights 224 192 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 16 3 17 17 --weights 64 3 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 2 64 19 19 --weights 510 64 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + // tensor larger than 4GB + {flags + " --input 2048 1 512 1024 --weights 1 1 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + // ho=wo=1 stride=2 + {flags + " --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + + //nhwc_fwd_nchw + {flags + " --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + {flags + " --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + {flags + " --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + {flags + " --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + {flags + " --input 2 64 59 57 --weights 12 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + {flags + " --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei}, + {flags + " --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_bk_data + dis_bk_wei}, + {flags + " --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + {flags + " --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei}, + {flags + " --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + {flags + " --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei}, + {flags + " --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_bk_data + dis_bk_wei}, + {flags + " --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_bk_data + dis_bk_wei}, + {flags + " --input 64 3 78 78 --weights 64 3 7 7 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei}, + {flags + " --input 16 192 17 17 --weights 224 192 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_bk_data + dis_bk_wei}, + {flags + " --input 16 3 17 17 --weights 64 3 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + {flags + " --input 2 64 19 19 --weights 510 64 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_bk_data + dis_bk_wei}, + // TODO: disabled for WORKAROUND_ISSUE_1979 + //{flags + " --input 16 3 224 224 --weights 63 1 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 3" + dis_bk_data + dis_bk_wei}, + + //nhwc_bwd + {flags + " --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 2 64 32 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 16 16 25 25 --weights 64 16 3 3 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 15 256 1 1 --weights 340 256 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 15 128 10 10 --weights 340 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + {flags + " --input 2 64 19 19 --weights 510 64 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + // tensor larger than 4GB + {flags + " --input 2048 1 512 1024 --weights 1 1 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + // ho=wo=1 stride=2 + {flags + " --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_wei + in_nhwc + fil_nhwc + out_nhwc}, + + //nhwc_bwd_nchw + {flags + " --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 " + dis_fwd + dis_bk_wei}, + {flags + " --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1 " + dis_fwd + dis_bk_wei}, + {flags + " --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 " + dis_fwd + dis_bk_wei}, + {flags + " --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 " + dis_fwd + dis_bk_wei}, + {flags + " --input 2 64 32 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 " + dis_fwd + dis_bk_wei}, + {flags + " --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1 " + dis_fwd + dis_bk_wei}, + {flags + " --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1 " + dis_fwd + dis_bk_wei}, + {flags + " --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1 " + dis_fwd + dis_bk_wei}, + {flags + " --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1 " + dis_fwd + dis_bk_wei}, + {flags + " --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1 " + dis_fwd + dis_bk_wei}, + {flags + " --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1 " + dis_fwd + dis_bk_wei}, + {flags + " --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1 " + dis_fwd + dis_bk_wei}, + {flags + " --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 " + dis_fwd + dis_bk_wei}, + {flags + " --input 16 16 25 25 --weights 64 16 3 3 --pads_strides_dilations 0 0 1 1 1 1 " + dis_fwd + dis_bk_wei}, + {flags + " --input 15 256 1 1 --weights 340 256 3 3 --pads_strides_dilations 1 1 1 1 1 1 " + dis_fwd + dis_bk_wei}, + {flags + " --input 15 128 10 10 --weights 340 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 " + dis_fwd + dis_bk_wei}, + {flags + " --input 2 64 19 19 --weights 510 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 " + dis_fwd + dis_bk_wei}, + + //nhwc_wrw + {flags + " --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_wrw}, + {flags + " --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_wrw}, + {flags + " --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_wrw}, + {flags + " --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_wrw}, + {flags + " --input 2 64 32 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_wrw}, + {flags + " --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1" + args_nhwc_wrw}, + {flags + " --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1" + args_nhwc_wrw}, + {flags + " --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1" + args_nhwc_wrw}, + {flags + " --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1" + args_nhwc_wrw}, + {flags + " --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1" + args_nhwc_wrw}, + {flags + " --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1" + args_nhwc_wrw}, + {flags + " --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1" + args_nhwc_wrw}, + {flags + " --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1" + args_nhwc_wrw}, + {flags + " --input 16 16 25 25 --weights 64 16 3 3 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_wrw}, + {flags + " --input 4 32 79 141 --weights 64 32 5 10 --pads_strides_dilations 0 0 2 2 1 1" + args_nhwc_wrw}, + {flags + " --input 400 256 7 7 --weights 1024 256 7 7 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_wrw}, + {flags + " --input 400 256 1 1 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + args_nhwc_wrw}, + {flags + " --input 1 3 32 32 --weights 1 3 11 11 --pads_strides_dilations 1 1 2 2 2 1" + args_nhwc_wrw}, + {flags + " --input 1 3 224 224 --weights 1 3 3 3 --pads_strides_dilations 0 0 1 1 2 2" + args_nhwc_wrw}, + {flags + " --input 1 1 8 8 --weights 1 1 2 2 --pads_strides_dilations 0 0 1 1 2 2" + args_nhwc_wrw}, + {flags + " --input 1 128 56 56 --weights 1 128 5 5 --pads_strides_dilations 0 0 2 2 1 1" + args_nhwc_wrw}, + {flags + " --input 2 64 19 19 --weights 510 64 3 3 --pads_strides_dilations 1 1 1 1 1 1" + args_nhwc_wrw}, + // ho=wo=1 stride=2 + {flags + " --input 256 2048 2 2 --weights 1024 2048 1 1 --pads_strides_dilations 0 0 2 2 1 1" + args_nhwc_wrw}, + + //nhwc_wrw_nchw + {flags + " --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + {flags + " --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + {flags + " --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + {flags + " --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + {flags + " --input 2 64 32 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + {flags + " --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_data}, + {flags + " --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_fwd + dis_bk_data}, + {flags + " --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1" + dis_fwd + dis_bk_data}, + {flags + " --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_data}, + {flags + " --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1" + dis_fwd + dis_bk_data}, + {flags + " --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_data}, + {flags + " --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_data}, + {flags + " --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_data}, + {flags + " --input 16 16 25 25 --weights 64 16 3 3 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + {flags + " --input 4 32 79 141 --weights 64 32 5 10 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_data}, + {flags + " --input 400 256 7 7 --weights 1024 256 7 7 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + {flags + " --input 400 256 1 1 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + {flags + " --input 1 3 32 32 --weights 1 3 11 11 --pads_strides_dilations 1 1 2 2 2 1" + dis_fwd + dis_bk_data}, + {flags + " --input 1 3 224 224 --weights 1 3 3 3 --pads_strides_dilations 0 0 1 1 2 2" + dis_fwd + dis_bk_data}, + {flags + " --input 1 1 8 8 --weights 1 1 2 2 --pads_strides_dilations 0 0 1 1 2 2" + dis_fwd + dis_bk_data}, + {flags + " --input 1 128 56 56 --weights 1 128 5 5 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_data}, + {flags + " --input 2 64 19 19 --weights 510 64 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_data} + // clang-format on + }; + return test_cases; +} + +} // namespace conv_igemm_dynamic_xdlops_nhwc_nchw + +using namespace conv_igemm_dynamic_xdlops_nhwc_nchw; + +TEST_P(Conv2dFloat_conv_igemm_dynamic_xdlops_nhwc_nchw, + FloatTest_conv_igemm_dynamic_xdlops_nhwc_nchw) +{ + const auto& handle = get_handle(); + if(IsTestSupportedForDevice(handle) && !SkipTest("--float")) + { + Run2dDriver(miopenFloat); + } + else + { + GTEST_SKIP(); + } +}; + +TEST_P(Conv2dHalf_conv_igemm_dynamic_xdlops_nhwc_nchw, HalfTest_conv_igemm_dynamic_xdlops_nhwc_nchw) +{ + const auto& handle = get_handle(); + if(IsTestSupportedForDevice(handle) && !SkipTest("--half")) + { + Run2dDriver(miopenHalf); + } + else + { + GTEST_SKIP(); + } +}; + +INSTANTIATE_TEST_SUITE_P(ConvIgemmDynamicXdlopsNhwcNchw, + Conv2dFloat_conv_igemm_dynamic_xdlops_nhwc_nchw, + testing::Values(GetTestCases("--float"))); + +INSTANTIATE_TEST_SUITE_P(ConvIgemmDynamicXdlopsNhwcNchw, + Conv2dHalf_conv_igemm_dynamic_xdlops_nhwc_nchw, + testing::Values(GetTestCases("--half"))); diff --git a/test/gtest/smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops.cpp b/test/gtest/smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops.cpp index b3804c752a..de2eb9311a 100644 --- a/test/gtest/smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops.cpp +++ b/test/gtest/smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops.cpp @@ -70,19 +70,24 @@ bool IsTestSupportedForDevice() } // namespace -class Conv2dFloat : public FloatTestCase> +class Conv2dFloat_smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops + : public FloatTestCase> { }; -class Conv2dHalf : public HalfTestCase> +class Conv2dHalf_smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops + : public HalfTestCase> { }; -TEST_P(Conv2dFloat, FloatTest_smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops) +TEST_P(Conv2dFloat_smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops, + FloatTest_smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops) { if(IsTestSupportedForDevice() && !SkipTest()) { - invoke_with_params(default_check); + invoke_with_params( + default_check); } else { @@ -90,11 +95,14 @@ TEST_P(Conv2dFloat, FloatTest_smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops) } }; -TEST_P(Conv2dHalf, HalfTest_smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops) +TEST_P(Conv2dHalf_smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops, + HalfTest_smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops) { if(IsTestSupportedForDevice() && !SkipTest()) { - invoke_with_params(default_check); + invoke_with_params( + default_check); } else { @@ -103,9 +111,9 @@ TEST_P(Conv2dHalf, HalfTest_smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops) }; INSTANTIATE_TEST_SUITE_P(SmokeSolverConvAsmImplicitGemmV4R1Dynamic, - Conv2dFloat, + Conv2dFloat_smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops, testing::Values(GetTestCases())); INSTANTIATE_TEST_SUITE_P(SmokeSolverConvAsmImplicitGemmV4R1Dynamic, - Conv2dHalf, + Conv2dHalf_smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops, testing::Values(GetTestCases()));