diff --git a/Dockerfile b/Dockerfile index d958879d3c..2c4a52af8d 100755 --- a/Dockerfile +++ b/Dockerfile @@ -89,8 +89,9 @@ ARG CCACHE_DIR="/tmp" RUN env # RUN cget -p $PREFIX install https://github.com/ccache/ccache/archive/7f1572ae9ca958fa923a66235f6a64a360b03523.tar.gz -DZSTD_FROM_INTERNET=ON -DHIREDIS_FROM_INTERNET=ON ARG CCACHE_COMMIT=7f1572ae9ca958fa923a66235f6a64a360b03523 -RUN rm -rf /tmp/ccache* && mkdir /tmp/ccache && wget https://github.com/ccache/ccache/archive/${CCACHE_COMMIT}.tar.gz -O /tmp/ccache.tar.gz && \ - tar zxvf /tmp/ccache.tar.gz -C /tmp/ && mkdir /tmp/ccache-${CCACHE_COMMIT}/build && \ +RUN rm -rf /tmp/ccache* && mkdir /tmp/ccache +ADD https://github.com/ccache/ccache/archive/${CCACHE_COMMIT}.tar.gz /tmp/ccache.tar.gz +RUN tar zxvf /tmp/ccache.tar.gz -C /tmp/ && mkdir /tmp/ccache-${CCACHE_COMMIT}/build && \ cd /tmp/ccache-${CCACHE_COMMIT}/build && \ cmake -DZSTD_FROM_INTERNET=ON -DHIREDIS_FROM_INTERNET=ON .. && make -j install && rm -rf /tmp/* RUN ccache -s diff --git a/Jenkinsfile b/Jenkinsfile index c3bb134be3..335585ed1b 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -469,6 +469,10 @@ pipeline { name: "TARGET_GFX90A", defaultValue: true, description: "") + booleanParam( + name: "TARGET_GFX94X", + defaultValue: false, + description: "") booleanParam( name: "TARGET_NAVI21", defaultValue: false, @@ -678,6 +682,19 @@ pipeline { buildHipClangJobAndReboot(build_type: 'debug', config_targets: Smoke_targets) } } + stage('Fp32 Hip Debug gfx94X') { + when { + beforeAgent true + expression { params.TARGET_GFX94X } + } + options { + retry(2) + } + agent{ label rocmnode("gfx94X") } + steps{ + buildHipClangJobAndReboot(build_type: 'debug', config_targets: Smoke_targets) + } + } } } stage("Smoke Aux 1") { @@ -863,6 +880,32 @@ pipeline { buildHipClangJobAndReboot(setup_flags: Bf16_flags, config_targets: Smoke_targets) } } + stage('Fp16 Hip gfx94X') { + when { + beforeAgent true + expression { params.TARGET_GFX94X && params.DATATYPE_FP16 } + } + options { + retry(2) + } + agent{ label rocmnode("gfx94X") } + steps{ + buildHipClangJobAndReboot( setup_flags: Fp16_flags, config_targets: Smoke_targets) + } + } + stage('Bf16 Hip gfx94X') { + when { + beforeAgent true + expression { params.TARGET_GFX94X && params.DATATYPE_BF16 } + } + options { + retry(2) + } + agent{ label rocmnode("gfx94X") } + steps{ + buildHipClangJobAndReboot(setup_flags: Bf16_flags, config_targets: Smoke_targets) + } + } } } stage("Full Tests") { @@ -914,6 +957,19 @@ pipeline { buildHipClangJobAndReboot(setup_flags: Bf16_flags + Full_test, build_install: "true") } } + stage('Bf16 Hip Install All gfx94X') { + when { + beforeAgent true + expression { params.TARGET_GFX94X && params.DATATYPE_BF16 } + } + options { + retry(2) + } + agent{ label rocmnode("gfx94X") } + steps{ + buildHipClangJobAndReboot(setup_flags: Bf16_flags + Full_test, build_install: "true") + } + } stage('Fp16 Hip All gfx1030') { when { beforeAgent true @@ -963,6 +1019,19 @@ pipeline { // buildHipClangJobAndReboot(setup_flags: Full_test, enforce_xnack_on: true) // } // } + stage('Fp32 Hip All gfx94X') { + when { + beforeAgent true + expression { params.TARGET_GFX94X && params.DATATYPE_FP32 } + } + options { + retry(2) + } + agent{ label rocmnode("gfx94X") } + steps{ + buildHipClangJobAndReboot(setup_flags: Full_test) + } + } stage('Fp16 Hip Install All Vega20') { when { beforeAgent true @@ -1028,6 +1097,19 @@ pipeline { buildHipClangJobAndReboot(setup_flags: Full_test + Fp16_flags, build_install: "true") } } + stage('Fp16 Hip All Install gfx94X') { + when { + beforeAgent true + expression { params.TARGET_GFX94X && params.DATATYPE_FP16 } + } + options { + retry(2) + } + agent{ label rocmnode("gfx94X") } + steps{ + buildHipClangJobAndReboot(setup_flags: Full_test + Fp16_flags, build_install: "true") + } + } } } stage("Performance Tests - gfx90a") { diff --git a/src/include/miopen/sqlite_db.hpp b/src/include/miopen/sqlite_db.hpp index 69d449fda6..a1beef32e1 100644 --- a/src/include/miopen/sqlite_db.hpp +++ b/src/include/miopen/sqlite_db.hpp @@ -230,8 +230,7 @@ class SQLiteBase dbInvalid = true; return; } - - if(!is_system && !filename.empty()) + else if(!is_system) { auto file = boost::filesystem::path(filename_); const auto directory = file.remove_filename(); diff --git a/src/kernel_warnings.cpp b/src/kernel_warnings.cpp index 2b698bfac7..253cdb9c17 100644 --- a/src/kernel_warnings.cpp +++ b/src/kernel_warnings.cpp @@ -68,6 +68,7 @@ static std::vector HipKernelWarnings() "-Wno-double-promotion", "-Wno-exit-time-destructors", "-Wno-extra-semi", + "-Wno-extra-semi-stmt", "-Wno-float-conversion", "-Wno-gnu-anonymous-struct", "-Wno-gnu-zero-variadic-macro-arguments", diff --git a/src/sqlite_db.cpp b/src/sqlite_db.cpp index 169b5163eb..5b42ea13a6 100644 --- a/src/sqlite_db.cpp +++ b/src/sqlite_db.cpp @@ -78,12 +78,12 @@ class SQLite::impl { sqlite3* ptr_tmp = nullptr; int rc = 0; -#ifdef __clang__ +#if defined(__clang__) || defined(__llvm__) #pragma clang diagnostic push #pragma clang diagnostic ignored "-Wcast-function-type-strict" #endif sqlite3_auto_extension(reinterpret_cast(miopen_sqlite3_memvfs_init)); -#ifdef __clang__ +#if defined(__clang__) || defined(__llvm__) #pragma clang diagnostic pop #endif // Open an in-memory database to use as a handle for loading the memvfs extension diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index d721367b38..9960c1c38e 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -117,6 +117,7 @@ if(NOT (MIOPEN_TEST_GFX900 OR MIOPEN_TEST_GFX906 OR MIOPEN_TEST_GFX908 OR MIOPEN OUTPUT_VARIABLE ROCMINFO_OUTPUT RESULT_VARIABLE ROCMINFO_EXIT_STATUS ) + message(STATUS "${ROCMINFO_OUTPUT}") if(ROCMINFO_OUTPUT MATCHES "no GPU devices") message(WARNING "ROCk module is NOT loaded, possibly no GPU devices") set(MIOPEN_NO_GPU TRUE) @@ -1220,7 +1221,7 @@ set(DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_BWD_ENVS MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC) # gfx90a is disabled due to WORKAROUND_ISSUE_1187 -add_custom_test(test_conv_igemm_dynamic_xdlops_bwd SKIP_UNLESS_ALL HALF_ENABLED GFX90A_DISABLED GFX94X_ENABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON +add_custom_test(test_conv_igemm_dynamic_xdlops_bwd SKIP_UNLESS_ALL HALF_ENABLED GFX90A_DISABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 28 28 --weights 16 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 128 36 36 --weights 32 128 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights @@ -1246,13 +1247,13 @@ COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIO #) # gfx90a is disabled due to WORKAROUND_ISSUE_1187 -add_custom_test(test_conv_igemm_dynamic_xdlops_bwd_float SKIP_UNLESS_ALL HALF_DISABLED FLOAT_ENABLED GFX90A_DISABLED GFX94X_ENABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON +add_custom_test(test_conv_igemm_dynamic_xdlops_bwd_float SKIP_UNLESS_ALL HALF_DISABLED FLOAT_ENABLED GFX90A_DISABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 4 512 128 128 --weights 12 512 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-weights ) # gfx90a is disabled due to WORKAROUND_ISSUE_1187 # Be careful to add testings for (x=1, y=1, c % 8 != 0) due to WORKAROUND_SWDEV_306318 -add_custom_test(test_conv_igemm_dynamic_xdlops_fwd SKIP_UNLESS_ALL HALF_ENABLED GFX90A_DISABLED GFX94X_ENABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON +add_custom_test(test_conv_igemm_dynamic_xdlops_fwd SKIP_UNLESS_ALL HALF_ENABLED GFX90A_DISABLED GFX900_DISABLED GFX906_DISABLED SKIP_XNACK_ON COMMAND ${DYNAMIC_IMPLICITGEMM_FWD_GTC_DYNAMIC_XDLOPS_ENVS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 1024 14 14 --weights 1024 1024 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights COMMAND ${DYNAMIC_IMPLICITGEMM_FWD_GTC_DYNAMIC_XDLOPS_ENVS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 256 56 56 --weights 512 256 1 1 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights COMMAND ${DYNAMIC_IMPLICITGEMM_FWD_GTC_DYNAMIC_XDLOPS_ENVS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 2048 7 7 --weights 2048 2048 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights @@ -1270,13 +1271,13 @@ COMMAND ${DYNAMIC_IMPLICITGEMM_FWD_GTC_DYNAMIC_XDLOPS_ENVS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 3 224 224 --weights 64 3 7 7 --pads_strides_dilations 3 3 2 2 1 1 --disable-backward-data --disable-backward-weights COMMAND ${DYNAMIC_IMPLICITGEMM_FWD_GTC_DYNAMIC_XDLOPS_ENVS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 3 230 230 --weights 64 3 7 7 --pads_strides_dilations 0 0 2 2 1 1 --disable-backward-data --disable-backward-weights ) # gfx90a is disabled due to WORKAROUND_ISSUE_1187 -add_custom_test(test_conv_igemm_dynamic_xdlops_wrw SKIP_UNLESS_ALL GFX90A_DISABLED GFX94X_ENABLED GFX900_DISABLED GFX906_DISABLED HALF_ENABLED SKIP_XNACK_ON +add_custom_test(test_conv_igemm_dynamic_xdlops_wrw SKIP_UNLESS_ALL GFX90A_DISABLED GFX900_DISABLED GFX906_DISABLED HALF_ENABLED SKIP_XNACK_ON COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 28 28 --weights 32 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 16 128 36 36 --weights 32 128 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 64 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-forward --disable-backward-data @@ -1299,7 +1300,7 @@ COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS_XDLOPS} $ ${MIO ) # gfx90a is disabled due to WORKAROUND_ISSUE_1187 -add_custom_test(test_conv_igemm_dynamic_xdlops_wrw_half SKIP_UNLESS_ALL GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED GFX94X_ENABLED HALF_ENABLED FLOAT_DISABLED SKIP_XNACK_ON +add_custom_test(test_conv_igemm_dynamic_xdlops_wrw_half SKIP_UNLESS_ALL GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED HALF_ENABLED FLOAT_DISABLED SKIP_XNACK_ON COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 3 32 32 --weights 1 3 11 11 --pads_strides_dilations 1 1 2 2 2 1 --disable-forward --disable-backward-data COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 3 224 224 --weights 1 3 3 3 --pads_strides_dilations 0 0 1 1 2 2 --disable-forward --disable-backward-data COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS_XDLOPS} $ ${MIOPEN_TEST_FLOAT_ARG} --verbose --input 1 1 8 8 --weights 1 1 2 2 --pads_strides_dilations 0 0 1 1 2 2 --disable-forward --disable-backward-data @@ -1594,7 +1595,7 @@ COMMAND ${DYNAMIC_IMPLICITGEMM_DLOPS_NCHWC_FWD_ENVS} $ COMMAND ${DYNAMIC_IMPLICITGEMM_DLOPS_NCHWC_FWD_ENVS} $ ${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 GFX103X_ENABLED SKIP_XNACK_ON +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 COMMAND ${DYNAMIC_IMPLICITGEMM_DLOPS_NCHWC_FWD_ENVS} $ ${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 ${DYNAMIC_IMPLICITGEMM_DLOPS_NCHWC_FWD_ENVS} $ ${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 ${DYNAMIC_IMPLICITGEMM_DLOPS_NCHWC_FWD_ENVS} $ ${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} @@ -1634,7 +1635,7 @@ COMMAND ${DYNAMIC_IMPLICITGEMM_DLOPS_NCHWC_FWD_ENVS} $ COMMAND ${DYNAMIC_IMPLICITGEMM_DLOPS_NCHWC_FWD_ENVS} $ ${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 GFX103X_ENABLED SKIP_XNACK_ON +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 COMMAND ${DYNAMIC_IMPLICITGEMM_DLOPS_NCHWC_FWD_ENVS} $ ${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 ${DYNAMIC_IMPLICITGEMM_DLOPS_NCHWC_FWD_ENVS} $ ${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 ${DYNAMIC_IMPLICITGEMM_DLOPS_NCHWC_FWD_ENVS} $ ${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} @@ -1814,19 +1815,19 @@ 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 GFX103X_ENABLED GFX110X_ENABLED +add_custom_test(smoke_solver_ConvFFT GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=fft $ ${TEST_CONV_VERBOSE_F} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=fft $ ${TEST_CONV_VERBOSE_B} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_solver_ConvDirectNaiveConv_F GFX103X_ENABLED GFX110X_ENABLED HALF_ENABLED BF16_ENABLED INT8_ENABLED +add_custom_test(smoke_solver_ConvDirectNaiveConv_F GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED HALF_ENABLED BF16_ENABLED INT8_ENABLED COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DRIVER_USE_GPU_REFERENCE=0 MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvFwd $ ${TEST_CONV_VERBOSE_F} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_solver_ConvDirectNaiveConv_BW GFX103X_ENABLED GFX110X_ENABLED HALF_ENABLED BF16_ENABLED +add_custom_test(smoke_solver_ConvDirectNaiveConv_BW GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED HALF_ENABLED BF16_ENABLED COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DRIVER_USE_GPU_REFERENCE=0 MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvBwd $ ${TEST_CONV_VERBOSE_B} --input 1 16 14 14 --weights 48 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DRIVER_USE_GPU_REFERENCE=0 MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvDirectNaiveConvWrw $ @@ -1933,7 +1934,7 @@ add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlops GFX900_DISABLED ${TEST_CONV_VERBOSE_F} --input 64 512 7 7 --weights 128 128 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_fp32_fp16 GFX900_DISABLED GFX906_DISABLED +add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_fp32_fp16 GFX900_DISABLED GFX906_DISABLED GFX94X_ENABLED HALF_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC $ @@ -1949,7 +1950,7 @@ add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_fp32_fp16 G --in_layout NHWC --fil_layout NHWC --out_layout NHWC ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16 GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED +add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16 GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED GFX94X_ENABLED FLOAT_DISABLED BF16_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdXdlopsNHWC $ @@ -1965,7 +1966,7 @@ add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicXdlopsNHWC_bf16 GFX900 --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 +add_custom_test(smoke_solver_ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC GFX900_DISABLED GFX906_DISABLED GFX908_DISABLED GFX90A_DISABLED GFX94X_ENABLED GFX103X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicFwdDlopsNCHWC $ @@ -2095,7 +2096,7 @@ add_custom_test(smoke_solver_ConvHipImplicitGemmForwardV4R5Xdlops GFX900_DISABLE ${TEST_CONV_VERBOSE_F} --input 128 16 54 54 --weights 64 16 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_solver_ConvHipImplicitGemmFwdXdlops OCL_DISABLED GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED HALF_ENABLED INT8_ENABLED +add_custom_test(smoke_solver_ConvHipImplicitGemmFwdXdlops OCL_DISABLED GFX900_DISABLED GFX906_DISABLED GFX90A_DISABLED GFX94X_ENABLED HALF_ENABLED INT8_ENABLED SKIP_UNLESS_COMPOSABLEKERNEL TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvHipImplicitGemmFwdXdlops $ @@ -2129,7 +2130,7 @@ add_custom_test(smoke_solver_ConvBinWinogradRxS_fp32 GFX90A_DISABLED SKIP_XNACK_ ) # FP16 ALT attribute is disabled to enable the backward solver on MI200 for HALF. -add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3g1_f16 GFX103X_ENABLED GFX110X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3g1_f16 GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ ${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} @@ -2141,7 +2142,7 @@ add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3g1_f16 GFX103X_ENABLED GFX110 ${TEST_CONV_VERBOSE_W} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3g1_f32 GFX103X_ENABLED GFX110X_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3g1_f32 GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ @@ -2151,7 +2152,7 @@ add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3g1_f32 GFX103X_ENABLED GFX110 ) # FP16 ALT attribute is disabled to enable the backward solver on MI200 for HALF. -add_custom_test(smoke_solver_ConvBinWinogradRxSf3x2_f16 GFX103X_ENABLED GFX110X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvBinWinogradRxSf3x2_f16 GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED FLOAT_DISABLED HALF_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ ${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} @@ -2163,7 +2164,7 @@ add_custom_test(smoke_solver_ConvBinWinogradRxSf3x2_f16 GFX103X_ENABLED GFX110X_ ${TEST_CONV_VERBOSE_W} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_solver_ConvBinWinogradRxSf3x2_f32 GFX103X_ENABLED GFX110X_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvBinWinogradRxSf3x2_f32 GFX94X_ENABLED GFX103X_ENABLED GFX110X_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ ${TEST_CONV_VERBOSE_F} --input 1 20 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} COMMAND MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $ @@ -2182,7 +2183,7 @@ add_custom_test(smoke_solver_ConvWinogradFuryRxSf2x3_f16 GFX900_DISABLED GFX906_ ) # FP16 ALT attribute is disabled to enable the backward solver on MI200 for HALF. -add_custom_test(smoke_solver_ConvWinograd3x3MultipassWrW GFX94X_DISABLED HALF_ENABLED BF16_ENABLED SKIP_XNACK_ON OCL_DISABLED +add_custom_test(smoke_solver_ConvWinograd3x3MultipassWrW HALF_ENABLED BF16_ENABLED SKIP_XNACK_ON OCL_DISABLED COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER='ConvWinograd3x3MultipassWrW<3-2>' $ ${TEST_CONV_VERBOSE_W} --input 1 16 24 24 --weights 16 16 3 3 --pads_strides_dilations 1 1 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} @@ -2224,20 +2225,20 @@ add_custom_test(smoke_solver_ConvWinograd3x3MultipassWrW GFX94X_DISABLED HALF_EN ${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 GFX103X_ENABLED HALF_ENABLED SKIP_XNACK_ON TEST_TUNING +add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3 GFX900_DISABLED GFX94X_ENABLED GFX103X_ENABLED HALF_ENABLED SKIP_XNACK_ON TEST_TUNING COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3 $ --input 1 40 20 20 --weights 20 20 3 3 --pads_strides_dilations 1 1 1 1 1 1 --group-count 2 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3g1 GFX900_DISABLED GFX103X_ENABLED HALF_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvBinWinogradRxSf2x3g1 GFX900_DISABLED GFX94X_ENABLED GFX103X_ENABLED HALF_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf2x3g1 $ --input 1 40 20 20 --weights 20 40 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) -add_custom_test(smoke_solver_ConvBinWinogradRxSf3x2 GFX900_DISABLED GFX103X_ENABLED HALF_ENABLED SKIP_XNACK_ON +add_custom_test(smoke_solver_ConvBinWinogradRxSf3x2 GFX900_DISABLED GFX94X_ENABLED GFX103X_ENABLED HALF_ENABLED SKIP_XNACK_ON COMMAND MIOPEN_FIND_ENFORCE=SEARCH_DB_UPDATE MIOPEN_DEBUG_TUNING_ITERATIONS_MAX=5 MIOPEN_DEBUG_CONVOLUTION_ATTRIB_FP16_ALT_IMPL=0 MIOPEN_FIND_MODE=normal MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvBinWinogradRxSf3x2 $