Roialign (#952)

Implementation of the roialign operator. For now, we have only the ref implementation. When we run a model on the GPU, we fall back the execution to use the ref implementation.

src/CMakeLists.txt

@@ -155,6 +155,7 @@ register_migraphx_ops(
     rnn_last_cell_output
     rnn_last_hs_output
     rnn_var_sl_last_output
+    roialign
     round
     rsqrt
     scalar

src/eliminate_data_type.cpp

@@ -11,7 +11,7 @@ inline namespace MIGRAPHX_INLINE_NS {
 void eliminate_data_type::apply(module& m) const
 {
     static const std::vector<std::string> skip_op_names = {
-        "convert", "get_tuple_elem", "if", "loop"};
+        "convert", "get_tuple_elem", "if", "loop", "roialign"};
     for(auto ins : iterator_for(m))
     {
         if(ins->name()[0] == '@')

src/include/migraphx/op/roialign.hpp

+#ifndef MIGRAPHX_GUARD_OPERATORS_ROIALIGN_HPP
+#define MIGRAPHX_GUARD_OPERATORS_ROIALIGN_HPP
+
+#include <limits>
+#include <migraphx/check_shapes.hpp>
+#include <migraphx/config.hpp>
+#include <migraphx/argument.hpp>
+#include <migraphx/par_for.hpp>
+#include <migraphx/dfor.hpp>
+#include <migraphx/ranges.hpp>
+#include <migraphx/shape_for_each.hpp>
+#include <cmath>
+#include <numeric>
+#include <utility>
+#include <vector>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace op {
+
+struct roialign
+{
+    std::string coord_trans_mode = "half_pixel";
+    std::string mode             = "avg";
+    int64_t output_height        = 1;
+    int64_t output_width         = 1;
+    int64_t sampling_ratio       = 0;
+    float spatial_scale          = 1.0f;
+
+    template <class Self, class F>
+    static auto reflect(Self& self, F f)
+    {
+        return pack(f(self.coord_trans_mode, "coordinate_transformation_mode"),
+                    f(self.mode, "mode"),
+                    f(self.output_height, "output_height"),
+                    f(self.output_width, "output_width"),
+                    f(self.sampling_ratio, "sampling_ratio"),
+                    f(self.spatial_scale, "spatial_scale"));
+    }
+
+    std::string name() const { return "roialign"; }
+
+    shape compute_shape(std::vector<shape> inputs) const
+    {
+        check_shapes{inputs, *this}.has(3).standard();
+        auto x_lens   = inputs.at(0).lens();
+        auto roi_lens = inputs.at(1).lens();
+        auto bi_lens  = inputs.at(2).lens();
+        auto type     = inputs.at(0).type();
+
+        // check input correct
+        if(bi_lens.size() != 1)
+        {
+            MIGRAPHX_THROW("ROIALIGN: batch indices should be 1 dimension!");
+        }
+
+        if(roi_lens.size() != 2 or roi_lens.at(1) != 4)
+        {
+            MIGRAPHX_THROW(
+                "ROIALIGN: rois should be 2 dimensions, and the second dim should be 4!");
+        }
+
+        if(roi_lens.front() != bi_lens.front())
+        {
+            MIGRAPHX_THROW("ROIALIGN: rois and batch indices inputs should have the same number!");
+        }
+
+        std::vector<std::size_t> out_lens = x_lens;
+        out_lens[0]                       = roi_lens[0];
+        out_lens[2]                       = output_height;
+        out_lens[3]                       = output_width;
+
+        return {type, out_lens};
+    }
+
+    struct pos_weight
+    {
+        // neighbor indices for the bilinear interpolation
+        std::array<std::size_t, 4> pos = {0, 0, 0, 0};
+        // neighbor weights for the bilinear interpolation
+        std::array<float, 4> w = {0.0f, 0.0f, 0.0f, 0.0f};
+    };
+
+    auto calc_pos_weight(const std::array<std::size_t, 2>& dims,
+                         const shape& comp_s,
+                         const std::array<float, 2>& roi_start,
+                         const std::array<float, 2>& bin_size,
+                         const std::array<std::size_t, 2>& bin_grid_size) const
+    {
+        std::vector<pos_weight> results(bin_grid_size[0] * bin_grid_size[1] * output_height *
+                                        output_width);
+        shape_for_each(comp_s, [&](auto idx) {
+            std::array<std::size_t, 2> p = {idx[0], idx[1]};
+            std::array<std::size_t, 2> i = {idx[2], idx[3]};
+            auto index                   = comp_s.index(idx);
+
+            std::array<float, 2> xy{};
+            std::array<int64_t, 2> low{};
+            std::array<int64_t, 2> high{};
+            for(auto ii : range(p.size()))
+            {
+                xy[ii] = roi_start[ii] + p[ii] * bin_size[ii] +
+                         (i[ii] + .5f) * bin_size[ii] / bin_grid_size[ii];
+                xy[ii] = (coord_trans_mode == "output_half_pixel") ? (xy[ii] - 0.5f) : xy[ii];
+                if(xy[ii] < -1.0 or xy[ii] > dims[ii])
+                {
+                    results[index] = pos_weight{};
+                    return;
+                }
+
+                xy[ii]   = std::max(xy[ii], 0.0f);
+                low[ii]  = xy[ii];
+                high[ii] = low[ii] + 1;
+                if(low[ii] >= dims[ii] - 1)
+                {
+                    xy[ii] = high[ii] = low[ii] = dims[ii] - 1;
+                }
+            }
+
+            results[index].pos = {low[0] * dims[1] + low[1],
+                                  low[0] * dims[1] + high[1],
+                                  high[0] * dims[1] + low[1],
+                                  high[0] * dims[1] + high[1]};
+
+            float ly = xy[0] - low[0];
+            float lx = xy[1] - low[1];
+            float hy = 1.0f - ly;
+            float hx = 1.0f - lx;
+
+            // save weights and indeces
+            results[index].w = {hy * hx, hy * lx, ly * hx, ly * lx};
+        });
+
+        return results;
+    }
+
+    struct max_pool
+    {
+        double init() { return std::numeric_limits<double>::lowest(); }
+
+        double operator()(double x, double y) { return std::max(x, y); }
+
+        double final(double x, std::size_t) { return (x); }
+    };
+
+    struct avg_pool
+    {
+        double init() { return 0.0; }
+
+        double operator()(double x, double y) { return x + y; }
+
+        double final(double x, std::size_t y) { return (y == 0) ? 0.0 : (x / y); }
+    };
+
+    template <class T, class Op>
+    std::tuple<double, int64_t> calc_pooling(const T& data,
+                                             const std::array<std::size_t, 2>& bin_grid_size,
+                                             const std::vector<pos_weight>& pos_weights,
+                                             int64_t index,
+                                             Op op) const
+    {
+        double output_val   = op.init();
+        const int64_t count = bin_grid_size[0] * bin_grid_size[1];
+        dfor(bin_grid_size[0], bin_grid_size[1])([&](auto, auto) {
+            const auto& pc = pos_weights[index];
+            std::array<double, 4> wv;
+            std::transform(
+                pc.w.begin(), pc.w.end(), pc.pos.begin(), wv.begin(), [&](auto w, auto pos) {
+                    return *(data + pos) * w;
+                });
+            output_val = std::accumulate(wv.begin(), wv.end(), output_val, op);
+            index += 1;
+        });
+
+        output_val = op.final(output_val, count);
+
+        return {output_val, index};
+    }
+
+    argument compute(const shape& output_shape, std::vector<argument> args) const
+    {
+        argument result{output_shape};
+        const auto& out_lens = output_shape.lens();
+        int64_t n_rois       = out_lens[0];
+        std::size_t channels = out_lens[1];
+        // output dims of height and width, in all 2-dim arrays, the first dim
+        // is for height and second dim is for width
+        std::array<std::size_t, 2> out_dims = {out_lens[2], out_lens[3]};
+        const auto& x_lens                  = args.at(0).get_shape().lens();
+        // input dims of height and width
+        std::array<std::size_t, 2> in_dims = {x_lens[2], x_lens[3]};
+        auto roi_s                         = args.at(1).get_shape();
+
+        visit_all(result, args.at(0), args.at(1))([&](auto output, auto x, auto roi) {
+            const auto* batch_indices = args.at(2).cast<int64_t>();
+            par_for(n_rois, [&](auto n) {
+                const auto bottom_data   = x.begin();
+                const auto roi_batch_ind = batch_indices[n];
+                // Do not using rounding; this implementation detail is critical
+                std::array<float, 2> roi_starts = {
+                    static_cast<float>(roi[roi_s.index({n, 1})] * spatial_scale),
+                    static_cast<float>(roi[roi_s.index({n, 0})] * spatial_scale)};
+                std::array<float, 2> roi_ends = {
+                    static_cast<float>(roi[roi_s.index({n, 3})] * spatial_scale),
+                    static_cast<float>(roi[roi_s.index({n, 2})] * spatial_scale)};
+
+                // Force malformed ROIs to be 1x1
+                std::array<float, 2> roi_size{};
+                std::array<float, 2> bin_size{};
+                std::array<std::size_t, 2> bin_grid_size{};
+
+                for(auto ii : range(roi_size.size()))
+                {
+                    roi_size[ii] = roi_ends[ii] - roi_starts[ii];
+                    roi_size[ii] = std::max(roi_size[ii], 1.0f);
+
+                    bin_size[ii]      = roi_size[ii] / out_dims[ii];
+                    bin_grid_size[ii] = (sampling_ratio > 0)
+                                            ? sampling_ratio
+                                            : std::ceil(roi_size[ii] / out_dims[ii]);
+                }
+
+                // we want to precalculate indices and weights shared by all channels,
+                // this is the key point of optimization
+                std::vector<std::size_t> comp_lens = {
+                    out_dims[0], out_dims[1], bin_grid_size[0], bin_grid_size[1]};
+                shape comp_s{shape::float_type, comp_lens};
+                auto pre_calc =
+                    this->calc_pos_weight(in_dims, comp_s, roi_starts, bin_size, bin_grid_size);
+
+                std::vector<std::size_t> comp_lens1 = {channels, out_dims[0], out_dims[1]};
+                shape comp_s1{migraphx::shape::float_type, comp_lens1};
+                std::vector<int64_t> vec_index(channels, 0);
+                shape_for_each(comp_s1, [&](auto idx) {
+                    auto c  = idx[0];
+                    auto ph = idx[1];
+                    auto pw = idx[2];
+
+                    const auto offset_bottom_data =
+                        bottom_data + static_cast<int64_t>((roi_batch_ind * channels + c) *
+                                                           in_dims[0] * in_dims[1]);
+                    double output_val;
+                    std::tie(output_val, vec_index[c]) =
+                        (mode == "avg") ? this->calc_pooling(offset_bottom_data,
+                                                             bin_grid_size,
+                                                             pre_calc,
+                                                             vec_index[c],
+                                                             avg_pool{})
+                                        : this->calc_pooling(offset_bottom_data,
+                                                             bin_grid_size,
+                                                             pre_calc,
+                                                             vec_index[c],
+                                                             max_pool{});
+                    output(n, c, ph, pw) = output_val;
+                });
+
+            });
+        });
+
+        return result;
+    }
+};
+
+} // namespace op
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

src/include/migraphx/operators.hpp

@@ -80,6 +80,7 @@
 #include <migraphx/op/rnn_last_hs_output.hpp>
 #include <migraphx/op/rnn_variable_seq_lens.hpp>
 #include <migraphx/op/rnn_var_sl_last_output.hpp>
+#include <migraphx/op/roialign.hpp>
 #include <migraphx/op/round.hpp>
 #include <migraphx/op/rsqrt.hpp>
 #include <migraphx/op/scalar.hpp>

src/onnx/parse_roialign.cpp

+#include <migraphx/onnx/op_parser.hpp>
+#include <migraphx/onnx/checks.hpp>
+#include <migraphx/ranges.hpp>
+#include <migraphx/instruction.hpp>
+#include <migraphx/make_op.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace onnx {
+
+struct parse_roialign : op_parser<parse_roialign>
+{
+    std::vector<op_desc> operators() const { return {{"RoiAlign"}}; }
+
+    instruction_ref parse(const op_desc& /*opd*/,
+                          const onnx_parser& /*parser*/,
+                          onnx_parser::node_info info,
+                          const std::vector<instruction_ref>& args) const
+    {
+        std::string coord_trans_mode = "half_pixel";
+        if(contains(info.attributes, "coordinate_transformation_mode"))
+        {
+            coord_trans_mode = info.attributes.at("coordinate_transformation_mode").s();
+        }
+        if(not contains({"half_pixel", "output_half_pixel"}, coord_trans_mode))
+        {
+            MIGRAPHX_THROW("coordinate_transformation_mode \"" + coord_trans_mode +
+                           "\": invalid value!");
+        }
+
+        std::string mode = "avg";
+        if(contains(info.attributes, "mode"))
+        {
+            mode = info.attributes.at("mode").s();
+        }
+
+        int64_t output_height = 1;
+        if(contains(info.attributes, "output_height"))
+        {
+            output_height = info.attributes.at("output_height").i();
+        }
+
+        int64_t output_width = 1;
+        if(contains(info.attributes, "output_width"))
+        {
+            output_width = info.attributes.at("output_width").i();
+        }
+
+        int64_t sampling_ratio = 0;
+        if(contains(info.attributes, "sampling_ratio"))
+        {
+            sampling_ratio = info.attributes.at("sampling_ratio").i();
+        }
+
+        float spatial_scale = 1.0f;
+        if(contains(info.attributes, "spatial_scale"))
+        {
+            spatial_scale = info.attributes.at("spatial_scale").f();
+        }
+
+        return info.add_instruction(make_op("roialign",
+                                            {{"coordinate_transformation_mode", coord_trans_mode},
+                                             {"mode", mode},
+                                             {"output_height", output_height},
+                                             {"output_width", output_width},
+                                             {"sampling_ratio", sampling_ratio},
+                                             {"spatial_scale", spatial_scale}}),
+                                    args);
+    }
+};
+
+} // namespace onnx
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx

src/targets/gpu/lowering.cpp

@@ -191,6 +191,7 @@ struct miopen_apply
         add_loop_op();
         add_neg_op();
         add_quant_convolution_op();
+        add_roialign();
     }
 
     void copy_params()
@@ -469,6 +470,26 @@ struct miopen_apply
         });
     }
 
