Merge branch 'develop' into fix_parse_if

test/operators.cpp

@@ -29,7 +29,6 @@
 #include <migraphx/module.hpp>
 #include <sstream>
 #include <string>
-#include <migraphx/make_op.hpp>
 
 #include <migraphx/serialize.hpp>
 

test/simplify_qdq_test.cpp

@@ -33,7 +33,6 @@
 #include <migraphx/matcher.hpp>
 #include <migraphx/generate.hpp>
 #include <migraphx/verify.hpp>
-#include <migraphx/ref/target.hpp>
 #include <migraphx/apply_alpha_beta.hpp>
 
 bool is_convolution(const migraphx::instruction& ins) { return ins.name() == "convolution"; }

test/verify/test_elu.cpp → test/verify/quant_conv_1d.cpp

@@ -27,14 +27,21 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct test_elu : verify_program<test_elu>
+struct quant_conv_1d : verify_program<quant_conv_1d>
 {
     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, {4, 3, 3, 3}});
-        mm->add_instruction(migraphx::make_op("leaky_relu", {{"alpha", 1.0}}), x);
+        migraphx::shape a_shape{migraphx::shape::int8_type, {2, 3, 4}};
+        auto pa = mm->add_parameter("a", a_shape);
+        migraphx::shape c_shape{migraphx::shape::int8_type, {2, 3, 3}};
+        auto pc = mm->add_parameter("c", c_shape);
+        mm->add_instruction(
+            migraphx::make_op("quant_convolution",
+                              {{"padding", {0}}, {"stride", {1}}, {"dilation", {1}}}),
+            pa,
+            pc);
         return p;
     }
 };

test/verify/test_leaky_relu.cpp → test/verify/test_pad_large.cpp

@@ -27,14 +27,16 @@
 #include <migraphx/generate.hpp>
 #include <migraphx/make_op.hpp>
 
-struct test_leaky_relu : verify_program<test_leaky_relu>
+struct test_pad_large : verify_program<test_pad_large>
 {
     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, {4, 3, 3, 3}});
-        mm->add_instruction(migraphx::make_op("leaky_relu", {{"alpha", 0.01}}), x);
+        migraphx::shape s0{migraphx::shape::float_type, {586, 3, 224, 224}};
+        std::vector<int64_t> pads0 = {0, 0, 1, 1, 0, 0, 1, 1};
+        auto l0                    = mm->add_parameter("x", s0);
+        mm->add_instruction(migraphx::make_op("pad", {{"pads", pads0}}), l0);
         return p;
     }
 };

src/targets/gpu/include/migraphx/gpu/deconvolution.hpp → test/verify/test_shape_alloc.cpp

@@ -21,47 +21,41 @@
  * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
  * THE SOFTWARE.
  */
-#ifndef MIGRAPHX_GUARD_RTGLIB_DECONVOLUTION_HPP
-#define MIGRAPHX_GUARD_RTGLIB_DECONVOLUTION_HPP
-
-#include <migraphx/shape.hpp>
-#include <migraphx/op/deconvolution.hpp>
-#include <migraphx/gpu/miopen.hpp>
-
-namespace migraphx {
-inline namespace MIGRAPHX_INLINE_NS {
-namespace gpu {
-
-struct context;
-
-struct miopen_deconvolution
+#include "verify_program.hpp"
+#include <migraphx/program.hpp>
+#include <migraphx/generate.hpp>
+#include <migraphx/make_op.hpp>
+#include <migraphx/op/reduce_mean.hpp>
+
+/**
+ * @brief test_shape_alloc sets up a situation that could lead to an exception "convolution: Shapes
+ * are not in standard layout" if a "replace_allocate" compiler pass is not followed with
+ *   "adjust_allocation".  The last transpose instruction generates a shape with a stride of 1 in
+ *   the 2nd index, a non-standard layout that should be reallocated by adjust_allocation.
+ */
+struct test_shape_alloc : verify_program<test_shape_alloc>
 {
-    op::deconvolution op;
-    shared<convolution_descriptor> cd;
-    miopenConvFwdAlgorithm_t algo{};
-    uint64_t solution_id = 0;
-
-    template <class Self, class F>
-    static auto reflect(Self& self, F f)
+    migraphx::program create_program() const
     {
-        return pack_join(op::deconvolution::reflect(self.op, f),
-                         pack(f(self.solution_id, "solution_id")));
-    }
-
-    std::string name() const { return "gpu::deconv"; }
-    shape compute_shape(const std::vector<shape>& inputs) const;
-    argument
-    compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;
-    shape find(context& ctx, const shape& output_shape, std::vector<shape> inputs);
-    void finalize(context& ctx, const shape& output_shape, std::vector<shape> inputs);
-    std::ptrdiff_t output_alias(const std::vector<shape>& shapes) const
-    {
-        return shapes.size() - 1;
+        migraphx::program p;
+        auto* mm = p.get_main_module();
+
+        auto weights = mm->add_literal(migraphx::generate_literal(
+            migraphx::shape{migraphx::shape::float_type, {11, 8, 1, 1}, {8, 1, 1, 1}}));
+
+        auto x = mm->add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 8, 7, 7}});
+        auto transpose1 =
+            mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 2, 3, 1}}}),
+                                x); //  -> float_type, {1, 7, 7, 8}, {392, 7, 1, 49}
+        auto reduce_ins =
+            mm->add_instruction(migraphx::make_op("reduce_mean", {{"axes", {1, 2}}}),
+                                transpose1); //  -> float_type, {1, 1, 1, 8}, {8, 8, 8, 1}
+        auto transpose2 =
+            mm->add_instruction(migraphx::make_op("transpose", {{"permutation", {0, 3, 1, 2}}}),
+                                reduce_ins); //  -> float_type, {1, 8, 1, 1}, {8, 1, 8, 8}
+        auto conv_op = migraphx::make_op("convolution");
+        mm->add_instruction(conv_op, transpose2, weights);
+
+        return p;
     }
 };
