Merge

src/onnx/parse_gemm.cpp

@@ -39,10 +39,19 @@ struct parse_gemm : op_parser<parse_gemm>
                           onnx_parser::node_info info,
                           std::vector<instruction_ref> args) const
     {
-        float alpha = 1.0f;
-        float beta  = 1.0f;
-        bool transa = false;
-        bool transb = false;
+        auto a_arg = args[0];
+        auto b_arg = args[1];
+        if(a_arg->get_shape().ndim() != 2 or b_arg->get_shape().ndim() != 2)
+        {
+            MIGRAPHX_THROW("PARSE_GEMM: A and B should be rank 2, A is rank " +
+                           std::to_string(a_arg->get_shape().ndim()) + ", B is rank " +
+                           std::to_string(b_arg->get_shape().ndim()));
+        }
+
+        float alpha  = 1.0f;
+        float beta   = 1.0f;
+        bool trans_a = false;
+        bool trans_b = false;
         if(contains(info.attributes, "alpha"))
         {
             alpha = parser.parse_value(info.attributes.at("alpha")).at<float>();
@@ -53,65 +62,73 @@ struct parse_gemm : op_parser<parse_gemm>
         }
         if(contains(info.attributes, "transA"))
         {
-            transa = parser.parse_value(info.attributes.at("transA")).at<bool>();
+            trans_a = parser.parse_value(info.attributes.at("transA")).at<bool>();
         }
         if(contains(info.attributes, "transB"))
         {
-            transb = parser.parse_value(info.attributes.at("transB")).at<bool>();
+            trans_b = parser.parse_value(info.attributes.at("transB")).at<bool>();
         }
 
-        std::vector<int64_t> perm(args[0]->get_shape().lens().size());
-        std::iota(perm.begin(), perm.end(), int64_t{0});
-        // swap the last two elements
-        std::swap(*perm.rbegin(), *(perm.rbegin() + 1));
-
-        auto l1       = args[0];
-        auto dot_type = l1->get_shape().type();
-
+        std::vector<int64_t> perm = {1, 0};
+        auto dot_type             = a_arg->get_shape().type();
         if(alpha != 1.0f)
         {
             auto alpha_literal = info.add_literal(alpha);
-            l1                 = info.add_broadcastable_binary_op("mul", alpha_literal, l1);
-            if(l1->get_shape().type() != dot_type)
+            a_arg              = info.add_broadcastable_binary_op("mul", alpha_literal, a_arg);
+
+            if(a_arg->get_shape().type() != dot_type)
             {
-                l1 = info.add_instruction(make_op("convert", {{"target_type", dot_type}}), l1);
+                a_arg =
+                    info.add_instruction(make_op("convert", {{"target_type", dot_type}}), a_arg);
             }
         }
 
-        l1 =
-            (transa) ? info.add_instruction(make_op("transpose", {{"permutation", perm}}), l1) : l1;
-        auto l2 = (transb)
-                      ? info.add_instruction(make_op("transpose", {{"permutation", perm}}), args[1])
-                      : args[1];
+        a_arg = (trans_a)
+                    ? info.add_instruction(make_op("transpose", {{"permutation", perm}}), a_arg)
+                    : a_arg;
+        b_arg = (trans_b)
+                    ? info.add_instruction(make_op("transpose", {{"permutation", perm}}), args[1])
+                    : args[1];
 
-        auto ret = info.add_instruction(make_op("dot"), l1, l2);
+        auto dot_ins = info.add_instruction(make_op("dot"), a_arg, b_arg);
 
         if(args.size() == 3)
         {
-            if(not float_equal(beta, 0.0f) && args[2]->get_shape().elements() > 0)
+            if(not float_equal(beta, 0.0f))
             {
-                auto out_lens   = l1->get_shape().lens();
-                out_lens.back() = l2->get_shape().lens().back();
-                auto l3         = args[2];
-                auto l3_lens    = l3->get_shape().lens();
-                if(not std::equal(out_lens.begin(), out_lens.end(), l3_lens.begin(), l3_lens.end()))
+                auto c_arg = args[2];
+                if(dot_ins->get_shape().dynamic())
                 {
-                    l3 = info.add_instruction(make_op("multibroadcast", {{"out_lens", out_lens}}),
-                                              args[2]);
+                    c_arg = info.add_instruction(make_op("multibroadcast"), args[2], dot_ins);
                 }
-                auto beta_literal = info.add_literal(beta);
-                auto beta_l3      = info.add_broadcastable_binary_op("mul", l3, beta_literal);
-                if(beta_l3->get_shape().type() != dot_type)
+                else
                 {
-                    beta_l3 = info.add_instruction(make_op("convert", {{"target_type", dot_type}}),
-                                                   beta_l3);
+                    auto out_lens   = a_arg->get_shape().lens();
+                    out_lens.back() = b_arg->get_shape().lens().back();
+                    auto c_lens     = c_arg->get_shape().lens();
+                    if(not std::equal(
+                           out_lens.begin(), out_lens.end(), c_lens.begin(), c_lens.end()))
+                    {
+                        c_arg = info.add_instruction(
+                            make_op("multibroadcast", {{"out_lens", out_lens}}), args[2]);
+                    }
                 }
 
-                return info.add_instruction(make_op("add"), ret, beta_l3);
+                if(not float_equal(beta, 1.0f))
+                {
+                    auto beta_literal = info.add_literal(beta);
+                    c_arg = info.add_broadcastable_binary_op("mul", c_arg, beta_literal);
+                    if(c_arg->get_shape().type() != dot_type)
+                    {
+                        c_arg = info.add_instruction(
+                            make_op("convert", {{"target_type", dot_type}}), c_arg);
+                    }
+                }
+
+                return info.add_instruction(make_op("add"), dot_ins, c_arg);
             }
         }
-
-        return ret;
+        return dot_ins;
     }
 };
 

