From 223773d278e109f1938788cd34ba628adb009514 Mon Sep 17 00:00:00 2001 From: Pranav Sharma Date: Thu, 10 Jan 2019 11:34:55 -0800 Subject: [PATCH] Implement ROI Align for object detection. (#308) * Implement ROI Align for object detection. * Fix Mac build * Fix Mac build --- onnxruntime/contrib_ops/contrib_kernels.cc | 4 + onnxruntime/contrib_ops/cpu/roialign.cc | 312 ++++++++++++++++++ onnxruntime/contrib_ops/cpu/roialign.h | 63 ++++ .../core/graph/contrib_ops/contrib_defs.cc | 105 ++++-- onnxruntime/test/contrib_ops/roialign_test.cc | 253 ++++++++++++++ 5 files changed, 712 insertions(+), 25 deletions(-) create mode 100644 onnxruntime/contrib_ops/cpu/roialign.cc create mode 100644 onnxruntime/contrib_ops/cpu/roialign.h create mode 100644 onnxruntime/test/contrib_ops/roialign_test.cc diff --git a/onnxruntime/contrib_ops/contrib_kernels.cc b/onnxruntime/contrib_ops/contrib_kernels.cc index 47ccb903af..052d9e691b 100644 --- a/onnxruntime/contrib_ops/contrib_kernels.cc +++ b/onnxruntime/contrib_ops/contrib_kernels.cc @@ -25,6 +25,8 @@ class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, WordC class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, GatherND); class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, MurmurHash3); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, float, MaxpoolWithMask); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, float, ROIAlign); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kMSDomain, 1, double, ROIAlign); void RegisterContribKernels(std::function fn) { fn(BuildKernel()); @@ -48,6 +50,8 @@ void RegisterContribKernels(std::function fn) { fn(BuildKernel()); fn(BuildKernel()); fn(BuildKernel()); + fn(BuildKernel()); + fn(BuildKernel()); } } // namespace contrib diff --git a/onnxruntime/contrib_ops/cpu/roialign.cc b/onnxruntime/contrib_ops/cpu/roialign.cc new file mode 100644 index 0000000000..bf28acf3d7 --- /dev/null +++ b/onnxruntime/contrib_ops/cpu/roialign.cc @@ -0,0 +1,312 @@ +/** +* Copyright (c) 2016-present, Facebook, Inc. +* +* Licensed under the Apache License, Version 2.0 (the "License"); +* you may not use this file except in compliance with the License. +* You may obtain a copy of the License at +* +* http://www.apache.org/licenses/LICENSE-2.0 +* +* Unless required by applicable law or agreed to in writing, software +* distributed under the License is distributed on an "AS IS" BASIS, +* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +* See the License for the specific language governing permissions and +* limitations under the License. +*/ +/* Modifications Copyright (c) Microsoft. */ + +#include "roialign.h" +#include "core/util/math_cpuonly.h" +#include "core/common/common.h" +#include "core/framework/tensor.h" + +namespace onnxruntime { +namespace contrib { +const int64_t EXPECTED_NUM_ROI_DIMS = 2; +const int64_t EXPECTED_SECOND_ROI_DIM = 5; + +#define ADD_TYPED_ROIALIGN_OP(data_type) \ + ONNX_CPU_OPERATOR_TYPED_MS_KERNEL( \ + ROIAlign, \ + 1, \ + data_type, \ + KernelDefBuilder() \ + .TypeConstraint("T", DataTypeImpl::GetTensorType()), \ + ROIAlign); + +ADD_TYPED_ROIALIGN_OP(float); +ADD_TYPED_ROIALIGN_OP(double); + +namespace { +template +struct PreCalc { + int64_t pos1; + int64_t pos2; + int64_t pos3; + int64_t pos4; + T w1; + T w2; + T w3; + T w4; +}; + +template +void pre_calc_for_bilinear_interpolate( + const int64_t height, + const int64_t width, + const int64_t pooled_height, + const int64_t pooled_width, + const int64_t iy_upper, + const int64_t ix_upper, + T roi_start_h, + T roi_start_w, + T bin_size_h, + T bin_size_w, + int64_t roi_bin_grid_h, + int64_t roi_bin_grid_w, + std::vector>& pre_calc) { + int64_t pre_calc_index = 0; + for (int64_t ph = 0; ph < pooled_height; ph++) { + for (int64_t pw = 0; pw < pooled_width; pw++) { + for (int64_t iy = 0; iy < iy_upper; iy++) { + const T yy = roi_start_h + ph * bin_size_h + + static_cast(iy + .5f) * bin_size_h / + static_cast(roi_bin_grid_h); // e.g., 0.5, 1.5 + for (int64_t ix = 0; ix < ix_upper; ix++) { + const T xx = roi_start_w + pw * bin_size_w + + static_cast(ix + .5f) * bin_size_w / + static_cast(roi_bin_grid_w); + + T x = xx; + T y = yy; + // deal with: inverse elements are out of feature map boundary + if (y < -1.0 || y > height || x < -1.0 || x > width) { + // empty + PreCalc pc; + pc.pos1 = 0; + pc.pos2 = 0; + pc.pos3 = 0; + pc.pos4 = 0; + pc.w1 = 0; + pc.w2 = 0; + pc.w3 = 0; + pc.w4 = 0; + pre_calc[pre_calc_index] = pc; + pre_calc_index += 1; + continue; + } + + if (y <= 0) { + y = 0; + } + if (x <= 0) { + x = 0; + } + + int64_t y_low = static_cast(y); + int64_t x_low = static_cast(x); + int64_t y_high; + int64_t x_high; + + if (y_low >= height - 1) { + y_high = y_low = height - 1; + y = (T)y_low; + } else { + y_high = y_low + 1; + } + + if (x_low >= width - 1) { + x_high = x_low = width - 1; + x = (T)x_low; + } else { + x_high = x_low + 1; + } + + T ly = y - y_low; + T lx = x - x_low; + T hy = static_cast(1.) - ly, hx = static_cast(1.) - lx; + T w1 = hy * hx, w2 = hy * lx, w3 = ly * hx, w4 = ly * lx; + + // save weights and indeces + PreCalc pc; + pc.pos1 = y_low * width + x_low; + pc.pos2 = y_low * width + x_high; + pc.pos3 = y_high * width + x_low; + pc.pos4 = y_high * width + x_high; + pc.w1 = w1; + pc.w2 = w2; + pc.w3 = w3; + pc.w4 = w4; + pre_calc[pre_calc_index] = pc; + + pre_calc_index += 1; + } + } + } + } +} + +template +void ROIAlignForward( + int64_t nthreads, + const T* bottom_data, + float spatial_scale, + int64_t channels, + int64_t height, + int64_t width, + int64_t pooled_height, + int64_t pooled_width, + int64_t sampling_ratio, + const T* bottom_rois, + int64_t roi_cols, + T* top_data, + const std::string& mode) { + int64_t n_rois = nthreads / channels / pooled_width / pooled_height; + +#ifdef USE_OPENMP +#pragma omp parallel for +#endif + for (int n = 0; n < n_rois; n++) { + int64_t index_n = n * channels * pooled_width * pooled_height; + + const T* offset_bottom_rois = bottom_rois + n * roi_cols; + const T roi_batch_ind = offset_bottom_rois[0]; + offset_bottom_rois++; + + // Do not using rounding; this implementation detail is critical + T roi_start_w = offset_bottom_rois[0] * spatial_scale; + T roi_start_h = offset_bottom_rois[1] * spatial_scale; + T roi_end_w = offset_bottom_rois[2] * spatial_scale; + T roi_end_h = offset_bottom_rois[3] * spatial_scale; + + // Force malformed ROIs to be 1x1 + T roi_width = std::max(roi_end_w - roi_start_w, (T)1.); + T roi_height = std::max(roi_end_h - roi_start_h, (T)1.); + T bin_size_h = static_cast(roi_height) / static_cast(pooled_height); + T bin_size_w = static_cast(roi_width) / static_cast(pooled_width); + + // We use roi_bin_grid to sample the grid and mimic integral + int64_t roi_bin_grid_h = (sampling_ratio > 0) + ? sampling_ratio + : static_cast(ceil(roi_height / pooled_height)); // e.g., = 2 + int64_t roi_bin_grid_w = + (sampling_ratio > 0) ? sampling_ratio : static_cast(ceil(roi_width / pooled_width)); + + // We do average (integral) pooling inside a bin + const int64_t count = roi_bin_grid_h * roi_bin_grid_w; // e.g. = 4 + + // we want to precalculate indices and weights shared by all channels, + // this is the key point of optimization + std::vector> pre_calc( + roi_bin_grid_h * roi_bin_grid_w * pooled_width * pooled_height); + pre_calc_for_bilinear_interpolate( + height, + width, + pooled_height, + pooled_width, + roi_bin_grid_h, + roi_bin_grid_w, + roi_start_h, + roi_start_w, + bin_size_h, + bin_size_w, + roi_bin_grid_h, + roi_bin_grid_w, + pre_calc); + + for (int64_t c = 0; c < channels; c++) { + int64_t index_n_c = index_n + c * pooled_width * pooled_height; + const T* offset_bottom_data = + bottom_data + static_cast((roi_batch_ind * channels + c) * height * width); + int64_t pre_calc_index = 0; + + for (int64_t ph = 0; ph < pooled_height; ph++) { + for (int64_t pw = 0; pw < pooled_width; pw++) { + int64_t index = index_n_c + ph * pooled_width + pw; + + T output_val = 0.; + if (mode == "avg") { // avg pooling + for (int64_t iy = 0; iy < roi_bin_grid_h; iy++) { + for (int64_t ix = 0; ix < roi_bin_grid_w; ix++) { + PreCalc pc = pre_calc[pre_calc_index]; + output_val += pc.w1 * offset_bottom_data[pc.pos1] + + pc.w2 * offset_bottom_data[pc.pos2] + + pc.w3 * offset_bottom_data[pc.pos3] + + pc.w4 * offset_bottom_data[pc.pos4]; + + pre_calc_index += 1; + } + } + output_val /= count; + } else { // max pooling + bool max_flag = false; + for (int64_t iy = 0; iy < roi_bin_grid_h; iy++) { + for (int64_t ix = 0; ix < roi_bin_grid_w; ix++) { + PreCalc pc = pre_calc[pre_calc_index]; + if (!max_flag) { + output_val = pc.w1 * offset_bottom_data[pc.pos1]; + max_flag = true; + } else { + output_val = std::max(std::max(std::max(output_val, pc.w2 * offset_bottom_data[pc.pos2]), + pc.w3 * offset_bottom_data[pc.pos3]), + pc.w4 * offset_bottom_data[pc.pos4]); + } + + pre_calc_index += 1; + } + } + } + + top_data[index] = output_val; + } // for pw + } // for ph + } // for c + } // for n +} +} // namespace + +template +Status ROIAlign::Compute(OpKernelContext* context) const { + const Tensor* X_ptr = context->Input(0); + if (!X_ptr) { + return Status(common::ONNXRUNTIME, common::INVALID_ARGUMENT, "Null input X ptr"); + } + + const Tensor* rois_ptr = context->Input(1); + if (!rois_ptr) { + return Status(common::ONNXRUNTIME, common::INVALID_ARGUMENT, "Null rois_ptr"); + } + + auto& x_dims = X_ptr->Shape(); + + auto& rois_dims = rois_ptr->Shape(); + + // validate rois_dims + if (rois_dims.NumDimensions() != EXPECTED_NUM_ROI_DIMS) { + return Status(common::ONNXRUNTIME, common::INVALID_ARGUMENT, "Number of dimensions for rois should be exactly 2"); + } + if (rois_dims[1] != EXPECTED_SECOND_ROI_DIM) { + return Status(common::ONNXRUNTIME, common::INVALID_ARGUMENT, "Second dimension for rois should be exactly 5"); + } + + auto& Y = *context->Output(0, {rois_dims[0], x_dims[1], pooled_h_, pooled_w_}); + int64_t output_size = Y.Shape().Size(); + ROIAlignForward( + output_size, + X_ptr->Data(), + spatial_scale_, + x_dims[1], + x_dims[2], + x_dims[3], + pooled_h_, + pooled_w_, + sampling_ratio_, + rois_ptr->Data(), + rois_dims[1], + Y.template MutableData(), + mode_); + + return Status::OK(); +} +} // namespace contrib +} // namespace onnxruntime diff --git a/onnxruntime/contrib_ops/cpu/roialign.h b/onnxruntime/contrib_ops/cpu/roialign.h new file mode 100644 index 0000000000..5e52b7ac98 --- /dev/null +++ b/onnxruntime/contrib_ops/cpu/roialign.h @@ -0,0 +1,63 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#pragma once + +#include "core/framework/op_kernel.h" +#include + +namespace onnxruntime { +namespace contrib { +template +class ROIAlign final : public OpKernel { + public: + explicit ROIAlign(const OpKernelInfo& info) : OpKernel(info) { + // mode + std::string mode_tmp; + if (info.GetAttr("mode", &mode_tmp).IsOK()) { + mode_ = mode_tmp; + std::transform(mode_.begin(), mode_.end(), mode_.begin(), [](auto& i) { return static_cast(::tolower(i)); }); + if (mode_ != "avg" && mode_ != "max") { + ORT_THROW("Invalid mode of value ", mode_, " specified. It should be either avg or max"); + } + } + + // pooled_h + int64_t pooled_h_tmp; + if (info.GetAttr("pooled_h", &pooled_h_tmp).IsOK()) { + pooled_h_ = pooled_h_tmp; + } + + // pooled_w + int64_t pooled_w_tmp; + if (info.GetAttr("pooled_w", &pooled_w_tmp).IsOK()) { + pooled_w_ = pooled_w_tmp; + } + + // sampling_ratio + int64_t sampling_ratio_tmp; + if (info.GetAttr("sampling_ratio", &sampling_ratio_tmp).IsOK()) { + sampling_ratio_ = sampling_ratio_tmp; + ORT_ENFORCE(sampling_ratio_ >= 0, "Sampling ratio should be >=0, but it was ", sampling_ratio_); + } + + // spatial_scale + float spatial_scale_tmp; + if (info.GetAttr("spatial_scale", &spatial_scale_tmp).IsOK()) { + spatial_scale_ = spatial_scale_tmp; + } + } + + Status Compute(OpKernelContext* context) const override; + + private: + std::string mode_{"avg"}; + int64_t pooled_h_{1}; + int64_t pooled_w_{1}; + int64_t sampling_ratio_{0}; + float spatial_scale_{1.0f}; + + ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(ROIAlign); +}; +} // namespace contrib +} // namespace onnxruntime diff --git a/onnxruntime/core/graph/contrib_ops/contrib_defs.cc b/onnxruntime/core/graph/contrib_ops/contrib_defs.cc index bd2589729f..183109d01a 100644 --- a/onnxruntime/core/graph/contrib_ops/contrib_defs.cc +++ b/onnxruntime/core/graph/contrib_ops/contrib_defs.cc @@ -18,8 +18,8 @@ void convPoolTypeAndShapeInference(ONNX_NAMESPACE::InferenceContext& ctx, bool u namespace onnxruntime { namespace contrib { using ::ONNX_NAMESPACE::AttributeProto; -using ::ONNX_NAMESPACE::OPTIONAL; using ::ONNX_NAMESPACE::OpSchema; +using ::ONNX_NAMESPACE::OPTIONAL; void RegisterContribSchemas() { ONNX_CONTRIB_OPERATOR_SCHEMA(SampleOp) @@ -80,7 +80,7 @@ Sample echo operator.)DOC"); .SetDomain(kMSDomain) .SinceVersion(1) .SetDoc(R"DOC( -The fused convolution operator schema is the same as Conv besides it includes an attribute +The fused convolution operator schema is the same as Conv besides it includes an attribute activation.)DOC") .Attr( "auto_pad", @@ -318,33 +318,33 @@ activation.)DOC") } }) .SetDoc(R"DOC( -This transform extracts n-grams from the input sequence and save them as a vector. Input can -be either a 1-D or 2-D tensor. For 1-D input, output is the n-gram representation of that input. -For 2-D input, the output is also a 2-D tensor whose i-th row is the n-gram representation of the i-th input row. -More specifically, if input shape is [C], the corresponding output shape would be [max(ngram_indexes) + 1]. -If input shape is [N, C], this operator produces a [N, max(ngram_indexes) + 1]-tensor. - -In contrast to standard n-gram extraction, here, the indexes of extracting an n-gram from the original -sequence are not necessarily consecutive numbers. The discontinuity between indexes are controlled by the number of skips. -If the number of skips is 2, we should skip two tokens when scanning through the original sequence. -Let's consider an example. Assume that input sequence is [94, 17, 36, 12, 28] and the number of skips is 2. -The associated 2-grams are [94, 12] and [17, 28] respectively indexed by [0, 3] and [1, 4]. -If the number of skips becomes 0, the 2-grams generated are [94, 17], [17, 36], [36, 12], [12, 28] +This transform extracts n-grams from the input sequence and save them as a vector. Input can +be either a 1-D or 2-D tensor. For 1-D input, output is the n-gram representation of that input. +For 2-D input, the output is also a 2-D tensor whose i-th row is the n-gram representation of the i-th input row. +More specifically, if input shape is [C], the corresponding output shape would be [max(ngram_indexes) + 1]. +If input shape is [N, C], this operator produces a [N, max(ngram_indexes) + 1]-tensor. + +In contrast to standard n-gram extraction, here, the indexes of extracting an n-gram from the original +sequence are not necessarily consecutive numbers. The discontinuity between indexes are controlled by the number of skips. +If the number of skips is 2, we should skip two tokens when scanning through the original sequence. +Let's consider an example. Assume that input sequence is [94, 17, 36, 12, 28] and the number of skips is 2. +The associated 2-grams are [94, 12] and [17, 28] respectively indexed by [0, 3] and [1, 4]. +If the number of skips becomes 0, the 2-grams generated are [94, 17], [17, 36], [36, 12], [12, 28] indexed by [0, 1], [1, 2], [2, 3], [3, 4], respectively. -The output vector stores the count of each n-gram; -Y[i] indicates the times that the i-th n-gram is found. The attribute ngram_indexes is used to determine the mapping +The output vector stores the count of each n-gram; +Y[i] indicates the times that the i-th n-gram is found. The attribute ngram_indexes is used to determine the mapping between index i and the corresponding n-gram. If pool_int64s is [94 , 17 ,17, 36], ngram_indexes is [1, 0], ngram_counts=[0, 0], then the Y[0] (first element in Y) and Y[1] (second element in Y) are the counts of [17, 36] and [94, 17], -respectively. An n-gram which cannot be found in pool_strings/pool_int64s should be ignored and has no effect on the output. -Note that we may consider all skips up to S when generating the n-grams. - -The examples used above are true if mode is "TF". If mode is "IDF", all the counts larger than 1 would be truncated to 1 and -the i-th element in weights would be used to scale (by multiplication) the count of the i-th n-gram in pool. If mode is "TFIDF", -this operator first computes the counts of all n-grams and then scale them by the associated values in the weights attribute. - -Only one of pool_strings and pool_int64s can be set. If pool_int64s is set, the input should be an integer tensor. -If pool_strings is set, the input must be a string tensor. +respectively. An n-gram which cannot be found in pool_strings/pool_int64s should be ignored and has no effect on the output. +Note that we may consider all skips up to S when generating the n-grams. + +The examples used above are true if mode is "TF". If mode is "IDF", all the counts larger than 1 would be truncated to 1 and +the i-th element in weights would be used to scale (by multiplication) the count of the i-th n-gram in pool. If mode is "TFIDF", +this operator first computes the counts of all n-grams and then scale them by the associated values in the weights attribute. + +Only one of pool_strings and pool_int64s can be set. If pool_int64s is set, the input should be an integer tensor. +If pool_strings is set, the input must be a string tensor. )DOC"); // Operators for linear 8 bit quanitzation support. @@ -893,6 +893,61 @@ Example 4: "Constrain to tensor(float).") .SetDoc(R"DOC(The WordConvEmbedding takes in a batch of sequence words and embed each word to a vector.)DOC"); + ONNX_CONTRIB_OPERATOR_SCHEMA(ROIAlign) + .SetDomain(kMSDomain) + .SinceVersion(1) + .Attr( + "spatial_scale", + "Multiplicative spatial scale factor to translate ROI coordinates " + "from their input spatial scale to the scale used when pooling, " + "i.e., spatial scale of the input feature map X relative to the " + "input image. E.g.; default is 1.0f. ", + AttributeProto::FLOAT, + 1.f) + .Attr( + "pooled_h", + "default 1; Pooled output Y's height.", + AttributeProto::INT, + static_cast(1)) + .Attr( + "pooled_w", + "default 1; Pooled output Y's width.", + AttributeProto::INT, + static_cast(1)) + .Attr( + "sampling_ratio", + "Number of sampling points in the interpolation grid used to compute " + "the output value of each pooled output bin. If > 0, then exactly " + "sampling_ratio x sampling_ratio grid points are used. If == 0, then " + "an adaptive number of grid points are used (computed as " + "ceil(roi_width / pooled_w), and likewise for height). Default is 0.", + AttributeProto::INT, + static_cast(0)) + .Attr( + "mode", + "The pooling method. Two modes are supported: 'avg' and 'max'. " + "Default is 'avg'.", + AttributeProto::STRING, + std::string("avg")) + .Input(0, "X", "Input data tensor from the previous operator; 4-D feature map of shape (N x C x H x W), where N is the batch size, C is the number of channels, and H and W are the height and the width of the data.", "T") + .Input(1, "rois", "RoIs (Regions of Interest2) to pool over; rois is 2-D input of shape (num_rois, 5) given as [[batch_id, x1, y1, x2, y2], ...]. The RoIs' coordinates are in the coordinate system of the input image.", "T") + .Output(0, "Y", "RoI pooled output, 4-D tesnor of shape (num_rois, C, pooled_h, pooled_w). The r-th batch element Y[r-1] is a pooled feature map corresponding to the r-th RoI X[r-1].", "T") + .TypeConstraint( + "T", + {"tensor(float16)", "tensor(float)", "tensor(double)"}, + "Constrain to float, float16 and double tensors.") + .SetDoc(R"DOC(Region of Interest (RoI) align operation described in the + [Mask R-CNN paper](https://arxiv.org/abs/1703.06870). + RoIAlign consumes an input tensor X and region of interests (rois) + to apply pooling across each RoI; it produces a 4-D tensor of shape + (num_rois, C, pooled_h, pooled_w). + + RoIAlign is proposed to avoid the misalignment by removing + quantizations while converting from original image into feature + map and from feature map into RoI feature; in each ROI bin, + the value of the sampled locations are computed directly + through bilinear interpolation.)DOC"); + #ifdef MICROSOFT_INTERNAL // register internal ops RegisterInternalSchemas(); diff --git a/onnxruntime/test/contrib_ops/roialign_test.cc b/onnxruntime/test/contrib_ops/roialign_test.cc new file mode 100644 index 0000000000..3dd52b8b0c --- /dev/null +++ b/onnxruntime/test/contrib_ops/roialign_test.cc @@ -0,0 +1,253 @@ +// Copyright (c) Microsoft Corporation. All rights reserved. +// Licensed under the MIT License. + +#include "gtest/gtest.h" +#include "test/providers/provider_test_utils.h" + +namespace onnxruntime { +namespace test { + +TEST(ROIAlignTest, AvgModePositive) { + OpTester test("ROIAlign", 1, onnxruntime::kMSDomain); + test.AddAttribute("pooled_h", 3); + test.AddAttribute("pooled_w", 4); + test.AddAttribute("sampling_ratio", 2); + test.AddAttribute("spatial_scale", 1.0f / 16.0f); + + const int N = 1; + const int C = 3; + const int H = 5; + const int W = 5; + + std::vector rois{0.,7.,5.,7.,5.,0.,-15.,-15.,-15.,-15.,0.,-10.,21.,-10.,21.,0.,13.,8.,13.,8.,0.,-14.,19.,-14.,19.}; + test.AddInput("X", {N, C, H, W}, {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., + 25.,26.,27.,28.,29.,30.,31.,32.,33.,34.,35.,36.,37.,38.,39.,40.,41.,42.,43.,44.,45.,46., + 47.,48.,49.,50.,51.,52.,53.,54.,55.,56.,57.,58.,59.,60.,61.,62.,63.,64.,65.,66.,67.,68., + 69.,70.,71.,72.,73.,74.}); + test.AddInput("rois", {5, 5}, {0.,7.,5.,7.,5.,0.,-15.,-15.,-15.,-15.,0.,-10.,21.,-10.,21.,0.,13.,8.,13.,8.,0.,-14.,19.,-14.,19.}); + test.AddOutput("Y", {5,3,3,4}, {2.95833f,3.20833f,3.45833f,3.70833f,4.625f,4.875f,5.125f,5.375f, + 6.29167f,6.54167f,6.79167f,7.04167f,27.9583f,28.2083f,28.4583f, + 28.7083f,29.625f,29.875f,30.125f,30.375f,31.2917f,31.5417f,31.7917f, + 32.0417f,52.9583f,53.2083f,53.4583f,53.7083f,54.625f,54.875f,55.125f, + 55.375f,56.2917f,56.5417f,56.7917f,57.0417f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f, + 25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f, + 50.f,50.f,50.f,7.39583f,7.39583f,7.42708f,7.64583f,9.0625f,9.0625f,9.09375f, + 9.3125f,10.7292f,10.7292f,10.7604f,10.9792f,32.3958f,32.3958f,32.4271f, + 32.6458f,34.0625f,34.0625f,34.0938f,34.3125f,35.7292f,35.7292f,35.7604f, + 35.9792f,57.3958f,57.3958f,57.4271f,57.6458f,59.0625f,59.0625f,59.0938f, + 59.3125f,60.7292f,60.7292f,60.7604f,60.9792f,4.27083f,4.52083f,4.77083f, + 5.02083f,5.9375f,6.1875f,6.4375f,6.6875f,7.60417f,7.85417f,8.10417f,8.35417f, + 29.2708f,29.5208f,29.7708f,30.0208f,30.9375f,31.1875f,31.4375f,31.6875f,32.6042f, + 32.8542f,33.1042f,33.3542f,54.2708f,54.5208f,54.7708f,55.0208f,55.9375f,56.1875f, + 56.4375f,56.6875f,57.6042f,57.8542f,58.1042f,58.3542f,6.77083f,6.77083f,6.77083f, + 6.80208f,8.4375f,8.4375f,8.4375f,8.46875f,10.1042f,10.1042f,10.1042f,10.1354f,31.7708f, + 31.7708f,31.7708f,31.8021f,33.4375f,33.4375f,33.4375f,33.4688f,35.1042f,35.1042f,35.1042f, + 35.1354f,56.7708f,56.7708f,56.7708f,56.8021f,58.4375f,58.4375f,58.4375f,58.4688f,60.1042f, + 60.1042f,60.1042f,60.1354f}); + + test.Run(); +} + +TEST(ROIAlignTest, MaxModePositive) { + OpTester test("ROIAlign", 1, onnxruntime::kMSDomain); + test.AddAttribute("mode", "max"); + test.AddAttribute("pooled_h", 3); + test.AddAttribute("pooled_w", 4); + test.AddAttribute("sampling_ratio", 2); + test.AddAttribute("spatial_scale", 1.0f / 16.0f); + + const int N = 1; + const int C = 3; + const int H = 5; + const int W = 5; + + std::vector rois{0.,7.,5.,7.,5.,0.,-15.,-15.,-15.,-15.,0.,-10.,21.,-10.,21.,0.,13.,8.,13.,8.,0.,-14.,19.,-14.,19.}; + test.AddInput("X", {N, C, H, W}, {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., + 25.,26.,27.,28.,29.,30.,31.,32.,33.,34.,35.,36.,37.,38.,39.,40.,41.,42.,43.,44.,45.,46., + 47.,48.,49.,50.,51.,52.,53.,54.,55.,56.,57.,58.,59.,60.,61.,62.,63.,64.,65.,66.,67.,68., + 69.,70.,71.,72.,73.,74.}); + test.AddInput("rois", {5, 5}, {0.,7.,5.,7.,5.,0.,-15.,-15.,-15.,-15.,0.,-10.,21.,-10.,21.,0.,13.,8.,13.,8.,0.,-14.,19.,-14.,19.}); + test.AddOutput("Y", {5,3,3,4}, {2.10938f,2.95313f,3.375f,2.53125f,3.35938f,4.70313f,5.375f,4.03125f,3.51563f,4.92188f,5.625f, + 4.21875f,10.8984f,15.2578f,17.4375f,13.0781f,17.3568f,24.2995f,27.7708f,20.8281f,18.1641f, + 25.4297f,29.0625f,21.7969f,19.6875f,27.5625f,31.5f,23.625f,31.3542f,43.8958f,50.1667f,37.625f, + 32.8125f,45.9375f,52.5f,39.375f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f, + 25.f,25.f,25.f,25.f,25.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f,5.625f,5.625f,5.625f,4.57031f, + 8.95833f,8.95833f,8.95833f,7.27865f,9.375f,9.375f,9.375f,7.61719f,19.6875f,19.6875f,19.6875f,15.9961f, + 31.3542f,31.3542f,31.3542f,25.4753f,32.8125f,32.8125f,32.8125f,26.6602f,33.75f,33.75f,33.75f,27.4219f, + 53.75f,53.75f,53.75f,43.6719f,56.25f,56.25f,56.25f,45.7031f,4.5f,3.9375f,2.8125f,3.9375f,5.5f,4.125f, + 3.20833f,4.8125f,4.58333f,4.01042f,2.86458f,3.9375f,23.25f,20.3438f,14.5313f,18.f,28.4167f,21.3125f, + 14.6667f,22.f,15.8229f,20.3437f,14.5312f,18.f,42.f,36.75f,26.25f,32.0625f,51.3333f,38.5f,26.125f,39.1875f, + 28.5833f,36.75f,26.25f,32.0625f,4.375f,4.375f,4.375f,4.375f,7.70833f,7.70833f,7.70833f,7.70833f,9.375f, + 9.375f,9.375f,8.78906f,21.875f,21.875f,21.875f,21.875f,26.9792f,26.9792f,26.9792f,26.9792f,32.8125f, + 32.8125f,32.8125f,30.7617f,40.1042f,40.1042f,40.1042f,40.1042f,46.25f,46.25f,46.25f,46.25f,56.25f,56.25f, + 56.25f,52.7344f}); + + test.Run(); +} + +TEST(ROIAlignTest, AvgModeNegativeInvalidMode) { + OpTester test("ROIAlign", 1, onnxruntime::kMSDomain); + test.AddAttribute("mode", "foobar"); // <-- + test.AddAttribute("pooled_h", 3); + test.AddAttribute("pooled_w", 4); + test.AddAttribute("sampling_ratio", -2); + test.AddAttribute("spatial_scale", 1.0f / 16.0f); + + const int N = 1; + const int C = 3; + const int H = 5; + const int W = 5; + + std::vector rois{0.,7.,5.,7.,5.,0.,-15.,-15.,-15.,-15.,0.,-10.,21.,-10.,21.,0.,13.,8.,13.,8.,0.,-14.,19.,-14.,19.}; + test.AddInput("X", {N, C, H, W}, {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., + 25.,26.,27.,28.,29.,30.,31.,32.,33.,34.,35.,36.,37.,38.,39.,40.,41.,42.,43.,44.,45.,46., + 47.,48.,49.,50.,51.,52.,53.,54.,55.,56.,57.,58.,59.,60.,61.,62.,63.,64.,65.,66.,67.,68., + 69.,70.,71.,72.,73.,74.}); + test.AddInput("rois", {5, 5}, {0.,7.,5.,7.,5.,0.,-15.,-15.,-15.,-15.,0.,-10.,21.,-10.,21.,0.,13.,8.,13.,8.,0.,-14.,19.,-14.,19.}); + test.AddOutput("Y", {5,3,3,4}, {2.95833f,3.20833f,3.45833f,3.70833f,4.625f,4.875f,5.125f,5.375f, + 6.29167f,6.54167f,6.79167f,7.04167f,27.9583f,28.2083f,28.4583f, + 28.7083f,29.625f,29.875f,30.125f,30.375f,31.2917f,31.5417f,31.7917f, + 32.0417f,52.9583f,53.2083f,53.4583f,53.7083f,54.625f,54.875f,55.125f, + 55.375f,56.2917f,56.5417f,56.7917f,57.0417f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f, + 25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f, + 50.f,50.f,50.f,7.39583f,7.39583f,7.42708f,7.64583f,9.0625f,9.0625f,9.09375f, + 9.3125f,10.7292f,10.7292f,10.7604f,10.9792f,32.3958f,32.3958f,32.4271f, + 32.6458f,34.0625f,34.0625f,34.0938f,34.3125f,35.7292f,35.7292f,35.7604f, + 35.9792f,57.3958f,57.3958f,57.4271f,57.6458f,59.0625f,59.0625f,59.0938f, + 59.3125f,60.7292f,60.7292f,60.7604f,60.9792f,4.27083f,4.52083f,4.77083f, + 5.02083f,5.9375f,6.1875f,6.4375f,6.6875f,7.60417f,7.85417f,8.10417f,8.35417f, + 29.2708f,29.5208f,29.7708f,30.0208f,30.9375f,31.1875f,31.4375f,31.6875f,32.6042f, + 32.8542f,33.1042f,33.3542f,54.2708f,54.5208f,54.7708f,55.0208f,55.9375f,56.1875f, + 56.4375f,56.6875f,57.6042f,57.8542f,58.1042f,58.3542f,6.77083f,6.77083f,6.77083f, + 6.80208f,8.4375f,8.4375f,8.4375f,8.46875f,10.1042f,10.1042f,10.1042f,10.1354f,31.7708f, + 31.7708f,31.7708f,31.8021f,33.4375f,33.4375f,33.4375f,33.4688f,35.1042f,35.1042f,35.1042f, + 35.1354f,56.7708f,56.7708f,56.7708f,56.8021f,58.4375f,58.4375f,58.4375f,58.4688f,60.1042f, + 60.1042f,60.1042f,60.1354f}); + + test.Run(OpTester::ExpectResult::kExpectFailure, "Invalid mode"); +} + +TEST(ROIAlignTest, AvgModeNegativeSamplingRatio) { + OpTester test("ROIAlign", 1, onnxruntime::kMSDomain); + test.AddAttribute("pooled_h", 3); + test.AddAttribute("pooled_w", 4); + test.AddAttribute("sampling_ratio", -2); // <-- + test.AddAttribute("spatial_scale", 1.0f / 16.0f); + + const int N = 1; + const int C = 3; + const int H = 5; + const int W = 5; + + std::vector rois{0.,7.,5.,7.,5.,0.,-15.,-15.,-15.,-15.,0.,-10.,21.,-10.,21.,0.,13.,8.,13.,8.,0.,-14.,19.,-14.,19.}; + test.AddInput("X", {N, C, H, W}, {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., + 25.,26.,27.,28.,29.,30.,31.,32.,33.,34.,35.,36.,37.,38.,39.,40.,41.,42.,43.,44.,45.,46., + 47.,48.,49.,50.,51.,52.,53.,54.,55.,56.,57.,58.,59.,60.,61.,62.,63.,64.,65.,66.,67.,68., + 69.,70.,71.,72.,73.,74.}); + test.AddInput("rois", {5, 5}, {0.,7.,5.,7.,5.,0.,-15.,-15.,-15.,-15.,0.,-10.,21.,-10.,21.,0.,13.,8.,13.,8.,0.,-14.,19.,-14.,19.}); + test.AddOutput("Y", {5,3,3,4}, {2.95833f,3.20833f,3.45833f,3.70833f,4.625f,4.875f,5.125f,5.375f, + 6.29167f,6.54167f,6.79167f,7.04167f,27.9583f,28.2083f,28.4583f, + 28.7083f,29.625f,29.875f,30.125f,30.375f,31.2917f,31.5417f,31.7917f, + 32.0417f,52.9583f,53.2083f,53.4583f,53.7083f,54.625f,54.875f,55.125f, + 55.375f,56.2917f,56.5417f,56.7917f,57.0417f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f, + 25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f, + 50.f,50.f,50.f,7.39583f,7.39583f,7.42708f,7.64583f,9.0625f,9.0625f,9.09375f, + 9.3125f,10.7292f,10.7292f,10.7604f,10.9792f,32.3958f,32.3958f,32.4271f, + 32.6458f,34.0625f,34.0625f,34.0938f,34.3125f,35.7292f,35.7292f,35.7604f, + 35.9792f,57.3958f,57.3958f,57.4271f,57.6458f,59.0625f,59.0625f,59.0938f, + 59.3125f,60.7292f,60.7292f,60.7604f,60.9792f,4.27083f,4.52083f,4.77083f, + 5.02083f,5.9375f,6.1875f,6.4375f,6.6875f,7.60417f,7.85417f,8.10417f,8.35417f, + 29.2708f,29.5208f,29.7708f,30.0208f,30.9375f,31.1875f,31.4375f,31.6875f,32.6042f, + 32.8542f,33.1042f,33.3542f,54.2708f,54.5208f,54.7708f,55.0208f,55.9375f,56.1875f, + 56.4375f,56.6875f,57.6042f,57.8542f,58.1042f,58.3542f,6.77083f,6.77083f,6.77083f, + 6.80208f,8.4375f,8.4375f,8.4375f,8.46875f,10.1042f,10.1042f,10.1042f,10.1354f,31.7708f, + 31.7708f,31.7708f,31.8021f,33.4375f,33.4375f,33.4375f,33.4688f,35.1042f,35.1042f,35.1042f, + 35.1354f,56.7708f,56.7708f,56.7708f,56.8021f,58.4375f,58.4375f,58.4375f,58.4688f,60.1042f, + 60.1042f,60.1042f,60.1354f}); + + test.Run(OpTester::ExpectResult::kExpectFailure, "Sampling ratio should be >=0"); +} + +TEST(ROIAlignTest, AvgModeNegativeInvalidNumRoiDims) { + OpTester test("ROIAlign", 1, onnxruntime::kMSDomain); + test.AddAttribute("pooled_h", 3); + test.AddAttribute("pooled_w", 4); + test.AddAttribute("sampling_ratio", 2); + test.AddAttribute("spatial_scale", 1.0f / 16.0f); + + const int N = 1; + const int C = 3; + const int H = 5; + const int W = 5; + + std::vector rois{0.,7.,5.,7.,5.,0.,-15.,-15.,-15.,-15.,0.,-10.,21.,-10.,21.,0.,13.,8.,13.,8.,0.,-14.,19.,-14.,19.}; + test.AddInput("X", {N, C, H, W}, {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., + 25.,26.,27.,28.,29.,30.,31.,32.,33.,34.,35.,36.,37.,38.,39.,40.,41.,42.,43.,44.,45.,46., + 47.,48.,49.,50.,51.,52.,53.,54.,55.,56.,57.,58.,59.,60.,61.,62.,63.,64.,65.,66.,67.,68., + 69.,70.,71.,72.,73.,74.}); + test.AddInput("rois", {5, 4, 1}, {0.,7.,5.,7.,5.,0.,-15.,-15.,-15.,-15.,0.,-10.,21.,-10.,21.,0.,13.,8.,13.,8.}); // <-- + test.AddOutput("Y", {5,3,3,4}, {2.95833f,3.20833f,3.45833f,3.70833f,4.625f,4.875f,5.125f,5.375f, + 6.29167f,6.54167f,6.79167f,7.04167f,27.9583f,28.2083f,28.4583f, + 28.7083f,29.625f,29.875f,30.125f,30.375f,31.2917f,31.5417f,31.7917f, + 32.0417f,52.9583f,53.2083f,53.4583f,53.7083f,54.625f,54.875f,55.125f, + 55.375f,56.2917f,56.5417f,56.7917f,57.0417f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f, + 25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f, + 50.f,50.f,50.f,7.39583f,7.39583f,7.42708f,7.64583f,9.0625f,9.0625f,9.09375f, + 9.3125f,10.7292f,10.7292f,10.7604f,10.9792f,32.3958f,32.3958f,32.4271f, + 32.6458f,34.0625f,34.0625f,34.0938f,34.3125f,35.7292f,35.7292f,35.7604f, + 35.9792f,57.3958f,57.3958f,57.4271f,57.6458f,59.0625f,59.0625f,59.0938f, + 59.3125f,60.7292f,60.7292f,60.7604f,60.9792f,4.27083f,4.52083f,4.77083f, + 5.02083f,5.9375f,6.1875f,6.4375f,6.6875f,7.60417f,7.85417f,8.10417f,8.35417f, + 29.2708f,29.5208f,29.7708f,30.0208f,30.9375f,31.1875f,31.4375f,31.6875f,32.6042f, + 32.8542f,33.1042f,33.3542f,54.2708f,54.5208f,54.7708f,55.0208f,55.9375f,56.1875f, + 56.4375f,56.6875f,57.6042f,57.8542f,58.1042f,58.3542f,6.77083f,6.77083f,6.77083f, + 6.80208f,8.4375f,8.4375f,8.4375f,8.46875f,10.1042f,10.1042f,10.1042f,10.1354f,31.7708f, + 31.7708f,31.7708f,31.8021f,33.4375f,33.4375f,33.4375f,33.4688f,35.1042f,35.1042f,35.1042f, + 35.1354f,56.7708f,56.7708f,56.7708f,56.8021f,58.4375f,58.4375f,58.4375f,58.4688f,60.1042f, + 60.1042f,60.1042f,60.1354f}); + + test.Run(OpTester::ExpectResult::kExpectFailure, "Number of dimensions for rois should be exactly 2"); +} + +TEST(ROIAlignTest, AvgModeNegativeInvalidSecondRoiDims) { + OpTester test("ROIAlign", 1, onnxruntime::kMSDomain); + test.AddAttribute("pooled_h", 3); + test.AddAttribute("pooled_w", 4); + test.AddAttribute("sampling_ratio", 2); + test.AddAttribute("spatial_scale", 1.0f / 16.0f); + + const int N = 1; + const int C = 3; + const int H = 5; + const int W = 5; + + std::vector rois{0.,7.,5.,7.,5.,0.,-15.,-15.,-15.,-15.,0.,-10.,21.,-10.,21.,0.,13.,8.,13.,8.,0.,-14.,19.,-14.,19.}; + test.AddInput("X", {N, C, H, W}, {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., + 25.,26.,27.,28.,29.,30.,31.,32.,33.,34.,35.,36.,37.,38.,39.,40.,41.,42.,43.,44.,45.,46., + 47.,48.,49.,50.,51.,52.,53.,54.,55.,56.,57.,58.,59.,60.,61.,62.,63.,64.,65.,66.,67.,68., + 69.,70.,71.,72.,73.,74.}); + test.AddInput("rois", {5, 4}, {0.,7.,5.,7.,5.,0.,-15.,-15.,-15.,-15.,0.,-10.,21.,-10.,21.,0.,13.,8.,13.,8.}); // <-- + test.AddOutput("Y", {5,3,3,4}, {2.95833f,3.20833f,3.45833f,3.70833f,4.625f,4.875f,5.125f,5.375f, + 6.29167f,6.54167f,6.79167f,7.04167f,27.9583f,28.2083f,28.4583f, + 28.7083f,29.625f,29.875f,30.125f,30.375f,31.2917f,31.5417f,31.7917f, + 32.0417f,52.9583f,53.2083f,53.4583f,53.7083f,54.625f,54.875f,55.125f, + 55.375f,56.2917f,56.5417f,56.7917f,57.0417f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f,0.f, + 25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,25.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f,50.f, + 50.f,50.f,50.f,7.39583f,7.39583f,7.42708f,7.64583f,9.0625f,9.0625f,9.09375f, + 9.3125f,10.7292f,10.7292f,10.7604f,10.9792f,32.3958f,32.3958f,32.4271f, + 32.6458f,34.0625f,34.0625f,34.0938f,34.3125f,35.7292f,35.7292f,35.7604f, + 35.9792f,57.3958f,57.3958f,57.4271f,57.6458f,59.0625f,59.0625f,59.0938f, + 59.3125f,60.7292f,60.7292f,60.7604f,60.9792f,4.27083f,4.52083f,4.77083f, + 5.02083f,5.9375f,6.1875f,6.4375f,6.6875f,7.60417f,7.85417f,8.10417f,8.35417f, + 29.2708f,29.5208f,29.7708f,30.0208f,30.9375f,31.1875f,31.4375f,31.6875f,32.6042f, + 32.8542f,33.1042f,33.3542f,54.2708f,54.5208f,54.7708f,55.0208f,55.9375f,56.1875f, + 56.4375f,56.6875f,57.6042f,57.8542f,58.1042f,58.3542f,6.77083f,6.77083f,6.77083f, + 6.80208f,8.4375f,8.4375f,8.4375f,8.46875f,10.1042f,10.1042f,10.1042f,10.1354f,31.7708f, + 31.7708f,31.7708f,31.8021f,33.4375f,33.4375f,33.4375f,33.4688f,35.1042f,35.1042f,35.1042f, + 35.1354f,56.7708f,56.7708f,56.7708f,56.8021f,58.4375f,58.4375f,58.4375f,58.4688f,60.1042f, + 60.1042f,60.1042f,60.1354f}); + + test.Run(OpTester::ExpectResult::kExpectFailure, "Second dimension for rois should be exactly 5"); +} +} // namespace test +} // namespace onnxruntime