diff --git a/onnxruntime/core/providers/cpu/nn/pool_base.h b/onnxruntime/core/providers/cpu/nn/pool_base.h index b6fa0843e3..11b70ac364 100644 --- a/onnxruntime/core/providers/cpu/nn/pool_base.h +++ b/onnxruntime/core/providers/cpu/nn/pool_base.h @@ -50,11 +50,11 @@ class AveragePool { static const PoolType type = PoolType::kAveragePool; }; -template +template class MaxPool; template <> -class MaxPool<1 /*VERSION*/> { +class MaxPool<1 /*START_VERSION*/> { public: static float Initialize() { return std::numeric_limits::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("kernel_shape", kernel_shape_).IsOK(), @@ -123,8 +125,12 @@ class PoolBase { ceil_mode_ = 0; } + default_dilations_ = false; if (!info.GetAttrs("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("storage_order", 0 /*default_value*/); } } @@ -261,6 +264,9 @@ class PoolBase { std::vector pads_; std::vector strides_; std::vector 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_; diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc index fc0767efff..ae53db2321 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc @@ -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, BuildKernelCreateInfo, BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -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( diff --git a/onnxruntime/core/providers/cuda/nn/max_pool_with_index.cu b/onnxruntime/core/providers/cuda/nn/max_pool_with_index.cu index 29b83ecdef..f61c454262 100644 --- a/onnxruntime/core/providers/cuda/nn/max_pool_with_index.cu +++ b/onnxruntime/core/providers/cuda/nn/max_pool_with_index.cu @@ -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(d_start + kernel_d, depth); - int64_t w_end = _Min(w_start + kernel_w, width); - int64_t h_end = _Min(h_start + kernel_h, height); + int64_t d_end = _Min(d_start + (kernel_d - 1) * dilation_d + 1, depth); + int64_t w_end = _Min(w_start + (kernel_w - 1) * dilation_w + 1, width); + int64_t h_end = _Min(h_start + (kernel_h - 1) * dilation_h + 1, height); d_start = _Max(d_start, 0); w_start = _Max(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 @@ -88,6 +93,7 @@ void MaxPoolWithIndex( const std::vector& kernel_shape, const std::vector& stride_shape, const std::vector& pads, + const std::vector& 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(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& kernel_shape, \ const std::vector& stride_shape, \ const std::vector& pads, \ + const std::vector& dilations, \ int64_t storage_order, \ const T* p_input, \ T* p_output, \ diff --git a/onnxruntime/core/providers/cuda/nn/max_pool_with_index.h b/onnxruntime/core/providers/cuda/nn/max_pool_with_index.h index 60fa7f806b..ec796c3d95 100644 --- a/onnxruntime/core/providers/cuda/nn/max_pool_with_index.h +++ b/onnxruntime/core/providers/cuda/nn/max_pool_with_index.h @@ -14,6 +14,7 @@ void MaxPoolWithIndex( const std::vector& kernel_shape, const std::vector& stride_shape, const std::vector& pads, + const std::vector& dilations, int64_t storage_order, const T* p_input, T* p_output, diff --git a/onnxruntime/core/providers/cuda/nn/pool.cc b/onnxruntime/core/providers/cuda/nn/pool.cc index 1d48b28b82..e26d3c4ff9 100644 --- a/onnxruntime/core/providers/cuda/nn/pool.cc +++ b/onnxruntime/core/providers/cuda/nn/pool.cc @@ -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>::ComputeInternal(OpKernelContext* context) const { auto y_data = reinterpret_cast(Y->template MutableData()); Tensor* I = context->Output(1, TensorShape(y_dims)); - if (nullptr != I) { - auto i_data = I->template MutableData(); + if (nullptr != I || !this->default_dilations_) { + auto i_data = nullptr == I ? nullptr : I->template MutableData(); MaxPoolWithIndex( x_shape, TensorShape(y_dims), kernel_shape, strides, pads, + this->dilations_, this->storage_order_, x_data, y_data,