Skip to content

GitLab

"ts/webui/src/static/style/overview/count.scss" did not exist on "30d2911662127156941ed075adc5c9c12cc82838"
Commit f9437603 authored by Khalique Ahmed's avatar Khalique Ahmed
Browse files

Merge branch 'develop' of https://github.com/ROCmSoftwarePlatform/AMDMIGraphX into mi100_opts

parents 781ce146 658cdab0
Hide whitespace changes
Inline Side-by-side
Showing with 536 additions and 3 deletions
test/simplify_reshapes_test.cpp
View file @ f9437603
......@@ -517,4 +517,423 @@ TEST_CASE(double_slice_multi_axes)
EXPECT(m1 == m2);
}
TEST_CASE(optimize_resize)
{
migraphx::shape sx{migraphx::shape::float_type, {1, 1, 2, 2}};
auto create_resize_module = [&] {
migraphx::module m;
auto inx = m.add_parameter("X", sx);
migraphx::shape si{migraphx::shape::int32_type, {1, 2, 4, 6}};
std::vector<int> ind = {0, 0, 0, 1, 1, 1, 0, 0, 0, 1, 1, 1, 2, 2, 2, 3,
3, 3, 2, 2, 2, 3, 3, 3, 0, 0, 0, 1, 1, 1, 0, 0,
0, 1, 1, 1, 2, 2, 2, 3, 3, 3, 2, 2, 2, 3, 3, 3};
auto li = m.add_literal(migraphx::literal(si, ind));
auto lrsp = m.add_instruction(migraphx::make_op("reshape", {{"dims", {4}}}), inx);
auto gr = m.add_instruction(migraphx::make_op("gather", {{"axis", 0}}), lrsp, li);
auto r = m.add_instruction(migraphx::make_op("softmax", {{"axis", 1}}), gr);
m.add_return({r});
return m;
};
auto m1 = create_resize_module();
run_pass(m1);
auto create_optimized_module = [&] {
migraphx::module m;
auto inx = m.add_parameter("X", sx);
std::vector<int64_t> dims = {1, 1, 2, 1, 2, 1};
auto rspx = m.add_instruction(migraphx::make_op("reshape", {{"dims", dims}}), inx);
std::vector<int64_t> mb_dims = {1, 2, 2, 2, 2, 3};
auto mbx = m.add_instruction(
migraphx::make_op("multibroadcast", {{"output_lens", mb_dims}}), rspx);
auto std_mb = m.add_instruction(migraphx::make_op("contiguous"), mbx);
std::vector<int64_t> orig_dims = {1, 2, 4, 6};
auto rmb = m.add_instruction(migraphx::make_op("reshape", {{"dims", orig_dims}}), std_mb);
auto r = m.add_instruction(migraphx::make_op("softmax", {{"axis", 1}}), rmb);
m.add_return({r});
return m;
};
EXPECT(m1 == create_optimized_module());
}
TEST_CASE(optimize_resize_ind_not_apply)
{
migraphx::shape sx{migraphx::shape::float_type, {1, 1, 2, 2}};
auto create_resize_module = [&] {
migraphx::module m;
auto inx = m.add_parameter("X", sx);
migraphx::shape si{migraphx::shape::int32_type, {1, 2, 4, 6}};
std::vector<int> ind = {0, 0, 0, 1, 1, 1, 0, 0, 0, 1, 0, 1, 2, 2, 2, 3,
3, 3, 2, 2, 2, 3, 3, 3, 0, 0, 0, 1, 1, 1, 0, 0,
0, 1, 1, 1, 2, 2, 2, 3, 3, 3, 2, 2, 2, 3, 3, 3};
auto li = m.add_literal(migraphx::literal(si, ind));
auto lrsp = m.add_instruction(migraphx::make_op("reshape", {{"dims", {4}}}), inx);
auto gr = m.add_instruction(migraphx::make_op("gather", {{"axis", 0}}), lrsp, li);
auto r = m.add_instruction(migraphx::make_op("softmax", {{"axis", 1}}), gr);
m.add_return({r});
return m;
};
auto m1 = create_resize_module();
run_pass(m1);
EXPECT(m1 == create_resize_module());
}
TEST_CASE(optimize_resize_rsp_dim_1)
{
migraphx::shape sx{migraphx::shape::float_type, {1, 1, 2, 2}};
auto create_resize_module = [&] {
migraphx::module m;
auto inx = m.add_parameter("X", sx);
migraphx::shape si{migraphx::shape::int32_type, {1, 1, 4, 3, 2}};
std::vector<int> ind = {0, 0, 0, 1, 1, 1, 0, 0, 0, 1, 1, 1,
2, 2, 2, 3, 3, 3, 2, 2, 2, 3, 3, 3};
auto li = m.add_literal(migraphx::literal(si, ind));
auto lrsp = m.add_instruction(migraphx::make_op("reshape", {{"dims", {2, 2}}}), inx);
auto r = m.add_instruction(migraphx::make_op("gather", {{"axis", 0}}), lrsp, li);
m.add_return({r});
return m;
};
auto m = create_resize_module();
run_pass(m);
EXPECT(m == create_resize_module());
}
TEST_CASE(optimize_resize_ndims_unequal)
{
migraphx::shape sx{migraphx::shape::float_type, {1, 1, 2, 2}};
migraphx::shape sy{migraphx::shape::float_type, {1, 1, 4, 3, 2}};
auto create_resize_module = [&] {
migraphx::module m;
auto inx = m.add_parameter("X", sx);
auto iny = m.add_parameter("Y", sy);
migraphx::shape si{migraphx::shape::int32_type, {1, 1, 4, 3, 2}};
std::vector<int> ind = {0, 0, 0, 1, 1, 1, 0, 0, 0, 1, 1, 1,
2, 2, 2, 3, 3, 3, 2, 2, 2, 3, 3, 3};
auto li = m.add_literal(migraphx::literal(si, ind));
auto lrsp = m.add_instruction(migraphx::make_op("reshape", {{"dims", {4}}}), inx);
auto gr = m.add_instruction(migraphx::make_op("gather", {{"axis", 0}}), lrsp, li);
auto r = m.add_instruction(migraphx::make_op("sub"), iny, gr);
m.add_return({r});
return m;
};
auto m = create_resize_module();
run_pass(m);
EXPECT(m == create_resize_module());
}
TEST_CASE(optimize_resize_ind_non_brcst)
{
migraphx::shape sx{migraphx::shape::float_type, {1, 1, 3, 2}};
migraphx::shape sy{migraphx::shape::float_type, {1, 1, 4, 6}};
auto create_resize_module = [&] {
migraphx::module m;
auto inx = m.add_parameter("X", sx);
auto iny = m.add_parameter("Y", sy);
migraphx::shape si{migraphx::shape::int32_type, {1, 1, 4, 6}};
std::vector<int> ind = {0, 0, 0, 1, 1, 1, 0, 0, 0, 1, 1, 1,
2, 2, 2, 3, 3, 3, 2, 2, 2, 3, 3, 3};
auto li = m.add_literal(migraphx::literal(si, ind));
auto lrsp = m.add_instruction(migraphx::make_op("reshape", {{"dims", {6}}}), inx);
auto gr = m.add_instruction(migraphx::make_op("gather", {{"axis", 0}}), lrsp, li);
auto r = m.add_instruction(migraphx::make_op("sub"), iny, gr);
m.add_return({r});
return m;
};
auto m = create_resize_module();
run_pass(m);
EXPECT(m == create_resize_module());
}
TEST_CASE(optimize_resize_ind_non_const)
{
migraphx::shape sx{migraphx::shape::float_type, {1, 1, 3, 2}};
migraphx::shape sy{migraphx::shape::float_type, {1, 1, 4, 6}};
auto create_resize_module = [&] {
migraphx::module m;
auto inx = m.add_parameter("X", sx);
auto iny = m.add_parameter("Y", sy);
migraphx::shape si{migraphx::shape::int32_type, {1, 1, 4, 6}};
auto li = m.add_parameter("ind", si);
auto lrsp = m.add_instruction(migraphx::make_op("reshape", {{"dims", {6}}}), inx);
auto gr = m.add_instruction(migraphx::make_op("gather", {{"axis", 0}}), lrsp, li);
auto r = m.add_instruction(migraphx::make_op("sub"), iny, gr);
m.add_return({r});
return m;
};
auto m = create_resize_module();
run_pass(m);
EXPECT(m == create_resize_module());
}
TEST_CASE(optimize_where_true)
{
migraphx::shape s{migraphx::shape::float_type, {1, 1, 3, 2}};
auto create_where_module = [&](bool cond) {
migraphx::module m;
auto inx = m.add_parameter("X", s);
auto iny = m.add_parameter("Y", s);
migraphx::shape si{migraphx::shape::bool_type, {1, 1, 3, 2}};
std::vector<char> idata(si.elements(), static_cast<char>(cond));
auto li = m.add_literal(migraphx::literal(si, idata));
auto data = m.add_instruction(migraphx::make_op("concat", {{"axis", 0}}), inx, iny);
auto data_1 = m.add_instruction(migraphx::make_op("reshape", {{"dims", {12}}}), data);
auto r = m.add_instruction(migraphx::make_op("gather", {{"axis", 0}}), data_1, li);
m.add_return({r});
return m;
};
auto create_opt_module = [&](std::string name) {
migraphx::module m;
auto in = m.add_parameter(std::move(name), s);
m.add_return({in});
return m;
};
auto m = create_where_module(true);
run_pass(m);
EXPECT(m == create_opt_module("X"));
auto m1 = create_where_module(false);
run_pass(m1);
EXPECT(m1 == create_opt_module("Y"));
}
TEST_CASE(where_different_cond_values)
{
auto create_where_module = [] {
migraphx::module m;
migraphx::shape s{migraphx::shape::float_type, {1, 1, 3, 2}};
auto inx = m.add_parameter("X", s);
auto iny = m.add_parameter("Y", s);
migraphx::shape si{migraphx::shape::bool_type, {1, 1, 3, 2}};
std::vector<char> idata = {1, 1, 0, 1, 0, 1};
auto li = m.add_literal(migraphx::literal(si, idata));
auto data = m.add_instruction(migraphx::make_op("concat", {{"axis", 0}}), inx, iny);
auto data_1 = m.add_instruction(migraphx::make_op("reshape", {{"dims", {12}}}), data);
auto r = m.add_instruction(migraphx::make_op("gather", {{"axis", 0}}), data_1, li);
m.add_return({r});
return m;
};
auto m = create_where_module();
run_pass(m);
EXPECT(m == create_where_module());
}
TEST_CASE(where_axis_nonzero)
{
auto create_where_module = [] {
migraphx::module m;
migraphx::shape s{migraphx::shape::float_type, {1, 1, 3, 2}};
auto inx = m.add_parameter("X", s);
auto iny = m.add_parameter("Y", s);
migraphx::shape si{migraphx::shape::bool_type, {1, 1, 3, 2}};
std::vector<char> idata(6, 1);
auto li = m.add_literal(migraphx::literal(si, idata));
auto data = m.add_instruction(migraphx::make_op("concat", {{"axis", 1}}), inx, iny);
auto data_1 = m.add_instruction(migraphx::make_op("reshape", {{"dims", {12}}}), data);
auto r = m.add_instruction(migraphx::make_op("gather", {{"axis", 0}}), data_1, li);
m.add_return({r});
return m;
};
auto m = create_where_module();
run_pass(m);
EXPECT(m == create_where_module());
}
TEST_CASE(where_three_concat_inputs)
{
auto create_where_module = [] {
migraphx::module m;
migraphx::shape s{migraphx::shape::float_type, {1, 1, 3, 2}};
auto inx = m.add_parameter("X", s);
auto iny = m.add_parameter("Y", s);
migraphx::shape si{migraphx::shape::bool_type, {1, 1, 3, 2}};
std::vector<char> idata(6, 1);
auto li = m.add_literal(migraphx::literal(si, idata));
auto data = m.add_instruction(migraphx::make_op("concat", {{"axis", 0}}), inx, iny, inx);
auto data_1 = m.add_instruction(migraphx::make_op("reshape", {{"dims", {18}}}), data);
auto r = m.add_instruction(migraphx::make_op("gather", {{"axis", 0}}), data_1, li);
m.add_return({r});
return m;
};
auto m = create_where_module();
run_pass(m);
EXPECT(m == create_where_module());
}
TEST_CASE(where_three_inputs_diff_shapes)
{
auto create_where_module = [] {
migraphx::module m;
migraphx::shape sx{migraphx::shape::float_type, {1, 1, 3, 2}};
migraphx::shape sy{migraphx::shape::float_type, {2, 1, 3, 2}};
auto inx = m.add_parameter("X", sx);
auto iny = m.add_parameter("Y", sy);
migraphx::shape si{migraphx::shape::bool_type, {1, 1, 3, 2}};
std::vector<char> idata(6, 1);
auto li = m.add_literal(migraphx::literal(si, idata));
auto data = m.add_instruction(migraphx::make_op("concat", {{"axis", 0}}), inx, iny);
auto data_1 = m.add_instruction(migraphx::make_op("reshape", {{"dims", {18}}}), data);
auto r = m.add_instruction(migraphx::make_op("gather", {{"axis", 0}}), data_1, li);
m.add_return({r});
return m;
};
auto m = create_where_module();
run_pass(m);
EXPECT(m == create_where_module());
}
TEST_CASE(where_three_lens_diff)
{
auto create_where_module = [] {
migraphx::module m;
migraphx::shape sx{migraphx::shape::float_type, {1, 1, 3, 2}};
migraphx::shape sy{migraphx::shape::float_type, {1, 1, 3, 2}};
auto inx = m.add_parameter("X", sx);
auto iny = m.add_parameter("Y", sy);
migraphx::shape si{migraphx::shape::bool_type, {1, 1, 6}};
std::vector<char> idata(6, 1);
auto li = m.add_literal(migraphx::literal(si, idata));
auto data = m.add_instruction(migraphx::make_op("concat", {{"axis", 0}}), inx, iny);
auto data_1 = m.add_instruction(migraphx::make_op("reshape", {{"dims", {12}}}), data);
auto r = m.add_instruction(migraphx::make_op("gather", {{"axis", 0}}), data_1, li);
m.add_return({r});
return m;
};
auto m = create_where_module();
run_pass(m);
EXPECT(m == create_where_module());
}
TEST_CASE(reshape_cont)
{
auto create_module = [] {
migraphx::module m;
migraphx::shape sx{migraphx::shape::float_type, {1, 4, 1}};
migraphx::shape sy{migraphx::shape::float_type, {2, 2, 2, 6}};
auto inx = m.add_parameter("x", sx);
auto iny = m.add_parameter("y", sy);
auto mb_inx = m.add_instruction(
migraphx::make_op("multibroadcast", {{"output_lens", {2, 4, 6}}}), inx);
auto std_inx = m.add_instruction(migraphx::make_op("contiguous"), mb_inx);
auto rsp =
m.add_instruction(migraphx::make_op("reshape", {{"dims", {2, 2, 2, 6}}}), std_inx);
auto r = m.add_instruction(migraphx::make_op("add"), rsp, iny);
m.add_return({r});
return m;
};
auto m1 = create_module();
run_pass(m1);
auto create_opt_module = [] {
migraphx::module m;
migraphx::shape sx{migraphx::shape::float_type, {1, 4, 1}};
migraphx::shape sy{migraphx::shape::float_type, {2, 2, 2, 6}};
auto inx = m.add_parameter("x", sx);
auto iny = m.add_parameter("y", sy);
auto mb_inx = m.add_instruction(
migraphx::make_op("multibroadcast", {{"output_lens", {2, 4, 6}}}), inx);
auto rsp_iny = m.add_instruction(migraphx::make_op("reshape", {{"dims", {2, 4, 6}}}), iny);
auto sum = m.add_instruction(migraphx::make_op("add"), mb_inx, rsp_iny);
auto r = m.add_instruction(migraphx::make_op("reshape", {{"dims", {2, 2, 2, 6}}}), sum);
m.add_return({r});
return m;
};
EXPECT(m1 == create_opt_module());
}
TEST_CASE(reshape_input_non_std)
{
auto create_module = [] {
migraphx::module m;
migraphx::shape sx{migraphx::shape::float_type, {1, 4, 1}};
migraphx::shape sy{migraphx::shape::float_type, {2, 6, 2, 2}};
auto inx = m.add_parameter("x", sx);
auto iny = m.add_parameter("y", sy);
auto mb_inx = m.add_instruction(
migraphx::make_op("multibroadcast", {{"output_lens", {2, 4, 6}}}), inx);
auto std_inx = m.add_instruction(migraphx::make_op("contiguous"), mb_inx);
auto rsp =
m.add_instruction(migraphx::make_op("reshape", {{"dims", {2, 2, 2, 6}}}), std_inx);
auto ty = m.add_instruction(migraphx::make_op("transpose", {{"dims", {0, 2, 3, 1}}}), iny);
auto r = m.add_instruction(migraphx::make_op("add"), rsp, ty);
m.add_return({r});
return m;
};
auto m1 = create_module();
run_pass(m1);
EXPECT(m1 == create_module());
}
TEST_CASE(reshape_cont_nonpw)
{
auto create_module = [] {
migraphx::module m;
migraphx::shape sx{migraphx::shape::float_type, {1, 4, 1}};
migraphx::shape sy{migraphx::shape::float_type, {2, 2, 2, 6}};
auto inx = m.add_parameter("x", sx);
auto iny = m.add_parameter("y", sy);
auto mb_inx = m.add_instruction(
migraphx::make_op("multibroadcast", {{"output_lens", {2, 4, 6}}}), inx);
auto std_inx = m.add_instruction(migraphx::make_op("contiguous"), mb_inx);
auto rsp =
m.add_instruction(migraphx::make_op("reshape", {{"dims", {2, 2, 2, 6}}}), std_inx);
auto r = m.add_instruction(migraphx::make_op("convolution"), rsp, iny);
m.add_return({r});
return m;
};
auto m1 = create_module();
run_pass(m1);
EXPECT(m1 == create_module());
}
int main(int argc, const char* argv[]) { test::run(argc, argv); }
test/value_test.cpp 100755 → 100644
View file @ f9437603
......@@ -486,6 +486,12 @@ TEST_CASE(value_emplace_object)
EXPECT(v["three"].get_key() == "three");
}
TEST_CASE(value_bracket_convert_throws)
{
migraphx::value v1;
EXPECT(test::throws([&] { v1["key"].to<std::string>(); }));
}
TEST_CASE(value_construct_object_string_value)
{
migraphx::value v = {{"one", "onev"}, {"two", "twov"}};
......@@ -798,4 +804,35 @@ TEST_CASE(value_binary_object_conv)
EXPECT(migraphx::equal(v["data"].get_binary(), data));
}
template <class T>
bool is_null_type(T)
{
return false;
}
bool is_null_type(std::nullptr_t) { return true; }
TEST_CASE(visit_null)
{
migraphx::value v;
EXPECT(v.is_null());
bool visited = false;
v.visit([&](auto&& x) { visited = is_null_type(x); });
EXPECT(visited);
}
TEST_CASE(value_or_convert)
{
migraphx::value v = 1;
EXPECT(v.is_int64());
EXPECT(v.value_or(3) == 1);
}
TEST_CASE(value_or_null)
{
migraphx::value v;
EXPECT(v.is_null());
EXPECT(v.value_or(3) == 3);
}
int main(int argc, const char* argv[]) { test::run(argc, argv); }
test/verify/run_verify.cpp
View file @ f9437603
......@@ -158,10 +158,12 @@ void run_verify::verify(const std::string& name, const migraphx::program& p) con
detach_async([=] { return run_target(t, p, m); }, ti.parallel));
}
assert(gold_f.valid());
auto gold = gold_f.get();
for(auto&& pp : results)
{
assert(pp.second.valid());
auto tname = pp.first;
auto x = pp.second.get();
auto cp = x.first;
......
test/verify/run_verify.hpp 100644 → 100755
View file @ f9437603
File mode changed from 100644 to 100755
test/verify/test_add_relu_add.cpp 0 → 100755
View file @ f9437603
#include "verify_program.hpp"
#include <migraphx/program.hpp>
#include <migraphx/generate.hpp>
#include <migraphx/make_op.hpp>
struct test_add_relu_add : verify_program<test_add_relu_add>
{
migraphx::program create_program() const
{
migraphx::program p;
auto* mm = p.get_main_module();
auto x = mm->add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 5, 2, 2}});
auto y = mm->add_parameter("y", migraphx::shape{migraphx::shape::float_type, {1, 5, 2, 2}});
auto z = mm->add_parameter("z", migraphx::shape{migraphx::shape::float_type, {1, 5, 2, 2}});
auto a = mm->add_instruction(migraphx::make_op("add"), x, y);
auto b = mm->add_instruction(migraphx::make_op("relu"), a);
auto c = mm->add_instruction(migraphx::make_op("add"), b, z);
mm->add_instruction(migraphx::make_op("relu"), c);
return p;
}
};
test/verify/test_log.cpp 100644 → 100755
View file @ f9437603
File mode changed from 100644 to 100755
tools/include/operation.hpp
View file @ f9437603
......@@ -341,6 +341,29 @@ auto has_finalize_op(const T&) -> decltype(has_finalize_op(rank<1>{},
return {};
}
template <class T>
auto compile_op(
rank<1>, T& x, context& ctx, const shape& output_shape, const std::vector<shape>& input)
-> decltype(x.compile(auto_any_cast(ctx), output_shape, input))
{
return x.compile(auto_any_cast(ctx), output_shape, input);
}
template <class T>
value compile_op(rank<0>, T&, context&, const shape&, const std::vector<shape>&)
{
return value::object{};
}
template <class T>
value compile_op(const T& x,
context& ctx,
const shape& output_shape,
const std::vector<shape>& input)
{
return compile_op(rank<1>{}, x, ctx, output_shape, input);
}
template <class T>
value attributes_op(const T&)
{
......@@ -361,6 +384,12 @@ void from_value_op(T& x, const value& v)
return migraphx::from_value(v, x);
}
template <class T>
bool is_borrowed_op(const T&)
{
return false;
}
} // namespace detail
<%
......@@ -374,11 +403,18 @@ void from_value_op(T& x, const value& v)
const = True,
default = 'detail::need_normalization_op'),
virtual('has_finalize', returns = 'bool', const = True, default = 'detail::has_finalize_op'),
virtual('is_borrowed', returns = 'bool', const = True, default = 'detail::is_borrowed_op'),
virtual('output_alias',
returns = 'std::ptrdiff_t',
input = 'const std::vector<shape>&',
const = True,
default = 'detail::output_alias_op'),
virtual('compile',
returns = 'value',
ctx = 'context&',
output = 'const shape&',
input = 'const std::vector<shape>&',
default = 'detail::compile_op'),
virtual('finalize',
ctx = 'context&',
output = 'const shape&',
......@@ -436,6 +472,24 @@ void from_value_op(T& x, const value& v)
return !(x == y);
}
inline value
compile(operation& op, context& ctx, const shape& output_shape, const std::vector<shape>& input)
{
return op.compile(ctx, output_shape, input);
}
template <class Context>
inline value
compile(operation& op, Context& ctx, const shape& output_shape, const std::vector<shape>& input)
{
dependent_type<context, Context> ctx2 = std::ref(ctx);
return compile(op, ctx2, output_shape, input);
}
template <class T, class Context>
inline auto compile(T& op, Context& ctx, const shape& output_shape, const std::vector<shape>& input)
-> decltype(op.compile(ctx, ctx, output_shape, input))
{
return op.compile(ctx, ctx, output_shape, input);
}
inline shape compute_shape(const operation& op, const std::vector<shape>& inputs)
{
return op.compute_shape(inputs);
......
tools/install_prereqs.sh
View file @ f9437603
......@@ -5,9 +5,6 @@
#install pip3, rocm-cmake, rocblas and miopen
apt update && apt install -y python3-pip rocm-cmake rocblas miopen-hip openmp-extras
# install onnx package for unit tests
pip3 install onnx==1.7.0 numpy==1.18.5 typing==3.7.4 pytest==6.0.1
# install rbuild to build dependencies
pip3 install https://github.com/RadeonOpenCompute/rbuild/archive/master.tar.gz
......@@ -33,4 +30,6 @@ cget -p $PREFIX init --cxx /opt/rocm/llvm/bin/clang++ --cc /opt/rocm/llvm/bin/cl
cget -p $PREFIX install -f ${REQ_FILE_DIR}dev-requirements.txt
cget -p $PREFIX install oneapi-src/oneDNN@v1.7
# install onnx package for unit tests
pip3 install onnx==1.8.1 numpy==1.18.5 typing==3.7.4 pytest==6.0.1 packaging==16.8
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment