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 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 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 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 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/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 { 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/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}); 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 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); } 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;