diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 43c14fcfe7..995d7057ad 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -412,8 +412,6 @@ set(LONG_TESTS test_conv3d_extra test_conv_3d test_pooling2d - test_conv_igemm_mlir - test_conv_igemm_mlir_xdlops test_activation test_conv_ck_igemm_fwd_v6r1_dlops_nchw ) @@ -772,18 +770,6 @@ set(IMPLICITGEMM_MLIR_ENV_F_XDLOPS ${IMPLICITGEMM_MLIR_ENV_BASE} MIOPEN_DEBUG_FI set(IMPLICITGEMM_MLIR_ENV_B_XDLOPS ${IMPLICITGEMM_MLIR_ENV_BASE} MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvMlirIgemmBwdXdlops) set(IMPLICITGEMM_MLIR_ENV_W_XDLOPS ${IMPLICITGEMM_MLIR_ENV_BASE} MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvMlirIgemmWrWXdlops) -add_custom_test(test_conv_igemm_mlir_xdlops_fwd SKIP_UNLESS_ALL HALF_ENABLED INT8_ENABLED SKIP_UNLESS_MLIR GFX900_DISABLED GFX906_DISABLED - ENVIRONMENT ${IMPLICITGEMM_MLIR_ENV_F_XDLOPS} - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 256 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1 - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 128 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 128 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 256 256 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --group-count 4 -) - set(IMPLICITGEMM_TESTING_ENV MIOPEN_DEBUG_CONV_WINOGRAD=0 MIOPEN_DEBUG_CONV_FFT=0 @@ -1553,102 +1539,6 @@ set(ARGS_NCHWC_CHWNC_FWD_FP16x8 --tensor_vect 1 --vector_length 8) -add_custom_test(test_conv_igemm_dynamic_dlops_nchwc_nchwc_fwd_fp16x4 SKIP_UNLESS_ALL HALF_ENABLED FLOAT_DISABLED BF16_DISABLED GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED GFX908_DISABLED GFX103X_ENABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_DLOPS_NCHWC_FWD_ENVS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 8 10 10 --weights 8 8 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NCHWC_NCHWC_FWD_FP16x4} - 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_NCHWC_NCHWC_FWD_FP16x4} - 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_NCHWC_NCHWC_FWD_FP16x4} - 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_NCHWC_NCHWC_FWD_FP16x4} - 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_NCHWC_NCHWC_FWD_FP16x4} - 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_NCHWC_NCHWC_FWD_FP16x4} - 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_NCHWC_NCHWC_FWD_FP16x4} - 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_NCHWC_NCHWC_FWD_FP16x4} - 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_NCHWC_NCHWC_FWD_FP16x4} - 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_NCHWC_NCHWC_FWD_FP16x4} - 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_NCHWC_NCHWC_FWD_FP16x4} - 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_NCHWC_NCHWC_FWD_FP16x4} - 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_NCHWC_NCHWC_FWD_FP16x4} - 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_NCHWC_NCHWC_FWD_FP16x4} - 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_NCHWC_NCHWC_FWD_FP16x4} - 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_NCHWC_NCHWC_FWD_FP16x4} - 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_NCHWC_NCHWC_FWD_FP16x4} -) - -add_custom_test(test_conv_igemm_dynamic_dlops_nchwc_chwnc_fwd_fp16x4 SKIP_UNLESS_ALL HALF_ENABLED FLOAT_DISABLED BF16_DISABLED GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED GFX908_DISABLED GFX94X_ENABLED GFX103X_ENABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_DLOPS_NCHWC_FWD_ENVS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 256 7 7 --weights 256 3 3 128 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x4} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 160 73 73 --weights 160 1 1 64 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x4} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 1 1 64 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x4} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 256 40 52 --weights 256 1 1 256 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x4} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 32 28 --weights 64 1 1 64 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x4} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 14 14 --weights 128 1 1 64 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x4} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 64 3 7 192 --pads_strides_dilations 0 3 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x4} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 64 7 1 192 --pads_strides_dilations 3 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x4} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 128 28 28 --weights 128 2 2 128 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x4} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 8 8 --weights 128 3 1 192 --pads_strides_dilations 1 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x4} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 192 17 17 --weights 192 3 3 160 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x4} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 32 73 73 --weights 32 3 3 64 --pads_strides_dilations 1 1 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x4} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 3 3 64 --pads_strides_dilations 1 1 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x4} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 16 25 25 --weights 16 3 3 64 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x4} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 32 79 141 --weights 32 5 10 64 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x4} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 400 256 7 7 --weights 256 7 7 1024 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x4} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 400 256 1 1 --weights 256 1 1 1024 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x4} -) - -add_custom_test(test_conv_igemm_dynamic_dlops_nchwc_nchwc_fwd_fp16x8 SKIP_UNLESS_ALL HALF_ENABLED FLOAT_DISABLED BF16_DISABLED GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED GFX908_DISABLED GFX103X_ENABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_DLOPS_NCHWC_FWD_ENVS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 8 10 10 --weights 8 8 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NCHWC_NCHWC_FWD_FP16x8} - 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_NCHWC_NCHWC_FWD_FP16x8} - 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_NCHWC_NCHWC_FWD_FP16x8} - 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_NCHWC_NCHWC_FWD_FP16x8} - 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_NCHWC_NCHWC_FWD_FP16x8} - 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_NCHWC_NCHWC_FWD_FP16x8} - 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_NCHWC_NCHWC_FWD_FP16x8} - 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_NCHWC_NCHWC_FWD_FP16x8} - 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_NCHWC_NCHWC_FWD_FP16x8} - 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_NCHWC_NCHWC_FWD_FP16x8} - 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_NCHWC_NCHWC_FWD_FP16x8} - 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_NCHWC_NCHWC_FWD_FP16x8} - 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_NCHWC_NCHWC_FWD_FP16x8} - 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_NCHWC_NCHWC_FWD_FP16x8} - 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_NCHWC_NCHWC_FWD_FP16x8} - 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_NCHWC_NCHWC_FWD_FP16x8} - 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_NCHWC_NCHWC_FWD_FP16x8} -) - -add_custom_test(test_conv_igemm_dynamic_dlops_nchwc_chwnc_fwd_fp16x8 SKIP_UNLESS_ALL HALF_ENABLED FLOAT_DISABLED BF16_DISABLED GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED GFX908_DISABLED GFX94X_ENABLED GFX103X_ENABLED SKIP_XNACK_ON - ENVIRONMENT ${DYNAMIC_IMPLICITGEMM_DLOPS_NCHWC_FWD_ENVS} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 256 7 7 --weights 256 1 1 128 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x8} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 160 73 73 --weights 160 1 1 64 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x8} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 1 1 64 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x8} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 256 40 52 --weights 256 1 1 256 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x8} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 2 64 32 28 --weights 64 1 1 64 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x8} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 14 14 --weights 128 1 1 64 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x8} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 64 1 7 192 --pads_strides_dilations 0 3 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x8} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 17 17 --weights 64 7 1 192 --pads_strides_dilations 3 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x8} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 128 28 28 --weights 128 2 2 128 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x8} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 128 8 8 --weights 128 3 1 192 --pads_strides_dilations 1 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x8} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 192 17 17 --weights 192 3 3 160 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x8} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 32 73 73 --weights 32 3 3 64 --pads_strides_dilations 1 1 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x8} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 64 56 56 --weights 64 3 3 64 --pads_strides_dilations 1 1 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x8} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 16 25 25 --weights 16 3 3 64 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x8} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 32 79 141 --weights 32 5 10 64 --pads_strides_dilations 0 0 2 2 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x8} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 400 256 7 7 --weights 256 7 7 1024 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x8} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 400 256 1 1 --weights 256 1 1 1024 --pads_strides_dilations 0 0 1 1 1 1 ${ARGS_NCHWC_CHWNC_FWD_FP16x8} -) - -add_custom_test(test_regression_half_mi100 SKIP_UNLESS_ALL FLOAT_DISABLED HALF_ENABLED GFX908_ENABLED GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED - # Regression test for SWDEV-291202 - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV4R1Xdlops - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 128 24 14 14 --weights 64 24 5 5 --pads_strides_dilations 2 2 1 1 1 1 --disable-forward --disable-backward-weights -) - -add_custom_test(test_regression_issue_1206 SKIP_UNLESS_ALL GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED - # Regression test for SWDEV-305815 (issue 1206) - ENVIRONMENT ${IMPLICITGEMM_TESTING_ENV} MIOPEN_LOG_LEVEL=5 - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 32 256 38 38 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights -) - add_custom_test(test_regression_issue_2624 SKIP_UNLESS_ALL GFX900_DISABLED GFX906_DISABLED # Regression test for SWDEV-433714 (issue 2624) ENVIRONMENT ${IMPLICITGEMM_TESTING_ENV} MIOPEN_LOG_LEVEL=5 @@ -1661,33 +1551,6 @@ set(CONV_CK_IGEMM_FWD_V6R1_DLOPS_NCHW_ENV MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvCkIgemmFwdV6r1DlopsNchw MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_V6R1_DLOPS_NCHW=1) -# gfx908 disabled as a workaround for https://github.com/ROCm/MIOpen/pull/1790/files?diff=split&w=1#r982923610 -add_custom_test(test_conv_ck_igemm_fwd_v6r1_dlops_nchw FLOAT_ENABLED HALF_ENABLED BF16_DISABLED GFX908_DISABLED GFX103X_ENABLED SKIP_UNLESS_ALL - ENVIRONMENT ${CONV_CK_IGEMM_FWD_V6R1_DLOPS_NCHW_ENV} - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 128 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 128 256 14 14 --weights 256 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 128 1024 14 14 --weights 512 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 128 128 28 28 --weights 128 1024 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 128 128 28 28 --weights 512 128 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 58 58 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 128 2048 7 7 --weights 512 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 256 14 14 --weights 1024 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 256 14 14 --weights 256 256 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 128 256 30 30 --weights 256 256 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 56 56 --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 128 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 128 256 56 56 --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 512 16 16 --weights 512 512 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 512 28 28 --weights 1024 512 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 128 512 28 28 --weights 128 512 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 512 28 28 --weights 256 512 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 512 7 7 --weights 2048 512 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 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 128 64 56 56 --weights 256 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 128 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 128 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 --disable-backward-data --disable-backward-weights -) - add_custom_test(test_reduce_custom_fp32 GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED SKIP_UNLESS_ALL COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --scales 1 0 --CompType 1 --D 1024 30528 1 --I 0 --N 1 ---ReduceOp 0 --R 0 1 2 ${MIOPEN_TEST_FLAGS_ARGS} ) @@ -1716,415 +1579,21 @@ COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --scales 1 0 -- COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --scales 1 0 --CompType 1 --D 64 3 280 81 --I 0 --N 0 --ReduceOp 0 --R 0 ${MIOPEN_TEST_FLAGS_ARGS} ) -if(MIOPEN_TEST_DEEPBENCH) - add_custom_test(test_deepbench_conv GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED - COMMAND $ --verbose --input 4 1 161 700 --weights 32 1 5 20 --pads_strides_dilations 0 0 2 2 1 1 - COMMAND $ --verbose --input 8 1 161 700 --weights 32 1 5 20 --pads_strides_dilations 0 0 2 2 1 1 - COMMAND $ --verbose --input 16 1 161 700 --weights 32 1 5 20 --pads_strides_dilations 0 0 2 2 1 1 - COMMAND $ --verbose --input 32 1 161 700 --weights 32 1 5 20 --pads_strides_dilations 0 0 2 2 1 1 - COMMAND $ --verbose --input 4 32 79 341 --weights 32 32 5 10 --pads_strides_dilations 0 0 2 2 1 1 - COMMAND $ --verbose --input 8 32 79 341 --weights 32 32 5 10 --pads_strides_dilations 0 0 2 2 1 1 - COMMAND $ --verbose --input 16 32 79 341 --weights 32 32 5 10 --pads_strides_dilations 0 0 2 2 1 1 - COMMAND $ --verbose --input 32 32 79 341 --weights 32 32 5 10 --pads_strides_dilations 0 0 2 2 1 1 - COMMAND $ --verbose --input 16 1 48 480 --weights 16 1 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 16 16 24 240 --weights 32 16 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 16 32 12 120 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 16 64 6 60 --weights 128 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 8 3 108 108 --weights 64 3 3 3 --pads_strides_dilations 1 1 2 2 1 1 - COMMAND $ --verbose --input 8 64 54 54 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 8 128 27 27 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 8 128 14 14 --weights 256 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 8 256 7 7 --weights 512 256 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 8 3 224 224 --weights 64 3 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 8 64 112 112 --weights 128 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 8 128 56 56 --weights 256 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 8 256 28 28 --weights 512 256 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 8 512 14 14 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 8 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 16 3 224 224 --weights 64 3 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 16 64 112 112 --weights 128 64 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 16 128 56 56 --weights 256 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 16 256 28 28 --weights 512 256 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 16 512 14 14 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 16 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 16 3 224 224 --weights 64 3 7 7 --pads_strides_dilations 3 3 2 2 1 1 - COMMAND $ --verbose --input 16 192 28 28 --weights 32 192 5 5 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 16 512 14 14 --weights 48 512 5 5 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 16 832 7 7 --weights 128 832 5 5 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 16 192 28 28 --weights 32 192 1 1 --pads_strides_dilations 0 0 1 1 1 1 - COMMAND $ --verbose --input 16 512 14 14 --weights 48 512 1 1 --pads_strides_dilations 0 0 1 1 1 1 - COMMAND $ --verbose --input 16 832 7 7 --weights 128 832 1 1 --pads_strides_dilations 0 0 1 1 1 1 - ) -endif() - -if(MIOPEN_TEST_CONV) - add_custom_test(test_miopen_conv GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED - COMMAND $ --verbose --input 1 3 32 32 --weights 1 3 7 7 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 1 3 227 227 --weights 1 3 7 7 --pads_strides_dilations 1 1 1 1 1 1 - COMMAND $ --verbose --input 1 64 56 56 --weights 1 64 1 1 --pads_strides_dilations 0 0 2 2 1 1 - COMMAND $ --verbose --input 1 3 32 32 --weights 1 3 3 3 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 1 3 224 224 --weights 1 3 3 3 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 1 3 227 227 --weights 1 3 3 3 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 1 3 231 231 --weights 1 3 3 3 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 1 3 224 224 --weights 1 3 5 5 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 1 3 227 227 --weights 1 3 5 5 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 1 3 231 231 --weights 1 3 5 5 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 1 3 32 32 --weights 1 3 7 7 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 1 3 224 224 --weights 1 3 7 7 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 1 3 227 227 --weights 1 3 7 7 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 1 3 231 231 --weights 1 3 7 7 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 1 64 56 56 --weights 1 64 3 3 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 1 64 112 112 --weights 1 64 3 3 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 1 64 512 1024 --weights 1 64 3 3 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 1 96 27 27 --weights 1 96 3 3 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 1 96 28 28 --weights 1 96 3 3 --pads_strides_dilations 2 2 1 1 1 1 - COMMAND $ --verbose --input 1 3 32 32 --weights 1 3 3 3 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 3 224 224 --weights 1 3 3 3 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 3 227 227 --weights 1 3 3 3 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 3 231 231 --weights 1 3 3 3 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 3 32 32 --weights 1 3 5 5 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 3 224 224 --weights 1 3 5 5 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 3 227 227 --weights 1 3 5 5 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 3 231 231 --weights 1 3 5 5 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 3 32 32 --weights 1 3 7 7 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 3 224 224 --weights 1 3 7 7 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 3 227 227 --weights 1 3 7 7 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 3 231 231 --weights 1 3 7 7 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 16 14 14 --weights 1 16 5 5 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 16 28 28 --weights 1 16 5 5 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 24 14 14 --weights 1 24 5 5 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 32 7 7 --weights 1 32 5 5 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 32 8 8 --weights 1 32 5 5 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 32 14 14 --weights 1 32 5 5 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 32 16 16 --weights 1 32 5 5 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 32 28 28 --weights 1 32 5 5 --pads_strides_dilations 0 0 4 4 1 1 - COMMAND $ --verbose --input 1 48 7 7 --weights 1 48 5 5 --pads_strides_dilations 0 0 4 4 1 1 - ) -endif() - if(MIOPEN_TEST_FLOAT) add_custom_test(test_reduce_double SKIP_UNLESS_ALL GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED COMMAND $ --double --all --verbose) endif() -add_custom_test(smoke_solver_ConvFFT GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=fft - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvDirectNaiveConv_Fwd GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED HALF_ENABLED BF16_ENABLED INT8_ENABLED - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DRIVER_USE_GPU_REFERENCE=0 MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvFwd - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvDirectNaiveConv_Bwd GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED HALF_ENABLED BF16_ENABLED - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DRIVER_USE_GPU_REFERENCE=0 MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvBwd - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvDirectNaiveConv_Wrw GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED HALF_ENABLED BF16_ENABLED - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DRIVER_USE_GPU_REFERENCE=0 MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvWrw - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsm_5x10u2v2f1 GFX90A_DISABLED SKIP_XNACK_ON - # GFX90A_DISABLED is because of WORKAROUND_ISSUE_1146 - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm5x10u2v2f1 - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 1 1 5 10 --weights 16 1 5 10 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsm_5x10u2v2b1 GFX90A_DISABLED SKIP_XNACK_ON - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm5x10u2v2b1 - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 1 1 16 160 --weights 16 16 5 10 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsm_7x7c3h224w224k64u2v2p3q3f1 GFX90A_DISABLED SKIP_XNACK_ON - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm7x7c3h224w224k64u2v2p3q3f1 - $ ${TEST_CONV_VERBOSE_F} --input 1 3 224 224 --weights 64 3 7 7 --pads_strides_dilations 3 3 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvOcl_Fwd11x11 GFX103X_ENABLED HALF_ENABLED BF16_ENABLED - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclDirectFwd11x11 - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 1 1 44 44 --weights 1 1 11 11 --pads_strides_dilations 0 0 4 4 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvOcl_FwdGen GFX103X_ENABLED HALF_ENABLED BF16_ENABLED - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclDirectFwdGen - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 1 1 6 6 --weights 1 1 3 3 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvOcl_BwdWrW53 GFX103X_ENABLED HALF_ENABLED BF16_ENABLED - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclBwdWrW53 - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 16 1 7 7 --weights 1 1 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - # NOTES ON WRITING TESTS FOR TUNABLE SOLVERS # * Enforce tuning (SEARCH_DB_UPDATE). # * Use TEST_TUNING. This flag leads to test failure in case of any "Error" # message output to the log, which happens if something is broken in the tuning machinery. # * Use MIOPEN_DEBUG_TUNING_ITERATIONS_MAX to save testing time. -# FP16 ALT attribute is disabled to enable the backward solver on MI200 for HALF. -add_custom_test(smoke_solver_ConvAsm1x1U HALF_ENABLED SKIP_XNACK_ON TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1U - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsm1x1UV2 SKIP_XNACK_ON TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm1x1UV2 - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 1 4 2 2 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsm3x3U SKIP_XNACK_ON TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsm3x3U - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 1 4 10 10 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 1 4 10 10 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsmBwdWrW1x1 HALF_ENABLED BF16_ENABLED SKIP_XNACK_ON TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmBwdWrW1x1 - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 1 4 5 5 --weights 4 4 1 1 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -# GFX90A_DISABLED for FP32 because of WORKAROUND_SWDEV_330460 -add_custom_test(smoke_solver_ConvAsmBwdWrW3x3_fp32 GFX90A_DISABLED SKIP_XNACK_ON TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmBwdWrW3x3 - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 2 4 3 3 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsmBwdWrW3x3_fp16 FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmBwdWrW3x3 - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 2 4 3 3 --weights 4 4 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -# GFX103X_DISABLED is due to WORKAROUND_SWDEV_266868 -add_custom_test(smoke_solver_ConvOclBwdWrW1x1 GFX103X_DISABLED HALF_ENABLED BF16_ENABLED - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclBwdWrW1x1 - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 1 16 14 14 --weights 16 16 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsmImplicitGemmV4R1Dynamic_Fwd GFX908_DISABLED GFX90A_DISABLED SKIP_XNACK_ON - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicFwd - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 16 16 16 16 --weights 16 16 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsmImplicitGemmV4R1Dynamic_Bwd GFX908_DISABLED GFX90A_DISABLED SKIP_XNACK_ON - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicBwd - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 64 64 14 14 --weights 16 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsmImplicitGemmV4R1Dynamic_Wrw GFX908_DISABLED GFX90A_DISABLED SKIP_XNACK_ON - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicWrw - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 1 32 28 28 --weights 32 32 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops_Wrw GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED HALF_ENABLED SKIP_XNACK_ON - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicWrwXdlops - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 2 256 12 18 --weights 256 256 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops_Bwd GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED HALF_ENABLED SKIP_XNACK_ON - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlops - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 64 64 28 28 --weights 16 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops_Fwd GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED HALF_ENABLED SKIP_XNACK_ON - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlops - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 64 512 7 7 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_fp32_fp16_Fwd GFX900_DISABLED GFX906_DISABLED GFX94X_ENABLED - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_fp32_fp16_Bwd GFX900_DISABLED GFX906_DISABLED GFX94X_ENABLED - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_fp32_fp16_Wrw GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED GFX94X_ENABLED - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16_Fwd GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED GFX94X_ENABLED - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16_Bwd GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED FLOAT_DISABLED BF16_ENABLED SKIP_XNACK_ON TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16_Wrw GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED FLOAT_DISABLED BF16_ENABLED SKIP_XNACK_ON TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicWrwXdlopsNHWC - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 64 256 7 7 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED GFX90A_DISABLED GFX103X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 64 256 7 7 --weights 256 3 3 128 --pads_strides_dilations 0 0 1 1 1 1 --in_layout NCHW --fil_layout CHWN --out_layout NCHW --tensor_vect 1 --vector_length 4 ${MIOPEN_TEST_FLAGS_ARGS} -) - -# MIOPEN_DEBUG_TUNING_ITERATIONS_MAX is set to 2 because kernels are very slow to build. -# MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_V6R1_DLOPS_NCHW is explicitly enabled due to the kernel is disabled by default via #2306 -add_custom_test(smoke_solver_ConvCkIgemmFwdV6r1DlopsNchw GFX103X_ENABLED HALF_ENABLED TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=2 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvCkIgemmFwdV6r1DlopsNchw MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_V6R1_DLOPS_NCHW=1 - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 128 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvHipImplicitGemmBwdDataV1R1 GFX103X_ENABLED TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV1R1 - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 32 128 32 32 --weights 12 128 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -# MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1=1 is necessary due to WORKAROUND_iGemm_936 in Jenkinsfile, -# which disables ConvHipImplicitGemmV4R1Fwd, but we still want to check that the solver is not broken. -add_custom_test(smoke_solver_ConvHipImplicitGemmV4R1Fwd_fp32 GFX103X_ENABLED TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1=1 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R1Fwd - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 256 32 27 27 --weights 128 32 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvHipImplicitGemmV4R1WrW GFX103X_ENABLED HALF_ENABLED BF16_ENABLED TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R1WrW - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 64 64 55 55 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -# MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1=1 is necessary due to WORKAROUND_iGemm_936 in Jenkinsfile, -# which disables ConvHipImplicitGemmV4R1Fwd, but we still want to check that the solver is not broken. -# smoke_solver_ConvHipImplicitGemmV4R1Fwd is split to BF16+FP16 and FP32 tests because of -# WORKAROUND_ISSUE_2038, which disables validation of FP16 and BF16 datatypes in this test, -# see https://github.com/ROCm/MIOpen/pull/2043#issuecomment-1482657160. -add_custom_test(smoke_solver_ConvHipImplicitGemmV4R1Fwd_fp16_bf16 GFX103X_ENABLED FLOAT_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1=1 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R1Fwd - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 256 32 27 27 --weights 128 32 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} --disable-validation -) - -# MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1=1 is necessary due to WORKAROUND_SWDEV_229277_227616_229195, -# which disables ConvHipImplicitGemmBwdDataV4R1, but we still want to check that the solver is not broken. -add_custom_test(smoke_solver_ConvHipImplicitGemmBwdDataV4R1 GFX103X_ENABLED TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1=1 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV4R1 - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 16 64 16 16 --weights 64 64 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvHipImplicitGemmV4R4_Fwd GFX103X_ENABLED TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R4Fwd - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 2 16 28 28 --weights 32 16 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvHipImplicitGemmV4R4_Wrw GFX103X_ENABLED TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmV4R4WrW - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 8 128 14 14 --weights 32 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -# WORKAROUND_SWDEV_251757 disables this solver due to precision issues. -# However we still want to check that solver is not broken and therefore use -# MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1_XDLOPS=1 to enable it. -add_custom_test(smoke_solver_ConvHipImplicitGemmBwdDataV1R1Xdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V1R1_XDLOPS=1 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV1R1Xdlops - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 32 128 32 32 --weights 12 128 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -# WORKAROUND_ISSUE_1206 disables this solver for FP32 due to precision issues. -# WORKAROUND_SWDEV_329642 disables this solver on MI200 for BF16. -# However we still want to check that these cases are not broken and therefore use -# MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1_XDLOPS=1 to enable the solver. -add_custom_test(smoke_solver_ConvHipImplicitGemmBwdDataV4R1Xdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_BWD_V4R1_XDLOPS=1 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmBwdDataV4R1Xdlops - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 1 160 28 28 --weights 128 160 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvHipImplicitGemmForwardV4R4Xdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmForwardV4R4Xdlops - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 128 48 13 13 --weights 192 48 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvHipImplicitGemmWrwV4R4Xdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmWrwV4R4Xdlops - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 1 192 28 28 --weights 16 192 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmForwardV4R4Xdlops_Padded_Gemm - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 16 1 7 7 --weights 1 1 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmWrwV4R4Xdlops_Padded_Gemm - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 256 2 5 5 --weights 1 2 3 3 --pads_strides_dilations 1 1 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvHipImplicitGemmForwardV4R5Xdlops GFX900_DISABLED GFX906_DISABLED HALF_ENABLED BF16_ENABLED TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmForwardV4R5Xdlops - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 128 16 54 54 --weights 64 16 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - add_custom_test(smoke_solver_ConvHipImplicitGemmFwdXdlops GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED GFX94X_ENABLED HALF_ENABLED INT8_ENABLED ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmFwdXdlops COMMAND $ ${TEST_CONV_VERBOSE_F} --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_CONV_INT8_OUTPUT_TYPE_INT8} --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_solver_ConvBinWinograd3x3U GFX90A_DISABLED SKIP_XNACK_ON - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinograd3x3U - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -# F16 is supported for 906 and 906 only, no WrW -add_custom_test(smoke_solver_ConvBinWinogradRxS_fp16 GFX900_DISABLED GFX90A_DISABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 1 40 20 20 --weights 20 40 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 1 20 20 20 --weights 40 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -# F32 is supported for 900, 906 and 908. -add_custom_test(smoke_solver_ConvBinWinogradRxS_fp32 GFX90A_DISABLED SKIP_XNACK_ON - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxS - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -# FP16 ALT attribute is disabled to enable the backward solver on MI200 for HALF. -add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3g1_f16 GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON - ENVIRONMENT MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 1 40 20 20 --weights 20 40 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 1 20 20 20 --weights 40 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3g1_f32 GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED SKIP_XNACK_ON - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -# FP16 ALT attribute is disabled to enable the backward solver on MI200 for HALF. -add_custom_test(smoke_solver_ConvBinWinogradRxSf3x2_f16 GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON - ENVIRONMENT MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 1 40 20 20 --weights 20 40 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 1 20 20 20 --weights 40 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvBinWinogradRxSf3x2_f32 GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED SKIP_XNACK_ON - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND $ ${TEST_CONV_VERBOSE_W} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvWinogradFuryRxSf2x3_f16 GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED GFX90A_DISABLED GFX110X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER='ConvWinoFuryRxS<2-3>' - COMMAND $ ${TEST_CONV_VERBOSE_F} --input 1 16 16 16 --weights 16 16 3 3 --pads_strides_dilations 1 1 1 1 1 1 --trans_output_pads 1 1 ${MIOPEN_TEST_FLAGS_ARGS} - COMMAND $ ${TEST_CONV_VERBOSE_B} --input 1 16 16 16 --weights 16 16 3 3 --pads_strides_dilations 1 1 1 1 1 1 --trans_output_pads 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - # FP16 ALT attribute is disabled to enable the backward solver on MI200 for HALF. add_custom_test(smoke_solver_ConvWinograd3x3MultipassWrW_3x2 HALF_ENABLED BF16_ENABLED SKIP_XNACK_ON ENVIRONMENT MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER='ConvWinograd3x3MultipassWrW<3-2>' @@ -2191,21 +1660,6 @@ add_custom_test(smoke_solver_ConvWinograd3x3MultipassWrW1x1x7x3 HALF_ENABLED BF1 COMMAND $ ${TEST_CONV_VERBOSE_W} --input 1 16 24 24 --weights 16 16 1 7 --pads_strides_dilations 0 3 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3 GFX900_DISABLED GFX94X_ENABLED GFX103X_ENABLED HALF_ENABLED SKIP_XNACK_ON TEST_TUNING - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3 - COMMAND $ --input 1 40 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 2 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3g1 GFX900_DISABLED GFX94X_ENABLED GFX103X_ENABLED HALF_ENABLED SKIP_XNACK_ON - ENVIRONMENT MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 - COMMAND $ --input 1 40 20 20 --weights 20 40 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - -add_custom_test(smoke_solver_ConvBinWinogradRxSf3x2 GFX900_DISABLED GFX94X_ENABLED GFX103X_ENABLED HALF_ENABLED SKIP_XNACK_ON - ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 - COMMAND $ --input 1 40 20 20 --weights 20 40 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} -) - add_custom_test(smoke_solver_ConvMlirIgemm_F GFX900_DISABLED GFX908_DISABLED GFX90A_DISABLED GFX103X_ENABLED HALF_ENABLED SKIP_UNLESS_MLIR TEST_TUNING ENVIRONMENT MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 ${IMPLICITGEMM_MLIR_ENV_F} COMMAND $ ${TEST_CONV_VERBOSE_F} --input 64 128 14 14 --weights 128 128 1 1 --pads_strides_dilations 0 0 2 2 1 1 --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} @@ -2221,14 +1675,6 @@ add_custom_test(smoke_solver_ConvMlirIgemm_W GFX900_DISABLED GFX908_DISABLED GFX COMMAND $ ${TEST_CONV_VERBOSE_W} --input 64 64 28 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -# Add here regression tests that should be run on Vega10/20 and GFX908 only with FP16. -add_custom_test(test_regression_half_vega_gfx908 FLOAT_DISABLED HALF_ENABLED GFX90A_DISABLED -# Issue #894. -# Can't be enabled for GFX10 due to WORKAROUND_SWDEV_271887 - ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvOclDirectFwd1x1 - COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --disable-backward-data --disable-backward-weights --disable-verification-cache --cmode conv --pmode default --group-count 1 --input 1 16 7 7 --weights 16 16 1 1 --pads_strides_dilations 0 0 1 1 1 1 -) - add_custom_test(test_regression_half_vega SKIP_UNLESS_ALL FLOAT_DISABLED HALF_ENABLED GFX908_DISABLED GFX90A_DISABLED # Issue #1956. ENVIRONMENT MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER='GemmBwdRest' diff --git a/test/gtest/CMakeLists.txt b/test/gtest/CMakeLists.txt index 415318a97c..69346a1c23 100644 --- a/test/gtest/CMakeLists.txt +++ b/test/gtest/CMakeLists.txt @@ -29,8 +29,7 @@ function(add_gtest TEST_NAME TEST_CPP) endif() if(NOT WIN32) # TODO: cannot run on Windows due to missing DLL dependencies # Enable CMake to discover the test binary - # Extend GTest DISCOVERY_TIMEOUT to 5 mins - gtest_discover_tests(${TEST_NAME} DISCOVERY_TIMEOUT 300 PROPERTIES ENVIRONMENT "MIOPEN_USER_DB_PATH=${CMAKE_CURRENT_BINARY_DIR};MIOPEN_TEST_FLOAT_ARG=${MIOPEN_TEST_FLOAT_ARG};MIOPEN_TEST_ALL=${MIOPEN_TEST_ALL};MIOPEN_TEST_MLIR=${MIOPEN_TEST_MLIR};MIOPEN_TEST_COMPOSABLEKERNEL=${MIOPEN_TEST_COMPOSABLEKERNEL};CODECOV_TEST=${CODECOV_TEST};MIOPEN_TEST_DBSYNC=${MIOPEN_TEST_DBSYNC}") + gtest_discover_tests(${TEST_NAME} DISCOVERY_TIMEOUT 300 PROPERTIES ENVIRONMENT "MIOPEN_USER_DB_PATH=${CMAKE_CURRENT_BINARY_DIR};MIOPEN_TEST_FLOAT_ARG=${MIOPEN_TEST_FLOAT_ARG};MIOPEN_TEST_ALL=${MIOPEN_TEST_ALL};MIOPEN_TEST_MLIR=${MIOPEN_TEST_MLIR};MIOPEN_TEST_COMPOSABLEKERNEL=${MIOPEN_TEST_COMPOSABLEKERNEL};CODECOV_TEST=${CODECOV_TEST};MIOPEN_TEST_DBSYNC=${MIOPEN_TEST_DBSYNC};MIOPEN_TEST_CONV=${MIOPEN_TEST_CONV};MIOPEN_TEST_DEEPBENCH=${MIOPEN_TEST_DEEPBENCH};MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=${MIOPEN_DEBUG_TUNING_ITERATIONS_MAX}") endif() target_link_libraries(${TEST_NAME} BZip2::BZip2) if(WIN32) diff --git a/test/gtest/conv_ck_igemm_fwd_v6r1_dlops_nchw.cpp b/test/gtest/conv_ck_igemm_fwd_v6r1_dlops_nchw.cpp new file mode 100644 index 0000000000..24bb7b925d --- /dev/null +++ b/test/gtest/conv_ck_igemm_fwd_v6r1_dlops_nchw.cpp @@ -0,0 +1,130 @@ +/******************************************************************************* + * + * 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 "gtest_common.hpp" + +#include "../conv2d.hpp" + +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL) + +namespace conv_ck_igemm_fwd_v6r1_dlops_nchw { + +auto GetTestCases() +{ + const auto env = std::tuple{ + std::pair{ENV(MIOPEN_FIND_MODE), std::string_view("normal")}, + std::pair{ENV(MIOPEN_DEBUG_FIND_ONLY_SOLVER), + std::string_view("ConvCkIgemmFwdV6r1DlopsNchw")}, + std::pair{ENV(MIOPEN_DEBUG_CONV_CK_IGEMM_FWD_V6R1_DLOPS_NCHW), std::string_view("1")}}; + + const std::string v = " --verbose"; + const std::string dis_bk_data = " --disable-backward-data"; + const std::string dis_bk_wei = " --disable-backward-weights"; + + return std::vector{ + // clang-format off + std::pair{env, v + " --input 128 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, v + " --input 128 256 14 14 --weights 256 1024 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env, v + " --input 128 1024 14 14 --weights 512 1024 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env, v + " --input 128 128 28 28 --weights 128 1024 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env, v + " --input 128 128 28 28 --weights 512 128 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env, v + " --input 128 128 58 58 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env, v + " --input 128 2048 7 7 --weights 512 2048 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env, v + " --input 128 256 14 14 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env, v + " --input 128 256 14 14 --weights 256 256 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env, v + " --input 128 256 30 30 --weights 256 256 3 3 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env, v + " --input 128 256 56 56 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env, v + " --input 128 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, v + " --input 128 256 56 56 --weights 64 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env, v + " --input 128 512 16 16 --weights 512 512 3 3 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env, v + " --input 128 512 28 28 --weights 1024 512 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env, v + " --input 128 512 28 28 --weights 128 512 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env, v + " --input 128 512 28 28 --weights 256 512 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env, v + " --input 128 512 7 7 --weights 2048 512 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env, v + " --input 128 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, v + " --input 128 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env, v + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env, v + " --input 128 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_bk_data + dis_bk_wei} + // clang-format on + }; +} + +using TestCase = decltype(GetTestCases())::value_type; + +bool SkipTest() { return miopen::IsDisabled(ENV(MIOPEN_TEST_ALL)); } +class Conv2dFloat_conv_ck_igemm_fwd_v6r1_dlops_nchw : public FloatTestCase> +{ +}; + +class Conv2dHalf_conv_ck_igemm_fwd_v6r1_dlops_nchw : public HalfTestCase> +{ +}; + +bool IsTestSupportedForDevice() +{ + using e_mask = enabled; + using d_mask = disabled; + return ::IsTestSupportedForDevMask(); +} + +} // namespace conv_ck_igemm_fwd_v6r1_dlops_nchw +using namespace conv_ck_igemm_fwd_v6r1_dlops_nchw; + +TEST_P(Conv2dFloat_conv_ck_igemm_fwd_v6r1_dlops_nchw, FloatTest) +{ + if(IsTestSupportedForDevice() && !SkipTest()) + { + invoke_with_params( + default_check); + } + else + { + GTEST_SKIP(); + } +}; + +TEST_P(Conv2dHalf_conv_ck_igemm_fwd_v6r1_dlops_nchw, HalfTest) +{ + if(IsTestSupportedForDevice() && !SkipTest()) + { + invoke_with_params( + default_check); + } + else + { + GTEST_SKIP(); + } +}; + +INSTANTIATE_TEST_SUITE_P(ConvCkIgemmFwdV6r1DlopsNchw, + Conv2dFloat_conv_ck_igemm_fwd_v6r1_dlops_nchw, + testing::Values(GetTestCases())); + +INSTANTIATE_TEST_SUITE_P(ConvCkIgemmFwdV6r1DlopsNchw, + Conv2dHalf_conv_ck_igemm_fwd_v6r1_dlops_nchw, + testing::Values(GetTestCases())); diff --git a/test/gtest/conv_igemm_dynamic.cpp b/test/gtest/conv_igemm_dynamic.cpp index 676ce1c352..8bc9390e4e 100644 --- a/test/gtest/conv_igemm_dynamic.cpp +++ b/test/gtest/conv_igemm_dynamic.cpp @@ -24,156 +24,104 @@ * *******************************************************************************/ #include +#include -#include -#include -#include -#include "get_handle.hpp" -#include "test_env.hpp" +#include "gtest_common.hpp" #include "../conv2d.hpp" -using TestCase = std::tuple, std::string>; - +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL) MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_GPU_XNACK_ENABLED) namespace conv_igemm_dynamic { -static bool SkipTest(void) { return miopen::IsEnabled(ENV(MIOPEN_TEST_GPU_XNACK_ENABLED)); } - -void GetArgs(const TestCase& param, std::vector& tokens) +auto GetTestCases() { - auto env_vars = std::get<0>(param); - for(auto& elem : env_vars) + const auto env = std::tuple{std::pair{ENV(MIOPEN_FIND_MODE), std::string_view("normal")}, + std::pair{ENV(MIOPEN_DEBUG_FIND_ONLY_SOLVER), + std::string_view("ConvAsmImplicitGemmV4R1DynamicFwd")}}; + const auto env_1x1 = + std::tuple{std::pair{ENV(MIOPEN_FIND_MODE), std::string_view("normal")}, + std::pair{ENV(MIOPEN_DEBUG_FIND_ONLY_SOLVER), + std::string_view("ConvAsmImplicitGemmV4R1DynamicFwd_1x1")}}; + const auto env_wrw = + std::tuple{std::pair{ENV(MIOPEN_FIND_MODE), std::string_view("normal")}, + std::pair{ENV(MIOPEN_DEBUG_FIND_ONLY_SOLVER), + std::string_view("ConvAsmImplicitGemmV4R1DynamicWrw")}}; + const auto env_bwd = + std::tuple{std::pair{ENV(MIOPEN_FIND_MODE), std::string_view("normal")}, + std::pair{ENV(MIOPEN_DEBUG_FIND_ONLY_SOLVER), + std::string_view("ConvAsmImplicitGemmV4R1DynamicBwd")}}; + + const std::string v = " --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"; + + auto basic_tests = std::vector { - putenv(elem.data()); - } - - auto cmd = std::get<1>(param); - - std::stringstream ss(cmd); - std::istream_iterator begin(ss); - std::istream_iterator end; - while(begin != end) - tokens.push_back(*begin++); -} - -class Conv2dFloatDynamic : public testing::TestWithParam> -{ -}; - -void Run2dDriver(miopenDataType_t prec) -{ + // clang-format off +#if CODECOV_TEST + std::pair{env , v + " --input 32 32 17 17 --weights 32 32 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_bk_data + dis_bk_wei + dis_vali}, + std::pair{env_wrw, 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 + dis_vali}, + std::pair{env_bwd, 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 + dis_vali}, +#else + std::pair{env , v + " --input 16 16 56 56 --weights 64 16 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env , v + " --input 16 64 34 34 --weights 64 64 3 3 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env , v + " --input 32 32 17 17 --weights 32 32 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env_1x1, v + " --input 16 384 8 8 --weights 64 384 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env_wrw, 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_wrw, 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_bwd, 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_bwd, 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} +#endif + // clang-format on + }; - std::vector params; - switch(prec) + if(miopen::IsEnabled(ENV(MIOPEN_TEST_ALL))) { - case miopenFloat: params = Conv2dFloatDynamic::GetParam(); break; - case miopenHalf: - case miopenInt8: - case miopenBFloat16: - case miopenInt32: - case miopenDouble: - case miopenFloat8: - case miopenBFloat8: - FAIL() << "miopenHalf, miopenInt8, miopenBFloat16, miopenInt32, " - "miopenDouble, miopenFloat8, miopenBFloat8 " - "data type not supported by conv_igemm_dynamic test"; - - default: params = Conv2dFloatDynamic::GetParam(); + basic_tests.insert(basic_tests.end(), + { + // clang-format off + std::pair{env , v + " --input 64 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env , v + " --input 64 256 34 34 --weights 256 256 3 3 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env , v + " --input 128 128 35 35 --weights 128 128 3 3 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env , v + " --input 64 1536 8 8 --weights 256 1536 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env , v + " --input 128 48 7 7 --weights 128 48 5 5 --pads_strides_dilations 2 2 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env , 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_1x1, v + " --input 128 256 28 28 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env_1x1, v + " --input 64 1536 8 8 --weights 256 1536 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env_1x1, v + " --input 128 768 17 17 --weights 128 768 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, + std::pair{env_wrw, 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_wrw, v + " --input 32 128 34 34 --weights 64 128 3 3 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + std::pair{env_wrw, v + " --input 128 128 35 35 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_data}, + std::pair{env_wrw, v + " --input 128 256 56 56 --weights 64 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + std::pair{env_wrw, v + " --input 64 512 28 28 --weights 256 512 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_data}, + std::pair{env_wrw, v + " --input 64 512 14 14 --weights 256 512 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, + std::pair{env_bwd, 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_bwd, v + " --input 32 128 34 34 --weights 64 128 3 3 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, + std::pair{env_bwd, v + " --input 128 128 35 35 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_wei}, + std::pair{env_bwd, v + " --input 128 256 56 56 --weights 64 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei} + // clang-format on + }); } + return basic_tests; +} - 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(); }); +using TestCase = decltype(GetTestCases())::value_type; - testing::internal::CaptureStderr(); - test_drive(ptrs.size(), ptrs.data()); - auto capture = testing::internal::GetCapturedStderr(); - std::cout << capture; - } -}; +bool SkipTest() { return miopen::IsEnabled(ENV(MIOPEN_TEST_GPU_XNACK_ENABLED)); } -bool IsTestSupportedForDevice(const miopen::Handle& handle) +class Conv2dFloatDynamic : public FloatTestCase> { - std::string devName = handle.GetDeviceName(); - if(devName == "gfx900" || devName == "gfx906") - return true; - else - return false; -} +}; -std::vector GetTestCases(const std::string& precision) +bool IsTestSupportedForDevice() { - - std::vector env = { - "MIOPEN_FIND_MODE=normal", - "MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicFwd"}; - std::vector env_1x1 = { - "MIOPEN_FIND_MODE=normal", - "MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicFwd_1x1"}; - std::vector env_wrw = { - "MIOPEN_FIND_MODE=normal", - "MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicWrw"}; - std::vector env_bwd = { - "MIOPEN_FIND_MODE=normal", - "MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmV4R1DynamicBwd"}; - - std::string v = " --verbose"; - std::string dis_bk_data = " --disable-backward-data"; - std::string dis_bk_wei = " --disable-backward-weights"; - std::string dis_fwd = " --disable-forward"; - std::string dis_vali = " --disable-validation"; - - const std::vector test_cases = { - // clang-format off -#if CODECOV_TEST - TestCase{env, precision + v + " --input 32 32 17 17 --weights 32 32 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_bk_data + dis_bk_wei + dis_vali}, - TestCase{env_wrw, precision + 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 + dis_vali}, - TestCase{env_bwd, precision + 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 + dis_vali}, -#else - TestCase{env, precision + v + " --input 16 16 56 56 --weights 64 16 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, - TestCase{env, precision + v + " --input 16 64 34 34 --weights 64 64 3 3 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, - TestCase{env, precision + v + " --input 32 32 17 17 --weights 32 32 1 7 --pads_strides_dilations 0 3 1 1 1 1" + dis_bk_data + dis_bk_wei}, - TestCase{env_1x1, precision + v + " --input 16 384 8 8 --weights 64 384 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, - TestCase{env_wrw, precision + 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}, - TestCase{env_wrw, precision + 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}, - TestCase{env_bwd, precision + 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}, - TestCase{env_bwd, precision + 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}, -#endif - -#if MIOPEN_TEST_ALL - //SKIP_UNLESS_ALL - TestCase{env, precision + v + " --input 64 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, - TestCase{env, precision + v + " --input 64 256 34 34 --weights 256 256 3 3 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, - TestCase{env, precision + v + " --input 128 128 35 35 --weights 128 128 3 3 --pads_strides_dilations 0 0 2 2 1 1" + dis_bk_data + dis_bk_wei}, - TestCase{env, precision + v + " --input 64 1536 8 8 --weights 256 1536 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, - TestCase{env, precision + v + " --input 128 48 7 7 --weights 128 48 5 5 --pads_strides_dilations 2 2 1 1 1 1" + dis_bk_data + dis_bk_wei}, - TestCase{env, precision + 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}, - TestCase{env_1x1, precision + v + " --input 128 256 28 28 --weights 128 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, - TestCase{env_1x1, precision + v + " --input 64 1536 8 8 --weights 256 1536 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, - TestCase{env_1x1, precision + v + " --input 128 768 17 17 --weights 128 768 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_bk_data + dis_bk_wei}, - TestCase{env_wrw, precision + 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}, - TestCase{env_wrw, precision + v + " --input 32 128 34 34 --weights 64 128 3 3 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, - TestCase{env_wrw, precision + v + " --input 128 128 35 35 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_data}, - TestCase{env_wrw, precision + v + " --input 128 256 56 56 --weights 64 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, - TestCase{env_wrw, precision + v + " --input 64 512 28 28 --weights 256 512 1 1 --pads_strides_dilations 0 0 2 2 1 1" + dis_fwd + dis_bk_data}, - TestCase{env_wrw, precision + v + " --input 64 512 14 14 --weights 256 512 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_data}, - TestCase{env_bwd, precision + 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}, - TestCase{env_bwd, precision + v + " --input 32 128 34 34 --weights 64 128 3 3 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei}, - TestCase{env_bwd, precision + v + " --input 128 128 35 35 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + dis_fwd + dis_bk_wei}, - TestCase{env_bwd, precision + v + " --input 128 256 56 56 --weights 64 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + dis_fwd + dis_bk_wei} -#endif - // clang-format on - }; - return test_cases; + using e_mask = enabled; + using d_mask = disabled; + return ::IsTestSupportedForDevMask(); } } // namespace conv_igemm_dynamic @@ -181,10 +129,9 @@ using namespace conv_igemm_dynamic; TEST_P(Conv2dFloatDynamic, FloatTest_conv_igemm_dynamic) { - const auto& handle = get_handle(); - if(IsTestSupportedForDevice(handle) && !SkipTest() && IsTestRunWith("--float")) + if(IsTestSupportedForDevice() && !SkipTest()) { - Run2dDriver(miopenFloat); + invoke_with_params(default_check); } else { @@ -192,6 +139,4 @@ TEST_P(Conv2dFloatDynamic, FloatTest_conv_igemm_dynamic) } }; -INSTANTIATE_TEST_SUITE_P(ConvIgemmDynamic, - Conv2dFloatDynamic, - testing::Values(GetTestCases("--float"))); +INSTANTIATE_TEST_SUITE_P(ConvIgemmDynamic, Conv2dFloatDynamic, testing::Values(GetTestCases())); diff --git a/test/gtest/conv_igemm_dynamic_dlops.cpp b/test/gtest/conv_igemm_dynamic_dlops.cpp new file mode 100644 index 0000000000..98ca156236 --- /dev/null +++ b/test/gtest/conv_igemm_dynamic_dlops.cpp @@ -0,0 +1,182 @@ +/******************************************************************************* + * + * 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 "gtest_common.hpp" + +#include "../conv2d.hpp" + +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL) +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_GPU_XNACK_ENABLED) + +namespace { + +auto GetTestCases() +{ + const auto env_fwd = + std::tuple{std::pair{ENV(MIOPEN_FIND_MODE), std::string_view("normal")}, + std::pair{ENV(MIOPEN_DEBUG_FIND_ONLY_SOLVER), + std::string_view("ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC")}}; + + const std::string v = " --verbose"; + const std::string dis_bk_data = " --disable-backward-data"; + const std::string dis_bk_wei = " --disable-backward-weights"; + const std::string in_nchw = " --in_layout NCHW"; + const std::string fil_nchw = " --fil_layout NCHW"; + const std::string fil_chwn = " --fil_layout CHWN"; + const std::string out_nchw = " --out_layout NCHW"; + const std::string tensor = " --tensor_vect 1"; + const std::string vlen4 = " --vector_length 4"; + const std::string vlen8 = " --vector_length 8"; + + const std::string common_base = " --cmode convfp16" + dis_bk_data + dis_bk_wei + in_nchw; + + const std::string nchwc_nchwc_base = common_base + fil_nchw + out_nchw + tensor; + const std::string nchwc_nchwc_fwd_fp16x4 = nchwc_nchwc_base + vlen4; + const std::string nchwc_nchwc_fwd_fp16x8 = nchwc_nchwc_base + vlen8; + + const std::string nchwc_chwnc_base = common_base + fil_chwn + out_nchw + tensor; + const std::string nchwc_chwnc_fwd_fp16x4 = nchwc_chwnc_base + vlen4; + const std::string nchwc_chwnc_fwd_fp16x8 = nchwc_chwnc_base + vlen8; + + return std::vector{ + // clang-format off + //nchwc_nchwc_fwd_fp16x4 + std::pair{env_fwd, v + " --input 1 8 10 10 --weights 8 8 3 3 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 2 64 32 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1" + nchwc_nchwc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1" + nchwc_nchwc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1" + nchwc_nchwc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1" + nchwc_nchwc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1" + nchwc_nchwc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1" + nchwc_nchwc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 16 16 25 25 --weights 64 16 3 3 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 4 32 79 141 --weights 64 32 5 10 --pads_strides_dilations 0 0 2 2 1 1" + nchwc_nchwc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 400 256 7 7 --weights 1024 256 7 7 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 400 256 1 1 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x4}, + + //nchwc_chwnc_fwd_fp16x4 + std::pair{env_fwd, v + " --input 64 256 7 7 --weights 256 3 3 128 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 32 160 73 73 --weights 160 1 1 64 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 16 64 56 56 --weights 64 1 1 64 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 2 256 40 52 --weights 256 1 1 256 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 2 64 32 28 --weights 64 1 1 64 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 32 128 14 14 --weights 128 1 1 64 --pads_strides_dilations 0 0 2 2 1 1" + nchwc_chwnc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 64 64 17 17 --weights 64 3 7 192 --pads_strides_dilations 0 3 1 1 1 1" + nchwc_chwnc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 64 64 17 17 --weights 64 7 1 192 --pads_strides_dilations 3 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 4 128 28 28 --weights 128 2 2 128 --pads_strides_dilations 0 0 2 2 1 1" + nchwc_chwnc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 32 128 8 8 --weights 128 3 1 192 --pads_strides_dilations 1 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 64 192 17 17 --weights 192 3 3 160 --pads_strides_dilations 0 0 2 2 1 1" + nchwc_chwnc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 64 32 73 73 --weights 32 3 3 64 --pads_strides_dilations 1 1 1 1 1 1" + nchwc_chwnc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 16 64 56 56 --weights 64 3 3 64 --pads_strides_dilations 1 1 1 1 1 1" + nchwc_chwnc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 16 16 25 25 --weights 16 3 3 64 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 4 32 79 141 --weights 32 5 10 64 --pads_strides_dilations 0 0 2 2 1 1" + nchwc_chwnc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 400 256 7 7 --weights 256 7 7 1024 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x4}, + std::pair{env_fwd, v + " --input 400 256 1 1 --weights 256 1 1 1024 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x4}, + + //nchwc_nchwc_fwd_fp16x8 + std::pair{env_fwd, v + " --input 1 8 10 10 --weights 8 8 3 3 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 32 160 73 73 --weights 64 160 1 1 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 16 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 2 256 40 52 --weights 256 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 2 64 32 28 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 32 128 14 14 --weights 64 128 1 1 --pads_strides_dilations 0 0 2 2 1 1" + nchwc_nchwc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 64 64 17 17 --weights 192 64 1 7 --pads_strides_dilations 0 3 1 1 1 1" + nchwc_nchwc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 64 64 17 17 --weights 192 64 7 1 --pads_strides_dilations 3 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 4 128 28 28 --weights 128 128 2 2 --pads_strides_dilations 0 0 2 2 1 1" + nchwc_nchwc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 32 128 8 8 --weights 192 128 3 1 --pads_strides_dilations 1 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 64 192 17 17 --weights 160 192 3 3 --pads_strides_dilations 0 0 2 2 1 1" + nchwc_nchwc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 64 32 73 73 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1" + nchwc_nchwc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 16 64 56 56 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1" + nchwc_nchwc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 16 16 25 25 --weights 64 16 3 3 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 4 32 79 141 --weights 64 32 5 10 --pads_strides_dilations 0 0 2 2 1 1" + nchwc_nchwc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 400 256 7 7 --weights 1024 256 7 7 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 400 256 1 1 --weights 1024 256 1 1 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_nchwc_fwd_fp16x8}, + + //nchwc_chwnc_fwd_fp16x8 + std::pair{env_fwd, v + " --input 64 256 7 7 --weights 256 1 1 128 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 32 160 73 73 --weights 160 1 1 64 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 16 64 56 56 --weights 64 1 1 64 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 2 256 40 52 --weights 256 1 1 256 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 2 64 32 28 --weights 64 1 1 64 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 32 128 14 14 --weights 128 1 1 64 --pads_strides_dilations 0 0 2 2 1 1" + nchwc_chwnc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 64 64 17 17 --weights 64 1 7 192 --pads_strides_dilations 0 3 1 1 1 1" + nchwc_chwnc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 64 64 17 17 --weights 64 7 1 192 --pads_strides_dilations 3 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 4 128 28 28 --weights 128 2 2 128 --pads_strides_dilations 0 0 2 2 1 1" + nchwc_chwnc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 32 128 8 8 --weights 128 3 1 192 --pads_strides_dilations 1 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 64 192 17 17 --weights 192 3 3 160 --pads_strides_dilations 0 0 2 2 1 1" + nchwc_chwnc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 64 32 73 73 --weights 32 3 3 64 --pads_strides_dilations 1 1 1 1 1 1" + nchwc_chwnc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 16 64 56 56 --weights 64 3 3 64 --pads_strides_dilations 1 1 1 1 1 1" + nchwc_chwnc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 16 16 25 25 --weights 16 3 3 64 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 4 32 79 141 --weights 32 5 10 64 --pads_strides_dilations 0 0 2 2 1 1" + nchwc_chwnc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 400 256 7 7 --weights 256 7 7 1024 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x8}, + std::pair{env_fwd, v + " --input 400 256 1 1 --weights 256 1 1 1024 --pads_strides_dilations 0 0 1 1 1 1" + nchwc_chwnc_fwd_fp16x8} + // clang-format on + }; +} + +using TestCase = decltype(GetTestCases())::value_type; + +bool SkipTest() +{ + return miopen::IsEnabled(ENV(MIOPEN_TEST_GPU_XNACK_ENABLED)) || + miopen::IsDisabled(ENV(MIOPEN_TEST_ALL)); +} + +bool IsTestSupportedForDevice() +{ + using e_mask = enabled; + using d_mask = disabled; + return ::IsTestSupportedForDevMask(); +} + +} // namespace + +class Conv2dDefaultHalf : public HalfTestCase> +{ +}; + +TEST_P(Conv2dDefaultHalf, HalfTest_conv_igemm_dynamic_dlops) +{ + if(IsTestSupportedForDevice() && !SkipTest()) + { + invoke_with_params(default_check); + } + else + { + GTEST_SKIP(); + } +}; + +INSTANTIATE_TEST_SUITE_P(ConvIgemmDynamicDlopsFwd, + Conv2dDefaultHalf, + testing::Values(GetTestCases())); diff --git a/test/gtest/conv_igemm_dynamic_xdlops.cpp b/test/gtest/conv_igemm_dynamic_xdlops.cpp new file mode 100644 index 0000000000..98a2e0c823 --- /dev/null +++ b/test/gtest/conv_igemm_dynamic_xdlops.cpp @@ -0,0 +1,158 @@ +/******************************************************************************* + * + * 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 "gtest_common.hpp" + +#include "../conv2d.hpp" + +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL) +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_GPU_XNACK_ENABLED) + +namespace { + +auto GetTestCases() +{ + const auto env_xdlops = + std::tuple{std::pair{ENV(MIOPEN_FIND_MODE), std::string_view("normal")}, + std::pair{ENV(MIOPEN_DEBUG_FIND_ONLY_SOLVER), + std::string_view("ConvAsmImplicitGemmGTCDynamicBwdXdlops;" + "ConvAsmImplicitGemmGTCDynamicFwdXdlops;" + "ConvAsmImplicitGemmGTCDynamicWrwXdlops")}}; + + const std::string v = " --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"; + + 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}, + //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}, + //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}, + //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}, + //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}, + //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}, + //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} + // clang-format on + }; +} + +using TestCase = decltype(GetTestCases())::value_type; + +bool SkipTest() +{ + return miopen::IsEnabled(ENV(MIOPEN_TEST_GPU_XNACK_ENABLED)) || + miopen::IsDisabled(ENV(MIOPEN_TEST_ALL)); +} + +bool IsTestSupportedForDevice() +{ + using e_mask = enabled; + using d_mask = disabled; + return ::IsTestSupportedForDevMask(); +} + +} // namespace + +class Conv2dDefaultFloat : public FloatTestCase> +{ +}; + +class Conv2dDefaultHalf : public HalfTestCase> +{ +}; + +TEST_P(Conv2dDefaultFloat, FloatTest_conv_igemm_dynamic_xdlops) +{ + if(IsTestSupportedForDevice() && !SkipTest()) + { + invoke_with_params(default_check); + } + else + { + GTEST_SKIP(); + } +}; + +TEST_P(Conv2dDefaultHalf, HalfTest_conv_igemm_dynamic_xdlops) +{ + if(IsTestSupportedForDevice() && !SkipTest()) + { + invoke_with_params(default_check); + } + else + { + GTEST_SKIP(); + } +}; + +INSTANTIATE_TEST_SUITE_P(ConvIgemmDynamic, Conv2dDefaultFloat, testing::Values(GetTestCases())); +INSTANTIATE_TEST_SUITE_P(ConvIgemmDynamic, Conv2dDefaultHalf, testing::Values(GetTestCases())); diff --git a/test/gtest/conv_igemm_mlir.cpp b/test/gtest/conv_igemm_mlir.cpp deleted file mode 100644 index 3b846d6e6c..0000000000 --- a/test/gtest/conv_igemm_mlir.cpp +++ /dev/null @@ -1,253 +0,0 @@ -/******************************************************************************* - * - * 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 -#include -#include "get_handle.hpp" -#include "test_env.hpp" - -#include "../conv2d.hpp" - -MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_MLIR) -MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL) - -namespace conv_igemm_mlir { - -using TestCase = std::tuple, std::string>; - -void GetArgs(const TestCase& param, std::vector& tokens) -{ - auto env_vars = std::get<0>(param); - for(auto& elem : env_vars) - { - putenv(elem.data()); - } - - auto cmd = std::get<1>(param); - - std::stringstream ss(cmd); - std::istream_iterator begin(ss); - std::istream_iterator end; - while(begin != end) - tokens.push_back(*begin++); -} - -class ConvIgemmMlirConfigFloat : public testing::TestWithParam> -{ -}; -class ConvIgemmMlirConfigHalf : public testing::TestWithParam> -{ -}; -class ConvIgemmMlirConfigInt8 : public testing::TestWithParam> -{ -}; - -void Run2dDriver(miopenDataType_t prec) -{ - - std::vector params; - switch(prec) - { - case miopenHalf: params = ConvIgemmMlirConfigHalf::GetParam(); break; - case miopenInt8: params = ConvIgemmMlirConfigInt8::GetParam(); break; - case miopenFloat: params = ConvIgemmMlirConfigFloat::GetParam(); break; - case miopenBFloat16: - case miopenInt32: - case miopenFloat8: - case miopenBFloat8: - case miopenDouble: - MIOPEN_THROW(miopenStatusBadParm, - "miopenBFloat16, miopenInt32, miopenFloat8, miopenBFloat8, " - "miopenDouble data type not supported by conv_igemm_mlir test"); - - default: params = ConvIgemmMlirConfigFloat::GetParam(); - } - - 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 auto& str) { - return str.data(); - }); - - testing::internal::CaptureStderr(); - test_drive(ptrs.size(), ptrs.data()); - auto capture = testing::internal::GetCapturedStderr(); - EXPECT_FALSE(capture.find("Perf Db: record not found") != std::string::npos); - } -}; - -std::vector GetTestCases(const std::string& precision) -{ - std::vector igemm_fwd = {"MIOPEN_FIND_MODE=normal", - "MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvMlirIgemmFwd"}; - std::string flags_fwd = " --verbose --disable-backward-data --disable-backward-weights"; - std::string layout = " --in_layout NHWC --fil_layout NHWC --out_layout NHWC"; - std::string groupCount_4 = " --group-count 4"; - - // FWD test cases for precision == "--int8" - std::vector test_cases = { - // clang-format off - TestCase{igemm_fwd, precision + flags_fwd + " --input 256 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1"}, - TestCase{igemm_fwd, precision + flags_fwd + " --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, - TestCase{igemm_fwd, precision + flags_fwd + " --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + layout}, - TestCase{igemm_fwd, precision + flags_fwd + " --input 128 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, - TestCase{igemm_fwd, precision + flags_fwd + " --input 128 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1" + layout}, - TestCase{igemm_fwd, precision + flags_fwd + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1"}, - TestCase{igemm_fwd, precision + flags_fwd + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, - TestCase{igemm_fwd, precision + flags_fwd + " --input 256 256 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + groupCount_4} - // clang-format on - }; - - std::vector igemm_bwd = {"MIOPEN_FIND_MODE=normal", - "MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvMlirIgemmBwd"}; - std::vector igemm_wrw = {"MIOPEN_FIND_MODE=normal", - "MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvMlirIgemmWrW"}; - - std::string flags_bwd = " --verbose --disable-forward --disable-backward-weights"; - std::string flags_wrw = " --verbose --disable-forward --disable-backward-data"; - std::string groupCount_32 = " --group-count 32"; - - // BWD WRW test cases - const std::vector test_cases_bwd_wrw = { - // clang-format off - TestCase{igemm_bwd, precision + flags_bwd + " --input 256 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1"}, - TestCase{igemm_bwd, precision + flags_bwd + " --input 256 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1" + layout}, - TestCase{igemm_bwd, precision + flags_bwd + " --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, - TestCase{igemm_bwd, precision + flags_bwd + " --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + layout}, - TestCase{igemm_bwd, precision + flags_bwd + " --input 128 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, - TestCase{igemm_bwd, precision + flags_bwd + " --input 128 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1" + layout}, - TestCase{igemm_bwd, precision + flags_bwd + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1"}, - TestCase{igemm_bwd, precision + flags_bwd + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, - - TestCase{igemm_wrw, precision + flags_wrw + " --input 64 1024 14 14 --weights 256 1024 1 1 --pads_strides_dilations 0 0 1 1 1 1"}, - TestCase{igemm_wrw, precision + flags_wrw + " --input 64 1024 14 14 --weights 256 1024 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, - TestCase{igemm_wrw, precision + flags_wrw + " --input 256 256 14 14 --weights 256 256 3 3 --pads_strides_dilations 0 0 2 2 1 1"}, - TestCase{igemm_wrw, precision + flags_wrw + " --input 256 256 14 14 --weights 256 256 3 3 --pads_strides_dilations 0 0 2 2 1 1" + layout}, - TestCase{igemm_wrw, precision + flags_wrw + " --input 128 2048 7 7 --weights 512 2048 1 1 --pads_strides_dilations 0 0 1 1 1 1"}, - TestCase{igemm_wrw, precision + flags_wrw + " --input 128 2048 7 7 --weights 512 2048 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, - TestCase{igemm_wrw, precision + flags_wrw + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, - TestCase{igemm_wrw, precision + flags_wrw + " --input 256 1024 14 14 --weights 1024 32 1 1 --pads_strides_dilations 0 0 1 1 1 1" + groupCount_32} - // clang-format on - }; - - // FWD BWD WRW cases in test_cases - if(precision == "--float" || precision == "--half") - { - test_cases.reserve(test_cases_bwd_wrw.size()); - test_cases.insert(test_cases.end(), test_cases_bwd_wrw.begin(), test_cases_bwd_wrw.end()); - } - - return test_cases; -} - -} // namespace conv_igemm_mlir -using namespace conv_igemm_mlir; - -TEST_P(ConvIgemmMlirConfigFloat, FloatTest_conv_igemm_mlir) -{ -#if MIOPEN_USE_MLIR - - const auto& handle = get_handle(); - if((miopen::StartsWith(handle.GetDeviceName(), "gfx103") || - miopen::StartsWith(handle.GetDeviceName(), "gfx906")) && - miopen::IsEnabled(ENV(MIOPEN_TEST_MLIR)) && miopen::IsEnabled(ENV(MIOPEN_TEST_ALL)) && - IsTestRunWith("--float")) - { - Run2dDriver(miopenFloat); - } - else - { - GTEST_SKIP(); - } - -#else - GTEST_SKIP(); -#endif -}; - -TEST_P(ConvIgemmMlirConfigHalf, HalfTest_conv_igemm_mlir) -{ -#if MIOPEN_USE_MLIR - - const auto& handle = get_handle(); - if((miopen::StartsWith(handle.GetDeviceName(), "gfx103") || - miopen::StartsWith(handle.GetDeviceName(), "gfx906")) && - miopen::IsEnabled(ENV(MIOPEN_TEST_MLIR)) && miopen::IsEnabled(ENV(MIOPEN_TEST_ALL)) && - IsTestRunWith("--half")) - { - Run2dDriver(miopenHalf); - } - else - { - GTEST_SKIP(); - } - -#else - GTEST_SKIP(); -#endif -}; - -TEST_P(ConvIgemmMlirConfigInt8, Int8Test_conv_igemm_mlir) -{ -#if MIOPEN_USE_MLIR - - const auto& handle = get_handle(); - if((miopen::StartsWith(handle.GetDeviceName(), "gfx103") || - miopen::StartsWith(handle.GetDeviceName(), "gfx906")) && - miopen::IsEnabled(ENV(MIOPEN_TEST_MLIR)) && miopen::IsEnabled(ENV(MIOPEN_TEST_ALL)) && - IsTestRunWith("--int8")) - { - Run2dDriver(miopenInt8); - } - else - { - GTEST_SKIP(); - } - -#else - GTEST_SKIP(); -#endif -}; - -// Float for FWD, BWD, WRW -INSTANTIATE_TEST_SUITE_P(ConvIgemmMlir, - ConvIgemmMlirConfigFloat, - testing::Values(GetTestCases("--float"))); -// Half for FWD, BWD, WRW -INSTANTIATE_TEST_SUITE_P(ConvIgemmMlir, - ConvIgemmMlirConfigHalf, - testing::Values(GetTestCases("--half"))); -// Int8 for FWD -INSTANTIATE_TEST_SUITE_P(ConvIgemmMlir, - ConvIgemmMlirConfigInt8, - testing::Values(GetTestCases("--int8"))); diff --git a/test/gtest/conv_igemm_mlir_bwd_wrw.cpp b/test/gtest/conv_igemm_mlir_bwd_wrw.cpp new file mode 100644 index 0000000000..d3e0494a6b --- /dev/null +++ b/test/gtest/conv_igemm_mlir_bwd_wrw.cpp @@ -0,0 +1,127 @@ +/******************************************************************************* + * + * 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 "gtest_common.hpp" + +#include "../conv2d.hpp" + +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_MLIR) +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL) + +namespace { + +auto GetTestCases() +{ + const auto igemm_bwd = std::tuple{ + std::pair{ENV(MIOPEN_FIND_MODE), std::string_view("normal")}, + std::pair{ENV(MIOPEN_DEBUG_FIND_ONLY_SOLVER), std::string_view("ConvMlirIgemmBwd")}}; + + const auto igemm_wrw = std::tuple{ + std::pair{ENV(MIOPEN_FIND_MODE), std::string_view("normal")}, + std::pair{ENV(MIOPEN_DEBUG_FIND_ONLY_SOLVER), std::string_view("ConvMlirIgemmWrW")}}; + + const std::string flags_bwd = " --verbose --disable-forward --disable-backward-weights"; + const std::string flags_wrw = " --verbose --disable-forward --disable-backward-data"; + const std::string layout = " --in_layout NHWC --fil_layout NHWC --out_layout NHWC"; + const std::string groupCount_32 = " --group-count 32"; + + // BWD WRW test cases + return std::vector{ + // clang-format off + std::pair{igemm_bwd, flags_bwd + " --input 256 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1"}, + std::pair{igemm_bwd, flags_bwd + " --input 256 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1" + layout}, + std::pair{igemm_bwd, flags_bwd + " --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{igemm_bwd, flags_bwd + " --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + layout}, + std::pair{igemm_bwd, flags_bwd + " --input 128 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{igemm_bwd, flags_bwd + " --input 128 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1" + layout}, + std::pair{igemm_bwd, flags_bwd + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1"}, + std::pair{igemm_bwd, flags_bwd + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, + + std::pair{igemm_wrw, flags_wrw + " --input 64 1024 14 14 --weights 256 1024 1 1 --pads_strides_dilations 0 0 1 1 1 1"}, + std::pair{igemm_wrw, flags_wrw + " --input 64 1024 14 14 --weights 256 1024 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, + std::pair{igemm_wrw, flags_wrw + " --input 256 256 14 14 --weights 256 256 3 3 --pads_strides_dilations 0 0 2 2 1 1"}, + std::pair{igemm_wrw, flags_wrw + " --input 256 256 14 14 --weights 256 256 3 3 --pads_strides_dilations 0 0 2 2 1 1" + layout}, + std::pair{igemm_wrw, flags_wrw + " --input 128 2048 7 7 --weights 512 2048 1 1 --pads_strides_dilations 0 0 1 1 1 1"}, + std::pair{igemm_wrw, flags_wrw + " --input 128 2048 7 7 --weights 512 2048 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, + std::pair{igemm_wrw, flags_wrw + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, + std::pair{igemm_wrw, flags_wrw + " --input 256 1024 14 14 --weights 1024 32 1 1 --pads_strides_dilations 0 0 1 1 1 1" + groupCount_32} + // clang-format on + }; +} + +using TestCase = decltype(GetTestCases())::value_type; + +bool SkipTest() +{ + return !(miopen::IsEnabled(ENV(MIOPEN_TEST_MLIR))) || miopen::IsDisabled(ENV(MIOPEN_TEST_ALL)); +} + +bool IsTestSupportedForDevice() +{ + using e_mask = enabled; + using d_mask = disabled; + return ::IsTestSupportedForDevMask(); +} + +} // namespace + +class Conv2dDefaultFloat : public FloatTestCase> +{ +}; +class Conv2dDefaultHalf : public HalfTestCase> +{ +}; + +TEST_P(Conv2dDefaultFloat, FloatTest_conv_igemm_mlir_bwd_wrw) +{ + if(IsTestSupportedForDevice() && !SkipTest()) + { + invoke_with_params(db_check); + } + else + { + GTEST_SKIP(); + } +}; + +TEST_P(Conv2dDefaultHalf, HalfTest_conv_igemm_mlir_bwd_wrw) +{ + if(IsTestSupportedForDevice() && !SkipTest()) + { + invoke_with_params(db_check); + } + else + { + GTEST_SKIP(); + } +}; + +// Float for FWD, BWD, WRW +INSTANTIATE_TEST_SUITE_P(ConvIgemmMlir, Conv2dDefaultFloat, testing::Values(GetTestCases())); +// Half for FWD, BWD, WRW +INSTANTIATE_TEST_SUITE_P(ConvIgemmMlir, Conv2dDefaultHalf, testing::Values(GetTestCases())); diff --git a/test/gtest/conv_igemm_mlir_fwd.cpp b/test/gtest/conv_igemm_mlir_fwd.cpp new file mode 100644 index 0000000000..c0d57b8ab6 --- /dev/null +++ b/test/gtest/conv_igemm_mlir_fwd.cpp @@ -0,0 +1,130 @@ +/******************************************************************************* + * + * 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 "gtest_common.hpp" + +#include "../conv2d.hpp" + +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_MLIR) +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL) + +namespace { + +auto GetTestCases() +{ + const auto igemm_fwd = std::tuple{ + std::pair{ENV(MIOPEN_FIND_MODE), std::string_view("normal")}, + std::pair{ENV(MIOPEN_DEBUG_FIND_ONLY_SOLVER), std::string_view("ConvMlirIgemmFwd")}}; + + const std::string vf = " --verbose --disable-backward-data --disable-backward-weights"; + const std::string layout = " --in_layout NHWC --fil_layout NHWC --out_layout NHWC"; + const std::string groupCount_4 = " --group-count 4"; + + // FWD test cases for precision == "--int8" + return std::vector{ + // clang-format off + std::pair{igemm_fwd, vf + " --input 256 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1"}, + std::pair{igemm_fwd, vf + " --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{igemm_fwd, vf + " --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + layout}, + std::pair{igemm_fwd, vf + " --input 128 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{igemm_fwd, vf + " --input 128 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1" + layout}, + std::pair{igemm_fwd, vf + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1"}, + std::pair{igemm_fwd, vf + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, + std::pair{igemm_fwd, vf + " --input 256 256 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + groupCount_4} + // clang-format on + }; +} + +using TestCase = decltype(GetTestCases())::value_type; + +bool SkipTest() +{ + return !(miopen::IsEnabled(ENV(MIOPEN_TEST_MLIR))) || miopen::IsDisabled(ENV(MIOPEN_TEST_ALL)); +} + +bool IsTestSupportedForDevice() +{ + using e_mask = enabled; + using d_mask = disabled; + return ::IsTestSupportedForDevMask(); +} + +} // namespace + +class Conv2dDefaultFloat : public FloatTestCase> +{ +}; +class Conv2dDefaultHalf : public HalfTestCase> +{ +}; +class Conv2dDefaultInt8 : public Int8TestCase> +{ +}; + +TEST_P(Conv2dDefaultFloat, FloatTest_conv_igemm_mlir_fwd) +{ + if(IsTestSupportedForDevice() && !SkipTest()) + { + invoke_with_params(db_check); + } + else + { + GTEST_SKIP(); + } +}; + +TEST_P(Conv2dDefaultHalf, HalfTest_conv_igemm_mlir_fwd) +{ + if(IsTestSupportedForDevice() && !SkipTest()) + { + invoke_with_params(db_check); + } + else + { + GTEST_SKIP(); + } +}; + +TEST_P(Conv2dDefaultInt8, Int8Test_conv_igemm_mlir_fwd) +{ + if(IsTestSupportedForDevice() && !SkipTest()) + { + invoke_with_params(db_check); + } + else + { + GTEST_SKIP(); + } +}; + +// Float for FWD, BWD, WRW +INSTANTIATE_TEST_SUITE_P(ConvIgemmMlir, Conv2dDefaultFloat, testing::Values(GetTestCases())); +// Half for FWD, BWD, WRW +INSTANTIATE_TEST_SUITE_P(ConvIgemmMlir, Conv2dDefaultHalf, testing::Values(GetTestCases())); +// Int8 for FWD +INSTANTIATE_TEST_SUITE_P(ConvIgemmMlir, Conv2dDefaultInt8, testing::Values(GetTestCases())); diff --git a/test/gtest/conv_igemm_mlir_xdlops.cpp b/test/gtest/conv_igemm_mlir_xdlops.cpp deleted file mode 100644 index ff43a91896..0000000000 --- a/test/gtest/conv_igemm_mlir_xdlops.cpp +++ /dev/null @@ -1,201 +0,0 @@ -#include - -#include -#include -#include -#include -#include "get_handle.hpp" -#include "test_env.hpp" - -#include "conv2d.hpp" - -MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_MLIR) -MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL) - -namespace conv_igemm_mlir_xdlops { - -using TestCase = std::tuple, std::string>; - -void GetArgs(const TestCase& param, std::vector& tokens) -{ - auto env_vars = std::get<0>(param); - for(auto& elem : env_vars) - { - putenv(elem.data()); - } - - auto cmd = std::get<1>(param); - - std::stringstream ss(cmd); - std::istream_iterator begin(ss); - std::istream_iterator end; - while(begin != end) - tokens.push_back(*begin++); -} - -class ConvIgemmMlirXdlopsConfigHalf : public testing::TestWithParam> -{ -}; -class ConvIgemmMlirXdlopsConfigInt8 : public testing::TestWithParam> -{ -}; - -void Run2dDriver(miopenDataType_t prec) -{ - - std::vector params; - switch(prec) - { - case miopenHalf: params = ConvIgemmMlirXdlopsConfigHalf::GetParam(); break; - case miopenInt8: params = ConvIgemmMlirXdlopsConfigInt8::GetParam(); break; - case miopenBFloat16: - case miopenFloat: - case miopenInt32: - case miopenDouble: - case miopenFloat8: - case miopenBFloat8: - MIOPEN_THROW(miopenStatusBadParm, - "miopenBFloat16, miopenFloat, miopenInt32, miopenDouble data " - "type not supported by " - "conv_igemm_mlir_xdlops test"); - - default: params = ConvIgemmMlirXdlopsConfigHalf::GetParam(); - } - - 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 auto& str) { - return str.data(); - }); - - testing::internal::CaptureStderr(); - test_drive(ptrs.size(), ptrs.data()); - auto capture = testing::internal::GetCapturedStderr(); - EXPECT_FALSE(capture.find("Perf Db: record not found") != std::string::npos); - } -}; - -std::vector GetTestCases(const std::string& precision) -{ - std::vector fwd = {"MIOPEN_FIND_MODE=normal", - "MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvMlirIgemmFwdXdlops"}; - std::string flags_fwd = " --verbose --disable-backward-data --disable-backward-weights"; - std::string layout = " --in_layout NHWC --fil_layout NHWC --out_layout NHWC"; - std::string groupCount_4 = " --group-count 4"; - - // FWD test cases for precision == "--int8" - std::vector test_cases = { - // clang-format off - TestCase{fwd, precision + flags_fwd + " --input 256 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1"}, - TestCase{fwd, precision + flags_fwd + " --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, - TestCase{fwd, precision + flags_fwd + " --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + layout}, - TestCase{fwd, precision + flags_fwd + " --input 128 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, - TestCase{fwd, precision + flags_fwd + " --input 128 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1" + layout}, - TestCase{fwd, precision + flags_fwd + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1"}, - TestCase{fwd, precision + flags_fwd + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, - TestCase{fwd, precision + flags_fwd + " --input 256 256 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + groupCount_4} - // clang-format on - }; - - std::vector bwd = {"MIOPEN_FIND_MODE=normal", - "MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvMlirIgemmBwdXdlops"}; - std::vector wrw = {"MIOPEN_FIND_MODE=normal", - "MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvMlirIgemmWrWXdlops"}; - - std::string flags_bwd = " --verbose --disable-forward --disable-backward-weights"; - std::string flags_wrw = " --verbose --disable-forward --disable-backward-data"; - std::string groupCount_32 = " --group-count 32"; - - // BWD WRW test cases - const std::vector test_cases_bwd_wrw = { - // clang-format off - TestCase{bwd, precision + flags_bwd + " --input 256 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1"}, - TestCase{bwd, precision + flags_bwd + " --input 256 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1" + layout}, - TestCase{bwd, precision + flags_bwd + " --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, - TestCase{bwd, precision + flags_bwd + " --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + layout}, - TestCase{bwd, precision + flags_bwd + " --input 128 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, - TestCase{bwd, precision + flags_bwd + " --input 128 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1" + layout}, - TestCase{bwd, precision + flags_bwd + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1"}, - TestCase{bwd, precision + flags_bwd + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, - - TestCase{wrw, precision + flags_wrw + " --input 64 1024 14 14 --weights 256 1024 1 1 --pads_strides_dilations 0 0 1 1 1 1"}, - TestCase{wrw, precision + flags_wrw + " --input 64 1024 14 14 --weights 256 1024 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, - TestCase{wrw, precision + flags_wrw + " --input 256 256 14 14 --weights 256 256 3 3 --pads_strides_dilations 0 0 2 2 1 1"}, - TestCase{wrw, precision + flags_wrw + " --input 256 256 14 14 --weights 256 256 3 3 --pads_strides_dilations 0 0 2 2 1 1" + layout}, - TestCase{wrw, precision + flags_wrw + " --input 128 2048 7 7 --weights 512 2048 1 1 --pads_strides_dilations 0 0 1 1 1 1"}, - TestCase{wrw, precision + flags_wrw + " --input 128 2048 7 7 --weights 512 2048 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, - TestCase{wrw, precision + flags_wrw + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, - TestCase{wrw, precision + flags_wrw + " --input 256 1024 14 14 --weights 1024 32 1 1 --pads_strides_dilations 0 0 1 1 1 1" + groupCount_32}, - TestCase{wrw, precision + flags_wrw + " --input 64 1024 14 14 --weights 1024 1024 1 1 --pads_strides_dilations 0 0 1 1 1 1"} - // clang-format on - }; - - // FWD BWD WRW cases in test_cases for precision == "--half" - if(precision == "--half") - { - test_cases.reserve(test_cases_bwd_wrw.size()); - test_cases.insert(test_cases.end(), test_cases_bwd_wrw.begin(), test_cases_bwd_wrw.end()); - } - - return test_cases; -} - -} // namespace conv_igemm_mlir_xdlops -using namespace conv_igemm_mlir_xdlops; - -TEST_P(ConvIgemmMlirXdlopsConfigHalf, HalfTest_conv_igemm_mlir_xdlops) -{ -#if MIOPEN_USE_MLIR - - const auto& handle = get_handle(); - if((miopen::StartsWith(handle.GetDeviceName(), "gfx908") || - miopen::StartsWith(handle.GetDeviceName(), "gfx90a")) && - miopen::IsEnabled(ENV(MIOPEN_TEST_MLIR)) && miopen::IsEnabled(ENV(MIOPEN_TEST_ALL)) && - IsTestRunWith("--half")) - { - Run2dDriver(miopenHalf); - } - else - { - GTEST_SKIP(); - } - -#else - GTEST_SKIP(); -#endif -}; - -TEST_P(ConvIgemmMlirXdlopsConfigInt8, Int8Test_conv_igemm_mlir_xdlops) -{ -#if MIOPEN_USE_MLIR - - const auto& handle = get_handle(); - if((miopen::StartsWith(handle.GetDeviceName(), "gfx908") || - miopen::StartsWith(handle.GetDeviceName(), "gfx90a")) && - miopen::IsEnabled(ENV(MIOPEN_TEST_MLIR)) && miopen::IsEnabled(ENV(MIOPEN_TEST_ALL)) && - IsTestRunWith("--int8")) - { - Run2dDriver(miopenInt8); - } - else - { - GTEST_SKIP(); - } - -#else - GTEST_SKIP(); -#endif -}; - -// Half for FWD, BWD, WRW -INSTANTIATE_TEST_SUITE_P(ConvIgemmMlirXdlops, - ConvIgemmMlirXdlopsConfigHalf, - testing::Values(GetTestCases("--half"))); -// Int8 for FWD -INSTANTIATE_TEST_SUITE_P(ConvIgemmMlirXdlops, - ConvIgemmMlirXdlopsConfigInt8, - testing::Values(GetTestCases("--int8"))); diff --git a/test/gtest/conv_igemm_mlir_xdlops_bwd_wrw.cpp b/test/gtest/conv_igemm_mlir_xdlops_bwd_wrw.cpp new file mode 100644 index 0000000000..5e65f956d7 --- /dev/null +++ b/test/gtest/conv_igemm_mlir_xdlops_bwd_wrw.cpp @@ -0,0 +1,113 @@ +/******************************************************************************* + * + * 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 "gtest_common.hpp" + +#include "../conv2d.hpp" + +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_MLIR) +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL) + +namespace { + +auto GetTestCases() +{ + const auto bwd = std::tuple{ + std::pair{ENV(MIOPEN_FIND_MODE), std::string_view("normal")}, + std::pair{ENV(MIOPEN_DEBUG_FIND_ONLY_SOLVER), std::string_view("ConvMlirIgemmBwdXdlops")}}; + + const auto wrw = std::tuple{ + std::pair{ENV(MIOPEN_FIND_MODE), std::string_view("normal")}, + std::pair{ENV(MIOPEN_DEBUG_FIND_ONLY_SOLVER), std::string_view("ConvMlirIgemmWrWXdlops")}}; + + const std::string flags_bwd = " --verbose --disable-forward --disable-backward-weights"; + const std::string flags_wrw = " --verbose --disable-forward --disable-backward-data"; + const std::string layout = " --in_layout NHWC --fil_layout NHWC --out_layout NHWC"; + + const std::string groupCount_32 = " --group-count 32"; + + // BWD WRW test cases + return std::vector{ + // clang-format off + std::pair{bwd, flags_bwd + " --input 256 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1"}, + std::pair{bwd, flags_bwd + " --input 256 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1" + layout}, + std::pair{bwd, flags_bwd + " --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{bwd, flags_bwd + " --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + layout}, + std::pair{bwd, flags_bwd + " --input 128 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{bwd, flags_bwd + " --input 128 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1" + layout}, + std::pair{bwd, flags_bwd + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1"}, + std::pair{bwd, flags_bwd + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, + + std::pair{wrw, flags_wrw + " --input 64 1024 14 14 --weights 256 1024 1 1 --pads_strides_dilations 0 0 1 1 1 1"}, + std::pair{wrw, flags_wrw + " --input 64 1024 14 14 --weights 256 1024 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, + std::pair{wrw, flags_wrw + " --input 256 256 14 14 --weights 256 256 3 3 --pads_strides_dilations 0 0 2 2 1 1"}, + std::pair{wrw, flags_wrw + " --input 256 256 14 14 --weights 256 256 3 3 --pads_strides_dilations 0 0 2 2 1 1" + layout}, + std::pair{wrw, flags_wrw + " --input 128 2048 7 7 --weights 512 2048 1 1 --pads_strides_dilations 0 0 1 1 1 1"}, + std::pair{wrw, flags_wrw + " --input 128 2048 7 7 --weights 512 2048 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, + std::pair{wrw, flags_wrw + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, + std::pair{wrw, flags_wrw + " --input 256 1024 14 14 --weights 1024 32 1 1 --pads_strides_dilations 0 0 1 1 1 1" + groupCount_32}, + std::pair{wrw, flags_wrw + " --input 64 1024 14 14 --weights 1024 1024 1 1 --pads_strides_dilations 0 0 1 1 1 1"} + // clang-format on + }; +} + +using TestCase = decltype(GetTestCases())::value_type; + +bool SkipTest() +{ + return !(miopen::IsEnabled(ENV(MIOPEN_TEST_MLIR))) || miopen::IsDisabled(ENV(MIOPEN_TEST_ALL)); +} + +bool IsTestSupportedForDevice() +{ + using e_mask = enabled; + using d_mask = disabled; + return ::IsTestSupportedForDevMask(); +} + +} // namespace + +class Conv2dDefaultHalf : public FloatTestCase> +{ +}; + +TEST_P(Conv2dDefaultHalf, HalfTest_conv_igemm_mlir_xdlops_bwd_wrw) +{ + if(IsTestSupportedForDevice() && !SkipTest()) + { + invoke_with_params(db_check); + } + else + { + GTEST_SKIP(); + } +}; + +// Half for FWD, BWD, WRW +INSTANTIATE_TEST_SUITE_P(ConvIgemmMlirXdlops, Conv2dDefaultHalf, testing::Values(GetTestCases())); diff --git a/test/gtest/conv_igemm_mlir_xdlops_fwd.cpp b/test/gtest/conv_igemm_mlir_xdlops_fwd.cpp new file mode 100644 index 0000000000..fb913a9346 --- /dev/null +++ b/test/gtest/conv_igemm_mlir_xdlops_fwd.cpp @@ -0,0 +1,115 @@ +/******************************************************************************* + * + * 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 "gtest_common.hpp" + +#include "../conv2d.hpp" + +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_MLIR) +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_ALL) + +namespace { + +auto GetTestCases() +{ + const auto fwd = std::tuple{ + std::pair{ENV(MIOPEN_FIND_MODE), std::string_view("normal")}, + std::pair{ENV(MIOPEN_DEBUG_FIND_ONLY_SOLVER), std::string_view("ConvMlirIgemmFwdXdlops")}}; + + const std::string flags_fwd = " --verbose --disable-backward-data --disable-backward-weights"; + const std::string layout = " --in_layout NHWC --fil_layout NHWC --out_layout NHWC"; + const std::string groupCount_4 = " --group-count 4"; + + // FWD test cases for precision == "--int8" + return std::vector{ + // clang-format off + std::pair{fwd, flags_fwd + " --input 256 1024 14 14 --weights 2048 1024 1 1 --pads_strides_dilations 0 0 2 2 1 1"}, + std::pair{fwd, flags_fwd + " --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{fwd, flags_fwd + " --input 256 128 28 28 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1" + layout}, + std::pair{fwd, flags_fwd + " --input 128 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{fwd, flags_fwd + " --input 128 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1" + layout}, + std::pair{fwd, flags_fwd + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1"}, + std::pair{fwd, flags_fwd + " --input 128 64 56 56 --weights 64 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + layout}, + std::pair{fwd, flags_fwd + " --input 256 256 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1" + groupCount_4} + // clang-format on + }; +} + +using TestCase = decltype(GetTestCases())::value_type; + +bool SkipTest() +{ + return !(miopen::IsEnabled(ENV(MIOPEN_TEST_MLIR))) || miopen::IsDisabled(ENV(MIOPEN_TEST_ALL)); +} + +bool IsTestSupportedForDevice() +{ + using e_mask = enabled; + using d_mask = disabled; + return ::IsTestSupportedForDevMask(); +} + +} // namespace + +class Conv2dDefaultHalf : public FloatTestCase> +{ +}; + +class Conv2dDefaultInt8 : public Int8TestCase> +{ +}; + +TEST_P(Conv2dDefaultHalf, HalfTest_conv_igemm_mlir_xdlops_fwd) +{ + if(IsTestSupportedForDevice() && !SkipTest()) + { + invoke_with_params(db_check); + } + else + { + GTEST_SKIP(); + } +}; + +TEST_P(Conv2dDefaultInt8, Int8Test_conv_igemm_mlir_xdlops_fwd) +{ + if(IsTestSupportedForDevice() && !SkipTest()) + { + invoke_with_params(db_check); + } + else + { + GTEST_SKIP(); + } +}; + +// Half for FWD, BWD, WRW +INSTANTIATE_TEST_SUITE_P(ConvIgemmMlirXdlops, Conv2dDefaultHalf, testing::Values(GetTestCases())); +// Int8 for FWD +INSTANTIATE_TEST_SUITE_P(ConvIgemmMlirXdlops, Conv2dDefaultInt8, testing::Values(GetTestCases())); diff --git a/test/gtest/deepbench_conv.cpp b/test/gtest/deepbench_conv.cpp new file mode 100644 index 0000000000..c22d6d70a5 --- /dev/null +++ b/test/gtest/deepbench_conv.cpp @@ -0,0 +1,112 @@ +/******************************************************************************* + * + * 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 "gtest_common.hpp" + +#include "../conv2d.hpp" + +MIOPEN_DECLARE_ENV_VAR_BOOL(MIOPEN_TEST_DEEPBENCH) + +namespace deepbench_conv { + +auto GetTestCases() +{ + const std::string v = " --verbose"; + + return std::vector{ + // clang-format off + std::pair{std::tuple<>{}, v + " --input 4 1 161 700 --weights 32 1 5 20 --pads_strides_dilations 0 0 2 2 1 1"}, + std::pair{std::tuple<>{}, v + " --input 8 1 161 700 --weights 32 1 5 20 --pads_strides_dilations 0 0 2 2 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 1 161 700 --weights 32 1 5 20 --pads_strides_dilations 0 0 2 2 1 1"}, + std::pair{std::tuple<>{}, v + " --input 32 1 161 700 --weights 32 1 5 20 --pads_strides_dilations 0 0 2 2 1 1"}, + std::pair{std::tuple<>{}, v + " --input 4 32 79 341 --weights 32 32 5 10 --pads_strides_dilations 0 0 2 2 1 1"}, + std::pair{std::tuple<>{}, v + " --input 8 32 79 341 --weights 32 32 5 10 --pads_strides_dilations 0 0 2 2 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 32 79 341 --weights 32 32 5 10 --pads_strides_dilations 0 0 2 2 1 1"}, + std::pair{std::tuple<>{}, v + " --input 32 32 79 341 --weights 32 32 5 10 --pads_strides_dilations 0 0 2 2 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 1 48 480 --weights 16 1 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 16 24 240 --weights 32 16 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 32 12 120 --weights 64 32 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 64 6 60 --weights 128 64 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 8 3 108 108 --weights 64 3 3 3 --pads_strides_dilations 1 1 2 2 1 1"}, + std::pair{std::tuple<>{}, v + " --input 8 64 54 54 --weights 64 64 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 8 128 27 27 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 8 128 14 14 --weights 256 128 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 8 256 7 7 --weights 512 256 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 8 3 224 224 --weights 64 3 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 8 64 112 112 --weights 128 64 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 8 128 56 56 --weights 256 128 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 8 256 28 28 --weights 512 256 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 8 512 14 14 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 8 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 3 224 224 --weights 64 3 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 64 112 112 --weights 128 64 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 128 56 56 --weights 256 128 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 256 28 28 --weights 512 256 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 512 14 14 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 512 7 7 --weights 512 512 3 3 --pads_strides_dilations 1 1 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 3 224 224 --weights 64 3 7 7 --pads_strides_dilations 3 3 2 2 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 192 28 28 --weights 32 192 5 5 --pads_strides_dilations 2 2 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 512 14 14 --weights 48 512 5 5 --pads_strides_dilations 2 2 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 832 7 7 --weights 128 832 5 5 --pads_strides_dilations 2 2 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 192 28 28 --weights 32 192 1 1 --pads_strides_dilations 0 0 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 512 14 14 --weights 48 512 1 1 --pads_strides_dilations 0 0 1 1 1 1"}, + std::pair{std::tuple<>{}, v + " --input 16 832 7 7 --weights 128 832 1 1 --pads_strides_dilations 0 0 1 1 1 1"} + // clang-format on + }; +} + +using TestCase = decltype(GetTestCases())::value_type; + +bool SkipTest() { return miopen::IsDisabled(ENV(MIOPEN_TEST_DEEPBENCH)); } + +class Conv2dFloat_deepbench : public FloatTestCase> +{ +}; + +bool IsTestSupportedForDevice() +{ + using e_mask = enabled; + using d_mask = disabled; + return ::IsTestSupportedForDevMask(); +} +} // namespace deepbench_conv +using namespace deepbench_conv; + +TEST_P(Conv2dFloat_deepbench, FloatTest_deepbench_conv) +{ + if(IsTestSupportedForDevice() && !SkipTest()) + { + invoke_with_params(default_check); + } + else + { + GTEST_SKIP(); + } +}; + +INSTANTIATE_TEST_SUITE_P(DeepbenchConv, Conv2dFloat_deepbench, testing::Values(GetTestCases())); diff --git a/test/gtest/gtest_common.hpp b/test/gtest/gtest_common.hpp index 2100a2dd1c..3a52e4cfe3 100644 --- a/test/gtest/gtest_common.hpp +++ b/test/gtest/gtest_common.hpp @@ -38,15 +38,21 @@ #include "../driver.hpp" -void default_check(const std::string& err) { std::cout << err; } +inline void default_check(const std::string& err) { std::cout << err; } -void tuning_check(const std::string& err) +inline void tuning_check(const std::string& err) { // TEST_TUNING - the test should fail if output contains "Error" or "failed". EXPECT_FALSE(err.find("Error") != std::string::npos || err.find("failed") != std::string::npos); default_check(err); } +inline void db_check(const std::string& err) +{ + EXPECT_FALSE(err.find("Perf Db: record not found") != std::string::npos); + default_check(err); +}; + enum class Gpu : int { Default = 0, @@ -74,7 +80,7 @@ struct disabled }; template -bool IsTestSupportedForDevice() +bool IsTestSupportedForDevMask() { static_assert((~disabled_mask::val & enabled_mask::val) == 0, "Enabled and Disabled GPUs are overlapped"); @@ -108,6 +114,30 @@ bool IsTestSupportedForDevice() return res; } +template +struct FloatTestCase : public testing::TestWithParam +{ + static constexpr std::string_view fp_args{"--float"}; +}; + +template +struct HalfTestCase : public testing::TestWithParam +{ + static constexpr std::string_view fp_args{"--half"}; +}; + +template +struct Bf16TestCase : public testing::TestWithParam +{ + static constexpr std::string_view fp_args{"--bfloat16"}; +}; + +template +struct Int8TestCase : public testing::TestWithParam +{ + static constexpr std::string_view fp_args{"--int8"}; +}; + template std::vector get_args(const Case& param) { @@ -123,14 +153,15 @@ std::vector get_args(const Case& param) return {begin, end}; } -template