saymrwulf/onnxruntime
1
0
Fork 0
mirror of https://github.com/saymrwulf/onnxruntime.git synced 2026-07-24 19:43:35 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

OpSet 12 ops (#3341)

Browse source 
Advance ONNX commit to pickup the latest ArgMax, ArgMin,
  ReduceMax/ReduceMin, MaxPool
  Declare new versions for CPU/CUDA.
  Implement infrastructure support for int8/uint8.
  Adust GatherOp test for a new error.
  Adjust Scan9.BadShape test.
  Add exclusions for index out of bounds checks.
  Rework result verification for SVDTransformer.
This commit is contained in:
Dmitri Smirnov 2020-03-31 15:31:06 -07:00 committed by GitHub
parent 044c466158
commit a4fe60c4d3
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
20 changed files with 991 additions and 148 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

2
cgmanifest.json
View file

@ -49,7 +49,7 @@
"component": {
"type": "git",
"git": {
"commitHash": "1facb4c1bb9cc2107d4dbaf9fd647fefdbbeb0ab",
"commitHash": "423f1977d314f05df4be7edb44428c1c0211341c",
"repositoryUrl": "https://github.com/onnx/onnx.git"
}
}

2
cmake/external/onnx vendored

@ -1 +1 @@
Subproject commit 1facb4c1bb9cc2107d4dbaf9fd647fefdbbeb0ab
Subproject commit 423f1977d314f05df4be7edb44428c1c0211341c

84
onnxruntime/core/providers/cpu/cpu_execution_provider.cc
View file

@ -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<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 6, 10,
@ -593,9 +615,7 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) {
AveragePool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 1, 7,
MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, 11, float,
MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, 11, double,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, 11,
MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 2, 10,
LpPool)>,
@ -1001,14 +1021,14 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) {
Round)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, uint8_t,
DynamicQuantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, float,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 11, float,
ArgMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, int32_t,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 11, int32_t,
ArgMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, float,
ArgMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, int32_t,
ArgMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 11,
float, ArgMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 11,
int32_t, ArgMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, Hardmax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, float, LogSoftmax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, double, LogSoftmax)>,
@ -1111,6 +1131,40 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) {
int32_t, Resize)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11,
uint8_t, Resize)>,
// OpSet 12
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, float,
ArgMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int32_t,
ArgMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, float,
ArgMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int32_t,
ArgMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, float,
ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int32_t,
ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int64_t,
ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int8_t,
ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, uint8_t,
ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, float,
ReduceMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int32_t,
ReduceMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int64_t,
ReduceMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int8_t,
ReduceMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, uint8_t,
ReduceMin)>,
};
for (auto& function_table_entry : function_table) {

62
onnxruntime/core/providers/cpu/nn/pool.cc
View file

@ -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 <typename T>
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<float, AveragePool>::Compute(OpKernelContext* context) const {
// version 8: Added storage_order And Indices
// version 10: Added ceil_mode
// version 11: Added dilations
template <typename T>
// version 12: Added int8/uint8 support
class MaxPoolV8 : public OpKernel, public PoolBase {
template <typename T>
struct ComputeHelper {
Status operator()(const MaxPoolV8* inst, OpKernelContext* context) const {
return inst->ComputeImpl<T>(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<Status, ComputeHelper, float, double, int8_t, uint8_t>
t_disp(context->Input<Tensor>(0)->GetElementType());
return t_disp.Invoke(this, context);
}
private:
template <typename T>
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<T, float>::value) {
if (OpKernel::Node().OutputDefs().size() == 1 && pool_attrs_.storage_order == 0 && !need_dilation) {
return PoolBase::Compute(context, MlasMaximumPooling);
}
}
const auto* X = context->Input<Tensor>(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<MaxPool3DTask<T>>(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<float>()),
Pool<float, MaxPool<1 /*VERSION*/>>);
ONNX_CPU_OPERATOR_VERSIONED_TYPED_KERNEL(MaxPool, 8, 11, float,
ONNX_CPU_OPERATOR_VERSIONED_KERNEL(MaxPool, 8, 11,
KernelDefBuilder()
.TypeConstraint("T", DataTypeImpl::GetTensorType<float>())
.TypeConstraint("T", {DataTypeImpl::GetTensorType<float>(),
DataTypeImpl::GetTensorType<double>()})
.TypeConstraint("I", DataTypeImpl::GetTensorType<int64_t>()),
MaxPoolV8<float>);
MaxPoolV8);
ONNX_CPU_OPERATOR_VERSIONED_TYPED_KERNEL(MaxPool, 8, 11, double,
KernelDefBuilder()
.TypeConstraint("T", DataTypeImpl::GetTensorType<double>())
.TypeConstraint("I", DataTypeImpl::GetTensorType<int64_t>()),
MaxPoolV8<double>);
ONNX_CPU_OPERATOR_KERNEL(MaxPool, 12,
KernelDefBuilder()
.TypeConstraint("T", {DataTypeImpl::GetTensorType<double>(),
DataTypeImpl::GetTensorType<float>(),
DataTypeImpl::GetTensorType<int8_t>(),
DataTypeImpl::GetTensorType<uint8_t>()})
.TypeConstraint("I", DataTypeImpl::GetTensorType<int64_t>()),
MaxPoolV8);
ONNX_CPU_OPERATOR_VERSIONED_KERNEL(LpPool, 2, 10,
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<float>()),

1
onnxruntime/core/providers/cpu/nn/pool.h
View file

@ -24,5 +24,4 @@ class Pool : public OpKernel, public PoolBase {
private:
PoolProcessContext pool_context_;
};
} // namespace onnxruntime

135
onnxruntime/core/providers/cpu/reduction/reduction_ops.cc
View file

@ -74,6 +74,22 @@ namespace onnxruntime {
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<int64_t>()), \
x<int64_t>);
#define REGISTER_UNARY_ELEMENTWISE_KERNEL_INT8_ONLY(x, sinceVersion) \
ONNX_CPU_OPERATOR_TYPED_KERNEL( \
x, \
sinceVersion, \
int8_t, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<int8_t>()), \
x<int8_t>);
#define REGISTER_UNARY_ELEMENTWISE_KERNEL_UINT8_ONLY(x, sinceVersion) \
ONNX_CPU_OPERATOR_TYPED_KERNEL( \
x, \
sinceVersion, \
uint8_t, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<uint8_t>()), \
x<uint8_t>);
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<int64_t> new_dims_(transposed_axes.size());
std::vector<int64_t> 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<int>(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<T>::Compute(OpKernelContext* ctx) const {
if (no_transpose) {
const T* input_data = ctx->Input<Tensor>(0)->template Data<T>();
for (int64_t i = 0; i < block_size; ++i) {
ConstEigenVectorMap<T>(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<const T> 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<T>(input_data + (i * blocks), blocks).maxCoeff(&maxIndex);
*(output_data++) = maxIndex;
}
}
} else {
auto matrixData = ConstEigenMatrixMap<T>(&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<T>::Compute(OpKernelContext* ctx) const {
if (no_transpose) {
const T* input_data = ctx->Input<Tensor>(0)->template Data<T>();
for (int64_t i = 0; i < block_size; ++i) {
ConstEigenVectorMap<T>(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<const T> 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<T>(input_data + (i * blocks), blocks).minCoeff(&minIndex);
*(output_data++) = minIndex;
}
}
} else {
auto matrixData = ConstEigenMatrixMap<T>(&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;
}
}
}

3
onnxruntime/core/providers/cpu/reduction/reduction_ops.h
View file

@ -27,10 +27,13 @@ class ReduceKernelBase {
ORT_ENFORCE(info.GetAttr("keepdims", &keepdims).IsOK());
}
keepdims_ = (keepdims == 1);
int64_t select_last_index = info.GetAttrOrDefault<int64_t>("select_last_index", 0);
select_last_index_ = (select_last_index != 0);
}
std::vector<int64_t> axes_;
bool keepdims_;
bool select_last_index_;
};
template <bool allow_multi_axes>

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

@ -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<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, MemcpyFromHost)>,
@ -1075,6 +1097,10 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, MLFloat16, ThresholdedRelu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, 10, TopK)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 1, 10, If)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, int8_t, QuantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, uint8_t, QuantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, int8_t, DequantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, uint8_t, DequantizeLinear)>,
// opset 11
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, ArgMax)>,
@ -1109,18 +1135,18 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, ReduceLogSumExp)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, ReduceLogSumExp)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, ReduceLogSumExp)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, int32_t, ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, float, ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, double, ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, MLFloat16, ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, int32_t, ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, ReduceMean)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, ReduceMean)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, ReduceMean)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, int32_t, ReduceMean)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, ReduceMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, ReduceMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, ReduceMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, int32_t, ReduceMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, float, ReduceMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, double, ReduceMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, MLFloat16, ReduceMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, int32_t, ReduceMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, ReduceProd)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, ReduceProd)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, ReduceProd)>,
@ -1153,9 +1179,9 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, AveragePool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, AveragePool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, AveragePool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, float, MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, double, MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, MLFloat16, MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, Resize)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, Resize)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, Resize)>,
@ -1171,11 +1197,29 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, Round)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, Round)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, Round)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, int8_t, QuantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, uint8_t, QuantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, int8_t, DequantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 10, uint8_t, DequantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, CumSum)>,
// OpSet 12
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int8_t, MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint8_t, MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int32_t, ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int8_t, ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint8_t, ReduceMax)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, ReduceMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, ReduceMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, ReduceMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int32_t, ReduceMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int8_t, ReduceMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint8_t, ReduceMin)>,
};
for (auto& function_table_entry : function_table) {

59
onnxruntime/core/providers/cuda/cudnn_common.cc
View file

@ -47,18 +47,6 @@ Status CudnnTensor::Set(const CudnnTensor& x_desc, cudnnBatchNormMode_t mode) {
return Status::OK();
}
template <typename ElemType>
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<int64_t>& filter_dims, cudnn
return Status::OK();
}
template cudnnDataType_t CudnnTensor::GetDataType<float>();
template cudnnDataType_t CudnnTensor::GetDataType<double>();
template cudnnDataType_t CudnnTensor::GetDataType<half>();
template <typename ElemType>
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<float>() {
return CUDNN_DATA_FLOAT;
}
template <>
cudnnDataType_t CudnnTensor::GetDataType<double>() {
return CUDNN_DATA_DOUBLE;
}
template <>
cudnnDataType_t CudnnTensor::GetDataType<half>() {
return CUDNN_DATA_HALF;
}
template <>
cudnnDataType_t CudnnTensor::GetDataType<int8_t>() {
return CUDNN_DATA_INT8;
}
template <>
cudnnDataType_t CudnnTensor::GetDataType<uint8_t>() {
return CUDNN_DATA_UINT8;
}
template <>
const float Consts<float>::One = 1;
@ -144,6 +161,18 @@ const float Consts<half>::Zero = 0;
const float Consts<half>::One = 1;
template <>
const int8_t Consts<int8_t>::Zero = 0;
template <>
const int8_t Consts<int8_t>::One = 1;
template <>
const uint8_t Consts<uint8_t>::Zero = 0;
template <>
const uint8_t Consts<uint8_t>::One = 1;
std::vector<cudaDeviceProp> DeviceProp::s_cachedDeviceProps;
std::once_flag DeviceProp::s_cachedDevicePropsInitFlag;

