Merge branch 'master' into fp16

src/targets/gpu/include/migraph/gpu/hip.hpp

@@ -2,9 +2,11 @@
 #define MIGRAPH_GUARD_MIGRAPHLIB_HIP_HPP
 
 #include <migraph/operators.hpp>
+#include <migraph/config.hpp>
 #include <utility>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 migraph::argument allocate_gpu(const migraph::shape& s, bool host = false);
@@ -67,6 +69,7 @@ struct hip_write
     {
         return to_gpu(args.front());
     }
+    int output_alias(const std::vector<shape>&) const { return 0; }
 };
 
 struct hip_copy
@@ -82,8 +85,11 @@ struct hip_copy
         copy_to_gpu(args[0], args[1]);
         return args[1];
     }
+    int output_alias(const std::vector<shape>&) const { return 1; }
 };
+
 } // namespace gpu
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph
 
 #endif

src/targets/gpu/include/migraph/gpu/leaky_relu.hpp

@@ -7,6 +7,7 @@
 #include <migraph/operators.hpp>
 #include <migraph/generate.hpp>
 #include <migraph/shape_for_each.hpp>
+#include <migraph/config.hpp>
 #include <migraph/gpu/miopen.hpp>
 #include <migraph/gpu/hip.hpp>
 #include <migraph/dfor.hpp>
@@ -18,6 +19,7 @@
 #include <utility>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 struct miopen_leaky_relu
@@ -27,10 +29,11 @@ struct miopen_leaky_relu
     shape compute_shape(const std::vector<shape>& inputs) const;
     argument
     compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;
+    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu
-
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph
 
 #endif

src/targets/gpu/include/migraph/gpu/lowering.hpp

@@ -2,9 +2,11 @@
 #define MIGRAPH_GUARD_RTGLIB_MIOPEN_LOWERING_HPP
 
 #include <migraph/program.hpp>
+#include <migraph/config.hpp>
 #include <migraph/gpu/context.hpp>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 struct lowering
@@ -15,7 +17,7 @@ struct lowering
 };
 
 } // namespace gpu
-
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph
 
 #endif

src/targets/gpu/include/migraph/gpu/miopen.hpp

@@ -4,8 +4,10 @@
 #include <migraph/manage_ptr.hpp>
 #include <migraph/operators.hpp>
 #include <miopen/miopen.h>
+#include <migraph/config.hpp>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 using miopen_handle          = MIGRAPH_MANAGE_PTR(miopenHandle_t, miopenDestroy);
@@ -115,7 +117,7 @@ inline fused_operator_args make_fused_args()
 }
 
 } // namespace gpu
-
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph
 
 #endif

src/targets/gpu/include/migraph/gpu/mul.hpp

@@ -7,6 +7,7 @@
 #include <migraph/operators.hpp>
 #include <migraph/generate.hpp>
 #include <migraph/shape_for_each.hpp>
+#include <migraph/config.hpp>
 #include <migraph/gpu/miopen.hpp>
 #include <migraph/gpu/hip.hpp>
 #include <migraph/dfor.hpp>
@@ -18,6 +19,7 @@
 #include <utility>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 struct hip_mul
@@ -25,10 +27,11 @@ struct hip_mul
     std::string name() const { return "gpu::mul"; }
     shape compute_shape(const std::vector<shape>& inputs) const;
     argument compute(context&, const shape&, const std::vector<argument>& args) const;
+    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu
-
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph
 
 #endif

src/targets/gpu/include/migraph/gpu/pooling.hpp

@@ -7,6 +7,7 @@
 #include <migraph/operators.hpp>
 #include <migraph/generate.hpp>
 #include <migraph/shape_for_each.hpp>
+#include <migraph/config.hpp>
 #include <migraph/gpu/miopen.hpp>
 #include <migraph/gpu/hip.hpp>
 #include <migraph/dfor.hpp>
@@ -18,6 +19,7 @@
 #include <utility>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 struct miopen_pooling
@@ -29,10 +31,11 @@ struct miopen_pooling
     shape compute_shape(const std::vector<shape>& inputs) const;
     argument
     compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;