src/onnx/parse_trilu.cpp

+/*
+ * The MIT License (MIT)
+ *
+ * Copyright (c) 2015-2022 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 <migraphx/onnx/op_parser.hpp>
+#include <migraphx/onnx/checks.hpp>
+#include <migraphx/ranges.hpp>
+#include <migraphx/instruction.hpp>
+#include <migraphx/make_op.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace onnx {
+
+struct parse_trilu : op_parser<parse_trilu>
+{
+    std::vector<op_desc> operators() const { return {{"Trilu"}}; }
+
+    instruction_ref parse(const op_desc&,
+                          const onnx_parser&,
+                          const onnx_parser::node_info& info,
+                          std::vector<instruction_ref> args) const
+    {
+        auto input_shape = args[0]->get_shape();
+        assert(input_shape.ndim() >= 2);
+        auto input_lens = input_shape.lens();
+
+        size_t num_rows = *(input_lens.rbegin() + 1);
+        size_t num_cols = input_lens.back();
+        int k           = 0;
+        bool upper      = true;
+
+        if(args.size() > 1)
+        {
+            auto arg_k = args[1]->eval();
+            check_arg_empty(arg_k, "PARSE_TRILU: dynamic k not supported");
+            k = arg_k.at<int>();
+        }
+
+        if(k < 0)
+            MIGRAPHX_THROW("PARSE_TRILU: negative k values not supported");
+
+        if(contains(info.attributes, "upper"))
+        {
+            upper = static_cast<bool>(info.attributes.at("upper").i());
+        }
+
+        shape::type_t output_type = args[0]->get_shape().type();
+
+        // when creating the mask, if upper == 1,
+        // the inner triangle will have values set to 0
+        std::vector<bool> mask_mat(num_rows * num_cols, upper);
+        for(size_t i = 0; i < num_rows; i++)
+        {
+            for(size_t j = 0; j < std::min(k, static_cast<int>(num_cols)); j++)
+            {
+                mask_mat[i * num_cols + j] = not upper;
+            }
+            k++;
+        }
+
+        auto mask = info.add_literal(
+            migraphx::literal{migraphx::shape{output_type, {num_rows, num_cols}}, mask_mat});
+
+        return info.add_broadcastable_binary_op("mul", mask, args[0]);
+    }
+};
+
+} // namespace onnx
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx

src/targets/gpu/jit/reduce.cpp

@@ -118,16 +118,14 @@ struct reduce_compiler : compiler<reduce_compiler>
         options.virtual_inputs = reduce_dims(inputs);
         auto faxis             = find_fast_axis({options.virtual_inputs.front()});
         vectorize vec{};
-        // Vectorize if the axis is a reduction axis
-        if(options.virtual_inputs.back().lens()[faxis] == 1)
-        {
-            vec = vectorize::elements(ctx, faxis, options.virtual_inputs);
-        }
-        auto relements = get_reduce_elements(options.virtual_inputs) / vec.size;
         auto nelements = options.virtual_inputs.back().elements();
         auto algo      = v.get("algo", get_reduce_algo(options.virtual_inputs));
         if(algo == "block")
         {
+            // Vectorize if the axis is a reduction axis
+            if(options.virtual_inputs.back().lens()[faxis] == 1)
+                vec = vectorize::elements(ctx, faxis, options.virtual_inputs);
+            auto relements  = get_reduce_elements(options.virtual_inputs) / vec.size;
             auto block_size = compute_block_size(relements, 256);
             options.set_launch_params(
                 v, compute_global_for(ctx, nelements * block_size, 256), block_size);
@@ -166,7 +164,7 @@ struct reduce_compiler : compiler<reduce_compiler>
             auto reduce_elements = get_reduce_elements(ins->inputs());
             auto reduce_type     = ins->inputs().front()->get_shape().type();
             v["reduction"]       = "op::sum{}";
-            std::string mean     = "op::mean{" + std::to_string(reduce_elements) + "}";
+            std::string mean     = "op::mean<" + std::to_string(reduce_elements) + ">{}";
             // Use float accumulator when reduction size is too large for half
             if(reduce_type == shape::half_type and reduce_elements > 16384)
                 v["read"] = "compose(" + mean + ", op::convert_to<float>{})";

src/targets/gpu/kernels/include/migraphx/kernels/debug.hpp

@@ -178,5 +178,9 @@ MIGRAPHX_HIP_NORETURN inline __host__ __device__ void assert_fail(const source_l
 #define MIGRAPHX_WARN(...)
 #endif
 
+#define MIGRAPHX_STATIC_ASSERT_FOR(...) \
+    static_assert(__VA_ARGS__);         \
+    if constexpr(__VA_ARGS__)
+
 } // namespace migraphx
 #endif // MIGRAPHX_GUARD_KERNELS_DEBUG_HPP

src/targets/gpu/kernels/include/migraphx/kernels/index.hpp

@@ -161,6 +161,30 @@ struct index
     }
 
     template <class F, class N, class Stride>
+    static constexpr void for_stride_loop_unroll(index_int start, N n, Stride stride, F f)
+    {
+        sequence(max_stride_iterations(n, stride), [&](auto... ks) {
+            fold([&](auto d, auto k) {
+                auto i = start + stride * k;
+                if(i < n)
+                    invoke_loop(f, i, d);
+                return d + _c<1>;
+            })(_c<0>, ks...);
+        });
+    }
+
+    template <class F, class N, class Stride>
+    static constexpr void for_stride_loop(index_int start, N n, Stride stride, F f)
+    {
+        index_int k = 0;
+        for(index_int i = start; i < n; i += stride)
+        {
+            invoke_loop(f, i, k);
+            k++;
+        }
+    }
+
+    template <bool Unroll, class F, class N, class Stride>
     static constexpr void for_stride(index_int start, N n, Stride stride, F f)
     {
         MIGRAPHX_ASSERT(start < stride);
@@ -178,40 +202,34 @@ struct index
                     invoke_loop(f, start, _c<0>);
                 }
             }
+            else if constexpr(Unroll)
+            {
+                MIGRAPHX_STATIC_ASSERT_FOR(max_stride_iterations(n, stride) < 256)
+                {
+                    for_stride_loop_unroll(start, n, stride, f);
+                }
+            }
             else
             {
-                // static_assert(max_stride_iterations(n, stride) < 64);
-                sequence(max_stride_iterations(n, stride), [&](auto... ks) {
-                    fold([&](auto d, auto k) {
-                        auto i = start + stride * k;
-                        if(i < n)
-                            invoke_loop(f, i, d);
-                        return d + _c<1>;
-                    })(_c<0>, ks...);
-                });
+                for_stride_loop(start, n, stride, f);
             }
         }
         else
         {
-            index_int k = 0;
-            for(index_int i = start; i < n; i += stride)
-            {
-                invoke_loop(f, i, k);
-                k++;
-            }
+            for_stride_loop(start, n, stride, f);
         }
     }
 
     template <class F, class N>
     __device__ void global_stride(N n, F f) const
     {
-        for_stride(global, n, nglobal(), f);
+        for_stride<false>(global, n, nglobal(), f);
     }
 
     template <class F, class N>
     __device__ void local_stride(N n, F f) const
     {
-        for_stride(local, n, nlocal(), f);
+        for_stride<true>(local, n, nlocal(), f);
     }
 };
 

src/targets/gpu/kernels/include/migraphx/kernels/ops.hpp

@@ -66,13 +66,22 @@ struct convert_to
     }
 };
 
+template <index_int N>
 struct mean
 {
-    index_int item_num = 1;
     template <class T>
-    MIGRAPHX_DEVICE_CONSTEXPR auto operator()(T x) const
+    MIGRAPHX_DEVICE_CONSTEXPR T operator()(T x) const
     {
-        return x / static_cast<T>(item_num);
+        using type = vec_type<T>;
+        if constexpr(is_floating_point<type>{})
+        {
+            constexpr type d = 1.0 / N;
+            return x * d;
+        }
+        else
+        {
+            return x / static_cast<type>(N);
+        }
     }
 };
 

src/targets/gpu/kernels/include/migraphx/kernels/reduce.hpp

@@ -391,22 +391,40 @@ struct block
 struct lane
 {
     template <class Slicer>
-    struct reducer
+    struct reducer : reducer_base<reducer<Slicer>>
     {
         index idx;
         Slicer slice;
-        template <class Op, class T, class Read>
-        __device__ auto reduce(Op op, T init, Read read) const
+
+        template <class Size, class F>
+        struct inner_storage : inner_storage_tag
         {
-            return sliced(slice, [=](auto x, auto... xs) {
-                using type = typename decltype(x)::type;
-                type r     = init;
-                for(index_int j = 0; j < x.get_shape().elements(); j++)
-                {
-                    r = op(r, read(x[j], xs[j]...));
-                }
-                return r;
-            });
+            using type = remove_reference_t<decltype(declval<F>()(0, _c<0>))>;
+            F f;
+            constexpr Size rsize() const { return {}; }
+            template <class U, class V>
+            constexpr auto operator()(U j, V d) const
+            {
+                return f(j, d);
+            }
+        };
+
+        template <class Size, class F>
+        constexpr inner_storage<Size, F> make_inner_storage(Size, F f)
+        {
+            return {f};
+        }
+
+        template <class Op, class T, class Read, class N, class U, class... Us>
+        __device__ auto reduce_impl(Op op, T init, Read read, N n, U&& x, Us&&... xs) const
+        {
+            using type = remove_reference_t<decltype(x(0, _c<0>))>;
+            type r     = init;
+            for(index_int j = 0; j < n; j++)
+            {
+                r = op(r, read(x(j, _c<0>), xs(j, _c<0>)...));
+            }
+            return r;
         }
 
         template <class F>
@@ -415,29 +433,25 @@ struct lane
             f();
         }
 
-        template <class F>
-        __device__ auto inner(F f) const
+        template <class F, class N, class... Ts>
+        __device__ void inner_void_impl(F f, N n, Ts&&... xs) const
         {
-            return sliced(slice, [=](auto x, auto... xs) {
-                for(index_int j = 0; j < x.get_shape().elements(); j++)
-                {
-                    f(x[j], xs[j]...);
-                }
-            });
+            for(index_int j = 0; j < n; j++)
+            {
+                f(xs(j, _c<0>)...);
+            }
         }
 
-        template <class Input>
-        constexpr auto elements() const
+        template <class R, class F, class N, class... Ts>
+        __device__ auto inner_impl(F f, N n, Ts&&... xs) const
         {
-            using reduce_type = decltype(slice(Input{}));
-            return get_shape_c<reduce_type>{}.elements();
+            return make_inner_storage(n, [=](auto j, auto d) { return f(xs(j, d)...); });
         }
     };
-
     template <class Slicer>
     static __device__ auto make(index idx, Slicer slicer)
     {
-        return reducer<Slicer>{idx, slicer};
+        return reducer<Slicer>{{}, idx, slicer};
     }
 
     template <class Output, class F>

src/targets/gpu/prefuse_ops.cpp

@@ -100,7 +100,8 @@ struct find_add_layernorm
 {
     auto matcher() const
     {
-        return match::layernorm()(match::var("x")(match::name("add").bind("add")));
+        return match::layernorm()(
+            match::var("x")(match::name("add")(match::used_once()).bind("add")));
     }
 
     void apply(module& m, const match::matcher_result& r) const

test/api/test_cpu.cpp

@@ -66,7 +66,7 @@ TEST_CASE(load_and_run_init_list)
 
 TEST_CASE(quantize_fp16)
 {
-    auto p1        = migraphx::parse_onnx("gemm_ex_test.onnx");
+    auto p1        = migraphx::parse_onnx("gemm_test.onnx");
     const auto& p2 = p1;
     const auto& p3 = p1;
     migraphx::quantize_fp16(p1);
@@ -82,7 +82,7 @@ TEST_CASE(quantize_fp16)
 
 TEST_CASE(quantize_int8)
 {
-    auto p1        = migraphx::parse_onnx("gemm_ex_test.onnx");
+    auto p1        = migraphx::parse_onnx("gemm_test.onnx");
     const auto& p2 = p1;
     auto t         = migraphx::target("ref");
     migraphx::quantize_int8_options options;

test/onnx/gen_onnx.py

@@ -2116,71 +2116,136 @@ def gathernd_batch_dims_test():
 
 @onnx_test()
 def gemm_test():
-    x = helper.make_tensor_value_info('0', TensorProto.FLOAT, [5, 7])
-    y = helper.make_tensor_value_info('1', TensorProto.FLOAT, [11, 5])
-    z = helper.make_tensor_value_info('2', TensorProto.FLOAT, [])
-    a = helper.make_tensor_value_info('3', TensorProto.FLOAT, [7, 11])
+    A = helper.make_tensor_value_info('A', TensorProto.FLOAT, [8, 6])
+    B = helper.make_tensor_value_info('B', TensorProto.FLOAT, [8, 7])
+    C = helper.make_tensor_value_info('C', TensorProto.FLOAT, [6, 7])
+    Y = helper.make_tensor_value_info('Y', TensorProto.FLOAT, [6, 7])
 
     node = onnx.helper.make_node('Gemm',
-                                 inputs=['0', '1', '2'],
-                                 outputs=['3'],
+                                 inputs=['A', 'B', 'C'],
+                                 outputs=['Y'],
+                                 alpha=0.5,
+                                 beta=0.8,
+                                 transA=1)
+
+    return ([node], [A, B, C], [Y])
+
+
+@onnx_test()
+def gemm_no_C_test():
+    A = helper.make_tensor_value_info('A', TensorProto.FLOAT, [5, 7])
+    B = helper.make_tensor_value_info('B', TensorProto.FLOAT, [11, 5])
+    C = helper.make_tensor_value_info('C', TensorProto.FLOAT, [])
+    Y = helper.make_tensor_value_info('Y', TensorProto.FLOAT, [7, 11])
+
+    node = onnx.helper.make_node('Gemm',
+                                 inputs=['A', 'B', 'C'],
+                                 outputs=['Y'],
                                  alpha=2.0,
                                  beta=2.0,
                                  transA=1,
                                  transB=1)
 
-    return ([node], [x, y, z], [a])
+    return ([node], [A, B, C], [Y])
 
 
 @onnx_test()
-def gemm_ex_test():
-    m1 = helper.make_tensor_value_info('1', TensorProto.FLOAT, [1, 1, 8, 6])
-    m2 = helper.make_tensor_value_info('2', TensorProto.FLOAT, [1, 1, 8, 7])
-    m3 = helper.make_tensor_value_info('3', TensorProto.FLOAT, [1, 1, 6, 7])
-    y = helper.make_tensor_value_info('y', TensorProto.FLOAT, [1, 1, 6, 7])
+def gemm_brcst_C_test():
+    A = helper.make_tensor_value_info('A', TensorProto.FLOAT, [5, 6])
+    B = helper.make_tensor_value_info('B', TensorProto.FLOAT, [5, 7])
+    C = helper.make_tensor_value_info('C', TensorProto.FLOAT, [6, 1])
+    Y = helper.make_tensor_value_info('Y', TensorProto.FLOAT, [6, 7])
 
     node = onnx.helper.make_node('Gemm',
-                                 inputs=['1', '2', '3'],
-                                 outputs=['y'],
+                                 inputs=['A', 'B', 'C'],
+                                 outputs=['Y'],
                                  alpha=0.5,
                                  beta=0.8,
                                  transA=1)
 
-    return ([node], [m1, m2, m3], [y])
+    return ([node], [A, B, C], [Y])
 
 
 @onnx_test()
-def gemm_ex_brcst_test():
-    m1 = helper.make_tensor_value_info('1', TensorProto.FLOAT, [1, 1, 5, 6])
-    m2 = helper.make_tensor_value_info('2', TensorProto.FLOAT, [1, 1, 5, 7])
-    m3 = helper.make_tensor_value_info('3', TensorProto.FLOAT, [1, 1, 6, 1])
-    y = helper.make_tensor_value_info('y', TensorProto.FLOAT, [1, 1, 6, 7])
+def gemm_half_test():
+    A = helper.make_tensor_value_info('A', TensorProto.FLOAT16, [8, 6])
+    B = helper.make_tensor_value_info('B', TensorProto.FLOAT16, [8, 7])
+    C = helper.make_tensor_value_info('C', TensorProto.FLOAT16, [6, 1])
+    Y = helper.make_tensor_value_info('Y', TensorProto.FLOAT16, [6, 7])
 
     node = onnx.helper.make_node('Gemm',
-                                 inputs=['1', '2', '3'],
-                                 outputs=['y'],
+                                 inputs=['A', 'B', 'C'],
+                                 outputs=['Y'],
                                  alpha=0.5,
                                  beta=0.8,
                                  transA=1)
 
-    return ([node], [m1, m2, m3], [y])
+    return ([node], [A, B, C], [Y])
 
 
 @onnx_test()
-def gemm_half_test():
-    m1 = helper.make_tensor_value_info('1', TensorProto.FLOAT16, [1, 1, 8, 6])
-    m2 = helper.make_tensor_value_info('2', TensorProto.FLOAT16, [1, 1, 8, 7])
-    m3 = helper.make_tensor_value_info('3', TensorProto.FLOAT16, [1, 1, 6, 1])
-    y = helper.make_tensor_value_info('y', TensorProto.FLOAT16, [1, 1, 6, 7])
+def gemm_dyn_inner_test():
+    A = helper.make_tensor_value_info('A', TensorProto.FLOAT, [None, 6])
+    B = helper.make_tensor_value_info('B', TensorProto.FLOAT, [None, 7])
+    Y = helper.make_tensor_value_info('Y', TensorProto.FLOAT, [6, 7])
 
     node = onnx.helper.make_node('Gemm',
-                                 inputs=['1', '2', '3'],
-                                 outputs=['y'],
+                                 inputs=['A', 'B'],
+                                 outputs=['Y'],
+                                 alpha=0.5,
+                                 transA=1)
+
+    return ([node], [A, B], [Y])
+
+
+@onnx_test()
+def gemm_dyn_outer_test():
+    A = helper.make_tensor_value_info('A', TensorProto.FLOAT, [5, None])
+    B = helper.make_tensor_value_info('B', TensorProto.FLOAT, [11, 5])
+    Y = helper.make_tensor_value_info('Y', TensorProto.FLOAT, [None, 11])
+
+    node = onnx.helper.make_node('Gemm',
+                                 inputs=['A', 'B'],
+                                 outputs=['Y'],
+                                 alpha=2.0,
+                                 transA=1,
+                                 transB=1)
+
+    return ([node], [A, B], [Y])
+
+
+@onnx_test()
+def gemm_dyn_bias_test():
+    A = helper.make_tensor_value_info('A', TensorProto.FLOAT, [8, None])
+    B = helper.make_tensor_value_info('B', TensorProto.FLOAT, [8, 7])
+    C = helper.make_tensor_value_info('C', TensorProto.FLOAT, [1, 7])
+    Y = helper.make_tensor_value_info('Y', TensorProto.FLOAT, [None, 7])
+
+    node = onnx.helper.make_node('Gemm',
+                                 inputs=['A', 'B', 'C'],
+                                 outputs=['Y'],
+                                 alpha=1.0,
+                                 beta=1.0,
+                                 transA=1)
+
+    return ([node], [A, B, C], [Y])
+
+
+@onnx_test()
+def gemm_rank_error():
+    A = helper.make_tensor_value_info('A', TensorProto.FLOAT, [4, 1, 8, 6])
+    B = helper.make_tensor_value_info('B', TensorProto.FLOAT, [4, 1, 8, 7])
+    C = helper.make_tensor_value_info('C', TensorProto.FLOAT, [6, 7])
+    Y = helper.make_tensor_value_info('Y', TensorProto.FLOAT, [4, 1, 6, 7])
+
+    node = onnx.helper.make_node('Gemm',
+                                 inputs=['A', 'B', 'C'],
+                                 outputs=['Y'],
                                  alpha=0.5,
                                  beta=0.8,
                                  transA=1)
 
-    return ([node], [m1, m2, m3], [y])
+    return ([node], [A, B, C], [Y])
 
 
 @onnx_test()
@@ -6706,6 +6771,92 @@ def transpose_gather_test():
     return ([td, ti, node], [x, i], [y])
 
 
+@onnx_test()
+def trilu_test():
+    x = helper.make_tensor_value_info('x', TensorProto.FLOAT, [3, 4])
+    y = helper.make_tensor_value_info('y', TensorProto.FLOAT, [3, 4])
+
+    node = onnx.helper.make_node(
+        'Trilu',
+        inputs=['x'],
+        outputs=['y'],
+    )
+    return ([node], [x], [y])
+
+
+@onnx_test()
+def trilu_batch_diff_k_test():
+    x = helper.make_tensor_value_info('x', TensorProto.FLOAT, [2, 2, 3])
+    k = np.array([2])
+    y = helper.make_tensor_value_info('y', TensorProto.FLOAT, [2, 2, 3])
+    k_tensor = helper.make_tensor(name='k',
+                                  data_type=TensorProto.INT64,
+                                  dims=k.shape,
+                                  vals=k.astype(np.int64))
+
+    node = onnx.helper.make_node(
+        'Trilu',
+        inputs=['x', 'k'],
+        outputs=['y'],
+    )
+    return ([node], [x], [y], [k_tensor])
+
+
+@onnx_test()
+def trilu_lower_test():
+    x = helper.make_tensor_value_info('x', TensorProto.FLOAT, [3, 4])
+    y = helper.make_tensor_value_info('y', TensorProto.FLOAT, [3, 4])
+
+    node = onnx.helper.make_node('Trilu', inputs=['x'], outputs=['y'], upper=0)
+    return ([node], [x], [y])
+
+
+@onnx_test()
+def trilu_neg_k_test():
+    x = helper.make_tensor_value_info('x', TensorProto.FLOAT, [3, 4])
+    k = np.array([-1])
+    y = helper.make_tensor_value_info('y', TensorProto.FLOAT, [3, 4])
+    k_tensor = helper.make_tensor(name='k',
+                                  data_type=TensorProto.INT64,
+                                  dims=k.shape,
+                                  vals=k.astype(np.int64))
+
+    node = onnx.helper.make_node('Trilu', inputs=['x', 'k'], outputs=['y'])
+    return ([node], [x], [y], [k_tensor])
+
+
+@onnx_test()
+def trilu_out_k_test():
+    x = helper.make_tensor_value_info('x', TensorProto.FLOAT, [3, 4])
+    k = np.array([5])
+    y = helper.make_tensor_value_info('y', TensorProto.FLOAT, [3, 4])
+    k_tensor = helper.make_tensor(name='k',
+                                  data_type=TensorProto.INT64,
+                                  dims=k.shape,
+                                  vals=k.astype(np.int64))
+
+    node = onnx.helper.make_node('Trilu', inputs=['x', 'k'], outputs=['y'])
+    return ([node], [x], [y], [k_tensor])
+
+
+@onnx_test()
+def trilu_row_one_test():
+    x = helper.make_tensor_value_info('x', TensorProto.FLOAT, [1, 4])
+    k = np.array([1])
+    y = helper.make_tensor_value_info('y', TensorProto.FLOAT, [1, 4])
+    k_tensor = helper.make_tensor(name='k',
+                                  data_type=TensorProto.INT64,
+                                  dims=k.shape,
+                                  vals=k.astype(np.int64))
+
+    node = onnx.helper.make_node(
+        'Trilu',
+        inputs=['x', 'k'],
+        outputs=['y'],
+    )
+    return ([node], [x], [y], [k_tensor])
+
+
 @onnx_test()
 def undefined_test():
     x = helper.make_tensor_value_info('0', TensorProto.FLOAT, [2, 3, 4, 5])