diff --git a/cgmanifest.json b/cgmanifest.json index b6f52eca84..0adbebf636 100644 --- a/cgmanifest.json +++ b/cgmanifest.json @@ -49,7 +49,7 @@ "component": { "type": "git", "git": { - "commitHash": "1facb4c1bb9cc2107d4dbaf9fd647fefdbbeb0ab", + "commitHash": "423f1977d314f05df4be7edb44428c1c0211341c", "repositoryUrl": "https://github.com/onnx/onnx.git" } } diff --git a/cmake/external/onnx b/cmake/external/onnx index 1facb4c1bb..423f1977d3 160000 --- a/cmake/external/onnx +++ b/cmake/external/onnx @@ -1 +1 @@ -Subproject commit 1facb4c1bb9cc2107d4dbaf9fd647fefdbbeb0ab +Subproject commit 423f1977d314f05df4be7edb44428c1c0211341c diff --git a/onnxruntime/core/providers/cpu/cpu_execution_provider.cc b/onnxruntime/core/providers/cpu/cpu_execution_provider.cc index 5169f6fdec..162b268e7e 100644 --- a/onnxruntime/core/providers/cpu/cpu_execution_provider.cc +++ b/onnxruntime/core/providers/cpu/cpu_execution_provider.cc @@ -124,8 +124,8 @@ class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 1, LpN class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 1, LRN); class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 7, 9, AveragePool); class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 1, 7, MaxPool); -class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, 11, float, MaxPool); -class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, 11, double, MaxPool); +class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, 11, MaxPool); +class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, MaxPool); class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 2, 10, LpPool); class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 2, GlobalLpPool); class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 1, GlobalAveragePool); @@ -331,10 +331,10 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, double, Round); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, MLFloat16, Round); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, uint8_t, DynamicQuantizeLinear); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, float, ArgMax); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, int32_t, ArgMax); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, float, ArgMin); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, int32_t, ArgMin); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 11, float, ArgMax); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 11, int32_t, ArgMax); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 11, float, ArgMin); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 11, int32_t, ArgMin); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, float, ReduceL1); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, int32_t, ReduceL1); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, float, ReduceL2); @@ -435,6 +435,28 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, int32_t, Resize); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, uint8_t, Resize); +// opset 12 +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, float, ArgMax); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int32_t, ArgMax); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, float, ArgMin); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int32_t, ArgMin); + +class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, MaxPool); + +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, float, ReduceMax); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int32_t, ReduceMax); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int64_t, ReduceMax); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int8_t, ReduceMax); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, uint8_t, ReduceMax); + +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, float, ReduceMin); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int32_t, ReduceMin); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int64_t, ReduceMin); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int8_t, ReduceMin); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, uint8_t, ReduceMin); + + + Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) { static const BuildKernelCreateInfoFn function_table[] = { BuildKernelCreateInfo, BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -1001,14 +1021,14 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) { Round)>, BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -1111,6 +1131,40 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) { int32_t, Resize)>, BuildKernelCreateInfo, + + // OpSet 12 + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + + BuildKernelCreateInfo, + + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, }; for (auto& function_table_entry : function_table) { diff --git a/onnxruntime/core/providers/cpu/nn/pool.cc b/onnxruntime/core/providers/cpu/nn/pool.cc index 2e0e3b5f9e..467cda1a38 100644 --- a/onnxruntime/core/providers/cpu/nn/pool.cc +++ b/onnxruntime/core/providers/cpu/nn/pool.cc @@ -2,6 +2,7 @@ // Licensed under the MIT License. #include "core/providers/cpu/nn/pool.h" +#include "core/framework/data_types_internal.h" #include "core/platform/threadpool.h" #include "core/util/eigen_common_wrapper.h" #include "pool_functors.h" @@ -11,7 +12,7 @@ using namespace ::onnxruntime::common; namespace onnxruntime { template -static void RunLoop(concurrency::ThreadPool* tp, Eigen::Index total_channels, T&& task) { +inline static void RunLoop(concurrency::ThreadPool* tp, Eigen::Index total_channels, T&& task) { concurrency::ThreadPool::TryParallelFor(tp, total_channels, task.Cost(), task); } @@ -134,13 +135,30 @@ Status Pool::Compute(OpKernelContext* context) const { // version 8: Added storage_order And Indices // version 10: Added ceil_mode // version 11: Added dilations -template +// version 12: Added int8/uint8 support + class MaxPoolV8 : public OpKernel, public PoolBase { + + template + struct ComputeHelper { + Status operator()(const MaxPoolV8* inst, OpKernelContext* context) const { + return inst->ComputeImpl(context); + } + }; + public: - MaxPoolV8(const OpKernelInfo& info) : OpKernel(info), PoolBase(info) { + explicit MaxPoolV8(const OpKernelInfo& info) : OpKernel(info), PoolBase(info) { } Status Compute(OpKernelContext* context) const override { + utils::MLTypeCallDispatcherRet + t_disp(context->Input(0)->GetElementType()); + return t_disp.Invoke(this, context); + } + + private: + template + Status ComputeImpl(OpKernelContext* context) const { concurrency::ThreadPool* tp = context->GetOperatorThreadPool(); // Use MLAS pooling if the index output tensor is not used // and also if dilation is not required @@ -150,8 +168,11 @@ class MaxPoolV8 : public OpKernel, public PoolBase { need_dilation |= n > 1; } - if (OpKernel::Node().OutputDefs().size() == 1 && pool_attrs_.storage_order == 0 && !need_dilation) { - return PoolBase::Compute(context, MlasMaximumPooling); + // MLAS implementation currently supports only floats + if (std::is_same::value) { + if (OpKernel::Node().OutputDefs().size() == 1 && pool_attrs_.storage_order == 0 && !need_dilation) { + return PoolBase::Compute(context, MlasMaximumPooling); + } } const auto* X = context->Input(0); @@ -210,10 +231,10 @@ class MaxPoolV8 : public OpKernel, public PoolBase { const int64_t dilation_w = pool_attrs_.dilations[1]; const int64_t dilation_d = pool_attrs_.dilations[2]; RunLoop>(tp, total_channels, - {X_data, Y_data, I_data, x_step, y_step, - dilation_h, dilation_w, dilation_d, pooled_height, pooled_width, - pooled_depth, stride_h(), stride_w(), stride_d(), height, - width, depth, kernel_shape, pads, pool_attrs_.storage_order}); + {X_data, Y_data, I_data, x_step, y_step, + dilation_h, dilation_w, dilation_d, pooled_height, pooled_width, + pooled_depth, stride_h(), stride_w(), stride_d(), height, + width, depth, kernel_shape, pads, pool_attrs_.storage_order}); break; } default: @@ -222,9 +243,6 @@ class MaxPoolV8 : public OpKernel, public PoolBase { return Status::OK(); } - - private: - PoolProcessContext pool_context_; }; ONNX_CPU_OPERATOR_VERSIONED_KERNEL(AveragePool, 7, 9, @@ -242,17 +260,21 @@ ONNX_CPU_OPERATOR_VERSIONED_KERNEL(MaxPool, 1, 7, KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType()), Pool>); -ONNX_CPU_OPERATOR_VERSIONED_TYPED_KERNEL(MaxPool, 8, 11, float, +ONNX_CPU_OPERATOR_VERSIONED_KERNEL(MaxPool, 8, 11, KernelDefBuilder() - .TypeConstraint("T", DataTypeImpl::GetTensorType()) + .TypeConstraint("T", {DataTypeImpl::GetTensorType(), + DataTypeImpl::GetTensorType()}) .TypeConstraint("I", DataTypeImpl::GetTensorType()), - MaxPoolV8); + MaxPoolV8); -ONNX_CPU_OPERATOR_VERSIONED_TYPED_KERNEL(MaxPool, 8, 11, double, - KernelDefBuilder() - .TypeConstraint("T", DataTypeImpl::GetTensorType()) - .TypeConstraint("I", DataTypeImpl::GetTensorType()), - MaxPoolV8); +ONNX_CPU_OPERATOR_KERNEL(MaxPool, 12, + KernelDefBuilder() + .TypeConstraint("T", {DataTypeImpl::GetTensorType(), + DataTypeImpl::GetTensorType(), + DataTypeImpl::GetTensorType(), + DataTypeImpl::GetTensorType()}) + .TypeConstraint("I", DataTypeImpl::GetTensorType()), + MaxPoolV8); ONNX_CPU_OPERATOR_VERSIONED_KERNEL(LpPool, 2, 10, KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType()), diff --git a/onnxruntime/core/providers/cpu/nn/pool.h b/onnxruntime/core/providers/cpu/nn/pool.h index 911339f736..5fcb1d5f0d 100644 --- a/onnxruntime/core/providers/cpu/nn/pool.h +++ b/onnxruntime/core/providers/cpu/nn/pool.h @@ -24,5 +24,4 @@ class Pool : public OpKernel, public PoolBase { private: PoolProcessContext pool_context_; }; - } // namespace onnxruntime diff --git a/onnxruntime/core/providers/cpu/reduction/reduction_ops.cc b/onnxruntime/core/providers/cpu/reduction/reduction_ops.cc index eda7948f70..c8d38518ea 100644 --- a/onnxruntime/core/providers/cpu/reduction/reduction_ops.cc +++ b/onnxruntime/core/providers/cpu/reduction/reduction_ops.cc @@ -74,6 +74,22 @@ namespace onnxruntime { KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType()), \ x); +#define REGISTER_UNARY_ELEMENTWISE_KERNEL_INT8_ONLY(x, sinceVersion) \ + ONNX_CPU_OPERATOR_TYPED_KERNEL( \ + x, \ + sinceVersion, \ + int8_t, \ + KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType()), \ + x); + +#define REGISTER_UNARY_ELEMENTWISE_KERNEL_UINT8_ONLY(x, sinceVersion) \ + ONNX_CPU_OPERATOR_TYPED_KERNEL( \ + x, \ + sinceVersion, \ + uint8_t, \ + KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType()), \ + x); + REGISTER_UNARY_ELEMENTWISE_VERSIONED_KERNEL(ReduceL1, 1, 10); REGISTER_UNARY_ELEMENTWISE_KERNEL(ReduceL1, 11); @@ -91,6 +107,11 @@ REGISTER_UNARY_ELEMENTWISE_KERNEL(ReduceMax, 11); REGISTER_UNARY_ELEMENTWISE_VERSIONED_KERNEL_INT64_ONLY(ReduceMax, 1, 10); REGISTER_UNARY_ELEMENTWISE_KERNEL_INT64_ONLY(ReduceMax, 11); +REGISTER_UNARY_ELEMENTWISE_KERNEL(ReduceMax, 12); +REGISTER_UNARY_ELEMENTWISE_KERNEL_INT64_ONLY(ReduceMax, 12); +REGISTER_UNARY_ELEMENTWISE_KERNEL_INT8_ONLY(ReduceMax, 12); +REGISTER_UNARY_ELEMENTWISE_KERNEL_UINT8_ONLY(ReduceMax, 12); + REGISTER_UNARY_ELEMENTWISE_VERSIONED_KERNEL(ReduceMean, 1, 10); REGISTER_UNARY_ELEMENTWISE_KERNEL(ReduceMean, 11); @@ -99,6 +120,11 @@ REGISTER_UNARY_ELEMENTWISE_KERNEL(ReduceMin, 11); REGISTER_UNARY_ELEMENTWISE_VERSIONED_KERNEL_INT64_ONLY(ReduceMin, 1, 10); REGISTER_UNARY_ELEMENTWISE_KERNEL_INT64_ONLY(ReduceMin, 11); +REGISTER_UNARY_ELEMENTWISE_KERNEL(ReduceMin, 12); +REGISTER_UNARY_ELEMENTWISE_KERNEL_INT64_ONLY(ReduceMin, 12); +REGISTER_UNARY_ELEMENTWISE_KERNEL_INT8_ONLY(ReduceMin, 12); +REGISTER_UNARY_ELEMENTWISE_KERNEL_UINT8_ONLY(ReduceMin, 12); + REGISTER_UNARY_ELEMENTWISE_VERSIONED_KERNEL(ReduceProd, 1, 10); REGISTER_UNARY_ELEMENTWISE_KERNEL(ReduceProd, 11); @@ -115,12 +141,15 @@ REGISTER_UNARY_ELEMENTWISE_VERSIONED_KERNEL_DOUBLE_ONLY(ReduceSumSquare, 1, 10); REGISTER_UNARY_ELEMENTWISE_KERNEL_DOUBLE_ONLY(ReduceSumSquare, 11); REGISTER_UNARY_ELEMENTWISE_VERSIONED_KERNEL(ArgMax, 1, 10); -REGISTER_UNARY_ELEMENTWISE_KERNEL(ArgMax, 11); +REGISTER_UNARY_ELEMENTWISE_VERSIONED_KERNEL(ArgMax, 11, 11); +REGISTER_UNARY_ELEMENTWISE_KERNEL(ArgMax, 12); REGISTER_UNARY_ELEMENTWISE_VERSIONED_KERNEL(ArgMin, 1, 10); -REGISTER_UNARY_ELEMENTWISE_KERNEL(ArgMin, 11); +REGISTER_UNARY_ELEMENTWISE_VERSIONED_KERNEL(ArgMin, 11, 11); +REGISTER_UNARY_ELEMENTWISE_KERNEL(ArgMin, 12); // When all reduce axises located at the tail of the dims, quite general cases, transpose and extra -// copy could be skipped to improve performance, if required by check_no_transpose = true; +// copy could be skipped to improve performance. If required by check_no_transpose = true, then +// the calling code will check if the data was transposed and act accordingly. // return value: true means transposedInputData is not created/copied, input tensor data could // be direct use as row major matrix [block_size, blocks], where blocks is the // size of each reduce. @@ -189,9 +218,9 @@ bool PrepareForReduce(OpKernelContext* ctx, } } - std::vector new_dims_(transposed_axes.size()); + std::vector new_dims(transposed_axes.size()); for (size_t i = 0; i < transposed_axes.size(); ++i) { - new_dims_[i] = input.Shape().GetDims().at(transposed_axes[i]); + new_dims[i] = input.Shape().GetDims().at(transposed_axes[i]); } int num_axes = static_cast(transposed_axes.size()); @@ -202,7 +231,7 @@ bool PrepareForReduce(OpKernelContext* ctx, int n_shared_idxs = 0; for (int i = num_axes - 1; i >= 0; --i) { if (transposed_axes[i] == i) { - blocksize *= new_dims_[i]; + blocksize *= new_dims[i]; ++n_shared_idxs; } else { break; @@ -290,7 +319,7 @@ bool PrepareForReduce(OpKernelContext* ctx, ++itr_idxs[itr_axes - 1]; for (int i = itr_axes - 1; i >= 1; --i) { - auto expected_dim = new_dims_[i]; + auto expected_dim = new_dims[i]; if (itr_idxs[i] < expected_dim) { break; } @@ -309,7 +338,7 @@ bool PrepareForReduce(OpKernelContext* ctx, ++itr_idxs[itr_axes - 1]; for (int i = itr_axes - 1; i >= 1; --i) { - auto expected_dim = new_dims_[i]; + auto expected_dim = new_dims[i]; if (itr_idxs[i] < expected_dim) { break; } @@ -591,16 +620,47 @@ Status ArgMax::Compute(OpKernelContext* ctx) const { if (no_transpose) { const T* input_data = ctx->Input(0)->template Data(); - - for (int64_t i = 0; i < block_size; ++i) { - ConstEigenVectorMap(input_data + (i * blocks), blocks).maxCoeff(&maxIndex); - *(output_data++) = maxIndex; + if (select_last_index_) { + assert(blocks > 0); + for (int64_t i = 0; i < block_size; ++i) { + gsl::span row(input_data, blocks); + auto first = row.cbegin(); + auto const end = row.cend(); + auto max_el = first; + while (++first < end) { + if (*first >= *max_el) { + max_el = first; + } + } + *(output_data++) = max_el - row.cbegin(); + input_data += blocks; + } + } else { + for (int64_t i = 0; i < block_size; ++i) { + ConstEigenVectorMap(input_data + (i * blocks), blocks).maxCoeff(&maxIndex); + *(output_data++) = maxIndex; + } } } else { auto matrixData = ConstEigenMatrixMap(&transposedInputData[0], block_size, blocks); - for (int i = 0; i < block_size; ++i) { - matrixData.row(i).maxCoeff(&maxIndex); - *(output_data++) = maxIndex; + if (select_last_index_) { + for (int i = 0; i < block_size; ++i) { + int idx = 0; + T max_val = matrixData(i, 0); + for (int c = 1; c < blocks; ++c) { + auto val = matrixData(i, c); + if (val >= max_val) { + idx = c; + max_val = val; + } + } + *(output_data++) = idx; + } + } else { + for (int i = 0; i < block_size; ++i) { + matrixData.row(i).maxCoeff(&maxIndex); + *(output_data++) = maxIndex; + } } } @@ -621,16 +681,47 @@ Status ArgMin::Compute(OpKernelContext* ctx) const { if (no_transpose) { const T* input_data = ctx->Input(0)->template Data(); - - for (int64_t i = 0; i < block_size; ++i) { - ConstEigenVectorMap(input_data + (i * blocks), blocks).minCoeff(&minIndex); - *(output_data++) = minIndex; + if (select_last_index_) { + assert(blocks > 0); + for (int64_t i = 0; i < block_size; ++i) { + gsl::span row(input_data, blocks); + auto first = row.cbegin(); + auto const end = row.cend(); + auto min_el = first; + while (++first < end) { + if (*first <= *min_el) { + min_el = first; + } + } + *(output_data++) = min_el - row.cbegin(); + input_data += blocks; + } + } else { + for (int64_t i = 0; i < block_size; ++i) { + ConstEigenVectorMap(input_data + (i * blocks), blocks).minCoeff(&minIndex); + *(output_data++) = minIndex; + } } } else { auto matrixData = ConstEigenMatrixMap(&transposedInputData[0], block_size, blocks); - for (int i = 0; i < block_size; ++i) { - matrixData.row(i).minCoeff(&minIndex); - *(output_data++) = minIndex; + if (select_last_index_) { + for (int i = 0; i < block_size; ++i) { + int idx = 0; + T min_val = matrixData(i, 0); + for (int c = 1; c < blocks; ++c) { + auto val = matrixData(i, c); + if (val <= min_val) { + idx = c; + min_val = val; + } + } + *(output_data++) = idx; + } + } else { + for (int i = 0; i < block_size; ++i) { + matrixData.row(i).minCoeff(&minIndex); + *(output_data++) = minIndex; + } } } diff --git a/onnxruntime/core/providers/cpu/reduction/reduction_ops.h b/onnxruntime/core/providers/cpu/reduction/reduction_ops.h index 7ee64149f0..1afcb69399 100644 --- a/onnxruntime/core/providers/cpu/reduction/reduction_ops.h +++ b/onnxruntime/core/providers/cpu/reduction/reduction_ops.h @@ -27,10 +27,13 @@ class ReduceKernelBase { ORT_ENFORCE(info.GetAttr("keepdims", &keepdims).IsOK()); } keepdims_ = (keepdims == 1); + int64_t select_last_index = info.GetAttrOrDefault("select_last_index", 0); + select_last_index_ = (select_last_index != 0); } std::vector axes_; bool keepdims_; + bool select_last_index_; }; template diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc index 73ab2f999d..55a7650580 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc @@ -638,18 +638,18 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, ReduceLogSumExp); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, ReduceLogSumExp); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, ReduceLogSumExp); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, ReduceMax); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, ReduceMax); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, ReduceMax); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, int32_t, ReduceMax); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, float, ReduceMax); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, double, ReduceMax); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, MLFloat16, ReduceMax); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, int32_t, ReduceMax); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, ReduceMean); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, ReduceMean); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, ReduceMean); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, int32_t, ReduceMean); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, ReduceMin); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, ReduceMin); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, ReduceMin); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, int32_t, ReduceMin); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, float, ReduceMin); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, double, ReduceMin); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, MLFloat16, ReduceMin); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, int32_t, ReduceMin); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, ReduceProd); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, ReduceProd); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, ReduceProd); @@ -682,9 +682,9 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, AveragePool); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, AveragePool); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, AveragePool); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, MaxPool); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, MaxPool); -class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, MaxPool); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, float, MaxPool); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, double, MaxPool); +class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, MLFloat16, MaxPool); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, Resize); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, Resize); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, Resize); @@ -706,6 +706,28 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, uint8_t, DequantizeLinear); class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, CumSum); +// OpSet 12 +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, MaxPool); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, MaxPool); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, MaxPool); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int8_t, MaxPool); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint8_t, MaxPool); + +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, ReduceMax); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, ReduceMax); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, ReduceMax); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int32_t, ReduceMax); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int8_t, ReduceMax); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint8_t, ReduceMax); + +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, ReduceMin); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, ReduceMin); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, ReduceMin); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int32_t, ReduceMin); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int8_t, ReduceMin); +class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint8_t, ReduceMin); + + static void RegisterCudaKernels(KernelRegistry& kernel_registry) { static const BuildKernelCreateInfoFn function_table[] = { BuildKernelCreateInfo, @@ -1075,6 +1097,10 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, // opset 11 BuildKernelCreateInfo, @@ -1109,18 +1135,18 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -1153,9 +1179,9 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -1171,11 +1197,29 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, - BuildKernelCreateInfo, BuildKernelCreateInfo, + + // OpSet 12 + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + }; for (auto& function_table_entry : function_table) { diff --git a/onnxruntime/core/providers/cuda/cudnn_common.cc b/onnxruntime/core/providers/cuda/cudnn_common.cc index 089313a383..d0d50718ec 100644 --- a/onnxruntime/core/providers/cuda/cudnn_common.cc +++ b/onnxruntime/core/providers/cuda/cudnn_common.cc @@ -47,18 +47,6 @@ Status CudnnTensor::Set(const CudnnTensor& x_desc, cudnnBatchNormMode_t mode) { return Status::OK(); } -template -cudnnDataType_t CudnnTensor::GetDataType() { - if (typeid(ElemType) == typeid(float)) - return CUDNN_DATA_FLOAT; - else if (typeid(ElemType) == typeid(double)) - return CUDNN_DATA_DOUBLE; - else if (typeid(ElemType) == typeid(half)) - return CUDNN_DATA_HALF; - else - ORT_THROW("cuDNN engine currently supports only single/double/half precision data types."); -} - CudnnDataTensor::CudnnDataTensor() : tensor_(nullptr) { } @@ -124,9 +112,38 @@ Status CudnnFilterDescriptor::Set(const std::vector& filter_dims, cudnn return Status::OK(); } -template cudnnDataType_t CudnnTensor::GetDataType(); -template cudnnDataType_t CudnnTensor::GetDataType(); -template cudnnDataType_t CudnnTensor::GetDataType(); +template +cudnnDataType_t CudnnTensor::GetDataType() { + ORT_THROW("cuDNN engine currently supports only single/double/half/int8/uint8 precision data types. Got:", + typeid(ElemType).name()); + // Not reachable but GCC complains + return 0; +} + +template<> +cudnnDataType_t CudnnTensor::GetDataType() { + return CUDNN_DATA_FLOAT; +} + +template <> +cudnnDataType_t CudnnTensor::GetDataType() { + return CUDNN_DATA_DOUBLE; +} + +template <> +cudnnDataType_t CudnnTensor::GetDataType() { + return CUDNN_DATA_HALF; +} + +template <> +cudnnDataType_t CudnnTensor::GetDataType() { + return CUDNN_DATA_INT8; +} + +template <> +cudnnDataType_t CudnnTensor::GetDataType() { + return CUDNN_DATA_UINT8; +} template <> const float Consts::One = 1; @@ -144,6 +161,18 @@ const float Consts::Zero = 0; const float Consts::One = 1; +template <> +const int8_t Consts::Zero = 0; + +template <> +const int8_t Consts::One = 1; + +template <> +const uint8_t Consts::Zero = 0; + +template <> +const uint8_t Consts::One = 1; + std::vector DeviceProp::s_cachedDeviceProps; std::once_flag DeviceProp::s_cachedDevicePropsInitFlag; diff --git a/onnxruntime/core/providers/cuda/cudnn_common.h b/onnxruntime/core/providers/cuda/cudnn_common.h index bfd233b68b..f5c3372c09 100644 --- a/onnxruntime/core/providers/cuda/cudnn_common.h +++ b/onnxruntime/core/providers/cuda/cudnn_common.h @@ -13,6 +13,7 @@ class CudnnTensor final { public: CudnnTensor(); ~CudnnTensor(); + ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(CudnnTensor); Status Set(const std::vector& input_dims, cudnnDataType_t dataType); Status Set(const CudnnTensor& x_desc, cudnnBatchNormMode_t mode); @@ -32,6 +33,7 @@ class CudnnDataTensor final { public: CudnnDataTensor(); ~CudnnDataTensor(); + ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(CudnnDataTensor); Status Set(cudnnDataType_t dataType, int64_t max_seq_length, @@ -51,6 +53,7 @@ class CudnnFilterDescriptor final { public: CudnnFilterDescriptor(); ~CudnnFilterDescriptor(); + ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(CudnnFilterDescriptor); Status Set(const std::vector& filter_dims, cudnnDataType_t data_typ); @@ -64,6 +67,7 @@ class CudnnDropout final { public: CudnnDropout() : dropout_desc_(nullptr) { } + ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(CudnnDropout); Status GetCudnnDropoutStatesSize(const cudnnHandle_t& cudnnHandle, size_t& stateSize) { CUDNN_RETURN_IF_ERROR(cudnnDropoutGetStatesSize(cudnnHandle, &stateSize)); 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 f61c454262..02335cad1d 100644 --- a/onnxruntime/core/providers/cuda/nn/max_pool_with_index.cu +++ b/onnxruntime/core/providers/cuda/nn/max_pool_with_index.cu @@ -98,7 +98,6 @@ void MaxPoolWithIndex( const T* p_input, T* p_output, int64_t* p_indices) { - size_t input_dim_count = input_shape.NumDimensions(); int64_t batchs = input_shape[0]; int64_t channels = input_shape[1]; @@ -179,6 +178,8 @@ void MaxPoolWithIndex( INSTANTIATEMAXPOOLWITHINDEX(float) INSTANTIATEMAXPOOLWITHINDEX(double) INSTANTIATEMAXPOOLWITHINDEX(half) +INSTANTIATEMAXPOOLWITHINDEX(int8_t) +INSTANTIATEMAXPOOLWITHINDEX(uint8_t) } // namespace cuda } // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/nn/pool.cc b/onnxruntime/core/providers/cuda/nn/pool.cc index 8dbdeca70c..367930d9b3 100644 --- a/onnxruntime/core/providers/cuda/nn/pool.cc +++ b/onnxruntime/core/providers/cuda/nn/pool.cc @@ -5,6 +5,8 @@ #include "core/providers/common.h" #include "core/providers/cuda/cudnn_common.h" #include "core/providers/cuda/nn/max_pool_with_index.h" +#include "core/providers/cuda/math/unary_elementwise_ops_impl.h" + using namespace onnxruntime::common; namespace onnxruntime { namespace cuda { @@ -53,9 +55,16 @@ POOLING_KERNEL_VERSIONED(MaxPool, MLFloat16, MaxPool<8>, 8, 9) POOLING_KERNEL_VERSIONED(MaxPool, float, MaxPool<8>, 10, 10) POOLING_KERNEL_VERSIONED(MaxPool, double, MaxPool<8>, 10, 10) POOLING_KERNEL_VERSIONED(MaxPool, MLFloat16, MaxPool<8>, 10, 10) -POOLING_KERNEL(MaxPool, float, MaxPool<8>, 11) -POOLING_KERNEL(MaxPool, double, MaxPool<8>, 11) -POOLING_KERNEL(MaxPool, MLFloat16, MaxPool<8>, 11) +POOLING_KERNEL_VERSIONED(MaxPool, float, MaxPool<8>, 11, 11) +POOLING_KERNEL_VERSIONED(MaxPool, double, MaxPool<8>, 11, 11) +POOLING_KERNEL_VERSIONED(MaxPool, MLFloat16, MaxPool<8>, 11, 11) +POOLING_KERNEL(MaxPool, float, MaxPool<8>, 12) +POOLING_KERNEL(MaxPool, double, MaxPool<8>, 12) +POOLING_KERNEL(MaxPool, MLFloat16, MaxPool<8>, 12) +POOLING_KERNEL(MaxPool, int8_t, MaxPool<8>, 12) +POOLING_KERNEL(MaxPool, uint8_t, MaxPool<8>, 12) + + POOLING_KERNEL(GlobalMaxPool, float, MaxPool<1>, 1) POOLING_KERNEL(GlobalMaxPool, double, MaxPool<1>, 1) POOLING_KERNEL(GlobalMaxPool, MLFloat16, MaxPool<1>, 1) @@ -72,6 +81,9 @@ class CudnnPoolingDescriptor final { } } + CudnnPoolingDescriptor(const CudnnPoolingDescriptor&) = delete; + CudnnPoolingDescriptor& operator=(const CudnnPoolingDescriptor&) = delete; + Status Set(cudnnPoolingMode_t mode, const std::vector& kernel_shape, const std::vector& pads, @@ -132,9 +144,10 @@ Status Pool::ComputeInternal(OpKernelContext* context) const { } std::vector y_dims = pool_attrs_.SetOutputSize(x_shape, x_shape[1], &pads); - Tensor* Y = context->Output(0, TensorShape(y_dims)); + TensorShape y_shape(y_dims); + Tensor* Y = context->Output(0, y_shape); // special case when there is a dim value of 0 in the shape. - if (Y->Shape().Size() == 0) + if (y_shape.Size() == 0) return Status::OK(); auto x_data = reinterpret_cast(X->template Data()); @@ -152,13 +165,6 @@ Status Pool::ComputeInternal(OpKernelContext* context) const { strides.push_back(1); } - const auto alpha = Consts::One; - const auto beta = Consts::Zero; - CudnnTensor x_tensor; - CudnnTensor y_tensor; - ORT_RETURN_IF_ERROR(x_tensor.Set(x_dims_cudnn, CudnnTensor::GetDataType())); - ORT_RETURN_IF_ERROR(y_tensor.Set(y_dims_cudnn, CudnnTensor::GetDataType())); - cudnnPoolingMode_t mode = CUDNN_POOLING_MAX; if (PoolType::type == onnxruntime::PoolType::kAveragePool) { mode = pool_attrs_.count_include_pad ? CUDNN_POOLING_AVERAGE_COUNT_INCLUDE_PADDING @@ -167,7 +173,33 @@ Status Pool::ComputeInternal(OpKernelContext* context) const { CudnnPoolingDescriptor pooling_desc; ORT_RETURN_IF_ERROR(pooling_desc.Set(mode, kernel_shape, pads, strides)); - CUDNN_RETURN_IF_ERROR(cudnnPoolingForward(CudnnHandle(), pooling_desc, &alpha, x_tensor, x_data, &beta, y_tensor, y_data)); + if (std::is_same::value || std::is_same::value) { + // Cast to float back and forth using temp buffer + const auto alpha = Consts::One; + const auto beta = Consts::Zero; + CudnnTensor x_tensor; + CudnnTensor y_tensor; + ORT_RETURN_IF_ERROR(x_tensor.Set(x_dims_cudnn, CudnnTensor::GetDataType())); + ORT_RETURN_IF_ERROR(y_tensor.Set(y_dims_cudnn, CudnnTensor::GetDataType())); + + const auto input_count = x_shape.Size(); + const auto output_count = y_shape.Size(); + + IAllocatorUniquePtr temp_X = GetScratchBuffer(input_count); + auto temp_Y = GetScratchBuffer(output_count); + Impl_Cast(reinterpret_cast(x_data), temp_X.get(), input_count); + CUDNN_RETURN_IF_ERROR(cudnnPoolingForward(CudnnHandle(), pooling_desc, &alpha, x_tensor, temp_X.get(), &beta, y_tensor, temp_Y.get())); + Impl_Cast(temp_Y.get(), y_data, output_count); + } else { + const auto alpha = Consts::One; + const auto beta = Consts::Zero; + CudnnTensor x_tensor; + CudnnTensor y_tensor; + ORT_RETURN_IF_ERROR(x_tensor.Set(x_dims_cudnn, CudnnTensor::GetDataType())); + ORT_RETURN_IF_ERROR(y_tensor.Set(y_dims_cudnn, CudnnTensor::GetDataType())); + + CUDNN_RETURN_IF_ERROR(cudnnPoolingForward(CudnnHandle(), pooling_desc, &alpha, x_tensor, x_data, &beta, y_tensor, y_data)); + } return Status::OK(); } diff --git a/onnxruntime/core/providers/cuda/reduction/reduction_ops.cc b/onnxruntime/core/providers/cuda/reduction/reduction_ops.cc index ea46326b59..f080672777 100644 --- a/onnxruntime/core/providers/cuda/reduction/reduction_ops.cc +++ b/onnxruntime/core/providers/cuda/reduction/reduction_ops.cc @@ -32,6 +32,34 @@ namespace cuda { KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType()), \ name); +// Register with the latest version 12 +#define REGISTER_KERNEL_TYPED_12(name, T) \ + ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \ + name, \ + kOnnxDomain, \ + 1, 10, \ + T, \ + kCudaExecutionProvider, \ + KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType()), \ + name); \ + ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \ + name, \ + kOnnxDomain, \ + 11, 11, \ + T, \ + kCudaExecutionProvider, \ + KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType()), \ + name); \ + ONNX_OPERATOR_TYPED_KERNEL_EX( \ + name, \ + kOnnxDomain, \ + 12, \ + T, \ + kCudaExecutionProvider, \ + KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType()), \ + name); + + // CUDA's reduction descriptor cudnnReduceTensorDescriptor_t is a pointer so // it's safer to wrap it with automatically memory deleter as CudnnReduceDescriptor. // An implicit caster from CudnnReduceDescriptor to cudnnReduceTensorDescriptor_t @@ -85,19 +113,18 @@ Status ReduceKernel::ReduceKernelShared( cudnnDataType_t cudnn_type_X = CudnnTensor::GetDataType(); const auto rank = input_shape.NumDimensions(); -// Block of fast matrix row reduction. + // Block of fast matrix row reduction. const auto stride = input_shape[input_shape.NumDimensions() - 1]; const auto reduction_size = input_shape.Size() / stride; if (fast_reduction_ && reduction_size <= std::numeric_limits::max() && stride <= std::numeric_limits::max() && is_matrix_row_reduction(cudnn_reduce_op, - static_cast(reduction_size), - static_cast(stride), rank, axes_)) { - + static_cast(reduction_size), + static_cast(stride), rank, axes_)) { reduce_matrix_rows( - reinterpret_cast(X), - reinterpret_cast(Y), - static_cast(reduction_size), - static_cast(stride)); + reinterpret_cast(X), + reinterpret_cast(Y), + static_cast(reduction_size), + static_cast(stride)); return Status::OK(); } @@ -402,17 +429,18 @@ Status ReduceKernel::ComputeImpl(OpKernelContext* ctx, cudnnRe // Block of fast matrix row reduction. // It relies on new atomicAdd for half type, so old CUDA can't use it. const auto reduction_size = input_count / stride; - if (fast_reduction_ && reduction_size <= std::numeric_limits::max() && stride <= std::numeric_limits::max() && - is_matrix_row_reduction(cudnn_reduce_op, - static_cast(reduction_size), - static_cast(stride), rank, axes_)) { - - reduce_matrix_rows( - reinterpret_cast(X->template Data()), - reinterpret_cast(Y->template MutableData()), - static_cast(reduction_size), - static_cast(stride)); - return Status::OK(); + if (!std::is_same::value && !std::is_same::value) { + if (fast_reduction_ && reduction_size <= std::numeric_limits::max() && stride <= std::numeric_limits::max() && + is_matrix_row_reduction(cudnn_reduce_op, + static_cast(reduction_size), + static_cast(stride), rank, axes_)) { + reduce_matrix_rows( + reinterpret_cast(X->template Data()), + reinterpret_cast(Y->template MutableData()), + static_cast(reduction_size), + static_cast(stride)); + return Status::OK(); + } } if (ReduceTensorIndices == CUDNN_REDUCE_TENSOR_FLATTENED_INDICES && std::is_same::value) { @@ -498,7 +526,7 @@ Status ReduceKernel::ComputeImpl(OpKernelContext* ctx, cudnnRe // This happens when the input is Scalar. We do not need to add anything in this case. if (input_count == output_count) { CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(reinterpret_cast(log_sum_result), exp_result, input_count * sizeof(T), cudaMemcpyDeviceToDevice)); - } else { + } else { // ReduceSum CUDNN_RETURN_IF_ERROR(cudnnReduceTensor( CudnnHandle(), reduce_desc, indices_cuda.get(), indices_bytes, workspace_cuda.get(), workspace_bytes, @@ -653,6 +681,170 @@ Status ReduceKernel::ComputeImpl( return Status::OK(); } +template <> +template <> +Status ReduceKernel::ComputeImpl(OpKernelContext* ctx, cudnnReduceTensorOp_t cudnn_reduce_op) const { + typedef typename ToCudaType::MappedType CudaT; + + const Tensor* X = nullptr; + Tensor* Y = nullptr; + + int64_t input_count = 0; + int64_t output_count = 0; + std::vector output_dims; + std::vector input_dims_cudnn; + std::vector output_dims_cudnn; + int64_t rank = 0; + int64_t stride = 0; + ORT_RETURN_IF_ERROR(PrepareForReduce(ctx, + keepdims_, + axes_, + &X, + &Y, + input_count, + output_count, + output_dims, + input_dims_cudnn, + output_dims_cudnn, + rank, stride)); + + // special case when there is a dim value of 0 in the shape. + if (input_count == 0) { + assert(Y->Shape().Size() == 0); + return Status::OK(); + } + + // cudnnReduceTensor has issue if input and output has same size, we just need to copy the data for this case + auto* const dst = Y->template MutableData(); + const auto* const src = X->template Data(); + if (input_count == output_count) { + if (src != dst) { + CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(dst, src, input_count * sizeof(int8_t), cudaMemcpyDeviceToDevice)); + } + return Status::OK(); + } + + size_t indices_bytes = 0; + size_t workspace_bytes = 0; + CudnnTensor input_tensor; + CudnnTensor output_tensor; + CudnnReduceDescriptor reduce_desc; + + cudnnDataType_t cudnn_type_X = CUDNN_DATA_FLOAT; + IAllocatorUniquePtr temp_X = GetScratchBuffer(input_count); + Impl_Cast(reinterpret_cast(src), temp_X.get(), X->Shape().Size()); + + ORT_RETURN_IF_ERROR(reduce_desc.Set(cudnn_reduce_op, cudnn_type_X, CUDNN_REDUCE_TENSOR_FLATTENED_INDICES)); + ORT_RETURN_IF_ERROR(input_tensor.Set(input_dims_cudnn, cudnn_type_X)); + ORT_RETURN_IF_ERROR(output_tensor.Set(output_dims_cudnn, cudnn_type_X)); + CUDNN_RETURN_IF_ERROR(cudnnGetReductionIndicesSize(CudnnHandle(), reduce_desc, input_tensor, output_tensor, &indices_bytes)); + CUDNN_RETURN_IF_ERROR(cudnnGetReductionWorkspaceSize(CudnnHandle(), reduce_desc, input_tensor, output_tensor, &workspace_bytes)); + IAllocatorUniquePtr indices_cuda = GetScratchBuffer(indices_bytes); + IAllocatorUniquePtr workspace_cuda = GetScratchBuffer(workspace_bytes); + + const auto one = Consts::One; + const auto zero = Consts::Zero; + auto temp_Y = GetScratchBuffer(output_count); + CUDNN_RETURN_IF_ERROR(cudnnReduceTensor(CudnnHandle(), + reduce_desc, + indices_cuda.get(), + indices_bytes, + workspace_cuda.get(), + workspace_bytes, + &one, + input_tensor, + temp_X.get(), + &zero, + output_tensor, + temp_Y.get())); + + Impl_Cast(temp_Y.get(), dst, output_count); + + return Status::OK(); +} + +template <> +template <> +Status ReduceKernel::ComputeImpl(OpKernelContext* ctx, cudnnReduceTensorOp_t cudnn_reduce_op) const { + typedef typename ToCudaType::MappedType CudaT; + + const Tensor* X = nullptr; + Tensor* Y = nullptr; + + int64_t input_count = 0; + int64_t output_count = 0; + std::vector output_dims; + std::vector input_dims_cudnn; + std::vector output_dims_cudnn; + int64_t rank = 0; + int64_t stride = 0; + ORT_RETURN_IF_ERROR(PrepareForReduce(ctx, + keepdims_, + axes_, + &X, + &Y, + input_count, + output_count, + output_dims, + input_dims_cudnn, + output_dims_cudnn, + rank, stride)); + + // special case when there is a dim value of 0 in the shape. + if (input_count == 0) { + assert(Y->Shape().Size() == 0); + return Status::OK(); + } + + // cudnnReduceTensor has issue if input and output has same size, we just need to copy the data for this case + auto* const dst = Y->template MutableData(); + const auto* const src = X->template Data(); + if (input_count == output_count) { + if (src != dst) { + CUDA_RETURN_IF_ERROR(cudaMemcpyAsync(dst, src, input_count * sizeof(uint8_t), cudaMemcpyDeviceToDevice)); + } + return Status::OK(); + } + + size_t indices_bytes = 0; + size_t workspace_bytes = 0; + CudnnTensor input_tensor; + CudnnTensor output_tensor; + CudnnReduceDescriptor reduce_desc; + + cudnnDataType_t cudnn_type_X = CUDNN_DATA_FLOAT; + IAllocatorUniquePtr temp_X = GetScratchBuffer(input_count); + Impl_Cast(reinterpret_cast(src), temp_X.get(), X->Shape().Size()); + + ORT_RETURN_IF_ERROR(reduce_desc.Set(cudnn_reduce_op, cudnn_type_X, CUDNN_REDUCE_TENSOR_FLATTENED_INDICES)); + ORT_RETURN_IF_ERROR(input_tensor.Set(input_dims_cudnn, cudnn_type_X)); + ORT_RETURN_IF_ERROR(output_tensor.Set(output_dims_cudnn, cudnn_type_X)); + CUDNN_RETURN_IF_ERROR(cudnnGetReductionIndicesSize(CudnnHandle(), reduce_desc, input_tensor, output_tensor, &indices_bytes)); + CUDNN_RETURN_IF_ERROR(cudnnGetReductionWorkspaceSize(CudnnHandle(), reduce_desc, input_tensor, output_tensor, &workspace_bytes)); + IAllocatorUniquePtr indices_cuda = GetScratchBuffer(indices_bytes); + IAllocatorUniquePtr workspace_cuda = GetScratchBuffer(workspace_bytes); + + const auto one = Consts::One; + const auto zero = Consts::Zero; + auto temp_Y = GetScratchBuffer(output_count); + CUDNN_RETURN_IF_ERROR(cudnnReduceTensor(CudnnHandle(), + reduce_desc, + indices_cuda.get(), + indices_bytes, + workspace_cuda.get(), + workspace_bytes, + &one, + input_tensor, + temp_X.get(), + &zero, + output_tensor, + temp_Y.get())); + + Impl_Cast(temp_Y.get(), dst, output_count); + + return Status::OK(); +} + #define REGISTER_KERNEL_HFD(name) \ REGISTER_KERNEL_TYPED(name, MLFloat16) \ REGISTER_KERNEL_TYPED(name, float) \ @@ -662,9 +854,23 @@ REGISTER_KERNEL_HFD(ArgMax) REGISTER_KERNEL_HFD(ArgMin) REGISTER_KERNEL_HFD(ReduceL1) REGISTER_KERNEL_HFD(ReduceL2) -REGISTER_KERNEL_HFD(ReduceMax) + +REGISTER_KERNEL_TYPED_12(ReduceMax, MLFloat16) +REGISTER_KERNEL_TYPED_12(ReduceMax, float) +REGISTER_KERNEL_TYPED_12(ReduceMax, double) +REGISTER_KERNEL_TYPED_12(ReduceMax, int32_t) +REGISTER_KERNEL_TYPED_12(ReduceMax, int8_t) +REGISTER_KERNEL_TYPED_12(ReduceMax, uint8_t) + REGISTER_KERNEL_HFD(ReduceMean) -REGISTER_KERNEL_HFD(ReduceMin) + +REGISTER_KERNEL_TYPED_12(ReduceMin, MLFloat16) +REGISTER_KERNEL_TYPED_12(ReduceMin, float) +REGISTER_KERNEL_TYPED_12(ReduceMin, double) +REGISTER_KERNEL_TYPED_12(ReduceMin, int32_t) +REGISTER_KERNEL_TYPED_12(ReduceMin, int8_t) +REGISTER_KERNEL_TYPED_12(ReduceMin, uint8_t) + REGISTER_KERNEL_HFD(ReduceProd) REGISTER_KERNEL_HFD(ReduceSum) REGISTER_KERNEL_HFD(ReduceLogSum) @@ -676,9 +882,8 @@ REGISTER_KERNEL_HFD(ReduceLogSumExp) REGISTER_KERNEL_INT32(ReduceL1) REGISTER_KERNEL_INT32(ReduceL2) -REGISTER_KERNEL_INT32(ReduceMax) REGISTER_KERNEL_INT32(ReduceMean) -REGISTER_KERNEL_INT32(ReduceMin) + REGISTER_KERNEL_INT32(ReduceProd) REGISTER_KERNEL_INT32(ReduceSum) diff --git a/onnxruntime/test/featurizers_ops/truncated_svdtransformer_test.cc b/onnxruntime/test/featurizers_ops/truncated_svdtransformer_test.cc index 6015cac261..bbd4305126 100644 --- a/onnxruntime/test/featurizers_ops/truncated_svdtransformer_test.cc +++ b/onnxruntime/test/featurizers_ops/truncated_svdtransformer_test.cc @@ -46,9 +46,27 @@ void TruncatedSVDTransformerTestRowMajStandard() { OpTester test("TruncatedSVDTransformer", 1, onnxruntime::kMSFeaturizersDomain); test.AddInput("State", {dim}, stream); test.AddInput("Input", {3, 3}, {-1, -1, 0, 0, -2, -1, -3, 0, -2}); - test.AddOutput("Output", {3, 3}, {1.009107f, 0.626315f, 0.767745f, 0.965105f, 1.995869f, -0.291682f, 3.529165f, -0.724887f, -0.139759f}); - test.SetOutputAbsErr("Output", 0.00016f); - test.Run(); + std::vector output = {-1.009107f, 0.626315f, -0.767745f, -0.965105f, 1.995869f, 0.291682f, -3.529165f, -0.724887f, 0.139759f}; + test.AddOutput("Output", {3, 3}, + output); + + Matrix verify_matrix(output.data(), 3, 3); + + // Custom verification function is necessary since the matrix output will vary from + // platform to platform enough so we choose to check max STD deviation. + OpTester::CustomOutputVerifierFn ver_fn = [&verify_matrix](const std::vector& fetches, const std::string& provider) { + std::cout << "Verifying TruncatedSVDTransformerTestRowMajStandard:" << provider << std::endl; + const float eps = (sizeof(size_t) == 8) ? 0.0001f : 0.0003f; + ASSERT_TRUE(fetches.size() == 1); + const auto& fetch = fetches.at(0); + const auto& tensor = fetch.Get(); + ASSERT_EQ(tensor.Shape().NumDimensions(), 2); + ASSERT_EQ(tensor.Shape().Size(), 9); + Matrix output_matrix(tensor.Data(), 3, 3); + ASSERT_LT((output_matrix.cwiseProduct(output_matrix) - verify_matrix.cwiseProduct(verify_matrix)).norm(), eps); + }; + + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kNupharExecutionProvider}, nullptr, {}, ORT_SEQUENTIAL, ver_fn); } TEST(FeaturizersTests, TruncatedSVDTransformer_double) { diff --git a/onnxruntime/test/onnx/main.cc b/onnxruntime/test/onnx/main.cc index 5fa7f0e8ff..2daa691303 100644 --- a/onnxruntime/test/onnx/main.cc +++ b/onnxruntime/test/onnx/main.cc @@ -479,7 +479,7 @@ int real_main(int argc, char* argv[], Ort::Env& env) { {"resize_upsample_sizes_nearest_round_prefer_ceil_asymmetric", "Bad onnx test output. Needs test fix."}, {"bitshift_right_uint16", "BitShift(11) uint16 support not enabled currently"}, {"bitshift_left_uint16", "BitShift(11) uint16 support not enabled currently"}, - {"maxunpool_export_with_output_shape", "Invalid output in ONNX test. See https://github.com/onnx/onnx/issues/2398"}, + {"maxunpool_export_with_output_shape", "Invalid output in ONNX test. See https://github.com/onnx/onnx/issues/2398"} }; if (enable_ngraph) { @@ -519,6 +519,7 @@ int real_main(int argc, char* argv[], Ort::Env& env) { broken_tests.insert({"maxpool_2d_dilations", "maxpool dilations not supported"}); broken_tests.insert({"mlperf_ssd_resnet34_1200", "test pass on dev box but fails on CI build"}); broken_tests.insert({"convtranspose_1d", "1d convtranspose not supported yet"}); + broken_tests.insert({"maxpool_2d_uint8", "Does not work on DNNL, NNAPI"}); } if (enable_openvino) { @@ -548,6 +549,7 @@ int real_main(int argc, char* argv[], Ort::Env& env) { broken_tests.insert({"range_float_type_positive_delta_expanded", "Temporarily disabled pending investigation"}); broken_tests.insert({"range_int32_type_negative_delta_expanded", "Temporarily disabled pending investigation"}); broken_tests.insert({"convtranspose_1d", "1d convtranspose not supported yet"}); + broken_tests.insert({"maxpool_2d_uint8", "Does not work on DNNL, NNAPI"}); } if (enable_tensorrt) { @@ -574,6 +576,7 @@ int real_main(int argc, char* argv[], Ort::Env& env) { broken_tests.insert({"mlperf_ssd_resnet34_1200", "unknown error"}); broken_tests.insert({"tf_inception_v1", "flaky test"}); //TODO: Investigate cause for flakiness broken_tests.insert({"convtranspose_1d", "1d convtranspose not supported yet"}); + broken_tests.insert({"faster_rcnn", "Linux: faster_rcnn:output=6383:shape mismatch, expect {77} got {57}"}); } if (enable_dml) { diff --git a/onnxruntime/test/providers/cpu/controlflow/scan_test.cc b/onnxruntime/test/providers/cpu/controlflow/scan_test.cc index 315012f0f1..7422b294b4 100644 --- a/onnxruntime/test/providers/cpu/controlflow/scan_test.cc +++ b/onnxruntime/test/providers/cpu/controlflow/scan_test.cc @@ -337,7 +337,7 @@ static void RunTest_v9(const std::string test_name, int64_t sequence_len, int64_ CreateSubgraph(graph, options, options.add_bad_shape ? failure_message : ""); auto& proto = graph.ToGraphProto(); - ScanOpTester test{11}; // use latest version - no significant change over 9 + ScanOpTester test{ (options.add_bad_shape) ? -1 : 11}; // use latest version - no significant change over 9 test.AddAttribute("body", proto); test.AddAttribute("num_scan_inputs", 2); diff --git a/onnxruntime/test/providers/cpu/nn/pool_op_test.cc b/onnxruntime/test/providers/cpu/nn/pool_op_test.cc index 99f47473d3..c050eeae89 100644 --- a/onnxruntime/test/providers/cpu/nn/pool_op_test.cc +++ b/onnxruntime/test/providers/cpu/nn/pool_op_test.cc @@ -215,6 +215,90 @@ TEST(PoolTest, MaxPool1D_8_With_Index) { MaxPool1D_8_WithIndexTest(1 /*storage_order*/); } +static void MaxPool1D_12_WithIndexTest_int8(int64_t storage_order) { + OpTester test("MaxPool", 12); + + test.AddAttribute("auto_pad", ""); + test.AddAttribute("strides", std::vector{2}); + test.AddAttribute("pads", vector{0, 0}); + test.AddAttribute("kernel_shape", vector{2}); + test.AddAttribute("storage_order", storage_order); + + std::vector x_vals = {1, 2, 3, 4, 5, 6, 7, 8}; + std::vector x_dims = {1, 2, 4}; + std::vector expected_dims = {1, 2, 2}; + std::vector expected_vals = {2, 4, 6, 8}; + std::vector expected_indices = {1, 3, 5, 7}; + + test.AddInput("X", x_dims, x_vals); + test.AddOutput("Y", expected_dims, expected_vals); + test.AddOutput("Indices", expected_dims, expected_indices); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kAclExecutionProvider}); +} + +static void MaxPool1D_12_WithIndexTest_uint8(int64_t storage_order) { + OpTester test("MaxPool", 12); + + test.AddAttribute("auto_pad", ""); + test.AddAttribute("strides", std::vector{2}); + test.AddAttribute("pads", vector{0, 0}); + test.AddAttribute("kernel_shape", vector{2}); + test.AddAttribute("storage_order", storage_order); + + std::vector x_vals = {1, 2, 3, 4, 5, 6, 7, 8}; + std::vector x_dims = {1, 2, 4}; + std::vector expected_dims = {1, 2, 2}; + std::vector expected_vals = {2, 4, 6, 8}; + std::vector expected_indices = {1, 3, 5, 7}; + + test.AddInput("X", x_dims, x_vals); + test.AddOutput("Y", expected_dims, expected_vals); + test.AddOutput("Indices", expected_dims, expected_indices); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kAclExecutionProvider}); +} + +TEST(PoolTest, MaxPool1D_12_With_Index_8bits) { + MaxPool1D_12_WithIndexTest_int8(0 /*storage_order*/); + MaxPool1D_12_WithIndexTest_int8(1 /*storage_order*/); + MaxPool1D_12_WithIndexTest_uint8(0 /*storage_order*/); + MaxPool1D_12_WithIndexTest_uint8(1 /*storage_order*/); +} + +// Used by MaxPool2D_uint8 +template +void print_vector(std::ostream& os, const std::string& txt, InputIter begin, InputIter end) { + os << txt; + while (begin != end) { + std::cout << uint16_t(*begin) << ", "; + ++begin; + } + os << std::endl; +} + +TEST(PoolTest, MaxPool2D_uint8) { + OpTester test("MaxPool", 12); + test.AddAttribute("kernel_shape", std::vector{5, 5}); + test.AddAttribute("pads", std::vector{2, 2, 2, 2}); + + std::vector output_shape = {1, 1, 5, 5}; + std::vector output = { + 13, 14, 15, 15, 15, + 18, 19, 20, 20, 20, + 23, 24, 25, 25, 25, + 23, 24, 25, 25, 25, + 23, 24, 25, 25, 25}; + + test.AddInput("Input", {1, 1, 5, 5}, { + 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}); + + test.AddOutput("Output", output_shape, output); + test.Run(); +} + TEST(PoolTest, MaxPool_10_Dilation_1d) { OpTester test("MaxPool", 10); @@ -238,7 +322,7 @@ TEST(PoolTest, MaxPool_10_Dilation_1d) { TEST(PoolTest, MaxPool_DefaultDilations) { OpTester test("MaxPool"); - test.AddAttribute("kernel_shape", vector{2}); + test.AddAttribute("kernel_shape", std::vector{2}); std::vector x_dims = {1, 3, 3}; std::vector x_vals = {0.f, 1.f, 2.f, @@ -255,6 +339,46 @@ TEST(PoolTest, MaxPool_DefaultDilations) { test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); } +TEST(PoolTest, MaxPool_DefaultDilations_int8) { + OpTester test("MaxPool", 12); + + test.AddAttribute("kernel_shape", std::vector{2}); + + std::vector x_dims = {1, 3, 3}; + std::vector x_vals = {0, 1, 2, + 3, 4, 5, + 6, 7, 8}; + + std::vector expected_dims = {1, 3, 2}; + std::vector expected_vals = {1, 2, + 4, 5, + 7, 8}; + + test.AddInput("X", x_dims, x_vals); + test.AddOutput("Y", expected_dims, expected_vals); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); +} + +TEST(PoolTest, MaxPool_DefaultDilations_uint8) { + OpTester test("MaxPool", 12); + + test.AddAttribute("kernel_shape", std::vector{2}); + + std::vector x_dims = {1, 3, 3}; + std::vector x_vals = {0, 1, 2, + 3, 4, 5, + 6, 7, 8}; + + std::vector expected_dims = {1, 3, 2}; + std::vector expected_vals = {1, 2, + 4, 5, + 7, 8}; + + test.AddInput("X", x_dims, x_vals); + test.AddOutput("Y", expected_dims, expected_vals); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); +} + TEST(PoolTest, MaxPool_10_DilationPadding_1d) { OpTester test("MaxPool", 10); @@ -298,6 +422,29 @@ TEST(PoolTest, MaxPool_10_Dilation_2d) { test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); } +TEST(PoolTest, MaxPool_10_Dilation_2d_int8) { + OpTester test("MaxPool", 12); + + test.AddAttribute("auto_pad", ""); + test.AddAttribute("strides", std::vector{1, 1}); + test.AddAttribute("pads", vector{0, 0, 0, 0}); + test.AddAttribute("kernel_shape", vector{2, 2}); + test.AddAttribute("dilations", vector{2, 2}); + + std::vector x_vals = { + 1, 3, 2, 4, -1, + 5, 7, 6, 8, -2, + 9, 11, 10, 12, -3, + 13, 15, 14, 16, -4}; + std::vector x_dims = {1, 1, 4, 5}; + std::vector expected_dims = {1, 1, 2, 3}; + std::vector expected_vals = {10, 12, 10, 14, 16, 14}; + + test.AddInput("X", x_dims, x_vals); + test.AddOutput("Y", expected_dims, expected_vals); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); +} + TEST(PoolTest, MaxPool_10_DilationPadding_2d) { OpTester test("MaxPool", 10); @@ -323,7 +470,7 @@ TEST(PoolTest, MaxPool_10_DilationPadding_2d) { test.AddInput("X", x_dims, x_vals); test.AddOutput("Y", expected_dims, expected_vals); test.Run(OpTester::ExpectResult::kExpectSuccess, "", - {kCudaExecutionProvider, kTensorrtExecutionProvider}); + {kCudaExecutionProvider, kTensorrtExecutionProvider}); } TEST(PoolTest, MaxPool_10_Dilation_Ceil0_2d) { @@ -349,6 +496,29 @@ TEST(PoolTest, MaxPool_10_Dilation_Ceil0_2d) { test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kAclExecutionProvider}); } +TEST(PoolTest, MaxPool_12_Dilation_Ceil0_2d_int8) { + OpTester test("MaxPool", 12); + + test.AddAttribute("auto_pad", ""); + test.AddAttribute("strides", std::vector{2, 1}); + test.AddAttribute("pads", vector{0, 0, 0, 0}); + test.AddAttribute("kernel_shape", vector{2, 2}); + test.AddAttribute("dilations", vector{2, 2}); + + std::vector x_vals = { + 1, 3, 2, 4, -1, + 5, 7, 6, 8, -2, + 9, 11, 10, 12, -3, + 13, 15, 14, 16, -4}; + std::vector x_dims = {1, 1, 4, 5}; + std::vector expected_dims = {1, 1, 1, 3}; + std::vector expected_vals = {10, 12, 10}; + + test.AddInput("X", x_dims, x_vals); + test.AddOutput("Y", expected_dims, expected_vals); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kAclExecutionProvider}); +} + TEST(PoolTest, MaxPool_10_Dilation_Ceil1_2d) { OpTester test("MaxPool", 10); diff --git a/onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc b/onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc index 69a1fb213f..607332d17a 100644 --- a/onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc +++ b/onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc @@ -610,6 +610,40 @@ TEST(ReductionOpTest, ReduceMax_int64) { test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: axis must be 0 } +TEST(ReductionOpTest, ReduceMax_int8) { + OpTester test("ReduceMax", 12); + test.AddAttribute("axes", std::vector{1, 2}); + test.AddAttribute("keepdims", (int64_t)1); + test.AddInput("data", {3, 2, 2}, + {1, 2, + 3, 4, + + 5, 6, + 7, 8, + + 9, 10, + 11, 12}); + test.AddOutput("reduced", {3, 1, 1}, {4, 8, 12}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: axis must be 0 +} + +TEST(ReductionOpTest, ReduceMax_uint8) { + OpTester test("ReduceMax", 12); + test.AddAttribute("axes", std::vector{1, 2}); + test.AddAttribute("keepdims", (int64_t)1); + test.AddInput("data", {3, 2, 2}, + {1, 2, + 3, 4, + + 5, 6, + 7, 8, + + 9, 10, + 11, 12}); + test.AddOutput("reduced", {3, 1, 1}, {4, 8, 12}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: axis must be 0 +} + #if !(defined USE_TENSORRT) && !(defined USE_TVM) TEST(ReductionOpTest, ReduceMax0DTensor) { OpTester test("ReduceMax"); @@ -858,6 +892,41 @@ TEST(ReductionOpTest, ReduceMin_int32) { test.Run(); } +TEST(ReductionOpTest, ReduceMin_int8) { + OpTester test("ReduceMin", 12); + test.AddAttribute("axes", std::vector{0, 2}); + test.AddAttribute("keepdims", (int64_t)1); + test.AddInput("data", {3, 2, 2}, + {1, 2, + 3, 4, + + 5, 6, + 7, 8, + + 9, 10, + 11, 12}); + test.AddOutput("reduced", {1, 2, 1}, {1, 3}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); +} + +TEST(ReductionOpTest, ReduceMin_uint8) { + OpTester test("ReduceMin", 12); + test.AddAttribute("axes", std::vector{0, 2}); + test.AddAttribute("keepdims", (int64_t)1); + test.AddInput("data", {3, 2, 2}, + {1, 2, + 3, 4, + + 5, 6, + 7, 8, + + 9, 10, + 11, 12}); + test.AddOutput("reduced", {1, 2, 1}, {1, 3}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); +} + + #if !(defined USE_TENSORRT) && !(defined USE_TVM) TEST(ReductionOpTest, ReduceMin0DTensor) { OpTester test("ReduceMin"); @@ -1478,6 +1547,50 @@ TEST(ReductionOpTest, ArgMax_int32) { test.Run(); } +TEST(ReductionOpTest, ArgMax_int32_last_index_nodups) { + OpTester test("ArgMax", 12); + test.AddAttribute("axis", (int64_t)1); + test.AddAttribute("keepdims", (int64_t)1); + test.AddAttribute("select_last_index", (int64_t)1); + + test.AddInput("data", {3, 2, 2}, + {2, 2, + 3, 4, + + 5, 6, + 7, 8, + + 10, 10, + 11, 12}); + test.AddOutput("reduced", {3, 1, 2}, + {1, 1, + 1, 1, + 1, 1}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kNGraphExecutionProvider}); +} + +TEST(ReductionOpTest, ArgMax_int32_last_index_dups) { + OpTester test("ArgMax", 12); + test.AddAttribute("axis", (int64_t)1); + test.AddAttribute("keepdims", (int64_t)1); + test.AddAttribute("select_last_index", (int64_t)1); + + test.AddInput("data", {3, 2, 2}, + {2, 4, + 3, 4, + + 8, 6, + 7, 8, + + 9, 13, + 11, 12}); + test.AddOutput("reduced", {3, 1, 2}, + {1, 1, + 0, 1, + 1, 0}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kNGraphExecutionProvider}); +} + TEST(ReductionOpTest, ArgMax_int32_neg_axis) { OpTester test("ArgMax"); test.AddAttribute("axis", (int64_t)(-2)); @@ -1512,6 +1625,20 @@ TEST(ReductionOpTest, ArgMax2D) { test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: axis must be 0 } +TEST(ReductionOpTest, ArgMax2D_select_last) { + OpTester test("ArgMax", 12); + test.AddAttribute("axis", (int64_t)1); + test.AddAttribute("keepdims", (int64_t)1); + test.AddAttribute("select_last_index", (int64_t)1); + test.AddInput("data", {3, 2}, + {1.0f, 2.0f, + 6.0f, 5.0f, + 9.0f, 10.0f}); + test.AddOutput("reduced", {3, 1}, + {1, 0, 1}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kNGraphExecutionProvider}); +} + TEST(ReductionOpTest, ArgMin) { OpTester test("ArgMin"); test.AddAttribute("axis", (int64_t)0); @@ -1560,6 +1687,18 @@ TEST(ReductionOpTest, ArgMin_do_not_keepdims_2) { test.Run(); } +TEST(ReductionOpTest, ArgMin_do_not_keepdims_2_select_last) { + OpTester test("ArgMin", 12); + test.AddAttribute("axis", (int64_t)0); + test.AddAttribute("keepdims", (int64_t)0); + test.AddAttribute("select_last_index", (int64_t)1); + test.AddInput("data", {3}, + {1.0f, 2.0f, 3.0f}); + test.AddOutput("reduced", {}, {0}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kNGraphExecutionProvider}); +} + + TEST(ReductionOpTest, ArgMin_int32) { OpTester test("ArgMin"); test.AddAttribute("axis", (int64_t)0); @@ -1579,6 +1718,27 @@ TEST(ReductionOpTest, ArgMin_int32) { test.Run(); } +TEST(ReductionOpTest, ArgMin_int32_select_last) { + OpTester test("ArgMin", 12); + test.AddAttribute("axis", (int64_t)0); + test.AddAttribute("keepdims", (int64_t)0); + test.AddAttribute("select_last_index", (int64_t)1); + test.AddInput("data", {3, 2, 2}, + {1, 2, + 3, 4, + + 1, 6, + 7, 8, + + 9, 10, + 11, 12}); + test.AddOutput("reduced", {2, 2}, + {1, 0, + 0, 0}); + + test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kNGraphExecutionProvider}); +} + #ifdef USE_CUDA void test_reduce_apis(size_t size) { diff --git a/onnxruntime/test/providers/cpu/tensor/gather_op_test.cc b/onnxruntime/test/providers/cpu/tensor/gather_op_test.cc index c8abb3ec15..198d4a3f12 100644 --- a/onnxruntime/test/providers/cpu/tensor/gather_op_test.cc +++ b/onnxruntime/test/providers/cpu/tensor/gather_op_test.cc @@ -76,7 +76,9 @@ TEST(GatherOpTest, Gather_invalid_index_cpu) { test.AddInput("indices", {3}, {0LL, 1L, 1000L}); test.AddOutput("output", {1}, {1.0f}); - test.Run(OpTester::ExpectResult::kExpectFailure, "Mismatch between number of source and target dimensions."); + // On Cuda it is impossible to dereference indecies memory on CPU so the check can not run + test.Run(OpTester::ExpectResult::kExpectFailure, "indices element out of data bounds, idx=1000 must be within the inclusive range [-3,2]", + {kCudaExecutionProvider, kNGraphExecutionProvider, kDnnlExecutionProvider, kNupharExecutionProvider, kTensorrtExecutionProvider}); } #ifdef USE_CUDA diff --git a/onnxruntime/test/python/onnx_backend_test_series.py b/onnxruntime/test/python/onnx_backend_test_series.py index 8eb488c294..00940b5050 100644 --- a/onnxruntime/test/python/onnx_backend_test_series.py +++ b/onnxruntime/test/python/onnx_backend_test_series.py @@ -112,7 +112,13 @@ def create_backend_test(testname=None): current_failing_tests += [ '^test_range_float_type_positive_delta_expanded_cpu', '^test_range_int32_type_negative_delta_expanded_cpu', '^test_averagepool_2d_ceil_cpu', - '^test_maxpool_2d_ceil_cpu', '^test_maxpool_2d_dilations_cpu' + '^test_maxpool_2d_ceil_cpu', '^test_maxpool_2d_dilations_cpu', + '^test_maxpool_2d_uint8' + ] + + if c2.supports_device('NNAPI'): + current_failing_tests += [ + '^test_maxpool_2d_uint8' ] if c2.supports_device('OPENVINO_GPU_FP32') or c2.supports_device('OPENVINO_GPU_FP16'):