+    void add_roialign()
+    {
+        apply_map.emplace("roialign", [=](instruction_ref ins) {
+            auto s      = ins->get_shape();
+            auto output = insert_allocation(ins, s);
+            std::vector<instruction_ref> cpu_inputs;
+            auto inputs = ins->inputs();
+            std::transform(
+                inputs.begin(), inputs.end(), std::back_inserter(cpu_inputs), [&](auto in) {
+                    return mod->insert_instruction(ins, make_op("hip::copy_from_gpu"), in);
+                });
+            cpu_inputs.front() =
+                mod->insert_instruction(ins, make_op("hip::sync_stream"), cpu_inputs);
+            auto cpu_out = mod->insert_instruction(ins, ins->get_operator(), cpu_inputs);
+            auto gpu_out =
+                mod->insert_instruction(ins, make_op("hip::copy_to_gpu"), cpu_out, output);
+            return mod->replace_instruction(ins, gpu_out);
+        });
+    }
+
     // replace the loop operator with gpu_loop operator
     void add_loop_op()
     {

test/onnx/gen_onnx.py

@@ -3835,6 +3835,7 @@ def resize_upsample_pf_test():
     return ([node], [X], [Y], [scale_tensor])
 
 
+@onnx_test
 def resize_upsample_pc_test():
     scales = np.array([1.0, 1.0, 2.0, 1.5], dtype=np.float32)
     scale_tensor = helper.make_tensor(name='scales',
@@ -3857,6 +3858,41 @@ def resize_upsample_pc_test():
     return ([node], [X], [Y], [scale_tensor])
 
 
+@onnx_test
+def roialign_default_test():
+    x = helper.make_tensor_value_info('x', TensorProto.FLOAT, [10, 4, 7, 8])
+    roi = helper.make_tensor_value_info('rois', TensorProto.FLOAT, [8, 4])
+    bi = helper.make_tensor_value_info('batch_ind', TensorProto.INT64, [8])
+    y = helper.make_tensor_value_info('y', TensorProto.FLOAT, [8, 4, 1, 1])
+
+    node = onnx.helper.make_node('RoiAlign',
+                                 inputs=['x', 'rois', 'batch_ind'],
+                                 outputs=['y'])
+
+    return ([node], [x, roi, bi], [y])
+
+
+@onnx_test
+def roialign_test():
+    x = helper.make_tensor_value_info('x', TensorProto.FLOAT, [10, 5, 4, 7])
+    roi = helper.make_tensor_value_info('rois', TensorProto.FLOAT, [8, 4])
+    bi = helper.make_tensor_value_info('batch_ind', TensorProto.INT64, [8])
+    y = helper.make_tensor_value_info('y', TensorProto.FLOAT, [8, 4, 5, 5])
+
+    node = onnx.helper.make_node(
+        'RoiAlign',
+        inputs=['x', 'rois', 'batch_ind'],
+        outputs=['y'],
+        spatial_scale=2.0,
+        output_height=5,
+        output_width=5,
+        sampling_ratio=3,
+        mode="avg",
+        coordinate_transformation_mode="output_half_pixel")
+
+    return ([node], [x, roi, bi], [y])
+
+
 @onnx_test
 def scatter_test():
     x = helper.make_tensor_value_info('data', TensorProto.FLOAT, [3, 4, 5, 6])

test/onnx/onnx_test.cpp

@@ -3640,6 +3640,55 @@ TEST_CASE(resize_upsample_pf_test)
     EXPECT(p == prog);
 }
 
+TEST_CASE(roialign_default_test)
+{
+    migraphx::shape sx{migraphx::shape::float_type, {10, 4, 7, 8}};
+    migraphx::shape srois{migraphx::shape::float_type, {8, 4}};
+    migraphx::shape sbi{migraphx::shape::int64_type, {8}};
+
+    migraphx::program p;
+    auto* mm  = p.get_main_module();
+    auto x    = mm->add_parameter("x", sx);
+    auto rois = mm->add_parameter("rois", srois);
+    auto bi   = mm->add_parameter("batch_ind", sbi);
+
+    auto r = mm->add_instruction(migraphx::make_op("roialign"), x, rois, bi);
+    mm->add_return({r});
+
+    auto prog = migraphx::parse_onnx("roialign_default_test.onnx");
+
+    EXPECT(p == prog);
+}
+
+TEST_CASE(roialign_test)
+{
+    migraphx::shape sx{migraphx::shape::float_type, {10, 5, 4, 7}};
+    migraphx::shape srois{migraphx::shape::float_type, {8, 4}};
+    migraphx::shape sbi{migraphx::shape::int64_type, {8}};
+
+    migraphx::program p;
+    auto* mm  = p.get_main_module();
+    auto x    = mm->add_parameter("x", sx);
+    auto rois = mm->add_parameter("rois", srois);
+    auto bi   = mm->add_parameter("batch_ind", sbi);
+
+    auto r = mm->add_instruction(
+        migraphx::make_op("roialign",
+                          {{"coordinate_transformation_mode", "output_half_pixel"},
+                           {"spatial_scale", 2.0f},
+                           {"output_height", 5},
+                           {"output_width", 5},
+                           {"sampling_ratio", 3}}),
+        x,
+        rois,
+        bi);
+    mm->add_return({r});
+
+    auto prog = migraphx::parse_onnx("roialign_test.onnx");
+
+    EXPECT(p == prog);
+}
+
 TEST_CASE(round_test)
 {
     migraphx::program p;

test/onnx/roialign_default_test.onnx

+roialign_default_test:

+!
+
x
+rois
+	batch_ind
y"RoiAlignroialign_default_testZ
+
x
+

+
+
+
+
+Z
+rois
+

+
+Z
+	batch_ind
+
+
+b
+
y
+

+
+
+

+
B
\ No newline at end of file

test/op_shape_test.cpp

@@ -1544,4 +1544,26 @@ TEST_CASE(where_broadcast_input)
     expect_shape(s2, migraphx::make_op("where"), s3, s1, s2);
 }
 
+TEST_CASE(roialign_test)
+{
+    migraphx::shape sx{migraphx::shape::float_type, {3, 4, 5, 6}};
+    migraphx::shape srois{migraphx::shape::float_type, {2, 4}};
+    migraphx::shape sbi{migraphx::shape::int64_type, {2}};
+    migraphx::shape sout{migraphx::shape::float_type, {2, 4, 1, 1}};
+
+    expect_shape(sout, migraphx::make_op("roialign"), sx, srois, sbi);
+
+    migraphx::shape sbi1{migraphx::shape::int64_type, {2, 3}};
+    throws_shape(migraphx::make_op("roialign"), sx, srois, sbi1);
+
+    migraphx::shape sbi2{migraphx::shape::int64_type, {3}};
+    throws_shape(migraphx::make_op("roialign"), sx, srois, sbi2);
+
+    migraphx::shape srois1{migraphx::shape::float_type, {2, 4, 3}};
+    throws_shape(migraphx::make_op("roialign"), sx, srois1, sbi);
+
+    migraphx::shape srois2{migraphx::shape::float_type, {2, 3}};
+    throws_shape(migraphx::make_op("roialign"), sx, srois2, sbi);
+}
+
 int main(int argc, const char* argv[]) { test::run(argc, argv); }

test/py/onnx_backend_test.py

@@ -174,6 +174,8 @@ def create_backend_test(testname=None, target_device=None):
         backend_test.include(r'.*test_reduce.*')
         backend_test.include(r'.*test_ReLU*')
         backend_test.include(r'.*test_relu.*')
+        backend_test.include(r'.*test_RoiAlign*')
+        backend_test.include(r'.*test_roialign.*')
         backend_test.include(r'.*test_scatter.*')
         backend_test.include(r'.*test_Scatter.*')
         backend_test.include(r'.*test_selu.*')

test/ref_ops_test.cpp

@@ -1838,8 +1838,6 @@ TEST_CASE(if_param_test)
         auto res = p.eval(m).back();
         std::vector<float> ret;
         res.visit([&](auto v) { ret.assign(v.begin(), v.end()); });
-        std::cout << std::endl;
-
         return ret;
     };
 
