Merge branch 'bert-attention-no-transpose-ops' into attention-plus

examples/nlp/python_bert_squad/requirements_bertsquad.txt

-tensorflow==2.6.4
+tensorflow==2.7.2
 onnxruntime
 tokenizers
\ No newline at end of file

src/CMakeLists.txt

@@ -117,7 +117,6 @@ register_migraphx_ops(
     if_op
     im2col
     isnan
-    layernorm
     leaky_relu
     less
     load

src/include/migraphx/op/layernorm.hpp

-#ifndef MIGRAPHX_GUARD_OPERATORS_LAYERNORMALIZATION_HPP
-#define MIGRAPHX_GUARD_OPERATORS_LAYERNORMALIZATION_HPP
-
-#include <array>
-#include <migraphx/check_shapes.hpp>
-#include <migraphx/stringutils.hpp>
-#include <migraphx/streamutils.hpp>
-#include <migraphx/literal.hpp>
-#include <migraphx/shape_for_each.hpp>
-#include <migraphx/config.hpp>
-#include <migraphx/value.hpp>
-#include <migraphx/op/normalize_attribute.hpp>
-#include <migraphx/par_for.hpp>
-#include <cmath>
-#include <utility>
-
-namespace migraphx {
-inline namespace MIGRAPHX_INLINE_NS {
-namespace op {
-
-struct layernorm
-{
-    float epsilon = 1e-3;
-    int64_t axis  = -1;
-
-    template <class Self, class F>
-    static auto reflect(Self& self, F f)
-    {
-        return pack(f(self.epsilon, "epsilon"), f(self.axis, "axis"));
-    }
-
-    value attributes() const
-    {
-        value normalize;
-        normalize["axis"] = value::array{normalize_attribute::include_min};
-        return {{"normalize_axes", normalize}};
-    }
-
-    std::string name() const { return "layernorm"; }
-
-    shape normalize_compute_shape(std::vector<shape> inputs) const
-    {
-        if(inputs.size() == 2)
-        {
-            if(inputs.at(1).lens().front() != inputs.front().lens().at(axis))
-                MIGRAPHX_THROW("LAYERNORM: weights have wrong shape");
-        }
-        if(inputs.size() == 3)
-        {
-            if(inputs.at(2).lens().front() != inputs.front().lens().at(axis))
-                MIGRAPHX_THROW("LAYERNORM: bias has wrong shape");
-        }
-
-        return inputs.front();
-    }
-
-    argument compute(const shape& output_shape, std::vector<argument> args) const
-    {
-        argument result{output_shape};
-        auto x_lens     = args.front().get_shape().lens();
-        auto norm_count = std::accumulate(
-            x_lens.begin(), x_lens.begin() + axis, std::size_t{1}, std::multiplies<std::size_t>());
-        auto norm_size = std::accumulate(
-            x_lens.begin() + axis, x_lens.end(), std::size_t{1}, std::multiplies<std::size_t>());
-
-        if(args.size() == 3)
-        {
-            visit_all(result, args[0], args[1], args[2])(
-                [&](auto output, auto data, auto weights, auto bias) {
-                    par_for(norm_count, [&](auto idx) {
-                        auto offset        = idx * norm_size;
-                        double mean        = 0;
-                        double mean_square = 0;
-                        for(std::size_t i = 0; i < norm_size; ++i)
-                        {
-                            mean += data[offset + i];
-                            mean_square += data[offset + i] * data[offset + i];
-                        }
-                        mean /= norm_size;
-                        mean_square = sqrt(mean_square / norm_size - mean * mean + epsilon);
-                        for(std::size_t i = 0; i < norm_size; ++i)
-                        {
-                            if(args.size() == 3)
-                                output[offset + i] =
-                                    (data[offset + i] - mean) / mean_square * weights[i] + bias[i];
-                            else
-                                output[offset + i] =
-                                    (data[offset + i] - mean) / mean_square * weights[i];
-                        }
-                    });
-                });
-        }
-        else
-        {
-            visit_all(result, args[0])([&](auto output, auto data) {
-                par_for(norm_count, [&](auto idx) {
-                    auto offset        = idx * norm_size;
-                    double mean        = 0;
-                    double mean_square = 0;
-                    for(std::size_t i = 0; i < norm_size; ++i)
-                    {
-                        mean += data[offset + i];
-                        mean_square += data[offset + i] * data[offset + i];
-                    }
-                    mean /= norm_size;
-                    mean_square = sqrt(mean_square / norm_size - mean * mean + epsilon);
-                    for(std::size_t i = 0; i < norm_size; ++i)
-                    {
-                        output[offset + i] = (data[offset + i] - mean) / mean_square;
-                        // scale and bias handled by pointwise ops
-                    }
-                });
-            });
-        }
-
-        return result;
-    }
-};
-
-} // namespace op
-} // namespace MIGRAPHX_INLINE_NS
-} // namespace migraphx
-
-#endif

src/include/migraphx/operators.hpp

@@ -43,7 +43,6 @@
 #include <migraphx/op/if_op.hpp>
 #include <migraphx/op/im2col.hpp>
 #include <migraphx/op/isnan.hpp>
-#include <migraphx/op/layernorm.hpp>
 #include <migraphx/op/leaky_relu.hpp>
 #include <migraphx/op/less.hpp>
 #include <migraphx/op/load.hpp>

src/onnx/parse_layernorm.cpp

 #include <migraphx/onnx/op_parser.hpp>
 #include <migraphx/ranges.hpp>
 #include <migraphx/make_op.hpp>
-#include <migraphx/op/layernorm.hpp>
 #include <migraphx/argument.hpp>
 #include <migraphx/instruction.hpp>
 
@@ -33,7 +32,7 @@ struct parse_layernorm : op_parser<parse_layernorm>
         }
         if(contains(info.attributes, "axis"))
         {
-            epsilon = parser.parse_value(info.attributes.at("axis")).at<int64_t>();
+            axis = parser.parse_value(info.attributes.at("axis")).at<int64_t>();
         }
 
         auto epsilon_lit = info.add_literal(literal{shape{x_type, {1}}, {epsilon}});

