Merge from develop

src/targets/gpu/include/migraphx/gpu/tan.hpp

 #ifndef MIGRAPHX_GUARD_RTGLIB_TAN_HPP
 #define MIGRAPHX_GUARD_RTGLIB_TAN_HPP
 
-#include <migraphx/gpu/lowering.hpp>
 #include <migraphx/gpu/oper.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/instruction.hpp>
-#include <migraphx/operators.hpp>
-#include <migraphx/generate.hpp>
-#include <migraphx/shape_for_each.hpp>
-#include <migraphx/gpu/miopen.hpp>
-#include <migraphx/gpu/hip.hpp>
-#include <migraphx/dfor.hpp>
-#include <migraphx/gpu/device/contiguous.hpp>
 #include <migraphx/gpu/device/tan.hpp>
-#include <migraphx/iterator_for.hpp>
-#include <migraphx/gpu/rocblas.hpp>
-#include <migraphx/gpu/context.hpp>
-#include <migraphx/config.hpp>
-#include <utility>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/targets/gpu/include/migraphx/gpu/tanh.hpp

 #ifndef MIGRAPHX_GUARD_RTGLIB_TANH_HPP
 #define MIGRAPHX_GUARD_RTGLIB_TANH_HPP
 
-#include <migraphx/gpu/lowering.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/instruction.hpp>
-#include <migraphx/operators.hpp>
-#include <migraphx/generate.hpp>
-#include <migraphx/shape_for_each.hpp>
-#include <migraphx/config.hpp>
+#include <migraphx/shape.hpp>
 #include <migraphx/gpu/miopen.hpp>
-#include <migraphx/gpu/hip.hpp>
-#include <migraphx/dfor.hpp>
-#include <migraphx/gpu/device/contiguous.hpp>
-#include <migraphx/gpu/device/add.hpp>
-#include <migraphx/iterator_for.hpp>
-#include <migraphx/gpu/rocblas.hpp>
-#include <migraphx/gpu/context.hpp>
-#include <utility>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
+struct context;
+
 struct miopen_tanh
 {
     shared<activation_descriptor> ad;

src/targets/gpu/leaky_relu.cpp

 #include <migraphx/gpu/leaky_relu.hpp>
-#include <migraphx/operators.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/gpu/miopen.hpp>
-#include <utility>
+#include <migraphx/gpu/context.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/targets/gpu/lowering.cpp

@@ -127,8 +127,7 @@ struct miopen_apply
         }
         else
         {
-            auto is     = prog->add_outline(s);
-            auto result = prog->insert_instruction(ins, hip_allocate{std::move(tag)}, is);
+            auto result = prog->insert_instruction(ins, hip_allocate{s, std::move(tag)});
             return result;
         }
     }

src/targets/gpu/lrn.cpp

 #include <migraphx/gpu/lrn.hpp>
-#include <migraphx/operators.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/gpu/miopen.hpp>
-#include <utility>
+#include <migraphx/gpu/context.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/targets/gpu/pad.cpp

 #include <migraphx/gpu/pad.hpp>
-#include <migraphx/operators.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/gpu/miopen.hpp>
+#include <migraphx/gpu/context.hpp>
 #include <migraphx/gpu/device/pad.hpp>
-#include <utility>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/targets/gpu/pooling.cpp

 #include <migraphx/gpu/pooling.hpp>
-#include <migraphx/operators.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/gpu/miopen.hpp>
-#include <utility>
+#include <migraphx/gpu/context.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/targets/gpu/relu.cpp

 #include <migraphx/gpu/relu.hpp>
-#include <migraphx/operators.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/gpu/miopen.hpp>
-#include <utility>
+#include <migraphx/gpu/context.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/targets/gpu/sigmoid.cpp

 #include <migraphx/gpu/sigmoid.hpp>
-#include <migraphx/operators.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/gpu/miopen.hpp>
-#include <utility>
+#include <migraphx/gpu/context.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/targets/gpu/softmax.cpp

 #include <migraphx/gpu/softmax.hpp>
-#include <migraphx/operators.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/gpu/miopen.hpp>
-#include <utility>
+#include <migraphx/gpu/context.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/targets/gpu/tanh.cpp

 #include <migraphx/gpu/tanh.hpp>
