Add output_alias functions

src/include/migraph/operators.hpp

@@ -296,6 +296,10 @@ struct transpose
     {
         return {std::move(output_shape), std::move(args.front().data)};
     }
+    int output_alias(const std::vector<shape>&) const
+    {
+        return 0;
+    }
 };
 
 struct contiguous
@@ -359,6 +363,10 @@ struct concat
         new_lens[axis] = new_dim_axis;
         return {type, new_lens};
     }
+    int output_alias(const std::vector<shape>&) const
+    {
+        return 0;
+    }
 };
 
 struct slice
@@ -440,6 +448,10 @@ struct slice
         auto offset = compute_offset(input.get_shape()) * output_shape.type_size();
         return {std::move(output_shape), [=] { return input.data() + offset; }};
     }
+    int output_alias(const std::vector<shape>&) const
+    {
+        return 0;
+    }
 };
 
 struct squeeze
@@ -487,6 +499,10 @@ struct squeeze
     {
         return {std::move(output_shape), std::move(args.front().data)};
     }
+    int output_alias(const std::vector<shape>&) const
+    {
+        return 0;
+    }
 };
 
 struct unsqueeze
@@ -525,6 +541,10 @@ struct unsqueeze
     {
         return {std::move(output_shape), std::move(args.front().data)};
     }
+    int output_alias(const std::vector<shape>&) const
+    {
+        return 0;
+    }
 };
 
 struct reshape
@@ -576,6 +596,10 @@ struct reshape
     {
         return {std::move(output_shape), std::move(args.front().data)};
     }
+    int output_alias(const std::vector<shape>&) const
+    {
+        return 0;
+    }
 };
 
 struct dot
@@ -713,6 +737,10 @@ struct flatten
     {
         return {std::move(output_shape), std::move(args.front().data)};
     }
+    int output_alias(const std::vector<shape>&) const
+    {
+        return 0;
+    }
 };
 struct broadcast
 {
@@ -755,6 +783,10 @@ struct broadcast
     {
         return {std::move(output_shape), std::move(args.at(0).data)};
     }
+    int output_alias(const std::vector<shape>&) const
+    {
+        return 0;
+    }
 };
 
 struct scalar
@@ -776,6 +808,10 @@ struct scalar
     {
         return {std::move(output_shape), std::move(args.at(0).data)};
     }
+    int output_alias(const std::vector<shape>&) const
+    {
+        return 0;
+    }
 };
 
 struct binary
@@ -828,6 +864,10 @@ struct load
     {
         return {s, args[0].data() + offset};
     }
+    int output_alias(const std::vector<shape>&) const
+    {
+        return 0;
+    }
 };
 
 struct outline

src/program.cpp

@@ -280,7 +280,7 @@ void program::compile(const target& t, tracer trace)
 {
     assert(this->validate() == impl->instructions.end());
     this->impl->ctx = t.get_context();
-    if(not trace.enabled() and enabled(MIGRAPH_TRACE_COMPILE{}))
+    if(not trace.enabled() or enabled(MIGRAPH_TRACE_COMPILE{}))
         trace = tracer{std::cout};
     trace(*this);
     trace();

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

@@ -25,6 +25,7 @@ struct hip_add
     std::string name() const { return "gpu::add"; }
     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; }
 };
 
 struct miopen_add
@@ -33,6 +34,7 @@ struct miopen_add
     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

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

@@ -27,6 +27,7 @@ struct miopen_batch_norm_inference
     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

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

@@ -28,6 +28,7 @@ struct hip_concat
     shape compute_shape(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

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

@@ -26,6 +26,7 @@ struct miopen_contiguous
     std::string name() const { return "gpu::contiguous"; }
     shape compute_shape(const std::vector<shape>& inputs) const;
     argument compute(context&, shape output_shape, const std::vector<argument>& args) const;
+    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu

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

@@ -38,6 +38,7 @@ struct miopen_convolution
     argument
     compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;
     shape compile(context& ctx, const shape& output_shape, std::vector<instruction_ref> inputs);
+    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu

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

@@ -27,6 +27,7 @@ struct miopen_gemm
     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

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

@@ -67,6 +67,7 @@ struct hip_write
     {
         return to_gpu(args.front());
     }
+    int output_alias(const std::vector<shape>&) const { return 0; }
 };
 
 struct hip_copy
@@ -82,6 +83,7 @@ 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

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

@@ -27,6 +27,7 @@ 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

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

@@ -25,6 +25,7 @@ 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

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

@@ -29,6 +29,7 @@ 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

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

@@ -27,6 +27,7 @@ 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

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

@@ -27,6 +27,7 @@ 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