+    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu
-
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph
 
 #endif

src/targets/gpu/include/migraph/gpu/relu.hpp

@@ -7,6 +7,7 @@
 #include <migraph/operators.hpp>
 #include <migraph/generate.hpp>
 #include <migraph/shape_for_each.hpp>
+#include <migraph/config.hpp>
 #include <migraph/gpu/miopen.hpp>
 #include <migraph/gpu/hip.hpp>
 #include <migraph/dfor.hpp>
@@ -18,6 +19,7 @@
 #include <utility>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 struct miopen_relu
@@ -27,10 +29,11 @@ struct miopen_relu
     shape compute_shape(const std::vector<shape>& inputs) const;
     argument
     compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;
+    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu
-
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph
 
 #endif

src/targets/gpu/include/migraph/gpu/rocblas.hpp

@@ -3,9 +3,11 @@
 
 #include <migraph/manage_ptr.hpp>
 #include <migraph/operators.hpp>
+#include <migraph/config.hpp>
 #include <rocblas.h>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 using rocblas_handle_ptr = MIGRAPH_MANAGE_PTR(rocblas_handle, rocblas_destroy_handle);
@@ -14,7 +16,7 @@ rocblas_handle_ptr create_rocblas_handle_ptr();
 rocblas_handle_ptr create_rocblas_handle_ptr(hipStream_t s);
 
 } // namespace gpu
-
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph
 
 #endif

src/targets/gpu/include/migraph/gpu/softmax.hpp

@@ -7,6 +7,7 @@
 #include <migraph/operators.hpp>
 #include <migraph/generate.hpp>
 #include <migraph/shape_for_each.hpp>
+#include <migraph/config.hpp>
 #include <migraph/gpu/miopen.hpp>
 #include <migraph/gpu/hip.hpp>
 #include <migraph/dfor.hpp>
@@ -18,6 +19,7 @@
 #include <utility>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 struct miopen_softmax
@@ -27,10 +29,11 @@ struct miopen_softmax
     shape compute_shape(const std::vector<shape>& inputs) const;
     argument
     compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;
+    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu
-
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph
 
 #endif

src/targets/gpu/include/migraph/gpu/target.hpp

@@ -2,8 +2,10 @@
 #define MIGRAPH_GUARD_MIGRAPHLIB_MIOPEN_TARGET_HPP
 
 #include <migraph/program.hpp>
+#include <migraph/config.hpp>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 struct target
@@ -12,7 +14,9 @@ struct target
     std::vector<pass> get_passes(migraph::context& gctx) const;
     migraph::context get_context() const;
 };
+
 } // namespace gpu
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph
 
 #endif

src/targets/gpu/include/migraph/gpu/write_literals.hpp

@@ -5,6 +5,7 @@
 #include <migraph/gpu/context.hpp>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 
 namespace gpu {
 
@@ -17,7 +18,7 @@ struct write_literals
 };
 
 } // namespace gpu
-
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph
 
 #endif

src/targets/gpu/leaky_relu.cpp

@@ -5,6 +5,7 @@
 #include <utility>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 shape miopen_leaky_relu::compute_shape(const std::vector<shape>& inputs) const
@@ -33,5 +34,5 @@ argument miopen_leaky_relu::compute(context& ctx,
 }
 
 } // namespace gpu
-
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph

src/targets/gpu/lowering.cpp

@@ -27,6 +27,7 @@
 #include <utility>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 struct miopen_apply
@@ -50,9 +51,9 @@ struct miopen_apply
             {
                 check_shape(s, apply_convolution(it));
             }
-            else if(it->name() == "activation")
+            else if(it->name() == "relu")
             {
-                check_shape(s, apply_activation(it));
+                check_shape(s, apply_relu(it));
             }
             else if(it->name() == "leaky_relu")
             {
@@ -131,17 +132,13 @@ struct miopen_apply
             ins, miopen_pooling{op, std::move(pd)}, ins->inputs().at(0), output);
     }
 
