Skip to content

GitLab

Commit 024de2a5 authored by Paul's avatar Paul
Browse files

Merge branch 'jit-layernorm' into jit-layernorm-merge

parents f568547e 6a8c231e
Hide whitespace changes
Inline Side-by-side
Showing with 382 additions and 109 deletions
src/include/migraphx/ranges.hpp
View file @ 024de2a5
......@@ -216,6 +216,12 @@ bool equal(R1&& r1, R2&& r2, Predicate... pred)
return std::equal(r1.begin(), r1.end(), r2.begin(), r2.end(), pred...);
}
template <class Range>
auto distance(Range&& r)
{
return std::distance(r.begin(), r.end());
}
template <class R>
using range_value = std::decay_t<decltype(*std::declval<R>().begin())>;
......
src/include/migraphx/stringutils.hpp
View file @ 024de2a5
......@@ -44,8 +44,8 @@ auto with_char(F f)
return [=](unsigned char c) -> bool { return f(c); };
}
inline std::string
replace_string(std::string subject, const std::string& search, const std::string& replace)
inline void
replace_string_inplace(std::string& subject, const std::string& search, const std::string& replace)
{
size_t pos = 0;
while((pos = subject.find(search, pos)) != std::string::npos)
......@@ -53,6 +53,12 @@ replace_string(std::string subject, const std::string& search, const std::string
subject.replace(pos, search.length(), replace);
pos += replace.length();
}
}
inline std::string
replace_string(std::string subject, const std::string& search, const std::string& replace)
{
replace_string_inplace(subject, search, replace);
return subject;
}
......
src/inline_module.cpp
View file @ 024de2a5
......@@ -35,7 +35,7 @@ static void inline_submodule(module& m, instruction_ref ins, bool cond)
{
const auto& mod_inputs = ins->module_inputs();
module_ref smod = cond ? mod_inputs.at(0) : mod_inputs.at(1);
auto mod_outputs = m.insert_module_instructions(ins, smod);
auto mod_outputs = m.insert_instructions(ins, smod);
auto ins_outputs = ins->outputs();
assert(mod_outputs.size() >= ins_outputs.size());
......
src/module.cpp
View file @ 024de2a5
......@@ -35,6 +35,7 @@
#include <migraphx/make_op.hpp>
#include <migraphx/register_target.hpp>
#include <migraphx/make_op.hpp>
#include <migraphx/json.hpp>
#include <iostream>
#include <sstream>
#include <algorithm>
......@@ -196,6 +197,62 @@ void module::assign(const module& m)
}
}
template <class Range>
static std::vector<instruction_ref>
insert_generic_instructions(module& m,
instruction_ref ins,
Range&& instructions,
std::unordered_map<instruction_ref, instruction_ref> map_ins)
{
assert(m.has_instruction(ins) or is_end(ins, m.end()));
std::vector<instruction_ref> mod_outputs;
instruction_ref last;
for(instruction_ref sins : instructions)
{
last = sins;
if(contains(map_ins, sins))
continue;
instruction_ref copy_ins;
if(sins->name() == "@literal")
{
auto l = sins->get_literal();
copy_ins = m.add_literal(l);
}
else if(sins->name() == "@param")
{
auto&& name = any_cast<builtin::param>(sins->get_operator()).parameter;
auto s = sins->get_shape();
copy_ins = m.add_parameter(name, s);
}
else if(sins->name() == "@outline")
{
auto s = sins->get_shape();
copy_ins = m.add_outline(s);
}
else
{
auto mod_args = sins->module_inputs();
auto inputs = sins->inputs();
std::vector<instruction_ref> copy_inputs(inputs.size());
std::transform(inputs.begin(), inputs.end(), copy_inputs.begin(), [&](auto i) {
return contains(map_ins, i) ? map_ins[i] : i;
});
if(sins->name() == "@return")
{
mod_outputs = copy_inputs;
break;
}
copy_ins = m.insert_instruction(ins, sins->get_operator(), copy_inputs, mod_args);
}
map_ins[sins] = copy_ins;
}
if(mod_outputs.empty() and instructions.begin() != instructions.end())
mod_outputs = {map_ins.at(last)};
return mod_outputs;
}
instruction_ref module::add_instruction(const operation& op, std::vector<instruction_ref> args)
{
return insert_instruction(impl->instructions.end(), op, std::move(args));
......@@ -334,53 +391,49 @@ instruction_ref module::move_instructions(instruction_ref src, instruction_ref d
return src;
}
std::vector<instruction_ref> module::insert_module_instructions(
instruction_ref ins, module_ref m, std::unordered_map<instruction_ref, instruction_ref> map_ins)
std::vector<instruction_ref>
module::add_instructions(const std::vector<instruction_ref>& instructions,
std::unordered_map<instruction_ref, instruction_ref> map_ins)
{
std::vector<instruction_ref> mod_outputs;
for(auto sins : iterator_for(*m))
{
if(contains(map_ins, sins))
continue;
instruction_ref copy_ins;
if(sins->name() == "@literal")
{
auto l = sins->get_literal();
copy_ins = this->add_literal(l);
}
else if(sins->name() == "@param")
{
auto&& name = any_cast<builtin::param>(sins->get_operator()).parameter;
auto s = sins->get_shape();
copy_ins = this->add_parameter(name, s);
}
else if(sins->name() == "@outline")
{
auto s = sins->get_shape();
copy_ins = this->add_outline(s);
}
else
{
auto mod_args = sins->module_inputs();
auto inputs = sins->inputs();
std::vector<instruction_ref> copy_inputs(inputs.size());
std::transform(inputs.begin(), inputs.end(), copy_inputs.begin(), [&](auto i) {
return contains(map_ins, i) ? map_ins[i] : i;
});
return this->insert_instructions(this->end(), instructions, std::move(map_ins));
}
if(sins->name() == "@return")
{
mod_outputs = copy_inputs;
break;
}
std::vector<instruction_ref>
module::add_instructions(module_ref m, std::unordered_map<instruction_ref, instruction_ref> map_ins)
{
return this->insert_instructions(this->end(), m, std::move(map_ins));
}
copy_ins = this->insert_instruction(ins, sins->get_operator(), copy_inputs, mod_args);
}
map_ins[sins] = copy_ins;
}
if(mod_outputs.empty())
mod_outputs = {map_ins.at(std::prev(m->end()))};
return mod_outputs;
std::vector<instruction_ref>
module::add_instructions(instruction_ref start,
instruction_ref last,
std::unordered_map<instruction_ref, instruction_ref> map_ins)
{
return this->insert_instructions(this->end(), start, last, std::move(map_ins));
}
std::vector<instruction_ref>
module::insert_instructions(instruction_ref ins,
const std::vector<instruction_ref>& instructions,
std::unordered_map<instruction_ref, instruction_ref> map_ins)
{
return insert_generic_instructions(*this, ins, instructions, std::move(map_ins));
}
std::vector<instruction_ref> module::insert_instructions(
instruction_ref ins, module_ref m, std::unordered_map<instruction_ref, instruction_ref> map_ins)
{
return insert_generic_instructions(*this, ins, iterator_for(*m), std::move(map_ins));
}
std::vector<instruction_ref>
module::insert_instructions(instruction_ref ins,
instruction_ref start,
instruction_ref last,
std::unordered_map<instruction_ref, instruction_ref> map_ins)
{
auto r = range(start, last);
return insert_generic_instructions(*this, ins, iterator_for(r), std::move(map_ins));
}
instruction_ref module::add_literal(literal l)
......@@ -706,44 +759,33 @@ void module::print_graph(std::ostream& os, bool brief) const
os << "}" << std::endl;
}
static std::string cpp_var_name(const std::string& name)
static std::string to_c_id(const std::string& name, char rep = '_')
{
return "m" + replace_string(name, "@", "x");
std::string id = transform_string(name, [&](auto c) {
if(with_char(::isalnum)(c) or c == '_')
return c;
return rep;
});
while(contains(id, "__"))
replace_string_inplace(id, "__", "_");
return id;
}
static std::string cpp_op_var(const std::string& name, instruction_ref ins)
static std::string cpp_var_name(const std::string& name)
{
return replace_string(name, "@", ins->name());
return to_c_id("x_" + replace_string(name, ":", "_module_"));
}
static void print_op_attributes(std::ostream& os, const std::string& name, const operation& op)
static void print_make_op(std::ostream& os, const operation& op)
{
std::string x = to_string(op);
if(contains(x, "["))
os << "migraphx::make_op(" << enclose_name(op.name());
auto v = op.to_value();
if(not v.empty())
{
auto start = x.find('[');
auto end = x.find(']');
std::string attribute_text = x.substr(start + 1, end - start - 1);
std::vector<std::string> attributes;
for(auto&& attribute : split_string(attribute_text, ','))
{
if(contains(attribute, '='))
attributes.push_back(attribute);
else
attributes.back() += "," + attribute;
}
for(auto&& attribute : attributes)
{
auto p = split_string(attribute, '=');
auto key = p.front();
auto value = p.back();
if(contains({"bn_mode", "padding_mode"}, key))
continue;
if(key == "mode")
value = enclose_name(trim(value));
os << name << "." << key << " = " << value << ";" << std::endl;
}
os << ", "
<< "migraphx::from_json_string(" << enclose_name(to_json_string(v)) << ")";
}
os << ")";
}
static void print_cpp_shape(std::ostream& os, const migraphx::shape& s)
......@@ -756,22 +798,25 @@ static void print_cpp_shape(std::ostream& os, const migraphx::shape& s)
}
std::unordered_map<instruction_ref, std::string>
module::print_cpp(std::ostream& os, std::unordered_map<instruction_ref, std::string> names) const
module::print_cpp(std::ostream& os,
const std::string& mname,
std::unordered_map<instruction_ref, std::string> names) const
{
os << "migraphx::module p;" << std::endl;
unsigned long seed = 0;
// cppcheck-suppress variableScope
unsigned long seed = names.size();
auto last = std::prev(this->end());
names = this->print(
[&](auto ins, auto ins_names) {
auto op = cpp_op_var(ins_names.at(ins), ins);
if(ins->name().front() != '@')
{
os << "migraphx::op::" << ins->name() << " " << op << ";" << std::endl;
print_op_attributes(os, op, ins->get_operator());
}
os << "auto " << cpp_var_name(ins_names.at(ins)) << " = ";
std::vector<std::string> input_vars;
std::transform(ins->inputs().begin(),
ins->inputs().end(),
std::back_inserter(input_vars),
[&](auto input) { return cpp_var_name(ins_names.at(input)); });
if(ins != last)
os << "auto " << cpp_var_name(ins_names.at(ins)) << " = ";
if(ins->name() == "@literal")
{
os << "p.add_literal(";
os << mname << "->add_literal(";
bool use_abs = false;
ins->get_literal().visit([&](auto v) {
use_abs = std::none_of(v.begin(), v.end(), [](auto x) { return x < 0; });
......@@ -789,17 +834,22 @@ module::print_cpp(std::ostream& os, std::unordered_map<instruction_ref, std::str
else if(ins->name() == "@param")
{
std::string name = any_cast<builtin::param>(ins->get_operator()).parameter;
os << "p.add_parameter(" << enclose_name(name) << ",";
os << mname << "->add_parameter(" << enclose_name(name) << ",";
print_cpp_shape(os, ins->get_shape());
os << ");" << std::endl;
}
else if(ins->name() == "@return")
{
os << mname << "->add_return({";
os << join_strings(input_vars, ", ");
os << "});" << std::endl;
}
else
{
os << "p.add_instruction(" << op;
for(auto input : ins->inputs())
{
os << ", " << cpp_var_name(ins_names.at(input));
}
assert(ins->name().front() != '@');
os << mname << "->add_instruction(";
print_make_op(os, ins->get_operator());
os << ", " << join_strings(input_vars, ", ");
os << ");" << std::endl;
}
},
......@@ -808,7 +858,7 @@ module::print_cpp(std::ostream& os, std::unordered_map<instruction_ref, std::str
return names;
}
void module::print_cpp(std::ostream& os) const { this->print_cpp(os, {}); }
void module::print_cpp(std::ostream& os) const { this->print_cpp(os, this->name(), {}); }
void module::annotate(std::ostream& os, std::function<void(instruction_ref)> a) const
{
......@@ -819,17 +869,20 @@ void module::annotate(std::ostream& os, std::function<void(instruction_ref)> a)
});
}
std::vector<module_ref> module::get_sub_modules() const
std::vector<module_ref> module::get_sub_modules(bool shallow) const
{
std::vector<module_ref> vec_modules;
for(auto ins : iterator_for(*this))
{
const auto& mod_args = ins->module_inputs();
vec_modules.insert(vec_modules.end(), mod_args.begin(), mod_args.end());
for(const auto& smod : mod_args)
if(not shallow)
{
auto sub_mods = smod->get_sub_modules();
vec_modules.insert(vec_modules.end(), sub_mods.begin(), sub_mods.end());
for(const auto& smod : mod_args)
{
auto sub_mods = smod->get_sub_modules();
vec_modules.insert(vec_modules.end(), sub_mods.begin(), sub_mods.end());
}
}
}
......
src/pass_manager.cpp
View file @ 024de2a5
......@@ -66,14 +66,12 @@ void run_pass(program& prog, const pass& p, tracer trace)
struct module_pm : module_pass_manager
{
module* mod;
program* prog;
tracer* t;
module* mod = nullptr;
tracer* t = nullptr;
module* common_parent = nullptr;
program* prog = nullptr;
module_pm(module* pmod = nullptr, program* pprog = nullptr, tracer* pt = nullptr)
: mod(pmod), prog(pprog), t(pt)
{
}
module_pm(module* pmod = nullptr, tracer* pt = nullptr) : mod(pmod), t(pt) {}
template <class... Ts>
void trace(Ts&&... xs) const
......@@ -92,6 +90,7 @@ struct module_pm : module_pass_manager
assert(prog);
return prog->create_module(name);
}
virtual module* get_common_parent() override { return common_parent; }
virtual void run_pass(const pass& p) override
{
assert(mod);
......@@ -111,7 +110,7 @@ void run_passes(module& mod, const std::vector<pass>& passes, tracer trace)
trace = tracer{std::cout};
for(const auto& p : passes)
{
module_pm{&mod, nullptr, &trace}.run_pass(p);
module_pm{&mod, &trace}.run_pass(p);
}
}
......@@ -119,14 +118,31 @@ void run_passes(program& prog, const std::vector<pass>& passes, tracer trace)
{
if(enabled(MIGRAPHX_TRACE_PASSES{}))
trace = tracer{std::cout};
std::unordered_set<module_ref> visited;
for(const auto& p : passes)
{
auto mods = prog.get_modules();
auto tree = prog.get_module_tree();
visited.clear();
for(const auto& mod : reverse(mods))
{
if(mod->bypass())
continue;
module_pm{mod, &prog, &trace}.run_pass(p);
if(not visited.insert(mod).second)
continue;
module_pm mpm{mod, &trace};
mpm.prog = &prog;
auto parents = range(tree.equal_range(mod));
auto nparents = distance(parents);
if(nparents == 0)
mpm.common_parent = nullptr;
else if(nparents == 1)
mpm.common_parent = parents.begin()->second;
else
// Just set common parent to main module when there is muliple parents for now
// TODO: Compute the common parent
mpm.common_parent = prog.get_main_module();
mpm.run_pass(p);
}
run_pass(prog, p, trace);
}
......
src/program.cpp
View file @ 024de2a5
......@@ -790,10 +790,17 @@ void program::print_cpp(std::ostream& os) const
{
auto vec_modules = this->get_modules();
std::unordered_map<instruction_ref, std::string> names;
os << "migraphx::program p;\n";
for(auto& mod : vec_modules)
{
os << "module: \"" << mod->name() << "\"" << std::endl;
names = mod->print_cpp(os, names);
std::string var_name = "m" + mod->name();
os << "migraphx::module_ref " << var_name << " = ";
if(mod->name() == "main")
os << "p.get_main_module();";
else
os << "p.create_module(\"" << mod->name() << "\");";
os << std::endl;
names = mod->print_cpp(os, var_name, names);
os << std::endl;
}
}
......@@ -869,6 +876,23 @@ std::vector<module*> program::get_modules()
return result;
}
template <class Module, class Map>
void generic_insert_module_tree(Module* pm, Map& m)
{
for(auto* sm : pm->get_sub_modules(true))
{
m.insert(std::make_pair(sm, pm));
generic_insert_module_tree(sm, m);
}
}
std::unordered_multimap<module_ref, module_ref> program::get_module_tree()
{
std::unordered_multimap<module_ref, module_ref> result;
generic_insert_module_tree(this->get_main_module(), result);
return result;
}
template <class Map, class T>
bool is_unused_module(Map& m, const std::vector<T*>& mods, const std::string& name)
{
......
src/targets/gpu/compile_hip.cpp
View file @ 024de2a5
......@@ -43,6 +43,7 @@ inline namespace MIGRAPHX_INLINE_NS {
namespace gpu {
MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_GPU_DEBUG);
MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_GPU_DEBUG_SYM);
MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_GPU_OPTIMIZE);
MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_GPU_DUMP_ASM);
MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_GPU_DUMP_SRC);
......@@ -227,6 +228,8 @@ compile_hip_src(const std::vector<src_file>& srcs, std::string params, const std
if(params.find("-std=") == std::string::npos)
params += " --std=c++17";
params += " -fno-gpu-rdc";
if(enabled(MIGRAPHX_GPU_DEBUG_SYM{}))
params += " -g";
params += " -c";
if(is_hcc_compiler())
{
......
src/targets/gpu/fuse_ops.cpp
View file @ 024de2a5
......@@ -701,6 +701,7 @@ struct miopen_fusion
return args.back();
}
};
MIGRAPHX_REGISTER_OP(miopen_fusion)
struct miopen_conv_bias
{
......
test/module_test.cpp
View file @ 024de2a5
......@@ -300,6 +300,96 @@ TEST_CASE(parameter_name_order)
EXPECT(param_names == names1);
}
TEST_CASE(insert_instructions_module)
{
migraphx::shape s{migraphx::shape::int32_type, {1}};
migraphx::module m1("m1");
auto x1 = m1.add_parameter("x1", s);
auto sqrt = m1.add_instruction(migraphx::make_op("sqrt"), {x1});
m1.add_instruction(migraphx::make_op("add"), {sqrt, x1});
migraphx::module m2("m2");
auto x2 = m2.add_parameter("x2", s);
m2.add_instruction(migraphx::make_op("sqrt"), {x2});
m1.insert_instructions(sqrt, &m2, {{x2, x1}});
EXPECT(std::prev(sqrt)->name() == "sqrt");
EXPECT(std::count_if(m1.begin(), m1.end(), [](auto&& ins) { return ins.name() == "sqrt"; }) ==
2);
EXPECT(std::count_if(m1.begin(), m1.end(), [](auto&& ins) { return ins.name() == "@param"; }) ==
1);
EXPECT(contains(m1.get_parameter_shapes(), "x1"));
EXPECT(not contains(m1.get_parameter_shapes(), "x2"));
}
TEST_CASE(add_instructions_module)
{
migraphx::shape s{migraphx::shape::int32_type, {1}};
migraphx::module m1("m1");
auto x1 = m1.add_parameter("x1", s);
m1.add_instruction(migraphx::make_op("sqrt"), {x1});
migraphx::module m2("m2");
auto x2 = m2.add_parameter("x2", s);
m2.add_instruction(migraphx::make_op("sqrt"), {x2});
m1.add_instructions(&m2, {{x2, x1}});
EXPECT(std::count_if(m1.begin(), m1.end(), [](auto&& ins) { return ins.name() == "sqrt"; }) ==
2);
EXPECT(std::count_if(m1.begin(), m1.end(), [](auto&& ins) { return ins.name() == "@param"; }) ==
1);
EXPECT(contains(m1.get_parameter_shapes(), "x1"));
EXPECT(not contains(m1.get_parameter_shapes(), "x2"));
}
TEST_CASE(add_instructions_range)
{
migraphx::shape s{migraphx::shape::int32_type, {1}};
migraphx::module m1("m1");
auto x1 = m1.add_parameter("x1", s);
m1.add_instruction(migraphx::make_op("sqrt"), {x1});
migraphx::module m2("m2");
auto x2 = m2.add_parameter("x2", s);
auto sqrt2 = m2.add_instruction(migraphx::make_op("sqrt"), {x2});
m1.add_instructions(sqrt2, m2.end(), {{x2, x1}});
EXPECT(std::any_of(
m1.begin(), m1.end(), [&](auto&& ins) { return migraphx::contains(ins.inputs(), x1); }));
EXPECT(std::count_if(m1.begin(), m1.end(), [](auto&& ins) { return ins.name() == "sqrt"; }) ==
2);
EXPECT(std::count_if(m1.begin(), m1.end(), [](auto&& ins) { return ins.name() == "@param"; }) ==
1);
EXPECT(contains(m1.get_parameter_shapes(), "x1"));
EXPECT(not contains(m1.get_parameter_shapes(), "x2"));
}
TEST_CASE(add_instructions_vector)
{
migraphx::shape s{migraphx::shape::int32_type, {1}};
migraphx::module m1("m1");
auto x1 = m1.add_parameter("x1", s);
m1.add_instruction(migraphx::make_op("sqrt"), {x1});
migraphx::module m2("m2");
auto x2 = m2.add_parameter("x2", s);
auto sqrt2 = m2.add_instruction(migraphx::make_op("sqrt"), {x2});
m1.add_instructions({sqrt2}, {{x2, x1}});
EXPECT(std::any_of(
m1.begin(), m1.end(), [&](auto&& ins) { return migraphx::contains(ins.inputs(), x1); }));
EXPECT(std::count_if(m1.begin(), m1.end(), [](auto&& ins) { return ins.name() == "sqrt"; }) ==
2);
EXPECT(std::count_if(m1.begin(), m1.end(), [](auto&& ins) { return ins.name() == "@param"; }) ==
1);
EXPECT(contains(m1.get_parameter_shapes(), "x1"));
EXPECT(not contains(m1.get_parameter_shapes(), "x2"));
}
struct check_for_pass_op
{
bool* found = nullptr;
......
test/ref_ops_test.cpp
View file @ 024de2a5
......@@ -3187,6 +3187,80 @@ TEST_CASE(nms_test)
EXPECT(migraphx::verify_range(result, gold));
}
TEST_CASE(nms_transpose1_test)
{
migraphx::program p;
auto* mm = p.get_main_module();
migraphx::shape boxes_s{migraphx::shape::float_type, {1, 4, 6}};
std::vector<float> boxes_vec = {
0.5, 0.5, 0.5, 0.5, 0.5, 0.5, 0.5, 0.6, 0.4, 10.5, 10.6, 100.5,
1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0,
};
migraphx::shape scores_s{migraphx::shape::float_type, {1, 1, 6}};
std::vector<float> scores_vec = {0.9, 0.75, 0.6, 0.95, 0.5, 0.3};
auto t_boxes_l = mm->add_literal(migraphx::literal(boxes_s, boxes_vec));
auto scores_l = mm->add_literal(migraphx::literal(scores_s, scores_vec));
auto max_out_l = mm->add_literal(int64_t{4});
auto iou_threshold = mm->add_literal(0.5f);
auto score_threshold = mm->add_literal(0.0f);
auto transpose_boxes = mm->add_instruction(
migraphx::make_op("transpose", {{"permutation", {0, 2, 1}}}), t_boxes_l);
auto r = mm->add_instruction(migraphx::make_op("nonmaxsuppression", {{"center_point_box", 1}}),
transpose_boxes,
scores_l,
max_out_l,
iou_threshold,
score_threshold);
mm->add_return({r});
p.compile(migraphx::ref::target{});
auto output = p.eval({}).back();
std::vector<int64_t> result;
output.visit([&](auto out) { result.assign(out.begin(), out.end()); });
std::vector<int64_t> gold = {0, 0, 3, 0, 0, 0, 0, 0, 5, 0, 0, 0, 0, 0, 0, 0, 0, 0};
EXPECT(migraphx::verify_range(result, gold));
}
TEST_CASE(nms_transpose2_test)
{
migraphx::program p;
auto* mm = p.get_main_module();
migraphx::shape boxes_s{migraphx::shape::float_type, {4, 1, 6}};
std::vector<float> boxes_vec = {
0.5, 0.5, 0.5, 0.5, 0.5, 0.5, 0.5, 0.6, 0.4, 10.5, 10.6, 100.5,
1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0,
};
migraphx::shape scores_s{migraphx::shape::float_type, {1, 1, 6}};
std::vector<float> scores_vec = {0.9, 0.75, 0.6, 0.95, 0.5, 0.3};
auto t_boxes_l = mm->add_literal(migraphx::literal(boxes_s, boxes_vec));
auto scores_l = mm->add_literal(migraphx::literal(scores_s, scores_vec));
auto max_out_l = mm->add_literal(int64_t{4});
auto iou_threshold = mm->add_literal(0.5f);
auto score_threshold = mm->add_literal(0.0f);
auto transpose_boxes = mm->add_instruction(
migraphx::make_op("transpose", {{"permutation", {1, 2, 0}}}), t_boxes_l);
auto r = mm->add_instruction(migraphx::make_op("nonmaxsuppression", {{"center_point_box", 1}}),
transpose_boxes,
scores_l,
max_out_l,
iou_threshold,
score_threshold);
mm->add_return({r});
p.compile(migraphx::ref::target{});
auto output = p.eval({}).back();
std::vector<int64_t> result;
output.visit([&](auto out) { result.assign(out.begin(), out.end()); });
std::vector<int64_t> gold = {0, 0, 3, 0, 0, 0, 0, 0, 5, 0, 0, 0, 0, 0, 0, 0, 0, 0};
EXPECT(migraphx::verify_range(result, gold));
}
TEST_CASE(nonzero_test)
{
migraphx::program p;
......
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