Merge branch 'develop' into workspace_size

src/program.cpp

@@ -78,11 +78,11 @@ program& program::operator=(program p)
 
 void program::assign(const program& p)
 {
-    if(!impl)
+    if(not impl)
     {
         impl = std::make_unique<program_impl>();
     }
-    else if(!impl->modules.empty())
+    else if(not impl->modules.empty())
     {
         impl->modules.clear();
     }

src/py/migraphx_py.cpp

@@ -83,7 +83,7 @@ void visit_py(T x, F f)
     {
         f(x.template cast<bool>());
     }
-    else if(py::isinstance<py::int_>(x) || py::hasattr(x, "__index__"))
+    else if(py::isinstance<py::int_>(x) or py::hasattr(x, "__index__"))
     {
         f(x.template cast<int>());
     }

src/quantization.cpp

@@ -70,7 +70,7 @@ void quantize_int8(program& prog,
 {
     std::set<std::string> op_names = {"convolution", "dot"};
     std::set<std::string> input_ins_names(ins_names.begin(), ins_names.end());
-    if(!std::includes(
+    if(not std::includes(
            op_names.begin(), op_names.end(), input_ins_names.begin(), input_ins_names.end()))
     {
         MIGRAPHX_THROW("QUANTIZE_INT8: only support DOT and CONVOLUTION operation");

src/rewrite_pooling.cpp

@@ -47,12 +47,12 @@ void rewrite_pooling::apply(module& m) const
         if(not s.standard())
             continue;
         auto&& op = any_cast<op::pooling>(ins->get_operator());
-        if(!std::all_of(op.padding.begin(), op.padding.end(), [](auto i) { return i == 0; }))
+        if(not std::all_of(op.padding.begin(), op.padding.end(), [](auto i) { return i == 0; }))
             continue;
-        if(!std::all_of(op.stride.begin(), op.stride.end(), [](auto i) { return i == 1; }))
+        if(not std::all_of(op.stride.begin(), op.stride.end(), [](auto i) { return i == 1; }))
             continue;
         auto lens = s.lens();
-        if(!std::equal(lens.begin() + 2, lens.end(), op.lengths.begin(), op.lengths.end()))
+        if(not std::equal(lens.begin() + 2, lens.end(), op.lengths.begin(), op.lengths.end()))
             continue;
         std::int64_t n = s.lens()[0];
         std::int64_t c = s.lens()[1];

src/rewrite_rnn.cpp

@@ -214,7 +214,7 @@ void rewrite_rnn::apply_vanilla_rnn(module& m, instruction_ref ins) const
             ih = m.add_literal(migraphx::literal{ih_shape, data});
         }
 
-        if(!is_forward and variable_seq_len)
+        if(not is_forward and variable_seq_len)
         {
             args[0] =
                 m.insert_instruction(ins, make_op("rnn_var_sl_shift_sequence"), args[0], seq_lens);
@@ -520,7 +520,7 @@ void rewrite_rnn::apply_gru(module& m, instruction_ref ins) const
             ih = m.add_literal(migraphx::literal{ih_shape, data});
         }
 
-        if(!is_forward and variable_seq_len)
+        if(not is_forward and variable_seq_len)
         {
             args[0] =
                 m.insert_instruction(ins, make_op("rnn_var_sl_shift_sequence"), args[0], seq_lens);
@@ -977,7 +977,7 @@ void rewrite_rnn::apply_lstm(module& m, instruction_ref ins) const
             pph = args[7];
         }
 
-        if(!is_forward and variable_seq_len)
+        if(not is_forward and variable_seq_len)
         {
             args[0] =
                 m.insert_instruction(ins, make_op("rnn_var_sl_shift_sequence"), args[0], seq_lens);
@@ -1294,11 +1294,11 @@ bool rewrite_rnn::is_variable_seq_lens(const module& m, instruction_ref seq_lens
             std::vector<int64_t> vec_lens;
             arg_lens.visit([&](auto l) { vec_lens.assign(l.begin(), l.end()); });
             int64_t l = 0;
-            if(!vec_lens.empty())
+            if(not vec_lens.empty())
             {
                 l = vec_lens[0];
             }
-            if(!std::all_of(vec_lens.begin(), vec_lens.end(), [&](auto v) { return v == l; }))
+            if(not std::all_of(vec_lens.begin(), vec_lens.end(), [&](auto v) { return v == l; }))
             {
                 is_var_lens = true;
             }
@@ -1318,7 +1318,7 @@ rewrite_rnn::get_seq_len(const module& m, instruction_ref input, instruction_ref
     bool is_var_lens = is_variable_seq_lens(m, seq_lens);
     auto input_shape = input->get_shape();
     auto length      = input_shape.lens()[0];
-    if(!is_var_lens and seq_lens != m.end())
+    if(not is_var_lens and seq_lens != m.end())
     {
         auto arg_len = seq_lens->eval();
         std::vector<std::size_t> vec_lens;
@@ -1387,7 +1387,7 @@ void rewrite_rnn::replace_last_cell_output(module& m,
 
     if(variable_seq_len)
     {
-        if(!ins_outputs.empty())
+        if(not ins_outputs.empty())
         {
             cell_outputs = m.insert_instruction(
                 std::next(ins),

src/shape.cpp

@@ -477,7 +477,7 @@ bool operator==(const shape::dynamic_dimension& x, const shape::dynamic_dimensio
 
 bool operator!=(const shape::dynamic_dimension& x, const shape::dynamic_dimension& y)
 {
-    return !(x == y);
+    return not(x == y);
 }
 std::ostream& operator<<(std::ostream& os, const shape::dynamic_dimension& x)
 {
@@ -497,7 +497,7 @@ bool operator==(const shape& x, const shape& y)
             x.strides() == y.strides() and x.sub_shapes() == y.sub_shapes());
 }
 
-bool operator!=(const shape& x, const shape& y) { return !(x == y); }
+bool operator!=(const shape& x, const shape& y) { return not(x == y); }
 
 std::ostream& operator<<(std::ostream& os, const shape& x)
 {

src/simplify_algebra.cpp

@@ -787,7 +787,7 @@ MIGRAPHX_PRED_MATCHER(horiz_conv_dot, instruction_ref ins)
     };
     auto dots  = std::count_if(ins->outputs().begin(), ins->outputs().end(), pred("dot"));
     auto convs = std::count_if(ins->outputs().begin(), ins->outputs().end(), pred("convolution"));
-    return !(dots < 2 and convs < 2);
+    return not(dots < 2 and convs < 2);
 }
 
 struct find_conv_dot_horiz_fusion
@@ -969,7 +969,7 @@ struct find_split_reshape
 
         // all outputs are reshape and of the same shape
         auto dims = any_cast<op::reshape>(rsp->get_operator()).dims;
-        if(!same_ops(vec_rsp))
+        if(not same_ops(vec_rsp))
         {
             return;
         }
@@ -1052,7 +1052,7 @@ struct find_split_transpose
 
         // all transpose are the same
         auto perm = any_cast<op::transpose>(trans->get_operator()).dims;
-        if(!same_ops(vec_trans))
+        if(not same_ops(vec_trans))
         {
             return;
         }

src/simplify_reshapes.cpp

@@ -99,7 +99,7 @@ struct find_reshaper
         std::vector<instruction_ref> reshapes{ins};
         while(is_reshaper(reshapes.back()))
         {
-            assert(!reshapes.back()->inputs().empty());
+            assert(not reshapes.back()->inputs().empty());
             assert(m.has_instruction(reshapes.back()->inputs().front()));
             auto input = reshapes.back()->inputs().front();
             reshapes.push_back(input);
@@ -288,7 +288,7 @@ struct find_concat_transpose
         auto permutation = find_permutation(s);
 
         // permutation should be the same for all inputs
-        if(!std::all_of(trans_inputs.begin(), trans_inputs.end(), [&](auto in) {
+        if(not std::all_of(trans_inputs.begin(), trans_inputs.end(), [&](auto in) {
                return (find_permutation(in->get_shape()) == permutation);
            }))
         {

src/targets/cpu/binary.cpp

@@ -49,7 +49,7 @@ struct dnnl_binary : dnnl_op<dnnl_binary, dnnl::binary>
         auto s0 = inputs.at(0);
         auto s1 = inputs.at(1);
         auto r  = s0;
-        if(s0 != s1 or !s0.packed())
+        if(s0 != s1 or not s0.packed())
         {
             r = shape{s0.type(), s0.lens()};
         }

src/targets/fpga/subgraph.cpp

@@ -95,7 +95,7 @@ void subgraph::apply(module_pass_manager& mpm) const
     for(auto it : iterator_for(mod))
     {
         // assuming we want all the params/literals as inputs to the FPGA submodule
-        if(migraphx::starts_with(it->name(), "@param") ||
+        if(migraphx::starts_with(it->name(), "@param") or
            migraphx::starts_with(it->name(), "@literal"))
         {
             literal_inputs.push_back(it);

src/targets/gpu/compile_gen.cpp

@@ -61,12 +61,19 @@ vectorize vectorize::elements(std::size_t axis, const std::vector<shape>& inputs
                    [&](const auto& input) -> std::size_t {
                        auto stride = input.strides()[axis];
                        auto len    = input.lens()[axis];
-                       if(stride != 0 and stride != 1)
+                       if(not contains({0, 1}, stride))
                            return 1;
                        if(len == 1 and input.elements() > sizes.front())
                            return sizes.front();
-                       auto it = std::find_if(
-                           sizes.begin(), sizes.end(), [&](auto i) { return (len % i) == 0; });
+                       auto it = std::find_if(sizes.begin(), sizes.end(), [&](auto vsize) {
+                           // The len is divisible by the size and all the strides are divisible by
+                           // the size
+                           return (len % vsize) == 0 and
+                                  std::all_of(
+                                      input.strides().begin(), input.strides().end(), [&](auto i) {
+                                          return contains({0, 1}, i) or i % vsize == 0;
+                                      });
+                       });
                        if(it != sizes.end())
                            return *it;
                        return 1;

src/targets/gpu/device/include/migraphx/gpu/device/array.hpp

@@ -131,7 +131,7 @@ struct hip_array
 
     friend MIGRAPHX_DEVICE_CONSTEXPR bool operator!=(const hip_array& x, const hip_array& y)
     {
-        return !(x == y);
+        return not(x == y);
     }
     // This uses the product order rather than lexical order
     friend MIGRAPHX_DEVICE_CONSTEXPR bool operator<(const hip_array& x, const hip_array& y)

src/targets/gpu/device/include/migraphx/gpu/device/visit.hpp

@@ -117,12 +117,13 @@ template <class V, class F, class... Ts>
 void hip_visit_all_impl(const shape& s, F f, V&& v, Ts&&... xs)
 {
     std::initializer_list<migraphx::shape::type_t> types = {get_shape(xs).type()...};
-    if(!std::all_of(
+    if(not std::all_of(
            types.begin(), types.end(), [&](migraphx::shape::type_t t) { return t == s.type(); }))
         MIGRAPHX_THROW("Types must be the same");
     std::initializer_list<index_int> ranks = {
         static_cast<index_int>(get_shape(xs).lens().size())...};
-    if(!std::all_of(ranks.begin(), ranks.end(), [&](index_int r) { return r == s.lens().size(); }))
+    if(not std::all_of(
+           ranks.begin(), ranks.end(), [&](index_int r) { return r == s.lens().size(); }))
         MIGRAPHX_THROW("Ranks must be the same");
     visit_tensor_size(s.lens().size(), [&](auto ndim) {
         s.visit_type(hip_visitor([&](auto as) { v(f(xs, ndim, as)...); }));
@@ -134,7 +135,8 @@ void hip_visit_views_impl(const shape& s, F f, V&& v, Ts&&... xs)
 {
     std::initializer_list<index_int> ranks = {
         static_cast<index_int>(get_shape(xs).lens().size())...};
-    if(!std::all_of(ranks.begin(), ranks.end(), [&](index_int r) { return r == s.lens().size(); }))
+    if(not std::all_of(
+           ranks.begin(), ranks.end(), [&](index_int r) { return r == s.lens().size(); }))
         MIGRAPHX_THROW("Ranks must be the same");
     visit_tensor_size(s.lens().size(), [&](auto ndim) { v(f(xs, ndim)...); });
 }

src/targets/gpu/device/multinomial.cpp

@@ -47,7 +47,7 @@ constexpr Iterator upper_bound(Iterator first, Iterator last, const T& value)
         it   = first;
         step = count / 2;
         std::advance(it, step);
-        if(!(value < *it))
+        if(not(value < *it))
         {
             first = ++it;
             count -= step + 1;

src/targets/gpu/fuse_ops.cpp

@@ -26,7 +26,6 @@
 #include <migraphx/gpu/fuse_ops.hpp>
 #include <migraphx/matcher.hpp>
 #include <migraphx/gpu/miopen.hpp>
-#include <migraphx/gpu/clip.hpp>
 #include <migraphx/gpu/convolution.hpp>
 #include <migraphx/gpu/device_name.hpp>
 #include <migraphx/gpu/oper.hpp>
@@ -50,8 +49,6 @@
 #include <migraphx/array.hpp>
 #include <migraphx/permutation.hpp>
 #include <migraphx/make_op.hpp>
-#include <migraphx/op/clip.hpp>
-#include <migraphx/op/contiguous.hpp>
 #include <cmath>
 #include <set>
 
@@ -1036,7 +1033,7 @@ struct find_gemm_pointwise
         // const-fold input if not standard shape since rocblas can't handle it
         if(not c_ins->get_shape().standard())
         {
-            auto c = op::contiguous{};
+            auto c = make_op("contiguous");
             auto l = c.compute(c.compute_shape({c_ins->get_shape()}), {c_ins->eval()});
             c_ins  = m.add_literal(l.get_shape(), l.data());
         }

src/targets/gpu/gemm_impl.cpp

@@ -112,7 +112,7 @@ void gemm_impl(context& ctx,
                bool compute_fp32)
 {
     const bool is_3inputs = (args.size() == 4);
-    if(!is_3inputs)
+    if(not is_3inputs)
     {
         beta = 0;
     }

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

@@ -27,7 +27,7 @@
 #include <migraphx/argument.hpp>
 #include <migraphx/reflect.hpp>
 #include <migraphx/op/gather.hpp>
-#include <migraphx/gpu/miopen.hpp>
+#include <migraphx/gpu/context.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

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

@@ -25,7 +25,6 @@
 #define MIGRAPHX_GUARD_RTGLIB_INT8_CONV_PACK_HPP
 
 #include <migraphx/argument.hpp>
-#include <migraphx/op/quant_dot.hpp>
 #include <migraphx/config.hpp>
 #include <utility>
 

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

@@ -25,7 +25,6 @@
 #define MIGRAPHX_GUARD_RTGLIB_INT8_GEMM_PACK_HPP
 
 #include <migraphx/argument.hpp>
-#include <migraphx/op/quant_dot.hpp>
 #include <migraphx/config.hpp>
 #include <utility>
 

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

@@ -24,22 +24,10 @@
 #ifndef MIGRAPHX_GUARD_RTGLIB_LOGSOFTMAX_HPP
 #define MIGRAPHX_GUARD_RTGLIB_LOGSOFTMAX_HPP
 
-#include <migraphx/gpu/lowering.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/instruction.hpp>
 #include <migraphx/op/logsoftmax.hpp>
-#include <migraphx/generate.hpp>
-#include <migraphx/shape_for_each.hpp>
-#include <migraphx/config.hpp>
-#include <migraphx/gpu/miopen.hpp>
-#include <migraphx/gpu/hip.hpp>
-#include <migraphx/dfor.hpp>
-#include <migraphx/gpu/device/contiguous.hpp>
-#include <migraphx/gpu/device/add.hpp>
-#include <migraphx/iterator_for.hpp>
-#include <migraphx/gpu/rocblas.hpp>
+#include <migraphx/shape.hpp>
+#include <migraphx/reflect.hpp>
 #include <migraphx/gpu/context.hpp>
-#include <utility>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {