manual merge

src/opt/memory_coloring.cpp

@@ -2,16 +2,16 @@
 #include "memory_coloring_impl.hpp"
 
 namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
+inline namespace MIGRAPHX_INLINE_NS {
 
 void memory_coloring::apply(program& p) const
 {
-    if(!enabled(MIGRAPH_DISABLE_MEMORY_COLORING{}))
+    if(!enabled(MIGRAPHX_DISABLE_MEMORY_COLORING{}))
     {
         memory_coloring_impl opt(&p, allocation_op, verify);
         opt.run();
     }
 }
 
-} // namespace MIGRAPH_INLINE_NS
+} // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx

src/opt/memory_coloring_impl.cpp

 #include "memory_coloring_impl.hpp"
 
 namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
+inline namespace MIGRAPHX_INLINE_NS {
 
 void memory_coloring_impl::run()
 {
-    MIGRAPH_DEBUG(dump("---Before memory coloring---"));
-    MIGRAPH_DEBUG(dump_program());
+    MIGRAPHX_DEBUG(dump("---Before memory coloring---"));
+    MIGRAPHX_DEBUG(dump_program());
     build();
     if(num_of_lives != 0)
     {
-        MIGRAPH_DEBUG(dump_intervals());
+        MIGRAPHX_DEBUG(dump_intervals());
         // Coloring
         while(!alloc_queue.empty())
         {
@@ -85,7 +85,7 @@ bool memory_coloring_impl::allocate(interval_ptr interval)
         conflict_queue.pop();
     }
     segment.offset = offset;
-    MIGRAPH_DEBUG(segment.dump());
+    MIGRAPHX_DEBUG(segment.dump());
     required_bytes = std::max(required_bytes, offset + segment.size);
     return true;
 }
@@ -218,8 +218,8 @@ void memory_coloring_impl::rewrite()
             }
         }
     }
-    MIGRAPH_DEBUG(dump("---After rewrite---"));
-    MIGRAPH_DEBUG(dump_program());
+    MIGRAPHX_DEBUG(dump("---After rewrite---"));
+    MIGRAPHX_DEBUG(dump_program());
 }
 
 void memory_coloring_impl::verify()
@@ -235,7 +235,7 @@ void memory_coloring_impl::verify()
             {
                 // TODO: This check breaks on the tests
                 // if(!interval.is_live_on_entry)
-                // MIGRAPH_THROW("interval is not live on entry");
+                // MIGRAPHX_THROW("interval is not live on entry");
                 continue;
             }
 
@@ -253,14 +253,14 @@ void memory_coloring_impl::verify()
                     if(range->offset == invalid_offset)
                         continue;
                     if(!is_disjoin(*range, segment))
-                        MIGRAPH_THROW("range and segment is not disjoined");
+                        MIGRAPHX_THROW("range and segment is not disjoined");
                 }
             }
         }
     }
 }
 
-#ifdef MIGRAPH_DEBUG_OPT
+#ifdef MIGRAPHX_DEBUG_OPT
 
 void memory_coloring_impl::dump(const std::string& str) { std::cout << str << std::endl; }
 
@@ -334,5 +334,5 @@ void live_interval::dump()
 
 #endif
 
-} // namespace MIGRAPH_INLINE_NS
+} // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx

src/opt/memory_coloring_impl.hpp

-#ifndef MIGRAPH_GUARD_RTGLIB_MEMORY_COLORING_IMPL_HPP
-#define MIGRAPH_GUARD_RTGLIB_MEMORY_COLORING_IMPL_HPP
+#ifndef MIGRAPHX_GUARD_RTGLIB_MEMORY_COLORING_IMPL_HPP
+#define MIGRAPHX_GUARD_RTGLIB_MEMORY_COLORING_IMPL_HPP
 #include "common_header.hpp"
 #include <migraphx/config.hpp>
 
 namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
+inline namespace MIGRAPHX_INLINE_NS {
 
 static const int invalid_offset = -1;
 
@@ -15,7 +15,7 @@ struct live_range
     long long offset; // offset to base pointer of allocated memory trunk.
     int vn;           // value number that identifies this live_range.
     long long size;   // size of required memory in bytes
-#ifdef MIGRAPH_DEBUG_OPT
+#ifdef MIGRAPHX_DEBUG_OPT
     void dump();
 #endif
 };
@@ -35,7 +35,7 @@ struct live_interval
     int get_end() const { return segment.end; }
     long long get_offset() const { return segment.offset; }
 