@@ -2768,7 +2766,6 @@ TEST_CASE(nonzero_test)
     mm->add_return({ret});
     p.compile(migraphx::ref::target{});
     auto result = p.eval({}).back();
-    std::cout << "result = " << result << std::endl;
     std::vector<int64_t> result_vector;
     result.visit([&](auto output) { result_vector.assign(output.begin(), output.end()); });
     std::vector<int64_t> gold = {0, 0, 0, 0, 1, 1, 1, 1, 0, 0, 0, 0, 0, 0, 1, 1, 0, 0,
@@ -3852,6 +3849,170 @@ TEST_CASE(reverse_test_axis10)
     EXPECT(migraphx::verify_range(results_vector, target_data));
 }
 
+TEST_CASE(roialign_out_of_bound_test)
+{
+    auto create_program = [](const std::string& trans_mode = "half_pixel") {
+        migraphx::program p;
+        auto* mm = p.get_main_module();
+        migraphx::shape x_s{migraphx::shape::float_type, {1, 1, 10, 10}};
+        std::vector<float> x_vec = {
+            0.2764, 0.7150, 0.1958, 0.3416, 0.4638, 0.0259, 0.2963, 0.6518, 0.4856, 0.7250,
+            0.9637, 0.0895, 0.2919, 0.6753, 0.0234, 0.6132, 0.8085, 0.5324, 0.8992, 0.4467,
+            0.3265, 0.8479, 0.9698, 0.2471, 0.9336, 0.1878, 0.4766, 0.4308, 0.3400, 0.2162,
+            0.0206, 0.1720, 0.2155, 0.4394, 0.0653, 0.3406, 0.7724, 0.3921, 0.2541, 0.5799,
+            0.4062, 0.2194, 0.4473, 0.4687, 0.7109, 0.9327, 0.9815, 0.6320, 0.1728, 0.6119,
+            0.3097, 0.1283, 0.4984, 0.5068, 0.4279, 0.0173, 0.4388, 0.0430, 0.4671, 0.7119,
+            0.1011, 0.8477, 0.4726, 0.1777, 0.9923, 0.4042, 0.1869, 0.7795, 0.9946, 0.9689,
+            0.1366, 0.3671, 0.7011, 0.6234, 0.9867, 0.5585, 0.6985, 0.5609, 0.8788, 0.9928,
+            0.5697, 0.8511, 0.6711, 0.9406, 0.8751, 0.7496, 0.1650, 0.1049, 0.1559, 0.2514,
+            0.7012, 0.4056, 0.7879, 0.3461, 0.0415, 0.2998, 0.5094, 0.3727, 0.5482, 0.0502};
+
+        migraphx::shape roi_s{migraphx::shape::float_type, {3, 4}};
+        std::vector<float> roi_vec = {0, 0, 9.99, 9.99, 0, 5, 4, 9, 5, 5, 9.9, 9.9};
+
+        migraphx::shape ind_s{migraphx::shape::int64_type, {3}};
+        std::vector<int64_t> ind_vec = {0, 0, 0};
+
+        auto x   = mm->add_literal(migraphx::literal(x_s, x_vec));
+        auto roi = mm->add_literal(migraphx::literal(roi_s, roi_vec));
+        auto ind = mm->add_literal(migraphx::literal(ind_s, ind_vec));
+        auto r =
+            mm->add_instruction(migraphx::make_op("roialign",
+                                                  {{"coordinate_transformation_mode", trans_mode},
+                                                   {"spatial_scale", 5.0},
+                                                   {"output_height", 1},
+                                                   {"output_width", 1},
+                                                   {"sampling_ratio", 1}}),
+                                x,
+                                roi,
+                                ind);
+        mm->add_return({r});
+        return p;
+    };
+
+    {
+        auto p = create_program("output_half_pixel");
+        p.compile(migraphx::ref::target{});
+        auto result = p.eval({}).back();
+        std::vector<float> results_vector;
+        result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
+        std::vector<float> gold = {0.0f, 0.0f, 0.0f};
+
+        EXPECT(migraphx::verify_range(results_vector, gold));
+    }
+}
+
+TEST_CASE(roialign_test)
+{
+    auto create_program = [](const std::string& trans_mode   = "half_pixel",
+                             const std::string& pooling_mode = "avg",
+                             int64_t sampling_ratio          = 2) {
+        migraphx::program p;
+        auto* mm = p.get_main_module();
+        migraphx::shape x_s{migraphx::shape::float_type, {1, 1, 10, 10}};
+        std::vector<float> x_vec = {
+            0.2764, 0.7150, 0.1958, 0.3416, 0.4638, 0.0259, 0.2963, 0.6518, 0.4856, 0.7250,
+            0.9637, 0.0895, 0.2919, 0.6753, 0.0234, 0.6132, 0.8085, 0.5324, 0.8992, 0.4467,
+            0.3265, 0.8479, 0.9698, 0.2471, 0.9336, 0.1878, 0.4766, 0.4308, 0.3400, 0.2162,
+            0.0206, 0.1720, 0.2155, 0.4394, 0.0653, 0.3406, 0.7724, 0.3921, 0.2541, 0.5799,
+            0.4062, 0.2194, 0.4473, 0.4687, 0.7109, 0.9327, 0.9815, 0.6320, 0.1728, 0.6119,
+            0.3097, 0.1283, 0.4984, 0.5068, 0.4279, 0.0173, 0.4388, 0.0430, 0.4671, 0.7119,
+            0.1011, 0.8477, 0.4726, 0.1777, 0.9923, 0.4042, 0.1869, 0.7795, 0.9946, 0.9689,
+            0.1366, 0.3671, 0.7011, 0.6234, 0.9867, 0.5585, 0.6985, 0.5609, 0.8788, 0.9928,
+            0.5697, 0.8511, 0.6711, 0.9406, 0.8751, 0.7496, 0.1650, 0.1049, 0.1559, 0.2514,
+            0.7012, 0.4056, 0.7879, 0.3461, 0.0415, 0.2998, 0.5094, 0.3727, 0.5482, 0.0502};
+
+        migraphx::shape roi_s{migraphx::shape::float_type, {3, 4}};
+        std::vector<float> roi_vec = {0, 0, 9, 9, 0, 5, 4, 9, 5, 5, 9, 9};
+
+        migraphx::shape ind_s{migraphx::shape::int64_type, {3}};
+        std::vector<int64_t> ind_vec = {0, 0, 0};
+
+        auto x   = mm->add_literal(migraphx::literal(x_s, x_vec));
+        auto roi = mm->add_literal(migraphx::literal(roi_s, roi_vec));
+        auto ind = mm->add_literal(migraphx::literal(ind_s, ind_vec));
+        auto r =
+            mm->add_instruction(migraphx::make_op("roialign",
+                                                  {{"coordinate_transformation_mode", trans_mode},
+                                                   {"spatial_scale", 1.0},
+                                                   {"output_height", 5},
+                                                   {"output_width", 5},
+                                                   {"sampling_ratio", sampling_ratio},
+                                                   {"mode", pooling_mode}}),
+                                x,
+                                roi,
+                                ind);
+        mm->add_return({r});
+        return p;
+    };
+
+    {
+        auto p = create_program();
+        p.compile(migraphx::ref::target{});
+        auto result = p.eval({}).back();
+        std::vector<float> results_vector;
+        result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
+        std::vector<float> gold = {
+            0.466421425, 0.446552634, 0.340521216, 0.568848491, 0.606780827, 0.371379346,
+            0.429571986, 0.383519977, 0.556241512, 0.351050019, 0.27680251,  0.488286227,
+            0.522200167, 0.552770197, 0.417057365, 0.471240699, 0.4844096,   0.690457463,
+            0.492039412, 0.877398551, 0.623889625, 0.712461948, 0.628926516, 0.335504025,
+            0.349469036, 0.302179992, 0.43046391,  0.469585985, 0.39774403,  0.542259991,
+            0.365552008, 0.704923987, 0.516481996, 0.317131996, 0.701444089, 0.291239977,
+            0.505897999, 0.647610962, 0.623489916, 0.829879999, 0.591567993, 0.738860011,
+            0.704825997, 0.837148011, 0.889315963, 0.622680008, 0.615276039, 0.709713995,
+            0.615356028, 0.458524048, 0.238451958, 0.337952018, 0.371693879, 0.609999895,
+            0.760059953, 0.376724035, 0.378532052, 0.71468991,  0.924308002, 0.972783983,
+            0.574903965, 0.582623959, 0.570936024, 0.761904061, 0.876998067, 0.535508037,
+            0.256580025, 0.214098021, 0.279604018, 0.360000014, 0.436488032, 0.350427985,
+            0.288755983, 0.366139978, 0.234920025};
+
+        EXPECT(migraphx::verify_range(results_vector, gold));
+    }
+
+    {
+        auto p = create_program("output_half_pixel");
+        p.compile(migraphx::ref::target{});
+        auto result = p.eval({}).back();
+        std::vector<float> results_vector;
+        result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
+        std::vector<float> gold = {
+            0.517783, 0.343411, 0.322905, 0.447362, 0.634375, 0.40308,  0.536647, 0.442791,
+            0.486144, 0.402313, 0.251194, 0.400154, 0.515524, 0.695369, 0.346537, 0.33504,
+            0.460099, 0.588069, 0.343863, 0.684932, 0.49319,  0.714058, 0.821744, 0.471935,
+            0.403946, 0.306955, 0.218678, 0.33369,  0.488001, 0.486962, 0.18709,  0.49142,
+            0.55611,  0.419167, 0.368608, 0.143278, 0.460835, 0.597125, 0.53096,  0.498207,
+            0.278818, 0.438569, 0.6022,   0.700038, 0.752436, 0.577385, 0.702383, 0.725097,
+            0.733754, 0.816304, 0.23933,  0.407514, 0.337893, 0.252521, 0.474335, 0.367075,
+            0.270168, 0.41051,  0.64189,  0.830777, 0.55564,  0.454295, 0.55645,  0.75015,
+            0.929997, 0.66257,  0.561664, 0.481275, 0.495449, 0.666306, 0.663573, 0.372107,
+            0.205603, 0.192776, 0.247849};
+
+        EXPECT(migraphx::verify_range(results_vector, gold));
+    }
+
+    {
+        auto p = create_program("output_half_pixel", "max", 0);
+        p.compile(migraphx::ref::target{});
+        auto result = p.eval({}).back();
+        std::vector<float> results_vector;
+        result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
+        std::vector<float> gold = {
+            0.819145, 0.373103, 0.258302,  0.515419, 0.726104, 0.540536, 0.545512,  0.38511,
+            0.376545, 0.274635, 0.22341,   0.184511, 0.230843, 0.404869, 0.29546,   0.540409,
+            0.265838, 0.409324, 0.213915,  0.708654, 0.687264, 0.580821, 0.461283,  0.462879,
+            0.709632, 0.27873,  0.083619,  0.22428,  0.313992, 0.410508, 0.0929099, 0.415373,
+            0.296695, 0.231574, 0.136836,  0.0683,   0.296695, 0.211925, 0.245385,  0.28053,
+            0.17091,  0.179879, 0.245385,  0.343539, 0.392742, 0.51273,  0.536193,  0.382995,
+            0.422793, 0.761886, 0.0839429, 0.276444, 0.19746,  0.126117, 0.378351,  0.254646,
+            0.092148, 0.272825, 0.381955,  0.626599, 0.251325, 0.244475, 0.194875,  0.272825,
+            0.44757,  0.351855, 0.342265,  0.244475, 0.274841, 0.553644, 0.607176,  0.202392,
+            0.07425,  0.066087, 0.126279};
+
+        EXPECT(migraphx::verify_range(results_vector, gold));
+    }
+}
+
 TEST_CASE(round_test)
 {
     migraphx::program p;

test/verify/test_roialign.cpp

+
+#include "verify_program.hpp"
+#include <migraphx/program.hpp>
+#include <migraphx/generate.hpp>
+#include <migraphx/make_op.hpp>
+
+struct test_roialign : verify_program<test_roialign>
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        auto* mm = p.get_main_module();
+        migraphx::shape x_s{migraphx::shape::float_type, {5, 4, 10, 10}};
+
+        migraphx::shape roi_s{migraphx::shape::float_type, {5, 4}};
+
+        migraphx::shape ind_s{migraphx::shape::int64_type, {5}};
+        std::vector<int64_t> ind_vec = {0, 2, 3, 4, 1};
+
+        auto x   = mm->add_parameter("x", x_s);
+        auto roi = mm->add_parameter("roi", roi_s);
+        auto ind = mm->add_literal(migraphx::literal(ind_s, ind_vec));
+        auto r   = mm->add_instruction(migraphx::make_op("roialign",
+                                                       {{"spatial_scale", 1.0},
+                                                        {"output_height", 5},
+                                                        {"output_width", 5},
+                                                        {"sampling_ratio", 2}}),
+                                     x,
+                                     roi,
+                                     ind);
+        mm->add_return({r});
+
+        return p;
+    }
+};