Code cleanup

src/include/migraphx/module.hpp

@@ -54,7 +54,9 @@ using ins_dep_map   = std::unordered_map<instruction_ref, std::unordered_set<ins
  */
 struct module
 {
-    bool use_local_alloc = true;
+    // used by replace_allocate pass
+    // allocate memory in this module rather than using output parmaeters
+    bool use_local_alloc = false;
 
     module(const std::string& name = "");
 

src/targets/gpu/loop.cpp

@@ -76,11 +76,6 @@ struct gpu_loop
         }
     }
 
-    /**
-     * This finds the output parameters for a module and returns a map between the parameter name
-     * and output argument indicies. Needs to have the module names mapped to the correct input
-     * parameters to begin with; not sure where those indices are set.
-     */
     std::unordered_map<std::string, int> get_output_params(const module& m) const
     {
         auto get_output_index = [](const std::string& name) {

src/targets/gpu/lowering.cpp

@@ -361,76 +361,10 @@ struct miopen_apply
         });
     }
 
-    // This might work, but it's going to have many arguments.
-    // Will need a map between submodules and argument index
-    // Also allocates seperate memory for each batch size...
-    // void add_select_module_op()
-    //{
-    //    apply_map.emplace("select_module", [=](instruction_ref ins) {
-    //        std::vector<instruction_ref> inputs = ins->inputs();
-    //        auto mod_args = ins->module_inputs();
-    //        for(const auto* smod : mod_args)
-    //        {
-    //            auto pn_list = smod->get_parameter_names();
-    //            std::transform(pn_list.begin(),
-    //                           pn_list.end(),
-    //                           std::back_inserter(inputs),
-    //                           [&](auto pn) { return insert_allocation(ins,
-    //                           smod->get_parameter_shape(pn)); });
-    //        }
-    //        return mod->replace_instruction(ins, ins->get_operator(), inputs, mod_args);
-    //    });
-    //}
-
-    // Wrong output parameter shape error with this
-    // Try instead to have allocates occur in the submodule
-    // void add_select_module_op()
-    //{
-    //    // make maximum buffer size allocation for output parameters
-    //    apply_map.emplace("select_module", [=](instruction_ref ins) {
-    //        std::vector<instruction_ref> inputs = ins->inputs();
-    //        auto output_sub_shapes              = ins->get_shape().sub_shapes();
-    //        std::transform(output_sub_shapes.begin(),
-    //                       output_sub_shapes.end(),
-    //                       std::back_inserter(inputs),
-    //                       [&](auto s) {
-    //                           shape max_shape{s.type(), s.max_lens()};
-    //                           return insert_allocation(ins, max_shape);
-    //                       });
-    //        return mod->replace_instruction(ins, ins->get_operator(), inputs,
-    //        ins->module_inputs());
-    //    });
-    //}
-
-    // do a copy to cpu of the submodule alloc
-    // void add_select_module_op()
-    //{
-    //    apply_map.emplace("select_module", [=](instruction_ref ins) {
-    //        std::vector<instruction_ref> inputs = ins->inputs();
-    //        auto mod_args = ins->module_inputs();
-    //        for(auto smod : mod_args)
-    //        {
-    //            auto last_ins = std::prev(smod->end());
-    //            if(last_ins->name() == "@return")
-    //            {
-    //                const auto& prog_outputs = last_ins->inputs();
-    //                std::vector<instruction_ref> outputs_alias(prog_outputs.size());
-    //                std::transform(prog_outputs.begin(),
-    //                        prog_outputs.end(),
-    //                        outputs_alias.begin(),
-    //                        [](const auto& i) { return instruction::get_output_alias(i); });
-    //                for(auto out_ins : outputs_alias)
-    //                {
-    //                    std::cout << "output_alias_ins: " << out_ins->name() << std::endl;
-    //                    smod->insert_instruction(out_ins->outputs()[0],
-    //                    make_op("hip::copy_from_gpu"), out_ins->inputs()[1]);
-    //                }
-    //            }
-    //        }
-    //        return ins;
-    //    });
-    //}
-
+    /**
+     * Turns on use_local_alloc in the select_module submodules.
+     * Changes the submodule returns to a hip::sync_stream.
+     */
     void add_select_module_op()
     {
         apply_map.emplace("select_module", [=](instruction_ref ins) {
@@ -438,7 +372,8 @@ struct miopen_apply
             auto mod_args                       = ins->module_inputs();
             for(auto smod : mod_args)
             {
-                auto last_ins = std::prev(smod->end());
+                smod->use_local_alloc = true;
+                auto last_ins         = std::prev(smod->end());
                 if(last_ins->name() == "@return")
                 {
                     for(auto out_ins : last_ins->inputs())

src/targets/gpu/select_module.cpp

-/*
- * The MIT License (MIT)
- *
- * Copyright (c) 2015-2023 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 <migraphx/gpu/select_module.hpp>
-#include <migraphx/gpu/context.hpp>
-
-namespace migraphx {
-inline namespace MIGRAPHX_INLINE_NS {
-namespace gpu {
-
-} // namespace gpu
-} // namespace MIGRAPHX_INLINE_NS
-} // namespace migraphx

test/ref_ops_test.cpp

@@ -6989,7 +6989,7 @@ TEST_CASE(scatternd_reduction_test)
     }
 }
 
-TEST_CASE(select_module_test0)
+TEST_CASE(select_module_reduce_test0)
 {
     migraphx::program p;
 
@@ -7035,7 +7035,7 @@ TEST_CASE(select_module_test0)
     EXPECT(migraphx::verify_range(results_vector, gold));
 }
 
-TEST_CASE(select_module_test1)
+TEST_CASE(select_module_reduce_test1)
 {
     migraphx::program p;