From d1627ecf92f59c5ee09e7e0f0acb6db3102a81ae Mon Sep 17 00:00:00 2001 From: Igor Mirosavljevic <104914075+IgorMirosavljevicHTEC@users.noreply.github.com> Date: Thu, 23 May 2024 15:45:44 +0200 Subject: [PATCH 01/13] Ci improvements (#3114) --- .github/workflows/config.md | 29 ++++++++ .github/workflows/performance.yaml | 57 +++++++++++--- .github/workflows/sync_rocMLIR.yaml | 90 +++++++++++++++++++++++ .github/workflows/weekly_master_sync.yaml | 25 +++++++ 4 files changed, 190 insertions(+), 11 deletions(-) create mode 100644 .github/workflows/config.md create mode 100644 .github/workflows/sync_rocMLIR.yaml create mode 100644 .github/workflows/weekly_master_sync.yaml diff --git a/.github/workflows/config.md b/.github/workflows/config.md new file mode 100644 index 00000000000..326c99ca483 --- /dev/null +++ b/.github/workflows/config.md @@ -0,0 +1,29 @@ +#=====ROCM INFO===== +ROCM_VERSION : '6.0.2' +#default ROCm version to be used +ROCM_BASE_IMAGE : 'rocm/dev-ubuntu-20.04' +#base image from dockerhub to be used +ROCM_BUILT_IMAGE : 'rocm-migraphx' +#name of the docker image built upon ROCm base +USE_NAVI : '0' +#disable NAVI in image build +OVERWRITE_EXISTING : 'true' +#building new ROCm image overwrites old with same version + +#=====REPOS INFO===== +ORGANIZATION_REPO : 'AMD' +BENCHMARK_UTILS_REPO : 'ROCm/migraphx-benchmark-utils' +PERFORMANCE_REPORTS_REPO : 'ROCm/migraphx-reports' +PERFORMANCE_BACKUP_REPO : 'migraphx-benchmark/performance-backup' + +#=====PERFORMANCE SCRIPT PARAMETERS===== +RESULTS_TO_COMPARE : '10' +#number of previous performance results to be used in calculations +CALCULATION_METHOD_FLAG : '-r' +#calculation method used in reporting, -m for Max value; -s for Std dev; -r for Threshold file +PERFORMANCE_TEST_TIMEOUT : '30m' +#timeout for each model after which test is aborted + +#===== W A R N I N G ===== +#VARIABLE NAMES NOT TO BE CHANGED, VALUES ONLY! +#VALUES MUST BE ENGLOSED IN SINGLE QUOTES! \ No newline at end of file diff --git a/.github/workflows/performance.yaml b/.github/workflows/performance.yaml index 42e92e6dc2b..98e79ac3efe 100644 --- a/.github/workflows/performance.yaml +++ b/.github/workflows/performance.yaml @@ -5,7 +5,7 @@ on: branches: [develop] types: [opened, synchronize, closed] schedule: - - cron: "0 6 * * 1-6" + - cron: "0 7 * * 1-6" workflow_dispatch: inputs: @@ -47,18 +47,53 @@ concurrency: cancel-in-progress: true jobs: - release: + get_config: + runs-on: ubuntu-latest + outputs: + rocm_version: ${{ steps.read_config.outputs.rocm_version }} + utils_repo: ${{ steps.read_config.outputs.utils_repo }} + reports_repo: ${{ steps.read_config.outputs.reports_repo }} + backup_repo: ${{ steps.read_config.outputs.backup_repo }} + repo_org: ${{ steps.read_config.outputs.repo_org }} + perf_number: ${{ steps.read_config.outputs.perf_number }} + perf_flag: ${{ steps.read_config.outputs.perf_flag }} + perf_timeout: ${{ steps.read_config.outputs.perf_timeout }} + steps: + - name: checkout + uses: actions/checkout@v4.1.1 + - name: read_config + id: read_config + run: | + ROCM_VERSION=$(grep 'ROCM_VERSION' .github/workflows/config.md | cut -d "'" -f2) + BENCHMARK_UTILS_REPO=$(grep 'BENCHMARK_UTILS_REPO' .github/workflows/config.md | cut -d "'" -f2) + PERFORMANCE_REPORTS_REPO=$(grep 'PERFORMANCE_REPORTS_REPO' .github/workflows/config.md | cut -d "'" -f2) + PERFORMANCE_BACKUP_REPO=$(grep 'PERFORMANCE_BACKUP_REPO' .github/workflows/config.md | cut -d "'" -f2) + ORGANIZATION_REPO=$(grep 'ORGANIZATION_REPO' .github/workflows/config.md | cut -d "'" -f2) + RESULTS_TO_COMPARE=$(grep 'RESULTS_TO_COMPARE' .github/workflows/config.md | cut -d "'" -f2) + CALCULATION_METHOD_FLAG=$(grep 'CALCULATION_METHOD_FLAG' .github/workflows/config.md | cut -d "'" -f2) + PERFORMANCE_TEST_TIMEOUT=$(grep 'PERFORMANCE_TEST_TIMEOUT' .github/workflows/config.md | cut -d "'" -f2) + echo "rocm_version=$ROCM_VERSION" >> $GITHUB_OUTPUT + echo "utils_repo=$BENCHMARK_UTILS_REPO" >> $GITHUB_OUTPUT + echo "reports_repo=$PERFORMANCE_REPORTS_REPO" >> $GITHUB_OUTPUT + echo "backup_repo=$PERFORMANCE_BACKUP_REPO" >> $GITHUB_OUTPUT + echo "repo_org=$ORGANIZATION_REPO" >> $GITHUB_OUTPUT + echo "perf_number=$RESULTS_TO_COMPARE" >> $GITHUB_OUTPUT + echo "perf_flag=$CALCULATION_METHOD_FLAG" >> $GITHUB_OUTPUT + echo "perf_timeout=$PERFORMANCE_TEST_TIMEOUT" >> $GITHUB_OUTPUT + + call_reusable: + needs: get_config uses: ROCm/migraphx-benchmark/.github/workflows/perf-test.yml@main with: - rocm_release: ${{ github.event.inputs.rocm_release || '6.0.2' }} - result_number: ${{ github.event.inputs.result_number || '10' }} - flags: ${{ github.event.inputs.flags || '-r' }} - performance_reports_repo: ${{ github.event.inputs.performance_reports_repo || 'ROCm/migraphx-reports' }} - performance_backup_repo: ${{ github.event.inputs.performance_backup_repo || 'migraphx-benchmark/performance-backup' }} - benchmark_utils_repo: ${{ github.event.inputs.benchmark_utils_repo || 'ROCm/migraphx-benchmark-utils' }} - organization: ${{ github.event.inputs.organization || 'AMD' }} - model_timeout: ${{ github.event.inputs.model_timeout || '30m' }} + rocm_release: ${{ github.event.inputs.rocm_release || needs.get_config.outputs.rocm_version }} + benchmark_utils_repo: ${{ github.event.inputs.benchmark_utils_repo || needs.get_config.outputs.utils_repo }} + performance_reports_repo: ${{ github.event.inputs.performance_reports_repo || needs.get_config.outputs.reports_repo }} + performance_backup_repo: ${{ github.event.inputs.performance_backup_repo || needs.get_config.outputs.backup_repo }} + organization: ${{ github.event.inputs.organization || needs.get_config.outputs.repo_org }} + result_number: ${{ github.event.inputs.result_number || needs.get_config.outputs.perf_number }} + flags: ${{ github.event.inputs.flags || needs.get_config.outputs.perf_flag }} + model_timeout: ${{ github.event.inputs.model_timeout || needs.get_config.outputs.perf_timeout }} secrets: gh_token: ${{ secrets.MIGRAPHX_BOT_TOKEN }} mail_user: ${{ secrets.MAIL_USERNAME }} - mail_pass: ${{ secrets.MAIL_PASSWORD }} + mail_pass: ${{ secrets.MAIL_PASSWORD }} \ No newline at end of file diff --git a/.github/workflows/sync_rocMLIR.yaml b/.github/workflows/sync_rocMLIR.yaml new file mode 100644 index 00000000000..08001f4175a --- /dev/null +++ b/.github/workflows/sync_rocMLIR.yaml @@ -0,0 +1,90 @@ +name: rocMLIR sync with extended accuracy + +on: + schedule: + - cron: '0 7 * * sun' + pull_request: + branches: [develop] + types: [synchronize, closed] + workflow_dispatch: + inputs: + rocm_release: + type: string + description: ROCm release version + required: true + default: '6.0.2' + base_image: + type: string + description: Base image for ROCm Docker build + required: true + default: 'rocm/dev-ubuntu-20.04' + docker_image: + type: string + description: Docker image name for rocm Docker build + required: true + default: 'rocm-migraphx' + build_navi: + type: string + description: Build navi number + required: true + default: '0' + benchmark_utils_repo: + type: string + description: Repository where benchmark utils are stored + required: true + default: 'ROCm/migraphx-benchmark-utils' + performance_reports_repo: + description: Repository where performance reports are stored + required: true + default: 'ROCm/migraphx-reports' + organization: + type: string + description: Organization based on which location of files will be different + required: true + default: 'AMD' + +jobs: + get_config: + runs-on: ubuntu-latest + outputs: + rocm_version: ${{ steps.read_config.outputs.rocm_version }} + rocm_base_image: ${{ steps.read_config.outputs.rocm_base_image }} + rocm_built_image: ${{ steps.read_config.outputs.rocm_built_image }} + use_navi: ${{ steps.read_config.outputs.use_navi }} + utils_repo: ${{ steps.read_config.outputs.utils_repo }} + reports_repo: ${{ steps.read_config.outputs.reports_repo }} + repo_org: ${{ steps.read_config.outputs.repo_org }} + steps: + - name: checkout + uses: actions/checkout@v4.1.1 + - name: read_config + id: read_config + run: | + ROCM_VERSION=$(grep 'ROCM_VERSION' .github/workflows/config.md | cut -d "'" -f2) + ROCM_BASE_IMAGE=$(grep 'ROCM_BASE_IMAGE' .github/workflows/config.md | cut -d "'" -f2) + ROCM_BUILT_IMAGE=$(grep 'ROCM_BUILT_IMAGE' .github/workflows/config.md | cut -d "'" -f2) + BENCHMARK_UTILS_REPO=$(grep 'BENCHMARK_UTILS_REPO' .github/workflows/config.md | cut -d "'" -f2) + PERFORMANCE_REPORTS_REPO=$(grep 'PERFORMANCE_REPORTS_REPO' .github/workflows/config.md | cut -d "'" -f2) + ORGANIZATION_REPO=$(grep 'ORGANIZATION_REPO' .github/workflows/config.md | cut -d "'" -f2) + USE_NAVI=$(grep 'USE_NAVI' .github/workflows/config.ymd | cut -d "'" -f2) + echo "rocm_version=$ROCM_VERSION" >> $GITHUB_OUTPUT + echo "rocm_base_image=$ROCM_BASE_IMAGE" >> $GITHUB_OUTPUT + echo "rocm_built_image=$ROCM_BUILT_IMAGE" >> $GITHUB_OUTPUT + echo "use_navi=$USE_NAVI" >> $GITHUB_OUTPUT + echo "utils_repo=$BENCHMARK_UTILS_REPO" >> $GITHUB_OUTPUT + echo "reports_repo=$PERFORMANCE_REPORTS_REPO" >> $GITHUB_OUTPUT + echo "repo_org=$ORGANIZATION_REPO" >> $GITHUB_OUTPUT + + call_reusable: + needs: get_config + uses: ROCm/migraphx-benchmark/.github/workflows/rocMLIR_sync.yml@main + with: + rocm_release: ${{ github.event.inputs.rocm_release || needs.get_config.outputs.rocm_version }} + base_image: ${{ github.event.inputs.base_image || needs.get_config.outputs.rocm_base_image }} + docker_image: ${{ github.event.inputs.docker_image || needs.get_config.outputs.rocm_built_image }} + build_navi: ${{ github.event.inputs.build_navi || needs.get_config.outputs.use_navi }} + benchmark_utils_repo: ${{ github.event.inputs.benchmark_utils_repo || needs.get_config.outputs.utils_repo }} + performance_reports_repo: ${{ github.event.inputs.performance_reports_repo || needs.get_config.outputs.reports_repo }} + organization: ${{ github.event.inputs.organization || needs.get_config.outputs.repo_org }} + secrets: + gh_token: ${{ secrets.MIGRAPHX_BOT_TOKEN }} \ No newline at end of file diff --git a/.github/workflows/weekly_master_sync.yaml b/.github/workflows/weekly_master_sync.yaml new file mode 100644 index 00000000000..9273e663c22 --- /dev/null +++ b/.github/workflows/weekly_master_sync.yaml @@ -0,0 +1,25 @@ +name: Master weekly sync + +on: + schedule: + - cron: '0 15 * * sun' + workflow_dispatch: + +jobs: + SyncAndMerge: + name: Sync master and merge develop + runs-on: ubuntu-latest + steps: + + - uses: actions/checkout@v4.1.1 + with: + ref: develop + fetch-depth: '0' + + - name: Merge Fast Forward Only + uses: IgorMirosavljevicHTEC/github-action-merge-fast-forward@v1.0 + with: + branchtomerge: develop + branch: master + env: + GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }} From bceef13f72d7e36a16973646cd6abfda6e8821ad Mon Sep 17 00:00:00 2001 From: Ahsan Saghir <142340507+ahsan-ca@users.noreply.github.com> Date: Thu, 23 May 2024 10:47:24 -0400 Subject: [PATCH 02/13] MIGraphx private lib changes (#3108) --- dev-requirements.txt | 2 +- src/CMakeLists.txt | 1 + src/targets/cpu/CMakeLists.txt | 1 + src/targets/fpga/CMakeLists.txt | 1 + src/targets/gpu/CMakeLists.txt | 1 + src/targets/ref/CMakeLists.txt | 1 + src/tf/CMakeLists.txt | 1 + 7 files changed, 7 insertions(+), 1 deletion(-) diff --git a/dev-requirements.txt b/dev-requirements.txt index cf520e3cb59..c9fce73e910 100644 --- a/dev-requirements.txt +++ b/dev-requirements.txt @@ -26,5 +26,5 @@ facebook/zstd@v1.4.5 -X subdir -DCMAKE_DIR=build/cmake ccache@v4.1 -DENABLE_TESTING=OFF pcre,pfultz2/pcre@8.45 -H sha256:d6f7182602a775a7d500a0cedca6449af0400c6493951513046d17615ed0bf11 danmar/cppcheck@bb2711c22a0be09efe7f1a8da3030876471026c8 -DHAVE_RULES=1 # 2.11 -RadeonOpenCompute/rocm-cmake@5a34e72d9f113eb5d028e740c2def1f944619595 --build +RadeonOpenCompute/rocm-cmake@a83c5075d85f1fd28d657a9277eb21c834d76f3f --build -f requirements.txt diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index eedc7a27501..62197380cc4 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -269,6 +269,7 @@ register_op(migraphx HEADER migraphx/builtin.hpp OPERATORS builtin::literal buil rocm_clang_tidy_check(migraphx) migraphx_generate_export_header(migraphx) rocm_install_targets( + PRIVATE TARGETS migraphx INCLUDE ${CMAKE_CURRENT_SOURCE_DIR}/include diff --git a/src/targets/cpu/CMakeLists.txt b/src/targets/cpu/CMakeLists.txt index 4c08d60123c..e8a57422b01 100755 --- a/src/targets/cpu/CMakeLists.txt +++ b/src/targets/cpu/CMakeLists.txt @@ -97,6 +97,7 @@ else() endif() rocm_install_targets( + PRIVATE TARGETS migraphx_cpu INCLUDE ${CMAKE_CURRENT_SOURCE_DIR}/include diff --git a/src/targets/fpga/CMakeLists.txt b/src/targets/fpga/CMakeLists.txt index 39df3be6400..304c92851c8 100644 --- a/src/targets/fpga/CMakeLists.txt +++ b/src/targets/fpga/CMakeLists.txt @@ -36,6 +36,7 @@ rocm_clang_tidy_check(migraphx_fpga) target_link_libraries(migraphx_fpga migraphx) rocm_install_targets( + PRIVATE TARGETS migraphx_fpga INCLUDE ${CMAKE_CURRENT_SOURCE_DIR}/include diff --git a/src/targets/gpu/CMakeLists.txt b/src/targets/gpu/CMakeLists.txt index d2e73ef1de7..da733498f7d 100644 --- a/src/targets/gpu/CMakeLists.txt +++ b/src/targets/gpu/CMakeLists.txt @@ -332,6 +332,7 @@ add_subdirectory(driver) add_subdirectory(hiprtc) rocm_install_targets( + PRIVATE TARGETS migraphx_gpu migraphx_device compile_for_gpu INCLUDE ${CMAKE_CURRENT_SOURCE_DIR}/include diff --git a/src/targets/ref/CMakeLists.txt b/src/targets/ref/CMakeLists.txt index 0d72afba107..3f442406a49 100644 --- a/src/targets/ref/CMakeLists.txt +++ b/src/targets/ref/CMakeLists.txt @@ -36,6 +36,7 @@ target_link_libraries(migraphx_ref PUBLIC migraphx) migraphx_generate_export_header(migraphx_ref) rocm_install_targets( + PRIVATE TARGETS migraphx_ref INCLUDE ${CMAKE_CURRENT_SOURCE_DIR}/include diff --git a/src/tf/CMakeLists.txt b/src/tf/CMakeLists.txt index 316bd7371ed..6257be2afa3 100644 --- a/src/tf/CMakeLists.txt +++ b/src/tf/CMakeLists.txt @@ -60,6 +60,7 @@ endif() target_link_libraries(migraphx_tf PUBLIC migraphx) rocm_install_targets( + PRIVATE TARGETS migraphx_tf ) From ca6512908d44cd56c622b523dc59cb5e79cd39a1 Mon Sep 17 00:00:00 2001 From: Zakor Gyula <126694206+gyulaz-htec@users.noreply.github.com> Date: Mon, 27 May 2024 21:01:17 +0200 Subject: [PATCH 03/13] Fix mod compile issue on GPU (#2268) (#3086) --- src/include/migraphx/op/mod.hpp | 3 +-- .../kernels/include/migraphx/kernels/math.hpp | 18 +++++++++++++++++- test/py/onnx_backend_test.py | 2 -- 3 files changed, 18 insertions(+), 5 deletions(-) diff --git a/src/include/migraphx/op/mod.hpp b/src/include/migraphx/op/mod.hpp index f1a48e3c58f..38f947a3587 100644 --- a/src/include/migraphx/op/mod.hpp +++ b/src/include/migraphx/op/mod.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -38,7 +38,6 @@ struct mod : binary { auto a = base_attributes(); a["commutative"] = false; - a["point_op"] = "${function:fmod}((${function:remainder}(${0}, ${1})) + ${1}, ${1})"; return a; } auto apply() const diff --git a/src/targets/gpu/kernels/include/migraphx/kernels/math.hpp b/src/targets/gpu/kernels/include/migraphx/kernels/math.hpp index 5a6cca7bc24..da00ff9c781 100644 --- a/src/targets/gpu/kernels/include/migraphx/kernels/math.hpp +++ b/src/targets/gpu/kernels/include/migraphx/kernels/math.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -256,6 +256,21 @@ constexpr auto min(const T& a, const U& b) return min>(a, b); } +template ())> +constexpr T mod(const T& a, const T& b) +{ + if constexpr(is_integral{}) + // onnx mod operator requires numpy style modulus + return ((a % b) + b) % b; + return static_cast(fmod(remainder(a, b) + b, b)); +} + +template {} and not is_any_vec())> +constexpr auto mod(const T& a, const U& b) +{ + return mod>(a, b); +} + MIGRAPHX_DEVICE_MATH_VEC(abs) MIGRAPHX_DEVICE_MATH_VEC(acos) MIGRAPHX_DEVICE_MATH_VEC(acosh) @@ -275,6 +290,7 @@ MIGRAPHX_DEVICE_MATH_VEC(isnan) MIGRAPHX_DEVICE_MATH_VEC(log) MIGRAPHX_DEVICE_MATH_VEC(max) MIGRAPHX_DEVICE_MATH_VEC(min) +MIGRAPHX_DEVICE_MATH_VEC(mod) MIGRAPHX_DEVICE_MATH_VEC(nearbyint) MIGRAPHX_DEVICE_MATH_VEC(pow) MIGRAPHX_DEVICE_MATH_VEC(remainder) diff --git a/test/py/onnx_backend_test.py b/test/py/onnx_backend_test.py index 353bcea3944..2d847c97300 100644 --- a/test/py/onnx_backend_test.py +++ b/test/py/onnx_backend_test.py @@ -119,8 +119,6 @@ def disabled_tests_onnx_1_7_0(backend_test): backend_test.exclude(r'test_einsum_transpose_cpu') backend_test.exclude(r'test_maxunpool_export_with_output_shape_cpu') backend_test.exclude(r'test_maxunpool_export_without_output_shape_cpu') - backend_test.exclude(r'test_mod_mixed_sign_int32_cpu') - backend_test.exclude(r'test_mod_mixed_sign_int8_cpu') backend_test.exclude(r'test_qlinearmatmul_2D_cpu') backend_test.exclude(r'test_qlinearmatmul_3D_cpu') backend_test.exclude(r'test_range_float_type_positive_delta_expanded_cpu') From 2dcca47322b967fc99093ec74d940d261d2713a4 Mon Sep 17 00:00:00 2001 From: Chris Austen Date: Mon, 27 May 2024 15:03:01 -0400 Subject: [PATCH 04/13] Move up pybind to support new builds of ONNX Runtime (#3121) --- requirements.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/requirements.txt b/requirements.txt index 49b42573add..ff97144acf2 100755 --- a/requirements.txt +++ b/requirements.txt @@ -24,7 +24,7 @@ google/protobuf@v3.19.0 -DCMAKE_POSITION_INDEPENDENT_CODE=On -X subdir -Dprotobuf_BUILD_TESTS=Off nlohmann/json@v3.8.0 ROCm/half@rocm-5.6.0 -pybind/pybind11@d159a563383d10c821ba7b2a71905d1207db6de4 --build +pybind/pybind11@3e9dfa2866941655c56877882565e7577de6fc7b --build msgpack/msgpack-c@cpp-3.3.0 -DMSGPACK_BUILD_TESTS=Off sqlite3@3.43.2 -DCMAKE_POSITION_INDEPENDENT_CODE=On ROCm/composable_kernel@57cdd70b7cb14e5e3b60cd9a5f96ba8dc343763e -DCK_BUILD_JIT_LIB=On -DCMAKE_POSITION_INDEPENDENT_CODE=On From 4e13bc1ec36ad7121e6814fb9d4933f53e061e75 Mon Sep 17 00:00:00 2001 From: Igor Mirosavljevic <104914075+IgorMirosavljevicHTEC@users.noreply.github.com> Date: Tue, 28 May 2024 16:11:02 +0200 Subject: [PATCH 05/13] Replace master sync action with commands (#3126) --- .github/workflows/weekly_master_sync.yaml | 10 ++++------ 1 file changed, 4 insertions(+), 6 deletions(-) diff --git a/.github/workflows/weekly_master_sync.yaml b/.github/workflows/weekly_master_sync.yaml index 9273e663c22..732d6a08d63 100644 --- a/.github/workflows/weekly_master_sync.yaml +++ b/.github/workflows/weekly_master_sync.yaml @@ -17,9 +17,7 @@ jobs: fetch-depth: '0' - name: Merge Fast Forward Only - uses: IgorMirosavljevicHTEC/github-action-merge-fast-forward@v1.0 - with: - branchtomerge: develop - branch: master - env: - GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }} + run: | + git checkout master + git merge origin/develop --ff-only + git push origin HEAD From 9a9994557ac1d8630e670a7b12abaa3d4ad63e57 Mon Sep 17 00:00:00 2001 From: "dependabot[bot]" <49699333+dependabot[bot]@users.noreply.github.com> Date: Tue, 28 May 2024 10:12:06 -0400 Subject: [PATCH 06/13] Bump rocm-docs-core from 1.1.2 to 1.2.0 in /docs/sphinx (#3127) --- docs/sphinx/requirements.in | 2 +- docs/sphinx/requirements.txt | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/docs/sphinx/requirements.in b/docs/sphinx/requirements.in index 173c33a0bb9..19a63badcf7 100644 --- a/docs/sphinx/requirements.in +++ b/docs/sphinx/requirements.in @@ -1,2 +1,2 @@ -rocm-docs-core==1.1.2 +rocm-docs-core==1.2.0 sphinx-collapse diff --git a/docs/sphinx/requirements.txt b/docs/sphinx/requirements.txt index 33dfde7f4b0..b95758e80a4 100644 --- a/docs/sphinx/requirements.txt +++ b/docs/sphinx/requirements.txt @@ -92,7 +92,7 @@ requests==2.32.0 # via # pygithub # sphinx -rocm-docs-core==1.1.2 +rocm-docs-core==1.2.0 # via -r requirements.in smmap==5.0.0 # via gitdb From eb731aec9bdb9e2ea65d75a53a64026c0469e83b Mon Sep 17 00:00:00 2001 From: Zakor Gyula <126694206+gyulaz-htec@users.noreply.github.com> Date: Tue, 28 May 2024 16:12:31 +0200 Subject: [PATCH 07/13] Enable test_mod verify test (#3125) --- test/verify/test_fmod_mod.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/test/verify/test_fmod_mod.cpp b/test/verify/test_fmod_mod.cpp index 8c968857f1f..992d88f2760 100644 --- a/test/verify/test_fmod_mod.cpp +++ b/test/verify/test_fmod_mod.cpp @@ -80,6 +80,5 @@ struct test_mod : verify_program> }; template struct test_mod; -// TODO: Fix half type test -// template struct test_mod; +template struct test_mod; template struct test_mod; From e182a61855eccfa8754797be808c69cc2d8b172c Mon Sep 17 00:00:00 2001 From: "github-actions[bot]" <41898282+github-actions[bot]@users.noreply.github.com> Date: Tue, 28 May 2024 10:13:08 -0400 Subject: [PATCH 08/13] Update onnxruntime main 33a68d221f28bd8d412f2e9188e50bac8a255b71 (#3120) --- test/onnx/.onnxrt-commit | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test/onnx/.onnxrt-commit b/test/onnx/.onnxrt-commit index cb6c41e130e..843156e909b 100644 --- a/test/onnx/.onnxrt-commit +++ b/test/onnx/.onnxrt-commit @@ -1 +1 @@ -156d52163dd487ca0a1afaa87a12a3e172a45200 +33a68d221f28bd8d412f2e9188e50bac8a255b71 From 78f55ed3a0947b6f5b5577663ce3bfc98470dd7a Mon Sep 17 00:00:00 2001 From: Brian Pickrell <95253842+bpickrel@users.noreply.github.com> Date: Tue, 28 May 2024 07:13:37 -0700 Subject: [PATCH 09/13] added set_bypass() call (#3116) --- src/fuse_reduce.cpp | 3 +++ test/fuse_reduce.cpp | 12 ++++++++++++ 2 files changed, 15 insertions(+) diff --git a/src/fuse_reduce.cpp b/src/fuse_reduce.cpp index e533856c38a..37dcfc0e028 100644 --- a/src/fuse_reduce.cpp +++ b/src/fuse_reduce.cpp @@ -59,6 +59,8 @@ struct fused_reduce const auto* sm = mods.front(); if(sm->get_output_shapes().size() != 1) MIGRAPHX_THROW("Only one output supported"); + if(not sm->bypass()) + MIGRAPHX_THROW("fused_reduce: bypass flag is not set"); auto names = sm->get_parameter_names(); check_shapes{inputs, *this}.has(names.size()).same_ndims(); std::sort(names.begin(), names.end()); @@ -426,6 +428,7 @@ struct reduce_reshape : rewrite_reshapes_base auto dims = base_dims(inputs); auto* oldm = ins->module_inputs().front(); auto* sm = mpm.create_module(oldm->name() + "_reshape"); + sm->set_bypass(); insert_module_in_submodule(sm, inputs, oldm, transform_op([&](const operation& sop) { if(contains(sop.name(), "reduce")) return make_op(sop.name(), {{"axes", axes}}); diff --git a/test/fuse_reduce.cpp b/test/fuse_reduce.cpp index bc138ff1225..ba6cf4579b5 100644 --- a/test/fuse_reduce.cpp +++ b/test/fuse_reduce.cpp @@ -173,6 +173,12 @@ TEST_CASE(scalar_multibroadcast) reduce_mod->add_instruction(migraphx::make_op("reduce_sum", {{"axes", {1, 2}}}), x0); reduce_mod->add_return({sqrtbc}); + EXPECT(test::throws([&] { + mm->add_instruction( + migraphx::make_op("fused_reduce", {{"axes", {1, 2}}}), {pow}, {reduce_mod}); + })); + // reduce modules must be flagged for bypass when running subsequent passes + reduce_mod->set_bypass(); auto bip = mm->add_instruction( migraphx::make_op("fused_reduce", {{"axes", {1, 2}}}), {pow}, {reduce_mod}); mm->add_return({bip}); @@ -217,6 +223,12 @@ TEST_CASE(scalar_multibroadcast_contiguous) reduce_mod->add_instruction(migraphx::make_op("reduce_sum", {{"axes", {1, 2}}}), x0); reduce_mod->add_return({sqrtbc}); + EXPECT(test::throws([&] { + mm->add_instruction( + migraphx::make_op("fused_reduce", {{"axes", {1, 2}}}), {pow}, {reduce_mod}); + })); + // reduce modules must be flagged for bypass when running subsequent passes + reduce_mod->set_bypass(); auto bip = mm->add_instruction( migraphx::make_op("fused_reduce", {{"axes", {1, 2}}}), {pow}, {reduce_mod}); mm->add_return({bip}); From d7520dd14514b18de6c2e3deeeba093b22e6728c Mon Sep 17 00:00:00 2001 From: "github-actions[bot]" <41898282+github-actions[bot]@users.noreply.github.com> Date: Tue, 28 May 2024 10:15:14 -0400 Subject: [PATCH 10/13] Update rocMLIR main 3612396bca1139abf25e2ed0085fe481d275af89 (#3122) --- requirements.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/requirements.txt b/requirements.txt index ff97144acf2..915e644c4c7 100755 --- a/requirements.txt +++ b/requirements.txt @@ -28,4 +28,4 @@ pybind/pybind11@3e9dfa2866941655c56877882565e7577de6fc7b --build msgpack/msgpack-c@cpp-3.3.0 -DMSGPACK_BUILD_TESTS=Off sqlite3@3.43.2 -DCMAKE_POSITION_INDEPENDENT_CODE=On ROCm/composable_kernel@57cdd70b7cb14e5e3b60cd9a5f96ba8dc343763e -DCK_BUILD_JIT_LIB=On -DCMAKE_POSITION_INDEPENDENT_CODE=On -ROCm/rocMLIR@e50d72fc6ab9a7a792d92a1ba7db6db45e4c508c -DBUILD_FAT_LIBROCKCOMPILER=On +ROCm/rocMLIR@3612396bca1139abf25e2ed0085fe481d275af89 -DBUILD_FAT_LIBROCKCOMPILER=On From 263509b10d5c6887e1396c0fef88c1243903ea45 Mon Sep 17 00:00:00 2001 From: Eddie Liao <54926923+eddieliao@users.noreply.github.com> Date: Tue, 28 May 2024 09:16:54 -0500 Subject: [PATCH 11/13] Split tf tests (#3112) --- test/api/test_tf_parser.cpp | 8 +- test/tf/CMakeLists.txt | 18 +- test/tf/gen_tf_pb.py | 4 +- test/tf/include/tf_conv_utils.hpp | 51 + test/tf/include/tf_test.hpp | 96 ++ test/tf/{ => models}/add_bcast_test.pb | 0 test/tf/{ => models}/add_test.pb | 6 +- test/tf/{ => models}/addv2_test.pb | 0 test/tf/{ => models}/argmax_test.pb | Bin test/tf/{ => models}/argmin_test.pb | Bin .../tf/{ => models}/assert_less_equal_test.pb | Bin test/tf/{ => models}/batchmatmul_test.pb | 0 test/tf/{ => models}/batchnorm_half_test.pb | Bin test/tf/{ => models}/batchnorm_test.pb | Bin test/tf/{ => models}/batchnormv3_test.pb | Bin test/tf/{ => models}/biasadd_scalar_test.pb | Bin test/tf/{ => models}/biasadd_test.pb | 0 test/tf/{ => models}/cast_test.pb | Bin test/tf/{ => models}/concat_test.pb | Bin test/tf/{ => models}/constant_test.pb | Bin test/tf/{ => models}/conv_add_test.pb | Bin test/tf/{ => models}/conv_batch_test.pb | Bin test/tf/{ => models}/conv_nchw_test.pb | Bin test/tf/{ => models}/conv_relu6_test.pb | Bin test/tf/{ => models}/conv_relu_test.pb | Bin test/tf/{ => models}/conv_test.pb | Bin test/tf/{ => models}/depthwise_conv_test.pb | Bin test/tf/{ => models}/expanddims_neg_test.pb | Bin test/tf/{ => models}/expanddims_test.pb | Bin test/tf/{ => models}/gather_test.pb | Bin test/tf/{ => models}/identity_test.pb | 0 test/tf/{ => models}/matmul_test.pb | 0 test/tf/{ => models}/mean_test.pb | Bin test/tf/{ => models}/mean_test_nhwc.pb | Bin test/tf/{ => models}/mul_test.pb | 0 test/tf/{ => models}/multi_output_test.pb | 0 test/tf/{ => models}/noop_test.pb | 0 test/tf/{ => models}/onehot_test.pb | Bin test/tf/{ => models}/pack_test.pb | 0 test/tf/{ => models}/pack_test_nhwc.pb | 0 test/tf/{ => models}/pad_test.pb | Bin test/tf/{ => models}/pooling_test.pb | 0 test/tf/{ => models}/pow_test.pb | 0 test/tf/{ => models}/relu6_half_test.pb | 0 test/tf/{ => models}/relu6_test.pb | 0 test/tf/{ => models}/relu_test.pb | 0 test/tf/{ => models}/reshape_test.pb | Bin test/tf/{ => models}/rsqrt_test.pb | 0 test/tf/{ => models}/shape_test.pb | 0 test/tf/{ => models}/slice_test.pb | Bin test/tf/{ => models}/softmax_test.pb | 0 test/tf/{ => models}/split_test.pb | Bin test/tf/{ => models}/split_test_one_output.pb | Bin .../split_test_vector_as_input.pb | Bin test/tf/{ => models}/sqdiff_test.pb | 0 test/tf/{ => models}/squeeze_test.pb | Bin test/tf/{ => models}/stopgradient_test.pb | 0 .../{ => models}/stridedslice_masks_test.pb | Bin test/tf/{ => models}/stridedslice_test.pb | Bin test/tf/{ => models}/sub_test.pb | 0 test/tf/{ => models}/tanh_test.pb | 0 test/tf/{ => models}/transpose_test.pb | Bin test/tf/{ => models}/variable_batch_test.pb | Bin test/tf/tests/add_bcast_test.cpp | 43 + test/tf/tests/add_test.cpp | 39 + test/tf/tests/addv2_test.cpp | 38 + test/tf/tests/argmax_test.cpp | 41 + test/tf/tests/argmin_test.cpp | 41 + test/tf/tests/assert_less_equal_test.cpp | 44 + test/tf/tests/batchmatmul_test.cpp | 45 + test/tf/tests/batchnorm_half_test.cpp | 56 + test/tf/tests/batchnorm_test.cpp | 56 + test/tf/tests/batchnormv3_test.cpp | 56 + test/tf/tests/biasadd_scalar_test.cpp | 44 + test/tf/tests/biasadd_test.cpp | 43 + test/tf/tests/cast_test.cpp | 41 + test/tf/tests/concat_test.cpp | 46 + test/tf/tests/constant_test.cpp | 37 + test/tf/tests/conv_add_test.cpp | 38 + test/tf/tests/conv_nchw_test.cpp | 35 + test/tf/tests/conv_relu6_test.cpp | 45 + test/tf/tests/conv_relu_test.cpp | 38 + test/tf/tests/conv_test.cpp | 35 + test/tf/tests/depthwise_conv_test.cpp | 54 + test/tf/tests/expanddims_neg_test.cpp | 41 + test/tf/tests/expanddims_test.cpp | 40 + test/tf/tests/gather_test.cpp | 44 + test/tf/tests/identity_test.cpp | 38 + test/tf/tests/main.cpp | 28 + test/tf/tests/matmul_test.cpp | 45 + test/tf/tests/mean_test.cpp | 44 + test/tf/tests/mean_test_nhwc.cpp | 43 + test/tf/tests/mul_test.cpp | 40 + test/tf/tests/multi_output_test.cpp | 42 + test/tf/tests/noop_test.cpp | 34 + test/tf/tests/onehot_test.cpp | 45 + test/tf/tests/pack_test.cpp | 52 + test/tf/tests/pack_test_nhwc.cpp | 58 + test/tf/tests/pad_test.cpp | 43 + test/tf/tests/pooling_test.cpp | 45 + test/tf/tests/pow_test.cpp | 39 + test/tf/tests/relu6_half_test.cpp | 47 + test/tf/tests/relu6_test.cpp | 45 + test/tf/tests/relu_test.cpp | 38 + test/tf/tests/reshape_test.cpp | 41 + test/tf/tests/rsqrt_test.cpp | 38 + test/tf/tests/shape_test.cpp | 39 + test/tf/tests/slice_test.cpp | 44 + test/tf/tests/softmax_test.cpp | 38 + test/tf/tests/split_test.cpp | 51 + test/tf/tests/split_test_one_output.cpp | 41 + test/tf/tests/split_test_vector_as_input.cpp | 53 + test/tf/tests/sqdiff_test.cpp | 39 + test/tf/tests/squeeze_test.cpp | 38 + test/tf/tests/stopgradient_test.cpp | 38 + test/tf/tests/stridedslice_masks_test.cpp | 54 + test/tf/tests/stridedslice_test.cpp | 45 + test/tf/tests/sub_test.cpp | 40 + test/tf/tests/tanh_test.cpp | 39 + test/tf/tests/transpose_test.cpp | 40 + test/tf/tests/variable_batch_test.cpp | 38 + test/tf/tf_test.cpp | 1080 ----------------- 122 files changed, 2650 insertions(+), 1095 deletions(-) mode change 100644 => 100755 test/tf/CMakeLists.txt create mode 100644 test/tf/include/tf_conv_utils.hpp create mode 100644 test/tf/include/tf_test.hpp rename test/tf/{ => models}/add_bcast_test.pb (100%) rename test/tf/{ => models}/add_test.pb (78%) rename test/tf/{ => models}/addv2_test.pb (100%) rename test/tf/{ => models}/argmax_test.pb (100%) rename test/tf/{ => models}/argmin_test.pb (100%) rename test/tf/{ => models}/assert_less_equal_test.pb (100%) rename test/tf/{ => models}/batchmatmul_test.pb (100%) rename test/tf/{ => models}/batchnorm_half_test.pb (100%) rename test/tf/{ => models}/batchnorm_test.pb (100%) rename test/tf/{ => models}/batchnormv3_test.pb (100%) rename test/tf/{ => models}/biasadd_scalar_test.pb (100%) rename test/tf/{ => models}/biasadd_test.pb (100%) rename test/tf/{ => models}/cast_test.pb (100%) rename test/tf/{ => models}/concat_test.pb (100%) rename test/tf/{ => models}/constant_test.pb (100%) rename test/tf/{ => models}/conv_add_test.pb (100%) rename test/tf/{ => models}/conv_batch_test.pb (100%) rename test/tf/{ => models}/conv_nchw_test.pb (100%) rename test/tf/{ => models}/conv_relu6_test.pb (100%) rename test/tf/{ => models}/conv_relu_test.pb (100%) rename test/tf/{ => models}/conv_test.pb (100%) rename test/tf/{ => models}/depthwise_conv_test.pb (100%) rename test/tf/{ => models}/expanddims_neg_test.pb (100%) rename test/tf/{ => models}/expanddims_test.pb (100%) rename test/tf/{ => models}/gather_test.pb (100%) rename test/tf/{ => models}/identity_test.pb (100%) rename test/tf/{ => models}/matmul_test.pb (100%) rename test/tf/{ => models}/mean_test.pb (100%) rename test/tf/{ => models}/mean_test_nhwc.pb (100%) rename test/tf/{ => models}/mul_test.pb (100%) rename test/tf/{ => models}/multi_output_test.pb (100%) rename test/tf/{ => models}/noop_test.pb (100%) rename test/tf/{ => models}/onehot_test.pb (100%) rename test/tf/{ => models}/pack_test.pb (100%) rename test/tf/{ => models}/pack_test_nhwc.pb (100%) rename test/tf/{ => models}/pad_test.pb (100%) rename test/tf/{ => models}/pooling_test.pb (100%) rename test/tf/{ => models}/pow_test.pb (100%) rename test/tf/{ => models}/relu6_half_test.pb (100%) rename test/tf/{ => models}/relu6_test.pb (100%) rename test/tf/{ => models}/relu_test.pb (100%) rename test/tf/{ => models}/reshape_test.pb (100%) rename test/tf/{ => models}/rsqrt_test.pb (100%) rename test/tf/{ => models}/shape_test.pb (100%) rename test/tf/{ => models}/slice_test.pb (100%) rename test/tf/{ => models}/softmax_test.pb (100%) rename test/tf/{ => models}/split_test.pb (100%) rename test/tf/{ => models}/split_test_one_output.pb (100%) rename test/tf/{ => models}/split_test_vector_as_input.pb (100%) rename test/tf/{ => models}/sqdiff_test.pb (100%) rename test/tf/{ => models}/squeeze_test.pb (100%) rename test/tf/{ => models}/stopgradient_test.pb (100%) rename test/tf/{ => models}/stridedslice_masks_test.pb (100%) rename test/tf/{ => models}/stridedslice_test.pb (100%) rename test/tf/{ => models}/sub_test.pb (100%) rename test/tf/{ => models}/tanh_test.pb (100%) rename test/tf/{ => models}/transpose_test.pb (100%) rename test/tf/{ => models}/variable_batch_test.pb (100%) create mode 100644 test/tf/tests/add_bcast_test.cpp create mode 100644 test/tf/tests/add_test.cpp create mode 100644 test/tf/tests/addv2_test.cpp create mode 100644 test/tf/tests/argmax_test.cpp create mode 100644 test/tf/tests/argmin_test.cpp create mode 100644 test/tf/tests/assert_less_equal_test.cpp create mode 100644 test/tf/tests/batchmatmul_test.cpp create mode 100644 test/tf/tests/batchnorm_half_test.cpp create mode 100644 test/tf/tests/batchnorm_test.cpp create mode 100644 test/tf/tests/batchnormv3_test.cpp create mode 100644 test/tf/tests/biasadd_scalar_test.cpp create mode 100644 test/tf/tests/biasadd_test.cpp create mode 100644 test/tf/tests/cast_test.cpp create mode 100644 test/tf/tests/concat_test.cpp create mode 100644 test/tf/tests/constant_test.cpp create mode 100644 test/tf/tests/conv_add_test.cpp create mode 100644 test/tf/tests/conv_nchw_test.cpp create mode 100644 test/tf/tests/conv_relu6_test.cpp create mode 100644 test/tf/tests/conv_relu_test.cpp create mode 100644 test/tf/tests/conv_test.cpp create mode 100644 test/tf/tests/depthwise_conv_test.cpp create mode 100644 test/tf/tests/expanddims_neg_test.cpp create mode 100644 test/tf/tests/expanddims_test.cpp create mode 100644 test/tf/tests/gather_test.cpp create mode 100644 test/tf/tests/identity_test.cpp create mode 100644 test/tf/tests/main.cpp create mode 100644 test/tf/tests/matmul_test.cpp create mode 100644 test/tf/tests/mean_test.cpp create mode 100644 test/tf/tests/mean_test_nhwc.cpp create mode 100644 test/tf/tests/mul_test.cpp create mode 100644 test/tf/tests/multi_output_test.cpp create mode 100644 test/tf/tests/noop_test.cpp create mode 100644 test/tf/tests/onehot_test.cpp create mode 100644 test/tf/tests/pack_test.cpp create mode 100644 test/tf/tests/pack_test_nhwc.cpp create mode 100644 test/tf/tests/pad_test.cpp create mode 100644 test/tf/tests/pooling_test.cpp create mode 100644 test/tf/tests/pow_test.cpp create mode 100644 test/tf/tests/relu6_half_test.cpp create mode 100644 test/tf/tests/relu6_test.cpp create mode 100644 test/tf/tests/relu_test.cpp create mode 100644 test/tf/tests/reshape_test.cpp create mode 100644 test/tf/tests/rsqrt_test.cpp create mode 100644 test/tf/tests/shape_test.cpp create mode 100644 test/tf/tests/slice_test.cpp create mode 100644 test/tf/tests/softmax_test.cpp create mode 100644 test/tf/tests/split_test.cpp create mode 100644 test/tf/tests/split_test_one_output.cpp create mode 100644 test/tf/tests/split_test_vector_as_input.cpp create mode 100644 test/tf/tests/sqdiff_test.cpp create mode 100644 test/tf/tests/squeeze_test.cpp create mode 100644 test/tf/tests/stopgradient_test.cpp create mode 100644 test/tf/tests/stridedslice_masks_test.cpp create mode 100644 test/tf/tests/stridedslice_test.cpp create mode 100644 test/tf/tests/sub_test.cpp create mode 100644 test/tf/tests/tanh_test.cpp create mode 100644 test/tf/tests/transpose_test.cpp create mode 100644 test/tf/tests/variable_batch_test.cpp delete mode 100644 test/tf/tf_test.cpp diff --git a/test/api/test_tf_parser.cpp b/test/api/test_tf_parser.cpp index 4cdc23fb9f9..b440e6b64f9 100644 --- a/test/api/test_tf_parser.cpp +++ b/test/api/test_tf_parser.cpp @@ -27,7 +27,7 @@ TEST_CASE(load_tf) { - auto p = migraphx::parse_tf("add_test.pb"); + auto p = migraphx::parse_tf("models/add_test.pb"); auto shapes = p.get_output_shapes(); CHECK(shapes.size() == 1); } @@ -38,7 +38,7 @@ TEST_CASE(load_tf_default_dim) size_t batch = 2; tf_options.set_default_dim_value(batch); tf_options.set_nhwc(); - auto p = migraphx::parse_tf("conv_batch_test.pb", tf_options); + auto p = migraphx::parse_tf("models/conv_batch_test.pb", tf_options); auto shapes = p.get_output_shapes(); CHECK(shapes.size() == 1); CHECK(shapes.front().lengths().front() == batch); @@ -50,7 +50,7 @@ TEST_CASE(load_tf_param_shape) std::vector new_shape{1, 3}; tf_options.set_input_parameter_shape("0", new_shape); tf_options.set_input_parameter_shape("1", new_shape); - auto p = migraphx::parse_tf("add_test.pb", tf_options); + auto p = migraphx::parse_tf("models/add_test.pb", tf_options); auto shapes = p.get_output_shapes(); CHECK(shapes.size() == 1); CHECK(shapes.front().lengths() == new_shape); @@ -60,7 +60,7 @@ TEST_CASE(load_tf_multi_outputs) { migraphx::tf_options tf_options; tf_options.set_output_names({"relu", "tanh"}); - auto p = migraphx::parse_tf("multi_output_test.pb", tf_options); + auto p = migraphx::parse_tf("models/multi_output_test.pb", tf_options); auto shapes = p.get_output_shapes(); CHECK(shapes.size() == 2); } diff --git a/test/tf/CMakeLists.txt b/test/tf/CMakeLists.txt old mode 100644 new mode 100755 index 7d59303b558..d99b10b9cf9 --- a/test/tf/CMakeLists.txt +++ b/test/tf/CMakeLists.txt @@ -22,11 +22,17 @@ # THE SOFTWARE. ##################################################################################### +function(add_tf_test TEST_NAME) + rocm_add_test_executable(${TEST_NAME} ${ARGN}) + rocm_clang_tidy_check(${TEST_NAME}) + target_link_libraries(${TEST_NAME} migraphx_tf pb_files) + target_include_directories(${TEST_NAME} PUBLIC ../include include) +endfunction() + include(Embed) -file(GLOB_RECURSE PB_FILES CONFIGURE_DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/*.pb) -add_embed_library(pb_files ${PB_FILES} RELATIVE ${CMAKE_CURRENT_SOURCE_DIR}) +file(GLOB_RECURSE PB_FILES CONFIGURE_DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/models/*.pb) +add_embed_library(pb_files ${PB_FILES} RELATIVE ${CMAKE_CURRENT_SOURCE_DIR}/models) + +file(GLOB TF_TESTS CONFIGRE_DEPENDS tests/*.cpp) -rocm_add_test_executable(test_tf tf_test.cpp) -rocm_clang_tidy_check(test_tf) -target_link_libraries(test_tf migraphx_tf pb_files) -target_include_directories(test_tf PUBLIC ../include) +add_tf_test(test_tf ${TF_TESTS}) diff --git a/test/tf/gen_tf_pb.py b/test/tf/gen_tf_pb.py index b1cc59ad0f0..b177194dabc 100644 --- a/test/tf/gen_tf_pb.py +++ b/test/tf/gen_tf_pb.py @@ -1,7 +1,7 @@ ##################################################################################### # The MIT License (MIT) # -# Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. +# Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. # # Permission is hereby granted, free of charge, to any person obtaining a copy # of this software and associated documentation files (the "Software"), to deal @@ -32,7 +32,7 @@ def run_test(): g1 = tf.Graph() op_test(g1) tf.io.write_graph(g1, - '.', + './models', '{}.pb'.format(op_test.__name__), as_text=False) diff --git a/test/tf/include/tf_conv_utils.hpp b/test/tf/include/tf_conv_utils.hpp new file mode 100644 index 00000000000..13f825c99f1 --- /dev/null +++ b/test/tf/include/tf_conv_utils.hpp @@ -0,0 +1,51 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#ifndef MIGRAPHX_GUARD_TEST_TF_TF_CONV_UTILS_HPP +#define MIGRAPHX_GUARD_TEST_TF_TF_CONV_UTILS_HPP + +inline migraphx::program create_conv() +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); + std::vector weight_data(3 * 3 * 3 * 32); + std::fill(weight_data.begin(), weight_data.end(), 1.0f); + auto l1 = + mm->add_literal(migraphx::shape{migraphx::shape::float_type, {3, 3, 3, 32}}, weight_data); + + migraphx::op::convolution op; + op.padding = {1, 1, 1, 1}; + op.stride = {1, 1}; + op.dilation = {1, 1}; + auto l2 = + mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {3, 2, 0, 1}}}), l1); + mm->add_instruction(op, l0, l2); + return p; +} + +#endif diff --git a/test/tf/include/tf_test.hpp b/test/tf/include/tf_test.hpp new file mode 100644 index 00000000000..d477befead1 --- /dev/null +++ b/test/tf/include/tf_test.hpp @@ -0,0 +1,96 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ + +#ifndef MIGRAPHX_GUARD_TEST_TF_TF_TEST_HPP +#define MIGRAPHX_GUARD_TEST_TF_TF_TEST_HPP + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +#include "test.hpp" + +inline migraphx::program read_pb_file(const std::string& name, const migraphx::tf_options& options) +{ + static auto pb_files{::pb_files()}; + if(pb_files.find(name) == pb_files.end()) + { + std::cerr << "Can not find TensorFlow Protobuf file by name: " << name + << " , aborting the program\n" + << std::endl; + std::abort(); + } + return migraphx::parse_tf_buffer(std::string{pb_files.at(name)}, options); +} + +inline migraphx::program +parse_tf(const std::string& name, + bool is_nhwc, + const std::unordered_map>& dim_params = {}, + const std::vector& output_node_names = {}) +{ + + return read_pb_file(name, migraphx::tf_options{is_nhwc, 1, dim_params, output_node_names}); +} + +inline migraphx::program optimize_tf(const std::string& name, bool is_nhwc) +{ + auto prog = read_pb_file(name, migraphx::tf_options{is_nhwc, 1}); + auto* mm = prog.get_main_module(); + if(is_nhwc) + migraphx::run_passes(*mm, + {migraphx::simplify_reshapes{}, + migraphx::dead_code_elimination{}, + migraphx::eliminate_identity{}}); + + // remove the last return instruction + + if(mm->size() > 0) + { + auto last_ins = std::prev(mm->end()); + if(last_ins->name() == "@return") + { + mm->remove_instruction(last_ins); + } + } + return prog; +} + +#endif diff --git a/test/tf/add_bcast_test.pb b/test/tf/models/add_bcast_test.pb similarity index 100% rename from test/tf/add_bcast_test.pb rename to test/tf/models/add_bcast_test.pb diff --git a/test/tf/add_test.pb b/test/tf/models/add_test.pb similarity index 78% rename from test/tf/add_test.pb rename to test/tf/models/add_test.pb index f176c1b2b93..41cc3fef3ed 100644 --- a/test/tf/add_test.pb +++ b/test/tf/models/add_test.pb @@ -7,6 +7,6 @@ 1 Placeholder* dtype0* shape: - -add1Add01* -T0"¸ \ No newline at end of file + +add1AddV201* +T0"æ \ No newline at end of file diff --git a/test/tf/addv2_test.pb b/test/tf/models/addv2_test.pb similarity index 100% rename from test/tf/addv2_test.pb rename to test/tf/models/addv2_test.pb diff --git a/test/tf/argmax_test.pb b/test/tf/models/argmax_test.pb similarity index 100% rename from test/tf/argmax_test.pb rename to test/tf/models/argmax_test.pb diff --git a/test/tf/argmin_test.pb b/test/tf/models/argmin_test.pb similarity index 100% rename from test/tf/argmin_test.pb rename to test/tf/models/argmin_test.pb diff --git a/test/tf/assert_less_equal_test.pb b/test/tf/models/assert_less_equal_test.pb similarity index 100% rename from test/tf/assert_less_equal_test.pb rename to test/tf/models/assert_less_equal_test.pb diff --git a/test/tf/batchmatmul_test.pb b/test/tf/models/batchmatmul_test.pb similarity index 100% rename from test/tf/batchmatmul_test.pb rename to test/tf/models/batchmatmul_test.pb diff --git a/test/tf/batchnorm_half_test.pb b/test/tf/models/batchnorm_half_test.pb similarity index 100% rename from test/tf/batchnorm_half_test.pb rename to test/tf/models/batchnorm_half_test.pb diff --git a/test/tf/batchnorm_test.pb b/test/tf/models/batchnorm_test.pb similarity index 100% rename from test/tf/batchnorm_test.pb rename to test/tf/models/batchnorm_test.pb diff --git a/test/tf/batchnormv3_test.pb b/test/tf/models/batchnormv3_test.pb similarity index 100% rename from test/tf/batchnormv3_test.pb rename to test/tf/models/batchnormv3_test.pb diff --git a/test/tf/biasadd_scalar_test.pb b/test/tf/models/biasadd_scalar_test.pb similarity index 100% rename from test/tf/biasadd_scalar_test.pb rename to test/tf/models/biasadd_scalar_test.pb diff --git a/test/tf/biasadd_test.pb b/test/tf/models/biasadd_test.pb similarity index 100% rename from test/tf/biasadd_test.pb rename to test/tf/models/biasadd_test.pb diff --git a/test/tf/cast_test.pb b/test/tf/models/cast_test.pb similarity index 100% rename from test/tf/cast_test.pb rename to test/tf/models/cast_test.pb diff --git a/test/tf/concat_test.pb b/test/tf/models/concat_test.pb similarity index 100% rename from test/tf/concat_test.pb rename to test/tf/models/concat_test.pb diff --git a/test/tf/constant_test.pb b/test/tf/models/constant_test.pb similarity index 100% rename from test/tf/constant_test.pb rename to test/tf/models/constant_test.pb diff --git a/test/tf/conv_add_test.pb b/test/tf/models/conv_add_test.pb similarity index 100% rename from test/tf/conv_add_test.pb rename to test/tf/models/conv_add_test.pb diff --git a/test/tf/conv_batch_test.pb b/test/tf/models/conv_batch_test.pb similarity index 100% rename from test/tf/conv_batch_test.pb rename to test/tf/models/conv_batch_test.pb diff --git a/test/tf/conv_nchw_test.pb b/test/tf/models/conv_nchw_test.pb similarity index 100% rename from test/tf/conv_nchw_test.pb rename to test/tf/models/conv_nchw_test.pb diff --git a/test/tf/conv_relu6_test.pb b/test/tf/models/conv_relu6_test.pb similarity index 100% rename from test/tf/conv_relu6_test.pb rename to test/tf/models/conv_relu6_test.pb diff --git a/test/tf/conv_relu_test.pb b/test/tf/models/conv_relu_test.pb similarity index 100% rename from test/tf/conv_relu_test.pb rename to test/tf/models/conv_relu_test.pb diff --git a/test/tf/conv_test.pb b/test/tf/models/conv_test.pb similarity index 100% rename from test/tf/conv_test.pb rename to test/tf/models/conv_test.pb diff --git a/test/tf/depthwise_conv_test.pb b/test/tf/models/depthwise_conv_test.pb similarity index 100% rename from test/tf/depthwise_conv_test.pb rename to test/tf/models/depthwise_conv_test.pb diff --git a/test/tf/expanddims_neg_test.pb b/test/tf/models/expanddims_neg_test.pb similarity index 100% rename from test/tf/expanddims_neg_test.pb rename to test/tf/models/expanddims_neg_test.pb diff --git a/test/tf/expanddims_test.pb b/test/tf/models/expanddims_test.pb similarity index 100% rename from test/tf/expanddims_test.pb rename to test/tf/models/expanddims_test.pb diff --git a/test/tf/gather_test.pb b/test/tf/models/gather_test.pb similarity index 100% rename from test/tf/gather_test.pb rename to test/tf/models/gather_test.pb diff --git a/test/tf/identity_test.pb b/test/tf/models/identity_test.pb similarity index 100% rename from test/tf/identity_test.pb rename to test/tf/models/identity_test.pb diff --git a/test/tf/matmul_test.pb b/test/tf/models/matmul_test.pb similarity index 100% rename from test/tf/matmul_test.pb rename to test/tf/models/matmul_test.pb diff --git a/test/tf/mean_test.pb b/test/tf/models/mean_test.pb similarity index 100% rename from test/tf/mean_test.pb rename to test/tf/models/mean_test.pb diff --git a/test/tf/mean_test_nhwc.pb b/test/tf/models/mean_test_nhwc.pb similarity index 100% rename from test/tf/mean_test_nhwc.pb rename to test/tf/models/mean_test_nhwc.pb diff --git a/test/tf/mul_test.pb b/test/tf/models/mul_test.pb similarity index 100% rename from test/tf/mul_test.pb rename to test/tf/models/mul_test.pb diff --git a/test/tf/multi_output_test.pb b/test/tf/models/multi_output_test.pb similarity index 100% rename from test/tf/multi_output_test.pb rename to test/tf/models/multi_output_test.pb diff --git a/test/tf/noop_test.pb b/test/tf/models/noop_test.pb similarity index 100% rename from test/tf/noop_test.pb rename to test/tf/models/noop_test.pb diff --git a/test/tf/onehot_test.pb b/test/tf/models/onehot_test.pb similarity index 100% rename from test/tf/onehot_test.pb rename to test/tf/models/onehot_test.pb diff --git a/test/tf/pack_test.pb b/test/tf/models/pack_test.pb similarity index 100% rename from test/tf/pack_test.pb rename to test/tf/models/pack_test.pb diff --git a/test/tf/pack_test_nhwc.pb b/test/tf/models/pack_test_nhwc.pb similarity index 100% rename from test/tf/pack_test_nhwc.pb rename to test/tf/models/pack_test_nhwc.pb diff --git a/test/tf/pad_test.pb b/test/tf/models/pad_test.pb similarity index 100% rename from test/tf/pad_test.pb rename to test/tf/models/pad_test.pb diff --git a/test/tf/pooling_test.pb b/test/tf/models/pooling_test.pb similarity index 100% rename from test/tf/pooling_test.pb rename to test/tf/models/pooling_test.pb diff --git a/test/tf/pow_test.pb b/test/tf/models/pow_test.pb similarity index 100% rename from test/tf/pow_test.pb rename to test/tf/models/pow_test.pb diff --git a/test/tf/relu6_half_test.pb b/test/tf/models/relu6_half_test.pb similarity index 100% rename from test/tf/relu6_half_test.pb rename to test/tf/models/relu6_half_test.pb diff --git a/test/tf/relu6_test.pb b/test/tf/models/relu6_test.pb similarity index 100% rename from test/tf/relu6_test.pb rename to test/tf/models/relu6_test.pb diff --git a/test/tf/relu_test.pb b/test/tf/models/relu_test.pb similarity index 100% rename from test/tf/relu_test.pb rename to test/tf/models/relu_test.pb diff --git a/test/tf/reshape_test.pb b/test/tf/models/reshape_test.pb similarity index 100% rename from test/tf/reshape_test.pb rename to test/tf/models/reshape_test.pb diff --git a/test/tf/rsqrt_test.pb b/test/tf/models/rsqrt_test.pb similarity index 100% rename from test/tf/rsqrt_test.pb rename to test/tf/models/rsqrt_test.pb diff --git a/test/tf/shape_test.pb b/test/tf/models/shape_test.pb similarity index 100% rename from test/tf/shape_test.pb rename to test/tf/models/shape_test.pb diff --git a/test/tf/slice_test.pb b/test/tf/models/slice_test.pb similarity index 100% rename from test/tf/slice_test.pb rename to test/tf/models/slice_test.pb diff --git a/test/tf/softmax_test.pb b/test/tf/models/softmax_test.pb similarity index 100% rename from test/tf/softmax_test.pb rename to test/tf/models/softmax_test.pb diff --git a/test/tf/split_test.pb b/test/tf/models/split_test.pb similarity index 100% rename from test/tf/split_test.pb rename to test/tf/models/split_test.pb diff --git a/test/tf/split_test_one_output.pb b/test/tf/models/split_test_one_output.pb similarity index 100% rename from test/tf/split_test_one_output.pb rename to test/tf/models/split_test_one_output.pb diff --git a/test/tf/split_test_vector_as_input.pb b/test/tf/models/split_test_vector_as_input.pb similarity index 100% rename from test/tf/split_test_vector_as_input.pb rename to test/tf/models/split_test_vector_as_input.pb diff --git a/test/tf/sqdiff_test.pb b/test/tf/models/sqdiff_test.pb similarity index 100% rename from test/tf/sqdiff_test.pb rename to test/tf/models/sqdiff_test.pb diff --git a/test/tf/squeeze_test.pb b/test/tf/models/squeeze_test.pb similarity index 100% rename from test/tf/squeeze_test.pb rename to test/tf/models/squeeze_test.pb diff --git a/test/tf/stopgradient_test.pb b/test/tf/models/stopgradient_test.pb similarity index 100% rename from test/tf/stopgradient_test.pb rename to test/tf/models/stopgradient_test.pb diff --git a/test/tf/stridedslice_masks_test.pb b/test/tf/models/stridedslice_masks_test.pb similarity index 100% rename from test/tf/stridedslice_masks_test.pb rename to test/tf/models/stridedslice_masks_test.pb diff --git a/test/tf/stridedslice_test.pb b/test/tf/models/stridedslice_test.pb similarity index 100% rename from test/tf/stridedslice_test.pb rename to test/tf/models/stridedslice_test.pb diff --git a/test/tf/sub_test.pb b/test/tf/models/sub_test.pb similarity index 100% rename from test/tf/sub_test.pb rename to test/tf/models/sub_test.pb diff --git a/test/tf/tanh_test.pb b/test/tf/models/tanh_test.pb similarity index 100% rename from test/tf/tanh_test.pb rename to test/tf/models/tanh_test.pb diff --git a/test/tf/transpose_test.pb b/test/tf/models/transpose_test.pb similarity index 100% rename from test/tf/transpose_test.pb rename to test/tf/models/transpose_test.pb diff --git a/test/tf/variable_batch_test.pb b/test/tf/models/variable_batch_test.pb similarity index 100% rename from test/tf/variable_batch_test.pb rename to test/tf/models/variable_batch_test.pb diff --git a/test/tf/tests/add_bcast_test.cpp b/test/tf/tests/add_bcast_test.cpp new file mode 100644 index 00000000000..84e71d3d4d2 --- /dev/null +++ b/test/tf/tests/add_bcast_test.cpp @@ -0,0 +1,43 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(add_bcast_test) +{ + + migraphx::program p; + + auto* mm = p.get_main_module(); + migraphx::shape s0{migraphx::shape::float_type, {2, 3}}; + auto l0 = mm->add_parameter("0", s0); + auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {2, 1}}); + auto l2 = + mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", s0.lens()}}), l1); + mm->add_instruction(migraphx::make_op("add"), l0, l2); + auto prog = optimize_tf("add_bcast_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/add_test.cpp b/test/tf/tests/add_test.cpp new file mode 100644 index 00000000000..eab59fe9c0c --- /dev/null +++ b/test/tf/tests/add_test.cpp @@ -0,0 +1,39 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(add_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); + auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); + mm->add_instruction(migraphx::make_op("add"), l0, l1); + auto prog = optimize_tf("add_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/addv2_test.cpp b/test/tf/tests/addv2_test.cpp new file mode 100644 index 00000000000..18d162b0ba7 --- /dev/null +++ b/test/tf/tests/addv2_test.cpp @@ -0,0 +1,38 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(addv2_test) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); + auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); + mm->add_instruction(migraphx::make_op("add"), l0, l1); + auto prog = optimize_tf("addv2_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/argmax_test.cpp b/test/tf/tests/argmax_test.cpp new file mode 100644 index 00000000000..683bd750a0d --- /dev/null +++ b/test/tf/tests/argmax_test.cpp @@ -0,0 +1,41 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(argmax_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {4, 5, 6, 7}}); + mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::int32_type}, {2}}); + auto ins = mm->add_instruction(migraphx::make_op("argmax", {{"axis", 2}}), l0); + auto l1 = mm->add_instruction(migraphx::make_op("squeeze", {{"axes", {2}}}), ins); + mm->add_return({l1}); + auto prog = parse_tf("argmax_test.pb", false, {{"0", {4, 5, 6, 7}}}); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/argmin_test.cpp b/test/tf/tests/argmin_test.cpp new file mode 100644 index 00000000000..81f34538f2f --- /dev/null +++ b/test/tf/tests/argmin_test.cpp @@ -0,0 +1,41 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(argmin_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {3, 4, 5, 6}}); + mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::int32_type}, {2}}); + auto ins = mm->add_instruction(migraphx::make_op("argmin", {{"axis", 2}}), l0); + auto l1 = mm->add_instruction(migraphx::make_op("squeeze", {{"axes", {2}}}), ins); + mm->add_return({l1}); + auto prog = parse_tf("argmin_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/assert_less_equal_test.cpp b/test/tf/tests/assert_less_equal_test.cpp new file mode 100644 index 00000000000..327f02c4ebb --- /dev/null +++ b/test/tf/tests/assert_less_equal_test.cpp @@ -0,0 +1,44 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(assert_less_equal_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + migraphx::shape s0{migraphx::shape::float_type, {2, 3}}; + auto l0 = mm->add_parameter("0", s0); + auto l1 = mm->add_parameter("1", s0); + migraphx::literal l{migraphx::shape{migraphx::shape::int32_type, {2}}, {0, 1}}; + auto l2 = mm->add_literal(l); + mm->add_instruction(migraphx::make_op("add"), l0, l1); + auto l3 = mm->add_instruction(migraphx::make_op("identity"), l0, l1); + mm->add_instruction(migraphx::make_op("identity"), l3, l2); + auto prog = optimize_tf("assert_less_equal_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/batchmatmul_test.cpp b/test/tf/tests/batchmatmul_test.cpp new file mode 100644 index 00000000000..aa2871d5dc1 --- /dev/null +++ b/test/tf/tests/batchmatmul_test.cpp @@ -0,0 +1,45 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(batchmatmul_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 2, 8, 4}}); + auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {1, 2, 4, 8}}); + + auto trans_l0 = + mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 1, 3, 2}}}), l0); + auto trans_l1 = + mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 1, 3, 2}}}), l1); + + mm->add_instruction(migraphx::make_op("dot"), trans_l0, trans_l1); + auto prog = optimize_tf("batchmatmul_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/batchnorm_half_test.cpp b/test/tf/tests/batchnorm_half_test.cpp new file mode 100644 index 00000000000..1281fd3678e --- /dev/null +++ b/test/tf/tests/batchnorm_half_test.cpp @@ -0,0 +1,56 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(batchnorm_half_test) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + + auto x = mm->add_parameter("x", {migraphx::shape::half_type, {1, 32, 16, 16}}); + auto bias = mm->add_parameter("bias", {migraphx::shape::float_type, {32}}); + auto mean = mm->add_parameter("mean", {migraphx::shape::float_type, {32}}); + auto var = mm->add_parameter("variance", {migraphx::shape::float_type, {32}}); + + std::vector scale_data(32, 1.0); + auto scale = mm->add_literal(migraphx::shape{migraphx::shape::float_type, {32}}, scale_data); + auto eps = mm->add_literal(migraphx::literal{migraphx::shape::half_type, {1e-4f}}); + + auto usq_scale = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), scale); + auto usq_bias = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), bias); + auto usq_mean = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), mean); + auto usq_var = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), var); + + auto x_sub_mean = add_common_op(*mm, migraphx::make_op("sub"), {x, usq_mean}); + auto var_eps = add_common_op(*mm, migraphx::make_op("add"), {usq_var, eps}); + auto rsqrt = mm->add_instruction(migraphx::make_op("rsqrt"), var_eps); + auto mul0 = add_common_op(*mm, migraphx::make_op("mul"), {usq_scale, rsqrt}); + auto r0 = add_common_op(*mm, migraphx::make_op("mul"), {x_sub_mean, mul0}); + add_common_op(*mm, migraphx::make_op("add"), {r0, usq_bias}); + + auto prog = optimize_tf("batchnorm_half_test.pb", true); + EXPECT(p == prog); +} diff --git a/test/tf/tests/batchnorm_test.cpp b/test/tf/tests/batchnorm_test.cpp new file mode 100644 index 00000000000..4a549ccb53a --- /dev/null +++ b/test/tf/tests/batchnorm_test.cpp @@ -0,0 +1,56 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(batchnorm_test) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + + auto x = mm->add_parameter("x", {migraphx::shape::float_type, {1, 32, 16, 16}}); + auto bias = mm->add_parameter("bias", {migraphx::shape::float_type, {32}}); + auto mean = mm->add_parameter("mean", {migraphx::shape::float_type, {32}}); + auto var = mm->add_parameter("variance", {migraphx::shape::float_type, {32}}); + + std::vector scale_data(32, 1.0); + auto scale = mm->add_literal(migraphx::shape{migraphx::shape::float_type, {32}}, scale_data); + auto eps = mm->add_literal(migraphx::literal{migraphx::shape::float_type, {1e-4f}}); + + auto usq_scale = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), scale); + auto usq_bias = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), bias); + auto usq_mean = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), mean); + auto usq_var = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), var); + + auto x_sub_mean = add_common_op(*mm, migraphx::make_op("sub"), {x, usq_mean}); + auto var_eps = add_common_op(*mm, migraphx::make_op("add"), {usq_var, eps}); + auto rsqrt = mm->add_instruction(migraphx::make_op("rsqrt"), var_eps); + auto mul0 = add_common_op(*mm, migraphx::make_op("mul"), {usq_scale, rsqrt}); + auto r0 = add_common_op(*mm, migraphx::make_op("mul"), {x_sub_mean, mul0}); + add_common_op(*mm, migraphx::make_op("add"), {r0, usq_bias}); + + auto prog = optimize_tf("batchnorm_test.pb", true); + EXPECT(p == prog); +} diff --git a/test/tf/tests/batchnormv3_test.cpp b/test/tf/tests/batchnormv3_test.cpp new file mode 100644 index 00000000000..2450e9c4331 --- /dev/null +++ b/test/tf/tests/batchnormv3_test.cpp @@ -0,0 +1,56 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(batchnormv3_test) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + + auto x = mm->add_parameter("x", {migraphx::shape::float_type, {1, 32, 16, 16}}); + auto bias = mm->add_parameter("bias", {migraphx::shape::float_type, {32}}); + auto mean = mm->add_parameter("mean", {migraphx::shape::float_type, {32}}); + auto var = mm->add_parameter("variance", {migraphx::shape::float_type, {32}}); + + std::vector scale_data(32, 1.0); + auto scale = mm->add_literal(migraphx::shape{migraphx::shape::float_type, {32}}, scale_data); + auto eps = mm->add_literal(migraphx::literal{migraphx::shape::float_type, {1e-6f}}); + + auto usq_scale = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), scale); + auto usq_bias = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), bias); + auto usq_mean = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), mean); + auto usq_var = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), var); + + auto x_sub_mean = add_common_op(*mm, migraphx::make_op("sub"), {x, usq_mean}); + auto var_eps = add_common_op(*mm, migraphx::make_op("add"), {usq_var, eps}); + auto rsqrt = mm->add_instruction(migraphx::make_op("rsqrt"), var_eps); + auto mul0 = add_common_op(*mm, migraphx::make_op("mul"), {usq_scale, rsqrt}); + auto r0 = add_common_op(*mm, migraphx::make_op("mul"), {x_sub_mean, mul0}); + add_common_op(*mm, migraphx::make_op("add"), {r0, usq_bias}); + + auto prog = optimize_tf("batchnormv3_test.pb", true); + EXPECT(p == prog); +} diff --git a/test/tf/tests/biasadd_scalar_test.cpp b/test/tf/tests/biasadd_scalar_test.cpp new file mode 100644 index 00000000000..13af83d6027 --- /dev/null +++ b/test/tf/tests/biasadd_scalar_test.cpp @@ -0,0 +1,44 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(biasadd_scalar_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + migraphx::shape s0{migraphx::shape::float_type, {1, 1}}; + uint64_t axis = 1; + auto l0 = mm->add_parameter("0", s0); + auto l1 = mm->add_literal( + migraphx::literal{migraphx::shape{migraphx::shape::float_type, {1}, {0}}, {1.0}}); + auto l2 = mm->add_instruction( + migraphx::make_op("broadcast", {{"axis", axis}, {"out_lens", l0->get_shape().lens()}}), l1); + mm->add_instruction(migraphx::make_op("add"), l0, l2); + auto prog = optimize_tf("biasadd_scalar_test.pb", true); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/biasadd_test.cpp b/test/tf/tests/biasadd_test.cpp new file mode 100644 index 00000000000..caf09833230 --- /dev/null +++ b/test/tf/tests/biasadd_test.cpp @@ -0,0 +1,43 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(biasadd_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + migraphx::shape s0{migraphx::shape::float_type, {1, 500, 1, 1}}; + uint64_t axis = 1; + auto l0 = mm->add_parameter("0", s0); + auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {500}}); + auto l2 = mm->add_instruction( + migraphx::make_op("broadcast", {{"axis", axis}, {"out_lens", l0->get_shape().lens()}}), l1); + mm->add_instruction(migraphx::make_op("add"), l0, l2); + auto prog = optimize_tf("biasadd_test.pb", true); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/cast_test.cpp b/test/tf/tests/cast_test.cpp new file mode 100644 index 00000000000..c67f7b77754 --- /dev/null +++ b/test/tf/tests/cast_test.cpp @@ -0,0 +1,41 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(cast_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); + mm->add_instruction( + migraphx::make_op("convert", + {{"target_type", migraphx::to_value(migraphx::shape::int32_type)}}), + l0); + auto prog = optimize_tf("cast_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/concat_test.cpp b/test/tf/tests/concat_test.cpp new file mode 100644 index 00000000000..afac63368c0 --- /dev/null +++ b/test/tf/tests/concat_test.cpp @@ -0,0 +1,46 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(concat_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {4, 7, 3}}); + auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {4, 2, 3}}); + + int axis = 1; + // tf uses axis as the third input, and it is in int32 format + // add the literal using a vector in order to set stride to 1 (like in tf parser) + mm->add_literal(migraphx::shape{migraphx::shape::int32_type}, std::vector{axis}); + + mm->add_instruction(migraphx::make_op("concat", {{"axis", axis}}), l0, l1); + auto prog = optimize_tf("concat_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/constant_test.cpp b/test/tf/tests/constant_test.cpp new file mode 100644 index 00000000000..b173a48e5fc --- /dev/null +++ b/test/tf/tests/constant_test.cpp @@ -0,0 +1,37 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(const_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + mm->add_literal(migraphx::shape{migraphx::shape::float_type}, std::vector{1.0f}); + auto prog = optimize_tf("constant_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/conv_add_test.cpp b/test/tf/tests/conv_add_test.cpp new file mode 100644 index 00000000000..a29655d21ed --- /dev/null +++ b/test/tf/tests/conv_add_test.cpp @@ -0,0 +1,38 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include +#include + +TEST_CASE(conv_add_test) +{ + migraphx::program p = create_conv(); + auto* mm = p.get_main_module(); + auto l0 = std::prev(mm->end()); + mm->add_instruction(migraphx::make_op("add"), l0, l0); + auto prog = optimize_tf("conv_add_test.pb", true); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/conv_nchw_test.cpp b/test/tf/tests/conv_nchw_test.cpp new file mode 100644 index 00000000000..bc2c80e01fa --- /dev/null +++ b/test/tf/tests/conv_nchw_test.cpp @@ -0,0 +1,35 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include +#include + +TEST_CASE(conv_nchw_test) +{ + migraphx::program p = create_conv(); + auto prog = optimize_tf("conv_nchw_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/conv_relu6_test.cpp b/test/tf/tests/conv_relu6_test.cpp new file mode 100644 index 00000000000..ed01dd87e14 --- /dev/null +++ b/test/tf/tests/conv_relu6_test.cpp @@ -0,0 +1,45 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include +#include + +TEST_CASE(conv_relu6_test) +{ + migraphx::program p = create_conv(); + auto* mm = p.get_main_module(); + std::vector input_lens{1, 32, 16, 16}; + auto l0 = std::prev(mm->end()); + auto min_val = mm->add_literal(0.0f); + auto max_val = mm->add_literal(6.0f); + min_val = mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", input_lens}}), + min_val); + max_val = mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", input_lens}}), + max_val); + mm->add_instruction(migraphx::make_op("clip"), l0, min_val, max_val); + auto prog = optimize_tf("conv_relu6_test.pb", true); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/conv_relu_test.cpp b/test/tf/tests/conv_relu_test.cpp new file mode 100644 index 00000000000..295d66f22eb --- /dev/null +++ b/test/tf/tests/conv_relu_test.cpp @@ -0,0 +1,38 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include +#include + +TEST_CASE(conv_relu_test) +{ + migraphx::program p = create_conv(); + auto* mm = p.get_main_module(); + auto l0 = std::prev(mm->end()); + mm->add_instruction(migraphx::make_op("relu"), l0); + auto prog = optimize_tf("conv_relu_test.pb", true); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/conv_test.cpp b/test/tf/tests/conv_test.cpp new file mode 100644 index 00000000000..69249765197 --- /dev/null +++ b/test/tf/tests/conv_test.cpp @@ -0,0 +1,35 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include +#include + +TEST_CASE(conv_test) +{ + migraphx::program p = create_conv(); + auto prog = optimize_tf("conv_test.pb", true); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/depthwise_conv_test.cpp b/test/tf/tests/depthwise_conv_test.cpp new file mode 100644 index 00000000000..53839fe96d5 --- /dev/null +++ b/test/tf/tests/depthwise_conv_test.cpp @@ -0,0 +1,54 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include +#include + +TEST_CASE(depthwiseconv_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); + std::vector weight_data(3 * 3 * 3 * 1); + std::fill(weight_data.begin(), weight_data.end(), 1.0f); + auto l1 = + mm->add_literal(migraphx::shape{migraphx::shape::float_type, {3, 3, 3, 1}}, weight_data); + + migraphx::op::convolution op; + op.padding = {1, 1}; + op.stride = {1, 1}; + op.dilation = {1, 1}; + op.group = 3; + auto l3 = + mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {3, 2, 0, 1}}}), l1); + auto l4 = mm->add_instruction(migraphx::make_op("contiguous"), l3); + auto l5 = mm->add_instruction(migraphx::make_op("reshape", {{"dims", {3, 1, 3, 3}}}), l4); + mm->add_instruction(op, l0, l5); + auto prog = optimize_tf("depthwise_conv_test.pb", true); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/expanddims_neg_test.cpp b/test/tf/tests/expanddims_neg_test.cpp new file mode 100644 index 00000000000..3e659bc912d --- /dev/null +++ b/test/tf/tests/expanddims_neg_test.cpp @@ -0,0 +1,41 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(expanddims_test_neg_dims) +{ + // this check makes sure the pb parses negative dim value correctly + migraphx::program p; + + auto* mm = p.get_main_module(); + + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {2, 3, 4}}); + mm->add_literal(-1); + mm->add_instruction(migraphx::make_op("reshape", {{"dims", {2, 3, 4, 1}}}), l0); + auto prog = optimize_tf("expanddims_neg_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/expanddims_test.cpp b/test/tf/tests/expanddims_test.cpp new file mode 100644 index 00000000000..85295d83257 --- /dev/null +++ b/test/tf/tests/expanddims_test.cpp @@ -0,0 +1,40 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(expanddims_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {2, 3, 4}}); + mm->add_literal(0); + mm->add_instruction(migraphx::make_op("reshape", {{"dims", {1, 2, 3, 4}}}), l0); + auto prog = optimize_tf("expanddims_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/gather_test.cpp b/test/tf/tests/gather_test.cpp new file mode 100644 index 00000000000..ef96f51523d --- /dev/null +++ b/test/tf/tests/gather_test.cpp @@ -0,0 +1,44 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(gather_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {2, 4}}); + auto l1 = mm->add_literal( + migraphx::literal{migraphx::shape{migraphx::shape::int32_type, {2}}, {1, 1}}); + mm->add_literal(1); + + int axis = 1; + mm->add_instruction(migraphx::make_op("gather", {{"axis", axis}}), l0, l1); + auto prog = optimize_tf("gather_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/identity_test.cpp b/test/tf/tests/identity_test.cpp new file mode 100644 index 00000000000..ae79c3160ac --- /dev/null +++ b/test/tf/tests/identity_test.cpp @@ -0,0 +1,38 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(identity_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); + mm->add_instruction(migraphx::make_op("identity"), l0); + auto prog = optimize_tf("identity_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/main.cpp b/test/tf/tests/main.cpp new file mode 100644 index 00000000000..336c0391aa6 --- /dev/null +++ b/test/tf/tests/main.cpp @@ -0,0 +1,28 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +int main(int argc, const char* argv[]) { test::run(argc, argv); } diff --git a/test/tf/tests/matmul_test.cpp b/test/tf/tests/matmul_test.cpp new file mode 100644 index 00000000000..82f46bbf9ce --- /dev/null +++ b/test/tf/tests/matmul_test.cpp @@ -0,0 +1,45 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(matmul_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {8, 4}}); + auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {4, 8}}); + + auto trans_l0 = + mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {1, 0}}}), l0); + auto trans_l1 = + mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {1, 0}}}), l1); + + mm->add_instruction(migraphx::make_op("dot"), trans_l0, trans_l1); + auto prog = optimize_tf("matmul_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/mean_test.cpp b/test/tf/tests/mean_test.cpp new file mode 100644 index 00000000000..62a19631a37 --- /dev/null +++ b/test/tf/tests/mean_test.cpp @@ -0,0 +1,44 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(mean_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + migraphx::literal l{migraphx::shape{migraphx::shape::int32_type, {2}}, {2, 3}}; + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); + mm->add_literal(l); + mm->add_literal(l); + migraphx::op::reduce_mean op{{2, 3}}; + mm->add_instruction(op, l0); + auto l3 = mm->add_instruction(op, l0); + mm->add_instruction(migraphx::make_op("squeeze", {{"axes", {2, 3}}}), l3); + auto prog = optimize_tf("mean_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/mean_test_nhwc.cpp b/test/tf/tests/mean_test_nhwc.cpp new file mode 100644 index 00000000000..b3f8eed2b7e --- /dev/null +++ b/test/tf/tests/mean_test_nhwc.cpp @@ -0,0 +1,43 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(mean_test_nhwc) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + migraphx::literal l{migraphx::shape{migraphx::shape::int32_type, {2}}, {1, 2}}; + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); + auto l1 = + mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 2, 3, 1}}}), l0); + migraphx::op::reduce_mean op{{1, 2}}; + auto l2 = mm->add_instruction(op, l1); + mm->add_instruction(migraphx::make_op("squeeze", {{"axes", {1, 2}}}), l2); + auto prog = optimize_tf("mean_test_nhwc.pb", true); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/mul_test.cpp b/test/tf/tests/mul_test.cpp new file mode 100644 index 00000000000..101cd6b6b82 --- /dev/null +++ b/test/tf/tests/mul_test.cpp @@ -0,0 +1,40 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(mul_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 1, 1, 16}}); + auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {1, 1, 1, 16}}); + + mm->add_instruction(migraphx::make_op("mul"), l0, l1); + auto prog = optimize_tf("mul_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/multi_output_test.cpp b/test/tf/tests/multi_output_test.cpp new file mode 100644 index 00000000000..06039b122b6 --- /dev/null +++ b/test/tf/tests/multi_output_test.cpp @@ -0,0 +1,42 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(multi_output_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); + auto l1 = mm->add_instruction(migraphx::make_op("relu"), l0); + auto l2 = mm->add_instruction(migraphx::make_op("tanh"), l0); + mm->add_return({l1, l2}); + + EXPECT(test::throws([&] { parse_tf("multi_output_test.pb", false, {}, {"relu", "relu6"}); })); + auto prog = parse_tf("multi_output_test.pb", false, {}, {"relu", "tanh"}); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/noop_test.cpp b/test/tf/tests/noop_test.cpp new file mode 100644 index 00000000000..06e82269fd0 --- /dev/null +++ b/test/tf/tests/noop_test.cpp @@ -0,0 +1,34 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(noop_test) +{ + migraphx::program p; + auto prog = optimize_tf("noop_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/onehot_test.cpp b/test/tf/tests/onehot_test.cpp new file mode 100644 index 00000000000..2cb67411590 --- /dev/null +++ b/test/tf/tests/onehot_test.cpp @@ -0,0 +1,45 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(onehot_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_literal( + migraphx::literal{migraphx::shape{migraphx::shape::int32_type, {5}}, {1, 1, 1, 1, 1}}); + mm->add_literal(2); + mm->add_literal(1.0f); + mm->add_literal(0.0f); + auto l1 = mm->add_literal( + migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2, 2}}, {1, 0, 0, 1}}); + int axis = 0; + mm->add_instruction(migraphx::make_op("gather", {{"axis", axis}}), l1, l0); + auto prog = optimize_tf("onehot_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/pack_test.cpp b/test/tf/tests/pack_test.cpp new file mode 100644 index 00000000000..7c20df43a08 --- /dev/null +++ b/test/tf/tests/pack_test.cpp @@ -0,0 +1,52 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(pack_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {2}}); + auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {2}}); + auto l2 = mm->add_parameter("2", migraphx::shape{migraphx::shape::float_type, {2}}); + std::vector args{l0, l1, l2}; + std::vector unsqueezed_args; + int64_t axis = 1; + + std::transform( + args.begin(), + args.end(), + std::back_inserter(unsqueezed_args), + [&](migraphx::instruction_ref arg) { + return mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {axis}}}), arg); + }); + mm->add_instruction(migraphx::make_op("concat", {{"axis", static_cast(axis)}}), + unsqueezed_args); + auto prog = optimize_tf("pack_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/pack_test_nhwc.cpp b/test/tf/tests/pack_test_nhwc.cpp new file mode 100644 index 00000000000..d87e53a56ba --- /dev/null +++ b/test/tf/tests/pack_test_nhwc.cpp @@ -0,0 +1,58 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(pack_test_nhwc) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 2, 1, 1}}); + auto lt0 = + mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 2, 3, 1}}}), l0); + auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {1, 2, 1, 1}}); + auto lt1 = + mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 2, 3, 1}}}), l1); + auto l2 = mm->add_parameter("2", migraphx::shape{migraphx::shape::float_type, {1, 2, 1, 1}}); + auto lt2 = + mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 2, 3, 1}}}), l2); + std::vector args{lt0, lt1, lt2}; + std::vector unsqueezed_args; + int64_t nchw_axis = 3; + + std::transform(args.begin(), + args.end(), + std::back_inserter(unsqueezed_args), + [&](migraphx::instruction_ref arg) { + return mm->add_instruction( + migraphx::make_op("unsqueeze", {{"axes", {nchw_axis}}}), arg); + }); + mm->add_instruction(migraphx::make_op("concat", {{"axis", static_cast(nchw_axis)}}), + unsqueezed_args); + auto prog = optimize_tf("pack_test_nhwc.pb", true); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/pad_test.cpp b/test/tf/tests/pad_test.cpp new file mode 100644 index 00000000000..431eff47236 --- /dev/null +++ b/test/tf/tests/pad_test.cpp @@ -0,0 +1,43 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(pad_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {2, 4}}); + std::vector pad_literals{1, 1, 2, 2}; + std::vector pads{1, 2, 1, 2}; + mm->add_literal(migraphx::shape{migraphx::shape::int32_type, {2, 2}}, pad_literals); + + mm->add_instruction(migraphx::make_op("pad", {{"pads", pads}}), l0); + auto prog = optimize_tf("pad_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/pooling_test.cpp b/test/tf/tests/pooling_test.cpp new file mode 100644 index 00000000000..3105328eff2 --- /dev/null +++ b/test/tf/tests/pooling_test.cpp @@ -0,0 +1,45 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(pooling_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); + migraphx::op::pooling avg_pool_op{migraphx::op::pooling_mode::average}; + migraphx::op::pooling max_pool_op{migraphx::op::pooling_mode::max}; + avg_pool_op.stride = {2, 2}; + max_pool_op.stride = {2, 2}; + avg_pool_op.lengths = {2, 2}; + max_pool_op.lengths = {2, 2}; + mm->add_instruction(avg_pool_op, l0); + mm->add_instruction(max_pool_op, l0); + auto prog = optimize_tf("pooling_test.pb", true); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/pow_test.cpp b/test/tf/tests/pow_test.cpp new file mode 100644 index 00000000000..30f6a01f9ce --- /dev/null +++ b/test/tf/tests/pow_test.cpp @@ -0,0 +1,39 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(pow_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); + auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); + mm->add_instruction(migraphx::make_op("pow"), l0, l1); + auto prog = optimize_tf("pow_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/relu6_half_test.cpp b/test/tf/tests/relu6_half_test.cpp new file mode 100644 index 00000000000..84ade6d366f --- /dev/null +++ b/test/tf/tests/relu6_half_test.cpp @@ -0,0 +1,47 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(relu6_half_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + std::vector input_lens{1, 3, 16, 16}; + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::half_type, input_lens}); + auto min_val = + mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::half_type}, {0.0f}}); + auto max_val = + mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::half_type}, {6.0f}}); + min_val = mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", input_lens}}), + min_val); + max_val = mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", input_lens}}), + max_val); + mm->add_instruction(migraphx::make_op("clip"), l0, min_val, max_val); + auto prog = optimize_tf("relu6_half_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/relu6_test.cpp b/test/tf/tests/relu6_test.cpp new file mode 100644 index 00000000000..5da363d62b7 --- /dev/null +++ b/test/tf/tests/relu6_test.cpp @@ -0,0 +1,45 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(relu6_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + std::vector input_lens{1, 3, 16, 16}; + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, input_lens}); + auto min_val = mm->add_literal(0.0f); + auto max_val = mm->add_literal(6.0f); + min_val = mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", input_lens}}), + min_val); + max_val = mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", input_lens}}), + max_val); + mm->add_instruction(migraphx::make_op("clip"), l0, min_val, max_val); + auto prog = optimize_tf("relu6_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/relu_test.cpp b/test/tf/tests/relu_test.cpp new file mode 100644 index 00000000000..a549df6193d --- /dev/null +++ b/test/tf/tests/relu_test.cpp @@ -0,0 +1,38 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(relu_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); + mm->add_instruction(migraphx::make_op("relu"), l0); + auto prog = optimize_tf("relu_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/reshape_test.cpp b/test/tf/tests/reshape_test.cpp new file mode 100644 index 00000000000..2700421de05 --- /dev/null +++ b/test/tf/tests/reshape_test.cpp @@ -0,0 +1,41 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(reshape_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {16}}); + migraphx::shape s0{migraphx::shape::int32_type, {4}}; + // in tf, the second arg is a literal that contains new dimensions + mm->add_literal(migraphx::literal{s0, {1, 1, 1, 16}}); + mm->add_instruction(migraphx::make_op("reshape", {{"dims", {1, 1, 1, 16}}}), l0); + auto prog = optimize_tf("reshape_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/rsqrt_test.cpp b/test/tf/tests/rsqrt_test.cpp new file mode 100644 index 00000000000..5cbf5b287e6 --- /dev/null +++ b/test/tf/tests/rsqrt_test.cpp @@ -0,0 +1,38 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(rsqrt_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); + mm->add_instruction(migraphx::make_op("rsqrt"), l0); + auto prog = optimize_tf("rsqrt_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/shape_test.cpp b/test/tf/tests/shape_test.cpp new file mode 100644 index 00000000000..f749fbdefb4 --- /dev/null +++ b/test/tf/tests/shape_test.cpp @@ -0,0 +1,39 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(shape_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); + mm->add_literal( + migraphx::literal{migraphx::shape{migraphx::shape::int32_type, {4}}, {1, 3, 16, 16}}); + auto prog = optimize_tf("shape_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/slice_test.cpp b/test/tf/tests/slice_test.cpp new file mode 100644 index 00000000000..8db6b04e3d7 --- /dev/null +++ b/test/tf/tests/slice_test.cpp @@ -0,0 +1,44 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(slice_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + std::size_t num_axes = 2; + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {5, 10}}); + migraphx::shape s0{migraphx::shape::int32_type, {num_axes}}; + mm->add_literal(migraphx::literal{s0, {1, 0}}); + mm->add_literal(migraphx::literal{s0, {2, -1}}); + + mm->add_instruction( + migraphx::make_op("slice", {{"starts", {1, 0}}, {"ends", {3, 10}}, {"axes", {0, 1}}}), l0); + auto prog = optimize_tf("slice_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/softmax_test.cpp b/test/tf/tests/softmax_test.cpp new file mode 100644 index 00000000000..80b39b30dfb --- /dev/null +++ b/test/tf/tests/softmax_test.cpp @@ -0,0 +1,38 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(softmax_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3}}); + mm->add_instruction(migraphx::make_op("softmax", {{"axis", 1}}), l0); + auto prog = optimize_tf("softmax_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/split_test.cpp b/test/tf/tests/split_test.cpp new file mode 100644 index 00000000000..6df77ea18b2 --- /dev/null +++ b/test/tf/tests/split_test.cpp @@ -0,0 +1,51 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(split_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + std::vector axes{0, 1}; + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {5, 30}}); + mm->add_literal(3); // num_splits + mm->add_literal(1); // split axis + mm->add_literal(1); // concat axis + mm->add_literal(1); // concat axis + auto l1 = mm->add_instruction( + migraphx::make_op("slice", {{"axes", axes}, {"starts", {0, 0}}, {"ends", {5, 10}}}), l0); + auto l2 = mm->add_instruction( + migraphx::make_op("slice", {{"axes", axes}, {"starts", {0, 10}}, {"ends", {5, 20}}}), l0); + auto l3 = mm->add_instruction( + migraphx::make_op("slice", {{"axes", axes}, {"starts", {0, 20}}, {"ends", {5, 30}}}), l0); + auto l4 = mm->add_instruction(migraphx::make_op("concat", {{"axis", 1}}), l1, l2); + auto l5 = mm->add_instruction(migraphx::make_op("concat", {{"axis", 1}}), l2, l3); + mm->add_return({l4, l5}); + auto prog = parse_tf("split_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/split_test_one_output.cpp b/test/tf/tests/split_test_one_output.cpp new file mode 100644 index 00000000000..309ab55a4f0 --- /dev/null +++ b/test/tf/tests/split_test_one_output.cpp @@ -0,0 +1,41 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(split_test_one_output) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {5, 30}}); + mm->add_literal(1); // num_splits + mm->add_literal(1); // split axis + auto l1 = mm->add_instruction(migraphx::make_op("identity"), l0); + mm->add_return({l1}); + auto prog = parse_tf("split_test_one_output.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/split_test_vector_as_input.cpp b/test/tf/tests/split_test_vector_as_input.cpp new file mode 100644 index 00000000000..96d566beef5 --- /dev/null +++ b/test/tf/tests/split_test_vector_as_input.cpp @@ -0,0 +1,53 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(split_test_vector_as_input) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + std::vector axes{0, 1}; + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {5, 30}}); + // split sizes + mm->add_literal( + migraphx::literal{migraphx::shape{migraphx::shape::int32_type, {3}}, {4, 15, 11}}); + mm->add_literal(1); // split axis + mm->add_literal(1); // concat axis + mm->add_literal(1); // concat axis + auto l1 = mm->add_instruction( + migraphx::make_op("slice", {{"axes", axes}, {"starts", {0, 0}}, {"ends", {5, 4}}}), l0); + auto l2 = mm->add_instruction( + migraphx::make_op("slice", {{"axes", axes}, {"starts", {0, 4}}, {"ends", {5, 19}}}), l0); + auto l3 = mm->add_instruction( + migraphx::make_op("slice", {{"axes", axes}, {"starts", {0, 19}}, {"ends", {5, 30}}}), l0); + auto l4 = mm->add_instruction(migraphx::make_op("concat", {{"axis", 1}}), l1, l2); + auto l5 = mm->add_instruction(migraphx::make_op("concat", {{"axis", 1}}), l2, l3); + mm->add_return({l4, l5}); + auto prog = parse_tf("split_test_vector_as_input.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/sqdiff_test.cpp b/test/tf/tests/sqdiff_test.cpp new file mode 100644 index 00000000000..d1ca678a84c --- /dev/null +++ b/test/tf/tests/sqdiff_test.cpp @@ -0,0 +1,39 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(sqdiff_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); + auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); + mm->add_instruction(migraphx::make_op("sqdiff"), l0, l1); + auto prog = optimize_tf("sqdiff_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/squeeze_test.cpp b/test/tf/tests/squeeze_test.cpp new file mode 100644 index 00000000000..da6e6fbf1ee --- /dev/null +++ b/test/tf/tests/squeeze_test.cpp @@ -0,0 +1,38 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(squeeze_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 2, 3, 1}}); + mm->add_instruction(migraphx::make_op("squeeze", {{"axes", {0, 3}}}), l0); + auto prog = optimize_tf("squeeze_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/stopgradient_test.cpp b/test/tf/tests/stopgradient_test.cpp new file mode 100644 index 00000000000..4d3d6d15fcf --- /dev/null +++ b/test/tf/tests/stopgradient_test.cpp @@ -0,0 +1,38 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(stopgradient_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); + mm->add_instruction(migraphx::make_op("identity"), l0); + auto prog = optimize_tf("stopgradient_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/stridedslice_masks_test.cpp b/test/tf/tests/stridedslice_masks_test.cpp new file mode 100644 index 00000000000..5d2466bd780 --- /dev/null +++ b/test/tf/tests/stridedslice_masks_test.cpp @@ -0,0 +1,54 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(stridedslice_masks_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 10, 3, 3}}); + // add literals for starts, ends, and strides in tf (NHWC format) + mm->add_literal(migraphx::shape{migraphx::shape::int32_type, {4}}, + std::vector{0, 1, 1, 0}); + mm->add_literal(migraphx::shape{migraphx::shape::int32_type, {4}}, + std::vector{0, 0, 0, 0}); + mm->add_literal(migraphx::shape{migraphx::shape::int32_type, {4}}, + std::vector{1, 1, 1, 1}); + + auto l1 = + mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 2, 3, 1}}}), l0); + auto l2 = mm->add_instruction( + migraphx::make_op( + "slice", {{"starts", {0, 1, 1, 0}}, {"ends", {1, 3, 3, 10}}, {"axes", {0, 1, 2, 3}}}), + l1); + auto l3 = + mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 3, 1, 2}}}), l2); + mm->add_return({l3}); + auto prog = parse_tf("stridedslice_masks_test.pb", true); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/stridedslice_test.cpp b/test/tf/tests/stridedslice_test.cpp new file mode 100644 index 00000000000..514ea53fb7a --- /dev/null +++ b/test/tf/tests/stridedslice_test.cpp @@ -0,0 +1,45 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(stridedslice_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 10, 1, 1}}); + auto l1 = + mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 2, 3, 1}}}), l0); + auto l2 = mm->add_instruction( + migraphx::make_op( + "slice", {{"starts", {0, 0, 0, 0}}, {"ends", {1, 1, 1, 5}}, {"axes", {0, 1, 2, 3}}}), + l1); + auto shrink_axis = 1; + mm->add_instruction(migraphx::make_op("squeeze", {{"axes", {shrink_axis}}}), l2); + auto prog = optimize_tf("stridedslice_test.pb", true); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/sub_test.cpp b/test/tf/tests/sub_test.cpp new file mode 100644 index 00000000000..1830fff3f0b --- /dev/null +++ b/test/tf/tests/sub_test.cpp @@ -0,0 +1,40 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(sub_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); + auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); + auto l2 = mm->add_instruction(migraphx::make_op("sub"), l0, l1); + mm->add_return({l2}); + auto prog = parse_tf("sub_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/tanh_test.cpp b/test/tf/tests/tanh_test.cpp new file mode 100644 index 00000000000..b686312d8c2 --- /dev/null +++ b/test/tf/tests/tanh_test.cpp @@ -0,0 +1,39 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(tanh_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); + auto l1 = mm->add_instruction(migraphx::make_op("tanh"), l0); + mm->add_return({l1}); + auto prog = parse_tf("tanh_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/transpose_test.cpp b/test/tf/tests/transpose_test.cpp new file mode 100644 index 00000000000..5f290520d7f --- /dev/null +++ b/test/tf/tests/transpose_test.cpp @@ -0,0 +1,40 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(transpose_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); + migraphx::shape s0{migraphx::shape::int32_type, {4}}; + mm->add_literal(migraphx::literal{s0, {0, 2, 3, 1}}); + mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 2, 3, 1}}}), l0); + auto prog = optimize_tf("transpose_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tests/variable_batch_test.cpp b/test/tf/tests/variable_batch_test.cpp new file mode 100644 index 00000000000..fe4e56d6d40 --- /dev/null +++ b/test/tf/tests/variable_batch_test.cpp @@ -0,0 +1,38 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * + */ + +#include + +TEST_CASE(variable_batch_test) +{ + migraphx::program p; + + auto* mm = p.get_main_module(); + auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); + mm->add_instruction(migraphx::make_op("identity"), l0); + auto prog = optimize_tf("variable_batch_test.pb", false); + + EXPECT(p == prog); +} diff --git a/test/tf/tf_test.cpp b/test/tf/tf_test.cpp deleted file mode 100644 index 70e58b9400c..00000000000 --- a/test/tf/tf_test.cpp +++ /dev/null @@ -1,1080 +0,0 @@ -/* - * The MIT License (MIT) - * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to deal - * in the Software without restriction, including without limitation the rights - * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - * copies of the Software, and to permit persons to whom the Software is - * furnished to do so, subject to the following conditions: - * - * The above copyright notice and this permission notice shall be included in - * all copies or substantial portions of the Software. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER - * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - * THE SOFTWARE. - */ -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include - -#include - -#include "test.hpp" - -migraphx::program read_pb_file(const std::string& name, const migraphx::tf_options& options) -{ - static auto pb_files{::pb_files()}; - if(pb_files.find(name) == pb_files.end()) - { - std::cerr << "Can not find TensorFlow Protobuf file by name: " << name - << " , aborting the program\n" - << std::endl; - std::abort(); - } - return migraphx::parse_tf_buffer(std::string{pb_files.at(name)}, options); -} - -migraphx::program -parse_tf(const std::string& name, - bool is_nhwc, - const std::unordered_map>& dim_params = {}, - const std::vector& output_node_names = {}) -{ - - return read_pb_file(name, migraphx::tf_options{is_nhwc, 1, dim_params, output_node_names}); -} - -migraphx::program optimize_tf(const std::string& name, bool is_nhwc) -{ - auto prog = read_pb_file(name, migraphx::tf_options{is_nhwc, 1}); - auto* mm = prog.get_main_module(); - if(is_nhwc) - migraphx::run_passes(*mm, - {migraphx::simplify_reshapes{}, - migraphx::dead_code_elimination{}, - migraphx::eliminate_identity{}}); - - // remove the last return instruction - - if(mm->size() > 0) - { - auto last_ins = std::prev(mm->end()); - if(last_ins->name() == "@return") - { - mm->remove_instruction(last_ins); - } - } - return prog; -} - -TEST_CASE(add_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); - auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); - mm->add_instruction(migraphx::make_op("add"), l0, l1); - auto prog = optimize_tf("add_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(addv2_test) -{ - migraphx::program p; - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); - auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); - mm->add_instruction(migraphx::make_op("add"), l0, l1); - auto prog = optimize_tf("addv2_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(add_bcast_test) -{ - - migraphx::program p; - - auto* mm = p.get_main_module(); - migraphx::shape s0{migraphx::shape::float_type, {2, 3}}; - auto l0 = mm->add_parameter("0", s0); - auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {2, 1}}); - auto l2 = - mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", s0.lens()}}), l1); - mm->add_instruction(migraphx::make_op("add"), l0, l2); - auto prog = optimize_tf("add_bcast_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(argmax_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {4, 5, 6, 7}}); - mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::int32_type}, {2}}); - auto ins = mm->add_instruction(migraphx::make_op("argmax", {{"axis", 2}}), l0); - auto l1 = mm->add_instruction(migraphx::make_op("squeeze", {{"axes", {2}}}), ins); - mm->add_return({l1}); - auto prog = parse_tf("argmax_test.pb", false, {{"0", {4, 5, 6, 7}}}); - - EXPECT(p == prog); -} - -TEST_CASE(argmin_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {3, 4, 5, 6}}); - mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::int32_type}, {2}}); - auto ins = mm->add_instruction(migraphx::make_op("argmin", {{"axis", 2}}), l0); - auto l1 = mm->add_instruction(migraphx::make_op("squeeze", {{"axes", {2}}}), ins); - mm->add_return({l1}); - auto prog = parse_tf("argmin_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(assert_less_equal_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - migraphx::shape s0{migraphx::shape::float_type, {2, 3}}; - auto l0 = mm->add_parameter("0", s0); - auto l1 = mm->add_parameter("1", s0); - migraphx::literal l{migraphx::shape{migraphx::shape::int32_type, {2}}, {0, 1}}; - auto l2 = mm->add_literal(l); - mm->add_instruction(migraphx::make_op("add"), l0, l1); - auto l3 = mm->add_instruction(migraphx::make_op("identity"), l0, l1); - mm->add_instruction(migraphx::make_op("identity"), l3, l2); - auto prog = optimize_tf("assert_less_equal_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(batchmatmul_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 2, 8, 4}}); - auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {1, 2, 4, 8}}); - - auto trans_l0 = - mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 1, 3, 2}}}), l0); - auto trans_l1 = - mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 1, 3, 2}}}), l1); - - mm->add_instruction(migraphx::make_op("dot"), trans_l0, trans_l1); - auto prog = optimize_tf("batchmatmul_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(batchnorm_test) -{ - migraphx::program p; - auto* mm = p.get_main_module(); - - auto x = mm->add_parameter("x", {migraphx::shape::float_type, {1, 32, 16, 16}}); - auto bias = mm->add_parameter("bias", {migraphx::shape::float_type, {32}}); - auto mean = mm->add_parameter("mean", {migraphx::shape::float_type, {32}}); - auto var = mm->add_parameter("variance", {migraphx::shape::float_type, {32}}); - - std::vector scale_data(32, 1.0); - auto scale = mm->add_literal(migraphx::shape{migraphx::shape::float_type, {32}}, scale_data); - auto eps = mm->add_literal(migraphx::literal{migraphx::shape::float_type, {1e-4f}}); - - auto usq_scale = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), scale); - auto usq_bias = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), bias); - auto usq_mean = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), mean); - auto usq_var = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), var); - - auto x_sub_mean = add_common_op(*mm, migraphx::make_op("sub"), {x, usq_mean}); - auto var_eps = add_common_op(*mm, migraphx::make_op("add"), {usq_var, eps}); - auto rsqrt = mm->add_instruction(migraphx::make_op("rsqrt"), var_eps); - auto mul0 = add_common_op(*mm, migraphx::make_op("mul"), {usq_scale, rsqrt}); - auto r0 = add_common_op(*mm, migraphx::make_op("mul"), {x_sub_mean, mul0}); - add_common_op(*mm, migraphx::make_op("add"), {r0, usq_bias}); - - auto prog = optimize_tf("batchnorm_test.pb", true); - EXPECT(p == prog); -} - -TEST_CASE(batchnorm_half_test) -{ - migraphx::program p; - auto* mm = p.get_main_module(); - - auto x = mm->add_parameter("x", {migraphx::shape::half_type, {1, 32, 16, 16}}); - auto bias = mm->add_parameter("bias", {migraphx::shape::float_type, {32}}); - auto mean = mm->add_parameter("mean", {migraphx::shape::float_type, {32}}); - auto var = mm->add_parameter("variance", {migraphx::shape::float_type, {32}}); - - std::vector scale_data(32, 1.0); - auto scale = mm->add_literal(migraphx::shape{migraphx::shape::float_type, {32}}, scale_data); - auto eps = mm->add_literal(migraphx::literal{migraphx::shape::half_type, {1e-4f}}); - - auto usq_scale = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), scale); - auto usq_bias = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), bias); - auto usq_mean = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), mean); - auto usq_var = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), var); - - auto x_sub_mean = add_common_op(*mm, migraphx::make_op("sub"), {x, usq_mean}); - auto var_eps = add_common_op(*mm, migraphx::make_op("add"), {usq_var, eps}); - auto rsqrt = mm->add_instruction(migraphx::make_op("rsqrt"), var_eps); - auto mul0 = add_common_op(*mm, migraphx::make_op("mul"), {usq_scale, rsqrt}); - auto r0 = add_common_op(*mm, migraphx::make_op("mul"), {x_sub_mean, mul0}); - add_common_op(*mm, migraphx::make_op("add"), {r0, usq_bias}); - - auto prog = optimize_tf("batchnorm_half_test.pb", true); - EXPECT(p == prog); -} - -TEST_CASE(batchnormv3_test) -{ - migraphx::program p; - auto* mm = p.get_main_module(); - - auto x = mm->add_parameter("x", {migraphx::shape::float_type, {1, 32, 16, 16}}); - auto bias = mm->add_parameter("bias", {migraphx::shape::float_type, {32}}); - auto mean = mm->add_parameter("mean", {migraphx::shape::float_type, {32}}); - auto var = mm->add_parameter("variance", {migraphx::shape::float_type, {32}}); - - std::vector scale_data(32, 1.0); - auto scale = mm->add_literal(migraphx::shape{migraphx::shape::float_type, {32}}, scale_data); - auto eps = mm->add_literal(migraphx::literal{migraphx::shape::float_type, {1e-6f}}); - - auto usq_scale = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), scale); - auto usq_bias = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), bias); - auto usq_mean = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), mean); - auto usq_var = mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {1, 2}}}), var); - - auto x_sub_mean = add_common_op(*mm, migraphx::make_op("sub"), {x, usq_mean}); - auto var_eps = add_common_op(*mm, migraphx::make_op("add"), {usq_var, eps}); - auto rsqrt = mm->add_instruction(migraphx::make_op("rsqrt"), var_eps); - auto mul0 = add_common_op(*mm, migraphx::make_op("mul"), {usq_scale, rsqrt}); - auto r0 = add_common_op(*mm, migraphx::make_op("mul"), {x_sub_mean, mul0}); - add_common_op(*mm, migraphx::make_op("add"), {r0, usq_bias}); - - auto prog = optimize_tf("batchnormv3_test.pb", true); - EXPECT(p == prog); -} - -TEST_CASE(biasadd_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - migraphx::shape s0{migraphx::shape::float_type, {1, 500, 1, 1}}; - uint64_t axis = 1; - auto l0 = mm->add_parameter("0", s0); - auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {500}}); - auto l2 = mm->add_instruction( - migraphx::make_op("broadcast", {{"axis", axis}, {"out_lens", l0->get_shape().lens()}}), l1); - mm->add_instruction(migraphx::make_op("add"), l0, l2); - auto prog = optimize_tf("biasadd_test.pb", true); - - EXPECT(p == prog); -} - -TEST_CASE(biasadd_scalar_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - migraphx::shape s0{migraphx::shape::float_type, {1, 1}}; - uint64_t axis = 1; - auto l0 = mm->add_parameter("0", s0); - auto l1 = mm->add_literal( - migraphx::literal{migraphx::shape{migraphx::shape::float_type, {1}, {0}}, {1.0}}); - auto l2 = mm->add_instruction( - migraphx::make_op("broadcast", {{"axis", axis}, {"out_lens", l0->get_shape().lens()}}), l1); - mm->add_instruction(migraphx::make_op("add"), l0, l2); - auto prog = optimize_tf("biasadd_scalar_test.pb", true); - - EXPECT(p == prog); -} - -TEST_CASE(cast_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); - mm->add_instruction( - migraphx::make_op("convert", - {{"target_type", migraphx::to_value(migraphx::shape::int32_type)}}), - l0); - auto prog = optimize_tf("cast_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(concat_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {4, 7, 3}}); - auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {4, 2, 3}}); - - int axis = 1; - // tf uses axis as the third input, and it is in int32 format - // add the literal using a vector in order to set stride to 1 (like in tf parser) - mm->add_literal(migraphx::shape{migraphx::shape::int32_type}, std::vector{axis}); - - mm->add_instruction(migraphx::make_op("concat", {{"axis", axis}}), l0, l1); - auto prog = optimize_tf("concat_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(const_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - mm->add_literal(migraphx::shape{migraphx::shape::float_type}, std::vector{1.0f}); - auto prog = optimize_tf("constant_test.pb", false); - - EXPECT(p == prog); -} - -migraphx::program create_conv() -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); - std::vector weight_data(3 * 3 * 3 * 32); - std::fill(weight_data.begin(), weight_data.end(), 1.0f); - auto l1 = - mm->add_literal(migraphx::shape{migraphx::shape::float_type, {3, 3, 3, 32}}, weight_data); - - migraphx::op::convolution op; - op.padding = {1, 1, 1, 1}; - op.stride = {1, 1}; - op.dilation = {1, 1}; - auto l2 = - mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {3, 2, 0, 1}}}), l1); - mm->add_instruction(op, l0, l2); - return p; -} - -TEST_CASE(conv_test) -{ - migraphx::program p = create_conv(); - auto prog = optimize_tf("conv_test.pb", true); - - EXPECT(p == prog); -} - -TEST_CASE(conv_add_test) -{ - migraphx::program p = create_conv(); - auto* mm = p.get_main_module(); - auto l0 = std::prev(mm->end()); - mm->add_instruction(migraphx::make_op("add"), l0, l0); - auto prog = optimize_tf("conv_add_test.pb", true); - - EXPECT(p == prog); -} - -TEST_CASE(conv_nchw_test) -{ - migraphx::program p = create_conv(); - auto prog = optimize_tf("conv_nchw_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(conv_relu_test) -{ - migraphx::program p = create_conv(); - auto* mm = p.get_main_module(); - auto l0 = std::prev(mm->end()); - mm->add_instruction(migraphx::make_op("relu"), l0); - auto prog = optimize_tf("conv_relu_test.pb", true); - - EXPECT(p == prog); -} - -TEST_CASE(conv_relu6_test) -{ - migraphx::program p = create_conv(); - auto* mm = p.get_main_module(); - std::vector input_lens{1, 32, 16, 16}; - auto l0 = std::prev(mm->end()); - auto min_val = mm->add_literal(0.0f); - auto max_val = mm->add_literal(6.0f); - min_val = mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", input_lens}}), - min_val); - max_val = mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", input_lens}}), - max_val); - mm->add_instruction(migraphx::make_op("clip"), l0, min_val, max_val); - auto prog = optimize_tf("conv_relu6_test.pb", true); - - EXPECT(p == prog); -} - -TEST_CASE(depthwiseconv_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); - std::vector weight_data(3 * 3 * 3 * 1); - std::fill(weight_data.begin(), weight_data.end(), 1.0f); - auto l1 = - mm->add_literal(migraphx::shape{migraphx::shape::float_type, {3, 3, 3, 1}}, weight_data); - - migraphx::op::convolution op; - op.padding = {1, 1}; - op.stride = {1, 1}; - op.dilation = {1, 1}; - op.group = 3; - auto l3 = - mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {3, 2, 0, 1}}}), l1); - auto l4 = mm->add_instruction(migraphx::make_op("contiguous"), l3); - auto l5 = mm->add_instruction(migraphx::make_op("reshape", {{"dims", {3, 1, 3, 3}}}), l4); - mm->add_instruction(op, l0, l5); - auto prog = optimize_tf("depthwise_conv_test.pb", true); - - EXPECT(p == prog); -} - -TEST_CASE(expanddims_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {2, 3, 4}}); - mm->add_literal(0); - mm->add_instruction(migraphx::make_op("reshape", {{"dims", {1, 2, 3, 4}}}), l0); - auto prog = optimize_tf("expanddims_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(expanddims_test_neg_dims) -{ - // this check makes sure the pb parses negative dim value correctly - migraphx::program p; - - auto* mm = p.get_main_module(); - - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {2, 3, 4}}); - mm->add_literal(-1); - mm->add_instruction(migraphx::make_op("reshape", {{"dims", {2, 3, 4, 1}}}), l0); - auto prog = optimize_tf("expanddims_neg_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(gather_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {2, 4}}); - auto l1 = mm->add_literal( - migraphx::literal{migraphx::shape{migraphx::shape::int32_type, {2}}, {1, 1}}); - mm->add_literal(1); - - int axis = 1; - mm->add_instruction(migraphx::make_op("gather", {{"axis", axis}}), l0, l1); - auto prog = optimize_tf("gather_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(identity_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); - mm->add_instruction(migraphx::make_op("identity"), l0); - auto prog = optimize_tf("identity_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(matmul_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {8, 4}}); - auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {4, 8}}); - - auto trans_l0 = - mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {1, 0}}}), l0); - auto trans_l1 = - mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {1, 0}}}), l1); - - mm->add_instruction(migraphx::make_op("dot"), trans_l0, trans_l1); - auto prog = optimize_tf("matmul_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(mean_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - migraphx::literal l{migraphx::shape{migraphx::shape::int32_type, {2}}, {2, 3}}; - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); - mm->add_literal(l); - mm->add_literal(l); - migraphx::op::reduce_mean op{{2, 3}}; - mm->add_instruction(op, l0); - auto l3 = mm->add_instruction(op, l0); - mm->add_instruction(migraphx::make_op("squeeze", {{"axes", {2, 3}}}), l3); - auto prog = optimize_tf("mean_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(mean_test_nhwc) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - migraphx::literal l{migraphx::shape{migraphx::shape::int32_type, {2}}, {1, 2}}; - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); - auto l1 = - mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 2, 3, 1}}}), l0); - migraphx::op::reduce_mean op{{1, 2}}; - auto l2 = mm->add_instruction(op, l1); - mm->add_instruction(migraphx::make_op("squeeze", {{"axes", {1, 2}}}), l2); - auto prog = optimize_tf("mean_test_nhwc.pb", true); - - EXPECT(p == prog); -} - -TEST_CASE(mul_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 1, 1, 16}}); - auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {1, 1, 1, 16}}); - - mm->add_instruction(migraphx::make_op("mul"), l0, l1); - auto prog = optimize_tf("mul_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(multi_output_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); - auto l1 = mm->add_instruction(migraphx::make_op("relu"), l0); - auto l2 = mm->add_instruction(migraphx::make_op("tanh"), l0); - mm->add_return({l1, l2}); - - EXPECT(test::throws([&] { parse_tf("multi_output_test.pb", false, {}, {"relu", "relu6"}); })); - auto prog = parse_tf("multi_output_test.pb", false, {}, {"relu", "tanh"}); - - EXPECT(p == prog); -} - -TEST_CASE(onehot_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_literal( - migraphx::literal{migraphx::shape{migraphx::shape::int32_type, {5}}, {1, 1, 1, 1, 1}}); - mm->add_literal(2); - mm->add_literal(1.0f); - mm->add_literal(0.0f); - auto l1 = mm->add_literal( - migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2, 2}}, {1, 0, 0, 1}}); - int axis = 0; - mm->add_instruction(migraphx::make_op("gather", {{"axis", axis}}), l1, l0); - auto prog = optimize_tf("onehot_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(noop_test) -{ - migraphx::program p; - auto prog = optimize_tf("noop_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(pack_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {2}}); - auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {2}}); - auto l2 = mm->add_parameter("2", migraphx::shape{migraphx::shape::float_type, {2}}); - std::vector args{l0, l1, l2}; - std::vector unsqueezed_args; - int64_t axis = 1; - - std::transform( - args.begin(), - args.end(), - std::back_inserter(unsqueezed_args), - [&](migraphx::instruction_ref arg) { - return mm->add_instruction(migraphx::make_op("unsqueeze", {{"axes", {axis}}}), arg); - }); - mm->add_instruction(migraphx::make_op("concat", {{"axis", static_cast(axis)}}), - unsqueezed_args); - auto prog = optimize_tf("pack_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(pack_test_nhwc) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 2, 1, 1}}); - auto lt0 = - mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 2, 3, 1}}}), l0); - auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {1, 2, 1, 1}}); - auto lt1 = - mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 2, 3, 1}}}), l1); - auto l2 = mm->add_parameter("2", migraphx::shape{migraphx::shape::float_type, {1, 2, 1, 1}}); - auto lt2 = - mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 2, 3, 1}}}), l2); - std::vector args{lt0, lt1, lt2}; - std::vector unsqueezed_args; - int64_t nchw_axis = 3; - - std::transform(args.begin(), - args.end(), - std::back_inserter(unsqueezed_args), - [&](migraphx::instruction_ref arg) { - return mm->add_instruction( - migraphx::make_op("unsqueeze", {{"axes", {nchw_axis}}}), arg); - }); - mm->add_instruction(migraphx::make_op("concat", {{"axis", static_cast(nchw_axis)}}), - unsqueezed_args); - auto prog = optimize_tf("pack_test_nhwc.pb", true); - - EXPECT(p == prog); -} - -TEST_CASE(pad_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {2, 4}}); - std::vector pad_literals{1, 1, 2, 2}; - std::vector pads{1, 2, 1, 2}; - mm->add_literal(migraphx::shape{migraphx::shape::int32_type, {2, 2}}, pad_literals); - - mm->add_instruction(migraphx::make_op("pad", {{"pads", pads}}), l0); - auto prog = optimize_tf("pad_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(pooling_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); - migraphx::op::pooling avg_pool_op{migraphx::op::pooling_mode::average}; - migraphx::op::pooling max_pool_op{migraphx::op::pooling_mode::max}; - avg_pool_op.stride = {2, 2}; - max_pool_op.stride = {2, 2}; - avg_pool_op.lengths = {2, 2}; - max_pool_op.lengths = {2, 2}; - mm->add_instruction(avg_pool_op, l0); - mm->add_instruction(max_pool_op, l0); - auto prog = optimize_tf("pooling_test.pb", true); - - EXPECT(p == prog); -} - -TEST_CASE(pow_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); - auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); - mm->add_instruction(migraphx::make_op("pow"), l0, l1); - auto prog = optimize_tf("pow_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(relu_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); - mm->add_instruction(migraphx::make_op("relu"), l0); - auto prog = optimize_tf("relu_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(relu6_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - std::vector input_lens{1, 3, 16, 16}; - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, input_lens}); - auto min_val = mm->add_literal(0.0f); - auto max_val = mm->add_literal(6.0f); - min_val = mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", input_lens}}), - min_val); - max_val = mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", input_lens}}), - max_val); - mm->add_instruction(migraphx::make_op("clip"), l0, min_val, max_val); - auto prog = optimize_tf("relu6_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(relu6_half_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - std::vector input_lens{1, 3, 16, 16}; - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::half_type, input_lens}); - auto min_val = - mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::half_type}, {0.0f}}); - auto max_val = - mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::half_type}, {6.0f}}); - min_val = mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", input_lens}}), - min_val); - max_val = mm->add_instruction(migraphx::make_op("multibroadcast", {{"out_lens", input_lens}}), - max_val); - mm->add_instruction(migraphx::make_op("clip"), l0, min_val, max_val); - auto prog = optimize_tf("relu6_half_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(reshape_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {16}}); - migraphx::shape s0{migraphx::shape::int32_type, {4}}; - // in tf, the second arg is a literal that contains new dimensions - mm->add_literal(migraphx::literal{s0, {1, 1, 1, 16}}); - mm->add_instruction(migraphx::make_op("reshape", {{"dims", {1, 1, 1, 16}}}), l0); - auto prog = optimize_tf("reshape_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(rsqrt_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); - mm->add_instruction(migraphx::make_op("rsqrt"), l0); - auto prog = optimize_tf("rsqrt_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(shape_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); - mm->add_literal( - migraphx::literal{migraphx::shape{migraphx::shape::int32_type, {4}}, {1, 3, 16, 16}}); - auto prog = optimize_tf("shape_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(slice_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - std::size_t num_axes = 2; - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {5, 10}}); - migraphx::shape s0{migraphx::shape::int32_type, {num_axes}}; - mm->add_literal(migraphx::literal{s0, {1, 0}}); - mm->add_literal(migraphx::literal{s0, {2, -1}}); - - mm->add_instruction( - migraphx::make_op("slice", {{"starts", {1, 0}}, {"ends", {3, 10}}, {"axes", {0, 1}}}), l0); - auto prog = optimize_tf("slice_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(softmax_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3}}); - mm->add_instruction(migraphx::make_op("softmax", {{"axis", 1}}), l0); - auto prog = optimize_tf("softmax_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(split_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - std::vector axes{0, 1}; - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {5, 30}}); - mm->add_literal(3); // num_splits - mm->add_literal(1); // split axis - mm->add_literal(1); // concat axis - mm->add_literal(1); // concat axis - auto l1 = mm->add_instruction( - migraphx::make_op("slice", {{"axes", axes}, {"starts", {0, 0}}, {"ends", {5, 10}}}), l0); - auto l2 = mm->add_instruction( - migraphx::make_op("slice", {{"axes", axes}, {"starts", {0, 10}}, {"ends", {5, 20}}}), l0); - auto l3 = mm->add_instruction( - migraphx::make_op("slice", {{"axes", axes}, {"starts", {0, 20}}, {"ends", {5, 30}}}), l0); - auto l4 = mm->add_instruction(migraphx::make_op("concat", {{"axis", 1}}), l1, l2); - auto l5 = mm->add_instruction(migraphx::make_op("concat", {{"axis", 1}}), l2, l3); - mm->add_return({l4, l5}); - auto prog = parse_tf("split_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(split_test_one_output) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {5, 30}}); - mm->add_literal(1); // num_splits - mm->add_literal(1); // split axis - auto l1 = mm->add_instruction(migraphx::make_op("identity"), l0); - mm->add_return({l1}); - auto prog = parse_tf("split_test_one_output.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(split_test_vector_as_input) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - std::vector axes{0, 1}; - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {5, 30}}); - // split sizes - mm->add_literal( - migraphx::literal{migraphx::shape{migraphx::shape::int32_type, {3}}, {4, 15, 11}}); - mm->add_literal(1); // split axis - mm->add_literal(1); // concat axis - mm->add_literal(1); // concat axis - auto l1 = mm->add_instruction( - migraphx::make_op("slice", {{"axes", axes}, {"starts", {0, 0}}, {"ends", {5, 4}}}), l0); - auto l2 = mm->add_instruction( - migraphx::make_op("slice", {{"axes", axes}, {"starts", {0, 4}}, {"ends", {5, 19}}}), l0); - auto l3 = mm->add_instruction( - migraphx::make_op("slice", {{"axes", axes}, {"starts", {0, 19}}, {"ends", {5, 30}}}), l0); - auto l4 = mm->add_instruction(migraphx::make_op("concat", {{"axis", 1}}), l1, l2); - auto l5 = mm->add_instruction(migraphx::make_op("concat", {{"axis", 1}}), l2, l3); - mm->add_return({l4, l5}); - auto prog = parse_tf("split_test_vector_as_input.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(sqdiff_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); - auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); - mm->add_instruction(migraphx::make_op("sqdiff"), l0, l1); - auto prog = optimize_tf("sqdiff_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(squeeze_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 2, 3, 1}}); - mm->add_instruction(migraphx::make_op("squeeze", {{"axes", {0, 3}}}), l0); - auto prog = optimize_tf("squeeze_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(stopgradient_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); - mm->add_instruction(migraphx::make_op("identity"), l0); - auto prog = optimize_tf("stopgradient_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(stridedslice_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 10, 1, 1}}); - auto l1 = - mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 2, 3, 1}}}), l0); - auto l2 = mm->add_instruction( - migraphx::make_op( - "slice", {{"starts", {0, 0, 0, 0}}, {"ends", {1, 1, 1, 5}}, {"axes", {0, 1, 2, 3}}}), - l1); - auto shrink_axis = 1; - mm->add_instruction(migraphx::make_op("squeeze", {{"axes", {shrink_axis}}}), l2); - auto prog = optimize_tf("stridedslice_test.pb", true); - - EXPECT(p == prog); -} - -TEST_CASE(stridedslice_masks_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 10, 3, 3}}); - // add literals for starts, ends, and strides in tf (NHWC format) - mm->add_literal(migraphx::shape{migraphx::shape::int32_type, {4}}, - std::vector{0, 1, 1, 0}); - mm->add_literal(migraphx::shape{migraphx::shape::int32_type, {4}}, - std::vector{0, 0, 0, 0}); - mm->add_literal(migraphx::shape{migraphx::shape::int32_type, {4}}, - std::vector{1, 1, 1, 1}); - - auto l1 = - mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 2, 3, 1}}}), l0); - auto l2 = mm->add_instruction( - migraphx::make_op( - "slice", {{"starts", {0, 1, 1, 0}}, {"ends", {1, 3, 3, 10}}, {"axes", {0, 1, 2, 3}}}), - l1); - auto l3 = - mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 3, 1, 2}}}), l2); - mm->add_return({l3}); - auto prog = parse_tf("stridedslice_masks_test.pb", true); - - EXPECT(p == prog); -} - -TEST_CASE(sub_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); - auto l1 = mm->add_parameter("1", migraphx::shape{migraphx::shape::float_type, {1, 2, 2, 3}}); - auto l2 = mm->add_instruction(migraphx::make_op("sub"), l0, l1); - mm->add_return({l2}); - auto prog = parse_tf("sub_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(tanh_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); - auto l1 = mm->add_instruction(migraphx::make_op("tanh"), l0); - mm->add_return({l1}); - auto prog = parse_tf("tanh_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(transpose_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); - migraphx::shape s0{migraphx::shape::int32_type, {4}}; - mm->add_literal(migraphx::literal{s0, {0, 2, 3, 1}}); - mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 2, 3, 1}}}), l0); - auto prog = optimize_tf("transpose_test.pb", false); - - EXPECT(p == prog); -} - -TEST_CASE(variable_batch_test) -{ - migraphx::program p; - - auto* mm = p.get_main_module(); - auto l0 = mm->add_parameter("0", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}}); - mm->add_instruction(migraphx::make_op("identity"), l0); - auto prog = optimize_tf("variable_batch_test.pb", false); - - EXPECT(p == prog); -} - -int main(int argc, const char* argv[]) { test::run(argc, argv); } From 1f8619e29995a8268d45eb149abcfae8e37ddfa9 Mon Sep 17 00:00:00 2001 From: Slobodan Josic <127323561+slojosic-amd@users.noreply.github.com> Date: Wed, 29 May 2024 00:55:05 +0200 Subject: [PATCH 12/13] Fix hipExtModuleLaunchKernel function (#3111) --- src/targets/gpu/kernel.cpp | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/src/targets/gpu/kernel.cpp b/src/targets/gpu/kernel.cpp index 1cbb45852b1..033adbd07a2 100644 --- a/src/targets/gpu/kernel.cpp +++ b/src/targets/gpu/kernel.cpp @@ -27,6 +27,9 @@ #include #include +#ifdef _WIN32 +#include +#else // extern declare the function since hip/hip_ext.h header is broken extern hipError_t hipExtModuleLaunchKernel(hipFunction_t, // NOLINT uint32_t, @@ -42,6 +45,7 @@ extern hipError_t hipExtModuleLaunchKernel(hipFunction_t, // NOLINT hipEvent_t = nullptr, hipEvent_t = nullptr, uint32_t = 0); +#endif namespace migraphx { inline namespace MIGRAPHX_INLINE_NS { From dc028dd96460c79400f24784354d1cde88643709 Mon Sep 17 00:00:00 2001 From: Joseph Macaranas <145489236+amd-jmacaran@users.noreply.github.com> Date: Tue, 28 May 2024 21:02:40 -0400 Subject: [PATCH 13/13] Enable external CI pipeline triggers (#3118) --- .azuredevops/rocm-ci.yml | 43 ++++++++++++++++++++++++++++++++++++++++ 1 file changed, 43 insertions(+) create mode 100644 .azuredevops/rocm-ci.yml diff --git a/.azuredevops/rocm-ci.yml b/.azuredevops/rocm-ci.yml new file mode 100644 index 00000000000..d923c27734b --- /dev/null +++ b/.azuredevops/rocm-ci.yml @@ -0,0 +1,43 @@ +resources: + repositories: + - repository: pipelines_repo + type: github + endpoint: ROCm + name: ROCm/ROCm + +variables: +- group: common +- template: /.azuredevops/variables-global.yml@pipelines_repo + +trigger: + batch: true + branches: + include: + - develop + paths: + exclude: + - .githooks + - .github + - docs + - '.*.y*ml' + - '*.md' + - Jenkinsfile + - LICENSE + +pr: + autoCancel: true + branches: + include: + - develop + paths: + exclude: + - .github + - docs + - '.*.y*ml' + - '*.md' + - Jenkinsfile + - LICENSE + drafts: false + +jobs: + - template: ${{ variables.CI_COMPONENT_PATH }}/AMDMIGraphX.yml@pipelines_repo