-#ifdef MIGRAPH_DEBUG_OPT
+#ifdef MIGRAPHX_DEBUG_OPT
     void dump();
 #endif
 
@@ -101,7 +101,7 @@ struct memory_coloring_impl
         return ((end1 < range2.offset) || (end2 < range1.offset));
     }
     void verify();
-#ifdef MIGRAPH_DEBUG_OPT
+#ifdef MIGRAPHX_DEBUG_OPT
     void dump(const std::string&);
     void dump_program();
     void dump_intervals();
@@ -154,6 +154,6 @@ struct memory_coloring_impl
     bool enable_verify;
 };
 
-} // namespace MIGRAPH_INLINE_NS
+} // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx
 #endif

src/program.cpp

@@ -11,10 +11,10 @@
 #include <utility>
 
 namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
+inline namespace MIGRAPHX_INLINE_NS {
 
-MIGRAPH_DECLARE_ENV_VAR(MIGRAPH_TRACE_COMPILE)
-MIGRAPH_DECLARE_ENV_VAR(MIGRAPH_TRACE_EVAL)
+MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_TRACE_COMPILE)
+MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_TRACE_EVAL)
 
 struct program_impl
 {
@@ -282,7 +282,7 @@ void program::compile(const target& t, tracer trace)
 {
     assert(this->validate() == impl->instructions.end());
     this->impl->ctx = t.get_context();
-    if(enabled(MIGRAPH_TRACE_COMPILE{}))
+    if(enabled(MIGRAPHX_TRACE_COMPILE{}))
         trace = tracer{std::cout};
     trace(*this);
     trace();
@@ -297,7 +297,7 @@ void program::compile(const target& t, tracer trace)
         if(invalid != impl->instructions.end())
         {
             auto index = std::distance(impl->instructions.begin(), invalid);
-            MIGRAPH_THROW(p.name() + " pass produces invalid program at instruction " +
+            MIGRAPHX_THROW(p.name() + " pass produces invalid program at instruction " +
                            std::to_string(index) + ": " + invalid->name());
         }
         trace();
@@ -307,7 +307,7 @@ void program::compile(const target& t, tracer trace)
     if(invalid != impl->instructions.end())
     {
         auto index = std::distance(impl->instructions.begin(), invalid);
-        MIGRAPH_THROW("Invalid program from compilation at instruction " + std::to_string(index));
+        MIGRAPHX_THROW("Invalid program from compilation at instruction " + std::to_string(index));
     }
 }
 
@@ -334,7 +334,7 @@ argument generic_eval(const program& p,
                                 auto param_name =
                                     any_cast<builtin::param>(ins->get_operator()).parameter;
                                 if(not contains(params, param_name))
-                                    MIGRAPH_THROW("Parameter not found: " + param_name);
+                                    MIGRAPHX_THROW("Parameter not found: " + param_name);
                                 return params.at(param_name);
                             }));
         }
@@ -361,7 +361,7 @@ argument generic_eval(const program& p,
 
 argument program::eval(std::unordered_map<std::string, argument> params) const
 {
-    if(enabled(MIGRAPH_TRACE_EVAL{}))
+    if(enabled(MIGRAPHX_TRACE_EVAL{}))
     {
         auto& ctx = this->impl->ctx;
         return generic_eval(*this, this->impl->ctx, std::move(params), [&](auto& ins, auto f) {
@@ -501,5 +501,5 @@ std::ostream& operator<<(std::ostream& os, const program& p)
     return os;
 }
 
-} // namespace MIGRAPH_INLINE_NS
+} // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx

src/shape.cpp

@@ -7,7 +7,7 @@
 #include <iostream>
 
 namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
+inline namespace MIGRAPHX_INLINE_NS {
 
 struct shape_impl
 {
@@ -169,12 +169,12 @@ std::string shape::type_string() const
 {
     switch(this->type())
     {
-#define MIGRAPH_SHAPE_TYPE_STRING_CASE(x, t) \
+#define MIGRAPHX_SHAPE_TYPE_STRING_CASE(x, t) \
     case x: return #x;
-        MIGRAPH_SHAPE_VISIT_TYPES(MIGRAPH_SHAPE_TYPE_STRING_CASE)
-#undef MIGRAPH_SHAPE_TYPE_STRING_CASE
+        MIGRAPHX_SHAPE_VISIT_TYPES(MIGRAPHX_SHAPE_TYPE_STRING_CASE)
+#undef MIGRAPHX_SHAPE_TYPE_STRING_CASE
     }
-    MIGRAPH_THROW("Invalid type");
+    MIGRAPHX_THROW("Invalid type");
 }
 
 bool operator==(const shape& x, const shape& y)
@@ -191,5 +191,5 @@ std::ostream& operator<<(std::ostream& os, const shape& x)
     return os;
 }
 
-} // namespace MIGRAPH_INLINE_NS
+} // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx

src/simplify_algebra.cpp

@@ -5,7 +5,7 @@
 #include <migraphx/literal.hpp>
 
 namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
+inline namespace MIGRAPHX_INLINE_NS {
 
 struct find_add_lit_broadcast
 {
@@ -61,5 +61,5 @@ struct find_add_lit_broadcast
 
 void simplify_algebra::apply(program& p) const { match::find_matches(p, find_add_lit_broadcast{}); }
 
-} // namespace MIGRAPH_INLINE_NS
+} // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx

src/simplify_reshapes.cpp

@@ -7,7 +7,7 @@
 #include <unordered_set>
 
 namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
+inline namespace MIGRAPHX_INLINE_NS {
 
 bool is_reshaper(const std::string& name)
 {
@@ -60,5 +60,5 @@ void simplify_reshapes::apply(program& p) const
     }
 }
 
-} // namespace MIGRAPH_INLINE_NS
+} // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx

src/targets/cpu/gemm.cpp

@@ -4,7 +4,7 @@
 #include <blaze/math/CustomMatrix.h>
 
 namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
+inline namespace MIGRAPHX_INLINE_NS {
 namespace cpu {
 
 template <class T>
@@ -94,5 +94,5 @@ void migemm(
 }
 
 } // namespace cpu
-} // namespace MIGRAPH_INLINE_NS
+} // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx

src/targets/cpu/include/migraphx/cpu/context.hpp

-#ifndef MIGRAPH_GUARD_RTGLIB_CONTEXT_HPP
-#define MIGRAPH_GUARD_RTGLIB_CONTEXT_HPP
+#ifndef MIGRAPHX_GUARD_RTGLIB_CONTEXT_HPP
+#define MIGRAPHX_GUARD_RTGLIB_CONTEXT_HPP
 
 #include <migraphx/config.hpp>
 
 namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
+inline namespace MIGRAPHX_INLINE_NS {
 namespace cpu {
 
 struct context
@@ -13,7 +13,7 @@ struct context
 };
 
 } // namespace cpu
-} // namespace MIGRAPH_INLINE_NS
+} // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx
 
 #endif

src/targets/cpu/include/migraphx/cpu/gemm.hpp

-#ifndef MIGRAPH_GUARD_RTGLIB_CPU_GEMM_HPP
-#define MIGRAPH_GUARD_RTGLIB_CPU_GEMM_HPP
+#ifndef MIGRAPHX_GUARD_RTGLIB_CPU_GEMM_HPP
+#define MIGRAPHX_GUARD_RTGLIB_CPU_GEMM_HPP
 
 #include <migraphx/argument.hpp>
 #include <migraphx/config.hpp>
 
 namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
+inline namespace MIGRAPHX_INLINE_NS {
 namespace cpu {
 
 void migemm(
     const argument& c_arg, const argument& a_arg, const argument& b_arg, float alpha, float beta);
 
 } // namespace cpu
-} // namespace MIGRAPH_INLINE_NS
+} // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx
 
 #endif

src/targets/cpu/include/migraphx/cpu/lowering.hpp

-#ifndef MIGRAPH_GUARD_RTGLIB_CPU_LOWERING_HPP
-#define MIGRAPH_GUARD_RTGLIB_CPU_LOWERING_HPP
+#ifndef MIGRAPHX_GUARD_RTGLIB_CPU_LOWERING_HPP
+#define MIGRAPHX_GUARD_RTGLIB_CPU_LOWERING_HPP
 
 #include <migraphx/program.hpp>
 #include <migraphx/config.hpp>
 
 namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
+inline namespace MIGRAPHX_INLINE_NS {
 namespace cpu {
 
 struct lowering
@@ -15,7 +15,7 @@ struct lowering
 };
 
 } // namespace cpu
-} // namespace MIGRAPH_INLINE_NS
+} // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx
 
 #endif

src/targets/cpu/include/migraphx/cpu/target.hpp