-
-} // namespace gpu
-} // namespace MIGRAPHX_INLINE_NS
-} // namespace migraphx
-
-#endif

tools/include/operation.hpp

@@ -32,6 +32,8 @@
 #include <utility>
 #include <unordered_map>
 #include <migraphx/reflect.hpp>
+#include <migraphx/dyn_output.hpp>
+#include <migraphx/functional.hpp>
 #include <migraphx/streamutils.hpp>
 #include <migraphx/normalize_attributes.hpp>
 #include <migraphx/argument.hpp>
@@ -199,9 +201,12 @@ auto compute_op(rank<1>,
                 context& ctx,
                 const shape& output_shape,
                 const std::vector<argument>& input)
-    -> decltype(x.compute(auto_any_cast(ctx), output_shape, input))
+    -> decltype(x.compute(auto_any_cast(ctx),
+                          make_compute_output_shape(pack(x, output_shape, input)),
+                          input))
 {
-    return x.compute(auto_any_cast(ctx), output_shape, input);
+    return x.compute(
+        auto_any_cast(ctx), make_compute_output_shape(pack(x, output_shape, input)), input);
 }
 
 template <class T>
@@ -220,9 +225,9 @@ compute_op(const T& x, context& ctx, const shape& output_shape, const std::vecto
 
 template <class T>
 auto compute_op(rank<1>, const T& x, const shape& output_shape, const std::vector<argument>& input)
-    -> decltype(x.compute(output_shape, input))
+    -> decltype(x.compute(make_compute_output_shape(pack(x, output_shape, input)), input))
 {
-    return x.compute(output_shape, input);
+    return x.compute(make_compute_output_shape(pack(x, output_shape, input)), input);
 }
 
 template <class T>
@@ -244,9 +249,11 @@ auto compute_op(rank<1>,
                 const shape& output,
                 const std::vector<argument>& inputs,
                 const std::vector<module_ref>& module_args,
-                F f) -> decltype(x.compute(output, inputs, module_args, f))
+                F f)
+    -> decltype(
+        x.compute(make_compute_output_shape(pack(x, output, inputs)), inputs, module_args, f))
 {
-    return x.compute(output, inputs, module_args, f);
+    return x.compute(make_compute_output_shape(pack(x, output, inputs)), inputs, module_args, f);
 }
 
 template <class T, class F>
@@ -278,9 +285,17 @@ auto compute_op(rank<4>,
                 const shape& output,
                 const std::vector<argument>& inputs,
                 const std::vector<module_ref>& module_args,
-                F f) -> decltype(x.compute(auto_any_cast(ctx), output, inputs, module_args, f))
+                F f) -> decltype(x.compute(auto_any_cast(ctx),
+                                           make_compute_output_shape(pack(x, output, inputs)),
+                                           inputs,
+                                           module_args,
+                                           f))
 {
-    return x.compute(auto_any_cast(ctx), output, inputs, module_args, f);
+    return x.compute(auto_any_cast(ctx),
+                     make_compute_output_shape(pack(x, output, inputs)),
+                     inputs,
+                     module_args,
+                     f);
 }
 
 template <class T, class F>
@@ -290,9 +305,11 @@ auto compute_op(rank<3>,
                 const shape& output,
                 const std::vector<argument>& inputs,
                 const std::vector<module_ref>& module_args,
-                F f) -> decltype(x.compute(output, inputs, module_args, f))
+                F f)
+    -> decltype(
+        x.compute(make_compute_output_shape(pack(x, output, inputs)), inputs, module_args, f))
 {
-    return x.compute(output, inputs, module_args, f);
+    return x.compute(make_compute_output_shape(pack(x, output, inputs)), inputs, module_args, f);
 }
 
 template <class T, class F>
@@ -302,9 +319,10 @@ auto compute_op(rank<2>,
                 const shape& output,
                 const std::vector<argument>& inputs,
                 const std::vector<module_ref>&,
-                F) -> decltype(x.compute(output, inputs))
+                F)
+    -> decltype(x.compute(make_compute_output_shape(pack(x, output, inputs)), inputs))
 {
-    return x.compute(output, inputs);
+    return x.compute(make_compute_output_shape(pack(x, output, inputs)), inputs);
 }
 
 template <class T, class F>
@@ -314,9 +332,12 @@ auto compute_op(rank<1>,
                 const shape& output,
                 const std::vector<argument>& inputs,
                 const std::vector<module_ref>&,
-                F) -> decltype(x.compute(auto_any_cast(ctx), output, inputs))
+                F) -> decltype(x.compute(auto_any_cast(ctx),
+                                         make_compute_output_shape(pack(x, output, inputs)),
+                                         inputs))
 {
-    return x.compute(auto_any_cast(ctx), output, inputs);
+    return x.compute(
+        auto_any_cast(ctx), make_compute_output_shape(pack(x, output, inputs)), inputs);
 }
 
 template <class T, class F>
@@ -348,7 +369,8 @@ auto is_context_free_op(rank<1>,
                         const T& x,
                         const shape& output_shape,
                         const std::vector<argument>& input)
-    -> decltype(x.compute(output_shape, input), std::true_type{});
+    -> decltype(x.compute(make_compute_output_shape(pack(x, output_shape, input)), input),
+                std::true_type{});
 
 template <class T>
 auto is_context_free_op(rank<0>, const T&, const shape&, const std::vector<argument>&)