4
onnxruntime/core/providers/cuda/cudnn_common.h
View file

@ -13,6 +13,7 @@ class CudnnTensor final {
public:
CudnnTensor();
~CudnnTensor();
ORT_DISALLOW_COPY_ASSIGNMENT_AND_MOVE(CudnnTensor);
Status Set(const std::vector<int64_t>& 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<int64_t>& 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));

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

@ -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

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

@ -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<int64_t>& kernel_shape,
const std::vector<int64_t>& pads,
@ -132,9 +144,10 @@ Status Pool<T, PoolType>::ComputeInternal(OpKernelContext* context) const {
}
std::vector<int64_t> 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<const CudaT*>(X->template Data<T>());
@ -152,13 +165,6 @@ Status Pool<T, PoolType>::ComputeInternal(OpKernelContext* context) const {
strides.push_back(1);
}
const auto alpha = Consts<CudaT>::One;
const auto beta = Consts<CudaT>::Zero;
CudnnTensor x_tensor;
CudnnTensor y_tensor;
ORT_RETURN_IF_ERROR(x_tensor.Set(x_dims_cudnn, CudnnTensor::GetDataType<CudaT>()));
ORT_RETURN_IF_ERROR(y_tensor.Set(y_dims_cudnn, CudnnTensor::GetDataType<CudaT>()));
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<T, PoolType>::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<T, int8_t>::value || std::is_same<T, uint8_t>::value) {
// Cast to float back and forth using temp buffer
const auto alpha = Consts<float>::One;
const auto beta = Consts<float>::Zero;
CudnnTensor x_tensor;
CudnnTensor y_tensor;
ORT_RETURN_IF_ERROR(x_tensor.Set(x_dims_cudnn, CudnnTensor::GetDataType<float>()));
ORT_RETURN_IF_ERROR(y_tensor.Set(y_dims_cudnn, CudnnTensor::GetDataType<float>()));
const auto input_count = x_shape.Size();
const auto output_count = y_shape.Size();
IAllocatorUniquePtr<float> temp_X = GetScratchBuffer<float>(input_count);
auto temp_Y = GetScratchBuffer<float>(output_count);
Impl_Cast<CudaT, float>(reinterpret_cast<const CudaT*>(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<float, CudaT>(temp_Y.get(), y_data, output_count);
} else {
const auto alpha = Consts<CudaT>::One;
const auto beta = Consts<CudaT>::Zero;
CudnnTensor x_tensor;
CudnnTensor y_tensor;
ORT_RETURN_IF_ERROR(x_tensor.Set(x_dims_cudnn, CudnnTensor::GetDataType<CudaT>()));
ORT_RETURN_IF_ERROR(y_tensor.Set(y_dims_cudnn, CudnnTensor::GetDataType<CudaT>()));
CUDNN_RETURN_IF_ERROR(cudnnPoolingForward(CudnnHandle(), pooling_desc, &alpha, x_tensor, x_data, &beta, y_tensor, y_data));
}
return Status::OK();
}

253
onnxruntime/core/providers/cuda/reduction/reduction_ops.cc
View file

@ -32,6 +32,34 @@ namespace cuda {
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
name<T>);
// 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<T>()), \
name<T>); \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
name, \
kOnnxDomain, \
11, 11, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
name<T>); \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
name, \
kOnnxDomain, \
12, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder().TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
name<T>);
// 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<allow_multi_axes>::ReduceKernelShared(
cudnnDataType_t cudnn_type_X = CudnnTensor::GetDataType<CudaT>();
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<int>::max() && stride <= std::numeric_limits<int>::max() &&
is_matrix_row_reduction(cudnn_reduce_op,
static_cast<int>(reduction_size),
static_cast<int>(stride), rank, axes_)) {
static_cast<int>(reduction_size),
static_cast<int>(stride), rank, axes_)) {
reduce_matrix_rows(
reinterpret_cast<const CudaT*>(X),
reinterpret_cast<CudaT*>(Y),
static_cast<int>(reduction_size),
static_cast<int>(stride));
reinterpret_cast<const CudaT*>(X),
reinterpret_cast<CudaT*>(Y),
static_cast<int>(reduction_size),
static_cast<int>(stride));
return Status::OK();
}
@ -402,17 +429,18 @@ Status ReduceKernel<allow_multi_axes>::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<int>::max() && stride <= std::numeric_limits<int>::max() &&
is_matrix_row_reduction(cudnn_reduce_op,
static_cast<int>(reduction_size),
static_cast<int>(stride), rank, axes_)) {
reduce_matrix_rows(
reinterpret_cast<const CudaT*>(X->template Data<T>()),
reinterpret_cast<CudaT*>(Y->template MutableData<T>()),
static_cast<int>(reduction_size),
static_cast<int>(stride));
return Status::OK();
if (!std::is_same<T, int8_t>::value && !std::is_same<T, uint8_t>::value) {
if (fast_reduction_ && reduction_size <= std::numeric_limits<int>::max() && stride <= std::numeric_limits<int>::max() &&
is_matrix_row_reduction(cudnn_reduce_op,
static_cast<int>(reduction_size),
static_cast<int>(stride), rank, axes_)) {
reduce_matrix_rows(
reinterpret_cast<const CudaT*>(X->template Data<T>()),
reinterpret_cast<CudaT*>(Y->template MutableData<T>()),
static_cast<int>(reduction_size),
static_cast<int>(stride));
return Status::OK();
}
}
if (ReduceTensorIndices == CUDNN_REDUCE_TENSOR_FLATTENED_INDICES && std::is_same<T, MLFloat16>::value) {
@ -498,7 +526,7 @@ Status ReduceKernel<allow_multi_axes>::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<CudaT*>(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<true>::ComputeImpl<int32_t, CUDNN_REDUCE_TENSOR_NO_INDICES>(
return Status::OK();
}
template <>
template <>
Status ReduceKernel<true>::ComputeImpl<int8_t, CUDNN_REDUCE_TENSOR_NO_INDICES>(OpKernelContext* ctx, cudnnReduceTensorOp_t cudnn_reduce_op) const {
typedef typename ToCudaType<int8_t>::MappedType CudaT;
const Tensor* X = nullptr;
Tensor* Y = nullptr;
int64_t input_count = 0;
int64_t output_count = 0;
std::vector<int64_t> output_dims;
std::vector<int64_t> input_dims_cudnn;
std::vector<int64_t> 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<int8_t>();
const auto* const src = X->template Data<int8_t>();
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<float> temp_X = GetScratchBuffer<float>(input_count);
Impl_Cast<CudaT, float>(reinterpret_cast<const CudaT*>(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<uint32_t> indices_cuda = GetScratchBuffer<uint32_t>(indices_bytes);
IAllocatorUniquePtr<CudaT> workspace_cuda = GetScratchBuffer<CudaT>(workspace_bytes);
const auto one = Consts<float>::One;
const auto zero = Consts<float>::Zero;
auto temp_Y = GetScratchBuffer<float>(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<float, int8_t>(temp_Y.get(), dst, output_count);
return Status::OK();
}
template <>
template <>
Status ReduceKernel<true>::ComputeImpl<uint8_t, CUDNN_REDUCE_TENSOR_NO_INDICES>(OpKernelContext* ctx, cudnnReduceTensorOp_t cudnn_reduce_op) const {
typedef typename ToCudaType<uint8_t>::MappedType CudaT;
const Tensor* X = nullptr;
Tensor* Y = nullptr;
int64_t input_count = 0;
int64_t output_count = 0;
std::vector<int64_t> output_dims;
std::vector<int64_t> input_dims_cudnn;
std::vector<int64_t> 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<uint8_t>();
const auto* const src = X->template Data<uint8_t>();
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<float> temp_X = GetScratchBuffer<float>(input_count);
Impl_Cast<CudaT, float>(reinterpret_cast<const CudaT*>(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<uint32_t> indices_cuda = GetScratchBuffer<uint32_t>(indices_bytes);
IAllocatorUniquePtr<CudaT> workspace_cuda = GetScratchBuffer<CudaT>(workspace_bytes);
const auto one = Consts<float>::One;
const auto zero = Consts<float>::Zero;
auto temp_Y = GetScratchBuffer<float>(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<float, uint8_t>(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)

24
onnxruntime/test/featurizers_ops/truncated_svdtransformer_test.cc
View file

@ -46,9 +46,27 @@ void TruncatedSVDTransformerTestRowMajStandard() {
OpTester test("TruncatedSVDTransformer", 1, onnxruntime::kMSFeaturizersDomain);
test.AddInput<uint8_t>("State", {dim}, stream);
test.AddInput<T>("Input", {3, 3}, {-1, -1, 0, 0, -2, -1, -3, 0, -2});
test.AddOutput<T>("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<T> output = {-1.009107f, 0.626315f, -0.767745f, -0.965105f, 1.995869f, 0.291682f, -3.529165f, -0.724887f, 0.139759f};
test.AddOutput<T>("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<OrtValue>& 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<Tensor>();
ASSERT_EQ(tensor.Shape().NumDimensions(), 2);
ASSERT_EQ(tensor.Shape().Size(), 9);
Matrix output_matrix(tensor.Data<T>(), 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) {

5
onnxruntime/test/onnx/main.cc
View file

@ -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) {

2
onnxruntime/test/providers/cpu/controlflow/scan_test.cc
View file

@ -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<int64_t>("num_scan_inputs", 2);

174
onnxruntime/test/providers/cpu/nn/pool_op_test.cc
View file

@ -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<int64_t>{2});
test.AddAttribute("pads", vector<int64_t>{0, 0});
test.AddAttribute("kernel_shape", vector<int64_t>{2});
test.AddAttribute("storage_order", storage_order);
std::vector<int8_t> x_vals = {1, 2, 3, 4, 5, 6, 7, 8};
std::vector<int64_t> x_dims = {1, 2, 4};
std::vector<int64_t> expected_dims = {1, 2, 2};
std::vector<int8_t> expected_vals = {2, 4, 6, 8};
std::vector<int64_t> expected_indices = {1, 3, 5, 7};
test.AddInput<int8_t>("X", x_dims, x_vals);
test.AddOutput<int8_t>("Y", expected_dims, expected_vals);
test.AddOutput<int64_t>("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<int64_t>{2});
test.AddAttribute("pads", vector<int64_t>{0, 0});
test.AddAttribute("kernel_shape", vector<int64_t>{2});
test.AddAttribute("storage_order", storage_order);
std::vector<uint8_t> x_vals = {1, 2, 3, 4, 5, 6, 7, 8};
std::vector<int64_t> x_dims = {1, 2, 4};
std::vector<int64_t> expected_dims = {1, 2, 2};
std::vector<uint8_t> expected_vals = {2, 4, 6, 8};
std::vector<int64_t> expected_indices = {1, 3, 5, 7};
test.AddInput<uint8_t>("X", x_dims, x_vals);
test.AddOutput<uint8_t>("Y", expected_dims, expected_vals);
test.AddOutput<int64_t>("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<typename InputIter>
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<int64_t>{5, 5});
test.AddAttribute("pads", std::vector<int64_t>{2, 2, 2, 2});
std::vector<int64_t> output_shape = {1, 1, 5, 5};
std::vector<uint8_t> 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<uint8_t>("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<uint8_t>("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<int64_t>{2});
test.AddAttribute("kernel_shape", std::vector<int64_t>{2});
std::vector<int64_t> x_dims = {1, 3, 3};
std::vector<float> 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<int64_t>{2});
std::vector<int64_t> x_dims = {1, 3, 3};
std::vector<int8_t> x_vals = {0, 1, 2,
3, 4, 5,
6, 7, 8};
std::vector<int64_t> expected_dims = {1, 3, 2};
std::vector<int8_t> expected_vals = {1, 2,
4, 5,
7, 8};
test.AddInput<int8_t>("X", x_dims, x_vals);
test.AddOutput<int8_t>("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<int64_t>{2});
std::vector<int64_t> x_dims = {1, 3, 3};
std::vector<uint8_t> x_vals = {0, 1, 2,
3, 4, 5,
6, 7, 8};
std::vector<int64_t> expected_dims = {1, 3, 2};
std::vector<uint8_t> expected_vals = {1, 2,
4, 5,
7, 8};
test.AddInput<uint8_t>("X", x_dims, x_vals);
test.AddOutput<uint8_t>("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<int64_t>{1, 1});
test.AddAttribute("pads", vector<int64_t>{0, 0, 0, 0});
test.AddAttribute("kernel_shape", vector<int64_t>{2, 2});
test.AddAttribute("dilations", vector<int64_t>{2, 2});
std::vector<int8_t> x_vals = {
1, 3, 2, 4, -1,
5, 7, 6, 8, -2,
9, 11, 10, 12, -3,
13, 15, 14, 16, -4};
std::vector<int64_t> x_dims = {1, 1, 4, 5};
std::vector<int64_t> expected_dims = {1, 1, 2, 3};
std::vector<int8_t> expected_vals = {10, 12, 10, 14, 16, 14};
test.AddInput<int8_t>("X", x_dims, x_vals);
test.AddOutput<int8_t>("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<float>("X", x_dims, x_vals);
test.AddOutput<float>("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<int64_t>{2, 1});
test.AddAttribute("pads", vector<int64_t>{0, 0, 0, 0});
test.AddAttribute("kernel_shape", vector<int64_t>{2, 2});
test.AddAttribute("dilations", vector<int64_t>{2, 2});
std::vector<int8_t> x_vals = {
1, 3, 2, 4, -1,
5, 7, 6, 8, -2,
9, 11, 10, 12, -3,
13, 15, 14, 16, -4};
std::vector<int64_t> x_dims = {1, 1, 4, 5};
std::vector<int64_t> expected_dims = {1, 1, 1, 3};
std::vector<int8_t> expected_vals = {10, 12, 10};
test.AddInput<int8_t>("X", x_dims, x_vals);
test.AddOutput<int8_t>("Y", expected_dims, expected_vals);
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kAclExecutionProvider});
}
TEST(PoolTest, MaxPool_10_Dilation_Ceil1_2d) {
OpTester test("MaxPool", 10);

160
onnxruntime/test/providers/cpu/reduction/reduction_ops_test.cc
View file

@ -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<int64_t>{1, 2});
test.AddAttribute("keepdims", (int64_t)1);
test.AddInput<int8_t>("data", {3, 2, 2},
{1, 2,
3, 4,
5, 6,
7, 8,
9, 10,
11, 12});
test.AddOutput<int8_t>("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<int64_t>{1, 2});
test.AddAttribute("keepdims", (int64_t)1);
test.AddInput<uint8_t>("data", {3, 2, 2},
{1, 2,
3, 4,
5, 6,
7, 8,
9, 10,
11, 12});
test.AddOutput<uint8_t>("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<int64_t>{0, 2});
test.AddAttribute("keepdims", (int64_t)1);
test.AddInput<int8_t>("data", {3, 2, 2},
{1, 2,
3, 4,
5, 6,
7, 8,
9, 10,
11, 12});
test.AddOutput<int8_t>("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<int64_t>{0, 2});
test.AddAttribute("keepdims", (int64_t)1);
test.AddInput<uint8_t>("data", {3, 2, 2},
{1, 2,
3, 4,
5, 6,
7, 8,
9, 10,
11, 12});
test.AddOutput<uint8_t>("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<int32_t>("data", {3, 2, 2},
{2, 2,
3, 4,
5, 6,
7, 8,
10, 10,
11, 12});
test.AddOutput<int64_t>("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<int32_t>("data", {3, 2, 2},
{2, 4,
3, 4,
8, 6,
7, 8,
9, 13,
11, 12});
test.AddOutput<int64_t>("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<float>("data", {3, 2},
{1.0f, 2.0f,
6.0f, 5.0f,
9.0f, 10.0f});
test.AddOutput<int64_t>("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<float>("data", {3},
{1.0f, 2.0f, 3.0f});
test.AddOutput<int64_t>("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<int32_t>("data", {3, 2, 2},
{1, 2,
3, 4,
1, 6,
7, 8,
9, 10,
11, 12});
test.AddOutput<int64_t>("reduced", {2, 2},
{1, 0,
0, 0});
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kNGraphExecutionProvider});
}
#ifdef USE_CUDA
void test_reduce_apis(size_t size) {

4
onnxruntime/test/providers/cpu/tensor/gather_op_test.cc
View file

@ -76,7 +76,9 @@ TEST(GatherOpTest, Gather_invalid_index_cpu) {
test.AddInput<int32_t>("indices", {3}, {0LL, 1L, 1000L});
test.AddOutput<float>("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

8
onnxruntime/test/python/onnx_backend_test_series.py
View file

@ -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'):