Merge branch 'develop' into grid_sample

.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

.github/workflows/weekly_master_sync.yaml

@@ -17,9 +17,7 @@ jobs:
           fetch-depth: '0'
 
       - name: Merge Fast Forward Only
-        uses: IgorMirosavljevicHTEC/[email protected]
-        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

docs/sphinx/requirements.in

@@ -1,2 +1,2 @@
-rocm-docs-core==1.1.2
+rocm-docs-core==1.2.0
 sphinx-collapse

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

requirements.txt

@@ -28,4 +28,4 @@ pybind/pybind11@3e9dfa2866941655c56877882565e7577de6fc7b --build
 msgpack/[email protected] -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

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

src/targets/gpu/kernel.cpp

@@ -27,6 +27,9 @@
 #include <migraphx/gpu/pack_args.hpp>
 #include <cassert>
 
+#ifdef _WIN32
+#include <hip/hip_ext.h>
+#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 {

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<size_t> 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);
 }

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

test/onnx/.onnxrt-commit

@@ -1 +1 @@
-156d52163dd487ca0a1afaa87a12a3e172a45200
+33a68d221f28bd8d412f2e9188e50bac8a255b71

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})

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)
 

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<float> 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

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 <iostream>
+#include <vector>
+#include <unordered_map>
+#include <pb_files.hpp>
+#include <migraphx/common.hpp>
+#include <migraphx/literal.hpp>
+#include <migraphx/pass_manager.hpp>
+#include <migraphx/simplify_reshapes.hpp>
+#include <migraphx/dead_code_elimination.hpp>
+#include <migraphx/eliminate_identity.hpp>
+#include <migraphx/program.hpp>
+#include <migraphx/instruction.hpp>
+#include <migraphx/tf.hpp>
+#include <migraphx/make_op.hpp>
+#include <migraphx/op/convolution.hpp>
+#include <migraphx/op/reduce_mean.hpp>
+#include <migraphx/op/pooling.hpp>
+
+#include <migraphx/serialize.hpp>
+
+#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<std::string, std::vector<std::size_t>>& dim_params = {},
+         const std::vector<std::string>& 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

test/tf/add_test.pb → test/tf/models/add_test.pb

@@ -7,6 +7,6 @@
 
1Placeholder*
 dtype0
*
 shape:

-
-add1Add
0
1*
-
T0
"�
+
+add1AddV2
0
1*
+
T0
"