From 0f2ba3b05076f121c88ba935bda3ae007c8968e7 Mon Sep 17 00:00:00 2001 From: Paul Date: Mon, 12 Jan 2026 02:03:37 +0000 Subject: [PATCH 01/23] Return vector for output alias --- src/api/api.cpp | 14 ++------ src/graphviz.cpp | 2 +- src/include/migraphx/op/as_shape.hpp | 2 +- src/include/migraphx/op/broadcast.hpp | 2 +- src/include/migraphx/op/broadcast_for_dot.hpp | 2 +- src/include/migraphx/op/capture.hpp | 2 +- src/include/migraphx/op/fill.hpp | 2 +- src/include/migraphx/op/get_tuple_elem.hpp | 2 +- src/include/migraphx/op/identity.hpp | 2 +- src/include/migraphx/op/load.hpp | 2 +- src/include/migraphx/op/multibroadcast.hpp | 2 +- src/include/migraphx/op/random_uniform.hpp | 2 +- src/include/migraphx/op/reshape_lazy.hpp | 2 +- src/include/migraphx/op/scalar.hpp | 2 +- src/include/migraphx/op/select_module.hpp | 4 +-- src/include/migraphx/op/slice.hpp | 2 +- src/include/migraphx/op/squeeze.hpp | 2 +- src/include/migraphx/op/step.hpp | 2 +- src/include/migraphx/op/transpose.hpp | 2 +- src/include/migraphx/op/unsqueeze.hpp | 2 +- src/include/migraphx/operation.hpp | 33 +++++++++---------- src/instruction.cpp | 5 +-- src/targets/cpu/copy.cpp | 4 +-- src/targets/cpu/gather.cpp | 4 +-- src/targets/cpu/include/migraphx/cpu/dnnl.hpp | 4 +-- .../cpu/include/migraphx/cpu/pointwise.hpp | 8 ++--- src/targets/gpu/compile_miopen.cpp | 4 +-- src/targets/gpu/compile_ops.cpp | 4 +-- src/targets/gpu/fuse_ops.cpp | 12 +++---- src/targets/gpu/include/migraphx/gpu/abs.hpp | 4 +-- .../gpu/include/migraphx/gpu/argmax.hpp | 4 +-- .../gpu/include/migraphx/gpu/argmin.hpp | 4 +-- .../include/migraphx/gpu/code_object_op.hpp | 4 +-- .../migraphx/gpu/compile_hipblaslt.hpp | 4 +-- .../gpu/include/migraphx/gpu/convolution.hpp | 4 +-- .../gpu/include/migraphx/gpu/fixed_pad.hpp | 4 +-- src/targets/gpu/include/migraphx/gpu/gemm.hpp | 4 +-- src/targets/gpu/include/migraphx/gpu/hip.hpp | 22 ++++++------- .../gpu/include/migraphx/gpu/hip_gemm.hpp | 4 +-- .../gpu/include/migraphx/gpu/logsoftmax.hpp | 4 +-- src/targets/gpu/include/migraphx/gpu/loop.hpp | 4 +-- src/targets/gpu/include/migraphx/gpu/lrn.hpp | 4 +-- .../gpu/include/migraphx/gpu/multinomial.hpp | 4 +-- .../gpu/include/migraphx/gpu/nonzero.hpp | 4 +-- src/targets/gpu/include/migraphx/gpu/oper.hpp | 4 +-- .../gpu/include/migraphx/gpu/pooling.hpp | 4 +-- .../include/migraphx/gpu/prefix_scan_sum.hpp | 4 +-- .../gpu/include/migraphx/gpu/reduce_op.hpp | 4 +-- .../gpu/include/migraphx/gpu/reverse.hpp | 4 +-- .../migraphx/gpu/rnn_variable_seq_lens.hpp | 12 +++---- src/targets/gpu/include/migraphx/gpu/topk.hpp | 4 +-- test/api/test_custom_op.cpp | 11 +++---- test/eliminate_concat_test.cpp | 2 +- test/eval_test.cpp | 4 +-- test/include/basic_ops.hpp | 16 ++++++--- test/operation.cpp | 2 +- test/replace_allocate.cpp | 2 +- test/run_loop_test.cpp | 2 +- test/schedule_test.cpp | 2 +- tools/api/api.cpp | 14 ++------ tools/include/operation.hpp | 12 +++---- 61 files changed, 148 insertions(+), 165 deletions(-) diff --git a/src/api/api.cpp b/src/api/api.cpp index 5c37462641d..512343bd0fc 100644 --- a/src/api/api.cpp +++ b/src/api/api.cpp @@ -377,19 +377,9 @@ struct custom_operation return op.compute(std::move(ctx), std::move(output_shape), std::move(inputs)); } - std::ptrdiff_t output_alias(std::vector inputs) const + std::vector output_alias(std::vector inputs) const { - auto alias_vec = op.output_alias(std::move(inputs)); - // TODO: For now, only support one output alias - if(alias_vec.empty()) - { - return -1; - } - if(alias_vec.size() > 1) - { - MIGRAPHX_THROW("Currently, CustomOps in MIGraphX only supports one output_alias"); - } - return alias_vec.front(); + return op.output_alias(std::move(inputs)); } bool runs_on_offload_target() const { return op.runs_on_offload_target(); } diff --git a/src/graphviz.cpp b/src/graphviz.cpp index e2d50ac720e..3661adc87ea 100644 --- a/src/graphviz.cpp +++ b/src/graphviz.cpp @@ -114,7 +114,7 @@ std::string get_graph_color(const instruction_ref& ins) const auto& attr = op.attributes(); bool context_free = is_context_free(op); - bool alias = op.output_alias(to_shapes(ins->inputs())) >= 0; + bool alias = not op.output_alias(to_shapes(ins->inputs())).empty(); if(ins->can_eval()) { diff --git a/src/include/migraphx/op/as_shape.hpp b/src/include/migraphx/op/as_shape.hpp index 162526de16a..a6398e06e58 100644 --- a/src/include/migraphx/op/as_shape.hpp +++ b/src/include/migraphx/op/as_shape.hpp @@ -53,7 +53,7 @@ struct as_shape { return args.front().reshape(output_shape); } - std::ptrdiff_t output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; } // namespace op diff --git a/src/include/migraphx/op/broadcast.hpp b/src/include/migraphx/op/broadcast.hpp index 9587fedd6d2..fd31e73b34f 100644 --- a/src/include/migraphx/op/broadcast.hpp +++ b/src/include/migraphx/op/broadcast.hpp @@ -148,7 +148,7 @@ struct broadcast { return args[0].reshape(dyn_out.computed_shape); } - std::ptrdiff_t output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } value attributes() const { return {{"fillcolor", "#9ACD32" /* yellowgreen */}}; } }; diff --git a/src/include/migraphx/op/broadcast_for_dot.hpp b/src/include/migraphx/op/broadcast_for_dot.hpp index e3432ac0dac..40a1f21e618 100644 --- a/src/include/migraphx/op/broadcast_for_dot.hpp +++ b/src/include/migraphx/op/broadcast_for_dot.hpp @@ -83,7 +83,7 @@ struct broadcast_for_dot return args[0].reshape(dyn_out.computed_shape); } - std::ptrdiff_t output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } value attributes() const { return {{"fillcolor", "#9ACD32" /* yellowgreen */}}; } }; diff --git a/src/include/migraphx/op/capture.hpp b/src/include/migraphx/op/capture.hpp index efb71c5c349..8a1c33a2c68 100644 --- a/src/include/migraphx/op/capture.hpp +++ b/src/include/migraphx/op/capture.hpp @@ -66,7 +66,7 @@ struct capture return args.front(); } - std::ptrdiff_t output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; } // namespace op diff --git a/src/include/migraphx/op/fill.hpp b/src/include/migraphx/op/fill.hpp index 38d608b9183..80279b5b6d9 100644 --- a/src/include/migraphx/op/fill.hpp +++ b/src/include/migraphx/op/fill.hpp @@ -60,7 +60,7 @@ struct fill return args[1]; } - std::ptrdiff_t output_alias(const std::vector&) const { return 1; } + std::vector output_alias(const std::vector&) const { return {1}; } }; } // namespace op diff --git a/src/include/migraphx/op/get_tuple_elem.hpp b/src/include/migraphx/op/get_tuple_elem.hpp index e8e24d11c53..ca1d14246a0 100644 --- a/src/include/migraphx/op/get_tuple_elem.hpp +++ b/src/include/migraphx/op/get_tuple_elem.hpp @@ -67,7 +67,7 @@ struct get_tuple_elem return vec_args.at(index); } - std::ptrdiff_t output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; } // namespace op diff --git a/src/include/migraphx/op/identity.hpp b/src/include/migraphx/op/identity.hpp index 9d80c5bfc9c..b76397e2778 100644 --- a/src/include/migraphx/op/identity.hpp +++ b/src/include/migraphx/op/identity.hpp @@ -39,7 +39,7 @@ struct identity value attributes() const { return {{"pointwise", true}, {"point_op", "${0}"}}; } - std::ptrdiff_t output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; } // namespace op diff --git a/src/include/migraphx/op/load.hpp b/src/include/migraphx/op/load.hpp index 0050e715089..e85e7fa8471 100644 --- a/src/include/migraphx/op/load.hpp +++ b/src/include/migraphx/op/load.hpp @@ -58,7 +58,7 @@ struct load return argument{s, args[0].data() + offset}; } lifetime get_lifetime() const { return lifetime::borrow; } - std::ptrdiff_t output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } friend std::ostream& operator<<(std::ostream& os, const load& op) { diff --git a/src/include/migraphx/op/multibroadcast.hpp b/src/include/migraphx/op/multibroadcast.hpp index 8a9d15c7d99..b1076661171 100644 --- a/src/include/migraphx/op/multibroadcast.hpp +++ b/src/include/migraphx/op/multibroadcast.hpp @@ -115,7 +115,7 @@ struct multibroadcast { return args[0].reshape(dyn_out.computed_shape); } - std::ptrdiff_t output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; } // namespace op diff --git a/src/include/migraphx/op/random_uniform.hpp b/src/include/migraphx/op/random_uniform.hpp index f873ae9d313..642d35ff082 100644 --- a/src/include/migraphx/op/random_uniform.hpp +++ b/src/include/migraphx/op/random_uniform.hpp @@ -107,7 +107,7 @@ struct random_uniform return result; } - std::ptrdiff_t output_alias(const std::vector&) const { return 1; } + std::vector output_alias(const std::vector&) const { return {1}; } }; } // namespace op diff --git a/src/include/migraphx/op/reshape_lazy.hpp b/src/include/migraphx/op/reshape_lazy.hpp index 4746015053e..6aab475140d 100644 --- a/src/include/migraphx/op/reshape_lazy.hpp +++ b/src/include/migraphx/op/reshape_lazy.hpp @@ -316,7 +316,7 @@ struct reshape_lazy return args[0].reshape(dyn_out.computed_shape); } - std::ptrdiff_t output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; } // namespace op diff --git a/src/include/migraphx/op/scalar.hpp b/src/include/migraphx/op/scalar.hpp index 9102e476d49..852aca4e551 100644 --- a/src/include/migraphx/op/scalar.hpp +++ b/src/include/migraphx/op/scalar.hpp @@ -56,7 +56,7 @@ struct scalar { return args[0].reshape(output_shape); } - std::ptrdiff_t output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; } // namespace op diff --git a/src/include/migraphx/op/select_module.hpp b/src/include/migraphx/op/select_module.hpp index a8064e80cfa..a16336413bb 100644 --- a/src/include/migraphx/op/select_module.hpp +++ b/src/include/migraphx/op/select_module.hpp @@ -138,9 +138,9 @@ struct select_module return argument{results}; } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; diff --git a/src/include/migraphx/op/slice.hpp b/src/include/migraphx/op/slice.hpp index 88c05269af9..5dd4eaf55c1 100644 --- a/src/include/migraphx/op/slice.hpp +++ b/src/include/migraphx/op/slice.hpp @@ -502,7 +502,7 @@ struct slice } } - std::ptrdiff_t output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; } // namespace op diff --git a/src/include/migraphx/op/squeeze.hpp b/src/include/migraphx/op/squeeze.hpp index 6802e352feb..829a4a7bcfe 100644 --- a/src/include/migraphx/op/squeeze.hpp +++ b/src/include/migraphx/op/squeeze.hpp @@ -141,7 +141,7 @@ struct squeeze { return args[0].reshape(dyn_out.computed_shape); } - std::ptrdiff_t output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; } // namespace op diff --git a/src/include/migraphx/op/step.hpp b/src/include/migraphx/op/step.hpp index d547ac9cffe..195402e877e 100644 --- a/src/include/migraphx/op/step.hpp +++ b/src/include/migraphx/op/step.hpp @@ -90,7 +90,7 @@ struct step return args[0].reshape(output_shape); } - std::ptrdiff_t output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; } // namespace op diff --git a/src/include/migraphx/op/transpose.hpp b/src/include/migraphx/op/transpose.hpp index 039bea6e4c2..60f7c37c213 100644 --- a/src/include/migraphx/op/transpose.hpp +++ b/src/include/migraphx/op/transpose.hpp @@ -92,7 +92,7 @@ struct transpose return args[0].reshape(dyn_out.computed_shape); } - std::ptrdiff_t output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; } // namespace op diff --git a/src/include/migraphx/op/unsqueeze.hpp b/src/include/migraphx/op/unsqueeze.hpp index d62c3250b0b..ca132317b97 100644 --- a/src/include/migraphx/op/unsqueeze.hpp +++ b/src/include/migraphx/op/unsqueeze.hpp @@ -148,7 +148,7 @@ struct unsqueeze { return args[0].reshape(dyn_out.computed_shape); } - std::ptrdiff_t output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; } // namespace op diff --git a/src/include/migraphx/operation.hpp b/src/include/migraphx/operation.hpp index 86fbe6ecd98..ac22cd6a425 100644 --- a/src/include/migraphx/operation.hpp +++ b/src/include/migraphx/operation.hpp @@ -79,9 +79,9 @@ struct operation * the same the `output` shape. */ argument compute(context& ctx, const shape& output, const std::vector& input) const; - /// An optional method to return which argument the output will alias. If - /// there is no aliased output then -1 can be returned. - std::ptrdiff_t output_alias(const std::vector& input) const; + /// An optional method to return which arguments the output will alias. If + /// there is no aliased output then an empty vector can be returned. + std::vector output_alias(const std::vector& input) const; /// An optional stream operator to print the operation. When this is not /// implemented, it will just print the operation's name. friend std::ostream& operator<<(std::ostream& os, const operation& op); @@ -407,9 +407,9 @@ auto need_normalization_op(const T& x) } template -std::ptrdiff_t output_alias_op(const T&, const std::vector&) +std::vector output_alias_op(const T&, const std::vector&) { - return -1; + return {}; } template @@ -517,7 +517,7 @@ struct MIGRAPHX_EXPORT operation // (optional) lifetime get_lifetime() const; // (optional) - std::ptrdiff_t output_alias(const std::vector& input) const; + std::vector output_alias(const std::vector& input) const; // (optional) value compile(context& ctx, const shape& output, const std::vector& input); // (optional) @@ -623,9 +623,8 @@ struct operation } template - static std::ptrdiff_t private_detail_te_default_output_alias(float, - T&& private_detail_te_self, - const std::vector& input) + static std::vector private_detail_te_default_output_alias( + float, T&& private_detail_te_self, const std::vector& input) { return detail::output_alias_op(private_detail_te_self, input); } @@ -1030,7 +1029,7 @@ struct operation return (*this).private_detail_te_get_handle().get_lifetime(); } - std::ptrdiff_t output_alias(const std::vector& input) const + std::vector output_alias(const std::vector& input) const { assert((*this).private_detail_te_handle_mem_var); return (*this).private_detail_te_get_handle().output_alias(input); @@ -1139,12 +1138,12 @@ struct operation virtual std::shared_ptr clone() const = 0; virtual const std::type_info& type() const = 0; - virtual std::string name() const = 0; - virtual bool is_context_free() const = 0; - virtual bool need_normalization() const = 0; - virtual bool has_finalize() const = 0; - virtual lifetime get_lifetime() const = 0; - virtual std::ptrdiff_t output_alias(const std::vector& input) const = 0; + virtual std::string name() const = 0; + virtual bool is_context_free() const = 0; + virtual bool need_normalization() const = 0; + virtual bool has_finalize() const = 0; + virtual lifetime get_lifetime() const = 0; + virtual std::vector output_alias(const std::vector& input) const = 0; virtual value compile(context& ctx, const shape& output, const std::vector& input) = 0; virtual void @@ -1229,7 +1228,7 @@ struct operation return private_detail_te_default_get_lifetime(char(0), private_detail_te_value); } - std::ptrdiff_t output_alias(const std::vector& input) const override + std::vector output_alias(const std::vector& input) const override { return private_detail_te_default_output_alias(char(0), private_detail_te_value, input); diff --git a/src/instruction.cpp b/src/instruction.cpp index b8c68a42ce9..5d117c9c1e1 100644 --- a/src/instruction.cpp +++ b/src/instruction.cpp @@ -472,9 +472,10 @@ void instruction::debug_print() const instruction_ref instruction::get_output_alias(instruction_ref ins, bool shallow) { - auto i = ins->get_operator().output_alias(to_shapes(ins->inputs())); - if(i < 0) + auto aliases = ins->get_operator().output_alias(to_shapes(ins->inputs())); + if(aliases.empty()) return ins; + auto i = aliases.front(); if(shallow) return ins->inputs().at(i); return get_output_alias(ins->inputs().at(i)); diff --git a/src/targets/cpu/copy.cpp b/src/targets/cpu/copy.cpp index 4c4af2b7164..03839f20e5c 100644 --- a/src/targets/cpu/copy.cpp +++ b/src/targets/cpu/copy.cpp @@ -54,9 +54,9 @@ struct cpu_copy : reduce_dims_base, auto_register_op return result.reshape(output_shape); } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; diff --git a/src/targets/cpu/gather.cpp b/src/targets/cpu/gather.cpp index 40bc556b961..36a4c77e02c 100644 --- a/src/targets/cpu/gather.cpp +++ b/src/targets/cpu/gather.cpp @@ -77,9 +77,9 @@ struct cpu_gather : auto_register_op return args.back(); } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; diff --git a/src/targets/cpu/include/migraphx/cpu/dnnl.hpp b/src/targets/cpu/include/migraphx/cpu/dnnl.hpp index b05cad85246..d4940769207 100644 --- a/src/targets/cpu/include/migraphx/cpu/dnnl.hpp +++ b/src/targets/cpu/include/migraphx/cpu/dnnl.hpp @@ -296,9 +296,9 @@ struct dnnl_op : auto_register_op return execute(ctx, args); } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } value compile(context&, const shape& output_shape, std::vector inputs) { diff --git a/src/targets/cpu/include/migraphx/cpu/pointwise.hpp b/src/targets/cpu/include/migraphx/cpu/pointwise.hpp index ece5498c839..4e93f10fe51 100644 --- a/src/targets/cpu/include/migraphx/cpu/pointwise.hpp +++ b/src/targets/cpu/include/migraphx/cpu/pointwise.hpp @@ -360,9 +360,9 @@ struct cpu_unary : reduce_dims_base, auto_register_op> return result.reshape(output_shape); } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; @@ -401,9 +401,9 @@ struct cpu_binary : reduce_dims_base, auto_register_op> return result.reshape(output_shape); } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; diff --git a/src/targets/gpu/compile_miopen.cpp b/src/targets/gpu/compile_miopen.cpp index 583601bdda1..6bb7fa2ccb7 100644 --- a/src/targets/gpu/compile_miopen.cpp +++ b/src/targets/gpu/compile_miopen.cpp @@ -52,9 +52,9 @@ struct miopen_op return op.compute_shape(inputs); } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; MIGRAPHX_REGISTER_OP(miopen_op); diff --git a/src/targets/gpu/compile_ops.cpp b/src/targets/gpu/compile_ops.cpp index 50c252ac08f..fb34e56b746 100644 --- a/src/targets/gpu/compile_ops.cpp +++ b/src/targets/gpu/compile_ops.cpp @@ -72,9 +72,9 @@ struct precompile_op return op.compute_shape(inputs, mods); } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; diff --git a/src/targets/gpu/fuse_ops.cpp b/src/targets/gpu/fuse_ops.cpp index 4d04394ce69..55684e9f07a 100644 --- a/src/targets/gpu/fuse_ops.cpp +++ b/src/targets/gpu/fuse_ops.cpp @@ -251,9 +251,9 @@ struct miopen_fusion return pack(f(self.ops, "ops")); } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } value compile(context& ctx, const shape&, std::vector inputs) @@ -383,9 +383,9 @@ struct miopen_conv_bias } shape get_workspace(context& ctx) { return fp.get_workspace(ctx); } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; MIGRAPHX_REGISTER_OP(miopen_conv_bias) @@ -431,9 +431,9 @@ struct miopen_conv_bias_relu } shape get_workspace(context& ctx) { return fp.get_workspace(ctx); } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; MIGRAPHX_REGISTER_OP(miopen_conv_bias_relu) diff --git a/src/targets/gpu/include/migraphx/gpu/abs.hpp b/src/targets/gpu/include/migraphx/gpu/abs.hpp index 1a9f4b87800..f8fc5755d25 100644 --- a/src/targets/gpu/include/migraphx/gpu/abs.hpp +++ b/src/targets/gpu/include/migraphx/gpu/abs.hpp @@ -52,9 +52,9 @@ struct miopen_abs argument compute(context& ctx, const shape& output_shape, const std::vector& args) const; void finalize(context&, const shape&, const std::vector&); - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; #endif diff --git a/src/targets/gpu/include/migraphx/gpu/argmax.hpp b/src/targets/gpu/include/migraphx/gpu/argmax.hpp index e05678fa873..a3750e8c7bb 100644 --- a/src/targets/gpu/include/migraphx/gpu/argmax.hpp +++ b/src/targets/gpu/include/migraphx/gpu/argmax.hpp @@ -48,9 +48,9 @@ struct hip_argmax std::string name() const { return "gpu::argmax"; } shape compute_shape(const std::vector& inputs) const; argument compute(context& ctx, const shape&, const std::vector& args) const; - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; diff --git a/src/targets/gpu/include/migraphx/gpu/argmin.hpp b/src/targets/gpu/include/migraphx/gpu/argmin.hpp index 071eb525e7f..0afd956965f 100644 --- a/src/targets/gpu/include/migraphx/gpu/argmin.hpp +++ b/src/targets/gpu/include/migraphx/gpu/argmin.hpp @@ -48,9 +48,9 @@ struct hip_argmin std::string name() const { return "gpu::argmin"; } shape compute_shape(const std::vector& inputs) const; argument compute(context& ctx, const shape&, const std::vector& args) const; - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; diff --git a/src/targets/gpu/include/migraphx/gpu/code_object_op.hpp b/src/targets/gpu/include/migraphx/gpu/code_object_op.hpp index 8186767289f..758c7a4c07b 100644 --- a/src/targets/gpu/include/migraphx/gpu/code_object_op.hpp +++ b/src/targets/gpu/include/migraphx/gpu/code_object_op.hpp @@ -72,9 +72,9 @@ struct code_object_op { return output_arg < 0 ? n + output_arg : output_arg; } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return get_output_arg(shapes.size()); + return {static_cast(get_output_arg(shapes.size()))}; } friend std::ostream& operator<<(std::ostream& os, const code_object_op& op) diff --git a/src/targets/gpu/include/migraphx/gpu/compile_hipblaslt.hpp b/src/targets/gpu/include/migraphx/gpu/compile_hipblaslt.hpp index 380fafa44ba..8c64fffc901 100644 --- a/src/targets/gpu/include/migraphx/gpu/compile_hipblaslt.hpp +++ b/src/targets/gpu/include/migraphx/gpu/compile_hipblaslt.hpp @@ -57,9 +57,9 @@ struct hipblaslt_op return op.compute_shape(inputs); } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; MIGRAPHX_REGISTER_OP(hipblaslt_op); diff --git a/src/targets/gpu/include/migraphx/gpu/convolution.hpp b/src/targets/gpu/include/migraphx/gpu/convolution.hpp index 1bca5573f71..25897f93d95 100644 --- a/src/targets/gpu/include/migraphx/gpu/convolution.hpp +++ b/src/targets/gpu/include/migraphx/gpu/convolution.hpp @@ -339,9 +339,9 @@ struct miopen_convolution #endif } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; #endif diff --git a/src/targets/gpu/include/migraphx/gpu/fixed_pad.hpp b/src/targets/gpu/include/migraphx/gpu/fixed_pad.hpp index ad78787c062..96ed5517462 100644 --- a/src/targets/gpu/include/migraphx/gpu/fixed_pad.hpp +++ b/src/targets/gpu/include/migraphx/gpu/fixed_pad.hpp @@ -48,9 +48,9 @@ struct hip_fixed_pad std::string name() const { return "gpu::fixed_pad"; } shape compute_shape(std::vector inputs) const; argument compute(context& ctx, const shape&, const std::vector& args) const; - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; diff --git a/src/targets/gpu/include/migraphx/gpu/gemm.hpp b/src/targets/gpu/include/migraphx/gpu/gemm.hpp index 332de8bf090..82f716726b8 100644 --- a/src/targets/gpu/include/migraphx/gpu/gemm.hpp +++ b/src/targets/gpu/include/migraphx/gpu/gemm.hpp @@ -121,9 +121,9 @@ struct rocblas_gemm return args.back(); } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } void finalize(context& ctx, const shape& output_shape, const std::vector& input_shapes) diff --git a/src/targets/gpu/include/migraphx/gpu/hip.hpp b/src/targets/gpu/include/migraphx/gpu/hip.hpp index 055594252c2..7869f84afbb 100644 --- a/src/targets/gpu/include/migraphx/gpu/hip.hpp +++ b/src/targets/gpu/include/migraphx/gpu/hip.hpp @@ -113,7 +113,7 @@ struct hip_fill gpu_fill(ctx, args.front(), value); return args.front(); } - std::ptrdiff_t output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; struct hip_sync_stream @@ -135,11 +135,11 @@ struct hip_sync_stream return args.front(); } - std::ptrdiff_t output_alias(const std::vector& args) const + std::vector output_alias(const std::vector& args) const { if(args.empty()) - return -1; - return 0; + return {}; + return {0}; } }; @@ -165,11 +165,11 @@ struct hip_copy_to_gpu // Associate the input since it was registered with hip return {result.get_shape(), [input, result]() mutable { return result.data(); }}; } - std::ptrdiff_t output_alias(const std::vector& args) const + std::vector output_alias(const std::vector& args) const { if(args.size() == 1) - return -1; - return 1; + return {}; + return {1}; } }; @@ -198,11 +198,11 @@ struct hip_copy_from_gpu copy_from_gpu(ctx, input, args[1]); return args[1]; } - std::ptrdiff_t output_alias(const std::vector& args) const + std::vector output_alias(const std::vector& args) const { if(args.size() == 1) - return -1; - return 1; + return {}; + return {1}; } }; @@ -224,7 +224,7 @@ struct hip_copy gpu_copy(ctx, args[0], result); return args[1]; } - std::ptrdiff_t output_alias(const std::vector&) const { return 1; } + std::vector output_alias(const std::vector&) const { return {1}; } }; MIGRAPHX_GPU_EXPORT void diff --git a/src/targets/gpu/include/migraphx/gpu/hip_gemm.hpp b/src/targets/gpu/include/migraphx/gpu/hip_gemm.hpp index e87ce9edcee..ae5151525c0 100644 --- a/src/targets/gpu/include/migraphx/gpu/hip_gemm.hpp +++ b/src/targets/gpu/include/migraphx/gpu/hip_gemm.hpp @@ -113,9 +113,9 @@ struct MIGRAPHX_GPU_EXPORT hip_gemm return args.back(); } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } void finalize(context& ctx, const shape& output_shape, const std::vector& input_shapes) diff --git a/src/targets/gpu/include/migraphx/gpu/logsoftmax.hpp b/src/targets/gpu/include/migraphx/gpu/logsoftmax.hpp index 5ea23ee2724..e7036aa695a 100644 --- a/src/targets/gpu/include/migraphx/gpu/logsoftmax.hpp +++ b/src/targets/gpu/include/migraphx/gpu/logsoftmax.hpp @@ -47,9 +47,9 @@ struct hip_logsoftmax shape compute_shape(const std::vector& inputs) const; argument compute(context& ctx, const shape& output_shape, const std::vector& args) const; - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; diff --git a/src/targets/gpu/include/migraphx/gpu/loop.hpp b/src/targets/gpu/include/migraphx/gpu/loop.hpp index 792c84b74f8..02b8681325d 100644 --- a/src/targets/gpu/include/migraphx/gpu/loop.hpp +++ b/src/targets/gpu/include/migraphx/gpu/loop.hpp @@ -53,9 +53,9 @@ struct hip_loop const std::vector& mods, const std::function( module_ref&, const std::unordered_map&)>& run) const; - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; diff --git a/src/targets/gpu/include/migraphx/gpu/lrn.hpp b/src/targets/gpu/include/migraphx/gpu/lrn.hpp index 8ccda7bba6a..19691a10704 100644 --- a/src/targets/gpu/include/migraphx/gpu/lrn.hpp +++ b/src/targets/gpu/include/migraphx/gpu/lrn.hpp @@ -50,9 +50,9 @@ struct miopen_lrn argument compute(context& ctx, const shape& output_shape, const std::vector& args) const; void finalize(context&, const shape&, const std::vector&); - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; #endif diff --git a/src/targets/gpu/include/migraphx/gpu/multinomial.hpp b/src/targets/gpu/include/migraphx/gpu/multinomial.hpp index c44d4808233..2be31fc35da 100644 --- a/src/targets/gpu/include/migraphx/gpu/multinomial.hpp +++ b/src/targets/gpu/include/migraphx/gpu/multinomial.hpp @@ -46,9 +46,9 @@ struct hip_multinomial shape compute_shape(std::vector inputs) const; argument compute(context& ctx, const shape& output_shape, const std::vector& args) const; - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; diff --git a/src/targets/gpu/include/migraphx/gpu/nonzero.hpp b/src/targets/gpu/include/migraphx/gpu/nonzero.hpp index cfc7e78dbe3..89f4c7c7885 100644 --- a/src/targets/gpu/include/migraphx/gpu/nonzero.hpp +++ b/src/targets/gpu/include/migraphx/gpu/nonzero.hpp @@ -49,9 +49,9 @@ struct hip_nonzero shape compute_shape(std::vector inputs) const; argument compute(context& ctx, const shape& output_shape, const std::vector& args) const; - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; diff --git a/src/targets/gpu/include/migraphx/gpu/oper.hpp b/src/targets/gpu/include/migraphx/gpu/oper.hpp index 13ac11a3d48..760494ba5ff 100644 --- a/src/targets/gpu/include/migraphx/gpu/oper.hpp +++ b/src/targets/gpu/include/migraphx/gpu/oper.hpp @@ -72,9 +72,9 @@ struct device_base : oper return {s0.type(), s0.lens()}; } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; diff --git a/src/targets/gpu/include/migraphx/gpu/pooling.hpp b/src/targets/gpu/include/migraphx/gpu/pooling.hpp index 7f6722b1130..b3ea6f22d8b 100644 --- a/src/targets/gpu/include/migraphx/gpu/pooling.hpp +++ b/src/targets/gpu/include/migraphx/gpu/pooling.hpp @@ -50,9 +50,9 @@ struct miopen_pooling void finalize(context&, const shape&, const std::vector&); argument compute(context& ctx, const shape& output_shape, const std::vector& args) const; - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; #endif diff --git a/src/targets/gpu/include/migraphx/gpu/prefix_scan_sum.hpp b/src/targets/gpu/include/migraphx/gpu/prefix_scan_sum.hpp index cca8efd6057..d26701897ce 100644 --- a/src/targets/gpu/include/migraphx/gpu/prefix_scan_sum.hpp +++ b/src/targets/gpu/include/migraphx/gpu/prefix_scan_sum.hpp @@ -67,9 +67,9 @@ struct hip_prefix_scan_sum : oper return args[1]; } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; diff --git a/src/targets/gpu/include/migraphx/gpu/reduce_op.hpp b/src/targets/gpu/include/migraphx/gpu/reduce_op.hpp index 10f3dcf8445..b52d70ee240 100644 --- a/src/targets/gpu/include/migraphx/gpu/reduce_op.hpp +++ b/src/targets/gpu/include/migraphx/gpu/reduce_op.hpp @@ -65,9 +65,9 @@ struct reduce_op : oper return args[1]; } - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } reduce_op() {} diff --git a/src/targets/gpu/include/migraphx/gpu/reverse.hpp b/src/targets/gpu/include/migraphx/gpu/reverse.hpp index 8ef82523526..1fb951f2c93 100644 --- a/src/targets/gpu/include/migraphx/gpu/reverse.hpp +++ b/src/targets/gpu/include/migraphx/gpu/reverse.hpp @@ -49,9 +49,9 @@ struct hip_reverse shape compute_shape(std::vector inputs) const; argument compute(context& ctx, const shape& output_shape, const std::vector& args) const; - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; diff --git a/src/targets/gpu/include/migraphx/gpu/rnn_variable_seq_lens.hpp b/src/targets/gpu/include/migraphx/gpu/rnn_variable_seq_lens.hpp index 7d811192da0..36778bd79dd 100644 --- a/src/targets/gpu/include/migraphx/gpu/rnn_variable_seq_lens.hpp +++ b/src/targets/gpu/include/migraphx/gpu/rnn_variable_seq_lens.hpp @@ -49,9 +49,9 @@ struct hip_rnn_var_sl_shift_sequence shape compute_shape(std::vector inputs) const; argument compute(context& ctx, const shape& output_shape, const std::vector& args) const; - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; @@ -69,9 +69,9 @@ struct hip_rnn_var_sl_shift_output shape compute_shape(std::vector inputs) const; argument compute(context& ctx, const shape& output_shape, const std::vector& args) const; - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; @@ -88,9 +88,9 @@ struct hip_rnn_var_sl_last_output std::string name() const { return "gpu::" + op.name(); } shape compute_shape(std::vector inputs) const; argument compute(context& ctx, const shape&, const std::vector& args) const; - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; diff --git a/src/targets/gpu/include/migraphx/gpu/topk.hpp b/src/targets/gpu/include/migraphx/gpu/topk.hpp index f1df9d469e5..5f8c70f3981 100644 --- a/src/targets/gpu/include/migraphx/gpu/topk.hpp +++ b/src/targets/gpu/include/migraphx/gpu/topk.hpp @@ -49,9 +49,9 @@ struct hip_topk shape compute_shape(std::vector inputs) const; argument compute(context& ctx, const shape& output_shape, const std::vector& args) const; - std::ptrdiff_t output_alias(const std::vector& shapes) const + std::vector output_alias(const std::vector& shapes) const { - return shapes.size() - 1; + return {shapes.size() - 1}; } }; diff --git a/test/api/test_custom_op.cpp b/test/api/test_custom_op.cpp index b4d4c0da0c3..9e7f6042760 100644 --- a/test/api/test_custom_op.cpp +++ b/test/api/test_custom_op.cpp @@ -133,10 +133,10 @@ struct identity_custom_op final : migraphx::experimental_custom_op_base return inputs.back(); } - virtual std::vector output_alias(migraphx::shapes) const override { return {0, 1}; } + virtual std::vector output_alias(migraphx::shapes) const override { return {0}; } }; -TEST_CASE(run_custom_op_with_invalid_output_alias) +TEST_CASE(run_custom_op_with_output_alias) { identity_custom_op i_op; migraphx::register_experimental_custom_op(i_op); @@ -147,11 +147,8 @@ TEST_CASE(run_custom_op_with_invalid_output_alias) migraphx::shape s{migraphx_shape_float_type, {12}}; migraphx::module m = p.get_main_module(); auto x = m.add_parameter("x", s); - auto i_ins = m.add_instruction(migraphx::operation("identity_custom_op"), {x}); - migraphx_test_private_disable_exception_catch(true); - EXPECT(test::throws( - [&] { p.compile(migraphx::target("ref")); }, - "Currently, CustomOps in MIGraphX only supports one output_alias")); + m.add_instruction(migraphx::operation("identity_custom_op"), {x}); + p.compile(migraphx::target("ref")); } int main(int argc, const char* argv[]) { test::run(argc, argv); } diff --git a/test/eliminate_concat_test.cpp b/test/eliminate_concat_test.cpp index 4fe59c0bac7..0f2138668b3 100644 --- a/test/eliminate_concat_test.cpp +++ b/test/eliminate_concat_test.cpp @@ -123,7 +123,7 @@ struct simple_op { return args.at(0); } - int output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; template diff --git a/test/eval_test.cpp b/test/eval_test.cpp index 0435cf61634..bd7f9463d72 100644 --- a/test/eval_test.cpp +++ b/test/eval_test.cpp @@ -65,7 +65,7 @@ struct id_ctx_op return {}; return inputs.front(); } - int output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; struct id_ctx_final_op @@ -88,7 +88,7 @@ struct id_ctx_final_op return {}; return inputs.front(); } - int output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; struct reverse_pass diff --git a/test/include/basic_ops.hpp b/test/include/basic_ops.hpp index fa57ec8b1ad..5a6acaf1e59 100644 --- a/test/include/basic_ops.hpp +++ b/test/include/basic_ops.hpp @@ -131,7 +131,10 @@ struct pass_op return {}; return inputs.front(); } - int output_alias(const std::vector& s) const { return s.empty() ? -1 : 0; } + std::vector output_alias(const std::vector& s) const + { + return s.empty() ? std::vector{} : std::vector{0}; + } }; struct non_const_pass_op @@ -151,7 +154,10 @@ struct non_const_pass_op return {}; return inputs.front(); } - int output_alias(const std::vector& s) const { return s.empty() ? -1 : 0; } + std::vector output_alias(const std::vector& s) const + { + return s.empty() ? std::vector{} : std::vector{0}; + } }; struct mod_pass_op @@ -176,7 +182,7 @@ struct mod_pass_op return {}; } - int output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; struct unary_pass_op @@ -196,7 +202,7 @@ struct unary_pass_op MIGRAPHX_THROW("Wrong inputs"); return inputs.front(); } - int output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; struct pass_standard_op @@ -221,7 +227,7 @@ struct pass_standard_op return {}; return inputs.front(); } - int output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; struct nop diff --git a/test/operation.cpp b/test/operation.cpp index 28b4f31f32c..ed17109edb1 100644 --- a/test/operation.cpp +++ b/test/operation.cpp @@ -87,7 +87,7 @@ struct compilable_op return inputs.front(); } - int output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } migraphx::value compile(migraphx::context&, const migraphx::shape&, const std::vector&) diff --git a/test/replace_allocate.cpp b/test/replace_allocate.cpp index 937920b9a5e..dc0eba17073 100644 --- a/test/replace_allocate.cpp +++ b/test/replace_allocate.cpp @@ -52,7 +52,7 @@ struct test_copy : migraphx::auto_register_op return inputs.back(); } - std::ptrdiff_t output_alias(const std::vector&) const { return 1; } + std::vector output_alias(const std::vector&) const { return {1}; } }; struct allocate_no_out : migraphx::auto_register_op diff --git a/test/run_loop_test.cpp b/test/run_loop_test.cpp index c371ef59c63..8c1238bb031 100644 --- a/test/run_loop_test.cpp +++ b/test/run_loop_test.cpp @@ -57,7 +57,7 @@ struct copy_op return args[1]; } - int output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; struct test_loop_op diff --git a/test/schedule_test.cpp b/test/schedule_test.cpp index 252d0e1c3cc..69edd4179fc 100644 --- a/test/schedule_test.cpp +++ b/test/schedule_test.cpp @@ -50,7 +50,7 @@ struct unary_op return {}; return inputs.front(); } - int output_alias(const std::vector&) const { return 0; } + std::vector output_alias(const std::vector&) const { return {0}; } }; struct nary_op diff --git a/tools/api/api.cpp b/tools/api/api.cpp index 0526175ddf3..0bdda69f47c 100644 --- a/tools/api/api.cpp +++ b/tools/api/api.cpp @@ -377,19 +377,9 @@ struct custom_operation return op.compute(std::move(ctx), std::move(output_shape), std::move(inputs)); } - std::ptrdiff_t output_alias(std::vector inputs) const + std::vector output_alias(std::vector inputs) const { - auto alias_vec = op.output_alias(std::move(inputs)); - // TODO: For now, only support one output alias - if(alias_vec.empty()) - { - return -1; - } - if(alias_vec.size() > 1) - { - MIGRAPHX_THROW("Currently, CustomOps in MIGraphX only supports one output_alias"); - } - return alias_vec.front(); + return op.output_alias(std::move(inputs)); } bool runs_on_offload_target() const { return op.runs_on_offload_target(); } diff --git a/tools/include/operation.hpp b/tools/include/operation.hpp index 769cc8f9be7..00abf231935 100644 --- a/tools/include/operation.hpp +++ b/tools/include/operation.hpp @@ -79,9 +79,9 @@ struct operation * the same the `output` shape. */ argument compute(context& ctx, const shape& output, const std::vector& input) const; - /// An optional method to return which argument the output will alias. If - /// there is no aliased output then -1 can be returned. - std::ptrdiff_t output_alias(const std::vector& input) const; + /// An optional method to return which arguments the output will alias. If + /// there is no aliased output then an empty vector can be returned. + std::vector output_alias(const std::vector& input) const; /// An optional stream operator to print the operation. When this is not /// implemented, it will just print the operation's name. friend std::ostream& operator<<(std::ostream& os, const operation& op); @@ -410,9 +410,9 @@ auto need_normalization_op(const T& x) } template -std::ptrdiff_t output_alias_op(const T&, const std::vector&) +std::vector output_alias_op(const T&, const std::vector&) { - return -1; + return {}; } template @@ -518,7 +518,7 @@ lifetime get_lifetime_op(const T&) virtual( 'get_lifetime', returns = 'lifetime', const = True, default = 'detail::get_lifetime_op'), virtual('output_alias', - returns = 'std::ptrdiff_t', + returns = 'std::vector', input = 'const std::vector&', const = True, default = 'detail::output_alias_op'), From 9f17a2d3ce6d483544956fa3ef447e3b4c653d47 Mon Sep 17 00:00:00 2001 From: Paul Date: Mon, 12 Jan 2026 02:03:41 +0000 Subject: [PATCH 02/23] Format --- tools/include/operation.hpp | 178 +++++++++++++++++++----------------- 1 file changed, 92 insertions(+), 86 deletions(-) diff --git a/tools/include/operation.hpp b/tools/include/operation.hpp index 00abf231935..6269a8a038b 100644 --- a/tools/include/operation.hpp +++ b/tools/include/operation.hpp @@ -505,92 +505,98 @@ lifetime get_lifetime_op(const T&) } // namespace detail <% - interface( - 'operation', - virtual('name', returns = 'std::string', const = True), - virtual( - 'is_context_free', returns = 'bool', const = True, default = 'detail::is_context_free_op'), - virtual('need_normalization', - returns = 'bool', - const = True, - default = 'detail::need_normalization_op'), - virtual('has_finalize', returns = 'bool', const = True, default = 'detail::has_finalize_op'), - virtual( - 'get_lifetime', returns = 'lifetime', const = True, default = 'detail::get_lifetime_op'), - virtual('output_alias', - returns = 'std::vector', - input = 'const std::vector&', - const = True, - default = 'detail::output_alias_op'), - virtual('compile', - returns = 'value', - ctx = 'context&', - output = 'const shape&', - input = 'const std::vector&', - default = 'detail::compile_op'), - virtual('finalize', - ctx = 'context&', - output = 'const shape&', - input = 'const std::vector&', - default = 'detail::finalize_op'), - virtual('compute_shape', - returns = 'shape', - input = 'const std::vector&', - const = True, - default = 'detail::compute_shape_op'), - virtual('compute_shape', - returns = 'shape', - inputs = 'const std::vector&', - mod_args = 'const std::vector&', - const = True, - default = 'detail::mod_compute_shape_op'), - virtual('compute', - returns = 'argument', - ctx = 'context&', - output = 'const shape&', - input = 'const std::vector&', - const = True, - default = 'detail::compute_op'), - virtual('compute', - returns = 'argument', - output = 'const shape&', - input = 'const std::vector&', - const = True, - default = 'detail::compute_op'), - virtual( - 'compute', - returns = 'argument', - output = 'const shape&', - input = 'const std::vector&', - module_args = 'const std::vector&', - run = - 'std::function(module_ref&, const std::unordered_map&)>', - const = True, - default = 'detail::compute_op'), - virtual( - 'compute', - returns = 'argument', - ctx = 'context&', - output = 'const shape&', - input = 'const std::vector&', - module_args = 'const std::vector&', - run = - 'std::function(module_ref&, const std::unordered_map&)>', - const = True, - default = 'detail::compute_op'), - virtual('to_value', returns = 'value', const = True, default = 'detail::to_value_op'), - virtual('from_value', v = 'const value&', default = 'detail::from_value_op'), - virtual('attributes', returns = 'value', const = True, default = 'detail::attributes_op'), - friend('operator<<', - returns = 'std::ostream &', - os = 'std::ostream &', - op = 'const operation &', - using = 'migraphx::detail::operation_operators::operator<<'), - friend('operator==', - returns = 'bool', - x = 'const operation &', - y = 'const operation &', - using = 'migraphx::detail::operation_operators::operator==')) %> + interface( + 'operation', + virtual('name', returns = 'std::string', const = True), + virtual('is_context_free', + returns = 'bool', + const = True, + default = 'detail::is_context_free_op'), + virtual('need_normalization', + returns = 'bool', + const = True, + default = 'detail::need_normalization_op'), + virtual( + 'has_finalize', returns = 'bool', const = True, default = 'detail::has_finalize_op'), + virtual('get_lifetime', + returns = 'lifetime', + const = True, + default = 'detail::get_lifetime_op'), + virtual('output_alias', + returns = 'std::vector', + input = 'const std::vector&', + const = True, + default = 'detail::output_alias_op'), + virtual('compile', + returns = 'value', + ctx = 'context&', + output = 'const shape&', + input = 'const std::vector&', + default = 'detail::compile_op'), + virtual('finalize', + ctx = 'context&', + output = 'const shape&', + input = 'const std::vector&', + default = 'detail::finalize_op'), + virtual('compute_shape', + returns = 'shape', + input = 'const std::vector&', + const = True, + default = 'detail::compute_shape_op'), + virtual('compute_shape', + returns = 'shape', + inputs = 'const std::vector&', + mod_args = 'const std::vector&', + const = True, + default = 'detail::mod_compute_shape_op'), + virtual('compute', + returns = 'argument', + ctx = 'context&', + output = 'const shape&', + input = 'const std::vector&', + const = True, + default = 'detail::compute_op'), + virtual('compute', + returns = 'argument', + output = 'const shape&', + input = 'const std::vector&', + const = True, + default = 'detail::compute_op'), + virtual( + 'compute', + returns = 'argument', + output = 'const shape&', + input = 'const std::vector&', + module_args = 'const std::vector&', + run = + 'std::function(module_ref&, const std::unordered_map&)>', + const = True, + default = 'detail::compute_op'), + virtual( + 'compute', + returns = 'argument', + ctx = 'context&', + output = 'const shape&', + input = 'const std::vector&', + module_args = 'const std::vector&', + run = + 'std::function(module_ref&, const std::unordered_map&)>', + const = True, + default = 'detail::compute_op'), + virtual('to_value', returns = 'value', const = True, default = 'detail::to_value_op'), + virtual('from_value', v = 'const value&', default = 'detail::from_value_op'), + virtual('attributes', returns = 'value', const = True, default = 'detail::attributes_op'), + friend('operator<<', + returns = 'std::ostream &', + os = 'std::ostream &', + op = 'const operation &', + using = 'migraphx::detail::operation_operators::operator<<'), + friend('operator==', + returns = 'bool', + x = 'const operation &', + y = 'const operation &', + using = 'migraphx::detail::operation_operators::operator==')) +%> inline bool operator!=(const operation& x, const operation& y) { From f90598519ec8c55bb405a2c108b843e45f9fb092 Mon Sep 17 00:00:00 2001 From: Paul Date: Mon, 12 Jan 2026 02:47:10 +0000 Subject: [PATCH 03/23] Handle multi-alias --- src/adjust_allocation.cpp | 36 +++++++++++++++------------- src/driver/perf.cpp | 14 +++++++---- src/driver/trim.cpp | 6 ++--- src/eliminate_concat.cpp | 12 ++++++---- src/include/migraphx/instruction.hpp | 2 +- src/include/migraphx/liveness.hpp | 13 ++++++---- src/instruction.cpp | 28 +++++++++++++++------- src/module.cpp | 21 +++++++++------- src/propagate_constant.cpp | 7 +++--- src/replace_allocate.cpp | 17 ++++++++----- src/simplify_algebra.cpp | 14 ++++++----- src/split_reduce.cpp | 6 +++-- src/targets/gpu/jit/mlir.cpp | 18 ++++++++------ src/targets/gpu/lowering.cpp | 11 +++++---- src/targets/gpu/mlir.cpp | 22 +++++++++-------- test/output_alias.cpp | 15 ++++++------ 16 files changed, 145 insertions(+), 97 deletions(-) diff --git a/src/adjust_allocation.cpp b/src/adjust_allocation.cpp index 3dad3b7fade..430c76b85ba 100644 --- a/src/adjust_allocation.cpp +++ b/src/adjust_allocation.cpp @@ -43,24 +43,28 @@ void adjust_allocation::apply(module& m) const if(ins->get_operator().is_context_free()) continue; - auto alias_ins = instruction::get_output_alias(ins, true); - if(alias_ins->name() != model.name() and alias_ins->name() != "@param") - continue; - // shape allocated is different from actual shape - // of the instruction, reallocate and replace the previous one - if(alias_ins->get_shape() == ins->get_shape()) - continue; - auto alloc_ins = m.insert_instruction(ins, model.allocate(ins->get_shape())); - m.replace_instruction(alias_ins, alloc_ins); - // If the memory is an output parameter then copy the memory to the parameter - if(alias_ins->name() == "@param") + auto aliases = instruction::get_output_alias(ins, true); + for(auto alias_ins : aliases) { - auto copy = m.insert_instruction(std::next(ins), make_op(model.copy()), ins, alias_ins); - auto tail = range(std::next(copy), m.end()); - for(auto i : iterator_for(tail)) + if(alias_ins->name() != model.name() and alias_ins->name() != "@param") + continue; + // shape allocated is different from actual shape + // of the instruction, reallocate and replace the previous one + if(alias_ins->get_shape() == ins->get_shape()) + continue; + auto alloc_ins = m.insert_instruction(ins, model.allocate(ins->get_shape())); + m.replace_instruction(alias_ins, alloc_ins); + // If the memory is an output parameter then copy the memory to the parameter + if(alias_ins->name() == "@param") { - if(contains(i->inputs(), ins)) - instruction::replace_argument(i, ins, copy); + auto copy = + m.insert_instruction(std::next(ins), make_op(model.copy()), ins, alias_ins); + auto tail = range(std::next(copy), m.end()); + for(auto i : iterator_for(tail)) + { + if(contains(i->inputs(), ins)) + instruction::replace_argument(i, ins, copy); + } } } } diff --git a/src/driver/perf.cpp b/src/driver/perf.cpp index c5f4f33c1fb..9d5dd046402 100644 --- a/src/driver/perf.cpp +++ b/src/driver/perf.cpp @@ -118,17 +118,21 @@ bool is_offload_copy_set(const program& p) { if(i.name() == "hip::copy_to_gpu") { - auto copy_arg = instruction::get_output_alias(i.inputs().front(), true); - param_ins.erase(copy_arg); + auto copy_args = instruction::get_output_alias(i.inputs().front(), true); + for(auto copy_arg : copy_args) + param_ins.erase(copy_arg); } else if(i.name() == "@return") { auto return_args = i.inputs(); for(const auto& j : return_args) { - auto alias_ins = instruction::get_output_alias(j, true); - if((alias_ins->name() == "@param" and param_ins.erase(alias_ins) == 0) or - (alias_ins->name() != "hip::copy_from_gpu")) + auto aliases = instruction::get_output_alias(j, true); + bool valid = std::any_of(aliases.begin(), aliases.end(), [&](instruction_ref alias_ins) { + return (alias_ins->name() == "@param" and param_ins.erase(alias_ins) > 0) or + alias_ins->name() == "hip::copy_from_gpu"; + }); + if(not valid) return false; } } diff --git a/src/driver/trim.cpp b/src/driver/trim.cpp index 8158a04f846..f7c81658fb3 100644 --- a/src/driver/trim.cpp +++ b/src/driver/trim.cpp @@ -37,11 +37,11 @@ inline namespace MIGRAPHX_INLINE_NS { static instruction_ref capture_arg(std::unordered_set& s, instruction_ref ins) { - auto alias = instruction::get_output_alias(ins, true); - if(alias != ins) + auto aliases = instruction::get_output_alias(ins, true); + if(not(aliases.size() == 1 and aliases.front() == ins)) { s.insert(ins); - return capture_arg(s, alias); + return capture_arg(s, aliases.front()); } if(contains({"reshape", "contiguous"}, ins->name())) { diff --git a/src/eliminate_concat.cpp b/src/eliminate_concat.cpp index 47a095659cc..49b8d902f52 100644 --- a/src/eliminate_concat.cpp +++ b/src/eliminate_concat.cpp @@ -70,11 +70,13 @@ void eliminate_concat::apply(module& m) const // Where are the allocations for the tensors to be concatenated? std::vector allocations; - std::transform( - ins->inputs().begin(), - std::prev(ins->inputs().end()), - std::back_inserter(allocations), - [&](instruction_ref x) { return instruction::get_output_alias(x, true); }); + std::transform(ins->inputs().begin(), + std::prev(ins->inputs().end()), + std::back_inserter(allocations), + [&](instruction_ref x) { + auto aliases = instruction::get_output_alias(x, true); + return aliases.front(); + }); if(std::any_of(allocations.begin(), allocations.end(), [&](auto x) { return x->name() != concat_opt.allocate(); diff --git a/src/include/migraphx/instruction.hpp b/src/include/migraphx/instruction.hpp index 0369bb6ab64..2691d46dbc6 100644 --- a/src/include/migraphx/instruction.hpp +++ b/src/include/migraphx/instruction.hpp @@ -140,7 +140,7 @@ struct MIGRAPHX_EXPORT instruction void finalize(context& ctx); - static instruction_ref get_output_alias(instruction_ref ins, bool shallow = false); + static std::vector get_output_alias(instruction_ref ins, bool shallow = false); void set_normalized(bool value = true); bool is_normalized() const; diff --git a/src/include/migraphx/liveness.hpp b/src/include/migraphx/liveness.hpp index 6d9715a8a10..eb98bed1cd1 100644 --- a/src/include/migraphx/liveness.hpp +++ b/src/include/migraphx/liveness.hpp @@ -53,11 +53,14 @@ void liveness(const module& m, F f) auto add_live_variables = [&](const auto& inputs) { for(auto input : inputs) { - auto i = instruction::get_output_alias(input); - // Skip if variable comes from parent - if(not m.has_instruction(i)) - continue; - live_set.insert(i); + auto aliases = instruction::get_output_alias(input); + for(auto i : aliases) + { + // Skip if variable comes from parent + if(not m.has_instruction(i)) + continue; + live_set.insert(i); + } } }; add_live_variables(ins->inputs()); diff --git a/src/instruction.cpp b/src/instruction.cpp index 5d117c9c1e1..c10cb9c2da7 100644 --- a/src/instruction.cpp +++ b/src/instruction.cpp @@ -470,15 +470,25 @@ void instruction::debug_print() const std::cout << " -> " << this->get_shape() << std::endl; } -instruction_ref instruction::get_output_alias(instruction_ref ins, bool shallow) -{ - auto aliases = ins->get_operator().output_alias(to_shapes(ins->inputs())); - if(aliases.empty()) - return ins; - auto i = aliases.front(); - if(shallow) - return ins->inputs().at(i); - return get_output_alias(ins->inputs().at(i)); +std::vector instruction::get_output_alias(instruction_ref ins, bool shallow) +{ + auto alias_indices = ins->get_operator().output_alias(to_shapes(ins->inputs())); + if(alias_indices.empty()) + return {ins}; + std::vector result; + for(auto i : alias_indices) + { + if(shallow) + { + result.push_back(ins->inputs().at(i)); + } + else + { + auto nested = get_output_alias(ins->inputs().at(i)); + result.insert(result.end(), nested.begin(), nested.end()); + } + } + return result; } void instruction::set_normalized(bool value) { normalized = value; } diff --git a/src/module.cpp b/src/module.cpp index 4838d241904..0b04444d920 100644 --- a/src/module.cpp +++ b/src/module.cpp @@ -779,19 +779,24 @@ instruction_ref module::validate() const static bool is_borrowed(instruction_ref ins) { - auto alias = instruction::get_output_alias(ins, true); - if(alias == ins) + auto aliases = instruction::get_output_alias(ins, true); + if(aliases.size() == 1 and aliases.front() == ins) return false; - lifetime l = alias->get_operator().get_lifetime(); - if(l == lifetime::borrow) - return true; - return is_borrowed(alias); + return std::any_of(aliases.begin(), aliases.end(), [](instruction_ref alias) { + lifetime l = alias->get_operator().get_lifetime(); + if(l == lifetime::borrow) + return true; + return is_borrowed(alias); + }); } static bool is_global(instruction_ref ins) { - const auto& op = instruction::get_output_alias(ins)->get_operator(); - return op.name() == "@param" or op.get_lifetime() == lifetime::global; + auto aliases = instruction::get_output_alias(ins); + return std::any_of(aliases.begin(), aliases.end(), [](instruction_ref alias) { + const auto& op = alias->get_operator(); + return op.name() == "@param" or op.get_lifetime() == lifetime::global; + }); } static bool is_dangling(instruction_ref ins) { return not is_global(ins) and is_borrowed(ins); } diff --git a/src/propagate_constant.cpp b/src/propagate_constant.cpp index 8b44064fa2f..57e6e5dd019 100644 --- a/src/propagate_constant.cpp +++ b/src/propagate_constant.cpp @@ -45,9 +45,10 @@ static bool skip_propagate(instruction_ref ins) auto&& s = ins->get_shape(); if(s.broadcasted() and s.element_space() < s.elements()) return true; - auto alias = instruction::get_output_alias(ins, true); - if(alias != ins) - return skip_propagate(alias); + auto aliases = instruction::get_output_alias(ins, true); + if(not(aliases.size() == 1 and aliases.front() == ins)) + return std::any_of( + aliases.begin(), aliases.end(), [](instruction_ref alias) { return skip_propagate(alias); }); if(ins->is_undefined()) return true; return false; diff --git a/src/replace_allocate.cpp b/src/replace_allocate.cpp index e2f01e1dc91..3c6c10e6b53 100644 --- a/src/replace_allocate.cpp +++ b/src/replace_allocate.cpp @@ -41,11 +41,14 @@ std::unordered_map create_output_names(const modul std::unordered_map mod_output_names; auto returns = mod.get_returns(); - std::vector outputs_alias(returns.size()); + std::vector outputs_alias; + outputs_alias.reserve(returns.size()); - std::transform(returns.begin(), returns.end(), outputs_alias.begin(), [](const auto& i) { - return instruction::get_output_alias(i); - }); + for(const auto& i : returns) + { + auto aliases = instruction::get_output_alias(i); + outputs_alias.push_back(aliases.front()); + } std::size_t index = 0; if(outputs_alias.size() == 1 and mod.name().empty()) @@ -78,8 +81,10 @@ void insert_copy(module& m, const allocation_model& model) { if(ins->get_shape().any_of_dynamic()) continue; - auto alias = instruction::get_output_alias(ins); - if(alias->get_shape() == ins->get_shape()) + auto aliases = instruction::get_output_alias(ins); + if(std::any_of(aliases.begin(), aliases.end(), [&](instruction_ref alias) { + return alias->get_shape() == ins->get_shape(); + })) continue; auto insert_ins = std::next(ins); auto alloc = m.insert_instruction( diff --git a/src/simplify_algebra.cpp b/src/simplify_algebra.cpp index 5199b850f4f..4174e71538b 100644 --- a/src/simplify_algebra.cpp +++ b/src/simplify_algebra.cpp @@ -68,12 +68,14 @@ static auto from_int4() { return match::make_predicate_matcher([](instruction_ref start) { return fix([&](auto self, instruction_ref ins) { - auto alias = instruction::get_output_alias(ins); - if(contains({"reshape", "dequantizelinear"}, alias->name())) - return self(alias->inputs().front()); - if(alias->name() == "concat") - return all_of(alias->inputs(), self); - return alias->name() == "unpack_int4"; + auto aliases = instruction::get_output_alias(ins); + return std::any_of(aliases.begin(), aliases.end(), [&](instruction_ref alias) { + if(contains({"reshape", "dequantizelinear"}, alias->name())) + return self(alias->inputs().front()); + if(alias->name() == "concat") + return all_of(alias->inputs(), self); + return alias->name() == "unpack_int4"; + }); })(start); }); } diff --git a/src/split_reduce.cpp b/src/split_reduce.cpp index 3188b00563f..bda5656ffa6 100644 --- a/src/split_reduce.cpp +++ b/src/split_reduce.cpp @@ -119,8 +119,10 @@ struct splitter if(rins == rm->begin()) return; // We want to know what instructions are live after the split instruction - auto ins = instruction::get_output_alias(std::prev(rins)); - if(not contains(splits, ins)) + auto aliases = instruction::get_output_alias(std::prev(rins)); + if(not std::any_of(aliases.begin(), aliases.end(), [&](instruction_ref ins) { + return contains(splits, ins); + })) return; std::copy_if(live_set.begin(), live_set.end(), diff --git a/src/targets/gpu/jit/mlir.cpp b/src/targets/gpu/jit/mlir.cpp index 7fc49f1ca50..39c93de3224 100644 --- a/src/targets/gpu/jit/mlir.cpp +++ b/src/targets/gpu/jit/mlir.cpp @@ -76,14 +76,18 @@ struct mlir_compiler : compiler std::optional input_is_param(const instruction_ref& ins) const { - auto cur = instruction::get_output_alias(ins); - while(contains({"reshape", "contiguous"}, cur->name())) + auto aliases = instruction::get_output_alias(ins); + for(auto cur : aliases) { - cur = instruction::get_output_alias(cur->inputs().at(0)); - } - if(cur->name() == "@param") - { - return cur; + while(contains({"reshape", "contiguous"}, cur->name())) + { + auto nested_aliases = instruction::get_output_alias(cur->inputs().at(0)); + cur = nested_aliases.front(); + } + if(cur->name() == "@param") + { + return cur; + } } return nullopt; } diff --git a/src/targets/gpu/lowering.cpp b/src/targets/gpu/lowering.cpp index 0e24c3cc936..d038477999d 100644 --- a/src/targets/gpu/lowering.cpp +++ b/src/targets/gpu/lowering.cpp @@ -197,11 +197,14 @@ struct miopen_apply void insert_fill(instruction_ref ins, value v) const { - instruction_ref alloc = instruction::get_output_alias(ins, true); - if(alloc == ins) + auto aliases = instruction::get_output_alias(ins, true); + if(aliases.size() == 1 and aliases.front() == ins) return; - auto fill = mod->insert_instruction(ins, make_op("hip::fill", {{"value", v}}), alloc); - instruction::replace_argument(ins, alloc, fill); + for(instruction_ref alloc : aliases) + { + auto fill = mod->insert_instruction(ins, make_op("hip::fill", {{"value", v}}), alloc); + instruction::replace_argument(ins, alloc, fill); + } } instruction_ref insert_custom_op(instruction_ref ins, const value& attrs) const diff --git a/src/targets/gpu/mlir.cpp b/src/targets/gpu/mlir.cpp index 9ca5c35330c..f9e914fccd8 100644 --- a/src/targets/gpu/mlir.cpp +++ b/src/targets/gpu/mlir.cpp @@ -699,16 +699,18 @@ struct mlir_program static bool input_is_unpack_fp4(instruction_ref ins) { - ins = instruction::get_output_alias(ins); - if(ins->name() == "reshape") - { - return input_is_unpack_fp4(ins->inputs().front()); - } - if(ins->name() == "unpack_fp4") - { - return true; - } - return false; + auto aliases = instruction::get_output_alias(ins); + return std::any_of(aliases.begin(), aliases.end(), [](instruction_ref alias) { + if(alias->name() == "reshape") + { + return input_is_unpack_fp4(alias->inputs().front()); + } + if(alias->name() == "unpack_fp4") + { + return true; + } + return false; + }); } static shape make_fp4_unpacked_shape(shape s) diff --git a/test/output_alias.cpp b/test/output_alias.cpp index be6b1eac102..322bb40002a 100644 --- a/test/output_alias.cpp +++ b/test/output_alias.cpp @@ -23,6 +23,7 @@ */ #include #include +#include #include #include @@ -32,8 +33,8 @@ TEST_CASE(simple_alias) auto* mm = p.get_main_module(); auto l = mm->add_literal(1); auto p1 = mm->add_instruction(pass_op{}, l); - EXPECT(migraphx::instruction::get_output_alias(l) == l); - EXPECT(migraphx::instruction::get_output_alias(p1) == l); + EXPECT(migraphx::contains(migraphx::instruction::get_output_alias(l), l)); + EXPECT(migraphx::contains(migraphx::instruction::get_output_alias(p1), l)); } TEST_CASE(cascade_alias) @@ -44,10 +45,10 @@ TEST_CASE(cascade_alias) auto p1 = mm->add_instruction(pass_op{}, l); auto p2 = mm->add_instruction(pass_op{}, p1); auto p3 = mm->add_instruction(pass_op{}, p2); - EXPECT(migraphx::instruction::get_output_alias(l) == l); - EXPECT(migraphx::instruction::get_output_alias(p1) == l); - EXPECT(migraphx::instruction::get_output_alias(p2) == l); - EXPECT(migraphx::instruction::get_output_alias(p3) == l); + EXPECT(migraphx::contains(migraphx::instruction::get_output_alias(l), l)); + EXPECT(migraphx::contains(migraphx::instruction::get_output_alias(p1), l)); + EXPECT(migraphx::contains(migraphx::instruction::get_output_alias(p2), l)); + EXPECT(migraphx::contains(migraphx::instruction::get_output_alias(p3), l)); } TEST_CASE(no_alias) @@ -57,7 +58,7 @@ TEST_CASE(no_alias) auto x = mm->add_literal(1); auto y = mm->add_literal(2); auto sum = mm->add_instruction(sum_op{}, x, y); - EXPECT(migraphx::instruction::get_output_alias(sum) == sum); + EXPECT(migraphx::contains(migraphx::instruction::get_output_alias(sum), sum)); } int main(int argc, const char* argv[]) { test::run(argc, argv); } From 7257328aedc0a61a334f8403f6d277d028a9f2e4 Mon Sep 17 00:00:00 2001 From: Paul Date: Mon, 12 Jan 2026 02:47:14 +0000 Subject: [PATCH 04/23] Format --- src/driver/perf.cpp | 9 +++++---- src/propagate_constant.cpp | 5 +++-- 2 files changed, 8 insertions(+), 6 deletions(-) diff --git a/src/driver/perf.cpp b/src/driver/perf.cpp index 9d5dd046402..aa3b8d16868 100644 --- a/src/driver/perf.cpp +++ b/src/driver/perf.cpp @@ -128,10 +128,11 @@ bool is_offload_copy_set(const program& p) for(const auto& j : return_args) { auto aliases = instruction::get_output_alias(j, true); - bool valid = std::any_of(aliases.begin(), aliases.end(), [&](instruction_ref alias_ins) { - return (alias_ins->name() == "@param" and param_ins.erase(alias_ins) > 0) or - alias_ins->name() == "hip::copy_from_gpu"; - }); + bool valid = + std::any_of(aliases.begin(), aliases.end(), [&](instruction_ref alias_ins) { + return (alias_ins->name() == "@param" and param_ins.erase(alias_ins) > 0) or + alias_ins->name() == "hip::copy_from_gpu"; + }); if(not valid) return false; } diff --git a/src/propagate_constant.cpp b/src/propagate_constant.cpp index 57e6e5dd019..9974b99e903 100644 --- a/src/propagate_constant.cpp +++ b/src/propagate_constant.cpp @@ -47,8 +47,9 @@ static bool skip_propagate(instruction_ref ins) return true; auto aliases = instruction::get_output_alias(ins, true); if(not(aliases.size() == 1 and aliases.front() == ins)) - return std::any_of( - aliases.begin(), aliases.end(), [](instruction_ref alias) { return skip_propagate(alias); }); + return std::any_of(aliases.begin(), aliases.end(), [](instruction_ref alias) { + return skip_propagate(alias); + }); if(ins->is_undefined()) return true; return false; From f61a10f48850525a63b0060b0f516a97ad3217d1 Mon Sep 17 00:00:00 2001 From: Paul Date: Mon, 12 Jan 2026 15:50:16 +0000 Subject: [PATCH 05/23] Add unit tests --- test/include/basic_ops.hpp | 19 +++++++++ test/output_alias.cpp | 73 ++++++++++++++++++++++++++++++++ test/propagate_constant_test.cpp | 30 +++++++++++++ test/replace_allocate.cpp | 23 ++++++++++ 4 files changed, 145 insertions(+) diff --git a/test/include/basic_ops.hpp b/test/include/basic_ops.hpp index 5a6acaf1e59..4e2664c6826 100644 --- a/test/include/basic_ops.hpp +++ b/test/include/basic_ops.hpp @@ -256,6 +256,25 @@ struct tuple_op } }; +// Operation that aliases multiple inputs (all inputs) +struct multi_alias_op +{ + std::string name() const { return "multi_alias"; } + migraphx::shape compute_shape(std::vector inputs) const + { + if(inputs.empty()) + MIGRAPHX_THROW("Need at least 1 input"); + return inputs.front(); + } + std::vector output_alias(const std::vector& s) const + { + std::vector result; + for(std::size_t i = 0; i < s.size(); ++i) + result.push_back(i); + return result; + } +}; + inline migraphx::literal get_2x2(int base = 0) { return migraphx::literal{{migraphx::shape::float_type, {2, 2}}, diff --git a/test/output_alias.cpp b/test/output_alias.cpp index 322bb40002a..e4d4830b940 100644 --- a/test/output_alias.cpp +++ b/test/output_alias.cpp @@ -61,4 +61,77 @@ TEST_CASE(no_alias) EXPECT(migraphx::contains(migraphx::instruction::get_output_alias(sum), sum)); } +TEST_CASE(multiple_aliases) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + auto x = mm->add_literal(1); + auto y = mm->add_literal(2); + auto ma = mm->add_instruction(multi_alias_op{}, x, y); + auto aliases = migraphx::instruction::get_output_alias(ma); + // multi_alias_op aliases both inputs, so we should get both literals back + EXPECT(aliases.size() == 2); + EXPECT(migraphx::contains(aliases, x)); + EXPECT(migraphx::contains(aliases, y)); +} + +TEST_CASE(multiple_aliases_shallow) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + auto x = mm->add_literal(1); + auto y = mm->add_literal(2); + auto p1 = mm->add_instruction(pass_op{}, x); + auto p2 = mm->add_instruction(pass_op{}, y); + auto ma = mm->add_instruction(multi_alias_op{}, p1, p2); + // shallow=true returns immediate inputs (p1, p2), not root aliases + auto shallow_aliases = migraphx::instruction::get_output_alias(ma, true); + EXPECT(shallow_aliases.size() == 2); + EXPECT(migraphx::contains(shallow_aliases, p1)); + EXPECT(migraphx::contains(shallow_aliases, p2)); + // shallow=false (default) returns root aliases (x, y) + auto deep_aliases = migraphx::instruction::get_output_alias(ma); + EXPECT(deep_aliases.size() == 2); + EXPECT(migraphx::contains(deep_aliases, x)); + EXPECT(migraphx::contains(deep_aliases, y)); +} + +TEST_CASE(multiple_aliases_cascade) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + auto x = mm->add_literal(1); + auto y = mm->add_literal(2); + auto z = mm->add_literal(3); + // First multi_alias aliases x and y + auto ma1 = mm->add_instruction(multi_alias_op{}, x, y); + // Second multi_alias aliases ma1 and z + auto ma2 = mm->add_instruction(multi_alias_op{}, ma1, z); + // Should recursively expand to get x, y, z + auto aliases = migraphx::instruction::get_output_alias(ma2); + EXPECT(aliases.size() == 3); + EXPECT(migraphx::contains(aliases, x)); + EXPECT(migraphx::contains(aliases, y)); + EXPECT(migraphx::contains(aliases, z)); +} + +TEST_CASE(alias_vector_size) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + auto l = mm->add_literal(1); + // No alias - returns vector with self + auto aliases_self = migraphx::instruction::get_output_alias(l); + EXPECT(aliases_self.size() == 1); + // Single alias - returns vector with one element + auto p1 = mm->add_instruction(pass_op{}, l); + auto aliases_single = migraphx::instruction::get_output_alias(p1); + EXPECT(aliases_single.size() == 1); + // Multiple aliases - returns vector with multiple elements + auto x = mm->add_literal(2); + auto ma = mm->add_instruction(multi_alias_op{}, l, x); + auto aliases_multi = migraphx::instruction::get_output_alias(ma); + EXPECT(aliases_multi.size() == 2); +} + int main(int argc, const char* argv[]) { test::run(argc, argv); } diff --git a/test/propagate_constant_test.cpp b/test/propagate_constant_test.cpp index 72011fd8d83..41537d789d9 100644 --- a/test/propagate_constant_test.cpp +++ b/test/propagate_constant_test.cpp @@ -562,4 +562,34 @@ TEST_CASE(pack_unpack_fp4) EXPECT(m1 == m2); } +// Test that propagate_constant correctly handles multi-alias operations +// when one of the aliases should be skipped (e.g., broadcasted) +TEST_CASE(skip_propagate_multi_alias) +{ + // When an instruction has multiple aliases and one should be skipped, + // propagation should be skipped for the entire instruction + migraphx::module m1; + { + // Create a broadcasted literal (should skip propagation) + auto broadcasted = m1.add_literal( + migraphx::literal{{migraphx::shape::float_type, {2, 1}, {1, 0}}, {1.0f, 2.0f}}); + // Create a normal literal + auto normal = m1.add_literal( + migraphx::literal{{migraphx::shape::float_type, {2}}, {3.0f, 4.0f}}); + // multi_alias_op aliases both inputs + auto ma = m1.add_instruction(multi_alias_op{}, broadcasted, normal); + // Add an operation that uses ma + auto neg = m1.add_instruction(migraphx::make_op("neg"), ma); + m1.add_return({neg}); + } + + migraphx::module m2 = m1; + run_pass(m1); + + // Since one alias (broadcasted) should skip propagation, + // the multi_alias instruction should not be propagated + // The modules should be equivalent (no propagation happened) + EXPECT(m1 == m2); +} + int main(int argc, const char* argv[]) { test::run(argc, argv); } diff --git a/test/replace_allocate.cpp b/test/replace_allocate.cpp index dc0eba17073..03a3358e331 100644 --- a/test/replace_allocate.cpp +++ b/test/replace_allocate.cpp @@ -279,4 +279,27 @@ TEST_CASE(allocate_copy_with_no_out) EXPECT(m1.sort() == m2.sort()); } +// Test that replace_allocate handles multi-alias operations correctly +// when checking for shape matches (insert_copy code path) +TEST_CASE(multi_alias_shape_check) +{ + migraphx::shape s{migraphx::shape::float_type, {5}}; + migraphx::module m1; + { + auto x = m1.add_parameter("x", s); + auto y = m1.add_parameter("y", s); + // multi_alias_op aliases both x and y (both have same shape) + auto ma = m1.add_instruction(multi_alias_op{}, x, y); + m1.add_return({ma}); + } + + // After pass, since the multi_alias aliases inputs with matching shapes, + // no copy should be inserted + migraphx::module m2 = m1; + run_pass(m1, allocation_no_out_model{}); + + // The module should remain unchanged since aliases have matching shapes + EXPECT(m1.sort() == m2.sort()); +} + int main(int argc, const char* argv[]) { test::run(argc, argv); } From ce51c3d17b027ceaf1301f89a79b758faa7a1f62 Mon Sep 17 00:00:00 2001 From: Paul Date: Mon, 12 Jan 2026 15:50:22 +0000 Subject: [PATCH 06/23] Format --- test/output_alias.cpp | 14 +++++++------- test/propagate_constant_test.cpp | 4 ++-- test/replace_allocate.cpp | 4 ++-- 3 files changed, 11 insertions(+), 11 deletions(-) diff --git a/test/output_alias.cpp b/test/output_alias.cpp index e4d4830b940..d56d8dfd59d 100644 --- a/test/output_alias.cpp +++ b/test/output_alias.cpp @@ -64,10 +64,10 @@ TEST_CASE(no_alias) TEST_CASE(multiple_aliases) { migraphx::program p; - auto* mm = p.get_main_module(); - auto x = mm->add_literal(1); - auto y = mm->add_literal(2); - auto ma = mm->add_instruction(multi_alias_op{}, x, y); + auto* mm = p.get_main_module(); + auto x = mm->add_literal(1); + auto y = mm->add_literal(2); + auto ma = mm->add_instruction(multi_alias_op{}, x, y); auto aliases = migraphx::instruction::get_output_alias(ma); // multi_alias_op aliases both inputs, so we should get both literals back EXPECT(aliases.size() == 2); @@ -124,12 +124,12 @@ TEST_CASE(alias_vector_size) auto aliases_self = migraphx::instruction::get_output_alias(l); EXPECT(aliases_self.size() == 1); // Single alias - returns vector with one element - auto p1 = mm->add_instruction(pass_op{}, l); + auto p1 = mm->add_instruction(pass_op{}, l); auto aliases_single = migraphx::instruction::get_output_alias(p1); EXPECT(aliases_single.size() == 1); // Multiple aliases - returns vector with multiple elements - auto x = mm->add_literal(2); - auto ma = mm->add_instruction(multi_alias_op{}, l, x); + auto x = mm->add_literal(2); + auto ma = mm->add_instruction(multi_alias_op{}, l, x); auto aliases_multi = migraphx::instruction::get_output_alias(ma); EXPECT(aliases_multi.size() == 2); } diff --git a/test/propagate_constant_test.cpp b/test/propagate_constant_test.cpp index 41537d789d9..c348f84a688 100644 --- a/test/propagate_constant_test.cpp +++ b/test/propagate_constant_test.cpp @@ -574,8 +574,8 @@ TEST_CASE(skip_propagate_multi_alias) auto broadcasted = m1.add_literal( migraphx::literal{{migraphx::shape::float_type, {2, 1}, {1, 0}}, {1.0f, 2.0f}}); // Create a normal literal - auto normal = m1.add_literal( - migraphx::literal{{migraphx::shape::float_type, {2}}, {3.0f, 4.0f}}); + auto normal = + m1.add_literal(migraphx::literal{{migraphx::shape::float_type, {2}}, {3.0f, 4.0f}}); // multi_alias_op aliases both inputs auto ma = m1.add_instruction(multi_alias_op{}, broadcasted, normal); // Add an operation that uses ma diff --git a/test/replace_allocate.cpp b/test/replace_allocate.cpp index 03a3358e331..8eac4d7e407 100644 --- a/test/replace_allocate.cpp +++ b/test/replace_allocate.cpp @@ -286,8 +286,8 @@ TEST_CASE(multi_alias_shape_check) migraphx::shape s{migraphx::shape::float_type, {5}}; migraphx::module m1; { - auto x = m1.add_parameter("x", s); - auto y = m1.add_parameter("y", s); + auto x = m1.add_parameter("x", s); + auto y = m1.add_parameter("y", s); // multi_alias_op aliases both x and y (both have same shape) auto ma = m1.add_instruction(multi_alias_op{}, x, y); m1.add_return({ma}); From d64b5ae01f2deb671368ad3ee684bb0e121ee183 Mon Sep 17 00:00:00 2001 From: Paul Date: Mon, 12 Jan 2026 12:17:56 -0600 Subject: [PATCH 07/23] Dont use any_of --- src/adjust_allocation.cpp | 38 +++++++++++++++++++------------------- src/propagate_constant.cpp | 6 ++---- 2 files changed, 21 insertions(+), 23 deletions(-) diff --git a/src/adjust_allocation.cpp b/src/adjust_allocation.cpp index 430c76b85ba..398bd75f479 100644 --- a/src/adjust_allocation.cpp +++ b/src/adjust_allocation.cpp @@ -44,27 +44,27 @@ void adjust_allocation::apply(module& m) const continue; auto aliases = instruction::get_output_alias(ins, true); - for(auto alias_ins : aliases) + if(aliases.size() != 1) + continue; + auto alias_ins = aliases.front(); + if(alias_ins->name() != model.name() and alias_ins->name() != "@param") + continue; + // shape allocated is different from actual shape + // of the instruction, reallocate and replace the previous one + if(alias_ins->get_shape() == ins->get_shape()) + continue; + auto alloc_ins = m.insert_instruction(ins, model.allocate(ins->get_shape())); + m.replace_instruction(alias_ins, alloc_ins); + // If the memory is an output parameter then copy the memory to the parameter + if(alias_ins->name() == "@param") { - if(alias_ins->name() != model.name() and alias_ins->name() != "@param") - continue; - // shape allocated is different from actual shape - // of the instruction, reallocate and replace the previous one - if(alias_ins->get_shape() == ins->get_shape()) - continue; - auto alloc_ins = m.insert_instruction(ins, model.allocate(ins->get_shape())); - m.replace_instruction(alias_ins, alloc_ins); - // If the memory is an output parameter then copy the memory to the parameter - if(alias_ins->name() == "@param") + auto copy = + m.insert_instruction(std::next(ins), make_op(model.copy()), ins, alias_ins); + auto tail = range(std::next(copy), m.end()); + for(auto i : iterator_for(tail)) { - auto copy = - m.insert_instruction(std::next(ins), make_op(model.copy()), ins, alias_ins); - auto tail = range(std::next(copy), m.end()); - for(auto i : iterator_for(tail)) - { - if(contains(i->inputs(), ins)) - instruction::replace_argument(i, ins, copy); - } + if(contains(i->inputs(), ins)) + instruction::replace_argument(i, ins, copy); } } } diff --git a/src/propagate_constant.cpp b/src/propagate_constant.cpp index 9974b99e903..9aed3268553 100644 --- a/src/propagate_constant.cpp +++ b/src/propagate_constant.cpp @@ -46,10 +46,8 @@ static bool skip_propagate(instruction_ref ins) if(s.broadcasted() and s.element_space() < s.elements()) return true; auto aliases = instruction::get_output_alias(ins, true); - if(not(aliases.size() == 1 and aliases.front() == ins)) - return std::any_of(aliases.begin(), aliases.end(), [](instruction_ref alias) { - return skip_propagate(alias); - }); + if(aliases.size() == 1 and aliases.front() != ins) + return skip_propagate(aliases.front()); if(ins->is_undefined()) return true; return false; From 4517550d9e1e57d7d8be43242c444cac7766ca65 Mon Sep 17 00:00:00 2001 From: Paul Date: Mon, 12 Jan 2026 12:18:01 -0600 Subject: [PATCH 08/23] Format --- src/adjust_allocation.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/src/adjust_allocation.cpp b/src/adjust_allocation.cpp index 398bd75f479..e95d7f53443 100644 --- a/src/adjust_allocation.cpp +++ b/src/adjust_allocation.cpp @@ -58,8 +58,7 @@ void adjust_allocation::apply(module& m) const // If the memory is an output parameter then copy the memory to the parameter if(alias_ins->name() == "@param") { - auto copy = - m.insert_instruction(std::next(ins), make_op(model.copy()), ins, alias_ins); + auto copy = m.insert_instruction(std::next(ins), make_op(model.copy()), ins, alias_ins); auto tail = range(std::next(copy), m.end()); for(auto i : iterator_for(tail)) { From 409bb9088e607864076177bed80c0225cf8034a5 Mon Sep 17 00:00:00 2001 From: Paul Date: Mon, 12 Jan 2026 19:31:09 +0000 Subject: [PATCH 09/23] Fix replace allocate --- src/replace_allocate.cpp | 30 ++++-- test/replace_allocate.cpp | 207 ++++++++++++++++++++++++++++++++++++++ 2 files changed, 229 insertions(+), 8 deletions(-) diff --git a/src/replace_allocate.cpp b/src/replace_allocate.cpp index 3c6c10e6b53..19d83ea5572 100644 --- a/src/replace_allocate.cpp +++ b/src/replace_allocate.cpp @@ -41,31 +41,45 @@ std::unordered_map create_output_names(const modul std::unordered_map mod_output_names; auto returns = mod.get_returns(); - std::vector outputs_alias; - outputs_alias.reserve(returns.size()); - + // Collect all allocation aliases from each return value + std::vector alloc_aliases; for(const auto& i : returns) { auto aliases = instruction::get_output_alias(i); - outputs_alias.push_back(aliases.front()); + for(auto ins : aliases) + { + if(ins->name() == "allocate") + alloc_aliases.push_back(ins); + } } std::size_t index = 0; - if(outputs_alias.size() == 1 and mod.name().empty()) + if(returns.size() == 1 and mod.name().empty()) { - mod_output_names[outputs_alias.front()] = "output"; + // Single return with empty module name: all aliases get "output" or "output_N" + if(alloc_aliases.size() == 1) + { + mod_output_names[alloc_aliases.front()] = "output"; + } + else + { + for(auto ins : alloc_aliases) + { + mod_output_names[ins] = "output_" + std::to_string(index++); + } + } } // Preserve main module output buffer naming across migraphx versions else if(mod.name() == "main") { - for(auto ins : outputs_alias) + for(auto ins : alloc_aliases) { mod_output_names[ins] = mod.name() + ":#output_" + std::to_string(index++); } } else { - for(auto ins : outputs_alias) + for(auto ins : alloc_aliases) { mod_output_names[ins] = param_name(index++, mod.name() + ":#output_"); } diff --git a/test/replace_allocate.cpp b/test/replace_allocate.cpp index 8eac4d7e407..e13c048a163 100644 --- a/test/replace_allocate.cpp +++ b/test/replace_allocate.cpp @@ -302,4 +302,211 @@ TEST_CASE(multi_alias_shape_check) EXPECT(m1.sort() == m2.sort()); } +// Test multi-alias with allocation replaced by output parameter +// When first alias is an allocation, it gets replaced with output parameter +TEST_CASE(multi_alias_alloc_out_param) +{ + migraphx::shape s{migraphx::shape::float_type, {5}}; + migraphx::module m1; + { + auto x = m1.add_parameter("x", s); + // Create allocation that is first in multi_alias - this will be used for output naming + auto alloc = + m1.add_instruction(migraphx::make_op("allocate", {{"shape", migraphx::to_value(s)}})); + auto p1 = m1.add_instruction(pass_op{}, alloc); + // Put allocation-based alias first so it gets used for output naming + auto ma = m1.add_instruction(multi_alias_op{}, p1, x); + m1.add_return({ma}); + } + run_pass(m1, allocation_with_out_model{}); + + migraphx::module m2; + { + auto x = m2.add_parameter("x", s); + // Only the allocation becomes an output parameter (named "output" since single alloc) + auto output = m2.add_parameter("output", s); + auto p1 = m2.add_instruction(pass_op{}, output); + auto ma = m2.add_instruction(multi_alias_op{}, p1, x); + m2.add_return({ma}); + } + EXPECT(m1.sort() == m2.sort()); +} + +// Test multi-alias where both inputs are allocations - both become output parameters +TEST_CASE(multi_alias_two_allocs_out_param) +{ + migraphx::shape s{migraphx::shape::float_type, {5}}; + migraphx::module m1; + { + // First allocation will be replaced with output parameter + auto alloc1 = + m1.add_instruction(migraphx::make_op("allocate", {{"shape", migraphx::to_value(s)}})); + auto p1 = m1.add_instruction(pass_op{}, alloc1); + // Second allocation will also be replaced with output parameter + auto alloc2 = + m1.add_instruction(migraphx::make_op("allocate", {{"shape", migraphx::to_value(s)}})); + auto p2 = m1.add_instruction(pass_op{}, alloc2); + auto ma = m1.add_instruction(multi_alias_op{}, p1, p2); + m1.add_return({ma}); + } + run_pass(m1, allocation_with_out_model{}); + + migraphx::module m2; + { + // Both aliases become output parameters + auto output0 = m2.add_parameter("output_0", s); + auto p1 = m2.add_instruction(pass_op{}, output0); + auto output1 = m2.add_parameter("output_1", s); + auto p2 = m2.add_instruction(pass_op{}, output1); + auto ma = m2.add_instruction(multi_alias_op{}, p1, p2); + m2.add_return({ma}); + } + EXPECT(m1.sort() == m2.sort()); +} + +// Test multi-alias with matching shapes - no copy needed when any alias matches +TEST_CASE(multi_alias_no_copy_when_any_matches) +{ + migraphx::shape s{migraphx::shape::float_type, {5}}; + migraphx::shape s2{migraphx::shape::float_type, {10}}; + migraphx::module m1; + { + // x has shape {5}, y has shape {10} + // multi_alias output has shape {5} (from first input) + // Since x's shape matches output shape, no copy is needed + auto x = m1.add_parameter("x", s); + auto y = m1.add_parameter("y", s2); + auto ma = m1.add_instruction(multi_alias_op{}, x, y); + m1.add_return({ma}); + } + + // Module should be unchanged - no copy inserted because first alias shape matches + migraphx::module m2 = m1; + run_pass(m1, allocation_no_out_model{}); + EXPECT(m1.sort() == m2.sort()); +} + +// Test multiple return values where each has a multi-alias with allocation as first alias +// Both allocations should be replaced with separate output parameters +TEST_CASE(multi_alias_multiple_outputs_out_params) +{ + migraphx::shape s{migraphx::shape::float_type, {5}}; + migraphx::module m1; + { + auto x = m1.add_parameter("x", s); + auto y = m1.add_parameter("y", s); + // First allocation -> pass -> multi_alias with x + auto alloc1 = + m1.add_instruction(migraphx::make_op("allocate", {{"shape", migraphx::to_value(s)}})); + auto p1 = m1.add_instruction(pass_op{}, alloc1); + auto ma1 = m1.add_instruction(multi_alias_op{}, p1, x); + // Second allocation -> pass -> multi_alias with y + auto alloc2 = + m1.add_instruction(migraphx::make_op("allocate", {{"shape", migraphx::to_value(s)}})); + auto p2 = m1.add_instruction(pass_op{}, alloc2); + auto ma2 = m1.add_instruction(multi_alias_op{}, p2, y); + // Return both multi_alias results - each should get its own output parameter + m1.add_return({ma1, ma2}); + } + run_pass(m1, allocation_with_out_model{}); + + migraphx::module m2; + { + auto x = m2.add_parameter("x", s); + auto y = m2.add_parameter("y", s); + // First output parameter replaces first allocation (named :#output_0) + auto output0 = m2.add_parameter(":#output_0", s); + auto p1 = m2.add_instruction(pass_op{}, output0); + auto ma1 = m2.add_instruction(multi_alias_op{}, p1, x); + // Second output parameter replaces second allocation (named :#output_1) + auto output1 = m2.add_parameter(":#output_1", s); + auto p2 = m2.add_instruction(pass_op{}, output1); + auto ma2 = m2.add_instruction(multi_alias_op{}, p2, y); + m2.add_return({ma1, ma2}); + } + EXPECT(m1.sort() == m2.sort()); +} + +// Test multi-alias with 3 allocations - all become output parameters +TEST_CASE(multi_alias_three_allocs) +{ + migraphx::shape s{migraphx::shape::float_type, {5}}; + migraphx::module m1; + { + // Three allocations wrapped in pass_op + auto alloc1 = + m1.add_instruction(migraphx::make_op("allocate", {{"shape", migraphx::to_value(s)}})); + auto p1 = m1.add_instruction(pass_op{}, alloc1); + auto alloc2 = + m1.add_instruction(migraphx::make_op("allocate", {{"shape", migraphx::to_value(s)}})); + auto p2 = m1.add_instruction(pass_op{}, alloc2); + auto alloc3 = + m1.add_instruction(migraphx::make_op("allocate", {{"shape", migraphx::to_value(s)}})); + auto p3 = m1.add_instruction(pass_op{}, alloc3); + // multi_alias aliases all three allocations + auto ma = m1.add_instruction(multi_alias_op{}, p1, p2, p3); + m1.add_return({ma}); + } + run_pass(m1, allocation_with_out_model{}); + + migraphx::module m2; + { + // All three aliases become output parameters + auto output0 = m2.add_parameter("output_0", s); + auto p1 = m2.add_instruction(pass_op{}, output0); + auto output1 = m2.add_parameter("output_1", s); + auto p2 = m2.add_instruction(pass_op{}, output1); + auto output2 = m2.add_parameter("output_2", s); + auto p3 = m2.add_instruction(pass_op{}, output2); + auto ma = m2.add_instruction(multi_alias_op{}, p1, p2, p3); + m2.add_return({ma}); + } + EXPECT(m1.sort() == m2.sort()); +} + +// Test where multiple multi_alias ops each contribute an allocation to multiple returns +// Each allocation becomes a separate output parameter +TEST_CASE(multi_alias_chain_multiple_out_params) +{ + migraphx::shape s{migraphx::shape::float_type, {5}}; + migraphx::module m1; + { + auto x = m1.add_parameter("x", s); + // Three allocations + auto alloc1 = + m1.add_instruction(migraphx::make_op("allocate", {{"shape", migraphx::to_value(s)}})); + auto p1 = m1.add_instruction(pass_op{}, alloc1); + auto alloc2 = + m1.add_instruction(migraphx::make_op("allocate", {{"shape", migraphx::to_value(s)}})); + auto p2 = m1.add_instruction(pass_op{}, alloc2); + auto alloc3 = + m1.add_instruction(migraphx::make_op("allocate", {{"shape", migraphx::to_value(s)}})); + auto p3 = m1.add_instruction(pass_op{}, alloc3); + // Each multi_alias puts an allocation first + auto ma1 = m1.add_instruction(multi_alias_op{}, p1, x); + auto ma2 = m1.add_instruction(multi_alias_op{}, p2, x); + auto ma3 = m1.add_instruction(multi_alias_op{}, p3, x); + // Return all three - each should get its own output parameter + m1.add_return({ma1, ma2, ma3}); + } + run_pass(m1, allocation_with_out_model{}); + + migraphx::module m2; + { + auto x = m2.add_parameter("x", s); + // Three output parameters - one for each return + auto output0 = m2.add_parameter(":#output_0", s); + auto p1 = m2.add_instruction(pass_op{}, output0); + auto output1 = m2.add_parameter(":#output_1", s); + auto p2 = m2.add_instruction(pass_op{}, output1); + auto output2 = m2.add_parameter(":#output_2", s); + auto p3 = m2.add_instruction(pass_op{}, output2); + auto ma1 = m2.add_instruction(multi_alias_op{}, p1, x); + auto ma2 = m2.add_instruction(multi_alias_op{}, p2, x); + auto ma3 = m2.add_instruction(multi_alias_op{}, p3, x); + m2.add_return({ma1, ma2, ma3}); + } + EXPECT(m1.sort() == m2.sort()); +} + int main(int argc, const char* argv[]) { test::run(argc, argv); } From 3ed20a5eaf66a824a58dd8b4d988ae3cc4e018ac Mon Sep 17 00:00:00 2001 From: Paul Date: Mon, 12 Jan 2026 20:08:50 +0000 Subject: [PATCH 10/23] Add split_reduce tests --- test/split_reduce.cpp | 109 ++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 109 insertions(+) diff --git a/test/split_reduce.cpp b/test/split_reduce.cpp index 4679f265753..25de25f96fd 100644 --- a/test/split_reduce.cpp +++ b/test/split_reduce.cpp @@ -26,6 +26,7 @@ #include #include #include +#include #include #include #include @@ -294,4 +295,112 @@ TEST_CASE(double_split_live) EXPECT(p1.sort() == p2.sort()); } +// Test multi-alias in parallel reduce scenario - both reduce outputs are aliased by multi_alias_op +// The pass should split both reduces and extract the multi_alias to the main module +TEST_CASE(parallel_reduce_multi_alias) +{ + migraphx::shape s{migraphx::shape::float_type, {2, 3, 327680}}; + migraphx::program p1; + { + auto* mm = p1.get_main_module(); + auto x = mm->add_parameter("x", s); + auto rsum = add_reduce( + p1, "fuse_reduce0", {x}, {2}, [&](auto* rm, const auto& inputs, const auto& axes) { + auto xx = add_pointwise(p1, rm, "main:pointwise0", {inputs[0]}, squared()); + auto rsum1 = rm->add_instruction(migraphx::make_op("reduce_sum", {{"axes", axes}}), + inputs[0]); + auto rsum2 = + rm->add_instruction(migraphx::make_op("reduce_sum", {{"axes", axes}}), xx); + // multi_alias_op aliases both reduce outputs + auto ma = rm->add_instruction(multi_alias_op{}, rsum1, rsum2); + return ma; + }); + mm->add_return({rsum}); + } + run_pass(p1); + migraphx::program p2; + { + auto* mm = p2.get_main_module(); + auto x = mm->add_parameter("x", s); + // The pointwise (squared) is extracted to main module + auto xx = add_pointwise(p2, mm, "main:pointwise0", {x}, squared()); + // Split module takes both xx and x as inputs + auto rsum = add_reduce( + p2, + "fuse_reduce0_split", + {xx, x}, + {2}, + "assign_add", + [&](auto* rm, + const auto& inputs, + const auto& axes) -> std::vector { + // inputs[0] is xx (squared), inputs[1] is x + // The pass returns (rsum1, rsum2) order based on the original fused module order + auto rsum1 = rm->add_instruction(migraphx::make_op("reduce_sum", {{"axes", axes}}), + inputs[1]); + auto rsum2 = rm->add_instruction(migraphx::make_op("reduce_sum", {{"axes", axes}}), + inputs[0]); + return {rsum1, rsum2}; + }); + auto rsum1 = mm->add_instruction(migraphx::make_op("get_tuple_elem", {{"index", 0}}), rsum); + auto rsum2 = mm->add_instruction(migraphx::make_op("get_tuple_elem", {{"index", 1}}), rsum); + // multi_alias_op is moved to main module after split + auto ma = mm->add_instruction(multi_alias_op{}, rsum1, rsum2); + mm->add_return({ma}); + } + EXPECT(p1.sort() == p2.sort()); +} + +// Test that find_alive correctly identifies live instructions through multi-alias chain +// sqrt is computed before reduce, used after reduce through multi_alias - should be split out +TEST_CASE(split_with_multi_alias_alive) +{ + migraphx::shape s{migraphx::shape::float_type, {2, 3, 327680}}; + migraphx::program p1; + { + auto* mm = p1.get_main_module(); + auto x = mm->add_parameter("x", s); + auto rsum = add_reduce( + p1, "fuse_reduce0", {x}, {2}, [&](auto* rm, const auto& inputs, const auto& axes) { + // Create a computation before the reduce + auto sqrt = + add_pointwise(p1, rm, "main:pointwise0", {inputs[0]}, single_pointwise("sqrt")); + auto rsum1 = + rm->add_instruction(migraphx::make_op("reduce_sum", {{"axes", axes}}), sqrt); + // multi_alias aliases sqrt and rsum1 - sqrt should be identified as alive + auto ma = rm->add_instruction(multi_alias_op{}, sqrt, rsum1); + auto rsumb = rm->add_instruction( + migraphx::make_op("multibroadcast", {{"out_lens", s.lens()}}), ma); + return add_pointwise( + p1, rm, "main:pointwise1", {rsumb, sqrt}, single_pointwise("mul")); + }); + mm->add_return({rsum}); + } + run_pass(p1); + migraphx::program p2; + { + auto* mm = p2.get_main_module(); + auto x = mm->add_parameter("x", s); + // sqrt is computed first, then passed to split module + auto sqrt = add_pointwise(p2, mm, "main:pointwise0", {x}, single_pointwise("sqrt")); + auto rsums = add_reduce( + p2, + "fuse_reduce0_split", + {sqrt}, + {2}, + "assign_add", + [&](auto* rm, const auto& inputs, const auto& axes) { + return rm->add_instruction(migraphx::make_op("reduce_sum", {{"axes", axes}}), + inputs[0]); + }); + // After split: multi_alias(sqrt, rsums) - shape is {2,3,327680} from sqrt + // multibroadcast is eliminated since multi_alias already has the right shape + auto ma = mm->add_instruction(multi_alias_op{}, sqrt, rsums); + // multiply multi_alias result with sqrt + auto result = add_pointwise(p2, mm, "main:pointwise1", {ma, sqrt}, single_pointwise("mul")); + mm->add_return({result}); + } + EXPECT(p1.sort() == p2.sort()); +} + int main(int argc, const char* argv[]) { test::run(argc, argv); } From 63bd6e88b4370cfbaaf6bd6dba6d46b32adcf6a1 Mon Sep 17 00:00:00 2001 From: Paul Date: Mon, 12 Jan 2026 20:08:53 +0000 Subject: [PATCH 11/23] Format --- test/split_reduce.cpp | 67 ++++++++++++++++++++++--------------------- 1 file changed, 34 insertions(+), 33 deletions(-) diff --git a/test/split_reduce.cpp b/test/split_reduce.cpp index 25de25f96fd..c15f716e06b 100644 --- a/test/split_reduce.cpp +++ b/test/split_reduce.cpp @@ -302,11 +302,11 @@ TEST_CASE(parallel_reduce_multi_alias) migraphx::shape s{migraphx::shape::float_type, {2, 3, 327680}}; migraphx::program p1; { - auto* mm = p1.get_main_module(); - auto x = mm->add_parameter("x", s); + auto* mm = p1.get_main_module(); + auto x = mm->add_parameter("x", s); auto rsum = add_reduce( p1, "fuse_reduce0", {x}, {2}, [&](auto* rm, const auto& inputs, const auto& axes) { - auto xx = add_pointwise(p1, rm, "main:pointwise0", {inputs[0]}, squared()); + auto xx = add_pointwise(p1, rm, "main:pointwise0", {inputs[0]}, squared()); auto rsum1 = rm->add_instruction(migraphx::make_op("reduce_sum", {{"axes", axes}}), inputs[0]); auto rsum2 = @@ -325,23 +325,24 @@ TEST_CASE(parallel_reduce_multi_alias) // The pointwise (squared) is extracted to main module auto xx = add_pointwise(p2, mm, "main:pointwise0", {x}, squared()); // Split module takes both xx and x as inputs - auto rsum = add_reduce( - p2, - "fuse_reduce0_split", - {xx, x}, - {2}, - "assign_add", - [&](auto* rm, - const auto& inputs, - const auto& axes) -> std::vector { - // inputs[0] is xx (squared), inputs[1] is x - // The pass returns (rsum1, rsum2) order based on the original fused module order - auto rsum1 = rm->add_instruction(migraphx::make_op("reduce_sum", {{"axes", axes}}), - inputs[1]); - auto rsum2 = rm->add_instruction(migraphx::make_op("reduce_sum", {{"axes", axes}}), - inputs[0]); - return {rsum1, rsum2}; - }); + auto rsum = + add_reduce(p2, + "fuse_reduce0_split", + {xx, x}, + {2}, + "assign_add", + [&](auto* rm, + const auto& inputs, + const auto& axes) -> std::vector { + // inputs[0] is xx (squared), inputs[1] is x + // The pass returns (rsum1, rsum2) order based on the original fused + // module order + auto rsum1 = rm->add_instruction( + migraphx::make_op("reduce_sum", {{"axes", axes}}), inputs[1]); + auto rsum2 = rm->add_instruction( + migraphx::make_op("reduce_sum", {{"axes", axes}}), inputs[0]); + return {rsum1, rsum2}; + }); auto rsum1 = mm->add_instruction(migraphx::make_op("get_tuple_elem", {{"index", 0}}), rsum); auto rsum2 = mm->add_instruction(migraphx::make_op("get_tuple_elem", {{"index", 1}}), rsum); // multi_alias_op is moved to main module after split @@ -368,7 +369,7 @@ TEST_CASE(split_with_multi_alias_alive) auto rsum1 = rm->add_instruction(migraphx::make_op("reduce_sum", {{"axes", axes}}), sqrt); // multi_alias aliases sqrt and rsum1 - sqrt should be identified as alive - auto ma = rm->add_instruction(multi_alias_op{}, sqrt, rsum1); + auto ma = rm->add_instruction(multi_alias_op{}, sqrt, rsum1); auto rsumb = rm->add_instruction( migraphx::make_op("multibroadcast", {{"out_lens", s.lens()}}), ma); return add_pointwise( @@ -379,20 +380,20 @@ TEST_CASE(split_with_multi_alias_alive) run_pass(p1); migraphx::program p2; { - auto* mm = p2.get_main_module(); - auto x = mm->add_parameter("x", s); + auto* mm = p2.get_main_module(); + auto x = mm->add_parameter("x", s); // sqrt is computed first, then passed to split module auto sqrt = add_pointwise(p2, mm, "main:pointwise0", {x}, single_pointwise("sqrt")); - auto rsums = add_reduce( - p2, - "fuse_reduce0_split", - {sqrt}, - {2}, - "assign_add", - [&](auto* rm, const auto& inputs, const auto& axes) { - return rm->add_instruction(migraphx::make_op("reduce_sum", {{"axes", axes}}), - inputs[0]); - }); + auto rsums = + add_reduce(p2, + "fuse_reduce0_split", + {sqrt}, + {2}, + "assign_add", + [&](auto* rm, const auto& inputs, const auto& axes) { + return rm->add_instruction( + migraphx::make_op("reduce_sum", {{"axes", axes}}), inputs[0]); + }); // After split: multi_alias(sqrt, rsums) - shape is {2,3,327680} from sqrt // multibroadcast is eliminated since multi_alias already has the right shape auto ma = mm->add_instruction(multi_alias_op{}, sqrt, rsums); From aafb7aa34bee0a50bd0cf55943e8a69b3f25994e Mon Sep 17 00:00:00 2001 From: Paul Date: Mon, 12 Jan 2026 17:25:27 -0600 Subject: [PATCH 12/23] Fix tidy issues --- src/driver/trim.cpp | 2 +- src/eliminate_concat.cpp | 1 + src/replace_allocate.cpp | 8 +++----- test/include/basic_ops.hpp | 5 ++--- test/split_reduce.cpp | 3 +-- 5 files changed, 8 insertions(+), 11 deletions(-) diff --git a/src/driver/trim.cpp b/src/driver/trim.cpp index f7c81658fb3..60355fbaa1b 100644 --- a/src/driver/trim.cpp +++ b/src/driver/trim.cpp @@ -38,7 +38,7 @@ inline namespace MIGRAPHX_INLINE_NS { static instruction_ref capture_arg(std::unordered_set& s, instruction_ref ins) { auto aliases = instruction::get_output_alias(ins, true); - if(not(aliases.size() == 1 and aliases.front() == ins)) + if(aliases.size() == 1 and aliases.front() != ins) { s.insert(ins); return capture_arg(s, aliases.front()); diff --git a/src/eliminate_concat.cpp b/src/eliminate_concat.cpp index 49b8d902f52..d523061e2f2 100644 --- a/src/eliminate_concat.cpp +++ b/src/eliminate_concat.cpp @@ -75,6 +75,7 @@ void eliminate_concat::apply(module& m) const std::back_inserter(allocations), [&](instruction_ref x) { auto aliases = instruction::get_output_alias(x, true); + // cppcheck-suppress returnDanglingLifetime return aliases.front(); }); diff --git a/src/replace_allocate.cpp b/src/replace_allocate.cpp index 19d83ea5572..3b2964420c3 100644 --- a/src/replace_allocate.cpp +++ b/src/replace_allocate.cpp @@ -46,11 +46,9 @@ std::unordered_map create_output_names(const modul for(const auto& i : returns) { auto aliases = instruction::get_output_alias(i); - for(auto ins : aliases) - { - if(ins->name() == "allocate") - alloc_aliases.push_back(ins); - } + std::copy_if(aliases.begin(), aliases.end(), std::back_inserter(alloc_aliases), [&](instruction_ref ins) { + return ins->name() == "allocate"; + }); } std::size_t index = 0; diff --git a/test/include/basic_ops.hpp b/test/include/basic_ops.hpp index 4e2664c6826..5424bba9562 100644 --- a/test/include/basic_ops.hpp +++ b/test/include/basic_ops.hpp @@ -268,9 +268,8 @@ struct multi_alias_op } std::vector output_alias(const std::vector& s) const { - std::vector result; - for(std::size_t i = 0; i < s.size(); ++i) - result.push_back(i); + std::vector result(s.size()); + std::iota(result.begin(), result.end(), 0); return result; } }; diff --git a/test/split_reduce.cpp b/test/split_reduce.cpp index c15f716e06b..09827d780e9 100644 --- a/test/split_reduce.cpp +++ b/test/split_reduce.cpp @@ -312,8 +312,7 @@ TEST_CASE(parallel_reduce_multi_alias) auto rsum2 = rm->add_instruction(migraphx::make_op("reduce_sum", {{"axes", axes}}), xx); // multi_alias_op aliases both reduce outputs - auto ma = rm->add_instruction(multi_alias_op{}, rsum1, rsum2); - return ma; + return rm->add_instruction(multi_alias_op{}, rsum1, rsum2); }); mm->add_return({rsum}); } From a8d1f637547ffa6d05c468e2a721e02cc3fd6a73 Mon Sep 17 00:00:00 2001 From: Paul Date: Mon, 12 Jan 2026 17:25:30 -0600 Subject: [PATCH 13/23] Format --- src/replace_allocate.cpp | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/src/replace_allocate.cpp b/src/replace_allocate.cpp index 3b2964420c3..39dfe0b5011 100644 --- a/src/replace_allocate.cpp +++ b/src/replace_allocate.cpp @@ -46,9 +46,10 @@ std::unordered_map create_output_names(const modul for(const auto& i : returns) { auto aliases = instruction::get_output_alias(i); - std::copy_if(aliases.begin(), aliases.end(), std::back_inserter(alloc_aliases), [&](instruction_ref ins) { - return ins->name() == "allocate"; - }); + std::copy_if(aliases.begin(), + aliases.end(), + std::back_inserter(alloc_aliases), + [&](instruction_ref ins) { return ins->name() == "allocate"; }); } std::size_t index = 0; From c2fcfe5502a3df2e06c39a2f967a576fdd098b5b Mon Sep 17 00:00:00 2001 From: Paul Date: Tue, 13 Jan 2026 11:01:24 -0600 Subject: [PATCH 14/23] Fix offload detection --- src/driver/perf.cpp | 18 +++++++++--------- 1 file changed, 9 insertions(+), 9 deletions(-) diff --git a/src/driver/perf.cpp b/src/driver/perf.cpp index aa3b8d16868..da1af7c3873 100644 --- a/src/driver/perf.cpp +++ b/src/driver/perf.cpp @@ -125,17 +125,17 @@ bool is_offload_copy_set(const program& p) else if(i.name() == "@return") { auto return_args = i.inputs(); - for(const auto& j : return_args) - { + return std::all_of(return_args.begin(), return_args.end(), [&](const auto& j) { + auto aliases = instruction::get_output_alias(j, true); - bool valid = - std::any_of(aliases.begin(), aliases.end(), [&](instruction_ref alias_ins) { - return (alias_ins->name() == "@param" and param_ins.erase(alias_ins) > 0) or - alias_ins->name() == "hip::copy_from_gpu"; - }); - if(not valid) + return std::all_of(aliases.begin(), aliases.end(), [&](instruction_ref alias_ins) { + if(alias_ins->name() == "hip::copy_from_gpu") + return true; + if(alias_ins->name() == "@param") + return not contains(param_ins, alias_ins); return false; - } + }); + }); } } return param_ins.empty(); From 5cdabadb50befaabb363e3b9a689782cdf1fb479 Mon Sep 17 00:00:00 2001 From: Paul Date: Tue, 13 Jan 2026 11:01:26 -0600 Subject: [PATCH 15/23] Format --- src/driver/perf.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/src/driver/perf.cpp b/src/driver/perf.cpp index da1af7c3873..32e8e483d21 100644 --- a/src/driver/perf.cpp +++ b/src/driver/perf.cpp @@ -126,13 +126,12 @@ bool is_offload_copy_set(const program& p) { auto return_args = i.inputs(); return std::all_of(return_args.begin(), return_args.end(), [&](const auto& j) { - auto aliases = instruction::get_output_alias(j, true); return std::all_of(aliases.begin(), aliases.end(), [&](instruction_ref alias_ins) { if(alias_ins->name() == "hip::copy_from_gpu") return true; if(alias_ins->name() == "@param") - return not contains(param_ins, alias_ins); + return not contains(param_ins, alias_ins); return false; }); }); From 91aac43eadde90f8773f47e2db692f8d8bfa7b76 Mon Sep 17 00:00:00 2001 From: Paul Date: Tue, 13 Jan 2026 11:08:27 -0600 Subject: [PATCH 16/23] Some cleanup --- src/include/migraphx/liveness.hpp | 12 +++++------- test/output_alias.cpp | 14 +++++++------- 2 files changed, 12 insertions(+), 14 deletions(-) diff --git a/src/include/migraphx/liveness.hpp b/src/include/migraphx/liveness.hpp index eb98bed1cd1..45622f9dc9d 100644 --- a/src/include/migraphx/liveness.hpp +++ b/src/include/migraphx/liveness.hpp @@ -54,13 +54,11 @@ void liveness(const module& m, F f) for(auto input : inputs) { auto aliases = instruction::get_output_alias(input); - for(auto i : aliases) - { - // Skip if variable comes from parent - if(not m.has_instruction(i)) - continue; - live_set.insert(i); - } + // Skip if variable comes from parent + std::copy_if(aliases.begin(), + aliases.end(), + std::inserter(live_set, live_set.end()), + [&](auto i) { return m.has_instruction(i); }); } }; add_live_variables(ins->inputs()); diff --git a/test/output_alias.cpp b/test/output_alias.cpp index d56d8dfd59d..346e2e07fb8 100644 --- a/test/output_alias.cpp +++ b/test/output_alias.cpp @@ -33,8 +33,8 @@ TEST_CASE(simple_alias) auto* mm = p.get_main_module(); auto l = mm->add_literal(1); auto p1 = mm->add_instruction(pass_op{}, l); - EXPECT(migraphx::contains(migraphx::instruction::get_output_alias(l), l)); - EXPECT(migraphx::contains(migraphx::instruction::get_output_alias(p1), l)); + EXPECT(migraphx::instruction::get_output_alias(l) == std::vector{l}); + EXPECT(migraphx::instruction::get_output_alias(p1) == std::vector{l}); } TEST_CASE(cascade_alias) @@ -45,10 +45,10 @@ TEST_CASE(cascade_alias) auto p1 = mm->add_instruction(pass_op{}, l); auto p2 = mm->add_instruction(pass_op{}, p1); auto p3 = mm->add_instruction(pass_op{}, p2); - EXPECT(migraphx::contains(migraphx::instruction::get_output_alias(l), l)); - EXPECT(migraphx::contains(migraphx::instruction::get_output_alias(p1), l)); - EXPECT(migraphx::contains(migraphx::instruction::get_output_alias(p2), l)); - EXPECT(migraphx::contains(migraphx::instruction::get_output_alias(p3), l)); + EXPECT(migraphx::instruction::get_output_alias(l) == std::vector{l}); + EXPECT(migraphx::instruction::get_output_alias(p1) == std::vector{l}); + EXPECT(migraphx::instruction::get_output_alias(p2) == std::vector{l}); + EXPECT(migraphx::instruction::get_output_alias(p3) == std::vector{l}); } TEST_CASE(no_alias) @@ -58,7 +58,7 @@ TEST_CASE(no_alias) auto x = mm->add_literal(1); auto y = mm->add_literal(2); auto sum = mm->add_instruction(sum_op{}, x, y); - EXPECT(migraphx::contains(migraphx::instruction::get_output_alias(sum), sum)); + EXPECT(migraphx::instruction::get_output_alias(sum) == std::vector{sum}); } TEST_CASE(multiple_aliases) From 11c34dba96521467b7538472bb293b63aa5552a0 Mon Sep 17 00:00:00 2001 From: Paul Date: Tue, 13 Jan 2026 11:39:11 -0600 Subject: [PATCH 17/23] Add liveness tests --- test/liveness_test.cpp | 133 +++++++++++++++++++++++++++++++++++++++++ 1 file changed, 133 insertions(+) create mode 100644 test/liveness_test.cpp diff --git a/test/liveness_test.cpp b/test/liveness_test.cpp new file mode 100644 index 00000000000..9bfe2b43308 --- /dev/null +++ b/test/liveness_test.cpp @@ -0,0 +1,133 @@ +/* + * The MIT License (MIT) + * + * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * + * Permission is hereby granted, free of charge, to any person obtaining a copy + * of this software and associated documentation files (the "Software"), to deal + * in the Software without restriction, including without limitation the rights + * to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + * copies of the Software, and to permit persons to whom the Software is + * furnished to do so, subject to the following conditions: + * + * The above copyright notice and this permission notice shall be included in + * all copies or substantial portions of the Software. + * + * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + */ +#include +#include +#include +#include +#include +#include + +TEST_CASE(liveness_single_alias) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + auto x = mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {1.0f, 2.0f}}); + auto p1 = mm->add_instruction(pass_op{}, x); + auto p2 = mm->add_instruction(pass_op{}, p1); + mm->add_return({p2}); + + std::vector consumed; + migraphx::liveness(*mm, [&](auto ins, const auto&) { + consumed.push_back(ins); + }); + + // With single alias ops, pass_ops alias to x, so liveness tracks x + // The callback is called when x is "consumed" (last usage) + EXPECT(migraphx::contains(consumed, x)); +} + +TEST_CASE(liveness_multi_alias) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + auto x = mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {1.0f, 2.0f}}); + auto y = mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {3.0f, 4.0f}}); + // multi_alias_op aliases both x and y + auto ma = mm->add_instruction(multi_alias_op{}, x, y); + auto p1 = mm->add_instruction(pass_op{}, ma); + mm->add_return({p1}); + + std::vector consumed; + migraphx::liveness(*mm, [&](auto ins, const auto&) { + consumed.push_back(ins); + }); + + // Both x and y should be tracked and consumed + // because multi_alias_op aliases both inputs + EXPECT(migraphx::contains(consumed, x)); + EXPECT(migraphx::contains(consumed, y)); +} + +TEST_CASE(liveness_multi_alias_cascade) +{ + migraphx::program p; + auto* mm = p.get_main_module(); + auto x = mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {1.0f, 2.0f}}); + auto y = mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {3.0f, 4.0f}}); + auto z = mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {5.0f, 6.0f}}); + // First multi_alias aliases x and y + auto ma1 = mm->add_instruction(multi_alias_op{}, x, y); + // Second multi_alias aliases ma1 (which aliases x,y) and z + auto ma2 = mm->add_instruction(multi_alias_op{}, ma1, z); + mm->add_return({ma2}); + + std::vector consumed; + migraphx::liveness(*mm, [&](auto ins, const auto&) { + consumed.push_back(ins); + }); + + // All three literals should be tracked and consumed + // ma2 transitively aliases x, y, z + EXPECT(migraphx::contains(consumed, x)); + EXPECT(migraphx::contains(consumed, y)); + EXPECT(migraphx::contains(consumed, z)); +} + +TEST_CASE(liveness_multi_alias_both_tracked) +{ + // This test verifies that when multi_alias_op aliases multiple inputs, + // ALL aliased instructions are properly tracked in liveness analysis. + migraphx::program p; + auto* mm = p.get_main_module(); + auto x = mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {1.0f, 2.0f}}); + auto y = mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {3.0f, 4.0f}}); + // multi_alias_op returns output_alias {0, 1} - it aliases both inputs + auto ma = mm->add_instruction(multi_alias_op{}, x, y); + mm->add_return({ma}); + + // Count how many times each literal appears in any live_set across all callbacks + std::size_t x_live_count = 0; + std::size_t y_live_count = 0; + migraphx::liveness(*mm, [&](auto, const auto& live_set) { + if(migraphx::contains(live_set, x)) + x_live_count++; + if(migraphx::contains(live_set, y)) + y_live_count++; + }); + + // Both x and y should appear as live at some point during liveness analysis + // (because multi_alias_op properly exposes both as aliases) + // Note: they might have count 0 if they're only processed when the live_set is already emptied + // The key test is that BOTH get consumed (callback called for both) + std::vector consumed; + migraphx::liveness(*mm, [&](auto ins, const auto&) { + consumed.push_back(ins); + }); + + EXPECT(migraphx::contains(consumed, x)); + EXPECT(migraphx::contains(consumed, y)); +} + +int main(int argc, const char* argv[]) { test::run(argc, argv); } + From ededd25f62c21003143af2344124ec98a7d948f3 Mon Sep 17 00:00:00 2001 From: Paul Date: Tue, 13 Jan 2026 11:39:15 -0600 Subject: [PATCH 18/23] Format --- test/liveness_test.cpp | 51 +++++++++++++++++++++--------------------- 1 file changed, 25 insertions(+), 26 deletions(-) diff --git a/test/liveness_test.cpp b/test/liveness_test.cpp index 9bfe2b43308..45e47b5381b 100644 --- a/test/liveness_test.cpp +++ b/test/liveness_test.cpp @@ -32,15 +32,14 @@ TEST_CASE(liveness_single_alias) { migraphx::program p; auto* mm = p.get_main_module(); - auto x = mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {1.0f, 2.0f}}); - auto p1 = mm->add_instruction(pass_op{}, x); - auto p2 = mm->add_instruction(pass_op{}, p1); + auto x = mm->add_literal( + migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {1.0f, 2.0f}}); + auto p1 = mm->add_instruction(pass_op{}, x); + auto p2 = mm->add_instruction(pass_op{}, p1); mm->add_return({p2}); std::vector consumed; - migraphx::liveness(*mm, [&](auto ins, const auto&) { - consumed.push_back(ins); - }); + migraphx::liveness(*mm, [&](auto ins, const auto&) { consumed.push_back(ins); }); // With single alias ops, pass_ops alias to x, so liveness tracks x // The callback is called when x is "consumed" (last usage) @@ -51,17 +50,17 @@ TEST_CASE(liveness_multi_alias) { migraphx::program p; auto* mm = p.get_main_module(); - auto x = mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {1.0f, 2.0f}}); - auto y = mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {3.0f, 4.0f}}); + auto x = mm->add_literal( + migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {1.0f, 2.0f}}); + auto y = mm->add_literal( + migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {3.0f, 4.0f}}); // multi_alias_op aliases both x and y - auto ma = mm->add_instruction(multi_alias_op{}, x, y); - auto p1 = mm->add_instruction(pass_op{}, ma); + auto ma = mm->add_instruction(multi_alias_op{}, x, y); + auto p1 = mm->add_instruction(pass_op{}, ma); mm->add_return({p1}); std::vector consumed; - migraphx::liveness(*mm, [&](auto ins, const auto&) { - consumed.push_back(ins); - }); + migraphx::liveness(*mm, [&](auto ins, const auto&) { consumed.push_back(ins); }); // Both x and y should be tracked and consumed // because multi_alias_op aliases both inputs @@ -73,9 +72,12 @@ TEST_CASE(liveness_multi_alias_cascade) { migraphx::program p; auto* mm = p.get_main_module(); - auto x = mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {1.0f, 2.0f}}); - auto y = mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {3.0f, 4.0f}}); - auto z = mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {5.0f, 6.0f}}); + auto x = mm->add_literal( + migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {1.0f, 2.0f}}); + auto y = mm->add_literal( + migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {3.0f, 4.0f}}); + auto z = mm->add_literal( + migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {5.0f, 6.0f}}); // First multi_alias aliases x and y auto ma1 = mm->add_instruction(multi_alias_op{}, x, y); // Second multi_alias aliases ma1 (which aliases x,y) and z @@ -83,9 +85,7 @@ TEST_CASE(liveness_multi_alias_cascade) mm->add_return({ma2}); std::vector consumed; - migraphx::liveness(*mm, [&](auto ins, const auto&) { - consumed.push_back(ins); - }); + migraphx::liveness(*mm, [&](auto ins, const auto&) { consumed.push_back(ins); }); // All three literals should be tracked and consumed // ma2 transitively aliases x, y, z @@ -100,10 +100,12 @@ TEST_CASE(liveness_multi_alias_both_tracked) // ALL aliased instructions are properly tracked in liveness analysis. migraphx::program p; auto* mm = p.get_main_module(); - auto x = mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {1.0f, 2.0f}}); - auto y = mm->add_literal(migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {3.0f, 4.0f}}); + auto x = mm->add_literal( + migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {1.0f, 2.0f}}); + auto y = mm->add_literal( + migraphx::literal{migraphx::shape{migraphx::shape::float_type, {2}}, {3.0f, 4.0f}}); // multi_alias_op returns output_alias {0, 1} - it aliases both inputs - auto ma = mm->add_instruction(multi_alias_op{}, x, y); + auto ma = mm->add_instruction(multi_alias_op{}, x, y); mm->add_return({ma}); // Count how many times each literal appears in any live_set across all callbacks @@ -121,13 +123,10 @@ TEST_CASE(liveness_multi_alias_both_tracked) // Note: they might have count 0 if they're only processed when the live_set is already emptied // The key test is that BOTH get consumed (callback called for both) std::vector consumed; - migraphx::liveness(*mm, [&](auto ins, const auto&) { - consumed.push_back(ins); - }); + migraphx::liveness(*mm, [&](auto ins, const auto&) { consumed.push_back(ins); }); EXPECT(migraphx::contains(consumed, x)); EXPECT(migraphx::contains(consumed, y)); } int main(int argc, const char* argv[]) { test::run(argc, argv); } - From 49120716909ecce9752a805c97fe7d2ba967a380 Mon Sep 17 00:00:00 2001 From: Paul Date: Tue, 13 Jan 2026 17:16:11 -0600 Subject: [PATCH 19/23] Fix naming of index --- src/replace_allocate.cpp | 23 +++++++---------------- test/replace_allocate.cpp | 39 +++++++++++++++++++++++++++++++-------- 2 files changed, 38 insertions(+), 24 deletions(-) diff --git a/src/replace_allocate.cpp b/src/replace_allocate.cpp index 39dfe0b5011..9d69241ceef 100644 --- a/src/replace_allocate.cpp +++ b/src/replace_allocate.cpp @@ -29,6 +29,7 @@ #include #include #include +#include #include #include @@ -43,17 +44,13 @@ std::unordered_map create_output_names(const modul // Collect all allocation aliases from each return value std::vector alloc_aliases; - for(const auto& i : returns) - { - auto aliases = instruction::get_output_alias(i); - std::copy_if(aliases.begin(), - aliases.end(), - std::back_inserter(alloc_aliases), - [&](instruction_ref ins) { return ins->name() == "allocate"; }); - } + // Use a join but perhaps a tuple output parameter might be better? + std::transform(returns.begin(), returns.end(), join_back_inserter(alloc_aliases), [](const auto& i) { + return instruction::get_output_alias(i); + }); std::size_t index = 0; - if(returns.size() == 1 and mod.name().empty()) + if(mod.name().empty()) { // Single return with empty module name: all aliases get "output" or "output_N" if(alloc_aliases.size() == 1) @@ -69,13 +66,6 @@ std::unordered_map create_output_names(const modul } } // Preserve main module output buffer naming across migraphx versions - else if(mod.name() == "main") - { - for(auto ins : alloc_aliases) - { - mod_output_names[ins] = mod.name() + ":#output_" + std::to_string(index++); - } - } else { for(auto ins : alloc_aliases) @@ -83,6 +73,7 @@ std::unordered_map create_output_names(const modul mod_output_names[ins] = param_name(index++, mod.name() + ":#output_"); } } + return mod_output_names; } diff --git a/test/replace_allocate.cpp b/test/replace_allocate.cpp index e13c048a163..cf72a1d7884 100644 --- a/test/replace_allocate.cpp +++ b/test/replace_allocate.cpp @@ -279,6 +279,29 @@ TEST_CASE(allocate_copy_with_no_out) EXPECT(m1.sort() == m2.sort()); } +TEST_CASE(allocate_out_multi_return_partial_alloc) +{ + migraphx::shape s{migraphx::shape::float_type, {5}}; + migraphx::module m1; + { + auto x = m1.add_parameter("x", s); + auto alloc = + m1.add_instruction(migraphx::make_op("allocate", {{"shape", migraphx::to_value(s)}})); + auto p1 = m1.add_instruction(pass_op{}, alloc); + m1.add_return({x, p1}); + } + run_pass(m1, allocation_with_out_model{}); + + migraphx::module m2; + { + auto x = m2.add_parameter("x", s); + auto output = m2.add_parameter("output_1", s); + auto p1 = m2.add_instruction(pass_op{}, output); + m2.add_return({x, p1}); + } + EXPECT(m1.sort() == m2.sort()); +} + // Test that replace_allocate handles multi-alias operations correctly // when checking for shape matches (insert_copy code path) TEST_CASE(multi_alias_shape_check) @@ -324,7 +347,7 @@ TEST_CASE(multi_alias_alloc_out_param) { auto x = m2.add_parameter("x", s); // Only the allocation becomes an output parameter (named "output" since single alloc) - auto output = m2.add_parameter("output", s); + auto output = m2.add_parameter("output_0", s); auto p1 = m2.add_instruction(pass_op{}, output); auto ma = m2.add_instruction(multi_alias_op{}, p1, x); m2.add_return({ma}); @@ -414,12 +437,12 @@ TEST_CASE(multi_alias_multiple_outputs_out_params) { auto x = m2.add_parameter("x", s); auto y = m2.add_parameter("y", s); - // First output parameter replaces first allocation (named :#output_0) - auto output0 = m2.add_parameter(":#output_0", s); + // First output parameter replaces first allocation (named output_0) + auto output0 = m2.add_parameter("output_0", s); auto p1 = m2.add_instruction(pass_op{}, output0); auto ma1 = m2.add_instruction(multi_alias_op{}, p1, x); - // Second output parameter replaces second allocation (named :#output_1) - auto output1 = m2.add_parameter(":#output_1", s); + // Second output parameter replaces second allocation (named output_1) + auto output1 = m2.add_parameter("output_2", s); auto p2 = m2.add_instruction(pass_op{}, output1); auto ma2 = m2.add_instruction(multi_alias_op{}, p2, y); m2.add_return({ma1, ma2}); @@ -495,11 +518,11 @@ TEST_CASE(multi_alias_chain_multiple_out_params) { auto x = m2.add_parameter("x", s); // Three output parameters - one for each return - auto output0 = m2.add_parameter(":#output_0", s); + auto output0 = m2.add_parameter("output_0", s); auto p1 = m2.add_instruction(pass_op{}, output0); - auto output1 = m2.add_parameter(":#output_1", s); + auto output1 = m2.add_parameter("output_2", s); auto p2 = m2.add_instruction(pass_op{}, output1); - auto output2 = m2.add_parameter(":#output_2", s); + auto output2 = m2.add_parameter("output_4", s); auto p3 = m2.add_instruction(pass_op{}, output2); auto ma1 = m2.add_instruction(multi_alias_op{}, p1, x); auto ma2 = m2.add_instruction(multi_alias_op{}, p2, x); From 2f32aae461d73c2d6e2425a4eeb6a5070674dac8 Mon Sep 17 00:00:00 2001 From: Paul Date: Tue, 13 Jan 2026 17:16:16 -0600 Subject: [PATCH 20/23] Format --- src/replace_allocate.cpp | 7 ++++--- test/replace_allocate.cpp | 2 +- 2 files changed, 5 insertions(+), 4 deletions(-) diff --git a/src/replace_allocate.cpp b/src/replace_allocate.cpp index 9d69241ceef..a832e21bdff 100644 --- a/src/replace_allocate.cpp +++ b/src/replace_allocate.cpp @@ -45,9 +45,10 @@ std::unordered_map create_output_names(const modul // Collect all allocation aliases from each return value std::vector alloc_aliases; // Use a join but perhaps a tuple output parameter might be better? - std::transform(returns.begin(), returns.end(), join_back_inserter(alloc_aliases), [](const auto& i) { - return instruction::get_output_alias(i); - }); + std::transform(returns.begin(), + returns.end(), + join_back_inserter(alloc_aliases), + [](const auto& i) { return instruction::get_output_alias(i); }); std::size_t index = 0; if(mod.name().empty()) diff --git a/test/replace_allocate.cpp b/test/replace_allocate.cpp index cf72a1d7884..ccf927f2fbb 100644 --- a/test/replace_allocate.cpp +++ b/test/replace_allocate.cpp @@ -294,7 +294,7 @@ TEST_CASE(allocate_out_multi_return_partial_alloc) migraphx::module m2; { - auto x = m2.add_parameter("x", s); + auto x = m2.add_parameter("x", s); auto output = m2.add_parameter("output_1", s); auto p1 = m2.add_instruction(pass_op{}, output); m2.add_return({x, p1}); From be5bc74e5c2ac90207da15df9881f8569994d704 Mon Sep 17 00:00:00 2001 From: Paul Date: Tue, 13 Jan 2026 19:45:17 -0600 Subject: [PATCH 21/23] Fix ctad warning --- test/output_alias.cpp | 16 +++++++++------- 1 file changed, 9 insertions(+), 7 deletions(-) diff --git a/test/output_alias.cpp b/test/output_alias.cpp index 346e2e07fb8..871169f2304 100644 --- a/test/output_alias.cpp +++ b/test/output_alias.cpp @@ -27,14 +27,16 @@ #include #include +using instruction_refs = std::vector; + TEST_CASE(simple_alias) { migraphx::program p; auto* mm = p.get_main_module(); auto l = mm->add_literal(1); auto p1 = mm->add_instruction(pass_op{}, l); - EXPECT(migraphx::instruction::get_output_alias(l) == std::vector{l}); - EXPECT(migraphx::instruction::get_output_alias(p1) == std::vector{l}); + EXPECT(migraphx::instruction::get_output_alias(l) == instruction_refs{l}); + EXPECT(migraphx::instruction::get_output_alias(p1) == instruction_refs{l}); } TEST_CASE(cascade_alias) @@ -45,10 +47,10 @@ TEST_CASE(cascade_alias) auto p1 = mm->add_instruction(pass_op{}, l); auto p2 = mm->add_instruction(pass_op{}, p1); auto p3 = mm->add_instruction(pass_op{}, p2); - EXPECT(migraphx::instruction::get_output_alias(l) == std::vector{l}); - EXPECT(migraphx::instruction::get_output_alias(p1) == std::vector{l}); - EXPECT(migraphx::instruction::get_output_alias(p2) == std::vector{l}); - EXPECT(migraphx::instruction::get_output_alias(p3) == std::vector{l}); + EXPECT(migraphx::instruction::get_output_alias(l) == instruction_refs{l}); + EXPECT(migraphx::instruction::get_output_alias(p1) == instruction_refs{l}); + EXPECT(migraphx::instruction::get_output_alias(p2) == instruction_refs{l}); + EXPECT(migraphx::instruction::get_output_alias(p3) == instruction_refs{l}); } TEST_CASE(no_alias) @@ -58,7 +60,7 @@ TEST_CASE(no_alias) auto x = mm->add_literal(1); auto y = mm->add_literal(2); auto sum = mm->add_instruction(sum_op{}, x, y); - EXPECT(migraphx::instruction::get_output_alias(sum) == std::vector{sum}); + EXPECT(migraphx::instruction::get_output_alias(sum) == instruction_refs{sum}); } TEST_CASE(multiple_aliases) From 352301e484fffe8f3d81a03c6339189b0ba766dd Mon Sep 17 00:00:00 2001 From: Paul Date: Tue, 13 Jan 2026 19:45:34 -0600 Subject: [PATCH 22/23] Update license --- src/adjust_allocation.cpp | 2 +- src/api/api.cpp | 2 +- src/driver/perf.cpp | 2 +- src/driver/trim.cpp | 2 +- src/eliminate_concat.cpp | 2 +- src/graphviz.cpp | 2 +- src/include/migraphx/instruction.hpp | 2 +- src/include/migraphx/liveness.hpp | 2 +- src/include/migraphx/op/as_shape.hpp | 2 +- src/include/migraphx/op/broadcast.hpp | 2 +- src/include/migraphx/op/broadcast_for_dot.hpp | 2 +- src/include/migraphx/op/capture.hpp | 2 +- src/include/migraphx/op/fill.hpp | 2 +- src/include/migraphx/op/get_tuple_elem.hpp | 2 +- src/include/migraphx/op/identity.hpp | 2 +- src/include/migraphx/op/load.hpp | 2 +- src/include/migraphx/op/multibroadcast.hpp | 2 +- src/include/migraphx/op/random_uniform.hpp | 2 +- src/include/migraphx/op/reshape_lazy.hpp | 2 +- src/include/migraphx/op/scalar.hpp | 2 +- src/include/migraphx/op/select_module.hpp | 2 +- src/include/migraphx/op/slice.hpp | 2 +- src/include/migraphx/op/squeeze.hpp | 2 +- src/include/migraphx/op/step.hpp | 2 +- src/include/migraphx/op/transpose.hpp | 2 +- src/include/migraphx/op/unsqueeze.hpp | 2 +- src/include/migraphx/operation.hpp | 2 +- src/instruction.cpp | 2 +- src/module.cpp | 2 +- src/propagate_constant.cpp | 2 +- src/replace_allocate.cpp | 2 +- src/simplify_algebra.cpp | 2 +- src/split_reduce.cpp | 2 +- src/targets/cpu/copy.cpp | 2 +- src/targets/cpu/gather.cpp | 2 +- src/targets/cpu/include/migraphx/cpu/dnnl.hpp | 2 +- src/targets/cpu/include/migraphx/cpu/pointwise.hpp | 2 +- src/targets/gpu/compile_miopen.cpp | 2 +- src/targets/gpu/compile_ops.cpp | 2 +- src/targets/gpu/fuse_ops.cpp | 2 +- src/targets/gpu/include/migraphx/gpu/abs.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/argmax.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/argmin.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/code_object_op.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/compile_hipblaslt.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/convolution.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/fixed_pad.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/gemm.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/hip.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/hip_gemm.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/logsoftmax.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/loop.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/lrn.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/multinomial.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/nonzero.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/oper.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/pooling.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/prefix_scan_sum.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/reduce_op.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/reverse.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/rnn_variable_seq_lens.hpp | 2 +- src/targets/gpu/include/migraphx/gpu/topk.hpp | 2 +- src/targets/gpu/jit/mlir.cpp | 2 +- src/targets/gpu/lowering.cpp | 2 +- test/api/test_custom_op.cpp | 2 +- test/eliminate_concat_test.cpp | 2 +- test/eval_test.cpp | 2 +- test/include/basic_ops.hpp | 2 +- test/liveness_test.cpp | 2 +- test/operation.cpp | 2 +- test/output_alias.cpp | 2 +- test/propagate_constant_test.cpp | 2 +- test/replace_allocate.cpp | 2 +- test/run_loop_test.cpp | 2 +- test/schedule_test.cpp | 2 +- test/split_reduce.cpp | 2 +- tools/api/api.cpp | 2 +- tools/include/operation.hpp | 2 +- 78 files changed, 78 insertions(+), 78 deletions(-) diff --git a/src/adjust_allocation.cpp b/src/adjust_allocation.cpp index e95d7f53443..eaa3689537c 100644 --- a/src/adjust_allocation.cpp +++ b/src/adjust_allocation.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/api/api.cpp b/src/api/api.cpp index 512343bd0fc..901222482d1 100644 --- a/src/api/api.cpp +++ b/src/api/api.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/driver/perf.cpp b/src/driver/perf.cpp index 32e8e483d21..fdd20331163 100644 --- a/src/driver/perf.cpp +++ b/src/driver/perf.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/driver/trim.cpp b/src/driver/trim.cpp index 60355fbaa1b..4286e5a44ff 100644 --- a/src/driver/trim.cpp +++ b/src/driver/trim.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/eliminate_concat.cpp b/src/eliminate_concat.cpp index d523061e2f2..7f9399777b3 100644 --- a/src/eliminate_concat.cpp +++ b/src/eliminate_concat.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/graphviz.cpp b/src/graphviz.cpp index 3661adc87ea..3726846a994 100644 --- a/src/graphviz.cpp +++ b/src/graphviz.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/instruction.hpp b/src/include/migraphx/instruction.hpp index 2691d46dbc6..a12f6508c81 100644 --- a/src/include/migraphx/instruction.hpp +++ b/src/include/migraphx/instruction.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/liveness.hpp b/src/include/migraphx/liveness.hpp index 45622f9dc9d..9a3418c8307 100644 --- a/src/include/migraphx/liveness.hpp +++ b/src/include/migraphx/liveness.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/op/as_shape.hpp b/src/include/migraphx/op/as_shape.hpp index a6398e06e58..7618451a317 100644 --- a/src/include/migraphx/op/as_shape.hpp +++ b/src/include/migraphx/op/as_shape.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/op/broadcast.hpp b/src/include/migraphx/op/broadcast.hpp index fd31e73b34f..f6da43940b0 100644 --- a/src/include/migraphx/op/broadcast.hpp +++ b/src/include/migraphx/op/broadcast.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/op/broadcast_for_dot.hpp b/src/include/migraphx/op/broadcast_for_dot.hpp index 40a1f21e618..96897b991a5 100644 --- a/src/include/migraphx/op/broadcast_for_dot.hpp +++ b/src/include/migraphx/op/broadcast_for_dot.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/op/capture.hpp b/src/include/migraphx/op/capture.hpp index 8a1c33a2c68..fe036b304af 100644 --- a/src/include/migraphx/op/capture.hpp +++ b/src/include/migraphx/op/capture.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/op/fill.hpp b/src/include/migraphx/op/fill.hpp index 80279b5b6d9..ae3d8e784b4 100644 --- a/src/include/migraphx/op/fill.hpp +++ b/src/include/migraphx/op/fill.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/op/get_tuple_elem.hpp b/src/include/migraphx/op/get_tuple_elem.hpp index ca1d14246a0..30d5da2b7d7 100644 --- a/src/include/migraphx/op/get_tuple_elem.hpp +++ b/src/include/migraphx/op/get_tuple_elem.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/op/identity.hpp b/src/include/migraphx/op/identity.hpp index b76397e2778..02a53f1d308 100644 --- a/src/include/migraphx/op/identity.hpp +++ b/src/include/migraphx/op/identity.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/op/load.hpp b/src/include/migraphx/op/load.hpp index e85e7fa8471..a37e9448c9a 100644 --- a/src/include/migraphx/op/load.hpp +++ b/src/include/migraphx/op/load.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/op/multibroadcast.hpp b/src/include/migraphx/op/multibroadcast.hpp index b1076661171..735a99d4b54 100644 --- a/src/include/migraphx/op/multibroadcast.hpp +++ b/src/include/migraphx/op/multibroadcast.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/op/random_uniform.hpp b/src/include/migraphx/op/random_uniform.hpp index 642d35ff082..73ae1232b65 100644 --- a/src/include/migraphx/op/random_uniform.hpp +++ b/src/include/migraphx/op/random_uniform.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/op/reshape_lazy.hpp b/src/include/migraphx/op/reshape_lazy.hpp index 6aab475140d..773f6c345b5 100644 --- a/src/include/migraphx/op/reshape_lazy.hpp +++ b/src/include/migraphx/op/reshape_lazy.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/op/scalar.hpp b/src/include/migraphx/op/scalar.hpp index 852aca4e551..005712bbe69 100644 --- a/src/include/migraphx/op/scalar.hpp +++ b/src/include/migraphx/op/scalar.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/op/select_module.hpp b/src/include/migraphx/op/select_module.hpp index a16336413bb..a87fb089941 100644 --- a/src/include/migraphx/op/select_module.hpp +++ b/src/include/migraphx/op/select_module.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/op/slice.hpp b/src/include/migraphx/op/slice.hpp index 5dd4eaf55c1..fa35d640849 100644 --- a/src/include/migraphx/op/slice.hpp +++ b/src/include/migraphx/op/slice.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/op/squeeze.hpp b/src/include/migraphx/op/squeeze.hpp index 829a4a7bcfe..9ea891c2b96 100644 --- a/src/include/migraphx/op/squeeze.hpp +++ b/src/include/migraphx/op/squeeze.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/op/step.hpp b/src/include/migraphx/op/step.hpp index 195402e877e..660c0b2744f 100644 --- a/src/include/migraphx/op/step.hpp +++ b/src/include/migraphx/op/step.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/op/transpose.hpp b/src/include/migraphx/op/transpose.hpp index 60f7c37c213..bb508ecec76 100644 --- a/src/include/migraphx/op/transpose.hpp +++ b/src/include/migraphx/op/transpose.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/op/unsqueeze.hpp b/src/include/migraphx/op/unsqueeze.hpp index ca132317b97..06446b23ffb 100644 --- a/src/include/migraphx/op/unsqueeze.hpp +++ b/src/include/migraphx/op/unsqueeze.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/include/migraphx/operation.hpp b/src/include/migraphx/operation.hpp index ac22cd6a425..82cb9f9cdd1 100644 --- a/src/include/migraphx/operation.hpp +++ b/src/include/migraphx/operation.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/instruction.cpp b/src/instruction.cpp index c10cb9c2da7..a9f032b1cf8 100644 --- a/src/instruction.cpp +++ b/src/instruction.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/module.cpp b/src/module.cpp index 0b04444d920..e7f3849d953 100644 --- a/src/module.cpp +++ b/src/module.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/propagate_constant.cpp b/src/propagate_constant.cpp index 9aed3268553..a09f4e44945 100644 --- a/src/propagate_constant.cpp +++ b/src/propagate_constant.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/replace_allocate.cpp b/src/replace_allocate.cpp index a832e21bdff..e4d454ac09b 100644 --- a/src/replace_allocate.cpp +++ b/src/replace_allocate.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/simplify_algebra.cpp b/src/simplify_algebra.cpp index 4174e71538b..d77fc63bf88 100644 --- a/src/simplify_algebra.cpp +++ b/src/simplify_algebra.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/split_reduce.cpp b/src/split_reduce.cpp index bda5656ffa6..d1d8930b956 100644 --- a/src/split_reduce.cpp +++ b/src/split_reduce.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/cpu/copy.cpp b/src/targets/cpu/copy.cpp index 03839f20e5c..9d917ca532c 100644 --- a/src/targets/cpu/copy.cpp +++ b/src/targets/cpu/copy.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/cpu/gather.cpp b/src/targets/cpu/gather.cpp index 36a4c77e02c..877874ba934 100644 --- a/src/targets/cpu/gather.cpp +++ b/src/targets/cpu/gather.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/cpu/include/migraphx/cpu/dnnl.hpp b/src/targets/cpu/include/migraphx/cpu/dnnl.hpp index d4940769207..e650e5084a1 100644 --- a/src/targets/cpu/include/migraphx/cpu/dnnl.hpp +++ b/src/targets/cpu/include/migraphx/cpu/dnnl.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/cpu/include/migraphx/cpu/pointwise.hpp b/src/targets/cpu/include/migraphx/cpu/pointwise.hpp index 4e93f10fe51..e0ca9c9771f 100644 --- a/src/targets/cpu/include/migraphx/cpu/pointwise.hpp +++ b/src/targets/cpu/include/migraphx/cpu/pointwise.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/compile_miopen.cpp b/src/targets/gpu/compile_miopen.cpp index 6bb7fa2ccb7..c726f8be686 100644 --- a/src/targets/gpu/compile_miopen.cpp +++ b/src/targets/gpu/compile_miopen.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/compile_ops.cpp b/src/targets/gpu/compile_ops.cpp index fb34e56b746..5fcea2b338c 100644 --- a/src/targets/gpu/compile_ops.cpp +++ b/src/targets/gpu/compile_ops.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/fuse_ops.cpp b/src/targets/gpu/fuse_ops.cpp index 55684e9f07a..ac925269fbc 100644 --- a/src/targets/gpu/fuse_ops.cpp +++ b/src/targets/gpu/fuse_ops.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/abs.hpp b/src/targets/gpu/include/migraphx/gpu/abs.hpp index f8fc5755d25..edeedaa6641 100644 --- a/src/targets/gpu/include/migraphx/gpu/abs.hpp +++ b/src/targets/gpu/include/migraphx/gpu/abs.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/argmax.hpp b/src/targets/gpu/include/migraphx/gpu/argmax.hpp index a3750e8c7bb..6870832f76d 100644 --- a/src/targets/gpu/include/migraphx/gpu/argmax.hpp +++ b/src/targets/gpu/include/migraphx/gpu/argmax.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/argmin.hpp b/src/targets/gpu/include/migraphx/gpu/argmin.hpp index 0afd956965f..dcb9f62b4b0 100644 --- a/src/targets/gpu/include/migraphx/gpu/argmin.hpp +++ b/src/targets/gpu/include/migraphx/gpu/argmin.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/code_object_op.hpp b/src/targets/gpu/include/migraphx/gpu/code_object_op.hpp index 758c7a4c07b..11d86dee5ff 100644 --- a/src/targets/gpu/include/migraphx/gpu/code_object_op.hpp +++ b/src/targets/gpu/include/migraphx/gpu/code_object_op.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/compile_hipblaslt.hpp b/src/targets/gpu/include/migraphx/gpu/compile_hipblaslt.hpp index 8c64fffc901..05bb19f9102 100644 --- a/src/targets/gpu/include/migraphx/gpu/compile_hipblaslt.hpp +++ b/src/targets/gpu/include/migraphx/gpu/compile_hipblaslt.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/convolution.hpp b/src/targets/gpu/include/migraphx/gpu/convolution.hpp index 25897f93d95..a25d6052812 100644 --- a/src/targets/gpu/include/migraphx/gpu/convolution.hpp +++ b/src/targets/gpu/include/migraphx/gpu/convolution.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/fixed_pad.hpp b/src/targets/gpu/include/migraphx/gpu/fixed_pad.hpp index 96ed5517462..70e1b89c4b8 100644 --- a/src/targets/gpu/include/migraphx/gpu/fixed_pad.hpp +++ b/src/targets/gpu/include/migraphx/gpu/fixed_pad.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/gemm.hpp b/src/targets/gpu/include/migraphx/gpu/gemm.hpp index 82f716726b8..7ada7dd05bc 100644 --- a/src/targets/gpu/include/migraphx/gpu/gemm.hpp +++ b/src/targets/gpu/include/migraphx/gpu/gemm.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/hip.hpp b/src/targets/gpu/include/migraphx/gpu/hip.hpp index 7869f84afbb..d04e81e218c 100644 --- a/src/targets/gpu/include/migraphx/gpu/hip.hpp +++ b/src/targets/gpu/include/migraphx/gpu/hip.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/hip_gemm.hpp b/src/targets/gpu/include/migraphx/gpu/hip_gemm.hpp index ae5151525c0..cfe5275d7f6 100644 --- a/src/targets/gpu/include/migraphx/gpu/hip_gemm.hpp +++ b/src/targets/gpu/include/migraphx/gpu/hip_gemm.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/logsoftmax.hpp b/src/targets/gpu/include/migraphx/gpu/logsoftmax.hpp index e7036aa695a..919d3afec34 100644 --- a/src/targets/gpu/include/migraphx/gpu/logsoftmax.hpp +++ b/src/targets/gpu/include/migraphx/gpu/logsoftmax.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/loop.hpp b/src/targets/gpu/include/migraphx/gpu/loop.hpp index 02b8681325d..ea6eb75e769 100644 --- a/src/targets/gpu/include/migraphx/gpu/loop.hpp +++ b/src/targets/gpu/include/migraphx/gpu/loop.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/lrn.hpp b/src/targets/gpu/include/migraphx/gpu/lrn.hpp index 19691a10704..e075fdfceb5 100644 --- a/src/targets/gpu/include/migraphx/gpu/lrn.hpp +++ b/src/targets/gpu/include/migraphx/gpu/lrn.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/multinomial.hpp b/src/targets/gpu/include/migraphx/gpu/multinomial.hpp index 2be31fc35da..0946ffb075b 100644 --- a/src/targets/gpu/include/migraphx/gpu/multinomial.hpp +++ b/src/targets/gpu/include/migraphx/gpu/multinomial.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/nonzero.hpp b/src/targets/gpu/include/migraphx/gpu/nonzero.hpp index 89f4c7c7885..a47eb138c75 100644 --- a/src/targets/gpu/include/migraphx/gpu/nonzero.hpp +++ b/src/targets/gpu/include/migraphx/gpu/nonzero.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/oper.hpp b/src/targets/gpu/include/migraphx/gpu/oper.hpp index 760494ba5ff..6ad559e4c0f 100644 --- a/src/targets/gpu/include/migraphx/gpu/oper.hpp +++ b/src/targets/gpu/include/migraphx/gpu/oper.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/pooling.hpp b/src/targets/gpu/include/migraphx/gpu/pooling.hpp index b3ea6f22d8b..0ea49b44b28 100644 --- a/src/targets/gpu/include/migraphx/gpu/pooling.hpp +++ b/src/targets/gpu/include/migraphx/gpu/pooling.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/prefix_scan_sum.hpp b/src/targets/gpu/include/migraphx/gpu/prefix_scan_sum.hpp index d26701897ce..12c7c2e1957 100644 --- a/src/targets/gpu/include/migraphx/gpu/prefix_scan_sum.hpp +++ b/src/targets/gpu/include/migraphx/gpu/prefix_scan_sum.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/reduce_op.hpp b/src/targets/gpu/include/migraphx/gpu/reduce_op.hpp index b52d70ee240..835cbe03da8 100644 --- a/src/targets/gpu/include/migraphx/gpu/reduce_op.hpp +++ b/src/targets/gpu/include/migraphx/gpu/reduce_op.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/reverse.hpp b/src/targets/gpu/include/migraphx/gpu/reverse.hpp index 1fb951f2c93..7c5e3aab5f5 100644 --- a/src/targets/gpu/include/migraphx/gpu/reverse.hpp +++ b/src/targets/gpu/include/migraphx/gpu/reverse.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/rnn_variable_seq_lens.hpp b/src/targets/gpu/include/migraphx/gpu/rnn_variable_seq_lens.hpp index 36778bd79dd..e29859287c3 100644 --- a/src/targets/gpu/include/migraphx/gpu/rnn_variable_seq_lens.hpp +++ b/src/targets/gpu/include/migraphx/gpu/rnn_variable_seq_lens.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/include/migraphx/gpu/topk.hpp b/src/targets/gpu/include/migraphx/gpu/topk.hpp index 5f8c70f3981..e07a7c4c24e 100644 --- a/src/targets/gpu/include/migraphx/gpu/topk.hpp +++ b/src/targets/gpu/include/migraphx/gpu/topk.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/jit/mlir.cpp b/src/targets/gpu/jit/mlir.cpp index 39c93de3224..31ec380dcc3 100644 --- a/src/targets/gpu/jit/mlir.cpp +++ b/src/targets/gpu/jit/mlir.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/src/targets/gpu/lowering.cpp b/src/targets/gpu/lowering.cpp index d038477999d..9b660beb407 100644 --- a/src/targets/gpu/lowering.cpp +++ b/src/targets/gpu/lowering.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/api/test_custom_op.cpp b/test/api/test_custom_op.cpp index 9e7f6042760..e6224ae6ca2 100644 --- a/test/api/test_custom_op.cpp +++ b/test/api/test_custom_op.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/eliminate_concat_test.cpp b/test/eliminate_concat_test.cpp index 0f2138668b3..c7fcf7859da 100644 --- a/test/eliminate_concat_test.cpp +++ b/test/eliminate_concat_test.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/eval_test.cpp b/test/eval_test.cpp index bd7f9463d72..018e601ef9a 100644 --- a/test/eval_test.cpp +++ b/test/eval_test.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/include/basic_ops.hpp b/test/include/basic_ops.hpp index 5424bba9562..4f634247d7b 100644 --- a/test/include/basic_ops.hpp +++ b/test/include/basic_ops.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/liveness_test.cpp b/test/liveness_test.cpp index 45e47b5381b..da206d239dc 100644 --- a/test/liveness_test.cpp +++ b/test/liveness_test.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/operation.cpp b/test/operation.cpp index ed17109edb1..ce79e3e4ac7 100644 --- a/test/operation.cpp +++ b/test/operation.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/output_alias.cpp b/test/output_alias.cpp index 871169f2304..a52927d93db 100644 --- a/test/output_alias.cpp +++ b/test/output_alias.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/propagate_constant_test.cpp b/test/propagate_constant_test.cpp index c348f84a688..01c37b80b29 100644 --- a/test/propagate_constant_test.cpp +++ b/test/propagate_constant_test.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/replace_allocate.cpp b/test/replace_allocate.cpp index ccf927f2fbb..1c7b5ca4882 100644 --- a/test/replace_allocate.cpp +++ b/test/replace_allocate.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/run_loop_test.cpp b/test/run_loop_test.cpp index 8c1238bb031..9b94c45effa 100644 --- a/test/run_loop_test.cpp +++ b/test/run_loop_test.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/schedule_test.cpp b/test/schedule_test.cpp index 69edd4179fc..addcabb80f3 100644 --- a/test/schedule_test.cpp +++ b/test/schedule_test.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/test/split_reduce.cpp b/test/split_reduce.cpp index 09827d780e9..cabdffb3ff3 100644 --- a/test/split_reduce.cpp +++ b/test/split_reduce.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/tools/api/api.cpp b/tools/api/api.cpp index 0bdda69f47c..88763e882f8 100644 --- a/tools/api/api.cpp +++ b/tools/api/api.cpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal diff --git a/tools/include/operation.hpp b/tools/include/operation.hpp index 6269a8a038b..a2e73c79940 100644 --- a/tools/include/operation.hpp +++ b/tools/include/operation.hpp @@ -1,7 +1,7 @@ /* * The MIT License (MIT) * - * Copyright (c) 2015-2025 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (c) 2015-2026 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal From 4d7c7821f2566206bb65d1578a4fd88dabb78881 Mon Sep 17 00:00:00 2001 From: charlie Date: Fri, 16 Jan 2026 15:58:43 -0600 Subject: [PATCH 23/23] add changelog entry --- CHANGELOG.md | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index dead3a4dbbf..e3f8d157971 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -3,6 +3,12 @@ Full documentation for MIGraphX is available at [https://rocmdocs.amd.com/projects/AMDMIGraphX/en/latest/](https://rocmdocs.amd.com/projects/AMDMIGraphX/en/latest/). +## Develop Branch + +### Changed + +* Refactored instruction output alias to return a vector of aliases (#4540). + ## MIGraphX 2.15 for ROCm 7.2.0 ### Added