src/targets/gpu/CMakeLists.txt

@@ -149,7 +149,6 @@ add_library(migraphx_gpu
     int8_conv_pack.cpp
     int8_gemm_pack.cpp
     kernel.cpp
-    layernorm.cpp
     lowering.cpp
     logsoftmax.cpp
     loop.cpp
@@ -206,7 +205,6 @@ register_migraphx_gpu_ops(hip_
     floor
     gather
     greater
-    layernorm
     less
     log
     logsoftmax

src/targets/gpu/include/migraphx/gpu/layernorm.hpp

-#ifndef MIGRAPHX_GUARD_RTGLIB_LAYERNORM_HPP
-#define MIGRAPHX_GUARD_RTGLIB_LAYERNORM_HPP
-
-#include <migraphx/op/layernorm.hpp>
-#include <migraphx/shape.hpp>
-#include <migraphx/reflect.hpp>
-#include <migraphx/argument.hpp>
-
-namespace migraphx {
-inline namespace MIGRAPHX_INLINE_NS {
-namespace gpu {
-
-struct context;
-
-struct hip_layernorm
-{
-    op::layernorm op;
-
-    template <class Self, class F>
-    static auto reflect(Self& self, F f)
-    {
-        return migraphx::reflect(self.op, f);
-    }
-
-    std::string name() const { return "gpu::layernorm"; }
-    shape compute_shape(std::vector<shape> inputs) const;
-    argument
-    compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;
-    void finalize(context&, const shape&, const std::vector<shape>&);
-    std::ptrdiff_t output_alias(const std::vector<shape>& shapes) const
-    {
-        return shapes.size() - 1;
-    }
-};
-
-} // namespace gpu
-} // namespace MIGRAPHX_INLINE_NS
-} // namespace migraphx
-
-#endif

src/targets/gpu/layernorm.cpp

-#include <migraphx/gpu/layernorm.hpp>
-#include <migraphx/gpu/context.hpp>
-#include <migraphx/gpu/device/layernorm.hpp>
-
-namespace migraphx {
-inline namespace MIGRAPHX_INLINE_NS {
-namespace gpu {
-
-shape hip_layernorm::compute_shape(std::vector<shape> inputs) const
-{
-    inputs.pop_back();
-    return op.normalize_compute_shape(inputs);
-}
-
-argument hip_layernorm::compute(context& ctx, const shape&, const std::vector<argument>& args) const
-{
-    device::layernorm(ctx.get_stream().get(), args.back(), args[0]);
-
-    return args.back();
-}
-
-} // namespace gpu
-} // namespace MIGRAPHX_INLINE_NS
-} // namespace migraphx

src/targets/gpu/lowering.cpp

@@ -11,7 +11,6 @@
 #include <migraphx/op/dot.hpp>
 #include <migraphx/op/elu.hpp>
 #include <migraphx/op/if_op.hpp>
-#include <migraphx/op/layernorm.hpp>
 #include <migraphx/op/leaky_relu.hpp>
 #include <migraphx/op/lrn.hpp>
 #include <migraphx/op/pooling.hpp>
@@ -30,7 +29,6 @@
 #include <migraphx/gpu/gemm.hpp>
 #include <migraphx/gpu/greater.hpp>
 #include <migraphx/gpu/int8_conv_pack.hpp>
-#include <migraphx/gpu/layernorm.hpp>
 #include <migraphx/gpu/leaky_relu.hpp>
 #include <migraphx/gpu/less.hpp>
 #include <migraphx/gpu/logical_and.hpp>
@@ -141,7 +139,6 @@ struct miopen_apply
         add_generic_op("exp");
         add_generic_op("floor");
         add_generic_op("greater");
-        add_generic_op("layernorm");
         add_generic_op("less");
         add_generic_op("log");
         add_generic_op("logical_and");

test/ref_ops_test.cpp

@@ -2435,50 +2435,6 @@ TEST_CASE(imagescaler_test)
     EXPECT(migraphx::verify_range(results_vector, gold));
 }
 
-TEST_CASE(layernorm_test)
-{
-    {
-        // with scale and bias
-        migraphx::program p;
-        auto* mm = p.get_main_module();
-        migraphx::shape sx{migraphx::shape::float_type, {1, 2, 3}};
-        migraphx::shape swb{migraphx::shape::float_type, {3}};
-        std::vector<float> x_vec{1.0, 2.0, 3.0, 4.0, 5.0, 6.0};
-        auto x = mm->add_literal(migraphx::literal{sx, x_vec});
-        auto w = mm->add_literal(migraphx::literal{swb, {1.0, 1.0, 1.0}});
-        auto b = mm->add_literal(migraphx::literal{swb, {0.0, 0.0, 0.0}});
-        mm->add_instruction(migraphx::make_op("layernorm", {{"epsilon", 1e-5}}), x, w, b);
-        p.compile(migraphx::ref::target{});
-        auto result = p.eval({}).back();
-        std::vector<float> results_vector(1 * 2 * 3);
-        result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
-        std::vector<float> gold = {-1.22474f, 0.0f, 1.22474f, -1.22474f, 0.0f, 1.22474f};
-        for(auto&& i : results_vector)
-            std::cout << i << ", ";
-        std::cout << std::endl;
-        EXPECT(migraphx::verify_range(results_vector, gold));
-    }
-
-    {
-        // without scale and bias
-        migraphx::program p;
-        auto* mm = p.get_main_module();
-        migraphx::shape sx{migraphx::shape::float_type, {1, 2, 3}};
-        std::vector<float> x_vec{1.0, 2.0, 3.0, 4.0, 5.0, 6.0};
-        auto x = mm->add_literal(migraphx::literal{sx, x_vec});
-        mm->add_instruction(migraphx::make_op("layernorm", {{"epsilon", 1e-5}}), x);
-        p.compile(migraphx::ref::target{});
-        auto result = p.eval({}).back();
-        std::vector<float> results_vector(1 * 2 * 3);
-        result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
-        std::vector<float> gold = {-1.22474f, 0.0f, 1.22474f, -1.22474f, 0.0f, 1.22474f};
-        for(auto&& i : results_vector)
-            std::cout << i << ", ";
-        std::cout << std::endl;
-        EXPECT(migraphx::verify_range(results_vector, gold));
-    }
-}
-
 TEST_CASE(leaky_relu_test)
 {
     migraphx::program p;

test/verify/test_layernorm_op.cpp

-#include "verify_program.hpp"
-#include <migraphx/program.hpp>
-#include <migraphx/generate.hpp>
-#include <migraphx/make_op.hpp>
-
-struct test_layernorm_op : verify_program<test_layernorm_op>
-{
-    migraphx::program create_program() const
-    {
-        migraphx::program p;
-        auto* mm = p.get_main_module();
-        auto x =
-            mm->add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 384, 768}});
-        mm->add_instruction(migraphx::make_op("layernorm", {{"axis", -1}, {"epsilon", 1e-12}}), x);
-
-        return p;
-    }
-};

tools/install_prereqs.sh

@@ -27,7 +27,7 @@ elif [ "$#" -eq 1 ]; then
   PREFIX=$1
 fi
 
-echo "Dependencies are install at $PREFIX"
+echo "Dependencies are installed at $PREFIX"
 
 # Install deps with rbuild
 rbuild prepare -d $PREFIX -s develop
@@ -35,3 +35,5 @@ rbuild prepare -d $PREFIX -s develop
 # install onnx package for unit tests
 pip3 install onnx==1.8.1 numpy==1.18.5 typing==3.7.4 pytest==6.0.1 packaging==16.8
 
+# pin version of protobuf in Python for onnx runtime unit tests
+pip3 install protobuf==3.20.0