-    instruction_ref apply_activation(instruction_ref ins)
+    instruction_ref apply_relu(instruction_ref ins)
     {
-        auto&& op = any_cast<op::activation>(ins->get_operator());
-        auto ad   = make_relu();
-        if(op.mode == "relu")
-        {
-            auto output = insert_allocation(ins, ins->get_shape());
-            return prog->replace_instruction(
-                ins, miopen_relu{std::move(ad)}, ins->inputs().at(0), output);
-        }
-        return ins;
+        auto ad = make_relu();
+
+        auto output = insert_allocation(ins, ins->get_shape());
+        return prog->replace_instruction(
+            ins, miopen_relu{std::move(ad)}, ins->inputs().at(0), output);
     }
 
     instruction_ref apply_leaky_relu(instruction_ref ins)
@@ -224,4 +221,5 @@ struct miopen_apply
 
 void lowering::apply(program& p) const { miopen_apply{&p, ctx}.apply(); }
 } // namespace gpu
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph

src/targets/gpu/mul.cpp

@@ -5,6 +5,7 @@
 #include <utility>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 shape hip_mul::compute_shape(const std::vector<shape>& inputs) const
@@ -21,5 +22,5 @@ argument hip_mul::compute(context& ctx, const shape&, const std::vector<argument
 }
 
 } // namespace gpu
-
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph

src/targets/gpu/pooling.cpp

@@ -5,6 +5,7 @@
 #include <utility>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 shape miopen_pooling::compute_shape(const std::vector<shape>& inputs) const
@@ -37,5 +38,5 @@ argument miopen_pooling::compute(context& ctx,
 }
 
 } // namespace gpu
-
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph

src/targets/gpu/relu.cpp

@@ -5,6 +5,7 @@
 #include <utility>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 shape miopen_relu::compute_shape(const std::vector<shape>& inputs) const
@@ -33,5 +34,5 @@ argument miopen_relu::compute(context& ctx,
 }
 
 } // namespace gpu
-
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph

src/targets/gpu/rocblas.cpp

 #include <migraph/gpu/rocblas.hpp>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 rocblas_handle_ptr create_rocblas_handle_ptr()
@@ -18,5 +19,5 @@ rocblas_handle_ptr create_rocblas_handle_ptr(hipStream_t s)
 }
 
 } // namespace gpu
-
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph

src/targets/gpu/softmax.cpp

@@ -5,6 +5,7 @@
 #include <utility>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 shape miopen_softmax::compute_shape(const std::vector<shape>& inputs) const
@@ -32,5 +33,5 @@ argument miopen_softmax::compute(context& ctx,
 }
 
 } // namespace gpu
-
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph

src/targets/gpu/target.cpp

@@ -15,8 +15,11 @@
 #include <migraph/eliminate_contiguous.hpp>
 #include <migraph/common_subexpression_elimination.hpp>
 #include <migraph/fwd_conv_batchnorm_rewrite.hpp>
+#include <migraph/eliminate_concat.hpp>
+#include <migraph/gpu/concat_gpu_opt.hpp>
 
 namespace migraph {
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 std::vector<pass> target::get_passes(migraph::context& gctx) const
@@ -38,6 +41,8 @@ std::vector<pass> target::get_passes(migraph::context& gctx) const
         simplify_reshapes{},
         dead_code_elimination{},
         lowering{ctx},
+        eliminate_concat{concat_gpu_optimization{}},
+        dead_code_elimination{},
         eliminate_contiguous{},
         dead_code_elimination{},
         fuse_ops{&ctx},
@@ -56,4 +61,5 @@ std::string target::name() const { return "miopen"; }
 
 migraph::context target::get_context() const { return context{}; }
 } // namespace gpu
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph

src/targets/gpu/write_literals.cpp

@@ -5,7 +5,7 @@
 #include <migraph/env.hpp>
 
 namespace migraph {
-
+inline namespace MIGRAPH_INLINE_NS {
 namespace gpu {
 
 MIGRAPH_DECLARE_ENV_VAR(MIGRAPH_COPY_LITERALS)
@@ -51,5 +51,7 @@ void write_literals::apply(program& p) const
         }
     }
 }
+
 } // namespace gpu
+} // namespace MIGRAPH_INLINE_NS
 } // namespace migraph