saymrwulf/onnxruntime
1
0
Fork 0
mirror of https://github.com/saymrwulf/onnxruntime.git synced 2026-07-27 20:02:15 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

Update MaxPool & AveragePool to support opset 10 (#1141)

Browse source 
* Update MaxPool & AveragePool to support opset 10

* fix build issue

* still use cudnn for MaxPool if dilation is not set or are default 1.

* fix build issue
This commit is contained in:
Hector Li 2019-06-02 20:42:57 -07:00 committed by GitHub
parent 10ea77a3d1
commit b8a699f70b
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
5 changed files with 58 additions and 20 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

20
onnxruntime/core/providers/cpu/nn/pool_base.h
View file

@ -50,11 +50,11 @@ class AveragePool {
static const PoolType type = PoolType::kAveragePool;
};
template <int VERSION>
template <int START_VERSION>
class MaxPool;
template <>
class MaxPool<1 /*VERSION*/> {
class MaxPool<1 /*START_VERSION*/> {
public:
static float Initialize() {
return std::numeric_limits<float>::lowest();
@ -74,7 +74,7 @@ class MaxPool<1 /*VERSION*/> {
};
template <>
class MaxPool<8 /*VERSION*/> {
class MaxPool<8 /*START_VERSION*/> {
public:
static const PoolType type = PoolType::kMaxPool;
};
@ -102,6 +102,8 @@ class PoolBase {
PoolBase(const OpKernelInfo& info) {
op_name_ = info.GetKernelDef().OpName();
global_pooling_ = (op_name_ == "GlobalAveragePool" || op_name_ == "GlobalMaxPool" || op_name_ == "GlobalLpPool");
int end;
info.GetKernelDef().SinceVersion(&start_version_, &end);
if (!global_pooling_) {
ORT_ENFORCE(info.GetAttrs<int64_t>("kernel_shape", kernel_shape_).IsOK(),
@ -123,8 +125,12 @@ class PoolBase {
ceil_mode_ = 0;
}
default_dilations_ = false;
if (!info.GetAttrs<int64_t>("dilations", dilations_).IsOK() || dilations_.empty()) {
dilations_.resize(kernel_shape_.size(), 1);
default_dilations_ = true;
} else {
default_dilations_ = std::all_of(dilations_.begin(), dilations_.end(), [](int64_t i) { return i == 1; });
}
if (op_name_ == "AveragePool") {
@ -134,10 +140,7 @@ class PoolBase {
}
if (op_name_ == "MaxPool") {
int start;
int end;
info.GetKernelDef().SinceVersion(&start, &end);
if (start == 8) {
if (start_version_ == 8) {
storage_order_ = info.GetAttrOrDefault<int64_t>("storage_order", 0 /*default_value*/);
}
}
@ -261,6 +264,9 @@ class PoolBase {
std::vector<int64_t> pads_;
std::vector<int64_t> strides_;
std::vector<int64_t> dilations_; // Introduced in MaxPool_10
int start_version_;
// default_dilations_ is true if dilations_ is not set or all dilations_ are 1
bool default_dilations_;
AutoPadType auto_pad_;

15
onnxruntime/core/providers/cuda/cuda_execution_provider.cc
View file

@ -437,9 +437,15 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain,
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 9, float, AveragePool);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 9, double, AveragePool);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 9, MLFloat16, AveragePool);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, float, AveragePool);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, double, AveragePool);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, MLFloat16, AveragePool);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float, GlobalAveragePool);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double, GlobalAveragePool);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16, GlobalAveragePool);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, float, MaxPool);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, double, MaxPool);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, MLFloat16, MaxPool);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 9, float, MaxPool);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 9, double, MaxPool);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 9, MLFloat16, MaxPool);
@ -759,9 +765,15 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 9, float, AveragePool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 9, double, AveragePool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 9, MLFloat16, AveragePool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, float, AveragePool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, double, AveragePool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, MLFloat16, AveragePool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, float, GlobalAveragePool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, double, GlobalAveragePool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MLFloat16, GlobalAveragePool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, float, MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, double, MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, MLFloat16, MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 9, float, MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 9, double, MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 9, MLFloat16, MaxPool)>,
@ -1059,8 +1071,7 @@ CUDAExecutionProvider::GetCapability(const onnxruntime::GraphViewer& graph,
if (!force_inside && (not_supported || force_outside)) {
defs_outside_cuda.insert(node.OutputDefs().cbegin(), node.OutputDefs().cend());
if (not_supported)
LOGS_DEFAULT(WARNING) << "Fallback to CPU execution provider for Op type: " << node.OpType() << " node name: " << node.Name();
LOGS_DEFAULT(WARNING) << "Fallback to CPU execution provider for Op type: " << node.OpType() << " node name: " << node.Name();
} else {
// for nodes placed on CUDA, check if its output is on CPU
node.ForEachWithIndex(

31
onnxruntime/core/providers/cuda/nn/max_pool_with_index.cu
View file

@ -29,6 +29,9 @@ __global__ void MaxPoolWithIndexKernel(
int64_t pad_h,
int64_t pad_w,
int64_t pad_d,
int64_t dilation_h,
int64_t dilation_w,
int64_t dilation_d,
fast_divmod fdm_c,
fast_divmod fdm_h,
fast_divmod fdm_w,
@ -51,9 +54,9 @@ __global__ void MaxPoolWithIndexKernel(
int64_t w_start = w_index * stride_w - pad_w;
int64_t h_start = h_index * stride_h - pad_h;
int64_t d_end = _Min<int64_t>(d_start + kernel_d, depth);
int64_t w_end = _Min<int64_t>(w_start + kernel_w, width);
int64_t h_end = _Min<int64_t>(h_start + kernel_h, height);
int64_t d_end = _Min<int64_t>(d_start + (kernel_d - 1) * dilation_d + 1, depth);
int64_t w_end = _Min<int64_t>(w_start + (kernel_w - 1) * dilation_w + 1, width);
int64_t h_end = _Min<int64_t>(h_start + (kernel_h - 1) * dilation_h + 1, height);
d_start = _Max<int64_t>(d_start, 0);
w_start = _Max<int64_t>(w_start, 0);
@ -64,9 +67,9 @@ __global__ void MaxPoolWithIndexKernel(
int64_t offset = (n_index * channels + c_index) * height * width * depth;
const T* p_slice = p_input + offset;
T maxval = p_slice[h_start * width * depth + w_start * depth + d_start] - (T)1;
for (int64_t d = d_start; d < d_end; ++d) {
for (int64_t w = w_start; w < w_end; ++w) {
for (int64_t h = h_start; h < h_end; ++h) {
for (int64_t d = d_start; d < d_end; d += dilation_d) {
for (int64_t w = w_start; w < w_end; w += dilation_w) {
for (int64_t h = h_start; h < h_end; h += dilation_h) {
if (p_slice[h * width * depth + w * depth + d] > maxval) {
h_index_max = h;
w_index_max = w;
@ -77,8 +80,10 @@ __global__ void MaxPoolWithIndexKernel(
}
}
p_output[id] = p_input[offset + h_index_max * width * depth + w_index_max * depth + d_index_max];
p_indices[id] = storage_order == 0 ? offset + h_index_max * width * depth + w_index_max * depth + d_index_max
: offset + h_index_max + w_index_max * height + d_index_max * width * height;
if (p_indices) {
p_indices[id] = storage_order == 0 ? offset + h_index_max * width * depth + w_index_max * depth + d_index_max
: offset + h_index_max + w_index_max * height + d_index_max * width * height;
}
}
template <typename T>
@ -88,6 +93,7 @@ void MaxPoolWithIndex(
const std::vector<int64_t>& kernel_shape,
const std::vector<int64_t>& stride_shape,
const std::vector<int64_t>& pads,
const std::vector<int64_t>& dilations,
int64_t storage_order,
const T* p_input,
T* p_output,
@ -112,8 +118,11 @@ void MaxPoolWithIndex(
//where xi_begin the number of pixels added at the beginning of axis i
//and xi_end, the number of pixels added at the end of axis i.
int64_t pad_h = pads[0];
int64_t pad_w = pads.size() == 4 ? pads[1] : 0;
int64_t pad_w = pads.size() >= 4 ? pads[1] : 0;
int64_t pad_d = pads.size() == 6 ? pads[2] : 0;
int64_t dilation_h = dilations[0];
int64_t dilation_w = dilations.size() >= 2 ? dilations[1] : 1;
int64_t dilation_d = dilations.size() == 3 ? dilations[2] : 1;
int64_t output_size = output_shape.Size();
fast_divmod fdm_c(static_cast<int>(channels));
@ -140,6 +149,9 @@ void MaxPoolWithIndex(
pad_h,
pad_w,
pad_d,
dilation_h,
dilation_w,
dilation_d,
fdm_c,
fdm_h,
fdm_w,
@ -158,6 +170,7 @@ void MaxPoolWithIndex(
const std::vector<int64_t>& kernel_shape, \
const std::vector<int64_t>& stride_shape, \
const std::vector<int64_t>& pads, \
const std::vector<int64_t>& dilations, \
int64_t storage_order, \
const T* p_input, \
T* p_output, \

1
onnxruntime/core/providers/cuda/nn/max_pool_with_index.h
View file

@ -14,6 +14,7 @@ void MaxPoolWithIndex(
const std::vector<int64_t>& kernel_shape,
const std::vector<int64_t>& stride_shape,
const std::vector<int64_t>& pads,
const std::vector<int64_t>& dilations,
int64_t storage_order,
const T* p_input,
T* p_output,

11
onnxruntime/core/providers/cuda/nn/pool.cc
View file

@ -33,6 +33,9 @@ namespace cuda {
POOLING_KERNEL_VERSIONED(AveragePool, float, AveragePool, 7, 9)
POOLING_KERNEL_VERSIONED(AveragePool, double, AveragePool, 7, 9)
POOLING_KERNEL_VERSIONED(AveragePool, MLFloat16, AveragePool, 7, 9)
POOLING_KERNEL(AveragePool, float, AveragePool, 10)
POOLING_KERNEL(AveragePool, double, AveragePool, 10)
POOLING_KERNEL(AveragePool, MLFloat16, AveragePool, 10)
POOLING_KERNEL(GlobalAveragePool, float, AveragePool, 1)
POOLING_KERNEL(GlobalAveragePool, double, AveragePool, 1)
POOLING_KERNEL(GlobalAveragePool, MLFloat16, AveragePool, 1)
@ -42,6 +45,9 @@ POOLING_KERNEL_VERSIONED(MaxPool, MLFloat16, MaxPool<1>, 1, 7)
POOLING_KERNEL_VERSIONED(MaxPool, float, MaxPool<8>, 8, 9)
POOLING_KERNEL_VERSIONED(MaxPool, double, MaxPool<8>, 8, 9)
POOLING_KERNEL_VERSIONED(MaxPool, MLFloat16, MaxPool<8>, 8, 9)
POOLING_KERNEL(MaxPool, float, MaxPool<8>, 10)
POOLING_KERNEL(MaxPool, double, MaxPool<8>, 10)
POOLING_KERNEL(MaxPool, MLFloat16, MaxPool<8>, 10)
POOLING_KERNEL(GlobalMaxPool, float, MaxPool<1>, 1)
POOLING_KERNEL(GlobalMaxPool, double, MaxPool<1>, 1)
POOLING_KERNEL(GlobalMaxPool, MLFloat16, MaxPool<1>, 1)
@ -182,14 +188,15 @@ Status Pool<T, MaxPool<8>>::ComputeInternal(OpKernelContext* context) const {
auto y_data = reinterpret_cast<CudaT*>(Y->template MutableData<T>());
Tensor* I = context->Output(1, TensorShape(y_dims));
if (nullptr != I) {
auto i_data = I->template MutableData<int64_t>();
if (nullptr != I || !this->default_dilations_) {
auto i_data = nullptr == I ? nullptr : I->template MutableData<int64_t>();
MaxPoolWithIndex<CudaT>(
x_shape,
TensorShape(y_dims),
kernel_shape,
strides,
pads,
this->dilations_,
this->storage_order_,
x_data,
y_data,