From f525e845367b1be986eb71ee110c3ae52e247727 Mon Sep 17 00:00:00 2001 From: Daniel Lowell Date: Tue, 6 Jul 2021 07:49:39 -0500 Subject: [PATCH] Reshuffle and reduce test cases (#147) - Limit the number of combinations for a single dimension for pooling and convolution tests - Resolves "[PR Testing] Get rid of test redundancy" #816 - Resolves "[COMGR] Code quality: reference binding to null pointer of type 'char'" #877 - Tests: Test generator now includes a batch greater than 1 and able to variate count of tests using --limit - Tests: Various improvements in tests/CMakeLists. Fixed LONG_TESTS, added information about skipped tests. - CI: Refactored Jenkinsfile, reshuffled test stage sequence - Added W/A: "[COMGR][debug][test_gpu_reference_kernel] compiler errors" #898 - Added W/A: "[iGemmfwd][test_conv2d][gfx906][half] Verification failed" #936 --- Jenkinsfile | 830 ++++++++++++--------------------------- src/ocl/batchnormocl.cpp | 2 +- test/CMakeLists.txt | 209 +++++++--- test/conv2d.cpp | 22 +- test/conv3d.cpp | 27 +- test/conv_common.hpp | 152 ++++++- test/driver.hpp | 39 ++ test/immed_conv2d.cpp | 26 +- test/immed_conv3d.cpp | 27 +- test/network_data.hpp | 17 - test/pooling2d.cpp | 29 +- test/pooling3d.cpp | 3 +- 12 files changed, 701 insertions(+), 682 deletions(-) diff --git a/Jenkinsfile b/Jenkinsfile index a2c120cb12..c9fee9e3f4 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -2,6 +2,7 @@ def rocmnode(name) { return 'rocmtest && miopen && ' + name } + def show_node_info() { sh """ echo "NODE_NAME = \$NODE_NAME" @@ -12,69 +13,115 @@ def show_node_info() { """ } -def cmake_build(compiler, flags, env4make, extradebugflags, prefixpath){ - def workspace_dir = pwd() - def vcache = "/var/jenkins/.cache/miopen/vcache" - def archive = (flags == '-DCMAKE_BUILD_TYPE=release') - def config_targets = "check doc MIOpenDriver" - def test_flags = "--disable-verification-cache" - def debug_flags = "-g ${extradebugflags} -fno-omit-frame-pointer -fsanitize=undefined -fno-sanitize-recover=undefined" - def compilerpath = "" - def configargs = "" +//default +// CXX=/opt/rocm/llvm/bin/clang++ CXXFLAGS='-Werror' cmake -DMIOPEN_GPU_SYNC=Off -DCMAKE_PREFIX_PATH=/usr/local -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release .. +// +def cmake_build(Map conf=[:]){ + + def compiler = conf.get("compiler","/opt/rocm/llvm/bin/clang++") + def config_targets = conf.get("config_targets","check") + def debug_flags = "-g -fno-omit-frame-pointer -fsanitize=undefined -fno-sanitize-recover=undefined " + conf.get("extradebugflags", "") + def build_envs = "CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 " + conf.get("build_env","") + def prefixpath = conf.get("prefixpath","/usr/local") + def setup_args = " -DMIOPEN_GPU_SYNC=Off " + conf.get("setup_flags","") + + if (prefixpath != "/usr/local"){ + setup_args = setup_args + " -DCMAKE_PREFIX_PATH=${prefixpath} " + } + + def build_type_debug = (conf.get("build_type",'release') == 'debug') - compilerpath = compiler; - if (prefixpath != "/usr/local") { - configargs = "-DCMAKE_PREFIX_PATH=${prefixpath}" + //cmake_env can overwrite default CXX variables. + def cmake_envs = "CXX=${compiler} CXXFLAGS='-Werror' " + conf.get("cmake_ex_env","") + + def package_build = (conf.get("package_build","") == "true") + + if (package_build == true) { + config_targets = "package" } - if(!flags.contains('-DBUILD_DEV=On')) + if(conf.get("build_install","") == "true") { - config_targets = 'install ' + config_targets - flags = '-DCMAKE_INSTALL_PREFIX=../install ' + flags + config_targets = 'install ' + config_targets + setup_args = ' -DBUILD_DEV=Off -DCMAKE_INSTALL_PREFIX=../install' + setup_args + } else{ + setup_args = ' -DBUILD_DEV=On' + setup_args } - if (archive == true) { - config_targets = "package" + // test_flags = ctest -> MIopen flags + def test_flags = conf.get("test_flags","") + + if (conf.get("vcache_enable","") == "true"){ + def vcache = conf.get(vcache_path,"/var/jenkins/.cache/miopen/vcache") + build_envs = " MIOPEN_VERIFY_CACHE_PATH='${vcache}' " + build_envs + } else{ + test_flags = " --disable-verification-cache " + test_flags } - def cmd = """ - echo \$HSA_ENABLE_SDMA - ulimit -c unlimited - cd build - CXX=${compilerpath} CXXFLAGS='-Werror' cmake ${configargs} -DMIOPEN_TEST_FLAGS='${test_flags}' -DCMAKE_CXX_FLAGS_DEBUG='${debug_flags}' ${flags} .. - MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_XDLOPS=1 CTEST_PARALLEL_LEVEL=4 MIOPEN_VERIFY_CACHE_PATH=${vcache} MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 ${env4make} dumb-init make -j\$(nproc) ${config_targets} - """ + + if(conf.get("codecov", false)){ //Need + setup_args = " -DCMAKE_BUILD_TYPE=debug -DCMAKE_CXX_FLAGS_DEBUG='${debug_flags} -fprofile-arcs -ftest-coverage' -DCODECOV_TEST=On " + setup_args + }else if(build_type_debug){ + setup_args = " -DCMAKE_BUILD_TYPE=debug -DCMAKE_CXX_FLAGS_DEBUG='${debug_flags}'" + setup_args + }else{ + setup_args = " -DCMAKE_BUILD_TYPE=release" + setup_args + } + + if(test_flags != ""){ + setup_args = "-DMIOPEN_TEST_FLAGS='${test_flags}'" + setup_args + } + + def pre_setup_cmd = """ + echo \$HSA_ENABLE_SDMA + ulimit -c unlimited + rm -rf build + mkdir build + rm -rf install + mkdir install + rm -f src/kernels/*.ufdb.txt + rm -f src/kernels/miopen*.udb + cd build + """ + def setup_cmd = conf.get("setup_cmd", "${cmake_envs} cmake ${setup_args} .. ") + def build_cmd = conf.get("build_cmd", "${build_envs} dumb-init make -j\$(nproc) ${config_targets}") + def execute_cmd = conf.get("execute_cmd", "") + + def cmd = conf.get("cmd", """ + ${pre_setup_cmd} + ${setup_cmd} + ${build_cmd} + ${execute_cmd} + """) + echo cmd sh cmd + // Only archive from master or develop - if (archive == true && (env.BRANCH_NAME == "develop" || env.BRANCH_NAME == "master")) { + if (package_build == true && (env.BRANCH_NAME == "develop" || env.BRANCH_NAME == "master")) { archiveArtifacts artifacts: "build/*.deb", allowEmptyArchive: true, fingerprint: true } } -def buildhipclangjob(Map conf){ +def buildHipClangJob(Map conf=[:]){ show_node_info() env.HSA_ENABLE_SDMA=0 env.CODECOV_TOKEN="aec031be-7673-43b5-9840-d8fb71a2354e" checkout scm - def compiler = conf.get("compiler", "/opt/rocm/llvm/bin/clang++") - def prefixpath = conf.get("prefixpath", "/usr/local") - def flags = conf.get("flags", "") - def env4make = conf.get("env4make", "") + def image = "miopen" - def cmd = conf.get("cmd", "") + def prefixpath = conf.get("prefixpath", "/usr/local") def gpu_arch = conf.get("gpu_arch", "gfx900;gfx906") + + def miotensile_version = conf.get("miotensile_version", "default") def target_id = conf.get("target_id", "OFF") def mlir_build = conf.get("mlir_build", "OFF") - def codecov = conf.get("codecov", false) - def miotensile_version = conf.get("miotensile_version", "default") def dockerOpts="--device=/dev/kfd --device=/dev/dri --group-add video --cap-add=SYS_PTRACE --security-opt seccomp=unconfined" def dockerArgs = "--build-arg PREFIX=${prefixpath} --build-arg GPU_ARCH='${gpu_arch}' --build-arg MIOTENSILE_VER='${miotensile_version}' --build-arg USE_TARGETID='${target_id}' --build-arg USE_MLIR='${mlir_build}' " - def extradebugflags = "" + def variant = env.STAGE_NAME - if (codecov) { - extradebugflags = "-fprofile-arcs -ftest-coverage" - } + + def codecov = conf.get("codecov", false) + def retimage gitStatusWrapper(credentialsId: '7126e5fe-eb51-4576-b52b-9aaf1de8f0fd', gitHubContext: "Jenkins - ${variant}", account: 'ROCmSoftwarePlatform', repo: 'MIOpen') { try { @@ -103,19 +150,8 @@ def buildhipclangjob(Map conf){ withDockerContainer(image: image, args: dockerOpts + ' -v=/var/jenkins/:/var/jenkins') { timeout(time: 5, unit: 'HOURS') { - sh ''' - rm -rf build - mkdir build - rm -rf install - mkdir install - rm -f src/kernels/*.ufdb.txt - rm -f src/kernels/miopen*.udb - ''' - if(cmd == ""){ - cmake_build(compiler, flags, env4make, extradebugflags, prefixpath) - }else{ - sh cmd - } + cmake_build(conf) + if (codecov) { sh ''' cd build @@ -136,10 +172,9 @@ def reboot(){ build job: 'reboot-slaves', propagate: false , parameters: [string(name: 'server', value: "${env.NODE_NAME}"),] } -def runDockerJob(Map conf){ +def buildHipClangJobAndReboot(Map conf=[:]){ try{ - def build_config = conf - buildhipclangjob(build_config) + buildHipClangJob(conf) } catch(e){ echo "throwing error exception for the stage" @@ -147,10 +182,13 @@ def runDockerJob(Map conf){ throw e } finally{ - reboot() + if (!conf.get("no_reboot", false)) { + reboot() + } } } + /// Stage name format: /// [DataType] Backend[/Compiler] BuildType [TestSet] [Target] /// @@ -173,12 +211,17 @@ def runDockerJob(Map conf){ /// Target := { gfx908 | Vega20 | Vega10 | Vega* } /// * "Vega" (gfx906 or gfx900) is the default and usually not specified. + pipeline { agent none options { parallelsAlwaysFailFast() } parameters { + booleanParam( + name: "DISABLE_ALL_STAGES", + defaultValue: false, + description: "Disables each stage in the pipline") booleanParam( name: "STATIC_CHECKS", defaultValue: true, @@ -215,51 +258,52 @@ pipeline { name: "PACKAGES", defaultValue: true, description: "") + booleanParam( + name: "BUILD_PACKAGES", + defaultValue: true, + description: "Run packages stage") + } + + environment{ + extra_log_env = " MIOPEN_LOG_LEVEL=5 " + gfx908_test = " -DMIOPEN_TEST_GFX908=On" + Fp16_flags = " -DMIOPEN_TEST_HALF=On" + Bf16_flags = " -DMIOPEN_TEST_BFLOAT16=On" + Int8_flags = " -DMIOPEN_TEST_INT8=On" + Full_test = " -DMIOPEN_TEST_ALL=On" + //limit greater than 2 leads to prolonged testing more than 5hrs per stage. + Full_test_limit = " -DMIOPEN_TEST_ALL=On -DMIOPEN_TEST_LIMIT=2" + Tensile_build_env = "MIOPEN_DEBUG_HIP_KERNELS=0 " + Tensile_setup = " -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF" } stages{ stage("Static checks"){ - when { expression { params.STATIC_CHECKS } } + when { expression { params.STATIC_CHECKS && !params.DISABLE_ALL_STAGES } } parallel{ stage('Hip Tidy') { agent{ label rocmnode("nogpu") } environment{ - cmd = "cd build; CXX='/opt/rocm/llvm/bin/clang++' cmake -DMIOPEN_BACKEND=HIP -DBUILD_DEV=On ..; make -j\$(nproc) -k analyze;" + setup_cmd = "CXX='/opt/rocm/llvm/bin/clang++' cmake -DMIOPEN_BACKEND=HIP -DBUILD_DEV=On .. " + build_cmd = "make -j\$(nproc) -k analyze" } steps{ - script{ - try{ - buildhipclangjob(compiler: 'clang++', flags: '-DCMAKE_BUILD_TYPE=release', cmd: cmd) - } - catch(e){ - echo "throwing error exception for the stage" - echo 'Exception occurred: ' + e.toString() - throw e - } - } + buildHipClangJobAndReboot(setup_cmd: setup_cmd, build_cmd: build_cmd, no_reboot:true) } } stage('OpenCL Tidy') { agent{ label rocmnode("nogpu") } environment{ - cmd = "cd build; cmake -DMIOPEN_BACKEND=OpenCL -DBUILD_DEV=On ..; make -j\$(nproc) -k analyze;" + setup_cmd = "cmake -DMIOPEN_BACKEND=OpenCL -DBUILD_DEV=On .." + build_cmd = "make -j\$(nproc) -k analyze" } steps{ - script{ - try{ - buildhipclangjob(compiler: 'g++', flags: '-DCMAKE_BUILD_TYPE=release', cmd: cmd) - } - catch(e){ - echo "throwing error exception for the stage" - echo 'Exception occurred: ' + e.toString() - throw e - } - } + buildHipClangJobAndReboot(setup_cmd: setup_cmd, build_cmd: build_cmd, no_reboot:true) } } stage('Clang Format') { agent{ label rocmnode("nogpu") } environment{ - cmd = "find . -iname \'*.h\' \ + execute_cmd = "find . -iname \'*.h\' \ -o -iname \'*.hpp\' \ -o -iname \'*.cpp\' \ -o -iname \'*.h.in\' \ @@ -270,778 +314,414 @@ pipeline { | xargs -n 1 -P 1 -I{} -t sh -c \'clang-format-3.8 -style=file {} | diff - {}\'" } steps{ - script{ - try{ - buildhipclangjob(compiler: 'clang++', flags: '-DCMAKE_BUILD_TYPE=release', cmd: cmd) - } - catch(e){ - echo "throwing error exception for the stage" - echo 'Exception occurred: ' + e.toString() - throw e - } - } + buildHipClangJobAndReboot(setup_cmd: "", build_cmd: "", execute_cmd: execute_cmd, no_reboot:true) } } } } stage("Smoke Fp32"){ - when { expression { params.SMOKE_FP32_AUX1 } } + when { expression { params.SMOKE_FP32_AUX1 && !params.DISABLE_ALL_STAGES } } + environment{ + Smoke_targets = "check doc MIOpenDriver" + } parallel{ - stage('Fp32 OpenCL Debug') { + stage('Fp32 OpenCL Debug + Codecov') { agent{ label rocmnode("vega") } steps{ - script{ - runDockerJob(compiler: 'g++', flags: '-DBUILD_DEV=On -DCMAKE_BUILD_TYPE=debug') - } + buildHipClangJobAndReboot(compiler: 'g++', build_type: 'debug', config_targets: Smoke_targets, codecov: true) } } stage('Fp32 OpenCL gfx908') { agent{ label rocmnode("gfx908") } steps{ - script{ - runDockerJob(compiler: 'g++', flags: '-DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_GFX908=On', gpu_arch: "gfx908") - } - } - } - stage('Fp32 Hip') { - agent{ label rocmnode("vega") } - steps{ - script{ - runDockerJob(flags: '-DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release') - } + buildHipClangJobAndReboot(compiler: 'g++', config_targets: Smoke_targets, setup_flags: gfx908_test, gpu_arch: "gfx908") } } stage('Fp32 Hip /opt/rocm') { agent{ label rocmnode("vega") } steps{ - script{ - runDockerJob(flags: '-DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release', prefixpath: '/opt/rocm') - } + buildHipClangJobAndReboot(prefixpath: '/opt/rocm', config_targets: Smoke_targets) } } stage('Fp32 Hip Debug') { agent{ label rocmnode("vega") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=debug -DMIOPEN_TEST_FLAGS=--disable-verification-cache .. - CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd) - } + buildHipClangJobAndReboot(build_type: 'debug', config_targets: Smoke_targets) } } stage('Fp32 Hip Debug gfx908 /opt/rocm') { agent{ label rocmnode("gfx908") } + environment{ + gpu_arch = "gfx908" + prefixpath = "/opt/rocm" + } steps{ - script{ - runDockerJob(flags: '-DMIOPEN_TEST_GFX908=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=debug', prefixpath: '/opt/rocm', gpu_arch: "gfx908") - } + buildHipClangJobAndReboot(prefixpath: prefixpath, build_type: 'debug', setup_flags: gfx908_test, config_targets: Smoke_targets, gpu_arch: gpu_arch) } } - } - } - stage("Smoke Aux 1"){ - when { expression { params.SMOKE_FP32_AUX1 } } - parallel{ stage('Fp32 HipNoGPU Debug') { agent{ label rocmnode("nogpu") } environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=debug -DMIOPEN_BACKEND=HIPNOGPU -DMIOPEN_INSTALL_CXX_HEADERS=On .. - make -j\$(nproc) - """ + HipNoGPU_flags = "-DMIOPEN_BACKEND=HIPNOGPU -DMIOPEN_INSTALL_CXX_HEADERS=On" + build_cmd = "make -j\$(nproc)" } steps{ - buildhipclangjob(env4make: "MIOPEN_LOG_LEVEL=5 MIOPEN_COMPILE_PARALLEL_LEVEL=1", cmd: cmd) + buildHipClangJob( build_type: 'debug', setup_flags: HipNoGPU_flags, build_cmd: build_cmd) } } + } + } + stage("Smoke Aux 1"){ + when { expression { params.SMOKE_FP32_AUX1 && !params.DISABLE_ALL_STAGES } } + parallel{ stage('Fp32 Hip Debug COMGR') { agent{ label rocmnode("vega") } environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_USE_COMGR=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=debug -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - CTEST_PARALLEL_LEVEL=2 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 MIOPEN_LOG_LEVEL=5 make -j\$(nproc) check - """ + COMGR_build_cmd = "CTEST_PARALLEL_LEVEL=2 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 MIOPEN_LOG_LEVEL=5 make -j\$(nproc) check" } steps{ - script{ - runDockerJob(cmd: cmd) - } + buildHipClangJobAndReboot( build_type: 'debug', setup_flags: "-DMIOPEN_USE_COMGR=On", build_cmd: COMGR_build_cmd, test_flags: ' --verbose ') } } stage('Fp32 Hip Debug Embedded Vega20') { agent{ label rocmnode("vega20") } environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_EMBED_DB="gfx906_60;gfx906_64" -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=debug -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_LOG_LEVEL=5 make -j\$(nproc) check - """ + Embedded_flags = "-DMIOPEN_EMBED_DB='gfx906_60;gfx906_64'" } steps{ - script{ - runDockerJob(cmd: cmd) - } + buildHipClangJobAndReboot( build_type: 'debug', setup_flags: Embedded_flags, build_env: extra_log_env, test_flags: ' --verbose ') } } stage('Fp32 Hip Static') { agent{ label rocmnode("vega") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DBUILD_DEV=On -DBUILD_SHARED_LIBS=Off -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_FLAGS=--disable-verification-cache .. - CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd) - } + buildHipClangJobAndReboot( setup_flags: "-DBUILD_SHARED_LIBS=Off") } } stage('Fp32 Hip Normal-Find') { agent{ label rocmnode("vega") } environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release .. - make -j test_conv2d - MIOPEN_FIND_MODE=normal CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 bin/test_conv2d --disable-verification-cache - """ + config_targets = "test_conv2d" + execute_cmd = "MIOPEN_FIND_MODE=1 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 bin/test_conv2d --disable-verification-cache" } steps{ - script{ - runDockerJob(cmd: cmd) - } + buildHipClangJobAndReboot(config_targets: config_targets, execute_cmd: execute_cmd) } } stage('Fp32 Hip Fast-Find') { agent{ label rocmnode("vega") } environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release .. - make -j test_conv2d - MIOPEN_FIND_MODE=fast CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 bin/test_conv2d --disable-verification-cache - """ + config_targets = "test_conv2d" + execute_cmd = "MIOPEN_FIND_MODE=2 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 bin/test_conv2d --disable-verification-cache" } steps{ - script{ - runDockerJob(cmd: cmd) - } + buildHipClangJobAndReboot( config_targets: config_targets, execute_cmd: execute_cmd) + } + } + stage('Fp32 Hip') { + agent{ label rocmnode("vega") } + steps{ + buildHipClangJobAndReboot() } } } } stage("Smoke MLIR"){ - when { expression { params.SMOKE_MLIR } } + when { expression { params.SMOKE_MLIR && !params.DISABLE_ALL_STAGES } } + environment{ MLIR_flags = "-DMIOPEN_USE_MLIR=On" } parallel{ stage('Fp32 Hip MLIR') { agent{ label rocmnode("vega") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_USE_MLIR=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - CTEST_PARALLEL_LEVEL=4 MIOPEN_LOG_LEVEL=5 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx900;gfx906", mlir_build: "ON") - } + buildHipClangJobAndReboot(setup_flags: MLIR_flags, build_env: extra_log_env, test_flags: ' --verbose ', mlir_build: "ON") } } stage('Fp16 Hip MLIR') { agent{ label rocmnode("vega") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_TEST_HALF=On -DMIOPEN_USE_MLIR=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - CTEST_PARALLEL_LEVEL=4 MIOPEN_LOG_LEVEL=5 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx900;gfx906", mlir_build: "ON") - } + buildHipClangJobAndReboot(setup_flags: MLIR_flags + Fp16_flags, build_env: extra_log_env, test_flags: ' --verbose ', mlir_build: "ON") } } - stage('Fp32 Hip MLIR Xdlops') { - agent{ label rocmnode("gfx908") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_USE_MLIR=On -DMIOPEN_TEST_GFX908=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - CTEST_PARALLEL_LEVEL=4 MIOPEN_LOG_LEVEL=5 make -j\$(nproc) check - """ - } + stage('Fp16 Hip MLIR Xdlops') { + agent{ label rocmnode("vega") } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx908", mlir_build: "ON") - } + buildHipClangJobAndReboot(setup_flags: MLIR_flags + Fp16_flags, build_env: extra_log_env, test_flags: ' --verbose ', gpu_arch: "gfx908", mlir_build: "ON") } } - stage('Fp16 Hip MLIR Xdlops') { - agent{ label rocmnode("gfx908") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_TEST_HALF=On -DMIOPEN_USE_MLIR=On -DMIOPEN_TEST_GFX908=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - CTEST_PARALLEL_LEVEL=4 MIOPEN_LOG_LEVEL=5 make -j\$(nproc) check - """ - } + stage('Fp32 Hip MLIR Xdlops') { + agent{ label rocmnode("vega") } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx908", mlir_build: "ON") - } + buildHipClangJobAndReboot(setup_flags: MLIR_flags, build_env: extra_log_env, test_flags: ' --verbose ', gpu_arch: "gfx908", mlir_build: "ON") } } } } stage("Smoke Fp16/Bf16/Int8"){ - when { expression { params.SMOKE_FP16_BF16_INT8 } } + when { expression { params.SMOKE_FP16_BF16_INT8 && !params.DISABLE_ALL_STAGES } } + environment{ + Smoke_targets = "check doc MIOpenDriver" + } parallel{ - stage('Fp16 Hip Vega20 /opt/rocm') { + stage('Fp16 OpenCL Vega20') { agent{ label rocmnode("vega20") } steps{ - script{ - runDockerJob(flags: '-DMIOPEN_TEST_HALF=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release', prefixpath: '/opt/rocm') - } + buildHipClangJobAndReboot(compiler: 'g++', setup_flags: Fp16_flags, config_targets: Smoke_targets) } } - stage('Fp16 OpenCL Vega20') { + stage('Int8 OpenCL Vega20') { agent{ label rocmnode("vega20") } steps{ - script{ - runDockerJob(compiler: 'g++', flags: '-DMIOPEN_TEST_HALF=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release') - } + buildHipClangJobAndReboot(compiler: 'g++', setup_flags: Int8_flags, config_targets: Smoke_targets) } } - stage('Int8 OpenCL Vega20') { + stage('Fp16 Hip Vega20 /opt/rocm') { agent{ label rocmnode("vega20") } steps{ - script{ - runDockerJob(compiler: 'g++', flags: '-DMIOPEN_TEST_INT8=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release') - } + buildHipClangJobAndReboot( setup_flags: Fp16_flags, prefixpath: '/opt/rocm', config_targets: Smoke_targets) } } stage('Bf16 Hip Vega20 /opt/rocm') { agent{ label rocmnode("vega20") } steps{ - script{ - runDockerJob(flags: '-DMIOPEN_TEST_BFLOAT16=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release', prefixpath: '/opt/rocm') - } + buildHipClangJobAndReboot(setup_flags: Bf16_flags, prefixpath: '/opt/rocm', config_targets: Smoke_targets) } } - stage('Bf16 Hip Debug gfx908 /opt/rocm') { + stage('Fp16 Hip gfx908 /opt/rocm') { agent{ label rocmnode("gfx908") } steps{ - script{ - runDockerJob(flags: '-DMIOPEN_TEST_BFLOAT16=On -DMIOPEN_TEST_GFX908=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=debug', prefixpath: '/opt/rocm', gpu_arch: "gfx908") - } + buildHipClangJobAndReboot( setup_flags: Fp16_flags + gfx908_test, prefixpath: '/opt/rocm', config_targets: Smoke_targets, gpu_arch: "gfx908") } } - stage('Fp16 Hip Debug gfx908 /opt/rocm') { + stage('Bf16 Hip gfx908 /opt/rocm') { agent{ label rocmnode("gfx908") } steps{ - script{ - runDockerJob(flags: '-DMIOPEN_TEST_HALF=On -DMIOPEN_TEST_GFX908=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=debug', prefixpath: '/opt/rocm', gpu_arch: "gfx908") - } + buildHipClangJobAndReboot(setup_flags: Bf16_flags + gfx908_test, prefixpath: '/opt/rocm', config_targets: Smoke_targets, gpu_arch: "gfx908") } } } } stage("Smoke MIOpenTensile Latest"){ - when { expression { params.SMOKE_MIOPENTENSILE_LATEST } } + environment{ + Tensile_version = "latest" + } parallel{ - stage('Fp16 Hip Tensile-Latest Vega20') { + stage('Fp16 Hip Tensile-Latest All Vega20') { + when { expression { params.SMOKE_MIOPENTENSILE_LATEST && !params.DISABLE_ALL_STAGES } } agent{ label rocmnode("vega20") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_HALF=On -DMIOPEN_GPU_SYNC=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx906:xnack-", miotensile_version: "latest", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + Fp16_flags, build_env: Tensile_build_env, test_flags: ' --verbose ', gpu_arch:"gfx906:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } - stage('Int8 Hip Tensile-Latest Vega20') { + stage('Int8 Hip Tensile-Latest All Vega20') { + when { expression { params.SMOKE_MIOPENTENSILE_LATEST && !params.DISABLE_ALL_STAGES } } agent{ label rocmnode("vega20") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_TEST_INT8=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_GPU_SYNC=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 MIOPEN_LOG_LEVEL=5 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx906:xnack-", miotensile_version: "latest", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + extra_log_env + Int8_flags, build_env: Tensile_build_env, test_flags: ' --verbose ', gpu_arch:"gfx906:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } - stage('Fp32 Hip Tensile-Latest gfx908') { + + stage('Fp32 Hip Tensile-Latest All gfx908') { + when { expression { params.SMOKE_MIOPENTENSILE_LATEST && !params.DISABLE_ALL_STAGES } } agent{ label rocmnode("gfx908") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_GFX908=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 MIOPEN_LOG_LEVEL=5 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx908:xnack-", miotensile_version: "latest", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + extra_log_env + gfx908_test , build_env: Tensile_build_env + extra_log_env, gpu_arch: "gfx908:xnack-", test_flags: ' --verbose ', miotensile_version: Tensile_version, target_id: "ON") } } - stage('Bf16 Hip Tensile-Latest gfx908') { + stage('Bf16 Hip Tensile-Latest All gfx908') { + when { expression { params.SMOKE_MIOPENTENSILE_LATEST && !params.DISABLE_ALL_STAGES } } agent{ label rocmnode("gfx908") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_TEST_BFLOAT16=On -DMIOPEN_TEST_GFX908=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_GPU_SYNC=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 MIOPEN_LOG_LEVEL=5 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx908:xnack-", miotensile_version: "latest", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + extra_log_env + gfx908_test + Bf16_flags, build_env: Tensile_build_env + extra_log_env, test_flags: ' --verbose ', gpu_arch: "gfx908:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } } } - stage("Full tests I"){ - when { expression { params.FULL_TESTS } } + stage("Full Tests I"){ + when { expression { params.FULL_TESTS && !params.DISABLE_ALL_STAGES } } parallel{ - stage('Fp32 OpenCL Debug + Codecov') { - agent{ label rocmnode("vega") } + stage('Int8 HIP All Vega20 /opt/rocm') { + agent{ label rocmnode("vega20") } steps{ - script{ - runDockerJob(compiler: 'g++', flags: '-DBUILD_DEV=On -DCMAKE_BUILD_TYPE=debug', codecov: true) - } + buildHipClangJobAndReboot( setup_flags: Int8_flags + Full_test_limit, prefixpath: '/opt/rocm') } } - stage('Int8 Hip All Vega20 /opt/rocm') { - agent{ label rocmnode("vega20") } + stage('Fp32 OpenCL Install All') { + agent{ label rocmnode("vega") } steps{ - script{ - runDockerJob(flags: '-DMIOPEN_TEST_INT8=On -DBUILD_DEV=On -DMIOPEN_TEST_ALL=On -DCMAKE_BUILD_TYPE=release', prefixpath: '/opt/rocm') - } + buildHipClangJobAndReboot(compiler: 'g++', setup_flags: Full_test_limit, build_install: "true") } } stage('Bf16 Hip Install All gfx908') { agent{ label rocmnode("gfx908") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_TEST_BFLOAT16=On -DMIOPEN_TEST_GFX908=On -DMIOPEN_TEST_ALL=On -DBUILD_DEV=Off -DCMAKE_INSTALL_PREFIX=../install -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_LOG_LEVEL=5 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) install check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx908") - } + buildHipClangJobAndReboot(setup_flags: Bf16_flags + Full_test_limit + gfx908_test, build_install: "true", gpu_arch: "gfx908") } } } } - stage("Full tests II"){ - when { expression { params.FULL_TESTS } } + + stage("Full Tests II"){ + when { expression { params.FULL_TESTS && !params.DISABLE_ALL_STAGES } } + environment{ + WORKAROUND_iGemm_936 = " MIOPEN_DEBUG_CONV_IMPLICIT_GEMM_HIP_FWD_V4R1=0" + } parallel{ - stage('Fp32 OpenCL Install All') { - agent{ label rocmnode("vega") } - steps{ - script{ - runDockerJob(compiler: 'g++', flags: '-DBUILD_DEV=Off -DMIOPEN_TEST_ALL=On -DCMAKE_BUILD_TYPE=release') - } - } - } stage('Fp32 Hip All gfx908') { agent{ label rocmnode("gfx908") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_TEST_GFX908=On -DMIOPEN_TEST_ALL=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_LOG_LEVEL=5 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx908") - } + buildHipClangJobAndReboot(setup_flags: Full_test_limit + gfx908_test, build_env: WORKAROUND_iGemm_936, gpu_arch: "gfx908") } } - stage('Fp16 Hip Install All gfx908') { - agent{ label rocmnode("gfx908") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_TEST_HALF=On -DMIOPEN_TEST_GFX908=On -DMIOPEN_TEST_ALL=On -DBUILD_DEV=Off -DCMAKE_INSTALL_PREFIX=../install -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_LOG_LEVEL=5 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) install check - """ - } - steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx908") - } - } - } - } - } - stage("Full tests III"){ - when { expression { params.FULL_TESTS } } - parallel{ stage('Fp16 Hip Install All Vega20') { agent{ label rocmnode("vega20") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DBUILD_DEV=Off -DCMAKE_INSTALL_PREFIX=../install -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_HALF=On -DMIOPEN_TEST_ALL=On -DMIOPEN_TEST_FLAGS=--disable-verification-cache .. - CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) install check - """ - } steps{ - script{ - runDockerJob(cmd: cmd) - } + buildHipClangJobAndReboot( setup_flags: Full_test_limit + Fp16_flags, build_env: WORKAROUND_iGemm_936, build_install: "true") } } stage('Fp32 Hip All Vega20') { agent{ label rocmnode("vega20") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_ALL=On -DMIOPEN_TEST_FLAGS=--disable-verification-cache .. - CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ + steps{ + buildHipClangJobAndReboot( setup_flags: Full_test_limit, build_env: WORKAROUND_iGemm_936) } + } + stage('Fp16 Hip All Install gfx908') { + agent{ label rocmnode("gfx908") } steps{ - script{ - runDockerJob(cmd: cmd) - } + buildHipClangJobAndReboot(setup_flags: Full_test_limit + gfx908_test + Fp16_flags, build_env: WORKAROUND_iGemm_936, build_install: "true", gpu_arch: "gfx908") } } } } + stage("MIOpenTensile"){ - when { expression { params.MIOPENTENSILE } } + when { expression { params.MIOPENTENSILE && !params.DISABLE_ALL_STAGES } } + environment{ + Tensile_version = "default" + } parallel{ stage('Fp32 Hip Tensile All Vega20') { agent{ label rocmnode("vega20") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_ALL=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx906:xnack-", miotensile_version: "default", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + Full_test_limit, build_env: Tensile_build_env, test_flags: ' --verbose ', gpu_arch:"gfx906:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } stage('Fp16 Hip Tensile All Vega20') { agent{ label rocmnode("vega20") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_HALF=On -DMIOPEN_GPU_SYNC=On -DMIOPEN_TEST_ALL=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx906:xnack-", miotensile_version: "default", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + Full_test_limit + Fp16_flags, build_env: Tensile_build_env, test_flags: ' --verbose ', gpu_arch:"gfx906:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } stage('Bf16 Hip Tensile All Vega20') { agent{ label rocmnode("vega20") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_TEST_BFLOAT16=On -DMIOPEN_TEST_ALL=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_GPU_SYNC=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 MIOPEN_LOG_LEVEL=5 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx906:xnack-", miotensile_version: "default", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + Full_test_limit + extra_log_env + Bf16_flags, build_env: Tensile_build_env, test_flags: ' --verbose ', gpu_arch:"gfx906:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } stage('Int8 Hip Tensile All Vega20') { agent{ label rocmnode("vega20") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_TEST_INT8=On -DMIOPEN_TEST_ALL=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_GPU_SYNC=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 MIOPEN_LOG_LEVEL=5 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx906:xnack-", miotensile_version: "default", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + Full_test_limit + extra_log_env + Int8_flags, build_env: Tensile_build_env, test_flags: ' --verbose ', gpu_arch:"gfx906:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } stage('Fp32 Hip Tensile All gfx908') { agent{ label rocmnode("gfx908") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_GFX908=On -DMIOPEN_TEST_ALL=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 MIOPEN_LOG_LEVEL=5 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx908:xnack-", miotensile_version: "default", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + Full_test_limit + extra_log_env + gfx908_test , build_env: Tensile_build_env + extra_log_env, test_flags: ' --verbose ', gpu_arch: "gfx908:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } stage('Fp16 Hip Tensile All gfx908') { agent{ label rocmnode("gfx908") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_TEST_HALF=On -DMIOPEN_TEST_GFX908=On -DMIOPEN_TEST_ALL=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_GPU_SYNC=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 MIOPEN_LOG_LEVEL=5 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx908:xnack-", miotensile_version: "default", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + Full_test_limit + extra_log_env + gfx908_test + Fp16_flags, build_env: Tensile_build_env + extra_log_env, test_flags: ' --verbose ', gpu_arch: "gfx908:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } stage('Bf16 Hip Tensile All gfx908') { agent{ label rocmnode("gfx908") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_TEST_BFLOAT16=On -DMIOPEN_TEST_GFX908=On -DMIOPEN_TEST_ALL=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_GPU_SYNC=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 MIOPEN_LOG_LEVEL=5 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx908:xnack-", miotensile_version: "default", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + Full_test_limit + extra_log_env + gfx908_test + Bf16_flags, build_env: Tensile_build_env + extra_log_env, test_flags: ' --verbose ', gpu_arch: "gfx908:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } stage('Int8 Hip Tensile All gfx908') { agent{ label rocmnode("gfx908") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_TEST_INT8=On -DMIOPEN_TEST_GFX908=On -DMIOPEN_TEST_ALL=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_GPU_SYNC=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 MIOPEN_LOG_LEVEL=5 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx908:xnack-", miotensile_version: "default", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + Full_test_limit + extra_log_env + gfx908_test + Int8_flags, build_env: Tensile_build_env + extra_log_env, test_flags: ' --verbose ', gpu_arch: "gfx908:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } } } stage("MIOpenTensile Latest"){ - when { expression { params.MIOPENTENSILE_LATEST } } + when { expression { params.MIOPENTENSILE_LATEST && !params.DISABLE_ALL_STAGES } } + environment{ + Tensile_version = "latest" + } parallel{ stage('Fp32 Hip Tensile-Latest All Vega20') { agent{ label rocmnode("vega20") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_ALL=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx906:xnack-", miotensile_version: "latest", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + Full_test_limit, build_env: Tensile_build_env, test_flags: ' --verbose ', gpu_arch:"gfx906:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } stage('Fp16 Hip Tensile-Latest All Vega20') { agent{ label rocmnode("vega20") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_HALF=On -DMIOPEN_GPU_SYNC=On -DMIOPEN_TEST_ALL=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx906:xnack-", miotensile_version: "latest", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + Full_test_limit + Fp16_flags, build_env: Tensile_build_env, test_flags: ' --verbose ', gpu_arch:"gfx906:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } stage('Bf16 Hip Tensile-Latest All Vega20') { agent{ label rocmnode("vega20") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_TEST_BFLOAT16=On -DMIOPEN_TEST_ALL=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_GPU_SYNC=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 MIOPEN_LOG_LEVEL=5 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx906:xnack-", miotensile_version: "latest", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + Full_test_limit + extra_log_env + Bf16_flags, build_env: Tensile_build_env, test_flags: ' --verbose ', gpu_arch:"gfx906:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } stage('Int8 Hip Tensile-Latest All Vega20') { agent{ label rocmnode("vega20") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_TEST_INT8=On -DMIOPEN_TEST_ALL=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_GPU_SYNC=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 MIOPEN_LOG_LEVEL=5 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx906:xnack-", miotensile_version: "latest", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + Full_test_limit + extra_log_env + Int8_flags, build_env: Tensile_build_env, test_flags: ' --verbose ', gpu_arch:"gfx906:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } stage('Fp32 Hip Tensile-Latest All gfx908') { agent{ label rocmnode("gfx908") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_TEST_GFX908=On -DMIOPEN_TEST_ALL=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 MIOPEN_LOG_LEVEL=5 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx908:xnack-", miotensile_version: "latest", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + Full_test_limit + extra_log_env + gfx908_test , build_env: Tensile_build_env + extra_log_env, test_flags: ' --verbose ', gpu_arch: "gfx908:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } stage('Fp16 Hip Tensile-Latest All gfx908') { agent{ label rocmnode("gfx908") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_TEST_HALF=On -DMIOPEN_TEST_GFX908=On -DMIOPEN_TEST_ALL=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_GPU_SYNC=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 MIOPEN_LOG_LEVEL=5 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx908:xnack-", miotensile_version: "latest", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + Full_test_limit + extra_log_env + gfx908_test + Fp16_flags, build_env: Tensile_build_env + extra_log_env, test_flags: ' --verbose ', gpu_arch: "gfx908:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } stage('Bf16 Hip Tensile-Latest All gfx908') { agent{ label rocmnode("gfx908") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_TEST_BFLOAT16=On -DMIOPEN_TEST_GFX908=On -DMIOPEN_TEST_ALL=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_GPU_SYNC=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 MIOPEN_LOG_LEVEL=5 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx908:xnack-", miotensile_version: "latest", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + Full_test_limit + extra_log_env + gfx908_test + Bf16_flags, build_env: Tensile_build_env + extra_log_env, test_flags: ' --verbose ', gpu_arch: "gfx908:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } stage('Int8 Hip Tensile-Latest All gfx908') { agent{ label rocmnode("gfx908") } - environment{ - cmd = """ - ulimit -c unlimited - cd build - CXX=/opt/rocm/llvm/bin/clang++ cmake -DMIOPEN_TEST_INT8=On -DMIOPEN_TEST_GFX908=On -DMIOPEN_TEST_ALL=On -DBUILD_DEV=On -DCMAKE_BUILD_TYPE=release -DMIOPEN_GPU_SYNC=On -DMIOPEN_TEST_MIOTENSILE=ON -DMIOPEN_USE_MIOPENTENSILE=ON -DMIOPEN_USE_ROCBLAS=OFF -DMIOPEN_TEST_FLAGS='--verbose --disable-verification-cache' .. - MIOPEN_DEBUG_HIP_KERNELS=0 MIOPEN_LOG_LEVEL=5 CTEST_PARALLEL_LEVEL=4 MIOPEN_CONV_PRECISE_ROCBLAS_TIMING=0 make -j\$(nproc) check - """ - } steps{ - script{ - runDockerJob(cmd: cmd, gpu_arch: "gfx908:xnack-", miotensile_version: "latest", target_id: "ON") - } + buildHipClangJobAndReboot( setup_flags: Tensile_setup + Full_test_limit + extra_log_env + gfx908_test + Int8_flags, build_env: Tensile_build_env + extra_log_env, test_flags: ' --verbose ', gpu_arch: "gfx908:xnack-", miotensile_version: Tensile_version, target_id: "ON") } } } } stage("Packages"){ - when { expression { params.PACKAGES } } + when { expression { params.PACKAGES && !params.DISABLE_ALL_STAGES } } parallel { stage('OpenCL Package') { agent{ label rocmnode("nogpu") } steps{ - script{ - runDockerJob(compiler: 'g++', flags: '-DCMAKE_BUILD_TYPE=release', gpu_arch: "gfx900;gfx906;gfx908") - } + buildHipClangJobAndReboot(compiler: 'g++', package_build: "true", gpu_arch: "gfx900;gfx906;gfx908") } } stage("HIP Package /opt/rocm"){ agent{ label rocmnode("nogpu") } steps{ - script{ - runDockerJob(flags: '-DCMAKE_BUILD_TYPE=release', prefixpath: '/opt/rocm', gpu_arch: "gfx900;gfx906;gfx908") - } + buildHipClangJobAndReboot( package_build: "true", prefixpath: '/opt/rocm', gpu_arch: "gfx900;gfx906;gfx908") } } } } } } + diff --git a/src/ocl/batchnormocl.cpp b/src/ocl/batchnormocl.cpp index 972596446b..d034a14ebe 100644 --- a/src/ocl/batchnormocl.cpp +++ b/src/ocl/batchnormocl.cpp @@ -924,7 +924,7 @@ void BatchNormBackward(Handle& handle, unsigned int ldsgcn = 0; unsigned int ldsnogcn = 0; bool single = true; - unsigned int variant = 1; + int variant = 1; //************************************************************************************************* // N*H*W < 32M and H*W > 1024, use batchnorm variant#1 implementation which parallelize diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 0858c56eec..f2c2956637 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -44,6 +44,8 @@ option( MIOPEN_TEST_MLIR "Add tests for MLIR -- EXPERIMENTAL" ${MIOPEN_USE_MLIR} option( WORKAROUND_ISSUE_898 "" ON) option( WORKAROUND_ISSUE_936 "" ON) +#Run the test suite to a depth limit +set(MIOPEN_TEST_LIMIT "0" CACHE STRING "") set(MIOPEN_TEST_FLAGS "" CACHE STRING "") set(MIOPEN_TEST_GDB On CACHE BOOL "") @@ -76,9 +78,6 @@ if(MIOPEN_TEST_GFX908) set(MIOPEN_TEST_VEGA FALSE) endif() -# The usage is non-trivial, see function add_test_command. -set(SKIP_ALL_EXCEPT_TESTS dummy) -set(SKIP_TESTS dummy) # dummy is for REMOVE_DUPLICATES set(MIOPEN_TEST_FLOAT_ARG) @@ -94,43 +93,67 @@ else() set(MIOPEN_TEST_FLOAT TRUE) endif() + if(NOT MIOPEN_TEST_MIOTENSILE) if(MIOPEN_TEST_HALF) if(MIOPEN_BACKEND_OPENCL) set(SKIP_TESTS test_gru test_rnn_vanilla test_lstm) endif() elseif(MIOPEN_TEST_INT8) - set(SKIP_ALL_EXCEPT_TESTS test_tensor_vec test_tensor_cast test_tensor_trans test_tensor_copy test_tensor_set test_tensor_transform test_conv2d) + set(SKIP_ALL_EXCEPT_TESTS + test_tensor_vec test_tensor_cast test_tensor_trans test_tensor_copy test_tensor_set + test_tensor_transform test_conv2d) elseif(MIOPEN_TEST_BFLOAT16) - set(SKIP_ALL_EXCEPT_TESTS test_conv2d test_tensor_copy test_tensor_set test_tensor_vec test_immed_conv2d test_check_numerics_test) - if(MIOPEN_TEST_GFX908) - list(APPEND SKIP_ALL_EXCEPT_TESTS test_conv_extra test_conv_for_implicit_gemm test_miopen_conv test_deepbench_conv) - endif() - endif() + set(SKIP_ALL_EXCEPT_TESTS + test_conv2d test_tensor_copy test_tensor_set test_tensor_vec test_immed_conv2d + test_check_numerics_test test_conv_extra test_conv_for_implicit_gemm test_miopen_conv + test_deepbench_conv) + endif() + if(${CODECOV_TEST}) + list(APPEND SKIP_TESTS test_conv3d test_immed_conv3d test_immed_conv2d test_pooling2d) + # replaced by smaller tests with suffix _codecov + endif() else() if(MIOPEN_TEST_HALF) - set(SKIP_ALL_EXCEPT_TESTS test_conv2d test_conv3d test_immed_conv2d test_immed_conv3d test_gru test_rnn_vanilla test_lstm test_gru_extra test_rnn_extra test_lstm_extra ) + set(SKIP_ALL_EXCEPT_TESTS test_conv2d test_conv3d test_conv3d_extra test_immed_conv2d + test_immed_conv3d test_gru test_rnn_vanilla test_lstm test_gru_extra test_rnn_extra + test_lstm_extra ) elseif(MIOPEN_TEST_INT8) set(SKIP_ALL_EXCEPT_TESTS test_conv2d) elseif(MIOPEN_TEST_BFLOAT16) set(SKIP_ALL_EXCEPT_TESTS test_conv2d test_immed_conv2d) else() - set(SKIP_ALL_EXCEPT_TESTS test_conv2d test_conv3d test_immed_conv2d test_immed_conv3d test_gru test_rnn_vanilla test_lstm test_gru_extra test_rnn_extra test_lstm_extra ) + set(SKIP_ALL_EXCEPT_TESTS test_conv2d test_conv3d test_conv3d_extra test_immed_conv2d + test_immed_conv3d test_gru test_rnn_vanilla test_lstm test_gru_extra + test_rnn_extra test_lstm_extra ) endif() endif() if(MIOPEN_TEST_GFX908) - list(APPEND SKIP_TESTS test_main test_tensor_scale test_tensor_set test_tensor_transform test_tensor_vec test_w_supertensor test_dropout test_immed_conv3d test_conv3d test_soft_max test_fusion_aux test_activation test_lrn_test test_ctc test_conv2d_bias test_conv3d_bias test_cba_inference test_cbna_inference test_pooling2d test_na_train test_na_inference test_bn_aux) + list(APPEND SKIP_TESTS test_main test_tensor_scale test_tensor_set test_tensor_transform + test_tensor_vec test_w_supertensor test_dropout test_immed_conv3d test_conv3d + test_soft_max test_fusion_aux test_activation test_lrn_test test_ctc test_conv2d_bias + test_conv3d_bias test_cba_inference test_cbna_inference test_pooling2d test_na_train + test_na_inference test_bn_aux) +endif() +# The usage is non-trivial, see function add_test_command. +if(SKIP_TESTS) + list(REMOVE_DUPLICATES SKIP_TESTS) endif() +if(SKIP_ALL_EXCEPT_TESTS) + list(REMOVE_DUPLICATES SKIP_ALL_EXCEPT_TESTS) +endif() +message(STATUS "SKIP_TESTS: ${SKIP_TESTS}") +message(STATUS "SKIP_ALL_EXCEPT_TESTS: ${SKIP_ALL_EXCEPT_TESTS}") -list(REMOVE_DUPLICATES SKIP_TESTS) -list(REMOVE_DUPLICATES SKIP_ALL_EXCEPT_TESTS) function(add_test_command NAME EXE) # Restrict the use of SKIP_ALL_EXCEPT_TESTS list in the Int8, BF16 and MIOpenTensile tests - if(((NOT (NAME IN_LIST SKIP_ALL_EXCEPT_TESTS)) AND (MIOPEN_TEST_INT8 OR MIOPEN_TEST_BFLOAT16 OR MIOPEN_TEST_MIOTENSILE)) OR (NAME IN_LIST SKIP_TESTS)) + if( (NOT (NAME IN_LIST SKIP_ALL_EXCEPT_TESTS) AND SKIP_ALL_EXCEPT_TESTS) + OR (NAME IN_LIST SKIP_TESTS) + ) add_test(NAME ${NAME} COMMAND echo skipped) set_tests_properties(${NAME} PROPERTIES DISABLED On) elseif(WIN32) @@ -165,6 +188,8 @@ function(add_test_command NAME EXE) endif() endfunction() +separate_arguments(MIOPEN_TEST_FLAGS_ARGS UNIX_COMMAND ${MIOPEN_TEST_FLAGS}) + function(add_test_executable TEST_NAME) add_executable (${TEST_NAME} EXCLUDE_FROM_ALL ${ARGN}) clang_tidy_check(${TEST_NAME}) @@ -173,13 +198,23 @@ function(add_test_executable TEST_NAME) if(CMAKE_CXX_COMPILER_ID MATCHES "GNU") set_target_properties(${TEST_NAME} PROPERTIES COMPILE_FLAGS -pthread LINK_FLAGS -pthread) endif() - separate_arguments(MIOPEN_TEST_FLAGS_ARGS UNIX_COMMAND ${MIOPEN_TEST_FLAGS}) + if(MIOPEN_TEST_ALL) set(TEST_COMMAND ${TEST_NAME} ${MIOPEN_TEST_FLOAT_ARG} --all ${MIOPEN_TEST_FLAGS_ARGS}) else() set(TEST_COMMAND ${TEST_NAME} ${MIOPEN_TEST_FLOAT_ARG} ${MIOPEN_TEST_FLAGS_ARGS}) endif() + + if(MIOPEN_TEST_LIMIT GREATER 0) + set(TEST_COMMAND ${TEST_COMMAND} --limit ${MIOPEN_TEST_LIMIT}) + endif() + + if(WORKAROUND_ISSUE_936 AND (${TEST_NAME} MATCHES "test_conv2d" OR ${TEST_NAME} MATCHES "test_immed_conv2d") ) + set(TEST_COMMAND ${TEST_COMMAND} --tolerance 130) #increased by 1.625 times + endif() + add_test_command(${TEST_NAME} ${TEST_COMMAND}) + rate_added_test(${TEST_NAME}) add_dependencies(tests ${TEST_NAME}) add_dependencies(check ${TEST_NAME}) set_tests_properties(${TEST_NAME} PROPERTIES FAIL_REGULAR_EXPRESSION "FAILED") @@ -219,29 +254,38 @@ file(GLOB TESTS *.cpp) # LONG_TESTS if it takes more than 800 seconds, or to SHORT_TESTS otherwise. # Please notice that each list is also sorted and it is highly recommended # to keep this sorting when adding new tests. +#For files NAME = test_ + (File name with neither the directory nor the longest extension) -set( LONG_TESTS - pooling2d.cpp - dropout.cpp - conv2d.cpp - pooling3d.cpp - soft_max.cpp - lrn_test.cpp +set(LONG_TESTS + test_dropout + test_conv2d + test_conv3d + test_conv_group + test_soft_max + test_lrn_test + test_conv_for_implicit_gemm + test_immed_conv3d + test_conv3d_extra + test_conv_3d + test_pooling2d ) +function(rate_added_test NAME) + if(${NAME} IN_LIST LONG_TESTS) + set_tests_properties(${NAME} PROPERTIES COST 800) + else() + set_tests_properties(${NAME} PROPERTIES COST 600) + endif() +endfunction() + + + foreach(TEST ${TESTS}) get_filename_component(BASE_NAME ${TEST} NAME_WE) add_test_executable(test_${BASE_NAME} ${TEST}) endforeach() -set_tests_properties( - test_pooling2d - test_dropout - test_conv2d - test_pooling3d - test_soft_max - test_lrn_test - PROPERTIES COST 800) + set_tests_properties(test_sqlite_perfdb test_perfdb PROPERTIES RUN_SERIAL On) @@ -425,12 +469,29 @@ function(add_custom_test NAME) else() set_tests_properties(${NAME} PROPERTIES FAIL_REGULAR_EXPRESSION "FAILED") endif() - set_tests_properties(${NAME} PROPERTIES COST 600) + rate_added_test(${NAME}) else() set_tests_properties(${NAME} PROPERTIES DISABLED On) endif() endfunction() +if(${CODECOV_TEST}) + add_custom_test(test_conv3d_codecov + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 2 4 4 4 4 --weights 2 4 1 1 1 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + ) + add_custom_test(test_immed_conv2d_codecov + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 2 2 14 14 --weights 8 2 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + ) + add_custom_test(test_immed_conv3d_codecov + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 1 4 4 4 4 --weights 2 4 3 3 3 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + ) + add_custom_test(test_pooling2d_codecov + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 1, 192, 28, 28 --lens 2 2 --strides 2 2 --pads 0 0 ${MIOPEN_TEST_FLAGS_ARGS} + ) + +endif() + + set(IMPLICITGEMM_ARGS ${MIOPEN_TEST_FLOAT_ARG}) # ./bin/MIOpenDriver conv -n 128 -c 1024 -H 14 -W 14 -k 2048 -y 1 -x 1 -p 0 -q 0 -u 2 -v 2 -l 1 -j 1 -m conv -g 1 -F 1 -t 1 @@ -706,6 +767,9 @@ COMMAND $ --verbose --input 8 3 108 108 --weights 63 1 COMMAND $ --verbose --input 8 3 108 108 --weights 63 1 3 3 --pads_strides_dilations 2 2 2 2 1 1 --group-count 3 ) + + + if(MIOPEN_TEST_DEEPBENCH) add_custom_test(test_deepbench_rnn MIOTENSILE_ENABLED COMMAND $ --verbose --batch-size 16 --seq-len 50 --vector-len 1760 --hidden-size 1760 --num-layers 1 --in-mode 1 --bias-mode 0 -dir-mode 0 --rnn-mode 0 --flat-batch-fill @@ -861,8 +925,56 @@ COMMAND $ --verbose --input 32 64 27 27 --weights 192 6 COMMAND $ --verbose --input 4 96 14 14 --weights 32 96 5 5 --pads_strides_dilations 2 2 1 1 1 1 COMMAND $ --verbose --input 4 16 14 14 --weights 4 16 5 5 --pads_strides_dilations 2 2 1 1 1 1 COMMAND $ --verbose --input 4 32 14 14 --weights 4 32 5 5 --pads_strides_dilations 2 2 1 1 1 1 + +COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 16 3 64 128 --weights 96 3 11 11 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 16 3 32 32 --weights 96 3 11 11 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 16 3 64 128 --weights 96 3 11 11 --pads_strides_dilations 5 5 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 16 3 32 32 --weights 96 3 11 11 --pads_strides_dilations 5 5 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + +COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 2 16 1024 2048 --weights 32 16 3 3 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 2 16 1024 2048 --weights 32 16 3 3 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 2 16 3072 3072 --weights 32 16 3 3 --pads_strides_dilations 0 0 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 2 16 3072 3072 --weights 32 16 3 3 --pads_strides_dilations 2 2 2 2 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + +COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 128 320 1 7 --weights 256 320 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 128 1024 1 7 --weights 2048 1024 1 1 --pads_strides_dilations 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 352 192 7 1 --weights 320 192 1 1 --pads_strides_dilations 0 0 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} +COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 352 16 7 1 --weights 32 16 1 1 --pads_strides_dilations 2 2 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} ) +if (0) #disabled too many errors +if(MIOPEN_TEST_LIMIT GREATER 0) +if(${MIOPEN_USE_MIOPENGEMM} AND (${MIOPEN_hip_VERSION_MAJOR} EQUAL 3) AND (${MIOPEN_hip_VERSION_MINOR} EQUAL 7)) + add_custom_test(test_conv3d_extra SKIP_UNLESS_ALL + COMMAND MIOPEN_LOG_LEVEL=6 $ ${MIOPEN_TEST_FLOAT_ARG} --input 2 16 50 50 50 --weights 32 16 5 5 5 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_LOG_LEVEL=6 $ ${MIOPEN_TEST_FLOAT_ARG} --input 2 16 50 50 50 --weights 32 16 5 5 5 --pads_strides_dilations 0 0 0 2 2 2 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_LOG_LEVEL=6 $ ${MIOPEN_TEST_FLOAT_ARG} --input 2 16 50 50 50 --weights 32 16 5 5 5 --pads_strides_dilations 2 2 2 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND MIOPEN_LOG_LEVEL=6 $ ${MIOPEN_TEST_FLOAT_ARG} --input 2 16 50 50 50 --weights 32 16 5 5 5 --pads_strides_dilations 0 0 0 1 1 1 2 2 2 ${MIOPEN_TEST_FLAGS_ARGS} + #COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 1 16 4 161 700 --weights 16 16 3 11 11 --pads_strides_dilations 1 1 1 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + + #ROCM3.7 compiler problems + #COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 1 16 4 161 700 --weights 16 16 3 11 11 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + #COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 1 16 4 161 700 --weights 16 16 3 11 11 --pads_strides_dilations 0 0 0 2 2 2 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + #COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 1 16 4 140 602 --weights 16 16 3 11 11 --pads_strides_dilations 1 1 1 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + #COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 1 16 4 140 602 --weights 16 16 3 11 11 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + ) + message(STATUS "test_conv3d_extra reduced set") +else() + add_custom_test(test_conv3d_extra SKIP_UNLESS_ALL + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 2 16 50 50 50 --weights 32 16 5 5 5 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 2 16 50 50 50 --weights 32 16 5 5 5 --pads_strides_dilations 0 0 0 2 2 2 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 2 16 50 50 50 --weights 32 16 5 5 5 --pads_strides_dilations 2 2 2 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 2 16 50 50 50 --weights 32 16 5 5 5 --pads_strides_dilations 0 0 0 1 1 1 2 2 2 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 1 16 4 161 700 --weights 16 16 3 11 11 --pads_strides_dilations 1 1 1 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 1 16 4 161 700 --weights 16 16 3 11 11 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 1 16 4 161 700 --weights 16 16 3 11 11 --pads_strides_dilations 0 0 0 2 2 2 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 1 16 4 140 602 --weights 16 16 3 11 11 --pads_strides_dilations 1 1 1 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + COMMAND $ ${MIOPEN_TEST_FLOAT_ARG} --input 1 16 4 140 602 --weights 16 16 3 11 11 --pads_strides_dilations 0 0 0 1 1 1 1 1 1 ${MIOPEN_TEST_FLAGS_ARGS} + ) +endif() +endif() +endif() + add_custom_test(test_conv_trans SKIP_UNLESS_ALL MIOTENSILE_ENABLED COMMAND $ --verbose --input 8 128 28 28 --weights 128 128 1 1 --pads_strides_dilations 0 0 1 1 1 1 --cmode trans --pmode default @@ -937,16 +1049,25 @@ set(DYNAMIC_IMPLICITGEMM_XDLOPS_NHWC_BWD_ENVS ${DYNAMIC_IMPLICITGEMM_COMMON} MIOPEN_DEBUG_FIND_ONLY_SOLVER=ConvAsmImplicitGemmGTCDynamicBwdXdlopsNHWC) -add_custom_test(test_conv_igemm_dynamic_small -COMMAND ${DYNAMIC_IMPLICITGEMM_ENVS} $ --verbose --input 16 16 56 56 --weights 64 16 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_ENVS} $ --verbose --input 16 64 34 34 --weights 64 64 3 3 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_ENVS} $ --verbose --input 32 32 17 17 --weights 32 32 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-backward-data --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_1X1_ENVS} $ --verbose --input 16 384 8 8 --weights 64 384 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights -COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS} $ --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} $ --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_BWD_ENVS} $ --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} $ --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 -) +if(${CODECOV_TEST}) + add_custom_test(test_conv_igemm_dynamic_small + COMMAND ${DYNAMIC_IMPLICITGEMM_ENVS} $ --verbose --input 32 32 17 17 --weights 32 32 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-backward-data --disable-backward-weights --disable-validation + COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS} $ --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 --disable-validation + COMMAND ${DYNAMIC_IMPLICITGEMM_BWD_ENVS} $ --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 --disable-validation + ) + set_tests_properties(test_conv_igemm_dynamic_small PROPERTIES COST 800) +else() + add_custom_test(test_conv_igemm_dynamic_small + COMMAND ${DYNAMIC_IMPLICITGEMM_ENVS} $ --verbose --input 16 16 56 56 --weights 64 16 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights + COMMAND ${DYNAMIC_IMPLICITGEMM_ENVS} $ --verbose --input 16 64 34 34 --weights 64 64 3 3 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights + COMMAND ${DYNAMIC_IMPLICITGEMM_ENVS} $ --verbose --input 32 32 17 17 --weights 32 32 1 7 --pads_strides_dilations 0 3 1 1 1 1 --disable-backward-data --disable-backward-weights + COMMAND ${DYNAMIC_IMPLICITGEMM_1X1_ENVS} $ --verbose --input 16 384 8 8 --weights 64 384 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights + COMMAND ${DYNAMIC_IMPLICITGEMM_WRW_ENVS} $ --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} $ --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_BWD_ENVS} $ --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} $ --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 + ) +endif() #if CODECOV_TEST add_custom_test(test_conv_igemm_dynamic SKIP_UNLESS_ALL COMMAND ${DYNAMIC_IMPLICITGEMM_ENVS} $ --verbose --input 64 64 56 56 --weights 256 64 1 1 --pads_strides_dilations 0 0 1 1 1 1 --disable-backward-data --disable-backward-weights @@ -1109,7 +1230,7 @@ if(MIOPEN_TEST_DEEPBENCH) 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) @@ -1154,7 +1275,7 @@ if(MIOPEN_TEST_CONV) 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) diff --git a/test/conv2d.cpp b/test/conv2d.cpp index f302802ba2..056ba503b2 100644 --- a/test/conv2d.cpp +++ b/test/conv2d.cpp @@ -30,12 +30,26 @@ struct conv2d_driver : conv_driver { conv2d_driver() : conv_driver() { - this->add(this->input, "input", this->get_tensor(get_inputs, tensor_elem_gen_integer())); - this->add( - this->weights, "weights", this->get_tensor(get_weights, tensor_elem_gen_integer())); + this->add(this->input_dims, "input"); + this->add(this->weight_tensor_dims, "weights"); + this->add(this->batch_size, + "batch_size", + this->generate_data_limited(this->get_batch_sizes(), 1)); + this->add(this->input_channels, + "input_channels", + this->generate_data_limited(this->get_input_channels(), 1, {32})); + this->add(this->output_channels, + "output_channels", + this->generate_data_limited(this->get_output_channels(), 1, {64})); + this->add(this->spatial_dim_elements, + "spatial_dim_elements", + this->generate_data_limited(this->get_2d_spatial_dims(), 1, {28, 28})); + this->add(this->filter_dims, + "filter_dims", + this->generate_data_limited(this->get_2d_filter_dims(), 2, {3, 3})); this->add(this->pads_strides_dilations, "pads_strides_dilations", - this->generate_data(this->get_2d_pads_strides_dilations())); + this->generate_data_limited(this->get_2d_pads_strides_dilations(), 2)); this->add(this->trans_output_pads, "trans_output_pads", this->generate_data(this->get_2d_trans_output_pads())); diff --git a/test/conv3d.cpp b/test/conv3d.cpp index 40b9fa9b29..aef7020a6a 100644 --- a/test/conv3d.cpp +++ b/test/conv3d.cpp @@ -30,18 +30,29 @@ struct conv3d_driver : conv_driver { conv3d_driver() : conv_driver() { - this->add(this->input, - "input", - this->get_tensor(get_3d_conv_input_shapes, tensor_elem_gen_integer())); - this->add(this->weights, - "weights", - this->get_tensor(get_3d_conv_weight_shapes, tensor_elem_gen_integer())); + this->add(this->input_dims, "input"); + this->add(this->weight_tensor_dims, "weights"); + this->add(this->batch_size, + "batch_size", + this->generate_data_limited(this->get_batch_sizes(), 1, {8})); + this->add(this->input_channels, + "input_channels", + this->generate_data_limited(this->get_input_channels(), 1, {32})); + this->add(this->output_channels, + "output_channels", + this->generate_data_limited(this->get_output_channels(), 1, {32})); + this->add(this->spatial_dim_elements, + "spatial_dim_elements", + this->generate_data_limited(this->get_3d_spatial_dims(), 1, {16, 16, 16})); + this->add(this->filter_dims, + "filter_dims", + this->generate_data_limited(this->get_3d_filter_dims(), 2, {3, 3, 3})); this->add(this->pads_strides_dilations, "pads_strides_dilations", - this->generate_data(this->get_3d_pads_strides_dilations())); + this->generate_data_limited(this->get_3d_pads_strides_dilations(), 2)); this->add(this->trans_output_pads, "trans_output_pads", - this->generate_data(this->get_3d_trans_output_pads())); + this->generate_data_limited(this->get_3d_trans_output_pads(), 1)); this->add(this->in_layout, "in_layout", this->generate_data({"NCDHW"})); this->add(this->fil_layout, "fil_layout", this->generate_data({"NCDHW"})); this->add(this->out_layout, "out_layout", this->generate_data({"NCDHW"})); diff --git a/test/conv_common.hpp b/test/conv_common.hpp index 6f692e375d..7c9cb3ec23 100644 --- a/test/conv_common.hpp +++ b/test/conv_common.hpp @@ -1571,6 +1571,13 @@ struct conv_driver : test_driver miopen::ConvolutionDescriptor filter; std::string conv_mode; std::string pad_mode; + std::vector spatial_dim_elements{}; + std::vector input_dims{}; + std::vector weight_tensor_dims{}; + std::vector filter_dims{}; + std::size_t batch_size{}; + std::size_t input_channels{}; + std::size_t output_channels{}; std::string in_layout; std::string fil_layout; // keep same as MIOpenDriver argument name std::string out_layout; @@ -1596,6 +1603,41 @@ struct conv_driver : test_driver {"VALID", miopenPaddingValid}, {"DEFAULT", miopenPaddingDefault}}; + std::vector get_batch_sizes() { return {1, 8, 2, 64, 30, 128, 352, 512}; } + + std::vector> get_2d_spatial_dims() + { + return {{14, 14}, + {28, 28}, + {32, 32}, + {7, 7}, + {17, 17}, + {56, 56}, + {55, 55}, + {64, 128}, + {224, 224}, + {1024, 2048}, + {3072, 3072}, + {1, 1}, + {1, 7}, + {7, 1}}; + } + + std::vector> get_2d_filter_dims() + { + return {{1, 1}, {3, 3}, {1, 7}, {5, 5}, {7, 1}, {7, 7}, {11, 11}, {2, 2}, {4, 4}}; + } + + std::vector get_output_channels() + { + return {32, 64, 16, 128, 96, 112, 192, 256, 320, 512, 1024}; + } + + std::vector get_input_channels() + { + return {16, 32, 3, 128, 96, 112, 192, 256, 320, 512, 1024}; + } + std::vector> get_2d_pads_strides_dilations() { return {{0, 0, 1, 1, 1, 1}, @@ -1611,6 +1653,32 @@ struct conv_driver : test_driver {1, 1, 2, 2, 2, 1}}; } + std::vector> get_3d_spatial_dims() + { + return {{3, 4, 4}, + {4, 9, 9}, + {3, 14, 14}, + {4, 28, 28}, + {4, 56, 56}, + {4, 161, 700}, + {4, 227, 227}, + {1, 1, 1}, + {1, 2, 2}}; + } + + std::vector> get_3d_filter_dims() + { + return {{1, 1, 1}, + {3, 3, 3}, + {3, 5, 5}, + {3, 7, 7}, + {5, 7, 7}, + {3, 11, 11}, + {3, 1, 7}, + {3, 7, 1}, + {3, 5, 20}}; + } + std::vector> get_2d_trans_output_pads() { return {{0, 0}}; } std::vector> get_3d_pads_strides_dilations() @@ -1636,11 +1704,12 @@ struct conv_driver : test_driver { for(int i = 2; i < 4; i++) { - if(input.desc.GetSize() == i + 2 and weights.desc.GetSize() == i + 2 and + if(input_dims.size() == i + 2 and weight_tensor_dims.size() == i + 2 and pads_strides_dilations.size() == i * 3 and trans_output_pads.size() == i) return i; } - return -1; + std::cout << "FAILED: get_spatial_dim() can't calculate dims count." << std::endl; + exit(-1); // NOLINT (concurrency-mt-unsafe) } conv_driver() @@ -1661,6 +1730,69 @@ struct conv_driver : test_driver void run() { + + if(!input_dims.empty()) + filter.spatialDim = get_spatial_dim(); + else + filter.spatialDim = filter_dims.size(); + + filter.mode = cmode_lookup[miopen::ToUpper(conv_mode)]; + filter.paddingMode = pmode_lookup[miopen::ToUpper(pad_mode)]; + std::size_t spatial_dim = filter.GetSpatialDimension(); + filter.group_count = std::max(static_cast(groupCount), 1); + + if(!input_dims.empty()) + { + input = tensor{input_dims}.generate(tensor_elem_gen_integer{17}); + batch_size = input_dims.at(0); + input_channels = input_dims.at(1); + std::copy(input_dims.begin() + 2, input_dims.end(), spatial_dim_elements.begin()); + } + else if(spatial_dim == 2) + { + input = tensor{ + batch_size, + input_channels, + spatial_dim_elements.at(0), + spatial_dim_elements.at( + 1)}.generate(tensor_elem_gen_integer{17}); + } + else if(spatial_dim == 3) + { + input = tensor{ + batch_size, + input_channels, + spatial_dim_elements.at(0), + spatial_dim_elements.at(1), + spatial_dim_elements.at( + 2)}.generate(tensor_elem_gen_integer{17}); + } + + if(!weight_tensor_dims.empty()) + { + weights = tensor{weight_tensor_dims}.generate(tensor_elem_gen_integer{17}); + output_channels = weight_tensor_dims.at(0); + } + else if(spatial_dim == 2) + { + weights = tensor{ + output_channels, + input_channels / filter.group_count, + filter_dims.at(0), + filter_dims.at( + 1)}.generate(tensor_elem_gen_integer{17}); + } + else if(spatial_dim == 3) + { + weights = tensor{ + output_channels, + input_channels / filter.group_count, + filter_dims.at(0), + filter_dims.at(1), + filter_dims.at( + 2)}.generate(tensor_elem_gen_integer{17}); + } + if(input.desc.GetSize() != in_layout.size() || weights.desc.GetSize() != fil_layout.size() || input.desc.GetSize() != out_layout.size()) { @@ -1695,11 +1827,6 @@ struct conv_driver : test_driver weights.desc = miopen::TensorDescriptor(miopen_type{}, dim_lens, dim_strides); } - filter.spatialDim = get_spatial_dim(); - filter.mode = cmode_lookup[miopen::ToUpper(conv_mode)]; - filter.paddingMode = pmode_lookup[miopen::ToUpper(pad_mode)]; - std::size_t spatial_dim = filter.GetSpatialDimension(); - if(input.desc.GetSize() != 2 + spatial_dim || weights.desc.GetSize() != 2 + spatial_dim || pads_strides_dilations.size() != 3 * spatial_dim || trans_output_pads.size() != spatial_dim) @@ -1721,8 +1848,6 @@ struct conv_driver : test_driver filter.dilations.begin()); std::copy_n(trans_output_pads.begin(), spatial_dim, filter.trans_output_pads.begin()); - filter.group_count = std::max(static_cast(groupCount), 1); - std::size_t in_c_len = input.desc.GetLengths()[1]; std::size_t wei_k_len = weights.desc.GetLengths()[0]; std::size_t wei_c_len = weights.desc.GetLengths()[1]; @@ -2113,6 +2238,15 @@ struct conv_bias_driver : test_driver tensor bias(bias_lens); + if(!(bias.desc.GetLengths()[0] == 1 && + bias.desc.GetLengths()[1] == output.desc.GetLengths()[0] && + std::all_of(bias.desc.GetLengths().begin() + 2, + bias.desc.GetLengths().end(), + [](auto v) { return v == 1; }))) + { + return; + } + size_t total_mem = bias.desc.GetNumBytes() + output.desc.GetNumBytes(); // estimate based on backward pass size_t device_mem = get_handle().GetGlobalMemorySize(); diff --git a/test/driver.hpp b/test/driver.hpp index 11e7a020f1..f0c7791dc2 100644 --- a/test/driver.hpp +++ b/test/driver.hpp @@ -175,6 +175,7 @@ struct test_driver std::string cache_path = compute_cache_path(); miopenDataType_t type = miopenFloat; bool full_set = false; + int limit_set = 0; bool verbose = false; double tolerance = 80; bool time = false; @@ -198,6 +199,7 @@ struct test_driver template void parse(Visitor v) { + v(limit_set, {"--limit"}, "Limits the number of generated test elements."); v(full_set, {"--all"}, "Run all tests"); v(verbose, {"--verbose", "-v"}, "Run verbose mode"); v(tolerance, {"--tolerance", "-t"}, "Set test tolerance"); @@ -477,6 +479,25 @@ struct test_driver }}; } + template + generate_data_t> + generate_data_limited(std::vector dims, int limit_multiplier, T single) + { + return {[=]() -> std::vector { + if(limit_set > 0) + { + auto endpoint = + std::min(static_cast(dims.size()), limit_set * limit_multiplier); + std::vector subvec(dims.cbegin(), dims.cbegin() + endpoint); + return subvec; + } + else if(full_set) + return dims; + else + return {single}; + }}; + } + template generate_data_t> generate_data(std::initializer_list dims) { @@ -501,6 +522,24 @@ struct test_driver }}; } + template + generate_data_t> generate_data_limited(std::vector dims, int limit_multiplier) + { + return {[=]() -> std::vector { + if(limit_set > 0) + { + auto endpoint = + std::min(static_cast(dims.size()), limit_set * limit_multiplier); + std::vector subvec(dims.cbegin(), dims.cbegin() + endpoint); + return subvec; + } + else if(full_set) + return dims; + else + return {dims.front()}; + }}; + } + template auto lazy_generate_data(F f, T single) -> generate_data_t { diff --git a/test/immed_conv2d.cpp b/test/immed_conv2d.cpp index 59112d251c..8ae0f2c841 100644 --- a/test/immed_conv2d.cpp +++ b/test/immed_conv2d.cpp @@ -30,17 +30,29 @@ struct conv2d_driver : conv_driver { conv2d_driver() : conv_driver() { - this->add( - this->input, "input", this->get_tensor(get_immed_inputs, tensor_elem_gen_integer())); - this->add(this->weights, - "weights", - this->get_tensor(get_immed_weights, tensor_elem_gen_integer())); + this->add(this->input_dims, "input"); + this->add(this->weight_tensor_dims, "weights"); + this->add(this->batch_size, + "batch_size", + this->generate_data_limited(this->get_batch_sizes(), 1, {16})); + this->add(this->input_channels, + "input_channels", + this->generate_data_limited(this->get_input_channels(), 1, {32})); + this->add(this->output_channels, + "output_channels", + this->generate_data_limited(this->get_output_channels(), 1, {32})); + this->add(this->spatial_dim_elements, + "spatial_dim_elements", + this->generate_data_limited(this->get_2d_spatial_dims(), 1, {56, 56})); + this->add(this->filter_dims, + "filter_dims", + this->generate_data_limited(this->get_2d_filter_dims(), 2, {3, 3})); this->add(this->pads_strides_dilations, "pads_strides_dilations", - this->generate_data(this->get_2d_pads_strides_dilations())); + this->generate_data_limited(this->get_2d_pads_strides_dilations(), 2)); this->add(this->trans_output_pads, "trans_output_pads", - this->generate_data(this->get_2d_trans_output_pads())); + this->generate_data_limited(this->get_2d_trans_output_pads(), 1)); this->add(this->in_layout, "in_layout", this->generate_data({"NCHW"})); this->add(this->fil_layout, "fil_layout", this->generate_data({"NCHW"})); this->add(this->out_layout, "out_layout", this->generate_data({"NCHW"})); diff --git a/test/immed_conv3d.cpp b/test/immed_conv3d.cpp index 9cc23e2456..4196fa3ad2 100644 --- a/test/immed_conv3d.cpp +++ b/test/immed_conv3d.cpp @@ -30,18 +30,29 @@ struct conv3d_driver : conv_driver { conv3d_driver() : conv_driver() { - this->add(this->input, - "input", - this->get_tensor(get_3d_conv_input_shapes, tensor_elem_gen_integer())); - this->add(this->weights, - "weights", - this->get_tensor(get_3d_conv_weight_shapes, tensor_elem_gen_integer())); + this->add(this->input_dims, "input"); + this->add(this->weight_tensor_dims, "weights"); + this->add(this->batch_size, + "batch_size", + this->generate_data_limited(this->get_batch_sizes(), 1, {8})); + this->add(this->input_channels, + "input_channels", + this->generate_data_limited(this->get_input_channels(), 1, {2})); + this->add(this->output_channels, + "output_channels", + this->generate_data_limited(this->get_output_channels(), 1, {16})); + this->add(this->spatial_dim_elements, + "spatial_dim_elements", + this->generate_data_limited(this->get_3d_spatial_dims(), 1, {16, 16, 16})); + this->add(this->filter_dims, + "filter_dims", + this->generate_data_limited(this->get_3d_filter_dims(), 2, {5, 5, 5})); this->add(this->pads_strides_dilations, "pads_strides_dilations", - this->generate_data(this->get_3d_pads_strides_dilations())); + this->generate_data_limited(this->get_3d_pads_strides_dilations(), 2)); this->add(this->trans_output_pads, "trans_output_pads", - this->generate_data(this->get_3d_trans_output_pads())); + this->generate_data_limited(this->get_3d_trans_output_pads(), 1)); this->add(this->in_layout, "in_layout", this->generate_data({"NCDHW"})); this->add(this->fil_layout, "fil_layout", this->generate_data({"NCDHW"})); this->add(this->out_layout, "out_layout", this->generate_data({"NCDHW"})); diff --git a/test/network_data.hpp b/test/network_data.hpp index 4dd6b56fd6..340de2ed9c 100644 --- a/test/network_data.hpp +++ b/test/network_data.hpp @@ -68,7 +68,6 @@ inline std::set> get_inputs(int n = MIOPEN_TEST_DEFAULT_BATCH_S { pick_batch_size(32, n), 192, 7, 7 }, { pick_batch_size(32, n), 256, 28, 28 }, { pick_batch_size(32, n), 3, 224, 224 }, - { pick_batch_size(32, n), 32, 14, 14 }, { pick_batch_size(32, n), 32, 28, 28 }, { pick_batch_size(32, n), 48, 7, 7 }, { pick_batch_size(32, n), 480, 128, 256 }, @@ -104,7 +103,6 @@ inline std::set> get_weights(int n = MIOPEN_TEST_DEFAULT_BATCH_ { { pick_batch_size(1024, n),1024, 3, 3 }, { pick_batch_size(1024, n),512, 3, 3 }, - { pick_batch_size(112, n), 512, 1, 1 }, { pick_batch_size(128, n), 256, 1, 1 }, { pick_batch_size(128, n), 528, 1, 1 }, { pick_batch_size(128, n), 96, 3, 3 }, @@ -145,8 +143,6 @@ inline std::set> get_immed_inputs(int n = MIOPEN_TEST_DEFAULT_B { pick_batch_size(256, n), 256, 13, 13 }, { pick_batch_size(256, n), 3, 227, 227 }, { pick_batch_size(32, n), 64, 56, 56 }, -// { pick_batch_size(32, n), 832, 64, 128 }, -// { pick_batch_size(32, n), 832, 7, 7 }, { pick_batch_size(32, n), 96, 14, 14 }, { pick_batch_size(32, n), 96, 28, 28 }, { pick_batch_size(64, n), 128, 56, 56 }, @@ -166,7 +162,6 @@ inline std::set> get_immed_weights(int n = MIOPEN_TEST_DEFAULT_ { pick_batch_size(24, n), 512, 1, 1 }, { pick_batch_size(256, n), 128, 3, 3 }, { pick_batch_size(256, n), 256, 3, 3 }, -// { pick_batch_size(256, n), 832, 1, 1 }, { pick_batch_size(256, n), 64, 5, 5 }, { pick_batch_size(288, n), 144, 3, 3 }, { pick_batch_size(96, n), 3, 11, 11 }, @@ -233,7 +228,6 @@ get_bn_peract_inputs(int n = MIOPEN_TEST_DEFAULT_BATCH_SIZE_FACTOR) { pick_batch_size(256, n), 3, 227, 227 }, { pick_batch_size(64, n), 64, 112, 112 },//Batch-norm ResNet 152 after this line { pick_batch_size(256, n), 1024, 14, 14 },// n is from the paper @ 256 - { pick_batch_size(256, n), 128, 28, 28 }, { pick_batch_size(256, n), 2048, 7, 7 }, { pick_batch_size(256, n), 256, 56, 56 }, { pick_batch_size(256, n), 256, 14, 14 }, @@ -246,7 +240,6 @@ get_bn_peract_inputs(int n = MIOPEN_TEST_DEFAULT_BATCH_SIZE_FACTOR) { pick_batch_size(32, n), 128, 28, 28 }, { pick_batch_size(32, n), 128, 4, 4 }, { pick_batch_size(32, n), 128, 7, 7 }, - { pick_batch_size(32, n), 160, 14, 14 }, { pick_batch_size(32, n), 160, 7, 7 }, { pick_batch_size(32, n), 192, 14, 14 }, { pick_batch_size(32, n), 192, 56, 56 }, @@ -254,11 +247,9 @@ get_bn_peract_inputs(int n = MIOPEN_TEST_DEFAULT_BATCH_SIZE_FACTOR) { pick_batch_size(32, n), 224, 14, 14 }, { pick_batch_size(32, n), 256, 7, 7 }, { pick_batch_size(32, n), 256, 14, 14 }, - { pick_batch_size(32, n), 32, 28, 28 }, { pick_batch_size(32, n), 352, 7, 7 }, { pick_batch_size(32, n), 64, 112, 112 }, { pick_batch_size(32, n), 64, 14, 14 }, - { pick_batch_size(32, n), 64, 28, 28 }, { pick_batch_size(32, n), 64, 56, 56 }, { pick_batch_size(32, n), 96, 28, 28 }, { pick_batch_size(32, n), 32, 256, 512 }, //Killing this config. Takes way too long on the CPU @@ -282,11 +273,8 @@ get_bn_spatial_inputs(int n = MIOPEN_TEST_DEFAULT_BATCH_SIZE_FACTOR) { pick_batch_size(256, n), 3, 227, 227 }, { pick_batch_size(256, n), 64, 112, 112 }, { pick_batch_size(512, n), 16, 32, 32 }, - { pick_batch_size(32, n), 64, 112, 112 }, - //{ pick_batch_size(100, n), 3, 32, 32 },// causing issues with Jenkins { pick_batch_size(100, n), 32, 8, 8 }, { pick_batch_size(128, n), 256, 12, 12 }, - { pick_batch_size(256, n), 1024, 14, 14 },// n is from the paper @ 256 { pick_batch_size(256, n), 128, 28, 28 }, { pick_batch_size(256, n), 2048, 7, 7 }, { pick_batch_size(256, n), 256, 56, 56 }, @@ -296,18 +284,13 @@ get_bn_spatial_inputs(int n = MIOPEN_TEST_DEFAULT_BATCH_SIZE_FACTOR) { pick_batch_size(256, n), 64, 56, 56 },//Batch-norm Inception_v3 after this { pick_batch_size(32, n), 1024, 1, 1 },// n is from the paper @ 32 { pick_batch_size(32, n), 128, 14, 14 }, - { pick_batch_size(32, n), 128, 28, 28 }, { pick_batch_size(32, n), 128, 4, 4 }, - { pick_batch_size(32, n), 128, 7, 7 }, - { pick_batch_size(32, n), 160, 14, 14 }, { pick_batch_size(32, n), 160, 7, 7 }, { pick_batch_size(32, n), 192, 14, 14 }, { pick_batch_size(32, n), 192, 56, 56 }, { pick_batch_size(32, n), 192, 7, 7 }, { pick_batch_size(32, n), 224, 14, 14 }, { pick_batch_size(32, n), 256, 7, 7 }, - { pick_batch_size(32, n), 256, 14, 14 }, - { pick_batch_size(32, n), 32, 28, 28 }, { pick_batch_size(32, n), 352, 7, 7 }, { pick_batch_size(32, n), 64, 14, 14 }, { pick_batch_size(32, n), 64, 28, 28 }, diff --git a/test/pooling2d.cpp b/test/pooling2d.cpp index 4e5e226e0e..4dca1cf5bc 100644 --- a/test/pooling2d.cpp +++ b/test/pooling2d.cpp @@ -26,19 +26,31 @@ #include "pooling_common.hpp" -#define TEST_GET_INPUT_TENSOR 1 +#define TEST_GET_INPUT_TENSOR 0 template struct pooling2d_driver : pooling_driver { std::vector> get_2d_pooling_input_shapes() { - return {{1, 19, 1024, 2048}, {100, 3, 32, 32}, {1, 32, 16, 16}, {5, 32, 8, 8}, - {2, 1024, 12, 12}, {4, 3, 231, 231}, {8, 3, 227, 227}, {1, 384, 13, 13}, - {1, 96, 27, 27}, {2, 112, 14, 14}, {2, 160, 7, 7}, {1, 192, 256, 512}, - {2, 192, 28, 28}, {1, 832, 64, 128}, {1, 256, 56, 56}, {4, 3, 224, 224}, - {2, 64, 112, 112}, {2, 608, 4, 4}, {1, 2048, 11, 11}, {1, 16, 2048, 2048}, - {1, 16, 3072, 3072}, {1, 16, 4096, 4096}}; + return {{1, 19, 1024, 2048}, + {10, 3, 32, 32}, + {5, 32, 8, 8}, + {2, 1024, 12, 12}, + {4, 3, 231, 231}, + {8, 3, 227, 227}, + {1, 384, 13, 13}, + {1, 96, 27, 27}, + {2, 160, 7, 7}, + {1, 192, 256, 512}, + {2, 192, 28, 28}, + {1, 832, 64, 128}, + {1, 256, 56, 56}, + {4, 3, 224, 224}, + {2, 64, 112, 112}, + {2, 608, 4, 4}, + {1, 2048, 11, 11}, + {1, 16, 4096, 4096}}; } pooling2d_driver() : pooling_driver() @@ -48,7 +60,8 @@ struct pooling2d_driver : pooling_driver std::vector> in_dim_vec(in_dim_set.begin(), in_dim_set.end()); this->add(this->in_shape, "input", this->generate_data(in_dim_vec, {16, 32, 8, 8})); #else - this->add(this->in_shape, "input", this->generate_data(get_2d_pooling_input_shapes())); + this->add( + this->in_shape, "input", this->generate_data_limited(get_2d_pooling_input_shapes(), 9)); #endif this->add(this->lens, "lens", this->generate_data({{2, 2}, {3, 3}})); this->add(this->strides, "strides", this->generate_data({{2, 2}, {1, 1}})); diff --git a/test/pooling3d.cpp b/test/pooling3d.cpp index 0870487f8c..90b37d5c75 100644 --- a/test/pooling3d.cpp +++ b/test/pooling3d.cpp @@ -42,7 +42,8 @@ struct pooling3d_driver : pooling_driver pooling3d_driver() : pooling_driver() { - this->add(this->in_shape, "input", this->generate_data(get_3d_pooling_input_shapes())); + this->add( + this->in_shape, "input", this->generate_data_limited(get_3d_pooling_input_shapes(), 4)); this->add(this->lens, "lens", this->generate_data({{2, 2, 2}, {3, 3, 3}})); this->add(this->strides, "strides", this->generate_data({{2, 2, 2}, {1, 1, 1}})); this->add(this->pads, "pads", this->generate_data({{0, 0, 0}, {1, 1, 1}}));