-#include <migraphx/operators.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/gpu/miopen.hpp>
-#include <utility>
+#include <migraphx/gpu/context.hpp>
 
 namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {

src/targets/gpu/write_literals.cpp

@@ -37,8 +37,7 @@ void write_literals::apply(program& p) const
             {
                 literal l  = ins->get_literal();
                 auto pre   = p.add_literal(l);
-                auto s     = p.add_outline(l.get_shape());
-                auto alloc = p.insert_instruction(std::next(pre), hip_allocate{}, s);
+                auto alloc = p.insert_instruction(std::next(pre), hip_allocate{l.get_shape()});
                 p.replace_instruction(ins, hip_copy{}, pre, alloc);
             }
             else

test/CMakeLists.txt

@@ -113,13 +113,18 @@ endif()
 
 # Onnx test
 set(TEST_ONNX_DIR ${CMAKE_CURRENT_SOURCE_DIR}/onnx)
-add_executable(test_onnx ${TEST_ONNX_DIR}/onnx_test.cpp)
-rocm_clang_tidy_check(test_onnx)
-target_link_libraries(test_onnx migraphx_onnx)
-target_include_directories(test_onnx PUBLIC include)
-add_test(NAME test_onnx COMMAND $<TARGET_FILE:test_onnx> WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/onnx) 
-add_dependencies(tests test_onnx)
-add_dependencies(check test_onnx)
+file (GLOB ONNX_TESTS ${TEST_ONNX_DIR}/*.cpp)
+foreach(ONNX_TEST ${ONNX_TESTS})
+    get_filename_component(BASE_NAME ${ONNX_TEST} NAME_WE)
+    set(TEST_NAME test_${BASE_NAME})
+    add_executable(${TEST_NAME} ${TES_ONNX_DIR}/${ONNX_TEST})
+    rocm_clang_tidy_check(${TEST_NAME})
+    target_link_libraries(${TEST_NAME} migraphx_onnx)
+    target_include_directories(${TEST_NAME} PUBLIC include)
+    add_test(NAME ${TEST_NAME} COMMAND $<TARGET_FILE:${TEST_NAME}> WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/onnx) 
+    add_dependencies(tests ${TEST_NAME})
+    add_dependencies(check ${TEST_NAME})
+endforeach()
 
 if(MIGRAPHX_ENABLE_PYTHON)
 add_subdirectory(py)

test/cpu_ops_test.cpp

@@ -7,6 +7,7 @@
 #include <migraphx/verify.hpp>
 #include <migraphx/onnx.hpp>
 #include "test.hpp"
+#include <migraphx/half.hpp>
 
 float sigmoid(float x) { return 1 / (1 + expf(-x)); }
 
@@ -163,6 +164,48 @@ TEST_CASE(gather_test)
         result.visit([&](auto output) { res_data.assign(output.begin(), output.end()); });
         EXPECT(migraphx::verify_range(res_data, golden));
     }
+
+    {
+        migraphx::program p;
+
+        std::vector<float> data(3 * 3);
+        std::iota(data.begin(), data.end(), 0.5);
+        migraphx::shape s{migraphx::shape::float_type, {3, 3}};
+        auto a0 = p.add_literal(migraphx::literal{s, data});
+        // scalar index
+        migraphx::shape s_indices{migraphx::shape::int32_type};
+        std::vector<int> indices{0};
+        auto a1  = p.add_literal(migraphx::literal{s_indices, indices});
+        int axis = -1;
+        p.add_instruction(migraphx::op::gather{axis}, a0, a1);
+        p.compile(migraphx::cpu::target{});
+        auto result = p.eval({});
+        std::vector<float> res_data{};
+        std::vector<float> golden = {0.5f, 3.5f, 6.5f};
+        result.visit([&](auto output) { res_data.assign(output.begin(), output.end()); });
+        EXPECT(migraphx::verify_range(res_data, golden));
+    }
+
+    {
+        migraphx::program p;
+
+        std::vector<float> data(3);
+        std::iota(data.begin(), data.end(), 0.5);
+        migraphx::shape s{migraphx::shape::float_type, {3}};
+        auto a0 = p.add_literal(migraphx::literal{s, data});
+        // scalar index
+        migraphx::shape s_indices{migraphx::shape::int32_type};
+        std::vector<int> indices{0};
+        auto a1  = p.add_literal(migraphx::literal{s_indices, indices});
+        int axis = -1;
+        p.add_instruction(migraphx::op::gather{axis}, a0, a1);
+        p.compile(migraphx::cpu::target{});
+        auto result = p.eval({});
+        std::vector<float> res_data{};
+        std::vector<float> golden = {0.5f};
+        result.visit([&](auto output) { res_data.assign(output.begin(), output.end()); });
+        EXPECT(migraphx::verify_range(res_data, golden));
+    }
 }
 
 TEST_CASE(squeeze_test)
@@ -1375,4 +1418,22 @@ TEST_CASE(pad_test)
     EXPECT(migraphx::verify_range(results_vector, gold));
 }
 
+TEST_CASE(fp16_test)
+{
+    migraphx::program p;
+    migraphx::shape s{migraphx::shape::half_type, {1}};
+    migraphx::half a{1.5};
+    migraphx::half b{2.5};
+    migraphx::half c{4.0};
+    auto l0 = p.add_literal(migraphx::literal{s, {a}});
+    auto l1 = p.add_literal(migraphx::literal{s, {b}});
+    p.add_instruction(migraphx::op::add{}, l0, l1);
+    p.compile(migraphx::cpu::target{});
+    auto result = p.eval({});
+    std::vector<migraphx::half> results_vector(1);
+    result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
+    std::vector<migraphx::half> gold{c};
+    EXPECT(migraphx::verify_range(results_vector, gold));
+}
+
 int main(int argc, const char* argv[]) { test::run(argc, argv); }

test/eval_test.cpp

@@ -128,7 +128,7 @@ TEST_CASE(print_test)
 {
     migraphx::program p;
 
-    auto x   = p.add_parameter("x", {migraphx::shape::int64_type});
+    auto x   = p.add_parameter("x", {migraphx::shape::int32_type});
     auto two = p.add_literal(2);
     p.add_instruction(sum_op{}, x, two);
 
@@ -142,8 +142,8 @@ TEST_CASE(param_test)
 {
     migraphx::program p;
 
-    auto x = p.add_parameter("x", {migraphx::shape::int64_type});
-    auto y = p.add_parameter("y", {migraphx::shape::int64_type});
+    auto x = p.add_parameter("x", {migraphx::shape::int32_type});
+    auto y = p.add_parameter("y", {migraphx::shape::int32_type});
 
     p.add_instruction(sum_op{}, x, y);
     auto result = p.eval(
@@ -156,8 +156,8 @@ TEST_CASE(param_error_test)
 {
     migraphx::program p;
 
-    auto x = p.add_parameter("x", {migraphx::shape::int64_type});
-    auto y = p.add_parameter("y", {migraphx::shape::int64_type});
+    auto x = p.add_parameter("x", {migraphx::shape::int32_type});
+    auto y = p.add_parameter("y", {migraphx::shape::int32_type});
 
     p.add_instruction(sum_op{}, x, y);
     EXPECT(test::throws<migraphx::exception>(
@@ -167,6 +167,22 @@ TEST_CASE(param_error_test)
         "Parameter not found: y"));
 }
 
+TEST_CASE(param_shape_error_test)
+{
+    migraphx::program p;
+
+    auto x = p.add_parameter("x", {migraphx::shape::int32_type, {1, 2}});
+    auto y = p.add_parameter("y", {migraphx::shape::int32_type, {1, 2}});
+
+    p.add_instruction(sum_op{}, x, y);
+    EXPECT(test::throws<migraphx::exception>(
+        [&] {
+            p.eval({{"x", migraphx::literal{1}.get_argument()},
+                    {"y", migraphx::literal{2}.get_argument()}});
+        },
+        "Incorrect shape"));
+}
+
 TEST_CASE(replace_test)
 {
     migraphx::program p;

test/memory_coloring_test.cpp

@@ -53,8 +53,8 @@ struct allocate
     std::string name() const { return "allocate"; }
     migraphx::shape compute_shape(const std::vector<migraphx::shape>& inputs) const
     {
-        migraphx::check_shapes{inputs, *this}.has(1);
-        return inputs.front();
+        migraphx::check_shapes{inputs, *this}.has(0);
+        return s;
     }
     migraphx::argument compute(migraphx::context&,
                                const migraphx::shape& output_shape,
@@ -66,8 +66,7 @@ struct allocate
 
 migraphx::instruction_ref add_alloc(migraphx::program& p, const migraphx::shape& s)
 {
-    auto a0 = p.add_outline(s);
-    return p.add_instruction(allocate{}, a0);
+    return p.add_instruction(allocate{s});
 }
 
 bool no_allocate(const migraphx::program& p)

test/onnx/add_fp16_test.onnx

+add-fp16-example:m
+
+
0
+
1
2"Add
test-add-fp16*

+*|B
0*

+*
B
1Z
+
0
+
+
+
+
Z
+
1
+
+
+
+
b
+
2
+
+
+
+
B
\ No newline at end of file

test/onnx/constant_scalar.onnx

+shape-gather-example:O
+2value"Constant*
+value**

Bconst_tensor 
constantb
+
z
+
+

+
B
\ No newline at end of file