-#ifndef MIGRAPH_GUARD_MIGRAPHLIB_CPU_TARGET_HPP
-#define MIGRAPH_GUARD_MIGRAPHLIB_CPU_TARGET_HPP
+#ifndef MIGRAPHX_GUARD_MIGRAPHLIB_CPU_TARGET_HPP
+#define MIGRAPHX_GUARD_MIGRAPHLIB_CPU_TARGET_HPP
 
 #include <migraphx/program.hpp>
 #include <migraphx/cpu/context.hpp>
 #include <migraphx/config.hpp>
 
 namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
+inline namespace MIGRAPHX_INLINE_NS {
 namespace cpu {
 
 struct target
@@ -17,7 +17,7 @@ struct target
 };
 
 } // namespace cpu
-} // namespace MIGRAPH_INLINE_NS
+} // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx
 
 #endif

src/targets/cpu/lowering.cpp

@@ -10,7 +10,7 @@
 #include <utility>
 
 namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
+inline namespace MIGRAPHX_INLINE_NS {
 namespace cpu {
 
 template <typename T>
@@ -699,5 +699,5 @@ struct cpu_apply
 void lowering::apply(program& p) const { cpu_apply{&p}.apply(); }
 
 } // namespace cpu
-} // namespace MIGRAPH_INLINE_NS
+} // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx

src/targets/cpu/target.cpp

@@ -4,7 +4,7 @@
 #include <migraphx/auto_contiguous.hpp>
 
 namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
+inline namespace MIGRAPHX_INLINE_NS {
 namespace cpu {
 
 std::string target::name() const { return "cpu"; }
@@ -15,5 +15,5 @@ std::vector<pass> target::get_passes(migraphx::context&) const
 }
 
 } // namespace cpu
-} // namespace MIGRAPH_INLINE_NS
+} // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx

src/targets/gpu/CMakeLists.txt

@@ -22,7 +22,7 @@ add_library(migraphx_device
 )
 set_target_properties(migraphx_device PROPERTIES EXPORT_NAME device)
 rocm_clang_tidy_check(migraphx_device)
-target_link_libraries(migraphx_device migraphx hip::device -Wno-invalid-command-line-argument -amdgpu-target=gfx803 -amdgpu-target=gfx900 -amdgpu-target=gfx903)
+target_link_libraries(migraphx_device migraphx hip::device -Wno-invalid-command-line-argument -amdgpu-target=gfx803 -amdgpu-target=gfx900 -amdgpu-target=gfx906)
 target_include_directories(migraphx_device PUBLIC $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>)
 target_include_directories(migraphx_device PRIVATE $<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/device/include>)
 

src/targets/gpu/add.cpp

@@ -6,7 +6,7 @@
 #include <utility>
 
 namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
+inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
 shape hip_add::compute_shape(const std::vector<shape>& inputs) const
@@ -51,5 +51,5 @@ argument miopen_add::compute(context& ctx,
 }
 
 } // namespace gpu
-} // namespace MIGRAPH_INLINE_NS
+} // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx

src/targets/gpu/batchnorm.cpp

@@ -5,7 +5,7 @@
 #include <utility>
 
 namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
+inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
 shape miopen_batch_norm_inference::compute_shape(const std::vector<shape>& inputs) const
@@ -43,5 +43,5 @@ argument miopen_batch_norm_inference::compute(context& ctx,
 }
 
 } // namespace gpu
-} // namespace MIGRAPH_INLINE_NS
+} // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx

src/targets/gpu/concat.cpp

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

src/targets/gpu/contiguous.cpp

@@ -5,7 +5,7 @@
 #include <utility>
 
 namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
+inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
 shape miopen_contiguous::compute_shape(const std::vector<shape>& inputs) const
@@ -25,5 +25,5 @@ argument miopen_contiguous::compute(context& ctx,
 }
 
 } // namespace gpu
-} // namespace MIGRAPH_INLINE_NS
+} // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx

src/targets/gpu/convolution.cpp

@@ -5,7 +5,7 @@
 #include <utility>
 
 namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
+inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
 shape miopen_convolution::compute_shape(const std::vector<shape>& inputs) const
@@ -82,5 +82,5 @@ shape miopen_convolution::compile(context& ctx,
 }
 
 } // namespace gpu
-} // namespace MIGRAPH_INLINE_NS
+} // namespace MIGRAPHX_INLINE_NS
 } // namespace migraphx