Skip to content

GitLab

Unverified Commit 500d9441 authored by Paul Fultz II's avatar Paul Fultz II Committed by GitHub
Browse files

Split cpu and reference implementation (#671) 



* Add all_targets cmake target

* Rename target

* Add ref target

* Rename tests

* Refactor compiler target

* Formatting

* Verify for every target

* Formatting

* Add verify test suite

* Formatting

* Add initial test programs

* Formatting

* Add rnn tests

* Formatting

* Validate gpu

* Formatting

* Remove old gpu tests

* Fix gpu tests

* Fix ref error

* Fix tidy issues

* Formatting

* Tidy fixes

* Fix header in python api

* Rename to ref

* Use ref in verify_onnx

* Fix tidy issue

* Build with verbose on

* Fix typo

* Remove verbose

* rename some cpu prefix to ref
Co-authored-by: default avatarShucai Xiao <Shucai.Xiao@amd.com>
parent ba33d25c
Hide whitespace changes
Inline Side-by-side
Showing with 1179 additions and 5231 deletions
src/targets/ref/include/migraphx/ref/gemm.hpp 0 → 100644
View file @ 500d9441
#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 MIGRAPHX_INLINE_NS {
namespace ref {
void migemm(
const argument& c_arg, const argument& a_arg, const argument& b_arg, float alpha, float beta);
void migemm(const argument& c_arg,
const argument& a_arg,
const argument& b_arg,
int32_t alpha,
int32_t beta);
} // namespace ref
} // namespace MIGRAPHX_INLINE_NS
} // namespace migraphx
#endif
src/targets/ref/include/migraphx/ref/lowering.hpp 0 → 100644
View file @ 500d9441
#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 MIGRAPHX_INLINE_NS {
namespace ref {
struct lowering
{
std::string name() const { return "ref::lowering"; }
void apply(program& p) const;
};
} // namespace ref
} // namespace MIGRAPHX_INLINE_NS
} // namespace migraphx
#endif
src/targets/ref/include/migraphx/ref/target.hpp 0 → 100644
View file @ 500d9441
#ifndef MIGRAPHX_GUARD_MIGRAPHLIB_CPU_TARGET_HPP
#define MIGRAPHX_GUARD_MIGRAPHLIB_CPU_TARGET_HPP
#include <migraphx/program.hpp>
#include <migraphx/register_target.hpp>
#include <migraphx/compile_options.hpp>
#include <migraphx/ref/context.hpp>
#include <migraphx/config.hpp>
namespace migraphx {
inline namespace MIGRAPHX_INLINE_NS {
struct pass;
namespace ref {
struct target
{
std::string name() const;
std::vector<pass> get_passes(migraphx::context& ctx, const compile_options&) const;
migraphx::context get_context() const { return context{}; }
argument copy_to(const argument& arg) const { return arg; }
argument copy_from(const argument& arg) const { return arg; }
argument allocate(const shape& s) const;
};
MIGRAPHX_REGISTER_TARGET(target);
} // namespace ref
} // namespace MIGRAPHX_INLINE_NS
} // namespace migraphx
#endif
src/targets/ref/lowering.cpp 0 → 100644
View file @ 500d9441
#include <migraphx/ref/lowering.hpp>
#include <migraphx/instruction.hpp>
#include <migraphx/dfor.hpp>
#include <migraphx/op/identity.hpp>
#include <migraphx/op/batch_norm_inference.hpp>
#include <migraphx/op/convolution.hpp>
#include <migraphx/op/deconvolution.hpp>
#include <migraphx/op/quant_convolution.hpp>
#include <migraphx/op/dot.hpp>
#include <migraphx/op/quant_dot.hpp>
#include <migraphx/op/elu.hpp>
#include <migraphx/op/im2col.hpp>
#include <migraphx/op/leaky_relu.hpp>
#include <migraphx/op/logsoftmax.hpp>
#include <migraphx/op/lrn.hpp>
#include <migraphx/op/pad.hpp>
#include <migraphx/op/pooling.hpp>
#include <migraphx/op/softmax.hpp>
#include <migraphx/op/argmax.hpp>
#include <migraphx/op/argmin.hpp>
#include <migraphx/op/rnn_var_sl_last_output.hpp>
#include <migraphx/shape_for_each.hpp>
#include <migraphx/iterator_for.hpp>
#include <migraphx/par_dfor.hpp>
#include <migraphx/clamp.hpp>
#include <migraphx/ref/gemm.hpp>
#include <migraphx/register_op.hpp>
#include <migraphx/make_op.hpp>
#include <unordered_map>
#include <utility>
#include <iostream>
namespace migraphx {
inline namespace MIGRAPHX_INLINE_NS {
namespace ref {
template <typename T>
T zero(const T&)
{
return T(0);
}
template <class T>
typename std::conditional_t<std::is_integral<T>{}, std::make_signed<T>, std::enable_if<true, T>>::
type
make_signed(T x)
{
return x;
}
//
// ref implemenataion of batch norm for inference
//
// inputs are:
// args[0] -> input data buffer
// args[1] -> mini batch mean
// args[2] -> mini batch variance
// args[3] -> gamma
// args[4] -> bias
//
// The equation to compute batch norm for inference is:
//
// output[i] = bias + gamma * (input[i] + mean) / sqrt(variance + epsilon)
//
// the input data format should be nchw
//
struct ref_batch_norm_inference
{
op::batch_norm_inference op;
template <class Self, class F>
static auto reflect(Self& self, F f)
{
return migraphx::reflect(self.op, f);
}
std::string name() const { return "ref::batch_norm_inference"; }
shape compute_shape(const std::vector<shape>& inputs) const { return op.compute_shape(inputs); }
argument compute(context&, const shape& output_shape, std::vector<argument> args) const
{
argument output{output_shape};
double epsilon = op.epsilon;
auto input = args[0];
auto arg_gamma = args[1];
auto arg_bias = args[2];
auto mini_batch_mean = args[3];
auto mini_batch_variance = args[4];
if(op.bn_mode == op::batch_norm_inference::spatial)
{
visit_all(output, input, mini_batch_mean, mini_batch_variance, arg_gamma, arg_bias)(
[&](auto result, auto buffer, auto mean, auto variance, auto gamma, auto bias) {
par_for(output_shape.elements(), [&](auto i) {
auto idx = output_shape.multi(i);
auto c = idx[1];
assert((variance[c] + epsilon) > 0);
result[i] =
gamma[c] * (buffer[i] - mean[c]) / std::sqrt(variance[c] + epsilon) +
bias[c];
});
});
}
if(op.bn_mode == op::batch_norm_inference::per_activation)
{
visit_all(output, input, mini_batch_mean, mini_batch_variance, arg_gamma, arg_bias)(
[&](auto result, auto buffer, auto mean, auto variance, auto gamma, auto bias) {
par_for(output_shape.elements(), [&](auto i) {
auto idx = output_shape.multi(i);
idx[0] = 0;
auto index = output_shape.index(idx);
assert((variance[index] + epsilon) > 0);
result[i] = gamma[index] * (buffer[i] - mean[index]) /
std::sqrt(variance[index] + epsilon) +
bias[index];
});
});
}
return output;
}
};
MIGRAPHX_REGISTER_OP(ref_batch_norm_inference)
struct ref_lrn
{
op::lrn op;
template <class Self, class F>
static auto reflect(Self& self, F f)
{
return migraphx::reflect(self.op, f);
}
std::string name() const { return "ref::lrn"; }
shape compute_shape(const std::vector<shape>& inputs) const { return op.compute_shape(inputs); }
argument compute(context&, shape output_shape, std::vector<argument> args) const
{
argument result{output_shape};
visit_all(result, args[0])([&](auto output, auto input) {
int n_batch = output_shape.lens()[0];
int channels = output_shape.lens()[1];
int height = output_shape.lens()[2];
int width = output_shape.lens()[3];
float alphaoverarea = op.alpha / float(op.size);
int radius_lower = (op.size - 1) / 2;
int radius_upper = op.size / 2 + 1;
par_dfor(n_batch, height, width)([&](int b, int h, int w) {
float scale = 0;
dfor(channels)([&](int c) {
auto start = (c - radius_lower) < 0 ? 0 : (c - radius_lower);
auto end = (c + radius_upper) > channels ? channels : (c + radius_upper);
for(auto k = start; k < end; ++k)
{
scale += std::pow(input(b, k, h, w), 2);
}
scale *= alphaoverarea;
scale += op.bias;
scale = std::pow(scale, -op.beta);
output(b, c, h, w) = input(b, c, h, w) * scale;
});
});
});
return result;
}
};
MIGRAPHX_REGISTER_OP(ref_lrn)
template <class V, class T, class... Ts>
void visit_quantize_impl(V&& v, T&& x, Ts&&... xs)
{
x.visit([&](auto y) { visit_all(xs...)([&](auto... ys) { v(y, ys...); }); });
}
template <class T, class... Ts>
auto visit_quantize(T&& x, Ts&&... xs)
{
return [&](auto v) {
// Workaround for https://gcc.gnu.org/bugzilla/show_bug.cgi?id=70100
visit_quantize_impl(v, x, xs...);
};
}
template <class Op>
struct ref_convolution : auto_register_op<ref_convolution<Op>>
{
ref_convolution() = default;
ref_convolution(Op pop) : op(std::move(pop)) {}
Op op;
template <class Self, class F>
static auto reflect(Self& self, F f)
{
return migraphx::reflect(self.op, f);
}
std::string name() const { return "ref::" + op.name(); }
shape compute_shape(const std::vector<shape>& inputs) const { return op.compute_shape(inputs); }
argument compute(context&, shape output_shape, std::vector<argument> args) const
{
argument result{output_shape};
visit_quantize(result, args[0], args[1])([&](auto output, auto input, auto weights) {
auto in_lens = input.get_shape().lens();
auto wei_lens = weights.get_shape().lens();
auto wei_n = wei_lens[0];
auto wei_c = wei_lens[1];
std::vector<std::size_t> win_size(wei_lens.begin() + 1, wei_lens.end());
par_for(output_shape.elements(), [&](auto i) {
auto idx_o = output_shape.multi(i);
auto w = idx_o[1];
auto n_dim = idx_o.size();
std::vector<std::ptrdiff_t> win_start;
for(std::size_t dim = 2; dim < n_dim; ++dim)
{
auto d_2 = dim - 2;
win_start.push_back(std::ptrdiff_t(idx_o[dim] * op.stride[d_2]) -
std::ptrdiff_t(op.padding[d_2]));
}
const auto group_id = w / (wei_n / op.group);
shape win_shape{output_shape.type(), win_size};
double acc = 0.0;
shape_for_each(win_shape, [&](auto idx_win) {
auto k = idx_win[0];
const auto in_ch = group_id * wei_c + k;
std::vector<std::ptrdiff_t> idx(idx_o.begin(), idx_o.end());
idx[1] = in_ch;
std::transform(idx_win.begin() + 1,
idx_win.end(),
win_start.begin(),
idx.begin() + 2,
[](std::ptrdiff_t ii, std::ptrdiff_t jj) { return ii + jj; });
std::vector<std::ptrdiff_t> idx_wei(idx_o.size());
idx_wei[0] = w;
std::copy(idx_win.begin(), idx_win.end(), idx_wei.begin() + 1);
if(std::all_of(idx.begin() + 2, idx.end(), [&](auto ii) { return ii >= 0; }) and
std::equal(idx.begin(),
idx.end(),
in_lens.begin(),
in_lens.end(),
std::less<std::ptrdiff_t>{}))
{
acc +=
input(idx.begin(), idx.end()) * weights(idx_wei.begin(), idx_wei.end());
}
});
output[i] = acc;
});
});
return result;
}
};
template <class Op>
struct ref_deconvolution : auto_register_op<ref_deconvolution<Op>>
{
ref_deconvolution() = default;
ref_deconvolution(Op pop) : op(std::move(pop)) {}
Op op;
template <class Self, class F>
static auto reflect(Self& self, F f)
{
return migraphx::reflect(self.op, f);
}
std::string name() const { return "ref::" + op.name(); }
shape compute_shape(const std::vector<shape>& inputs) const { return op.compute_shape(inputs); }
argument compute(context&, shape output_shape, std::vector<argument> args) const
{
argument result{output_shape};
visit_all(result, args[0], args[1])([&](auto output, auto input, auto weights) {
using type = typename decltype(output)::value_type;
std::fill(output.begin(), output.end(), type{0});
auto in_lens = input.get_shape().lens();
auto in_n = in_lens[0];
auto in_c = in_lens[1];
auto wei = weights.get_shape().lens();
auto wei_n = wei[0];
auto wei_c = wei[1];
auto out_lens = output_shape.lens();
auto kdims = op.kdims();
std::vector<std::size_t> win_size{in_c};
std::copy(in_lens.begin() + 2, in_lens.end(), std::back_inserter(win_size));
std::copy(wei.begin() + 2, wei.end(), std::back_inserter(win_size));
shape win_shape{output_shape.type(), win_size};
par_dfor(in_n, wei_c)([&](int o, int k) {
shape_for_each(win_shape, [&](auto idx_win) {
const int w = idx_win[0];
auto input_dims_start = idx_win.begin() + 1;
auto wei_dims_start = idx_win.begin() + kdims + 1;
std::vector<std::ptrdiff_t> win_start;
for(std::size_t n = 0; n < kdims; ++n)
{
win_start.push_back(std::ptrdiff_t(*(input_dims_start + n) * op.stride[n]) -
std::ptrdiff_t(op.padding[n]));
}
const int group_id = w / (wei_n / op.group);
const int in_ch = group_id * wei_c + k;
std::vector<std::ptrdiff_t> idx_out{o, in_ch};
for(size_t n = 0; n < kdims; n++)
{
idx_out.push_back(win_start[n] + *(wei_dims_start + n) * op.dilation[n]);
}
std::vector<std::ptrdiff_t> idx_wei{w, k};
std::copy(wei_dims_start, idx_win.end(), std::back_inserter(idx_wei));
std::vector<std::ptrdiff_t> idx_in{o, w};
std::copy(input_dims_start, wei_dims_start, std::back_inserter(idx_in));
if(std::all_of(
idx_out.begin() + 2, idx_out.end(), [&](auto ii) { return ii >= 0; }) and
std::equal(idx_out.begin() + 2,
idx_out.end(),
out_lens.begin() + 2,
out_lens.end(),
std::less<std::ptrdiff_t>{}))
{
output(idx_out.begin(), idx_out.end()) +=
input(idx_in.begin(), idx_in.end()) *
weights(idx_wei.begin(), idx_wei.end());
}
});
});
});
return result;
}
};
struct ref_im2col
{
op::im2col op;
template <class Self, class F>
static auto reflect(Self& self, F f)
{
return migraphx::reflect(self.op, f);
}
static std::string name() { return "ref::im2col"; }
shape compute_shape(const std::vector<shape>& inputs) const { return op.compute_shape(inputs); }
argument compute(context&, const shape& output_shape, std::vector<argument> args) const
{
argument result{output_shape};
auto input_shape = args[0].get_shape();
auto weights_shape = args[1].get_shape();
visit_all(result, args[0])([&](auto col, auto input) {
const std::size_t& height = input_shape.lens()[2];
const std::size_t& width = input_shape.lens()[3];
const std::size_t& channels = weights_shape.lens()[1];
const std::size_t& kernel_h = weights_shape.lens()[2];
const std::size_t& kernel_w = weights_shape.lens()[3];
const std::size_t& pad_h = op.padding[0];
const std::size_t& pad_w = op.padding[1];
const std::size_t& stride_h = op.stride[0];
const std::size_t& stride_w = op.stride[1];
long kdiv2_h = long(kernel_h) / 2;
long kdiv2_w = long(kernel_w) / 2;
// calculate output sizes
const std::size_t col_height = (height - kernel_h + 2 * pad_h) / stride_h + 1;
const std::size_t col_width = (width - kernel_w + 2 * pad_w) / stride_w + 1;
// account for padding for the starting position of the input pixels
long iinput = kdiv2_h - long(pad_h);
// loop over output pixels (ioutput, joutput)
for(std::size_t ioutput = 0; ioutput < col_height; ioutput++, iinput += stride_h)
{
long jinput = kdiv2_w - long(pad_w);
for(std::size_t joutput = 0; joutput < col_width; joutput++, jinput += stride_w)
{
// compute linear index for output
std::size_t ldx = ioutput * col_width + joutput;
std::size_t p = 0;
dfor(channels,
kernel_h,
kernel_w)([&](std::size_t c, std::size_t koffset, std::size_t loffset) {
auto idx = iinput + long(koffset) - kdiv2_h;
auto jdx = jinput + long(loffset) - kdiv2_w;
col(ldx, p) = ((idx >= 0) && (idx < height) && (jdx >= 0) && (jdx < width))
? input(0, c, idx, jdx)
: 0;
p++;
});
}
}
});
return result;
}
};
MIGRAPHX_REGISTER_OP(ref_im2col)
struct max_pool
{
static std::string name() { return "max"; }
template <class T>
static T start()
{
return std::numeric_limits<T>::lowest();
}
static double apply(double x, double y)
{
double m = std::max(x, y);
return (m);
}
static double final(double x, std::size_t) { return (x); }
};
struct avg_pool
{
static std::string name() { return "average"; }
template <class T>
static double start()
{
return 0.0;
}
static double apply(double x, double y) { return x + y; }
static double final(double x, std::size_t y) { return (y == 0) ? 0.0 : (x / y); }
};
template <class Op>
struct ref_pooling : auto_register_op<ref_pooling<Op>>
{
ref_pooling() = default;
ref_pooling(op::pooling pop) : op(std::move(pop)) {}
op::pooling op;
template <class Self, class F>
static auto reflect(Self& self, F f)
{
return migraphx::reflect(self.op, f);
}
std::string name() const { return "ref::pooling_" + Op::name(); }
shape compute_shape(const std::vector<shape>& inputs) const { return op.compute_shape(inputs); }
argument compute(context&, const shape& output_shape, std::vector<argument> args) const
{
argument result{output_shape};
visit_all(result, args[0])([&](auto output, auto input) {
using type = typename decltype(output)::value_type;
auto in_s = input.get_shape();
auto in_lens = in_s.lens();
std::vector<std::size_t> vec_len(in_lens.begin() + 2, in_lens.end());
par_for(output_shape.elements(), [&](auto i) {
auto idx_o = output_shape.multi(i);
auto n_dim = idx_o.size();
std::vector<std::size_t> win_start;
std::vector<std::size_t> win_size;
for(std::size_t dim = 2; dim < n_dim; ++dim)
{
auto d_2 = dim - 2;
int start = static_cast<int>(idx_o[dim] * op.stride[d_2]) -
static_cast<int>(op.padding[d_2]);
int end = std::min(start + op.lengths[d_2], in_lens[dim]);
start = std::max(start, 0);
win_start.push_back(start);
win_size.push_back(end - start);
}
shape win_shape{output_shape.type(), win_size};
auto pool_size = win_shape.elements();
double acc = Op::template start<type>();
shape_for_each(win_shape, [&](auto idx_w) {
auto idx = idx_o;
std::transform(idx_w.begin(),
idx_w.end(),
win_start.begin(),
idx.begin() + 2,
[](auto ii, auto jj) { return ii + jj; });
if(std::all_of(idx.begin() + 2, idx.end(), [&](auto ii) { return ii >= 0; }) and
idx < in_lens)
{
acc = Op::apply(acc, input[in_s.index(idx)]);
}
});
output[i] = type(Op::final(acc, pool_size));
});
});
return result;
}
};
struct ref_op
{
operation op = op::identity{};
template <class Self, class F>
static auto reflect(Self& self, F f)
{
return migraphx::reflect(self.op, f);
}
std::string name() const { return "ref::op"; }
shape compute_shape(const std::vector<shape>& inputs) const { return op.compute_shape(inputs); }
argument compute(context&, const shape& output_shape, const std::vector<argument>& args) const
{
return op.compute(output_shape, args);
}
value to_value() const
{
value v;
v["name"] = op.name();
v["operator"] = op.to_value();
return v;
}
void from_value(const value& v)
{
op = make_op(v.at("name").to<std::string>(), v.at("operator"));
}
friend std::ostream& operator<<(std::ostream& os, const ref_op& x)
{
os << "ref::" << x.op;
return os;
}
};
MIGRAPHX_REGISTER_OP(ref_op)
struct ref_pad
{
op::pad op;
template <class Self, class F>
static auto reflect(Self& self, F f)
{
return migraphx::reflect(self.op, f);
}
std::string name() const { return "ref::pad"; }
shape compute_shape(const std::vector<shape>& inputs) const { return op.compute_shape(inputs); }
argument compute(context&, const shape& output_shape, std::vector<argument> args) const
{
assert(output_shape.standard());
argument result{output_shape};
result.visit([&](auto output) {
using type = typename decltype(output)::value_type;
std::fill(output.begin(), output.end(), pad_clamp<type>(op.value));
});
visit_all(result, args[0])([&](auto output, auto input) {
shape_for_each(input.get_shape(), [&](const auto& idx) {
std::vector<std::size_t> new_idx(idx.size());
std::transform(
idx.begin(), idx.end(), op.pads.begin(), new_idx.begin(), [](auto i, auto j) {
return i + j;
});
output(new_idx.begin(), new_idx.end()) = input(idx.begin(), idx.end());
});
});
return result;
}
};
MIGRAPHX_REGISTER_OP(ref_pad)
struct ref_gemm
{
op::dot op;
template <class Self, class F>
static auto reflect(Self& self, F f)
{
return migraphx::reflect(self.op, f);
}
std::string name() const { return "ref::dot"; }
shape compute_shape(const std::vector<shape>& inputs) const
{
if(inputs.size() == 3)
{
auto c_shape = inputs.at(2);
check_shapes{{c_shape}, *this}.not_broadcasted();
}
return op.compute_shape(inputs);
}
argument compute(context&, const shape& output_shape, std::vector<argument> args) const
{
argument result{output_shape};
// 3 inputs, it is alpha * A * B + beta * C, then
// A and B are matrices, and C is of the same shape as A * B
if(args.size() == 3)
{
// no need to consider the value of args[2]
if(op.beta == 0.0f)
{
result.visit([&](auto output) { std::fill(output.begin(), output.end(), 0); });
}
else
{
visit_all(result, args[2])([&](auto output, auto input) {
std::copy(input.begin(), input.end(), output.begin());
});
}
migemm(result, args[0], args[1], op.alpha, op.beta);
return result;
}
// 2 input arguments
migemm(result, args[0], args[1], op.alpha, 0.0f);
return result;
}
};
MIGRAPHX_REGISTER_OP(ref_gemm)
struct ref_quant_gemm
{
op::quant_dot op;
template <class Self, class F>
static auto reflect(Self& self, F f)
{
return migraphx::reflect(self.op, f);
}
std::string name() const { return "ref::quant_dot"; }
shape compute_shape(const std::vector<shape>& inputs) const
{
if(inputs.size() == 3)
{
auto c_shape = inputs.at(2);
check_shapes{{c_shape}, *this}.not_broadcasted();
}
return op.compute_shape(inputs);
}
argument compute(context&, const shape& output_shape, std::vector<argument> args) const
{
argument result{output_shape};
// 3 inputs, it is alpha * A * B + beta * C, then
// A and B are matrices, and C is of the same shape to A * B
// first, convert the args[0] and args[1] from int8_t to int32_t
argument arg_0{{shape::int32_type, {args.at(0).get_shape().lens()}}};
argument arg_1{{shape::int32_type, {args.at(1).get_shape().lens()}}};
arg_0.visit([&](auto output) {
args.at(0).visit(
[&](auto input) { std::copy(input.begin(), input.end(), output.begin()); });
});
arg_1.visit([&](auto output) {
args.at(1).visit(
[&](auto input) { std::copy(input.begin(), input.end(), output.begin()); });
});
if(args.size() == 3)
{
// no need to consider the value of args[2]
if(op.beta == 0)
{
result.visit([&](auto output) { std::fill(output.begin(), output.end(), 0); });
}
else
{
visit_all(result, args[2])([&](auto output, auto input) {
std::copy(input.begin(), input.end(), output.begin());
});
}
migemm(result, arg_0, arg_1, op.alpha, op.beta);
return result;
}
// 2 input arguments
migemm(result, arg_0, arg_1, op.alpha, int32_t{0});
return result;
}
};
MIGRAPHX_REGISTER_OP(ref_gemm)
struct leaky_relu_op
{
op::leaky_relu op;
std::string name() const { return "ref::leaky_relu"; }
auto fcn() const
{
auto a = op.alpha;
return [a](auto x) { return x > 0 ? x : x * a; };
}
};
struct elu_op
{
op::elu op;
std::string name() const { return "ref::elu"; }
auto fcn() const
{
auto a = op.alpha;
return [a](auto x) { return x > 0 ? x : a * std::expm1(x); };
}
};
template <typename Op>
struct ref_unary : auto_register_op<ref_unary<Op>>
{
ref_unary() = default;
template <class T>
ref_unary(T pop) : op(Op{std::move(pop)})
{
}
Op op;
template <class Self, class F>
static auto reflect(Self& self, F f)
{
return migraphx::reflect(self.op.op, f);
}
std::string name() const { return op.name(); }
shape compute_shape(const std::vector<shape>& inputs) const
{
check_shapes{inputs, *this}.has(1);
auto s = inputs.at(0);
return {s.type(), s.lens()};
}
argument compute(context&, const shape& output_shape, std::vector<argument> args) const
{
argument result{output_shape};
visit_all(result, args[0])([&](auto output, auto input) {
assert(input.get_shape().standard());
std::transform(input.begin(), input.end(), output.begin(), op.fcn());
});
return result;
}
};
template <class Op>
struct ref_softmax : auto_register_op<ref_softmax<Op>>
{
ref_softmax() = default;
ref_softmax(Op pop) : op(std::move(pop)) {}
Op op;
template <class Self, class F>
static auto reflect(Self& self, F f)
{
return migraphx::reflect(self.op, f);
}
std::string name() const { return "ref::" + op.name(); }
shape compute_shape(const std::vector<shape>& inputs) const { return op.compute_shape(inputs); }
argument compute(context&, const shape& output_shape, std::vector<argument> args) const
{
argument result{output_shape};
auto batch_lens = output_shape.lens();
int64_t tuned_axis = (op.axis < 0) ? op.axis + args[0].get_shape().lens().size() : op.axis;
std::size_t n_dims = batch_lens[tuned_axis];
batch_lens[tuned_axis] = 1;
shape batch_shape{shape::int32_type, batch_lens};
visit_all(result, args[0])([&](auto output, auto input) {
using value_type = typename decltype(input)::value_type;
std::vector<value_type> batch_max(batch_shape.elements(),
std::numeric_limits<value_type>::lowest());
std::vector<value_type> batch_sum(batch_shape.elements(), value_type(0));
par_for(batch_shape.elements(), [&](auto i) {
auto idx = batch_shape.multi(i);
for(std::size_t j = 0; j < n_dims; ++j)
{
idx[tuned_axis] = j;
batch_max[i] = std::max(batch_max[i], input(idx.begin(), idx.end()));
}
for(std::size_t j = 0; j < n_dims; ++j)
{
idx[tuned_axis] = j;
std::size_t index = output_shape.index(idx);
output[index] = std::exp(input[index] - batch_max[i]);
}
for(std::size_t j = 0; j < n_dims; ++j)
{
idx[tuned_axis] = j;
batch_sum[i] += output(idx.begin(), idx.end());
}
for(std::size_t j = 0; j < n_dims; ++j)
{
idx[tuned_axis] = j;
output(idx.begin(), idx.end()) =
op.output()(output(idx.begin(), idx.end()), batch_sum[i]);
}
});
});
return result;
}
};
struct ref_rnn_var_sl_last_output
{
op::rnn_var_sl_last_output op;
template <class Self, class F>
static auto reflect(Self& self, F f)
{
return migraphx::reflect(self.op, f);
}
std::string name() const { return "ref::rnn_var_sl_last_output"; }
shape compute_shape(std::vector<shape> inputs) const
{
return op.compute_shape(std::move(inputs));
}
argument compute(const shape& output_shape, std::vector<argument> args) const
{
argument result{output_shape};
auto out_comp_lens = args[0].get_shape().lens();
out_comp_lens[0] = 1;
shape out_comp_s{output_shape.type(), out_comp_lens};
visit_all(result, args[0])([&](auto output, auto input) {
args[1].visit([&](auto seq_lens) {
par_for(output_shape.elements(), [&](auto i) {
auto idx = out_comp_s.multi(i);
auto b = idx[2];
if(op.direction == op::rnn_direction::reverse or idx[1] == 1)
{
idx[0] = 0;
}
else
{
idx[0] = seq_lens[b] - 1;
}
output[i] = input(idx.begin(), idx.end());
});
});
});
return result;
}
};
MIGRAPHX_REGISTER_OP(ref_rnn_var_sl_last_output)
struct ref_apply
{
program* prog;
std::unordered_map<std::string, std::function<void(instruction_ref)>> apply_map{};
template <class T>
auto simple_op()
{
return [this](instruction_ref ins) { apply_simple_op<T>(ins); };
}
template <class T, class Op>
auto extend_op()
{
return [this](instruction_ref ins) { apply_extend_op<T, Op>(ins); };
}
void init()
{
apply_map["batch_norm_inference"] =
extend_op<ref_batch_norm_inference, op::batch_norm_inference>();
apply_map["convolution"] = extend_op<ref_convolution<op::convolution>, op::convolution>();
apply_map["deconvolution"] =
extend_op<ref_deconvolution<op::deconvolution>, op::deconvolution>();
apply_map["dot"] = extend_op<ref_gemm, op::dot>();
apply_map["quant_dot"] = extend_op<ref_quant_gemm, op::quant_dot>();
apply_map["quant_convolution"] =
extend_op<ref_convolution<op::quant_convolution>, op::quant_convolution>();
apply_map["elu"] = extend_op<ref_unary<elu_op>, op::elu>();
apply_map["im2col"] = extend_op<ref_im2col, op::im2col>();
apply_map["leaky_relu"] = extend_op<ref_unary<leaky_relu_op>, op::leaky_relu>();
apply_map["logsoftmax"] = extend_op<ref_softmax<op::logsoftmax>, op::logsoftmax>();
apply_map["lrn"] = extend_op<ref_lrn, op::lrn>();
apply_map["pad"] = extend_op<ref_pad, op::pad>();
apply_map["softmax"] = extend_op<ref_softmax<op::softmax>, op::softmax>();
apply_map["rnn_var_sl_last_output"] =
extend_op<ref_rnn_var_sl_last_output, op::rnn_var_sl_last_output>();
}
void apply()
{
init();
for(auto it : iterator_for(*prog))
{
if(it->name() == "pooling")
{
apply_pooling(it);
}
else if(apply_map.count(it->name()) > 0)
{
apply_map.at(it->name())(it);
}
else if(is_context_free(it->get_operator()))
{
apply_ref_op(it);
}
}
}
void apply_ref_op(instruction_ref ins) const
{
prog->replace_instruction(ins, ref_op{ins->get_operator()}, ins->inputs());
}
template <class T>
void apply_simple_op(instruction_ref ins)
{
prog->replace_instruction(ins, T{}, ins->inputs());
}
template <class T, class Op>
void apply_extend_op(instruction_ref ins)
{
auto&& op = any_cast<Op>(ins->get_operator());
prog->replace_instruction(ins, T{op}, ins->inputs());
}
void apply_pooling(instruction_ref ins) const
{
auto&& op = any_cast<op::pooling>(ins->get_operator());
if(op.mode == "max")
prog->replace_instruction(ins, ref_pooling<max_pool>{op}, ins->inputs());
else if(op.mode == "average")
prog->replace_instruction(ins, ref_pooling<avg_pool>{op}, ins->inputs());
}
};
void lowering::apply(program& p) const { ref_apply{&p}.apply(); }
} // namespace ref
} // namespace MIGRAPHX_INLINE_NS
} // namespace migraphx
src/targets/ref/target.cpp 0 → 100644
View file @ 500d9441
#include <migraphx/ref/target.hpp>
#include <migraphx/ref/lowering.hpp>
#include <migraphx/register_target.hpp>
#include <migraphx/pass.hpp>
#include <migraphx/auto_contiguous.hpp>
#include <migraphx/rewrite_rnn.hpp>
#include <migraphx/dead_code_elimination.hpp>
#include <migraphx/generate.hpp>
namespace migraphx {
inline namespace MIGRAPHX_INLINE_NS {
namespace ref {
std::string target::name() const { return "ref"; }
std::vector<pass> target::get_passes(migraphx::context&, const compile_options&) const
{
return {rewrite_rnn{},
dead_code_elimination{},
auto_contiguous{},
dead_code_elimination{},
lowering{},
dead_code_elimination{}};
}
argument target::allocate(const shape& s) const { return fill_argument(s, 0); }
MIGRAPHX_REGISTER_TARGET(target);
} // namespace ref
} // namespace MIGRAPHX_INLINE_NS
} // namespace migraphx
src/verify_args.cpp
View file @ 500d9441
...@@ -5,75 +5,75 @@ namespace migraphx { ...@@ -5,75 +5,75 @@ namespace migraphx {
inline namespace MIGRAPHX_INLINE_NS { inline namespace MIGRAPHX_INLINE_NS {
bool verify_args(const std::string& name, bool verify_args(const std::string& name,
const argument& cpu_arg, const argument& ref_arg,
const argument& gpu_arg, const argument& target_arg,
double tolerance) double tolerance)
{ {
bool passed = true; bool passed = true;
visit_all(cpu_arg, gpu_arg)([&](auto cpu, auto gpu) { visit_all(ref_arg, target_arg)([&](auto ref, auto target) {
double error; double error;
passed = verify_range(cpu, gpu, tolerance, &error); passed = verify_range(ref, target, tolerance, &error);
if(not passed) if(not passed)
{ {
// TODO: Check for nans // TODO: Check for nans
std::cout << "FAILED: " << name << std::endl; std::cout << "FAILED: " << name << std::endl;
std::cout << "error: " << error << std::endl; std::cout << "error: " << error << std::endl;
if(cpu.size() < 32) if(ref.size() < 32)
std::cout << "cpu:" << cpu << std::endl; std::cout << "ref:" << ref << std::endl;
if(gpu.size() < 32) if(target.size() < 32)
std::cout << "gpu:" << gpu << std::endl; std::cout << "target:" << target << std::endl;
if(range_zero(cpu)) if(range_zero(ref))
std::cout << "Cpu data is all zeros" << std::endl; std::cout << "Ref data is all zeros" << std::endl;
if(range_zero(gpu)) if(range_zero(target))
std::cout << "Gpu data is all zeros" << std::endl; std::cout << "Target data is all zeros" << std::endl;
auto mxdiff = max_diff(cpu, gpu); auto mxdiff = max_diff(ref, target);
std::cout << "Max diff: " << mxdiff << std::endl; std::cout << "Max diff: " << mxdiff << std::endl;
auto idx = mismatch_idx(cpu, gpu, float_equal); auto idx = mismatch_idx(ref, target, float_equal);
if(idx < range_distance(cpu)) if(idx < range_distance(ref))
{ {
std::cout << "Mismatch at " << idx << ": " << cpu[idx] << " != " << gpu[idx] std::cout << "Mismatch at " << idx << ": " << ref[idx] << " != " << target[idx]
<< std::endl; << std::endl;
} }
auto cpu_nan_idx = find_idx(cpu, not_finite); auto ref_nan_idx = find_idx(ref, not_finite);
if(cpu_nan_idx >= 0) if(ref_nan_idx >= 0)
std::cout << "Non finite number found in cpu at " << cpu_nan_idx << ": " std::cout << "Non finite number found in ref at " << ref_nan_idx << ": "
<< cpu[cpu_nan_idx] << std::endl; << ref[ref_nan_idx] << std::endl;
auto gpu_nan_idx = find_idx(gpu, not_finite); auto target_nan_idx = find_idx(target, not_finite);
if(gpu_nan_idx >= 0) if(target_nan_idx >= 0)
std::cout << "Non finite number found in gpu at " << gpu_nan_idx << ": " std::cout << "Non finite number found in target at " << target_nan_idx << ": "
<< gpu[gpu_nan_idx] << std::endl; << target[target_nan_idx] << std::endl;
std::cout << std::endl; std::cout << std::endl;
} }
else else
{ {
if(range_zero(cpu)) if(range_zero(ref))
std::cout << "Cpu data is all zeros" << std::endl; std::cout << "Ref data is all zeros" << std::endl;
if(range_zero(gpu)) if(range_zero(target))
std::cout << "Gpu data is all zeros" << std::endl; std::cout << "Target data is all zeros" << std::endl;
// auto mxdiff = max_diff(cpu, gpu); // auto mxdiff = max_diff(ref, target);
// std::cout << "Max diff: " << mxdiff << std::endl; // std::cout << "Max diff: " << mxdiff << std::endl;
// auto idx = mismatch_idx(cpu, gpu, float_equal); // auto idx = mismatch_idx(ref, target, float_equal);
// if(idx < range_distance(cpu)) // if(idx < range_distance(ref))
// { // {
// std::cout << "Mismatch at " << idx << ": " << cpu[idx] << " != " << gpu[idx] // std::cout << "Mismatch at " << idx << ": " << ref[idx] << " != " << target[idx]
// << std::endl; // << std::endl;
// } // }
auto cpu_nan_idx = find_idx(cpu, not_finite); auto ref_nan_idx = find_idx(ref, not_finite);
if(cpu_nan_idx >= 0) if(ref_nan_idx >= 0)
std::cout << "Non finite number found in cpu at " << cpu_nan_idx << ": " std::cout << "Non finite number found in ref at " << ref_nan_idx << ": "
<< cpu[cpu_nan_idx] << std::endl; << ref[ref_nan_idx] << std::endl;
auto gpu_nan_idx = find_idx(gpu, not_finite); auto target_nan_idx = find_idx(target, not_finite);
if(gpu_nan_idx >= 0) if(target_nan_idx >= 0)
std::cout << "Non finite number found in gpu at " << gpu_nan_idx << ": " std::cout << "Non finite number found in target at " << target_nan_idx << ": "
<< gpu[gpu_nan_idx] << std::endl; << target[target_nan_idx] << std::endl;
// std::cout << std::endl; // std::cout << std::endl;
} }
}); });
......
test/CMakeLists.txt
View file @ 500d9441
...@@ -67,6 +67,7 @@ function(add_test_command NAME EXE) ...@@ -67,6 +67,7 @@ function(add_test_command NAME EXE)
add_test(NAME ${NAME} COMMAND ${EXE} ${ARGN}) add_test(NAME ${NAME} COMMAND ${EXE} ${ARGN})
endif() endif()
endif() endif()
set_tests_properties(${NAME} PROPERTIES FAIL_REGULAR_EXPRESSION "FAILED")
endfunction() endfunction()
function(add_test_executable TEST_NAME) function(add_test_executable TEST_NAME)
...@@ -85,8 +86,7 @@ function(add_test_executable TEST_NAME) ...@@ -85,8 +86,7 @@ function(add_test_executable TEST_NAME)
add_test_command(${TEST_NAME} ${TEST_COMMAND}) add_test_command(${TEST_NAME} ${TEST_COMMAND})
add_dependencies(tests ${TEST_NAME}) add_dependencies(tests ${TEST_NAME})
add_dependencies(check ${TEST_NAME}) add_dependencies(check ${TEST_NAME})
set_tests_properties(${TEST_NAME} PROPERTIES FAIL_REGULAR_EXPRESSION "FAILED") target_link_libraries(${TEST_NAME} migraphx migraphx_ref migraphx_onnx)
target_link_libraries(${TEST_NAME} migraphx migraphx_cpu migraphx_onnx)
target_include_directories(${TEST_NAME} PUBLIC include) target_include_directories(${TEST_NAME} PUBLIC include)
endfunction(add_test_executable) endfunction(add_test_executable)
...@@ -122,7 +122,7 @@ foreach(ONNX_TEST ${ONNX_TESTS}) ...@@ -122,7 +122,7 @@ foreach(ONNX_TEST ${ONNX_TESTS})
set(TEST_NAME test_${BASE_NAME}) set(TEST_NAME test_${BASE_NAME})
add_executable(${TEST_NAME} ${TES_ONNX_DIR}/${ONNX_TEST}) add_executable(${TEST_NAME} ${TES_ONNX_DIR}/${ONNX_TEST})
rocm_clang_tidy_check(${TEST_NAME}) rocm_clang_tidy_check(${TEST_NAME})
target_link_libraries(${TEST_NAME} migraphx_onnx migraphx_cpu) target_link_libraries(${TEST_NAME} migraphx_onnx migraphx_ref)
target_include_directories(${TEST_NAME} PUBLIC include) 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_test(NAME ${TEST_NAME} COMMAND $<TARGET_FILE:${TEST_NAME}> WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/onnx)
add_dependencies(tests ${TEST_NAME}) add_dependencies(tests ${TEST_NAME})
...@@ -132,13 +132,14 @@ endforeach() ...@@ -132,13 +132,14 @@ endforeach()
# tf test # tf test
add_executable(test_tf tf/tf_test.cpp) add_executable(test_tf tf/tf_test.cpp)
rocm_clang_tidy_check(test_tf) rocm_clang_tidy_check(test_tf)
target_link_libraries(test_tf migraphx_tf migraphx_cpu) target_link_libraries(test_tf migraphx_tf migraphx_ref)
target_include_directories(test_tf PUBLIC include) target_include_directories(test_tf PUBLIC include)
add_test(NAME test_tf COMMAND $<TARGET_FILE:test_tf> WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/tf) add_test(NAME test_tf COMMAND $<TARGET_FILE:test_tf> WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/tf)
add_dependencies(tests test_tf) add_dependencies(tests test_tf)
add_dependencies(check test_tf) add_dependencies(check test_tf)
add_subdirectory(api) add_subdirectory(api)
add_subdirectory(verify)
if(MIGRAPHX_ENABLE_PYTHON) if(MIGRAPHX_ENABLE_PYTHON)
add_subdirectory(py) add_subdirectory(py)
endif() endif()
...@@ -172,7 +173,7 @@ function(test_headers PREFIX) ...@@ -172,7 +173,7 @@ function(test_headers PREFIX)
endfunction() endfunction()
test_headers(migraphx ${CMAKE_SOURCE_DIR}/src/include/migraphx/*.hpp) test_headers(migraphx ${CMAKE_SOURCE_DIR}/src/include/migraphx/*.hpp)
test_headers(migraphx/cpu ${CMAKE_SOURCE_DIR}/src/targets/cpu/include/migraphx/cpu/*.hpp) test_headers(migraphx/ref ${CMAKE_SOURCE_DIR}/src/targets/ref/include/migraphx/ref/*.hpp)
if(MIGRAPHX_ENABLE_GPU) if(MIGRAPHX_ENABLE_GPU)
test_headers(migraphx/gpu ${CMAKE_SOURCE_DIR}/src/targets/gpu/include/migraphx/gpu/*.hpp) test_headers(migraphx/gpu ${CMAKE_SOURCE_DIR}/src/targets/gpu/include/migraphx/gpu/*.hpp)
endif() endif()
test/api/CMakeLists.txt
View file @ 500d9441
...@@ -11,7 +11,7 @@ function(add_api_test TEST_NAME TEST_SRC) ...@@ -11,7 +11,7 @@ function(add_api_test TEST_NAME TEST_SRC)
endfunction() endfunction()
add_api_test(cpu test_cpu.cpp) add_api_test(ref test_cpu.cpp)
add_api_test(save_load test_save_load.cpp) add_api_test(save_load test_save_load.cpp)
add_api_test(op test_op_construct.cpp) add_api_test(op test_op_construct.cpp)
if(MIGRAPHX_ENABLE_GPU) if(MIGRAPHX_ENABLE_GPU)
......
test/api/test_cpu.cpp
View file @ 500d9441
...@@ -6,7 +6,7 @@ TEST_CASE(load_and_run) ...@@ -6,7 +6,7 @@ TEST_CASE(load_and_run)
{ {
auto p = migraphx::parse_onnx("conv_relu_maxpool_test.onnx"); auto p = migraphx::parse_onnx("conv_relu_maxpool_test.onnx");
auto shapes_before = p.get_output_shapes(); auto shapes_before = p.get_output_shapes();
p.compile(migraphx::target("cpu")); p.compile(migraphx::target("ref"));
auto shapes_after = p.get_output_shapes(); auto shapes_after = p.get_output_shapes();
CHECK(shapes_before.size() == 1); CHECK(shapes_before.size() == 1);
CHECK(shapes_before.size() == shapes_after.size()); CHECK(shapes_before.size() == shapes_after.size());
...@@ -26,7 +26,7 @@ TEST_CASE(load_and_run_init_list) ...@@ -26,7 +26,7 @@ TEST_CASE(load_and_run_init_list)
{ {
auto p = migraphx::parse_onnx("conv_relu_maxpool_test.onnx"); auto p = migraphx::parse_onnx("conv_relu_maxpool_test.onnx");
auto shapes_before = p.get_output_shapes(); auto shapes_before = p.get_output_shapes();
p.compile(migraphx::target("cpu")); p.compile(migraphx::target("ref"));
auto shapes_after = p.get_output_shapes(); auto shapes_after = p.get_output_shapes();
CHECK(shapes_before.size() == 1); CHECK(shapes_before.size() == 1);
CHECK(shapes_before.size() == shapes_after.size()); CHECK(shapes_before.size() == shapes_after.size());
...@@ -61,7 +61,7 @@ TEST_CASE(quantize_int8) ...@@ -61,7 +61,7 @@ TEST_CASE(quantize_int8)
{ {
auto p1 = migraphx::parse_onnx("gemm_ex_test.onnx"); auto p1 = migraphx::parse_onnx("gemm_ex_test.onnx");
const auto& p2 = p1; const auto& p2 = p1;
auto t = migraphx::target("cpu"); auto t = migraphx::target("ref");
migraphx::quantize_int8_options options; migraphx::quantize_int8_options options;
migraphx::quantize_int8(p1, t, options); migraphx::quantize_int8(p1, t, options);
...@@ -84,7 +84,7 @@ TEST_CASE(load_and_run_user_input_shape) ...@@ -84,7 +84,7 @@ TEST_CASE(load_and_run_user_input_shape)
options.set_input_parameter_shape("0", {2, 3, 64, 64}); options.set_input_parameter_shape("0", {2, 3, 64, 64});
auto p = migraphx::parse_onnx("conv_relu_maxpool_test.onnx", options); auto p = migraphx::parse_onnx("conv_relu_maxpool_test.onnx", options);
auto shapes_before = p.get_output_shapes(); auto shapes_before = p.get_output_shapes();
p.compile(migraphx::target("cpu")); p.compile(migraphx::target("ref"));
auto shapes_after = p.get_output_shapes(); auto shapes_after = p.get_output_shapes();
CHECK(shapes_before.size() == 1); CHECK(shapes_before.size() == 1);
CHECK(shapes_before.size() == shapes_after.size()); CHECK(shapes_before.size() == shapes_after.size());
...@@ -104,7 +104,7 @@ TEST_CASE(zero_parameter) ...@@ -104,7 +104,7 @@ TEST_CASE(zero_parameter)
{ {
auto p = migraphx::parse_onnx("constant_fill_test.onnx"); auto p = migraphx::parse_onnx("constant_fill_test.onnx");
auto shapes_before = p.get_output_shapes(); auto shapes_before = p.get_output_shapes();
p.compile(migraphx::target("cpu")); p.compile(migraphx::target("ref"));
auto shapes_after = p.get_output_shapes(); auto shapes_after = p.get_output_shapes();
CHECK(shapes_before.size() == 1); CHECK(shapes_before.size() == 1);
CHECK(shapes_before.size() == shapes_after.size()); CHECK(shapes_before.size() == shapes_after.size());
......
test/context_test.cpp
View file @ 500d9441
#include <migraphx/serialize.hpp> #include <migraphx/serialize.hpp>
#include <migraphx/context.hpp> #include <migraphx/context.hpp>
#include <migraphx/cpu/context.hpp> #include <migraphx/ref/context.hpp>
#include <migraphx/functional.hpp> #include <migraphx/functional.hpp>
#include <test.hpp> #include <test.hpp>
TEST_CASE(context) TEST_CASE(context)
{ {
migraphx::context ctx = migraphx::cpu::context{}; migraphx::context ctx = migraphx::ref::context{};
migraphx::value v = ctx.to_value(); migraphx::value v = ctx.to_value();
EXPECT(v.empty()); EXPECT(v.empty());
migraphx::context cpu_ctx = migraphx::cpu::context{}; migraphx::context cpu_ctx = migraphx::ref::context{};
cpu_ctx.from_value(v); cpu_ctx.from_value(v);
} }
......
test/gpu/fast_math.cpp
View file @ 500d9441
#include <test.hpp> #include <test.hpp>
#include <migraphx/quantization.hpp> #include <migraphx/quantization.hpp>
#include <migraphx/iterator_for.hpp> #include <migraphx/iterator_for.hpp>
#include "test_utils.hpp" #include <migraphx/op/add.hpp>
#include "test.hpp" #include <migraphx/op/mul.hpp>
#include <migraphx/op/multibroadcast.hpp>
#include <migraphx/op/pow.hpp>
#include <migraphx/op/tanh.hpp>
#include <migraphx/gpu/target.hpp>
#include <migraphx/instruction.hpp>
migraphx::program create_gelu() migraphx::program create_gelu()
{ {
......
test/gpu/ops_test.cpp deleted 100644 → 0
View file @ ba33d25c
#include <test.hpp>
#include <migraphx/quantization.hpp>
#include "test_utils.hpp"
#ifdef __clang__
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wglobal-constructors"
#endif
struct test_literals : verify_program<test_literals>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input = p.add_literal(
generate_literal(migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}}));
auto weights = p.add_literal(
generate_literal(migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}}));
auto conv = p.add_instruction(migraphx::op::convolution{}, input, weights);
p.add_instruction(migraphx::op::relu{}, conv);
return p;
}
};
struct test_add : verify_program<test_add>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {3}};
auto x = p.add_parameter("x", s);
auto y = p.add_parameter("y", s);
p.add_instruction(migraphx::op::add{}, x, y);
return p;
}
};
struct test_add_half : verify_program<test_add_half>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::half_type, {3}};
auto x = p.add_parameter("x", s);
auto y = p.add_parameter("y", s);
p.add_instruction(migraphx::op::add{}, x, y);
return p;
}
};
struct test_mul : verify_program<test_mul>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {3}};
auto x = p.add_parameter("x", s);
auto y = p.add_parameter("y", s);
p.add_instruction(migraphx::op::mul{}, x, y);
return p;
}
};
struct test_exp : verify_program<test_exp>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {6}};
auto x = p.add_instruction(migraphx::op::abs{}, p.add_parameter("x", s));
p.add_instruction(migraphx::op::exp{}, x);
return p;
}
};
struct test_erf : verify_program<test_erf>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {2, 3, 4, 6}};
auto param = p.add_parameter("x", s);
p.add_instruction(migraphx::op::erf{}, param);
return p;
}
};
struct test_sqrt : verify_program<test_sqrt>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {2, 3, 4, 6}};
auto param = p.add_parameter("x", s);
auto param_abs = p.add_instruction(migraphx::op::abs{}, param);
p.add_instruction(migraphx::op::sqrt{}, param_abs);
return p;
}
};
struct test_sign : verify_program<test_sign>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::double_type, {2, 3, 4, 6}};
auto param = p.add_parameter("x", s);
p.add_instruction(migraphx::op::sign{}, param);
return p;
}
};
struct test_log : verify_program<test_log>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {6}};
auto x = p.add_instruction(migraphx::op::abs{}, p.add_parameter("x", s));
p.add_instruction(migraphx::op::log{}, x);
return p;
}
};
struct test_pow : verify_program<test_pow>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {6}};
std::vector<float> vec_e(s.elements(), 2.0f);
auto b = p.add_parameter("x", s);
auto e = p.add_literal(migraphx::literal(s, vec_e));
p.add_instruction(migraphx::op::pow{}, b, e);
return p;
}
};
struct test_prelu_brcst : verify_program<test_prelu_brcst>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {6}};
auto x = p.add_parameter("x", s);
auto slp = p.add_parameter("slp", s);
auto r = p.add_instruction(migraphx::op::prelu{}, x, slp);
p.add_return({r});
return p;
}
};
struct test_sin : verify_program<test_sin>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {10}};
auto x = p.add_parameter("x", s);
p.add_instruction(migraphx::op::sin{}, x);
return p;
}
};
struct test_cos : verify_program<test_cos>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::double_type, {8}};
auto x = p.add_parameter("x", s);
p.add_instruction(migraphx::op::cos{}, x);
return p;
}
};
struct test_tan : verify_program<test_tan>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {16}};
auto x = p.add_parameter("x", s);
p.add_instruction(migraphx::op::tan{}, x);
return p;
}
};
struct test_sinh : verify_program<test_sinh>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::double_type, {16}};
auto x = p.add_parameter("x", s);
p.add_instruction(migraphx::op::sinh{}, x);
return p;
}
};
struct test_cosh : verify_program<test_cosh>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::double_type, {16}};
auto x = p.add_parameter("x", s);
p.add_instruction(migraphx::op::cosh{}, x);
return p;
}
};
struct test_tanh : verify_program<test_tanh>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
p.add_instruction(migraphx::op::tanh{}, x);
return p;
}
};
struct test_trans_tanh : verify_program<test_trans_tanh>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto tx = p.add_instruction(migraphx::op::transpose{{0, 1, 3, 2}}, x);
auto tanhx = p.add_instruction(migraphx::op::tanh{}, tx);
auto r = p.add_instruction(migraphx::op::add{}, tanhx, tanhx);
p.add_instruction(migraphx::op::contiguous{}, r);
return p;
}
};
struct test_trans_tanh1 : verify_program<test_trans_tanh1>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto tx = p.add_instruction(migraphx::op::transpose{{0, 1, 3, 2}}, x);
auto tanhx = p.add_instruction(migraphx::op::tanh{}, tx);
auto r = p.add_instruction(migraphx::op::add{}, tanhx, tanhx);
p.add_return({tx, r});
return p;
}
};
struct test_slice_sin : verify_program<test_slice_sin>
{
migraphx::program create_program() const
{
migraphx::program p;
auto l = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {2, 2}});
auto t = p.add_instruction(migraphx::op::slice{{1}, {1}, {2}}, l);
p.add_instruction(migraphx::op::sin{}, t);
return p;
}
};
struct test_asin : verify_program<test_asin>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::double_type, {16}};
auto x = p.add_parameter("x", s);
p.add_instruction(migraphx::op::asin{}, x);
return p;
}
};
struct test_acos : verify_program<test_acos>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::double_type, {16}};
auto x = p.add_parameter("x", s);
p.add_instruction(migraphx::op::acos{}, x);
return p;
}
};
struct test_atan : verify_program<test_atan>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::double_type, {16}};
auto x = p.add_parameter("x", s);
p.add_instruction(migraphx::op::atan{}, x);
return p;
}
};
struct test_asinh : verify_program<test_asinh>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::double_type, {16}};
auto x = p.add_parameter("x", s);
p.add_instruction(migraphx::op::asinh{}, x);
return p;
}
};
struct test_acosh : verify_program<test_acosh>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {16}};
auto x = p.add_parameter("x", s);
auto min_val = p.add_literal(1.1f);
auto max_val = p.add_literal(100.0f);
min_val = p.add_instruction(migraphx::op::multibroadcast{{16}}, min_val);
max_val = p.add_instruction(migraphx::op::multibroadcast{{16}}, max_val);
auto cx = p.add_instruction(migraphx::op::clip{}, x, min_val, max_val);
p.add_instruction(migraphx::op::acosh{}, cx);
return p;
}
};
struct test_atanh : verify_program<test_atanh>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::double_type, {16}};
auto x = p.add_parameter("x", s);
auto min_val = p.add_literal(-0.95);
auto max_val = p.add_literal(0.95);
min_val = p.add_instruction(migraphx::op::multibroadcast{{16}}, min_val);
max_val = p.add_instruction(migraphx::op::multibroadcast{{16}}, max_val);
auto cx = p.add_instruction(migraphx::op::clip{}, x, min_val, max_val);
p.add_instruction(migraphx::op::atanh{}, cx);
return p;
}
};
struct test_scale : verify_program<test_scale>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {3}};
auto x = p.add_parameter("x", s);
auto y = p.add_parameter("y", migraphx::shape::float_type);
auto scale = p.add_instruction(migraphx::op::scalar{s.lens()}, y);
p.add_instruction(migraphx::op::mul{}, x, scale);
return p;
}
};
struct test_slice : verify_program<test_slice>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::int32_type, {2, 2, 4}};
auto x = p.add_parameter("x", s);
auto y = p.add_parameter("y", {migraphx::shape::int32_type, {2, 2, 2}});
auto slice0 = p.add_instruction(migraphx::op::slice{{2}, {0}, {2}}, x);
p.add_instruction(migraphx::op::add{}, y, slice0);
return p;
}
};
struct test_triadd : verify_program<test_triadd>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {3}};
auto x = p.add_parameter("x", s);
auto y = p.add_parameter("y", s);
auto z = p.add_parameter("z", s);
auto sum = p.add_instruction(migraphx::op::add{}, x, y);
p.add_instruction(migraphx::op::add{}, sum, z);
return p;
}
};
struct test_triadd2 : verify_program<test_triadd2>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {2, 3}};
migraphx::shape b{migraphx::shape::float_type, {3}};
auto x = p.add_parameter("x", s);
auto y = p.add_parameter("y", s);
auto z = p.add_parameter("z", b);
auto zb = p.add_instruction(migraphx::op::broadcast{1, s.lens()}, z);
auto sum = p.add_instruction(migraphx::op::add{}, x, y);
p.add_instruction(migraphx::op::add{}, sum, zb);
return p;
}
};
struct test_mul_add : verify_program<test_mul_add>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {2, 3}};
migraphx::shape bs{migraphx::shape::float_type, {3}};
auto x = p.add_parameter("x", s);
auto a = p.add_parameter("a", bs);
auto b = p.add_parameter("b", bs);
auto ab = p.add_instruction(migraphx::op::broadcast{1, s.lens()}, a);
auto bb = p.add_instruction(migraphx::op::broadcast{1, s.lens()}, b);
auto mul = p.add_instruction(migraphx::op::mul{}, x, ab);
p.add_instruction(migraphx::op::add{}, mul, bb);
return p;
}
};
struct test_add_broadcast : verify_program<test_add_broadcast>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {3}};
auto x = p.add_parameter("x", {migraphx::shape::float_type, {2, 2, 3}});
auto y = p.add_parameter("y", {migraphx::shape::float_type, {2, 2}});
auto by = p.add_instruction(migraphx::op::broadcast{0, x->get_shape().lens()}, y);
p.add_instruction(migraphx::op::add{}, x, by);
return p;
}
};
struct test_add_broadcast2 : verify_program<test_add_broadcast2>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {3}};
auto x = p.add_parameter("x", {migraphx::shape::float_type, {2, 3, 4}});
auto y = p.add_parameter("y", {migraphx::shape::float_type, {3}});
auto by = p.add_instruction(migraphx::op::broadcast{1, x->get_shape().lens()}, y);
p.add_instruction(migraphx::op::add{}, x, by);
return p;
}
};
struct test_add_broadcast3 : verify_program<test_add_broadcast3>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {3}};
auto x = p.add_parameter("x", {migraphx::shape::float_type, {2, 4, 5}});
auto y = p.add_parameter("y", {migraphx::shape::float_type, {4}});
auto by = p.add_instruction(migraphx::op::broadcast{1, x->get_shape().lens()}, y);
p.add_instruction(migraphx::op::add{}, x, by);
return p;
}
};
struct test_add_broadcast4 : verify_program<test_add_broadcast4>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {3}};
auto x = p.add_parameter("x", {migraphx::shape::float_type, {2, 3, 5}});
auto y = p.add_parameter("y", {migraphx::shape::float_type, {3}});
auto by = p.add_instruction(migraphx::op::broadcast{1, x->get_shape().lens()}, y);
p.add_instruction(migraphx::op::add{}, x, by);
return p;
}
};
struct test_add_broadcast5 : verify_program<test_add_broadcast5>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {3}};
auto x = p.add_parameter("x", {migraphx::shape::float_type, {2, 4, 8}});
auto y = p.add_parameter("y", {migraphx::shape::float_type, {4}});
auto by = p.add_instruction(migraphx::op::broadcast{1, x->get_shape().lens()}, y);
p.add_instruction(migraphx::op::add{}, x, by);
return p;
}
};
struct test_triadd_broadcast : verify_program<test_triadd_broadcast>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {3}};
auto x = p.add_parameter("x", {migraphx::shape::float_type, {2, 2, 3}});
auto y = p.add_parameter("y", {migraphx::shape::float_type, {2, 2}});
auto z = p.add_parameter("z", {migraphx::shape::float_type, {2, 2, 3}});
auto by = p.add_instruction(migraphx::op::broadcast{0, x->get_shape().lens()}, y);
auto sum = p.add_instruction(migraphx::op::add{}, x, by);
p.add_instruction(migraphx::op::add{}, sum, z);
return p;
}
};
struct test_gelu : verify_program<test_gelu>
{
migraphx::program create_program() const
{
migraphx::program p;
std::vector<size_t> input_lens{1, 1, 5};
auto x = p.add_parameter("x", {migraphx::shape::float_type, input_lens});
auto half = p.add_literal(0.5f);
auto one = p.add_literal(1.0f);
auto sqrt2 = p.add_literal(static_cast<float>(M_SQRT2));
auto half_mbcast = p.add_instruction(migraphx::op::multibroadcast{input_lens}, half);
auto mul_half = p.add_instruction(migraphx::op::mul{}, x, half_mbcast);
auto sqrt2_mbcast = p.add_instruction(migraphx::op::multibroadcast{input_lens}, sqrt2);
auto div = p.add_instruction(migraphx::op::div{}, x, sqrt2_mbcast);
auto erf = p.add_instruction(migraphx::op::erf{}, div);
auto one_mbcast = p.add_instruction(migraphx::op::multibroadcast{input_lens}, one);
auto add_one = p.add_instruction(migraphx::op::add{}, erf, one_mbcast);
p.add_instruction(migraphx::op::mul{}, mul_half, add_one);
return p;
}
};
struct test_add_gelu : verify_program<test_add_gelu>
{
migraphx::program create_program() const
{
migraphx::program p;
std::vector<size_t> input_lens{1, 1, 5};
auto x = p.add_parameter("x", {migraphx::shape::float_type, input_lens});
auto y = p.add_parameter("y", {migraphx::shape::float_type, input_lens});
auto half = p.add_literal(0.5f);
auto one = p.add_literal(1.0f);
auto sqrt2 = p.add_literal(static_cast<float>(M_SQRT2));
auto add = p.add_instruction(migraphx::op::add{}, x, y);
auto half_mbcast = p.add_instruction(migraphx::op::multibroadcast{input_lens}, half);
auto mul_half = p.add_instruction(migraphx::op::mul{}, add, half_mbcast);
auto sqrt2_mbcast = p.add_instruction(migraphx::op::multibroadcast{input_lens}, sqrt2);
auto div = p.add_instruction(migraphx::op::div{}, add, sqrt2_mbcast);
auto erf = p.add_instruction(migraphx::op::erf{}, div);
auto one_mbcast = p.add_instruction(migraphx::op::multibroadcast{input_lens}, one);
auto add_one = p.add_instruction(migraphx::op::add{}, erf, one_mbcast);
p.add_instruction(migraphx::op::mul{}, mul_half, add_one);
return p;
}
};
struct test_sub : verify_program<test_sub>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {3}};
auto x = p.add_parameter("x", s);
auto y = p.add_parameter("y", s);
auto z = p.add_parameter("z", s);
auto diff = p.add_instruction(migraphx::op::sub{}, x, y);
p.add_instruction(migraphx::op::sub{}, diff, z);
return p;
}
};
struct test_sub2 : verify_program<test_sub2>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {2, 3}};
migraphx::shape b{migraphx::shape::float_type, {3}};
auto x = p.add_parameter("x", s);
auto y = p.add_parameter("y", s);
auto z = p.add_parameter("z", b);
auto zb = p.add_instruction(migraphx::op::broadcast{1, s.lens()}, z);
auto diff = p.add_instruction(migraphx::op::sub{}, x, y);
p.add_instruction(migraphx::op::sub{}, diff, zb);
return p;
}
};
struct test_div : verify_program<test_div>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {3}};
auto x = p.add_parameter("x", s);
auto y = p.add_parameter("y", s);
auto z = p.add_parameter("z", s);
auto diff = p.add_instruction(migraphx::op::div{}, x, y);
p.add_instruction(migraphx::op::div{}, diff, z);
return p;
}
};
struct test_div2 : verify_program<test_div2>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {2, 3}};
migraphx::shape b{migraphx::shape::float_type, {3}};
auto x = p.add_parameter("x", s);
auto y = p.add_parameter("y", s);
auto z = p.add_parameter("z", b);
auto zb = p.add_instruction(migraphx::op::broadcast{1, s.lens()}, z);
auto diff = p.add_instruction(migraphx::op::div{}, x, y);
p.add_instruction(migraphx::op::div{}, diff, zb);
return p;
}
};
struct test_softmax1 : verify_program<test_softmax1>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {5, 3, 3, 4}});
p.add_instruction(migraphx::op::softmax{0}, x);
return p;
}
};
struct test_softmax2 : verify_program<test_softmax2>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x =
p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 1000, 1, 1}});
p.add_instruction(migraphx::op::softmax{}, x);
return p;
}
};
template <int Axis, migraphx::shape::type_t T>
struct test_softmax : verify_program<test_softmax<Axis, T>>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{T, {512, 4, 1067, 6}};
auto param = p.add_parameter("0", s);
p.add_instruction(migraphx::op::softmax{Axis}, param);
return p;
}
};
template struct test_softmax<0, migraphx::shape::float_type>;
template struct test_softmax<2, migraphx::shape::float_type>;
template struct test_softmax<1, migraphx::shape::double_type>;
template struct test_softmax<3, migraphx::shape::double_type>;
template struct test_softmax<0, migraphx::shape::half_type>;
template struct test_softmax<1, migraphx::shape::half_type>;
template struct test_softmax<2, migraphx::shape::half_type>;
template struct test_softmax<3, migraphx::shape::half_type>;
template <class T, int Axis>
struct test_arg_ops : verify_program<test_arg_ops<T, Axis>>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {2, 3, 4, 1025}};
auto param = p.add_parameter("data", s);
p.add_instruction(T{Axis}, param);
return p;
}
};
template struct test_arg_ops<migraphx::op::argmax, 0>;
template struct test_arg_ops<migraphx::op::argmax, 1>;
template struct test_arg_ops<migraphx::op::argmax, 2>;
template struct test_arg_ops<migraphx::op::argmax, 3>;
template struct test_arg_ops<migraphx::op::argmax, -1>;
template struct test_arg_ops<migraphx::op::argmax, -2>;
template struct test_arg_ops<migraphx::op::argmin, 0>;
template struct test_arg_ops<migraphx::op::argmin, 1>;
template struct test_arg_ops<migraphx::op::argmin, 2>;
template struct test_arg_ops<migraphx::op::argmin, 3>;
template struct test_arg_ops<migraphx::op::argmin, -3>;
template struct test_arg_ops<migraphx::op::argmin, -4>;
struct test_conv : verify_program<test_conv>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input =
p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto weights =
p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
p.add_instruction(migraphx::op::convolution{}, input, weights);
return p;
}
};
struct test_conv2 : verify_program<test_conv2>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input =
p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 512, 28, 28}});
auto weights =
p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {256, 512, 1, 1}});
p.add_instruction(migraphx::op::convolution{{0, 0}, {1, 1}, {1, 1}}, input, weights);
return p;
}
};
struct test_conv3d : verify_program<test_conv3d>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input =
p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3, 3}});
auto weights =
p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3, 3}});
p.add_instruction(
migraphx::op::convolution{{0, 0, 0}, {1, 1, 1}, {1, 1, 1}}, input, weights);
return p;
}
};
struct test_group_conv : verify_program<test_group_conv>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input =
p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 4, 16, 16}});
auto weights =
p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {4, 1, 3, 3}});
migraphx::op::convolution op;
op.group = 4;
p.add_instruction(op, input, weights);
return p;
}
};
struct test_deconv : verify_program<test_deconv>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input =
p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 1, 3, 3}});
auto weights =
p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {1, 1, 3, 3}});
p.add_instruction(migraphx::op::deconvolution{}, input, weights);
return p;
}
};
struct test_deconv_2x3 : verify_program<test_deconv_2x3>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input =
p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 6, 7}});
auto weights =
p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {3, 4, 3, 3}});
p.add_instruction(migraphx::op::deconvolution{{1, 1}, {2, 3}, {1, 1}}, input, weights);
return p;
}
};
struct test_deconv_1d : verify_program<test_deconv_1d>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 1, 3}});
auto weights =
p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {1, 1, 3}});
p.add_instruction(migraphx::op::deconvolution{{0}, {1}, {1}}, input, weights);
return p;
}
};
struct test_deconv_3d : verify_program<test_deconv_3d>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input =
p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 1, 3, 3, 3}});
auto weights =
p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {1, 1, 3, 3, 3}});
p.add_instruction(
migraphx::op::deconvolution{{0, 0, 0}, {1, 1, 1}, {1, 1, 1}}, input, weights);
return p;
}
};
struct test_conv_relu : verify_program<test_conv_relu>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input =
p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto weights =
p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto conv = p.add_instruction(migraphx::op::convolution{}, input, weights);
p.add_instruction(migraphx::op::relu{}, conv);
return p;
}
};
struct test_conv_relu_half : verify_program<test_conv_relu_half>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input =
p.add_parameter("x", migraphx::shape{migraphx::shape::half_type, {4, 3, 3, 3}});
auto weights =
p.add_parameter("w", migraphx::shape{migraphx::shape::half_type, {4, 3, 3, 3}});
auto conv = p.add_instruction(migraphx::op::convolution{}, input, weights);
p.add_instruction(migraphx::op::relu{}, conv);
return p;
}
};
struct test_conv_bias_clipped_relu : verify_program<test_conv_bias_clipped_relu>
{
migraphx::program create_program() const
{
migraphx::program p;
std::vector<size_t> input_lens{4, 3, 3, 3};
auto input =
p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto weights =
p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto l0 = migraphx::literal{migraphx::shape{migraphx::shape::float_type, {4}},
{2.0f, 2.0f, 2.0f, 2.0f}};
auto bias = p.add_literal(l0);
auto conv = p.add_instruction(migraphx::op::convolution{}, input, weights);
auto bcast_add =
p.add_instruction(migraphx::op::broadcast{1, conv->get_shape().lens()}, bias);
auto bias_add = p.add_instruction(migraphx::op::add{}, conv, bcast_add);
auto min_val = p.add_literal(0.0f);
auto max_val = p.add_literal(6.0f);
min_val = p.add_instruction(migraphx::op::multibroadcast{input_lens}, min_val);
max_val = p.add_instruction(migraphx::op::multibroadcast{input_lens}, max_val);
p.add_instruction(migraphx::op::clip{}, bias_add, min_val, max_val);
return p;
}
};
struct test_conv_add : verify_program<test_conv_add>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", {migraphx::shape::float_type, {1, 8, 4, 4}});
auto w = p.add_literal(
migraphx::generate_literal({migraphx::shape::float_type, {2, 8, 3, 3}}, 1));
auto y = p.add_parameter("y", {migraphx::shape::float_type, {1, 8, 4, 4}});
auto v = p.add_literal(
migraphx::generate_literal({migraphx::shape::float_type, {2, 8, 3, 3}}, 2));
auto conv1 = p.add_instruction(migraphx::op::convolution{}, x, w);
auto conv2 = p.add_instruction(migraphx::op::convolution{}, y, v);
auto sum = p.add_instruction(migraphx::op::add{}, conv1, conv2);
p.add_instruction(migraphx::op::exp{}, sum);
return p;
}
};
struct test_conv_add_1x1_diff_strides : verify_program<test_conv_add_1x1_diff_strides>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", {migraphx::shape::float_type, {1, 8, 2, 2}});
auto w = p.add_literal(
migraphx::generate_literal({migraphx::shape::float_type, {2, 8, 1, 1}}, 1));
auto y = p.add_parameter("y", {migraphx::shape::float_type, {1, 8, 4, 4}});
auto v = p.add_literal(
migraphx::generate_literal({migraphx::shape::float_type, {2, 8, 1, 1}}, 2));
auto conv1 = p.add_instruction(migraphx::op::convolution{}, x, w);
auto conv2 = p.add_instruction(migraphx::op::convolution{{0, 0}, {2, 2}}, y, v);
auto sum = p.add_instruction(migraphx::op::add{}, conv1, conv2);
p.add_instruction(migraphx::op::exp{}, sum);
return p;
}
};
// struct test_conv_bn_add : verify_program<test_conv_bn_add>
// {
// static migraphx::instruction_ref add_bn(migraphx::program& p,
// migraphx::instruction_ref x,
// std::size_t channels,
// std::size_t seed = 1)
// {
// migraphx::shape vars{migraphx::shape::float_type, {channels}};
// auto scale = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 1 + seed)));
// auto bias = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 2 + seed)));
// auto mean = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 3 + seed)));
// auto variance = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 4 + seed)));
// return p.add_instruction(
// migraphx::op::batch_norm_inference{}, x, scale, bias, mean, variance);
// }
// migraphx::program create_program() const
// {
// migraphx::program p;
// std::size_t ichannels = 64;
// std::size_t ochannels = 256;
// auto x = p.add_parameter("x", {migraphx::shape::float_type, {1, ichannels, 56, 56}});
// auto w = p.add_literal(migraphx::generate_literal(
// {migraphx::shape::float_type, {ochannels, ichannels, 1, 1}}, 1));
// auto y = p.add_parameter("y", {migraphx::shape::float_type, {1, ichannels, 56, 56}});
// auto v = p.add_literal(migraphx::generate_literal(
// {migraphx::shape::float_type, {ochannels, ichannels, 1, 1}}, 2));
// auto relu1 = p.add_instruction(migraphx::op::relu{}, x);
// auto conv1 = p.add_instruction(migraphx::op::convolution{}, relu1, w);
// auto bn1 = add_bn(p, conv1, ochannels, 1);
// auto relu2 = p.add_instruction(migraphx::op::relu{}, y);
// auto conv2 = p.add_instruction(migraphx::op::convolution{}, relu2, v);
// auto bn2 = add_bn(p, conv2, ochannels, 1);
// auto sum = p.add_instruction(migraphx::op::add{}, bn1, bn2);
// p.add_instruction(migraphx::op::relu{}, sum);
// return p;
// }
// };
struct test_add_relu : verify_program<test_add_relu>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto y = p.add_parameter("y", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto add = p.add_instruction(migraphx::op::add{}, x, y);
p.add_instruction(migraphx::op::relu{}, add);
return p;
}
};
struct test_add_sigmoid : verify_program<test_add_sigmoid>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto y = p.add_parameter("y", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto add = p.add_instruction(migraphx::op::add{}, x, y);
p.add_instruction(migraphx::op::sigmoid{}, add);
return p;
}
};
struct test_add_tanh : verify_program<test_add_tanh>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto y = p.add_parameter("y", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto add = p.add_instruction(migraphx::op::add{}, x, y);
p.add_instruction(migraphx::op::tanh{}, add);
return p;
}
};
struct test_triadd_relu : verify_program<test_triadd_relu>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto y = p.add_parameter("y", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto z = p.add_parameter("z", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto sum = p.add_instruction(migraphx::op::add{}, x, y);
auto triadd = p.add_instruction(migraphx::op::add{}, sum, z);
p.add_instruction(migraphx::op::relu{}, triadd);
return p;
}
};
struct test_triadd_sigmoid : verify_program<test_triadd_sigmoid>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto y = p.add_parameter("y", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto z = p.add_parameter("z", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto sum = p.add_instruction(migraphx::op::add{}, x, y);
auto triadd = p.add_instruction(migraphx::op::add{}, sum, z);
p.add_instruction(migraphx::op::sigmoid{}, triadd);
return p;
}
};
struct test_triadd_tanh : verify_program<test_triadd_tanh>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto y = p.add_parameter("y", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto z = p.add_parameter("z", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto sum = p.add_instruction(migraphx::op::add{}, x, y);
auto triadd = p.add_instruction(migraphx::op::add{}, sum, z);
p.add_instruction(migraphx::op::tanh{}, triadd);
return p;
}
};
migraphx::instruction_ref add_layernorm(migraphx::program& p, std::vector<size_t> dims)
{
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, dims});
auto scale =
p.add_parameter("scale", migraphx::shape{migraphx::shape::float_type, {dims.back()}});
auto bias =
p.add_parameter("bias", migraphx::shape{migraphx::shape::float_type, {dims.back()}});
auto epsilon = p.add_literal(1e-12f);
auto exponent = p.add_literal(2.0f);
auto mean = p.add_instruction(migraphx::op::reduce_mean({2}), x);
auto mean_mbcast = p.add_instruction(migraphx::op::multibroadcast{{dims}}, mean);
auto sub = p.add_instruction(migraphx::op::sub{}, x, mean_mbcast);
auto exponent_mbcast = p.add_instruction(migraphx::op::multibroadcast{{dims}}, exponent);
auto pow = p.add_instruction(migraphx::op::pow{}, sub, exponent_mbcast);
auto var = p.add_instruction(migraphx::op::reduce_mean({2}), pow);
auto epsilon_mbcast =
p.add_instruction(migraphx::op::multibroadcast{{1, dims.at(1), 1}}, epsilon);
auto add_epsilon = p.add_instruction(migraphx::op::add{}, var, epsilon_mbcast);
auto sqrt = p.add_instruction(migraphx::op::sqrt{}, add_epsilon);
auto sqrt_mbcast = p.add_instruction(migraphx::op::multibroadcast{dims}, sqrt);
auto div = p.add_instruction(migraphx::op::div{}, sub, sqrt_mbcast);
auto scale_mbcast = p.add_instruction(migraphx::op::multibroadcast{dims}, scale);
auto mul = p.add_instruction(migraphx::op::mul{}, scale_mbcast, div);
auto bias_mbcast = p.add_instruction(migraphx::op::multibroadcast{dims}, bias);
return p.add_instruction(migraphx::op::add{}, mul, bias_mbcast);
}
struct test_layernorm : verify_program<test_layernorm>
{
migraphx::program create_program() const
{
migraphx::program p;
add_layernorm(p, {1, 1, 5});
return p;
}
};
struct test_layernorm2 : verify_program<test_layernorm2>
{
migraphx::program create_program() const
{
migraphx::program p;
add_layernorm(p, {1, 4, 24});
return p;
}
};
struct test_sigmoid : verify_program<test_sigmoid>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
p.add_instruction(migraphx::op::sigmoid{}, x);
return p;
}
};
struct test_abs : verify_program<test_abs>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
p.add_instruction(migraphx::op::abs{}, x);
return p;
}
};
struct test_trans_abs : verify_program<test_trans_abs>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto tx = p.add_instruction(migraphx::op::transpose{{0, 1, 3, 2}}, x);
auto absx = p.add_instruction(migraphx::op::abs{}, tx);
auto r = p.add_instruction(migraphx::op::add{}, absx, absx);
p.add_instruction(migraphx::op::contiguous{}, r);
return p;
}
};
struct test_leaky_relu : verify_program<test_leaky_relu>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
p.add_instruction(migraphx::op::leaky_relu{0.01}, x);
return p;
}
};
struct test_elu : verify_program<test_elu>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
p.add_instruction(migraphx::op::leaky_relu{1.0}, x);
return p;
}
};
struct test_relu_lrn : verify_program<test_relu_lrn>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 5, 2, 2}});
auto y = p.add_instruction(migraphx::op::relu{}, x);
p.add_instruction(migraphx::op::lrn{0.0001, 0.75, 1.0, 5}, y);
return p;
}
};
struct test_conv_pooling : verify_program<test_conv_pooling>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input =
p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 32, 32}});
auto weights =
p.add_parameter("w", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto conv = p.add_instruction(migraphx::op::convolution{}, input, weights);
auto pooling = p.add_instruction(migraphx::op::pooling{"max"}, conv);
p.add_instruction(migraphx::op::relu{}, pooling);
return p;
}
};
struct test_concat_pooling : verify_program<test_concat_pooling>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input =
p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 256, 8, 8}});
auto transpose = p.add_instruction(migraphx::op::transpose{{0, 2, 3, 1}}, input);
auto concat = p.add_instruction(migraphx::op::concat{3}, transpose);
auto concat_t = p.add_instruction(migraphx::op::transpose{{0, 3, 1, 2}}, concat);
auto pooling =
p.add_instruction(migraphx::op::pooling{"average", {0, 0}, {1, 1}, {8, 8}}, concat_t);
p.add_instruction(migraphx::op::relu{}, pooling);
return p;
}
};
struct test_global_avg_pooling : verify_program<test_global_avg_pooling>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input =
p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}});
auto op = migraphx::op::pooling{"average"};
auto lens = input->get_shape().lens();
op.lengths = {lens[2], lens[3]};
p.add_instruction(op, input);
return p;
}
};
struct test_global_max_pooling : verify_program<test_global_max_pooling>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input =
p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 16, 16}});
auto op = migraphx::op::pooling{"max"};
auto lens = input->get_shape().lens();
op.lengths = {lens[2], lens[3]};
p.add_instruction(op, input);
return p;
}
};
struct test_max_pooling_ceil_3d : verify_program<test_max_pooling_ceil_3d>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input =
p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 5, 5, 5}});
auto op = migraphx::op::pooling{"max", {1, 1, 1}, {3, 3, 3}, {3, 3, 3}, true};
p.add_instruction(op, input);
return p;
}
};
struct test_avg_pooling_1d : verify_program<test_avg_pooling_1d>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 5}});
auto op = migraphx::op::pooling{"average", {0}, {1}, {3}};
p.add_instruction(op, input);
return p;
}
};
struct test_avg_pooling_3d : verify_program<test_avg_pooling_3d>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input =
p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 5, 5, 5}});
auto op = migraphx::op::pooling{"average", {1, 1, 1}, {3, 3, 3}, {3, 3, 3}};
p.add_instruction(op, input);
return p;
}
};
struct test_avg_pooling_ceil_3d : verify_program<test_avg_pooling_ceil_3d>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input =
p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {1, 3, 5, 5, 5}});
auto op = migraphx::op::pooling{"average", {1, 1, 1}, {3, 3, 3}, {3, 3, 3}, true};
p.add_instruction(op, input);
return p;
}
};
struct test_avg_pooling_3d_opt : verify_program<test_avg_pooling_3d_opt>
{
migraphx::program create_program() const
{
migraphx::program p;
auto input =
p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 2, 3, 3, 3}});
auto op = migraphx::op::pooling{"average", {0, 0, 0}, {1, 1, 1}, {3, 3, 3}};
p.add_instruction(op, input);
return p;
}
};
struct test_gemm : verify_program<test_gemm>
{
migraphx::program create_program() const
{
migraphx::program p;
auto a = p.add_parameter("a", migraphx::shape{migraphx::shape::float_type, {4, 5}});
auto b = p.add_parameter("b", migraphx::shape{migraphx::shape::float_type, {5, 3}});
p.add_instruction(migraphx::op::dot{}, a, b);
return p;
}
};
struct test_gemm_copy : verify_program<test_gemm_copy>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape sa{migraphx::shape::float_type, {2, 16}};
migraphx::shape sb{migraphx::shape::float_type, {16, 8}};
migraphx::shape sc{migraphx::shape::float_type, {2, 8}};
auto pa = p.add_parameter("a", sa);
auto pb = p.add_parameter("b", sb);
auto pc = p.add_parameter("c", sc);
auto dr = p.add_instruction(migraphx::op::dot{}, pa, pb, pc);
p.add_instruction(migraphx::op::add{}, dr, dr);
return p;
}
};
struct test_gemm_ex : verify_program<test_gemm_ex>
{
migraphx::program create_program() const
{
migraphx::program p;
auto a = p.add_parameter("a", migraphx::shape{migraphx::shape::float_type, {1, 1, 4, 5}});
auto b = p.add_parameter("b", migraphx::shape{migraphx::shape::float_type, {1, 1, 5, 3}});
p.add_instruction(migraphx::op::dot{}, a, b);
return p;
}
};
struct test_gemm_half : verify_program<test_gemm_half>
{
migraphx::program create_program() const
{
migraphx::program p;
auto a = p.add_parameter("a", migraphx::shape{migraphx::shape::half_type, {4, 5}});
auto b = p.add_parameter("b", migraphx::shape{migraphx::shape::half_type, {5, 3}});
p.add_instruction(migraphx::op::dot{}, a, b);
return p;
}
};
struct test_gemm_ld //: verify_program<test_gemm_ld>
{
migraphx::program create_program() const
{
migraphx::program p;
auto a =
p.add_parameter("a", migraphx::shape{migraphx::shape::float_type, {4, 5}, {10, 1}});
auto b =
p.add_parameter("b", migraphx::shape{migraphx::shape::float_type, {5, 3}, {20, 1}});
p.add_instruction(migraphx::op::dot{}, a, b);
return p;
}
};
struct test_gemm_transposeb : verify_program<test_gemm_transposeb>
{
migraphx::program create_program() const
{
migraphx::program p;
auto a = p.add_parameter("a", migraphx::shape{migraphx::shape::float_type, {4, 5}});
auto b = p.add_parameter("b", migraphx::shape{migraphx::shape::float_type, {3, 5}});
auto bt = p.add_instruction(migraphx::op::transpose{{1, 0}}, b);
p.add_instruction(migraphx::op::dot{}, a, bt);
return p;
}
};
struct test_gemm_transposeb_ex : verify_program<test_gemm_transposeb_ex>
{
migraphx::program create_program() const
{
migraphx::program p;
auto a = p.add_parameter("a", migraphx::shape{migraphx::shape::float_type, {1, 4, 5}});
auto b = p.add_parameter("b", migraphx::shape{migraphx::shape::float_type, {1, 3, 5}});
auto bt = p.add_instruction(migraphx::op::transpose{{0, 2, 1}}, b);
p.add_instruction(migraphx::op::dot{}, a, bt);
return p;
}
};
struct test_gemm_transposea : verify_program<test_gemm_transposea>
{
migraphx::program create_program() const
{
migraphx::program p;
auto a = p.add_parameter("a", migraphx::shape{migraphx::shape::float_type, {5, 4}});
auto b = p.add_parameter("b", migraphx::shape{migraphx::shape::float_type, {5, 3}});
auto at = p.add_instruction(migraphx::op::transpose{{1, 0}}, a);
p.add_instruction(migraphx::op::dot{}, at, b);
return p;
}
};
struct test_gemm_transposea_ex : verify_program<test_gemm_transposea_ex>
{
migraphx::program create_program() const
{
migraphx::program p;
auto a = p.add_parameter("a", migraphx::shape{migraphx::shape::float_type, {1, 1, 5, 4}});
auto b = p.add_parameter("b", migraphx::shape{migraphx::shape::float_type, {1, 1, 5, 3}});
auto at = p.add_instruction(migraphx::op::transpose{{0, 1, 3, 2}}, a);
p.add_instruction(migraphx::op::dot{}, at, b);
return p;
}
};
struct test_gemm_transposeab : verify_program<test_gemm_transposeab>
{
migraphx::program create_program() const
{
migraphx::program p;
auto a = p.add_parameter("a", migraphx::shape{migraphx::shape::float_type, {5, 4}});
auto b = p.add_parameter("b", migraphx::shape{migraphx::shape::float_type, {3, 5}});
auto at = p.add_instruction(migraphx::op::transpose{{1, 0}}, a);
auto bt = p.add_instruction(migraphx::op::transpose{{1, 0}}, b);
p.add_instruction(migraphx::op::dot{}, at, bt);
return p;
}
};
struct gemm_multi_dim_2 : verify_program<gemm_multi_dim_2>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {2, 2, 3}};
migraphx::shape m2_shape{migraphx::shape::float_type, {2, 3, 4}};
auto l1 = p.add_parameter("1", m1_shape);
auto l2 = p.add_parameter("2", m2_shape);
p.add_instruction(migraphx::op::dot{}, l1, l2);
return p;
}
};
struct gemm_2args_mm_1 : verify_program<gemm_2args_mm_1>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {2, 2, 3}};
migraphx::shape m2_shape{migraphx::shape::float_type, {1, 3, 4}};
auto l1 = p.add_parameter("1", m1_shape);
auto l2 = p.add_parameter("2", m2_shape);
auto bl2 = p.add_instruction(migraphx::op::multibroadcast{{2, 3, 4}}, l2);
p.add_instruction(migraphx::op::dot{}, l1, bl2);
return p;
}
};
struct gemm_2args_mm_2 : verify_program<gemm_2args_mm_2>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {2, 2, 3}};
migraphx::shape m2_shape{migraphx::shape::float_type, {3, 4}};
auto l1 = p.add_parameter("1", m1_shape);
auto l2 = p.add_parameter("2", m2_shape);
auto bl2 = p.add_instruction(migraphx::op::multibroadcast{{2, 3, 4}}, l2);
p.add_instruction(migraphx::op::dot{}, l1, bl2);
return p;
}
};
struct gemm_2args_mm_3 : verify_program<gemm_2args_mm_3>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {1, 2, 3}};
migraphx::shape m2_shape{migraphx::shape::float_type, {3, 3, 4}};
auto l1 = p.add_parameter("1", m1_shape);
auto bl1 = p.add_instruction(migraphx::op::multibroadcast{{3, 2, 3}}, l1);
auto l2 = p.add_parameter("2", m2_shape);
p.add_instruction(migraphx::op::dot{}, bl1, l2);
return p;
}
};
struct gemm_2args_mm_4 : verify_program<gemm_2args_mm_4>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {2, 3}};
migraphx::shape m2_shape{migraphx::shape::float_type, {3, 3, 4}};
auto l1 = p.add_parameter("1", m1_shape);
auto bl1 = p.add_instruction(migraphx::op::multibroadcast{{3, 2, 3}}, l1);
auto l2 = p.add_parameter("2", m2_shape);
p.add_instruction(migraphx::op::dot{}, bl1, l2);
return p;
}
};
struct gemm_2args_mm_5 : verify_program<gemm_2args_mm_5>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {2, 1, 2, 3}};
migraphx::shape m2_shape{migraphx::shape::float_type, {2, 3, 3, 4}};
auto l1 = p.add_parameter("1", m1_shape);
auto bl1 = p.add_instruction(migraphx::op::multibroadcast{{2, 3, 2, 3}}, l1);
auto l2 = p.add_parameter("2", m2_shape);
p.add_instruction(migraphx::op::dot{}, bl1, l2);
return p;
}
};
struct gemm_2args_mm_6 : verify_program<gemm_2args_mm_6>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {2, 1, 2, 3}};
migraphx::shape m2_shape{migraphx::shape::float_type, {1, 3, 3, 4}};
auto l1 = p.add_parameter("1", m1_shape);
auto bl1 = p.add_instruction(migraphx::op::multibroadcast{{2, 3, 2, 3}}, l1);
auto l2 = p.add_parameter("2", m2_shape);
auto bl2 = p.add_instruction(migraphx::op::multibroadcast{{2, 3, 3, 4}}, l2);
p.add_instruction(migraphx::op::dot{}, bl1, bl2);
return p;
}
};
struct gemm_2args_mm_7 : verify_program<gemm_2args_mm_7>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {2, 3}};
migraphx::shape m2_shape{migraphx::shape::float_type, {2, 3, 3, 4}};
auto l1 = p.add_parameter("1", m1_shape);
auto bl1 = p.add_instruction(migraphx::op::multibroadcast{{2, 3, 2, 3}}, l1);
auto l2 = p.add_parameter("2", m2_shape);
p.add_instruction(migraphx::op::dot{}, bl1, l2);
return p;
}
};
struct gemm_multi_dim_2_3 : verify_program<gemm_multi_dim_2_3>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {2, 3, 2, 3}};
migraphx::shape m2_shape{migraphx::shape::float_type, {2, 3, 3, 2}};
auto l1 = p.add_parameter("1", m1_shape);
auto l2 = p.add_parameter("2", m2_shape);
p.add_instruction(migraphx::op::dot{}, l1, l2);
return p;
}
};
struct gemm_2args_vv : verify_program<gemm_2args_vv>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {8}};
migraphx::shape m2_shape{migraphx::shape::float_type, {8}};
auto l1 = p.add_parameter("1", m1_shape);
auto ul1 = p.add_instruction(migraphx::op::unsqueeze{{0}}, l1);
auto l2 = p.add_parameter("2", m2_shape);
auto ul2 = p.add_instruction(migraphx::op::unsqueeze{{1}}, l2);
float alpha = 0.23f;
auto res = p.add_instruction(migraphx::op::dot{alpha}, ul1, ul2);
auto sres = p.add_instruction(migraphx::op::squeeze{{0}}, res);
p.add_instruction(migraphx::op::squeeze{{0}}, sres);
return p;
}
};
struct gemm_2args_mv : verify_program<gemm_2args_mv>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {3, 5}};
migraphx::shape m2_shape{migraphx::shape::float_type, {5}};
auto l1 = p.add_parameter("1", m1_shape);
auto l2 = p.add_parameter("2", m2_shape);
auto ul2 = p.add_instruction(migraphx::op::unsqueeze{{1}}, l2);
p.add_instruction(migraphx::op::dot{}, l1, ul2);
return p;
}
};
struct gemm_2args_bmv : verify_program<gemm_2args_bmv>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {2, 3, 3, 5}};
migraphx::shape m2_shape{migraphx::shape::float_type, {5}};
auto l1 = p.add_parameter("1", m1_shape);
auto l2 = p.add_parameter("2", m2_shape);
auto ul2 = p.add_instruction(migraphx::op::unsqueeze{{1}}, l2);
auto bul2 = p.add_instruction(migraphx::op::multibroadcast{{2, 3, 5, 1}}, ul2);
p.add_instruction(migraphx::op::dot{}, l1, bul2);
return p;
}
};
struct gemm_2args_vm : verify_program<gemm_2args_vm>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {5}};
migraphx::shape m2_shape{migraphx::shape::float_type, {5, 4}};
auto l1 = p.add_parameter("1", m1_shape);
auto ul1 = p.add_instruction(migraphx::op::unsqueeze{{0}}, l1);
auto l2 = p.add_parameter("2", m2_shape);
auto res = p.add_instruction(migraphx::op::dot{}, ul1, l2);
p.add_instruction(migraphx::op::squeeze{{0}}, res);
return p;
}
};
struct gemm_2args_vbm : verify_program<gemm_2args_vbm>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {5}};
migraphx::shape m2_shape{migraphx::shape::float_type, {2, 2, 5, 4}};
auto l1 = p.add_parameter("1", m1_shape);
auto ul1 = p.add_instruction(migraphx::op::unsqueeze{{0}}, l1);
auto bul1 = p.add_instruction(migraphx::op::multibroadcast{{2, 2, 1, 5}}, ul1);
auto l2 = p.add_parameter("2", m2_shape);
auto res = p.add_instruction(migraphx::op::dot{}, bul1, l2);
p.add_instruction(migraphx::op::squeeze{{2}}, res);
return p;
}
};
struct gemm_multi_3args : verify_program<gemm_multi_3args>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {2, 3, 2, 3}};
migraphx::shape m2_shape{migraphx::shape::float_type, {2, 3, 3, 2}};
migraphx::shape m3_shape{migraphx::shape::float_type, {2, 3, 2, 2}};
auto l1 = p.add_parameter("1", m1_shape);
auto l2 = p.add_parameter("2", m2_shape);
auto l3 = p.add_parameter("3", m3_shape);
float alpha = 0.35;
float beta = 0.41;
p.add_instruction(migraphx::op::dot{alpha, beta}, l1, l2, l3);
return p;
}
};
struct gemm_multi_3args_c25 : verify_program<gemm_multi_3args_c25>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {2, 3}};
migraphx::shape m2_shape{migraphx::shape::float_type, {3, 5}};
migraphx::shape m3_shape{migraphx::shape::float_type, {2, 5}};
auto l1 = p.add_parameter("1", m1_shape);
auto l2 = p.add_parameter("2", m2_shape);
auto l3 = p.add_parameter("3", m3_shape);
float alpha = 0.35;
float beta = 0.41;
p.add_instruction(migraphx::op::dot{alpha, beta}, l1, l2, l3);
return p;
}
};
struct gemm_multi_3args_beta0 : verify_program<gemm_multi_3args_beta0>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {1, 2, 3}};
migraphx::shape m2_shape{migraphx::shape::float_type, {1, 3, 4}};
migraphx::shape m3_shape{migraphx::shape::float_type, {1, 2, 4}};
auto l1 = p.add_parameter("1", m1_shape);
auto l2 = p.add_parameter("2", m2_shape);
auto l3 = p.add_parameter("3", m3_shape);
float alpha = 1.0f;
float beta = 0.0f;
p.add_instruction(migraphx::op::dot{alpha, beta}, l1, l2, l3);
return p;
}
};
struct gemm_multi_3args_alpha0 : verify_program<gemm_multi_3args_alpha0>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {1, 2, 3}};
migraphx::shape m2_shape{migraphx::shape::float_type, {1, 3, 4}};
migraphx::shape m3_shape{migraphx::shape::float_type, {1, 2, 4}};
auto l1 = p.add_parameter("1", m1_shape);
auto l2 = p.add_parameter("2", m2_shape);
auto l3 = p.add_parameter("3", m3_shape);
float alpha = 0.0f;
float beta = 1.0f;
p.add_instruction(migraphx::op::dot{alpha, beta}, l1, l2, l3);
return p;
}
};
struct gemm_multi_transpose : verify_program<gemm_multi_transpose>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::float_type, {2, 2, 3}};
migraphx::shape m2_shape{migraphx::shape::float_type, {3, 2, 4}};
auto l1 = p.add_parameter("1", m1_shape);
auto l2 = p.add_parameter("2", m2_shape);
auto tl2 = p.add_instruction(migraphx::op::transpose{{1, 0, 2}}, l2);
float alpha = 1.0f;
float beta = 1.0f;
p.add_instruction(migraphx::op::dot{alpha, beta}, l1, tl2);
return p;
}
};
struct quant_dot_3args_1 : verify_program<quant_dot_3args_1>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::int8_type, {2, 8}};
migraphx::shape m2_shape{migraphx::shape::int8_type, {8, 7}};
migraphx::shape m3_shape{migraphx::shape::int32_type, {2, 7}};
auto l1 = p.add_parameter("a", m1_shape);
auto l2 = p.add_parameter("b", m2_shape);
auto l3 = p.add_parameter("c", m3_shape);
p.add_instruction(migraphx::op::quant_dot{}, l1, l2, l3);
return p;
}
};
struct quant_dot_3args_2 : verify_program<quant_dot_3args_2>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::int8_type, {8, 2}};
migraphx::shape m2_shape{migraphx::shape::int8_type, {8, 7}};
migraphx::shape m3_shape{migraphx::shape::int32_type, {2, 7}};
auto l1 = p.add_parameter("a", m1_shape);
auto tl1 = p.add_instruction(migraphx::op::transpose{{1, 0}}, l1);
auto l2 = p.add_parameter("b", m2_shape);
auto l3 = p.add_parameter("c", m3_shape);
p.add_instruction(migraphx::op::quant_dot{1, 3}, tl1, l2, l3);
return p;
}
};
struct quant_dot_3args_3 : verify_program<quant_dot_3args_3>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::int8_type, {2, 8}};
migraphx::shape m2_shape{migraphx::shape::int8_type, {7, 8}};
migraphx::shape m3_shape{migraphx::shape::int32_type, {2, 7}};
auto l1 = p.add_parameter("a", m1_shape);
auto l2 = p.add_parameter("b", m2_shape);
auto tl2 = p.add_instruction(migraphx::op::transpose{{1, 0}}, l2);
auto l3 = p.add_parameter("c", m3_shape);
p.add_instruction(migraphx::op::quant_dot{2, 3}, l1, tl2, l3);
return p;
}
};
struct quant_dot_3args_4 : verify_program<quant_dot_3args_4>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::int8_type, {8, 2}};
migraphx::shape m2_shape{migraphx::shape::int8_type, {7, 8}};
migraphx::shape m3_shape{migraphx::shape::int32_type, {2, 7}};
auto l1 = p.add_parameter("a", m1_shape);
auto tl1 = p.add_instruction(migraphx::op::transpose{{1, 0}}, l1);
auto l2 = p.add_parameter("b", m2_shape);
auto tl2 = p.add_instruction(migraphx::op::transpose{{1, 0}}, l2);
auto l3 = p.add_parameter("c", m3_shape);
p.add_instruction(migraphx::op::quant_dot{3, 2}, tl1, tl2, l3);
return p;
}
};
struct batch_quant_dot_1 : verify_program<batch_quant_dot_1>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::int8_type, {3, 2, 8, 2}};
migraphx::shape m2_shape{migraphx::shape::int8_type, {3, 2, 7, 8}};
migraphx::shape m3_shape{migraphx::shape::int32_type, {3, 2, 2, 7}};
auto l1 = p.add_parameter("a", m1_shape);
auto tl1 = p.add_instruction(migraphx::op::transpose{{0, 1, 3, 2}}, l1);
auto l2 = p.add_parameter("b", m2_shape);
auto tl2 = p.add_instruction(migraphx::op::transpose{{0, 1, 3, 2}}, l2);
auto l3 = p.add_parameter("c", m3_shape);
p.add_instruction(migraphx::op::quant_dot{3, 2}, tl1, tl2, l3);
return p;
}
};
struct batch_quant_dot_2 : verify_program<batch_quant_dot_2>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape m1_shape{migraphx::shape::int8_type, {3, 2, 2, 8}};
migraphx::shape m2_shape{migraphx::shape::int8_type, {3, 2, 8, 7}};
migraphx::shape m3_shape{migraphx::shape::int32_type, {3, 2, 2, 7}};
auto l1 = p.add_parameter("a", m1_shape);
auto l2 = p.add_parameter("b", m2_shape);
auto l3 = p.add_parameter("c", m3_shape);
p.add_instruction(migraphx::op::quant_dot{1, 3}, l1, l2, l3);
return p;
}
};
struct test_contiguous : verify_program<test_contiguous>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {4, 4, 4, 3}, {48, 4, 1, 16}};
auto x = p.add_parameter("x", s);
p.add_instruction(migraphx::op::contiguous{}, x);
EXPECT(p.get_output_shapes().back().standard());
return p;
}
};
struct test_contiguous_broadcast : verify_program<test_contiguous_broadcast>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {1, 2}, {0, 1}};
auto x = p.add_parameter("x", s);
p.add_instruction(migraphx::op::contiguous{}, x);
EXPECT(p.get_output_shapes().back().standard());
return p;
}
};
struct test_contiguous_broadcast_transpose : verify_program<test_contiguous_broadcast_transpose>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {1, 3072, 768}, {0, 1, 3072}};
auto x = p.add_parameter("x", s);
p.add_instruction(migraphx::op::contiguous{}, x);
EXPECT(p.get_output_shapes().back().standard());
return p;
}
};
struct test_transpose : verify_program<test_transpose>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {4, 3, 4, 4}};
auto x = p.add_parameter("x", s);
std::vector<int64_t> perm = {0, 2, 3, 1};
auto l = p.add_instruction(migraphx::op::transpose{perm}, x);
p.add_instruction(migraphx::op::contiguous{}, l);
return p;
}
};
struct test_trans_ret : verify_program<test_trans_ret>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {4, 3, 3, 3}});
auto tx = p.add_instruction(migraphx::op::transpose{{0, 1, 3, 2}}, x);
p.add_return({tx});
return p;
}
};
struct test_batchnorm_inference_2 : verify_program<test_batchnorm_inference_2>
{
const size_t width = 14;
const size_t height = 14;
const size_t channels = 256;
const size_t batches = 1;
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {batches, channels, height, width}};
migraphx::shape vars{migraphx::shape::float_type, {channels}};
auto x = p.add_parameter("x", s);
auto scale = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 1)));
auto bias = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 2)));
auto mean = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 3)));
auto variance = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 4)));
p.add_instruction(migraphx::op::batch_norm_inference{}, x, scale, bias, mean, variance);
return p;
}
};
struct test_batchnorm_inference : verify_program<test_batchnorm_inference>
{
const size_t width = 3;
const size_t height = 3;
const size_t channels = 3;
const size_t batches = 4;
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {batches, channels, height, width}};
migraphx::shape vars{migraphx::shape::float_type, {channels}};
auto x = p.add_parameter("x", s);
auto scale = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 1)));
auto bias = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 2)));
auto mean = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 3)));
auto variance = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 4)));
p.add_instruction(migraphx::op::batch_norm_inference{}, x, scale, bias, mean, variance);
return p;
}
};
struct test_batchnorm_1d : verify_program<test_batchnorm_1d>
{
const size_t size = 3;
const size_t channels = 3;
const size_t batches = 4;
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {batches, channels, size}};
migraphx::shape vars{migraphx::shape::float_type, {channels}};
auto x = p.add_parameter("x", s);
auto scale = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 1)));
auto bias = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 2)));
auto mean = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 3)));
auto variance = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 4)));
p.add_instruction(migraphx::op::batch_norm_inference{}, x, scale, bias, mean, variance);
return p;
}
};
struct test_batchnorm_3d : verify_program<test_batchnorm_3d>
{
const size_t d1 = 2;
const size_t d2 = 2;
const size_t d3 = 2;
const size_t channels = 2;
const size_t batches = 2;
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {batches, channels, d1, d2, d3}};
migraphx::shape vars{migraphx::shape::float_type, {channels}};
auto x = p.add_parameter("x", s);
auto scale = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 1)));
auto bias = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 2)));
auto mean = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 3)));
auto variance = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 4)));
p.add_instruction(migraphx::op::batch_norm_inference{}, x, scale, bias, mean, variance);
return p;
}
};
struct test_batchnorm_1d_per_actv : verify_program<test_batchnorm_1d_per_actv>
{
const size_t d1 = 5;
const size_t channels = 2;
const size_t batches = 3;
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {batches, channels, d1}};
migraphx::shape vars{migraphx::shape::float_type, {channels, d1}};
auto x = p.add_parameter("x", s);
auto scale = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 1)));
auto bias = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 2)));
auto mean = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 3)));
auto variance = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 4)));
p.add_instruction(
migraphx::op::batch_norm_inference{
1.0e-5, 0.96f, migraphx::op::batch_norm_inference::per_activation},
x,
scale,
bias,
mean,
variance);
return p;
}
};
struct test_batchnorm_2d_per_actv : verify_program<test_batchnorm_2d_per_actv>
{
const size_t d1 = 2;
const size_t d2 = 4;
const size_t channels = 2;
const size_t batches = 3;
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {batches, channels, d1, d2}};
migraphx::shape vars{migraphx::shape::float_type, {channels, d1, d2}};
auto x = p.add_parameter("x", s);
auto scale = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 1)));
auto bias = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 2)));
auto mean = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 3)));
auto variance = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 4)));
p.add_instruction(
migraphx::op::batch_norm_inference{
1.0e-6, 0.9f, migraphx::op::batch_norm_inference::per_activation},
x,
scale,
bias,
mean,
variance);
return p;
}
};
struct test_batchnorm_3d_per_actv : verify_program<test_batchnorm_3d_per_actv>
{
const size_t d1 = 2;
const size_t d2 = 4;
const size_t d3 = 5;
const size_t channels = 2;
const size_t batches = 3;
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {batches, channels, d1, d2, d3}};
migraphx::shape vars{migraphx::shape::float_type, {channels, d1, d2, d3}};
auto x = p.add_parameter("x", s);
auto scale = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 1)));
auto bias = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 2)));
auto mean = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 3)));
auto variance = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 4)));
p.add_instruction(
migraphx::op::batch_norm_inference{
1.0e-6, 0.8f, migraphx::op::batch_norm_inference::per_activation},
x,
scale,
bias,
mean,
variance);
return p;
}
};
struct test_clip : verify_program<test_clip>
{
migraphx::program create_program() const
{
migraphx::program p;
auto x = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {3}});
auto min_val = p.add_literal(0.0f);
auto max_val = p.add_literal(6.0f);
min_val = p.add_instruction(migraphx::op::multibroadcast{{3}}, min_val);
max_val = p.add_instruction(migraphx::op::multibroadcast{{3}}, max_val);
p.add_instruction(migraphx::op::clip{}, x, min_val, max_val);
return p;
}
};
struct test_conv_bn : verify_program<test_conv_bn>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape xs{migraphx::shape::float_type, {1, 3, 224, 224}};
migraphx::shape ws{migraphx::shape::float_type, {64, 3, 7, 7}};
migraphx::shape vars{migraphx::shape::float_type, {64}};
auto x = p.add_parameter("x", xs);
auto w = p.add_parameter("w", ws);
auto conv = p.add_instruction(migraphx::op::convolution{{3, 3}, {2, 2}, {1, 1}}, x, w);
auto scale = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 1)));
auto bias = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 2)));
auto mean = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 3)));
auto variance = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 4)));
p.add_instruction(migraphx::op::batch_norm_inference{}, conv, scale, bias, mean, variance);
return p;
}
};
struct test_conv_bn_relu_pooling : verify_program<test_conv_bn_relu_pooling>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape xs{migraphx::shape::float_type, {1, 3, 224, 224}};
migraphx::shape ws{migraphx::shape::float_type, {64, 3, 7, 7}};
migraphx::shape vars{migraphx::shape::float_type, {64}};
auto x = p.add_parameter("x", xs);
auto w = p.add_parameter("w", ws);
auto conv = p.add_instruction(migraphx::op::convolution{{3, 3}, {2, 2}, {1, 1}}, x, w);
auto scale = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 1)));
auto bias = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 2)));
auto mean = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 3)));
auto variance = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 4)));
auto bn = p.add_instruction(
migraphx::op::batch_norm_inference{}, conv, scale, bias, mean, variance);
auto relu = p.add_instruction(migraphx::op::relu{}, bn);
p.add_instruction(migraphx::op::pooling{"average", {1, 1}, {2, 2}, {3, 3}}, relu);
return p;
}
};
struct quant_conv : verify_program<quant_conv>
{
migraphx::program create_program()
{
migraphx::program p;
migraphx::shape a_shape{migraphx::shape::int8_type, {2, 3, 4, 4}};
auto pa = p.add_parameter("a", a_shape);
migraphx::shape c_shape{migraphx::shape::int8_type, {2, 3, 3, 3}};
auto pc = p.add_parameter("c", c_shape);
p.add_instruction(migraphx::op::quant_convolution{}, pa, pc);
return p;
}
};
struct quant_conv_default_mode : verify_program<quant_conv_default_mode>
{
migraphx::program create_program()
{
migraphx::program p;
migraphx::shape a_shape{migraphx::shape::int8_type, {2, 3, 4, 4}};
auto pa = p.add_parameter("a", a_shape);
migraphx::shape c_shape{migraphx::shape::int8_type, {2, 3, 3, 3}};
auto pc = p.add_parameter("c", c_shape);
p.add_instruction(
migraphx::op::quant_convolution{{{0, 0}}, {{1, 1}}, {{1, 1}}, migraphx::op::same},
pa,
pc);
return p;
}
};
struct quant_conv_valid_mode : verify_program<quant_conv_valid_mode>
{
migraphx::program create_program()
{
migraphx::program p;
migraphx::shape a_shape{migraphx::shape::int8_type, {2, 3, 4, 4}};
auto pa = p.add_parameter("a", a_shape);
migraphx::shape c_shape{migraphx::shape::int8_type, {2, 3, 3, 3}};
auto pc = p.add_parameter("c", c_shape);
p.add_instruction(
migraphx::op::quant_convolution{{{0, 0}}, {{1, 1}}, {{1, 1}}, migraphx::op::valid},
pa,
pc);
return p;
}
};
struct quant_conv_padding : verify_program<quant_conv_padding>
{
migraphx::program create_program()
{
migraphx::program p;
migraphx::shape a_shape{migraphx::shape::int8_type, {2, 3, 4, 4}};
auto pa = p.add_parameter("a", a_shape);
migraphx::shape c_shape{migraphx::shape::int8_type, {2, 3, 3, 3}};
auto pc = p.add_parameter("c", c_shape);
p.add_instruction(migraphx::op::quant_convolution{{{1, 1}}, {{1, 1}}}, pa, pc);
return p;
}
};
struct quant_conv_padding_stride : verify_program<quant_conv_padding_stride>
{
migraphx::program create_program()
{
migraphx::program p;
migraphx::shape a_shape{migraphx::shape::int8_type, {2, 3, 4, 4}};
auto pa = p.add_parameter("a", a_shape);
migraphx::shape c_shape{migraphx::shape::int8_type, {2, 3, 3, 3}};
auto pc = p.add_parameter("c", c_shape);
p.add_instruction(migraphx::op::quant_convolution{{{1, 1}}, {{2, 2}}}, pa, pc);
return p;
}
};
struct test_concat_axis_1 : verify_program<test_concat_axis_1>
{
migraphx::program create_program() const
{
migraphx::program p;
int axis = 1;
migraphx::shape s0{migraphx::shape::int32_type, {2, 2}};
migraphx::shape s1{migraphx::shape::int32_type, {2, 3}};
migraphx::shape s2{migraphx::shape::int32_type, {2, 1}};
auto l0 = p.add_parameter("x", s0);
auto l1 = p.add_parameter("y", s1);
auto l2 = p.add_parameter("z", s2);
p.add_instruction(migraphx::op::concat{axis}, l0, l1, l2);
return p;
}
};
struct test_concat_axis_neg_1 : verify_program<test_concat_axis_neg_1>
{
migraphx::program create_program() const
{
migraphx::program p;
int axis = -1;
migraphx::shape s0{migraphx::shape::int32_type, {2, 2}};
migraphx::shape s1{migraphx::shape::int32_type, {2, 3}};
migraphx::shape s2{migraphx::shape::int32_type, {2, 1}};
auto l0 = p.add_parameter("x", s0);
auto l1 = p.add_parameter("y", s1);
auto l2 = p.add_parameter("z", s2);
p.add_instruction(migraphx::op::concat{axis}, l0, l1, l2);
return p;
}
};
struct test_concat_axis_0 : verify_program<test_concat_axis_0>
{
migraphx::program create_program() const
{
migraphx::program p;
int axis = 0;
migraphx::shape s0{migraphx::shape::int32_type, {2, 2}};
migraphx::shape s1{migraphx::shape::int32_type, {3, 2}};
migraphx::shape s2{migraphx::shape::int32_type, {1, 2}};
auto l0 = p.add_parameter("x", s0);
auto l1 = p.add_parameter("y", s1);
auto l2 = p.add_parameter("z", s2);
p.add_instruction(migraphx::op::concat{axis}, l0, l1, l2);
return p;
}
};
struct test_concat_transpose : verify_program<test_concat_transpose>
{
migraphx::program create_program() const
{
migraphx::program p;
int axis = 1;
migraphx::shape s0{migraphx::shape::int32_type, {2, 2}};
migraphx::shape s1{migraphx::shape::int32_type, {3, 2}};
migraphx::shape s2{migraphx::shape::int32_type, {2, 4}};
auto l0 = p.add_parameter("x", s0);
auto lp1 = p.add_parameter("y", s1);
auto l1 = p.add_instruction(migraphx::op::transpose{{1, 0}}, lp1);
auto l2 = p.add_parameter("z", s2);
p.add_instruction(migraphx::op::concat{axis}, l0, l1, l2);
return p;
}
};
struct test_concat_transpose2 : verify_program<test_concat_transpose2>
{
migraphx::program create_program() const
{
migraphx::program p;
int axis = 1;
migraphx::shape s0{migraphx::shape::int32_type, {2, 2}};
migraphx::shape s1{migraphx::shape::int32_type, {2, 3}};
migraphx::shape s2{migraphx::shape::int32_type, {5, 2}};
auto l0 = p.add_parameter("x", s0);
auto l1 = p.add_parameter("y", s1);
auto lp2 = p.add_parameter("z", s2);
auto l2 = p.add_instruction(migraphx::op::transpose{{1, 0}}, lp2);
p.add_instruction(migraphx::op::concat{axis}, l0, l1, l2);
return p;
}
};
struct test_concat_transpose3 : verify_program<test_concat_transpose3>
{
migraphx::program create_program() const
{
migraphx::program p;
int axis = 1;
migraphx::shape s0{migraphx::shape::int32_type, {2, 2}};
migraphx::shape s1{migraphx::shape::int32_type, {3, 2}};
migraphx::shape s2{migraphx::shape::int32_type, {5, 2}};
auto l0 = p.add_parameter("x", s0);
auto lp1 = p.add_parameter("y", s1);
auto l1 = p.add_instruction(migraphx::op::transpose{{1, 0}}, lp1);
auto lp2 = p.add_parameter("z", s2);
auto l2 = p.add_instruction(migraphx::op::transpose{{1, 0}}, lp2);
p.add_instruction(migraphx::op::concat{axis}, l0, l1, l2);
return p;
}
};
struct test_concat_relu : verify_program<test_concat_relu>
{
migraphx::program create_program() const
{
migraphx::program p;
int axis = 0;
migraphx::shape s0{migraphx::shape::float_type, {2, 2}};
migraphx::shape s1{migraphx::shape::float_type, {3, 2}};
migraphx::shape s2{migraphx::shape::float_type, {1, 2}};
auto l0 = p.add_parameter("x", s0);
auto l1 = p.add_parameter("y", s1);
auto l2 = p.add_parameter("z", s2);
auto r0 = p.add_instruction(migraphx::op::relu{}, l0);
auto r1 = p.add_instruction(migraphx::op::relu{}, l1);
auto r2 = p.add_instruction(migraphx::op::relu{}, l2);
auto c0 = p.add_instruction(migraphx::op::concat{axis}, r0, r1, r2);
p.add_instruction(migraphx::op::relu{}, c0);
return p;
}
};
struct test_pad : verify_program<test_pad>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s0{migraphx::shape::int32_type, {1, 96, 165, 165}};
std::vector<int64_t> pads0 = {0, 0, 0, 0, 0, 0, 1, 1};
std::vector<int64_t> pads1 = {0, 0, 0, 0, 1, 1, 1, 1};
std::vector<int64_t> pads2 = {1, 1, 1, 1, 0, 0, 0, 0};
std::vector<int64_t> pads3 = {1, 0, 1, 0, 1, 0, 2, 0};
auto l0 = p.add_parameter("x", s0);
p.add_instruction(migraphx::op::pad{pads0}, l0);
p.add_instruction(migraphx::op::pad{pads1}, l0);
p.add_instruction(migraphx::op::pad{pads2}, l0);
p.add_instruction(migraphx::op::pad{pads3}, l0);
return p;
}
};
struct test_pad_transposed : verify_program<test_pad_transposed>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::int32_type, {1, 224, 224, 3}};
auto x = p.add_parameter("x", s);
auto t = p.add_instruction(migraphx::op::transpose{{0, 3, 1, 2}}, x);
p.add_instruction(migraphx::op::pad{{0, 0, 2, 2, 0, 0, 3, 3}}, t);
return p;
}
};
struct test_pad_int8 : verify_program<test_pad_int8>
{
migraphx::program create_program() const
{
migraphx::program p;
std::vector<int8_t> data0 = {0, 1, 2, 3};
migraphx::shape s0{migraphx::shape::float_type, {2, 2}};
auto l0 = p.add_literal(migraphx::literal{s0, data0});
migraphx::op::pad op{};
op.value = std::numeric_limits<int8_t>::lowest();
op.pads = {0, 0, 1, 1};
p.add_instruction(op, l0);
return p;
}
};
struct test_pad_lowest : verify_program<test_pad_lowest>
{
migraphx::program create_program() const
{
migraphx::program p;
std::vector<migraphx::half> data0(4);
std::iota(data0.begin(), data0.end(), 0);
migraphx::shape s0{migraphx::shape::half_type, {2, 2}};
auto l0 = p.add_literal(migraphx::literal{s0, data0});
migraphx::op::pad op{};
op.value = std::numeric_limits<float>::lowest();
op.pads = {0, 0, 1, 1};
p.add_instruction(op, l0);
return p;
}
};
struct test_pad_highest : verify_program<test_pad_highest>
{
migraphx::program create_program() const
{
migraphx::program p;
std::vector<migraphx::half> data0(4);
std::iota(data0.begin(), data0.end(), 0);
migraphx::shape s0{migraphx::shape::half_type, {2, 2}};
auto l0 = p.add_literal(migraphx::literal{s0, data0});
migraphx::op::pad op{};
op.value = std::numeric_limits<float>::max();
op.pads = {0, 0, 1, 1};
p.add_instruction(op, l0);
return p;
}
};
struct test_pooling_autopad : verify_program<test_pooling_autopad>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s0{migraphx::shape::float_type, {1, 3, 63, 63}};
auto l0 = p.add_parameter("x", s0);
migraphx::op::pooling op{"max"};
op.lengths = {2, 2};
op.stride = {2, 2};
p.add_instruction(op, l0);
return p;
}
};
struct test_gather : verify_program<test_gather>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {3, 3}};
migraphx::shape s_indices{migraphx::shape::int32_type, {2, 2}};
std::vector<int> indices{1, 2, 2, 1};
auto a0 = p.add_parameter("data", s);
auto a1 = p.add_literal(migraphx::literal{s_indices, indices});
int axis = 0;
p.add_instruction(migraphx::op::gather{axis}, a0, a1);
return p;
}
};
struct test_gather_neg_axis : verify_program<test_gather_neg_axis>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {3, 3}};
migraphx::shape s_indices{migraphx::shape::int32_type, {2, 2}};
std::vector<int> indices{1, 2, 2, 1};
auto a0 = p.add_parameter("data", s);
auto a1 = p.add_literal(migraphx::literal{s_indices, indices});
int axis = -1;
p.add_instruction(migraphx::op::gather{axis}, a0, a1);
return p;
}
};
struct test_gather_neg_indices : verify_program<test_gather_neg_indices>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {3, 3}};
migraphx::shape s_indices{migraphx::shape::int32_type, {2, 2}};
std::vector<int> indices{-2, -1, -1, -2};
auto a0 = p.add_parameter("data", s);
auto a1 = p.add_literal(migraphx::literal{s_indices, indices});
int axis = -1;
p.add_instruction(migraphx::op::gather{axis}, a0, a1);
return p;
}
};
struct test_gather_scalar_output : verify_program<test_gather_scalar_output>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {3}};
migraphx::shape s_indices{migraphx::shape::int32_type};
std::vector<int> indices{1};
auto a0 = p.add_parameter("data", s);
auto a1 = p.add_literal(migraphx::literal{s_indices, indices});
int axis = 0;
p.add_instruction(migraphx::op::gather{axis}, a0, a1);
return p;
}
};
struct test_gather_scalar_index : verify_program<test_gather_scalar_index>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {3, 3}};
migraphx::shape s_indices{migraphx::shape::int32_type};
std::vector<int> indices{1};
auto a0 = p.add_parameter("data", s);
auto a1 = p.add_literal(migraphx::literal{s_indices, indices});
int axis = -1;
p.add_instruction(migraphx::op::gather{axis}, a0, a1);
return p;
}
};
struct test_gather_1d_index : verify_program<test_gather_1d_index>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {3, 3}};
migraphx::shape s_indices{migraphx::shape::int32_type, {1}};
std::vector<int> indices{1};
auto a0 = p.add_parameter("data", s);
auto a1 = p.add_literal(migraphx::literal{s_indices, indices});
int axis = -1;
p.add_instruction(migraphx::op::gather{axis}, a0, a1);
return p;
}
};
void manual_identity()
{
migraphx::program p;
std::vector<float> data0 = {0, 1, 2, 3};
migraphx::shape s0{migraphx::shape::float_type, {2, 2}};
auto l0 = p.add_literal(migraphx::literal{s0, data0});
p.add_instruction(migraphx::op::identity{}, l0);
p.compile(migraphx::gpu::target{});
migraphx::program::parameter_map m;
for(auto&& x : p.get_parameter_shapes())
{
m[x.first] = migraphx::gpu::to_gpu(migraphx::generate_argument(x.second));
}
auto result = migraphx::gpu::from_gpu(p.eval(m).back());
std::cout << result << std::endl;
}
void manual_test_concat_relu()
{
migraphx::program p;
int axis = 0;
std::vector<float> data0 = {0, 1, 2, 3};
std::vector<float> data1 = {4, 5, 6, 7, 8, 9};
std::vector<float> data2 = {10, 11};
migraphx::shape s0{migraphx::shape::float_type, {2, 2}};
migraphx::shape s1{migraphx::shape::float_type, {3, 2}};
migraphx::shape s2{migraphx::shape::float_type, {1, 2}};
auto l0 = p.add_literal(migraphx::literal{s0, data0});
auto l1 = p.add_literal(migraphx::literal{s1, data1});
auto l2 = p.add_literal(migraphx::literal{s2, data2});
auto r0 = p.add_instruction(migraphx::op::relu{}, l0);
auto r1 = p.add_instruction(migraphx::op::relu{}, l1);
auto r2 = p.add_instruction(migraphx::op::relu{}, l2);
auto c0 = p.add_instruction(migraphx::op::concat{axis}, r0, r1, r2);
p.add_instruction(migraphx::op::relu{}, c0);
p.compile(migraphx::gpu::target{});
migraphx::program::parameter_map m;
for(auto&& x : p.get_parameter_shapes())
{
m[x.first] = migraphx::gpu::to_gpu(migraphx::generate_argument(x.second));
}
auto result = migraphx::gpu::from_gpu(p.eval(m).back());
std::cout << result << std::endl;
}
struct test_conv_bn_relu_pooling2 : verify_program<test_conv_bn_relu_pooling2>
{
static migraphx::instruction_ref
add_bn(migraphx::program& p, migraphx::instruction_ref x, std::size_t channels)
{
migraphx::shape vars{migraphx::shape::float_type, {channels}};
auto scale = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 1 + channels)));
auto bias = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 2 + channels)));
auto mean = p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 3 + channels)));
auto variance =
p.add_literal(migraphx::abs(migraphx::generate_literal(vars, 4 + channels)));
return p.add_instruction(
migraphx::op::batch_norm_inference{}, x, scale, bias, mean, variance);
}
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape xs1{migraphx::shape::float_type, {1, 512, 7, 7}};
migraphx::shape xs2{migraphx::shape::float_type, {1, 1024, 14, 14}};
migraphx::shape ws1{migraphx::shape::float_type, {2048, 512, 1, 1}};
migraphx::shape ws2{migraphx::shape::float_type, {2048, 1024, 1, 1}};
auto x1 = p.add_parameter("x1", xs1);
auto w1 = p.add_parameter("w1", ws1);
auto conv1 = p.add_instruction(migraphx::op::convolution{{0, 0}, {1, 1}, {1, 1}}, x1, w1);
auto bn1 = add_bn(p, conv1, 2048);
auto x2 = p.add_parameter("x2", xs2);
auto w2 = p.add_parameter("w2", ws2);
auto conv2 = p.add_instruction(migraphx::op::convolution{{0, 0}, {2, 2}, {1, 1}}, x2, w2);
auto bn2 = add_bn(p, conv2, 2048);
auto add = p.add_instruction(migraphx::op::add{}, bn1, bn2);
auto relu = p.add_instruction(migraphx::op::relu{}, add);
p.add_instruction(migraphx::op::pooling{"average", {1, 1}, {2, 2}, {3, 3}}, relu);
return p;
}
};
template <int Axis, migraphx::shape::type_t T>
struct test_logsoftmax : verify_program<test_logsoftmax<Axis, T>>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{T, {10, 4, 2080, 6}};
auto param = p.add_parameter("0", s);
p.add_instruction(migraphx::op::logsoftmax{Axis}, param);
return p;
}
};
template struct test_logsoftmax<0, migraphx::shape::float_type>;
template struct test_logsoftmax<1, migraphx::shape::float_type>;
template struct test_logsoftmax<2, migraphx::shape::float_type>;
template struct test_logsoftmax<3, migraphx::shape::float_type>;
template struct test_logsoftmax<1, migraphx::shape::double_type>;
template struct test_logsoftmax<3, migraphx::shape::double_type>;
template struct test_logsoftmax<1, migraphx::shape::half_type>;
template struct test_logsoftmax<0, migraphx::shape::half_type>;
template struct test_logsoftmax<2, migraphx::shape::half_type>;
template struct test_logsoftmax<3, migraphx::shape::half_type>;
struct test_fp32_fp16_lall : verify_program<test_fp32_fp16_lall>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {2, 3}};
std::vector<float> data(2 * 3);
std::iota(data.begin(), data.end(), 1.0f);
auto l1 = p.add_literal(migraphx::literal(s, data));
auto l2 = p.add_parameter("p2", s);
p.add_instruction(migraphx::op::add{}, l1, l2);
migraphx::quantize_fp16(p, {"all"});
return p;
};
};
struct test_fp32_fp16_ladd : verify_program<test_fp32_fp16_ladd>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {2, 3}};
std::vector<float> data(2 * 3);
std::iota(data.begin(), data.end(), 1.0f);
auto l1 = p.add_literal(migraphx::literal(s, data));
auto l2 = p.add_parameter("p2", s);
p.add_instruction(migraphx::op::add{}, l1, l2);
migraphx::quantize_fp16(p, {"add"});
return p;
};
};
struct test_fp32_fp16_add : verify_program<test_fp32_fp16_add>
{
migraphx::program create_program()
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {2, 3}};
auto p1 = p.add_parameter("x", s);
auto p2 = p.add_parameter("y", s);
auto sum = p.add_instruction(migraphx::op::add{}, p1, p2);
auto diff = p.add_instruction(migraphx::op::sub{}, sum, p2);
p.add_instruction(migraphx::op::add{}, diff, p1);
migraphx::quantize_fp16(p, {"add"});
return p;
};
};
struct test_fp32_fp16_sub : verify_program<test_fp32_fp16_sub>
{
migraphx::program create_program()
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {2, 3}};
auto p1 = p.add_parameter("x", s);
auto p2 = p.add_parameter("y", s);
auto sum = p.add_instruction(migraphx::op::add{}, p1, p2);
auto diff = p.add_instruction(migraphx::op::sub{}, sum, p2);
p.add_instruction(migraphx::op::add{}, diff, p1);
migraphx::quantize_fp16(p, {"sub"});
return p;
};
};
template <class Op, int Axis, migraphx::shape::type_t T>
struct test_reduce_op_large : verify_program<test_reduce_op_large<Op, Axis, T>>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{T, {3, 1026, 4, 3}};
auto x = p.add_parameter("x", s);
p.add_instruction(Op{{1}}, x);
return p;
};
};
template struct test_reduce_op_large<migraphx::op::reduce_max, 1, migraphx::shape::float_type>;
template struct test_reduce_op_large<migraphx::op::reduce_mean, 1, migraphx::shape::float_type>;
template struct test_reduce_op_large<migraphx::op::reduce_min, 1, migraphx::shape::float_type>;
template struct test_reduce_op_large<migraphx::op::reduce_prod, 2, migraphx::shape::float_type>;
template struct test_reduce_op_large<migraphx::op::reduce_sum, 1, migraphx::shape::float_type>;
template <class Op, int Axis, migraphx::shape::type_t T>
struct test_reduce_op_small : verify_program<test_reduce_op_small<Op, Axis, T>>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{T, {3, 4, 8, 8}};
auto x = p.add_parameter("x", s);
p.add_instruction(Op{{1}}, x);
return p;
};
};
template struct test_reduce_op_small<migraphx::op::reduce_sum, 2, migraphx::shape::int32_type>;
template struct test_reduce_op_small<migraphx::op::reduce_mean, 2, migraphx::shape::int32_type>;
template struct test_reduce_op_small<migraphx::op::reduce_max, 2, migraphx::shape::int32_type>;
template struct test_reduce_op_small<migraphx::op::reduce_min, 2, migraphx::shape::int32_type>;
template struct test_reduce_op_small<migraphx::op::reduce_sum, 2, migraphx::shape::half_type>;
template struct test_reduce_op_small<migraphx::op::reduce_mean, 2, migraphx::shape::half_type>;
template struct test_reduce_op_small<migraphx::op::reduce_max, 2, migraphx::shape::half_type>;
template struct test_reduce_op_small<migraphx::op::reduce_min, 2, migraphx::shape::half_type>;
template struct test_reduce_op_small<migraphx::op::reduce_prod, -2, migraphx::shape::half_type>;
struct test_rsqrt : verify_program<test_rsqrt>
{
migraphx::program create_program() const
{
migraphx::program p;
std::vector<size_t> input_lens{1, 3, 16, 16};
migraphx::shape s{migraphx::shape::float_type, input_lens};
auto x = p.add_parameter("x", s);
auto min_val = p.add_literal(1.0f);
auto max_val = p.add_literal(std::numeric_limits<float>::max());
min_val = p.add_instruction(migraphx::op::multibroadcast{input_lens}, min_val);
max_val = p.add_instruction(migraphx::op::multibroadcast{input_lens}, max_val);
auto l0 = p.add_instruction(migraphx::op::clip{}, x, min_val, max_val);
p.add_instruction(migraphx::op::rsqrt{}, l0);
return p;
};
};
struct test_round : verify_program<test_round>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {2, 3, 4, 6}};
auto param = p.add_parameter("x", s);
p.add_instruction(migraphx::op::round{}, param);
return p;
};
};
struct test_ceil : verify_program<test_ceil>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::double_type, {2, 3, 4, 6}};
auto param = p.add_parameter("x", s);
p.add_instruction(migraphx::op::ceil{}, param);
return p;
};
};
struct test_floor : verify_program<test_floor>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::float_type, {2, 3, 4, 6}};
auto param = p.add_parameter("x", s);
p.add_instruction(migraphx::op::floor{}, param);
return p;
};
};
struct test_convert : verify_program<test_convert>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape sa{migraphx::shape::float_type, {8, 24}};
migraphx::shape sb{migraphx::shape::float_type, {24, 6}};
auto pa = p.add_parameter("a", sa);
auto pb = p.add_parameter("b", sb);
auto ia = p.add_instruction(migraphx::op::convert{migraphx::shape::int8_type}, pa);
auto ib = p.add_instruction(migraphx::op::convert{migraphx::shape::int8_type}, pb);
p.add_instruction(migraphx::op::quant_dot{}, ia, ib);
return p;
};
};
struct test_recip : verify_program<test_recip>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::double_type, {3}};
auto x = p.add_parameter("x", s);
p.add_instruction(migraphx::op::recip{}, x);
return p;
}
};
struct test_neg : verify_program<test_neg>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::double_type, {2, 3, 4, 6}};
auto input = p.add_parameter("x", s);
p.add_instruction(migraphx::op::neg{}, input);
return p;
};
};
struct test_equal : verify_program<test_equal>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::double_type, {2, 3, 4, 6}};
auto input1 = p.add_parameter("x", s);
auto input2 = p.add_parameter("y", s);
auto r = p.add_instruction(migraphx::op::equal{}, input1, input2);
p.add_return({r});
return p;
};
};
struct test_equal_brcst : verify_program<test_equal_brcst>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s0{migraphx::shape::float_type, {3, 3}};
auto l0 = p.add_parameter("x", s0);
migraphx::shape s1{migraphx::shape::float_type, {3, 1}};
auto l1 = p.add_parameter("y", s1);
auto bl1 = p.add_instruction(migraphx::op::multibroadcast{s0.lens()}, l1);
auto r = p.add_instruction(migraphx::op::equal{}, l0, bl1);
p.add_return({r});
return p;
};
};
struct test_greater : verify_program<test_greater>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::double_type, {2, 3, 4, 6}};
auto input1 = p.add_parameter("x", s);
auto input2 = p.add_parameter("y", s);
auto r = p.add_instruction(migraphx::op::greater{}, input1, input2);
p.add_return({r});
return p;
};
};
struct test_greater_brcst : verify_program<test_greater_brcst>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s0{migraphx::shape::float_type, {3, 3}};
auto l0 = p.add_parameter("x", s0);
migraphx::shape s1{migraphx::shape::float_type, {3, 1}};
auto l1 = p.add_parameter("y", s1);
auto bl1 = p.add_instruction(migraphx::op::multibroadcast{s0.lens()}, l1);
auto r = p.add_instruction(migraphx::op::greater{}, l0, bl1);
p.add_return({r});
return p;
};
};
struct test_less : verify_program<test_less>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s{migraphx::shape::double_type, {2, 3, 4, 6}};
auto input1 = p.add_parameter("x", s);
auto input2 = p.add_parameter("y", s);
auto r = p.add_instruction(migraphx::op::less{}, input1, input2);
p.add_return({r});
return p;
};
};
struct test_less_brcst : verify_program<test_less_brcst>
{
migraphx::program create_program() const
{
migraphx::program p;
migraphx::shape s0{migraphx::shape::float_type, {3, 3}};
auto l0 = p.add_parameter("x", s0);
migraphx::shape s1{migraphx::shape::float_type, {3, 1}};
auto l1 = p.add_parameter("y", s1);
auto bl1 = p.add_instruction(migraphx::op::multibroadcast{s0.lens()}, l1);
auto r = p.add_instruction(migraphx::op::less{}, l0, bl1);
p.add_return({r});
return p;
};
};
int main(int argc, const char* argv[]) { test::run(argc, argv); }
test/gpu/quantization.cpp
View file @ 500d9441
...@@ -4,7 +4,7 @@ ...@@ -4,7 +4,7 @@
#include <migraphx/instruction.hpp> #include <migraphx/instruction.hpp>
#include <migraphx/quantization.hpp> #include <migraphx/quantization.hpp>
#include <migraphx/generate.hpp> #include <migraphx/generate.hpp>
#include <migraphx/cpu/target.hpp> #include <migraphx/ref/target.hpp>
#include <migraphx/gpu/target.hpp> #include <migraphx/gpu/target.hpp>
#include <migraphx/verify.hpp> #include <migraphx/verify.hpp>
#include <migraphx/quantization.hpp> #include <migraphx/quantization.hpp>
...@@ -12,23 +12,23 @@ ...@@ -12,23 +12,23 @@
#include <migraphx/propagate_constant.hpp> #include <migraphx/propagate_constant.hpp>
#include <migraphx/pass_manager.hpp> #include <migraphx/pass_manager.hpp>
#include <migraphx/onnx.hpp> #include <migraphx/onnx.hpp>
#include "test.hpp" #include <test.hpp>
#include <migraphx/half.hpp> #include <migraphx/half.hpp>
TEST_CASE(gpu_target_copy) TEST_CASE(gpu_target_copy)
{ {
migraphx::target gpu_t = migraphx::gpu::target{}; migraphx::target gpu_t = migraphx::gpu::target{};
migraphx::target cpu_t = migraphx::cpu::target{}; migraphx::target ref_t = migraphx::ref::target{};
migraphx::shape s{migraphx::shape::int8_type, {2, 3, 4, 5}}; migraphx::shape s{migraphx::shape::int8_type, {2, 3, 4, 5}};
auto cpu_arg_orig = migraphx::generate_argument(s, 0x123456L); auto ref_arg_orig = migraphx::generate_argument(s, 0x123456L);
auto gpu_arg = gpu_t.copy_to(cpu_arg_orig); auto gpu_arg = gpu_t.copy_to(ref_arg_orig);
auto cpu_arg_final = gpu_t.copy_from(gpu_arg); auto ref_arg_final = gpu_t.copy_from(gpu_arg);
std::vector<int8_t> val_orig; std::vector<int8_t> val_orig;
cpu_arg_orig.visit([&](auto v) { val_orig.assign(v.begin(), v.end()); }); ref_arg_orig.visit([&](auto v) { val_orig.assign(v.begin(), v.end()); });
std::vector<int8_t> val_final; std::vector<int8_t> val_final;
cpu_arg_final.visit([&](auto v) { val_final.assign(v.begin(), v.end()); }); ref_arg_final.visit([&](auto v) { val_final.assign(v.begin(), v.end()); });
EXPECT(migraphx::verify_range(val_orig, val_final)); EXPECT(migraphx::verify_range(val_orig, val_final));
} }
...@@ -82,15 +82,15 @@ TEST_CASE(int8_quantization) ...@@ -82,15 +82,15 @@ TEST_CASE(int8_quantization)
m["a"] = migraphx::generate_argument(sa); m["a"] = migraphx::generate_argument(sa);
m["b"] = migraphx::generate_argument(sb); m["b"] = migraphx::generate_argument(sb);
m["c"] = migraphx::generate_argument(sc); m["c"] = migraphx::generate_argument(sc);
std::vector<float> cpu_result; std::vector<float> ref_result;
migraphx::target cpu_t = migraphx::cpu::target{}; migraphx::target ref_t = migraphx::ref::target{};
run_prog(p, cpu_t, m, cpu_result); run_prog(p, ref_t, m, ref_result);
std::vector<float> gpu_result; std::vector<float> gpu_result;
migraphx::target gpu_t = migraphx::gpu::target{}; migraphx::target gpu_t = migraphx::gpu::target{};
run_prog(p, gpu_t, m, gpu_result); run_prog(p, gpu_t, m, gpu_result);
EXPECT(migraphx::verify_range(cpu_result, gpu_result)); EXPECT(migraphx::verify_range(ref_result, gpu_result));
} }
} }
......
test/gpu/rnn_ops_test.cpp deleted 100644 → 0
View file @ ba33d25c
#include <test.hpp>
#include "test_utils.hpp"
#ifdef __clang__
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wglobal-constructors"
#endif
struct test_rnn_forward : verify_program<test_rnn_forward>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 1;
std::size_t hidden_size = 4;
std::size_t input_size = 3;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
auto hs = p.add_instruction(migraphx::op::rnn{hidden_size,
{migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::forward,
clip},
seq,
w,
r,
bias,
und,
ih);
auto lho = p.add_instruction(migraphx::op::rnn_last_hs_output{}, hs);
p.add_return({hs, lho});
return p;
}
};
struct test_rnn_forward10 : verify_program<test_rnn_forward10>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 10;
std::size_t hidden_size = 4;
std::size_t input_size = 3;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
auto hs = p.add_instruction(migraphx::op::rnn{hidden_size,
{migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::forward,
clip},
seq,
w,
r,
bias,
und,
ih);
auto lho = p.add_instruction(migraphx::op::rnn_last_hs_output{}, hs);
p.add_return({hs, lho});
return p;
}
};
struct test_rnn_sql_1 : verify_program<test_rnn_sql_1>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 10;
std::size_t hidden_size = 4;
std::size_t input_size = 3;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};
migraphx::shape s_shape{migraphx::shape::int32_type, {batch_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
std::vector<int> sl_data{5, 7};
auto sql = p.add_literal(migraphx::literal{s_shape, sl_data});
auto ih = p.add_parameter("ih", ih_shape);
auto hs = p.add_instruction(migraphx::op::rnn{hidden_size,
{migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::forward,
clip},
seq,
w,
r,
bias,
sql,
ih);
auto last_hs = p.add_instruction(migraphx::op::rnn_last_hs_output{}, hs);
p.add_return({hs, last_hs});
return p;
}
};
struct test_rnn_sql_2 : verify_program<test_rnn_sql_2>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 10;
std::size_t hidden_size = 4;
std::size_t input_size = 3;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};
migraphx::shape s_shape{migraphx::shape::int32_type, {batch_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
auto seq_orig = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
migraphx::shape pad_s{migraphx::shape::float_type, {2, batch_size, input_size}};
std::vector<float> pad_data(pad_s.elements(), 0.0f);
auto seq_pad = p.add_literal(migraphx::literal{pad_s, pad_data});
auto seq = p.add_instruction(migraphx::op::concat{0}, seq_orig, seq_pad);
std::vector<int> sl_data(batch_size, static_cast<int>(seq_len));
auto sql = p.add_literal(migraphx::literal{s_shape, sl_data});
auto ih = p.add_parameter("ih", ih_shape);
auto hs = p.add_instruction(migraphx::op::rnn{hidden_size,
{migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::forward,
clip},
seq,
w,
r,
bias,
sql,
ih);
auto last_hs = p.add_instruction(migraphx::op::rnn_last_hs_output{}, hs);
p.add_return({hs, last_hs});
return p;
}
};
struct test_rnn_reverse : verify_program<test_rnn_reverse>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 1;
std::size_t hidden_size = 4;
std::size_t input_size = 3;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
p.add_instruction(migraphx::op::rnn{hidden_size,
{migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::reverse,
clip},
seq,
w,
r,
bias,
und,
ih);
return p;
}
};
struct test_rnn_reverse2 : verify_program<test_rnn_reverse2>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 2;
std::size_t hidden_size = 4;
std::size_t input_size = 3;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
p.add_instruction(migraphx::op::rnn{hidden_size,
{migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::reverse,
clip},
seq,
w,
r,
bias,
und,
ih);
return p;
}
};
struct test_rnn_3args : verify_program<test_rnn_3args>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 1;
std::size_t hidden_size = 4;
std::size_t input_size = 3;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
p.add_instruction(migraphx::op::rnn{hidden_size,
{migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::reverse,
clip},
seq,
w,
r);
return p;
}
};
struct test_rnn_4args : verify_program<test_rnn_4args>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 5;
std::size_t hidden_size = 4;
std::size_t input_size = 3;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
p.add_instruction(migraphx::op::rnn{hidden_size,
{migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::reverse,
clip},
seq,
w,
r,
bias);
return p;
}
};
struct test_rnn_5args : verify_program<test_rnn_5args>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 10;
std::size_t hidden_size = 4;
std::size_t input_size = 3;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
auto output =
p.add_instruction(migraphx::op::rnn{hidden_size,
{migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::forward,
clip},
seq,
w,
r,
bias,
und);
p.add_instruction(migraphx::op::rnn_last_hs_output{}, output);
return p;
}
};
struct test_rnn_bidirectional : verify_program<test_rnn_bidirectional>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 1;
std::size_t hidden_size = 4;
std::size_t input_size = 3;
std::size_t num_dirct = 2;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
auto output =
p.add_instruction(migraphx::op::rnn{hidden_size,
{migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::bidirectional,
clip},
seq,
w,
r,
bias,
und,
ih);
p.add_instruction(migraphx::op::rnn_last_hs_output{}, output);
return p;
}
};
struct test_rnn_bidirectional10 : verify_program<test_rnn_bidirectional10>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 10;
std::size_t hidden_size = 4;
std::size_t input_size = 3;
std::size_t num_dirct = 2;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
auto output =
p.add_instruction(migraphx::op::rnn{hidden_size,
{migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::bidirectional,
clip},
seq,
w,
r,
bias,
und,
ih);
p.add_instruction(migraphx::op::rnn_last_hs_output{}, output);
return p;
}
};
struct test_rnn_bi_3args : verify_program<test_rnn_bi_3args>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 10;
std::size_t hidden_size = 4;
std::size_t input_size = 3;
std::size_t num_dirct = 2;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type, {num_dirct, hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type, {num_dirct, hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 2 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto output =
p.add_instruction(migraphx::op::rnn{hidden_size,
{migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::bidirectional,
clip},
seq,
w,
r);
p.add_instruction(migraphx::op::rnn_last_hs_output{}, output);
return p;
}
};
struct test_gru_forward : verify_program<test_gru_forward>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 6 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
auto hs =
p.add_instruction(migraphx::op::gru{hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::forward,
clip},
seq,
w,
r,
bias,
und,
ih);
auto lho = p.add_instruction(migraphx::op::rnn_last_hs_output{}, hs);
p.add_return({lho, hs});
return p;
}
};
struct test_gru_forward_3args_und : verify_program<test_gru_forward_3args_und>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
p.add_instruction(migraphx::op::gru{hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::forward,
clip},
seq,
w,
r,
und,
und,
und);
return p;
}
};
struct test_var_sl_gru_forward : verify_program<test_var_sl_gru_forward>
{
migraphx::program create_program() const
{
std::size_t batch_size = 3;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 6 * hidden_size}};
migraphx::shape sl_shape{migraphx::shape::int32_type, {batch_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
std::vector<int> sl_data{3, 2, 1};
auto sql = p.add_literal(migraphx::literal{sl_shape, sl_data});
auto hs =
p.add_instruction(migraphx::op::gru{hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::forward,
clip},
seq,
w,
r,
bias,
sql,
ih);
auto lho = p.add_instruction(migraphx::op::rnn_last_hs_output{}, hs);
p.add_return({lho, hs});
return p;
}
};
struct test_gru_forward_3args : verify_program<test_gru_forward_3args>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
p.add_instruction(migraphx::op::gru{hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::forward,
clip},
seq,
w,
r);
return p;
}
};
struct test_gru_forward_seq1 : verify_program<test_gru_forward_seq1>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 1;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
p.add_instruction(migraphx::op::gru{hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::forward,
clip},
seq,
w,
r);
return p;
}
};
struct test_gru_forward_default_actv : verify_program<test_gru_forward_default_actv>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 1;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
p.add_instruction(
migraphx::op::gru{hidden_size, {}, migraphx::op::rnn_direction::forward, clip},
seq,
w,
r);
return p;
}
};
struct test_gru_two_outputs : verify_program<test_gru_two_outputs>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 1;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto hs = p.add_instruction(
migraphx::op::gru{hidden_size, {}, migraphx::op::rnn_direction::forward, clip},
seq,
w,
r);
auto last_hs = p.add_instruction(migraphx::op::rnn_last_hs_output{}, hs);
p.add_return({hs, last_hs});
return p;
}
};
struct test_gru_forward_default_actv1 : verify_program<test_gru_forward_default_actv1>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 6 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
p.add_instruction(
migraphx::op::gru{
hidden_size, {migraphx::op::sigmoid{}}, migraphx::op::rnn_direction::forward, clip},
seq,
w,
r,
bias,
und,
ih);
return p;
}
};
struct test_gru_reverse_last : verify_program<test_gru_reverse_last>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 6 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
auto output =
p.add_instruction(migraphx::op::gru{hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::reverse,
clip},
seq,
w,
r,
bias,
und,
ih);
p.add_instruction(migraphx::op::rnn_last_hs_output{}, output);
return p;
}
};
struct test_gru_reverse_3args : verify_program<test_gru_reverse_3args>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
p.add_instruction(migraphx::op::gru{hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::reverse,
clip},
seq,
w,
r);
return p;
}
};
struct test_gru_bidirct : verify_program<test_gru_bidirct>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 2;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 6 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
auto hs =
p.add_instruction(migraphx::op::gru{hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::bidirectional,
clip},
seq,
w,
r,
bias,
und,
ih);
auto lho = p.add_instruction(migraphx::op::rnn_last_hs_output{}, hs);
p.add_return({hs, lho});
return p;
}
};
struct test_var_sl_gru_bidirct : verify_program<test_var_sl_gru_bidirct>
{
migraphx::program create_program() const
{
std::size_t batch_size = 3;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 2;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 6 * hidden_size}};
migraphx::shape sl_shape{migraphx::shape::int32_type, {batch_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
std::vector<int> sl_data{2, 1, 3};
auto sql = p.add_literal(migraphx::literal{sl_shape, sl_data});
auto hs =
p.add_instruction(migraphx::op::gru{hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::bidirectional,
clip},
seq,
w,
r,
bias,
sql,
ih);
auto lho = p.add_instruction(migraphx::op::rnn_last_hs_output{}, hs);
p.add_return({hs, lho});
return p;
}
};
struct test_gru_bidirct_3args_und : verify_program<test_gru_bidirct_3args_und>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 2;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
p.add_instruction(migraphx::op::gru{hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::bidirectional,
clip},
seq,
w,
r,
und,
und,
und);
return p;
}
};
struct test_gru_bidirct_3args : verify_program<test_gru_bidirct_3args>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 2;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
p.add_instruction(migraphx::op::gru{hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::bidirectional,
clip},
seq,
w,
r);
return p;
}
};
struct test_gru_bidirct_seq1 : verify_program<test_gru_bidirct_seq1>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 1;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 2;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
p.add_instruction(migraphx::op::gru{hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::bidirectional,
clip},
seq,
w,
r);
return p;
}
};
struct test_gru_bidirct_default_actv : verify_program<test_gru_bidirct_default_actv>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 1;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 2;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
p.add_instruction(
migraphx::op::gru{hidden_size, {}, migraphx::op::rnn_direction::bidirectional, clip},
seq,
w,
r);
return p;
}
};
struct test_gru_bidirct_default_actv1 : verify_program<test_gru_bidirct_default_actv1>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 2;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 3 * hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 6 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
p.add_instruction(migraphx::op::gru{hidden_size,
{migraphx::op::sigmoid{}},
migraphx::op::rnn_direction::bidirectional,
clip},
seq,
w,
r,
bias,
und,
ih);
return p;
}
};
struct test_lstm_forward_last : verify_program<test_lstm_forward_last>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 8 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
migraphx::shape l_shape{migraphx::shape::int32_type, {batch_size}};
migraphx::shape ic_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
migraphx::shape pph_shape{migraphx::shape::float_type, {num_dirct, 3 * hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
auto len = p.add_literal(migraphx::literal(l_shape, {1, 2}));
auto ic = p.add_parameter("ic", ic_shape);
auto pph = p.add_parameter("pph", pph_shape);
auto output = p.add_instruction(
migraphx::op::lstm{
hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::forward,
clip},
seq,
w,
r,
bias,
len,
ih,
ic,
pph);
p.add_instruction(migraphx::op::rnn_last_hs_output{}, output, len);
return p;
}
};
struct test_lstm_forward_hs : verify_program<test_lstm_forward_hs>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 8 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
migraphx::shape ic_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
migraphx::shape pph_shape{migraphx::shape::float_type, {num_dirct, 3 * hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
auto ic = p.add_parameter("ic", ic_shape);
auto pph = p.add_parameter("pph", pph_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
p.add_instruction(
migraphx::op::lstm{
hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::forward,
clip},
seq,
w,
r,
bias,
und,
ih,
ic,
pph);
return p;
}
};
struct test_lstm_forward_3args_und : verify_program<test_lstm_forward_3args_und>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
p.add_instruction(
migraphx::op::lstm{
hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::forward,
clip},
seq,
w,
r,
und,
und,
und,
und,
und);
return p;
}
};
struct test_lstm_forward_3args : verify_program<test_lstm_forward_3args>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
p.add_instruction(
migraphx::op::lstm{
hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::forward,
clip},
seq,
w,
r);
return p;
}
};
struct test_lstm_two_outputs : verify_program<test_lstm_two_outputs>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto hs = p.add_instruction(
migraphx::op::lstm{
hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::forward,
clip},
seq,
w,
r);
auto last_hs = p.add_instruction(migraphx::op::rnn_last_hs_output{}, hs);
p.add_return({hs, last_hs});
return p;
}
};
struct test_lstm_three_outputs : verify_program<test_lstm_three_outputs>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto hs = p.add_instruction(
migraphx::op::lstm{
hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::forward,
clip},
seq,
w,
r);
auto last_hs = p.add_instruction(migraphx::op::rnn_last_hs_output{}, hs);
auto last_cell = p.add_instruction(migraphx::op::rnn_last_cell_output{}, hs);
p.add_return({hs, last_hs, last_cell});
return p;
}
};
struct test_lstm_forward_seq1 : verify_program<test_lstm_forward_seq1>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 1;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
p.add_instruction(
migraphx::op::lstm{
hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::forward,
clip},
seq,
w,
r);
return p;
}
};
struct test_lstm_forward_default_actv : verify_program<test_lstm_forward_default_actv>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 1;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
p.add_instruction(
migraphx::op::lstm{hidden_size, {}, migraphx::op::rnn_direction::forward, clip},
seq,
w,
r);
return p;
}
};
struct test_lstm_forward_default_actv1 : verify_program<test_lstm_forward_default_actv1>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 8 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
p.add_instruction(
migraphx::op::lstm{
hidden_size, {migraphx::op::sigmoid{}}, migraphx::op::rnn_direction::forward, clip},
seq,
w,
r,
bias,
und,
ih);
return p;
}
};
struct test_lstm_reverse_last : verify_program<test_lstm_reverse_last>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 8 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
migraphx::shape ic_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
migraphx::shape pph_shape{migraphx::shape::float_type, {num_dirct, 3 * hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
auto ic = p.add_parameter("ic", ic_shape);
auto pph = p.add_parameter("pph", pph_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
auto output = p.add_instruction(
migraphx::op::lstm{
hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::reverse,
clip},
seq,
w,
r,
bias,
und,
ih,
ic,
pph);
p.add_instruction(migraphx::op::rnn_last_hs_output{}, output);
return p;
}
};
struct test_lstm_reverse_3args : verify_program<test_lstm_reverse_3args>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
p.add_instruction(
migraphx::op::lstm{
hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::reverse,
clip},
seq,
w,
r);
return p;
}
};
struct test_lstm_reverse_3args_cell_output : verify_program<test_lstm_reverse_3args_cell_output>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 1;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto hs = p.add_instruction(
migraphx::op::lstm{
hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::reverse,
clip},
seq,
w,
r);
p.add_instruction(migraphx::op::rnn_last_cell_output{}, hs);
return p;
}
};
struct test_lstm_bidirct_last : verify_program<test_lstm_bidirct_last>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 2;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 8 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
migraphx::shape ic_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
migraphx::shape pph_shape{migraphx::shape::float_type, {num_dirct, 3 * hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
auto ic = p.add_parameter("ic", ic_shape);
auto pph = p.add_parameter("pph", pph_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
auto output = p.add_instruction(
migraphx::op::lstm{
hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::bidirectional,
clip},
seq,
w,
r,
bias,
und,
ih,
ic,
pph);
p.add_instruction(migraphx::op::rnn_last_hs_output{}, output);
return p;
}
};
struct test_lstm_bidirct_hs : verify_program<test_lstm_bidirct_hs>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 2;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 8 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
migraphx::shape sl_shape{migraphx::shape::int32_type, {batch_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
std::vector<int> sl_data{3, 2};
auto sql = p.add_literal(migraphx::literal{migraphx::literal{sl_shape, sl_data}});
p.add_instruction(migraphx::op::lstm{hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::bidirectional,
clip},
seq,
w,
r,
bias,
sql,
ih);
return p;
}
};
struct test_lstm_bidirct_3args_und : verify_program<test_lstm_bidirct_3args_und>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 2;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
p.add_instruction(
migraphx::op::gru{hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::bidirectional,
clip},
seq,
w,
r,
und,
und,
und,
und,
und);
return p;
}
};
struct test_lstm_bidirct_3args : verify_program<test_lstm_bidirct_3args>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 2;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
p.add_instruction(migraphx::op::lstm{hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::bidirectional,
clip},
seq,
w,
r);
return p;
}
};
struct test_lstm_bidirct_seq1 : verify_program<test_lstm_bidirct_seq1>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 1;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 2;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
p.add_instruction(migraphx::op::lstm{hidden_size,
{migraphx::op::sigmoid{}, migraphx::op::tanh{}},
migraphx::op::rnn_direction::bidirectional,
clip},
seq,
w,
r);
return p;
}
};
struct test_lstm_bidirct_default_actv : verify_program<test_lstm_bidirct_default_actv>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 1;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 2;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
p.add_instruction(
migraphx::op::lstm{hidden_size, {}, migraphx::op::rnn_direction::bidirectional, clip},
seq,
w,
r);
return p;
}
};
struct test_lstm_bidirct_default_actv1 : verify_program<test_lstm_bidirct_default_actv1>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 2;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 8 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
migraphx::shape sl_shape{migraphx::shape::int32_type, {batch_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
std::vector<int> sl_data(batch_size, 2);
auto sql = p.add_literal(migraphx::literal{sl_shape, sl_data});
p.add_instruction(migraphx::op::lstm{hidden_size,
{migraphx::op::sigmoid{}},
migraphx::op::rnn_direction::bidirectional,
clip},
seq,
w,
r,
bias,
sql,
ih);
return p;
}
};
struct test_lstm_bidirct_default_actv2 : verify_program<test_lstm_bidirct_default_actv2>
{
migraphx::program create_program() const
{
std::size_t batch_size = 2;
std::size_t seq_len = 3;
std::size_t hidden_size = 5;
std::size_t input_size = 8;
std::size_t num_dirct = 2;
float clip = 0.0f;
migraphx::program p;
migraphx::shape in_shape{migraphx::shape::float_type, {seq_len, batch_size, input_size}};
migraphx::shape w_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, input_size}};
migraphx::shape r_shape{migraphx::shape::float_type,
{num_dirct, 4 * hidden_size, hidden_size}};
migraphx::shape b_shape{migraphx::shape::float_type, {num_dirct, 8 * hidden_size}};
migraphx::shape ih_shape{migraphx::shape::float_type, {num_dirct, batch_size, hidden_size}};
auto seq = p.add_parameter("seq", in_shape);
auto w = p.add_parameter("w", w_shape);
auto r = p.add_parameter("r", r_shape);
auto bias = p.add_parameter("bias", b_shape);
auto ih = p.add_parameter("ih", ih_shape);
auto und = p.add_instruction(migraphx::op::undefined{});
p.add_instruction(migraphx::op::lstm{hidden_size,
{migraphx::op::tanh{}, migraphx::op::sigmoid{}},
migraphx::op::rnn_direction::bidirectional,
clip},
seq,
w,
r,
bias,
und,
ih);
return p;
}
};
int main(int argc, const char* argv[]) { test::run(argc, argv); }
test/gpu/test_utils.hpp deleted 100644 → 0
View file @ ba33d25c
#ifndef MIGRAPHX_GUARD_TEST_GPU_TEST_UTILS_HPP
#define MIGRAPHX_GUARD_TEST_GPU_TEST_UTILS_HPP
#include <migraphx/env.hpp>
#include <migraphx/program.hpp>
#include <migraphx/operators.hpp>
#include <migraphx/generate.hpp>
#include <migraphx/ranges.hpp>
#include <migraphx/cpu/target.hpp>
#include <migraphx/gpu/target.hpp>
#include <migraphx/gpu/analyze_streams.hpp>
#include <migraphx/gpu/miopen.hpp>
#include <migraphx/gpu/hip.hpp>
#include <migraphx/manage_ptr.hpp>
#include <migraphx/type_name.hpp>
#include <migraphx/verify_args.hpp>
#include <migraphx/instruction.hpp>
#include <miopen/miopen.h>
#include <future>
#include <thread>
#include <cmath>
#include <numeric>
#include <test.hpp>
#ifdef __clang__
#pragma clang diagnostic push
#pragma clang diagnostic ignored "-Wglobal-constructors"
#endif
MIGRAPHX_DECLARE_ENV_VAR(MIGRAPHX_TRACE_GPU_COMPILE)
static std::array<std::function<void()>, 2>& handlers()
{
// NOLINTNEXTLINE
static std::array<std::function<void()>, 2> r = {};
return r;
};
// An improved async, that doesn't block
template <class Function>
std::future<typename std::result_of<Function()>::type> detach_async(Function&& f,
bool parallel = true)
{
if(parallel)
{
using result_type = typename std::result_of<Function()>::type;
std::packaged_task<result_type()> task(std::forward<Function>(f));
auto fut = task.get_future();
std::thread(std::move(task)).detach();
return std::move(fut);
}
return std::async(std::launch::deferred, std::forward<Function>(f));
}
struct auto_print
{
static void set_terminate_handler(const std::string& name)
{
// NOLINTNEXTLINE
static std::string pname;
pname = name;
std::set_terminate(+[] {
std::cout << "FAILED: " << pname << std::endl;
try
{
std::rethrow_exception(std::current_exception());
}
catch(const std::exception& e)
{
std::cout << " what(): " << e.what() << std::endl;
}
std::cout << std::endl;
auto hdlers = handlers();
for(auto&& handle : hdlers)
handle();
});
}
// static std::array<std::function<void()>, 2> handlers;
int index;
template <class T>
auto_print(T& x, int i) : index(i)
{
auto hdlers = handlers();
hdlers[index] = [&x] { std::cout << x << std::endl; };
}
~auto_print()
{
auto hdlers = handlers();
hdlers[index] = [] {};
}
};
template <class T>
auto get_hash(const T& x)
{
return std::hash<T>{}(x);
}
inline void compile_check(migraphx::program& p, const migraphx::target& t, bool show_trace = false)
{
auto name = t.name();
auto shapes = p.get_output_shapes();
std::stringstream ss;
migraphx::compile_options options;
options.trace = migraphx::tracer{ss};
p.compile(t, options);
if(shapes.size() != p.get_output_shapes().size())
{
std::cout << ss.str() << std::endl;
throw std::runtime_error("Compiling program with " + name +
" alters its number of outputs");
}
auto num = shapes.size();
for(std::size_t i = 0; i < num; ++i)
{
if(p.get_output_shapes()[i].lens() != shapes[i].lens())
{
std::cout << ss.str() << std::endl;
throw std::runtime_error("Compiling program with " + name + " alters its shape");
}
}
if(show_trace)
{
std::cout << ss.str() << std::endl;
}
}
inline void check_gpu_streams(const migraphx::program& p)
{
auto races = migraphx::gpu::analyze_streams(p);
for(auto&& race : races)
{
std::cout << "FAILED: " << std::endl;
std::cout << "Race condition detected for: ";
p.debug_print(race.ins);
std::cout << "Should happen after: ";
p.debug_print(race.before);
}
}
template <class V>
std::vector<migraphx::argument> run_cpu(migraphx::program& p)
{
V v;
p = v.create_program();
auto_print pp{p, 0};
compile_check(p, migraphx::cpu::target{});
migraphx::program::parameter_map m;
for(auto&& x : p.get_parameter_shapes())
{
m[x.first] = migraphx::generate_argument(x.second, get_hash(x.first));
}
return p.eval(m);
}
template <class V>
std::vector<migraphx::argument> run_gpu(migraphx::program& p)
{
V v;
p = v.create_program();
auto_print pp{p, 1};
compile_check(p, migraphx::gpu::target{}, migraphx::enabled(MIGRAPHX_TRACE_GPU_COMPILE{}));
check_gpu_streams(p);
migraphx::program::parameter_map m;
for(auto&& x : p.get_parameter_shapes())
{
m[x.first] =
migraphx::gpu::to_gpu(migraphx::generate_argument(x.second, get_hash(x.first)));
}
// Program should have an output parameter
EXPECT(std::any_of(
m.begin(), m.end(), [](auto& x) { return migraphx::contains(x.first, "output"); }));
// Ensure the program doesn't modify the context in a dry run
auto ctx = p.get_context();
assert(&ctx != &p.get_context());
EXPECT(is_shared(ctx, p.get_context()));
p.dry_run(m);
EXPECT(is_shared(ctx, p.get_context()));
p.eval(m);
auto gpu_res = p.eval(m);
std::vector<migraphx::argument> res(gpu_res.size());
std::transform(gpu_res.begin(), gpu_res.end(), res.begin(), [&](auto& argu) {
return migraphx::gpu::from_gpu(argu);
});
return res;
}
template <class V>
void run_verify_program()
{
auto_print::set_terminate_handler(migraphx::get_type_name<V>());
// std::cout << migraphx::get_type_name<V>() << std::endl;
migraphx::program cpu_prog;
migraphx::program gpu_prog;
auto cpu_arg_f = detach_async([&] { return run_cpu<V>(cpu_prog); });
auto gpu_arg = run_gpu<V>(gpu_prog);
auto cpu_arg = cpu_arg_f.get();
bool passed = true;
passed &= (cpu_arg.size() == gpu_arg.size());
std::size_t num = cpu_arg.size();
for(std::size_t i = 0; ((i < num) and passed); ++i)
{
passed &= verify_args(migraphx::get_type_name<V>(), cpu_arg[i], gpu_arg[i]);
}
if(not passed)
{
V v;
auto p = v.create_program();
std::cout << p << std::endl;
std::cout << "cpu:\n" << cpu_prog << std::endl;
std::cout << "gpu:\n" << gpu_prog << std::endl;
std::cout << std::endl;
}
std::set_terminate(nullptr);
}
template <class T>
int auto_register_verify_program()
{
test::add_test_case(migraphx::get_type_name<T>(), [] { run_verify_program<T>(); });
return 0;
}
template <class T>
struct verify_program
{
static const int static_register;
// This typedef ensures that the static member will be instantiated if
// the class itself is instantiated
using static_register_type =
std::integral_constant<decltype(&static_register), &static_register>;
};
template <class T>
const int verify_program<T>::static_register = auto_register_verify_program<T>(); // NOLINT
#endif
test/onnx/verify_onnx.cpp
View file @ 500d9441
...@@ -3,7 +3,7 @@ ...@@ -3,7 +3,7 @@
#include <migraphx/literal.hpp> #include <migraphx/literal.hpp>
#include <migraphx/operators.hpp> #include <migraphx/operators.hpp>
#include <migraphx/program.hpp> #include <migraphx/program.hpp>
#include <migraphx/cpu/target.hpp> #include <migraphx/ref/target.hpp>
#include <migraphx/pass_manager.hpp> #include <migraphx/pass_manager.hpp>
#include <migraphx/verify.hpp> #include <migraphx/verify.hpp>
#include <migraphx/onnx.hpp> #include <migraphx/onnx.hpp>
...@@ -12,7 +12,7 @@ ...@@ -12,7 +12,7 @@
TEST_CASE(averagepool_notset_test) TEST_CASE(averagepool_notset_test)
{ {
auto p = migraphx::parse_onnx("averagepool_notset_test.onnx"); auto p = migraphx::parse_onnx("averagepool_notset_test.onnx");
p.compile(migraphx::cpu::target{}); p.compile(migraphx::ref::target{});
std::vector<float> data_x = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, std::vector<float> data_x = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12,
13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24}; 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24};
migraphx::shape s_x{migraphx::shape::float_type, {1, 1, 5, 5}}; migraphx::shape s_x{migraphx::shape::float_type, {1, 1, 5, 5}};
...@@ -30,7 +30,7 @@ TEST_CASE(averagepool_notset_test) ...@@ -30,7 +30,7 @@ TEST_CASE(averagepool_notset_test)
TEST_CASE(averagepool_nt_cip_test) TEST_CASE(averagepool_nt_cip_test)
{ {
auto p = migraphx::parse_onnx("averagepool_nt_cip_test.onnx"); auto p = migraphx::parse_onnx("averagepool_nt_cip_test.onnx");
p.compile(migraphx::cpu::target{}); p.compile(migraphx::ref::target{});
std::vector<float> data_x = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, std::vector<float> data_x = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12,
13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24}; 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24};
migraphx::shape s_x{migraphx::shape::float_type, {1, 1, 5, 5}}; migraphx::shape s_x{migraphx::shape::float_type, {1, 1, 5, 5}};
...@@ -48,7 +48,7 @@ TEST_CASE(averagepool_nt_cip_test) ...@@ -48,7 +48,7 @@ TEST_CASE(averagepool_nt_cip_test)
TEST_CASE(gather_elements) TEST_CASE(gather_elements)
{ {
migraphx::program p = migraphx::parse_onnx("gather_elements_axis0_test.onnx"); migraphx::program p = migraphx::parse_onnx("gather_elements_axis0_test.onnx");
p.compile(migraphx::cpu::target{}); p.compile(migraphx::ref::target{});
migraphx::shape s_data{migraphx::shape::float_type, {3, 4}}; migraphx::shape s_data{migraphx::shape::float_type, {3, 4}};
std::vector<float> data = { std::vector<float> data = {
0.25, 0.75, 0.9375, 0.4375, 0.6875, 0.5625, -0.875, 0.1875, -0.125, 0.5, -0.9375, -0.0625}; 0.25, 0.75, 0.9375, 0.4375, 0.6875, 0.5625, -0.875, 0.1875, -0.125, 0.5, -0.9375, -0.0625};
...@@ -72,7 +72,7 @@ TEST_CASE(instance_norm_test) ...@@ -72,7 +72,7 @@ TEST_CASE(instance_norm_test)
{ {
migraphx::program p = migraphx::parse_onnx("instance_norm_val_test.onnx"); migraphx::program p = migraphx::parse_onnx("instance_norm_val_test.onnx");
p.compile(migraphx::cpu::target{}); p.compile(migraphx::ref::target{});
auto result = p.eval({}).back(); auto result = p.eval({}).back();
std::vector<float> result_vector(9); std::vector<float> result_vector(9);
result.visit([&](auto output) { result_vector.assign(output.begin(), output.end()); }); result.visit([&](auto output) { result_vector.assign(output.begin(), output.end()); });
...@@ -102,7 +102,7 @@ TEST_CASE(instance_norm_3d_test) ...@@ -102,7 +102,7 @@ TEST_CASE(instance_norm_3d_test)
{ {
migraphx::program p = migraphx::parse_onnx("instance_norm_val_3d_test.onnx"); migraphx::program p = migraphx::parse_onnx("instance_norm_val_3d_test.onnx");
p.compile(migraphx::cpu::target{}); p.compile(migraphx::ref::target{});
auto result = p.eval({}).back(); auto result = p.eval({}).back();
std::vector<float> result_vector(16); std::vector<float> result_vector(16);
result.visit([&](auto output) { result_vector.assign(output.begin(), output.end()); }); result.visit([&](auto output) { result_vector.assign(output.begin(), output.end()); });
...@@ -130,7 +130,7 @@ TEST_CASE(instance_norm_3d_test) ...@@ -130,7 +130,7 @@ TEST_CASE(instance_norm_3d_test)
TEST_CASE(resize_test) TEST_CASE(resize_test)
{ {
migraphx::program p = migraphx::parse_onnx("resize_upsample_pf_test.onnx"); migraphx::program p = migraphx::parse_onnx("resize_upsample_pf_test.onnx");
p.compile(migraphx::cpu::target{}); p.compile(migraphx::ref::target{});
migraphx::shape sx{migraphx::shape::float_type, {1, 1, 2, 2}}; migraphx::shape sx{migraphx::shape::float_type, {1, 1, 2, 2}};
std::vector<float> dx = {1.0f, 2.0f, 3.0f, 4.0f}; std::vector<float> dx = {1.0f, 2.0f, 3.0f, 4.0f};
...@@ -151,7 +151,7 @@ TEST_CASE(resize_test) ...@@ -151,7 +151,7 @@ TEST_CASE(resize_test)
TEST_CASE(selu_test) TEST_CASE(selu_test)
{ {
migraphx::program p = migraphx::parse_onnx("selu_test.onnx"); migraphx::program p = migraphx::parse_onnx("selu_test.onnx");
p.compile(migraphx::cpu::target{}); p.compile(migraphx::ref::target{});
migraphx::shape xs{migraphx::shape::double_type, {2, 3}}; migraphx::shape xs{migraphx::shape::double_type, {2, 3}};
std::vector<double> x_data = {1.1, 2.1, 0.0, -1.3, -5.3, 12.0}; std::vector<double> x_data = {1.1, 2.1, 0.0, -1.3, -5.3, 12.0};
...@@ -190,7 +190,7 @@ TEST_CASE(upsample_test) ...@@ -190,7 +190,7 @@ TEST_CASE(upsample_test)
TEST_CASE(where_test) TEST_CASE(where_test)
{ {
migraphx::program p = migraphx::parse_onnx("where_test.onnx"); migraphx::program p = migraphx::parse_onnx("where_test.onnx");
p.compile(migraphx::cpu::target{}); p.compile(migraphx::ref::target{});
migraphx::shape c_shape{migraphx::shape::bool_type, {2}}; migraphx::shape c_shape{migraphx::shape::bool_type, {2}};
std::vector<int8_t> c_data = {1, 0}; std::vector<int8_t> c_data = {1, 0};
......
test/perf_report.cpp
View file @ 500d9441
#include <migraphx/program.hpp> #include <migraphx/program.hpp>
#include <migraphx/op/add.hpp> #include <migraphx/op/add.hpp>
#include <migraphx/cpu/target.hpp> #include <migraphx/ref/target.hpp>
#include <migraphx/ranges.hpp> #include <migraphx/ranges.hpp>
#include "test.hpp" #include "test.hpp"
...@@ -12,7 +12,7 @@ TEST_CASE(perf_report) ...@@ -12,7 +12,7 @@ TEST_CASE(perf_report)
auto one = p.add_literal(1); auto one = p.add_literal(1);
auto two = p.add_literal(2); auto two = p.add_literal(2);
p.add_instruction(migraphx::op::add{}, one, two); p.add_instruction(migraphx::op::add{}, one, two);
p.compile(migraphx::cpu::target{}); p.compile(migraphx::ref::target{});
p.perf_report(ss, 2, {}); p.perf_report(ss, 2, {});
std::string output = ss.str(); std::string output = ss.str();
......
test/program_test.cpp
View file @ 500d9441
...@@ -5,7 +5,7 @@ ...@@ -5,7 +5,7 @@
#include <migraphx/op/add.hpp> #include <migraphx/op/add.hpp>
#include <migraphx/op/dot.hpp> #include <migraphx/op/dot.hpp>
#include <migraphx/op/mul.hpp> #include <migraphx/op/mul.hpp>
#include <migraphx/cpu/target.hpp> #include <migraphx/ref/target.hpp>
#include <sstream> #include <sstream>
#include "test.hpp" #include "test.hpp"
#include <basic_ops.hpp> #include <basic_ops.hpp>
...@@ -77,10 +77,10 @@ TEST_CASE(program_copy) ...@@ -77,10 +77,10 @@ TEST_CASE(program_copy)
migraphx::program p2{}; migraphx::program p2{};
p2 = p1; p2 = p1;
p2.compile(migraphx::cpu::target{}); p2.compile(migraphx::ref::target{});
EXPECT(p1 != p2); EXPECT(p1 != p2);
p1.compile(migraphx::cpu::target{}); p1.compile(migraphx::ref::target{});
EXPECT(p1 == p2); EXPECT(p1 == p2);
EXPECT(p1.get_parameter_names() == p2.get_parameter_names()); EXPECT(p1.get_parameter_names() == p2.get_parameter_names());
...@@ -91,7 +91,7 @@ TEST_CASE(program_copy) ...@@ -91,7 +91,7 @@ TEST_CASE(program_copy)
auto p2(p1); auto p2(p1);
EXPECT(p1 == p2); EXPECT(p1 == p2);
p1.compile(migraphx::cpu::target{}); p1.compile(migraphx::ref::target{});
EXPECT(p1 != p2); EXPECT(p1 != p2);
p2 = p1; p2 = p1;
...@@ -106,8 +106,8 @@ TEST_CASE(program_copy) ...@@ -106,8 +106,8 @@ TEST_CASE(program_copy)
p2 = p1; p2 = p1;
EXPECT(p1 == p2); EXPECT(p1 == p2);
p1.compile(migraphx::cpu::target{}); p1.compile(migraphx::ref::target{});
p2.compile(migraphx::cpu::target{}); p2.compile(migraphx::ref::target{});
EXPECT(p1 == p2); EXPECT(p1 == p2);
} }
...@@ -126,8 +126,8 @@ TEST_CASE(program_copy) ...@@ -126,8 +126,8 @@ TEST_CASE(program_copy)
p2 = p1; p2 = p1;
EXPECT(p2 == p1); EXPECT(p2 == p1);
p1.compile(migraphx::cpu::target{}); p1.compile(migraphx::ref::target{});
p2.compile(migraphx::cpu::target{}); p2.compile(migraphx::ref::target{});
EXPECT(p2 == p1); EXPECT(p2 == p1);
} }
} }
......
test/py/CMakeLists.txt
View file @ 500d9441
...@@ -23,7 +23,7 @@ foreach(PYTHON_VERSION ${PYTHON_VERSIONS}) ...@@ -23,7 +23,7 @@ foreach(PYTHON_VERSION ${PYTHON_VERSIONS})
add_dependencies(check migraphx_py_${PYTHON_VERSION}) add_dependencies(check migraphx_py_${PYTHON_VERSION})
endforeach() endforeach()
add_py_test(cpu test_cpu.py WORKING_DIRECTORY ${TEST_ONNX_DIR}) add_py_test(ref test_cpu.py WORKING_DIRECTORY ${TEST_ONNX_DIR})
add_py_test(save_load test_save_load.py WORKING_DIRECTORY ${TEST_ONNX_DIR}) add_py_test(save_load test_save_load.py WORKING_DIRECTORY ${TEST_ONNX_DIR})
if(MIGRAPHX_ENABLE_GPU) if(MIGRAPHX_ENABLE_GPU)
add_py_test(gpu_offload test_gpu_offload.py WORKING_DIRECTORY ${TEST_ONNX_DIR}) add_py_test(gpu_offload test_gpu_offload.py WORKING_DIRECTORY ${TEST_ONNX_DIR})
......
test/py/test_cpu.py
View file @ 500d9441
...@@ -6,7 +6,7 @@ def test_conv_relu(): ...@@ -6,7 +6,7 @@ def test_conv_relu():
print(p) print(p)
s1 = p.get_output_shapes()[-1] s1 = p.get_output_shapes()[-1]
print("Compiling ...") print("Compiling ...")
p.compile(migraphx.get_target("cpu")) p.compile(migraphx.get_target("ref"))
print(p) print(p)
s2 = p.get_output_shapes()[-1] s2 = p.get_output_shapes()[-1]
assert s1 == s2 assert s1 == s2
...@@ -35,7 +35,7 @@ def test_add_scalar(): ...@@ -35,7 +35,7 @@ def test_add_scalar():
print(p) print(p)
s1 = p.get_output_shapes()[-1] s1 = p.get_output_shapes()[-1]
print("Compiling ...") print("Compiling ...")
p.compile(migraphx.get_target("cpu")) p.compile(migraphx.get_target("ref"))
print(p) print(p)
s2 = p.get_output_shapes()[-1] s2 = p.get_output_shapes()[-1]
assert s1 == s2 assert s1 == s2
......
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