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

Implement Min/Max/Clip(12) (#3410)

Browse source 
Implement Max/Min for opset 12.
  Add CLip(12) CPU impl.
  Implement Clip(12) for CPU and CUDA add tests
This commit is contained in:
Dmitri Smirnov 2020-04-06 14:24:59 -07:00 committed by GitHub
parent 7c69b1703b
commit c8f5e6e632
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
18 changed files with 657 additions and 209 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

8
include/onnxruntime/core/framework/op_kernel.h
View file

@ -359,4 +359,12 @@ using BuildKernelCreateInfoFn = KernelCreateInfo (*)();
static_cast<KernelCreatePtrFn>([](const OpKernelInfo& info) -> OpKernel* { return new __VA_ARGS__(info); })); \
}
// Use within macro definitions to create a custom vector of constraints.
// Example: #define REG_KERNEL(OP, VERSION, KERNEL_CLASS, Type, ...)
// .TypeConstraint("T", BuildKernelDefConstraints<Type, __VA_ARGS_>())
template <typename T, typename... Types>
inline std::vector<MLDataType> BuildKernelDefConstraints() {
return {DataTypeImpl::GetTensorType<T>(), DataTypeImpl::GetTensorType<Types>()...};
}
} // namespace onnxruntime

2
onnxruntime/core/optimizer/conv_activation_fusion.cc
View file

@ -108,7 +108,7 @@ Status ConvActivationFusion::ApplyImpl(Graph& graph, bool& modified, int graph_l
!graph_utils::IsSupportedOptypeVersionAndDomain(next_node, "Tanh", {6})) {
if (graph_utils::IsSupportedOptypeVersionAndDomain(next_node, "LeakyRelu", {6})) {
activation_params.push_back(graph_utils::GetNodeAttribute(next_node, "alpha")->f());
} else if (graph_utils::IsSupportedOptypeVersionAndDomain(next_node, "Clip", {6, 11})) {
} else if (graph_utils::IsSupportedOptypeVersionAndDomain(next_node, "Clip", {6, 11, 12})) {
float min, max;
if (GetClipConstantMinMax(graph, next_node, min, max)) {
activation_params.push_back(min);

2
onnxruntime/core/optimizer/nchwc_transformer.cc
View file

@ -925,7 +925,7 @@ void NchwcTransformerImpl::Transform(Node& node) {
if (graph_utils::IsSupportedOptypeVersionAndDomain(node, "Conv", {1, 11}) ||
graph_utils::IsSupportedOptypeVersionAndDomain(node, "FusedConv", {1}, kMSDomain)) {
TransformConv(node);
} else if (graph_utils::IsSupportedOptypeVersionAndDomain(node, "MaxPool", {1, 8, 10, 11}) ||
} else if (graph_utils::IsSupportedOptypeVersionAndDomain(node, "MaxPool", {1, 8, 10, 11, 12}) ||
graph_utils::IsSupportedOptypeVersionAndDomain(node, "AveragePool", {1, 7, 10, 11}) ||
graph_utils::IsSupportedOptypeVersionAndDomain(node, "GlobalMaxPool", {1}) ||
graph_utils::IsSupportedOptypeVersionAndDomain(node, "GlobalAveragePool", {1})) {

2
onnxruntime/core/optimizer/relu_clip_fusion.cc
View file

@ -122,7 +122,7 @@ bool FuseReluClip::SatisfyCondition(const Graph& graph, const Node& node, const
// as Clip will apply the minimum. If the Clip 'min' value is < 0 we need
// to update it to 0 to apply what the Relu would have done. We do that in Apply.
const auto& next_node = *node.OutputNodesBegin();
if (!graph_utils::IsSupportedOptypeVersionAndDomain(next_node, "Clip", {6, 11}) ||
if (!graph_utils::IsSupportedOptypeVersionAndDomain(next_node, "Clip", {6, 11, 12}) ||
next_node.GetExecutionProviderType() != node.GetExecutionProviderType()) {
return false;
}

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

@ -78,10 +78,9 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain,
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 6, 7, float, Sum);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, float, Sum);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 6, 7, float, Min);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, float, Min);
class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, 11, Min);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 6, 7, float, Max);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, float, Max);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, double, Max);
class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, 11, Max);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 1, Not);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 7, And);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 7, Or);
@ -322,7 +321,7 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain,
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 10, ReverseSequence);
// opset 11
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, Clip);
class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 11, Clip);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, float, CumSum);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, double, CumSum);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, int32_t, CumSum);
@ -441,6 +440,10 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain,
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, Clip);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, Min);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, Max);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, MaxPool);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, float, ReduceMax);
@ -543,11 +546,10 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, float, Sum)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 6, 7,
float, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, float, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, 11, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 6, 7,
float, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, float, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, double, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 8, 11, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 1, Not)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 7, And)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 7, Or)>,
@ -1004,7 +1006,8 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 10, ReverseSequence)>,
//opset 11
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, Clip)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, 11, Clip)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, float,
CumSum)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 11, double,
@ -1142,6 +1145,12 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int32_t,
ArgMin)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12,Clip)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, MaxPool)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, float,

52
onnxruntime/core/providers/cpu/math/clip.cc
View file

@ -2,6 +2,7 @@
// Licensed under the MIT License.
#include "core/providers/cpu/math/clip.h"
#include "core/framework/data_types_internal.h"
namespace onnxruntime {
@ -12,10 +13,57 @@ ONNX_CPU_OPERATOR_VERSIONED_KERNEL(
KernelDefBuilder().MayInplace(0, 0).TypeConstraint("T", DataTypeImpl::GetTensorType<float>()),
Clip_6<float>);
ONNX_CPU_OPERATOR_KERNEL(
ONNX_CPU_OPERATOR_VERSIONED_KERNEL(
Clip,
11,
11,
KernelDefBuilder().MayInplace(0, 0).TypeConstraint("T", DataTypeImpl::GetTensorType<float>()),
Clip<float>);
Clip);
#define REG_KERNEL_NONTEMPL(OP_TYPE, VERSION, KERNEL_CLASS, ...) \
ONNX_CPU_OPERATOR_KERNEL( \
OP_TYPE, \
VERSION, \
KernelDefBuilder() \
.MayInplace(0, 0) \
.TypeConstraint("T", BuildKernelDefConstraints<__VA_ARGS__>()), \
KERNEL_CLASS);
REG_KERNEL_NONTEMPL(Clip, 12, Clip, float, double, int8_t, uint8_t, int64_t, uint64_t);
template <typename T>
struct Clip::ComputeImpl {
void operator()(const Tensor* X, const Tensor* min, const Tensor* max, Tensor* Y) const {
auto min_val = std::numeric_limits<T>::lowest();
auto max_val = std::numeric_limits<T>::max();
if (min) {
ORT_ENFORCE(min->Shape().NumDimensions() == 0, "min should be a scalar.");
min_val = *(min->template Data<T>());
}
if (max) {
ORT_ENFORCE(max->Shape().NumDimensions() == 0, "max should be a scalar.");
max_val = *(max->template Data<T>());
}
EigenVectorMap<T>(Y->template MutableData<T>(), Y->Shape().Size()) =
ConstEigenVectorMap<T>(X->template Data<T>(), X->Shape().Size())
.cwiseMax(min_val)
.cwiseMin(max_val);
}
};
Status Clip::Compute(OpKernelContext* ctx) const {
const auto* X = ctx->Input<Tensor>(0);
const auto* min = ctx->Input<Tensor>(1);
const auto* max = ctx->Input<Tensor>(2);
Tensor* Y = ctx->Output(0, X->Shape());
utils::MLTypeCallDispatcher<ComputeImpl, float, double, int8_t, uint8_t, int64_t, uint64_t>
t_disp(X->GetElementType());
t_disp.Invoke(X, min, max, Y);
return Status::OK();
}
} // namespace onnxruntime

66
onnxruntime/core/providers/cpu/math/clip.h
View file

@ -9,59 +9,53 @@
namespace onnxruntime {
template <typename T>
class Clip_6 final : public OpKernel {
namespace clip_internal {
template<typename T>
class Clip_6Base {
public:
Clip_6(const OpKernelInfo& info) : OpKernel(info) {
ORT_ENFORCE(info.GetAttr<T>("max", &max_).IsOK());
ORT_ENFORCE(info.GetAttr<T>("min", &min_).IsOK());
explicit Clip_6Base(const OpKernelInfo& info) {
auto min_val = std::numeric_limits<T>::lowest();
auto max_val = std::numeric_limits<T>::max();
info.GetAttrOrDefault("min", &min_, min_val);
info.GetAttrOrDefault("max", &max_, max_val);
ORT_ENFORCE(min_ <= max_);
}
Status Compute(OpKernelContext* ctx) const override {
const auto* X = ctx->Input<Tensor>(0);
Tensor* Y = ctx->Output(0, X->Shape());
EigenVectorMap<T>(Y->template MutableData<T>(), Y->Shape().Size()) =
ConstEigenVectorMap<T>(X->template Data<T>(), X->Shape().Size())
.cwiseMax(min_)
.cwiseMin(max_);
return Status::OK();
}
private:
protected:
T max_;
T min_;
};
} // namespace clip_internal
template <typename T>
class Clip final : public OpKernel {
class Clip_6 final : public clip_internal::Clip_6Base<T>, public OpKernel {
public:
Clip(const OpKernelInfo& info) : OpKernel(info) {
explicit Clip_6(const OpKernelInfo& info) : clip_internal::Clip_6Base<T>(info), OpKernel(info) {
}
Status Compute(OpKernelContext* ctx) const override {
const auto* X = ctx->Input<Tensor>(0);
const auto* min = ctx->Input<Tensor>(1);
const auto* max = ctx->Input<Tensor>(2);
Tensor* Y = ctx->Output(0, X->Shape());
auto min_val = -std::numeric_limits<T>::infinity();
auto max_val = std::numeric_limits<T>::infinity();
if (min) {
ORT_ENFORCE(min->Shape().NumDimensions() == 0, "min should be a scalar.");
min_val = *(min->template Data<T>());
}
if (max) {
ORT_ENFORCE(max->Shape().NumDimensions() == 0, "max should be a scalar.");
max_val = *(max->template Data<T>());
}
EigenVectorMap<T>(Y->template MutableData<T>(), Y->Shape().Size()) =
ConstEigenVectorMap<T>(X->template Data<T>(), X->Shape().Size())
.cwiseMax(min_val)
.cwiseMin(max_val);
.cwiseMax(this->min_)
.cwiseMin(this->max_);
return Status::OK();
}
};
// Since version 11. Min and Max are inputs
// version 12 adds type support
class Clip final : public OpKernel {
public:
explicit Clip(const OpKernelInfo& info) : OpKernel(info) {
}
Status Compute(OpKernelContext* ctx) const override;
private:
template<typename T>
struct ComputeImpl;
};
} // namespace onnxruntime

75
onnxruntime/core/providers/cpu/math/element_wise_ops.cc
View file

@ -1,6 +1,7 @@
// Copyright (c) Microsoft Corporation. All rights reserved.
// Licensed under the MIT License.
#include "core/framework/data_types_internal.h"
#include "core/providers/cpu/math/element_wise_ops.h"
#include <unsupported/Eigen/SpecialFunctions>
#include "core/util/math.h"
@ -46,6 +47,27 @@ namespace onnxruntime {
.TypeConstraint("T1", DataTypeImpl::GetTensorType<bool>()), \
KERNEL_CLASS<TYPE>);
// var args are type constraints
#define REG_ELEMENTWISE_KERNEL_NONT(OP_TYPE, VERSION, KERNEL_CLASS, ...) \
ONNX_CPU_OPERATOR_KERNEL( \
OP_TYPE, \
VERSION, \
KernelDefBuilder() \
.TypeConstraint("T", BuildKernelDefConstraints<__VA_ARGS__>()) \
.TypeConstraint("T1", BuildKernelDefConstraints<__VA_ARGS__>()), \
KERNEL_CLASS);
// var args are type constraints
#define REG_ELEMENTWISE_VERSIONED_KERNEL_NONT(OP_TYPE, VERSION_FROM, VERSION_TO, KERNEL_CLASS, ...) \
ONNX_CPU_OPERATOR_VERSIONED_KERNEL( \
OP_TYPE, \
VERSION_FROM, \
VERSION_TO, \
KernelDefBuilder() \
.TypeConstraint("T", BuildKernelDefConstraints<__VA_ARGS__>()) \
.TypeConstraint("T1", BuildKernelDefConstraints<__VA_ARGS__>()), \
KERNEL_CLASS);
REG_ELEMENTWISE_TYPED_KERNEL(Add, 7, float, Add);
REG_ELEMENTWISE_TYPED_KERNEL(Add, 7, double, Add);
REG_ELEMENTWISE_TYPED_KERNEL(Add, 7, int32_t, Add);
@ -103,12 +125,13 @@ REG_ELEMENTWISE_TYPED_KERNEL(Log, 6, float, Log);
REG_ELEMENTWISE_VERSIONED_TYPED_KERNEL(Sum, 6, 7, float, Sum_6);
REG_ELEMENTWISE_TYPED_KERNEL(Sum, 8, float, Sum_8);
REG_ELEMENTWISE_VERSIONED_TYPED_KERNEL(Min, 6, 7, float, Min_6);
REG_ELEMENTWISE_TYPED_KERNEL(Min, 8, float, Min_8);
REG_ELEMENTWISE_VERSIONED_TYPED_KERNEL(Max, 6, 7, float, Max_6);
REG_ELEMENTWISE_TYPED_KERNEL(Max, 8, float, Max_8);
REG_ELEMENTWISE_TYPED_KERNEL(Max, 8, double, Max_8);
REG_ELEMENTWISE_VERSIONED_KERNEL_NONT(Max, 8, 11, Max_8, float, double);
REG_ELEMENTWISE_KERNEL_NONT(Max, 12, Max_8, float, double, MLFloat16, int32_t, uint32_t, int64_t, uint64_t);
REG_ELEMENTWISE_VERSIONED_TYPED_KERNEL(Min, 6, 7, float, Min_6);
REG_ELEMENTWISE_VERSIONED_KERNEL_NONT(Min, 8, 11, Min_8, float);
REG_ELEMENTWISE_KERNEL_NONT(Min, 12, Min_8, float, double, MLFloat16, int32_t, uint32_t, int64_t, uint64_t);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Less, 7, 9, float, Less);
REG_ELEMENTWISE_LOGICALOP_VERSIONED_TYPED_KERNEL(Less, 7, 9, double, Less);
@ -343,13 +366,21 @@ Status Min_6<float>::Compute(OpKernelContext* ctx) const {
return Status::OK();
}
template <>
Status Min_8<float>::Compute(OpKernelContext* context) const {
return BroadcastVariadic<float, float>(
Node(), *context,
[](EigenVectorMap<float> output, float input0, ConstEigenVectorMap<float> input1) { output = input1.array().min(input0); },
[](EigenVectorMap<float> output, ConstEigenVectorMap<float> input0, float input1) { output = input0.array().min(input1); },
[](EigenVectorMap<float> output, ConstEigenVectorMap<float> input0, ConstEigenVectorMap<float> input1) { output = input0.array().min(input1.array()); });
template <typename T>
struct Min_8::ComputeImpl {
Status operator()(const Min_8* inst, OpKernelContext* context) const {
return BroadcastVariadic<T, T>(
inst->Node(), *context,
[](EigenVectorMap<T> output, T input0, ConstEigenVectorMap<T> input1) { output = input1.array().min(input0); },
[](EigenVectorMap<T> output, ConstEigenVectorMap<T> input0, T input1) { output = input0.array().min(input1); },
[](EigenVectorMap<T> output, ConstEigenVectorMap<T> input0, ConstEigenVectorMap<T> input1) { output = input0.array().min(input1.array()); });
}
};
Status Min_8::Compute(OpKernelContext* context) const {
utils::MLTypeCallDispatcherRet<Status, ComputeImpl, float, double, MLFloat16, int32_t, uint32_t, int64_t, uint64_t>
t_disp(context->Input<Tensor>(0)->GetElementType());
return t_disp.Invoke(this, context);
}
template <>
@ -371,12 +402,20 @@ Status Max_6<float>::Compute(OpKernelContext* ctx) const {
}
template <typename T>
Status Max_8<T>::Compute(OpKernelContext* context) const {
return BroadcastVariadic<T, T>(
Node(), *context,
[](EigenVectorMap<T> output, T input0, ConstEigenVectorMap<T> input1) { output = input1.array().max(input0); },
[](EigenVectorMap<T> output, ConstEigenVectorMap<T> input0, T input1) { output = input0.array().max(input1); },
[](EigenVectorMap<T> output, ConstEigenVectorMap<T> input0, ConstEigenVectorMap<T> input1) { output = input0.array().max(input1.array()); });
struct Max_8::ComputeImpl {
Status operator()(const Max_8* inst, OpKernelContext* context) const {
return BroadcastVariadic<T, T>(
inst->Node(), *context,
[](EigenVectorMap<T> output, T input0, ConstEigenVectorMap<T> input1) { output = input1.array().max(input0); },
[](EigenVectorMap<T> output, ConstEigenVectorMap<T> input0, T input1) { output = input0.array().max(input1); },
[](EigenVectorMap<T> output, ConstEigenVectorMap<T> input0, ConstEigenVectorMap<T> input1) { output = input0.array().max(input1.array()); });
}
};
Status Max_8::Compute(OpKernelContext* context) const {
utils::MLTypeCallDispatcherRet<Status, ComputeImpl, float, double, MLFloat16, int32_t, uint32_t, int64_t, uint64_t>
t_disp(context->Input<Tensor>(0)->GetElementType());
return t_disp.Invoke(this, context);
}
Status Not::Compute(OpKernelContext* context) const {

17
onnxruntime/core/providers/cpu/math/element_wise_ops.h
View file

@ -165,13 +165,19 @@ class Min_6 final : public OpKernel {
Status Compute(OpKernelContext* context) const override;
};
template <typename T>
// Max versions 8 - 12
// Version 8 added broadcast
// Version 12 added types support
class Min_8 final : public OpKernel {
public:
Min_8(const OpKernelInfo& info) : OpKernel(info) {
}
Status Compute(OpKernelContext* context) const override;
private:
template <typename T>
struct ComputeImpl;
};
template <typename T>
@ -183,13 +189,20 @@ class Max_6 final : public OpKernel {
Status Compute(OpKernelContext* context) const override;
};
template <typename T>
// Max versions 8 - 12
// Version 8 added broadcast
// Version 12 added types support
class Max_8 final : public OpKernel {
public:
Max_8(const OpKernelInfo& info) : OpKernel(info) {
}
Status Compute(OpKernelContext* context) const override;
private:
template<typename T>
struct ComputeImpl;
};
class Not final : public OpKernel {

83
onnxruntime/core/providers/cuda/cu_inc/common.cuh
View file

@ -129,7 +129,7 @@ template <>
__device__ __inline__ double _Round(double a) { return rint(a); }
template <>
__device__ __inline__ half _Round(half a) {
__device__ __inline__ half _Round(half a) {
#if __CUDA_ARCH__ < 530
return half(rintf((float)a));
#else
@ -196,9 +196,33 @@ __device__ __inline__ half _Pow(half a, half b) { return half(powf((float)a, (fl
template <typename T>
__device__ __inline__ T _Min(T a, T b) { return a < b ? a : b; }
template <>
__device__ __inline__ int _Min(int a, int b) { return a < b ? a : b; }
template <>
__device__ __inline__ unsigned int _Min(unsigned int a, unsigned int b) { return a < b ? a : b; }
template <>
__device__ __inline__ long long _Min(long long a, long long b) { return a < b ? a : b; }
template <>
__device__ __inline__ unsigned long long _Min(unsigned long long a, unsigned long long b) { return a < b ? a : b; }
template <typename T>
__device__ __inline__ T _Max(T a, T b) { return a > b ? a : b; }
template <>
__device__ __inline__ int _Max(int a, int b) { return a > b ? a : b; }
template <>
__device__ __inline__ unsigned int _Max(unsigned int a, unsigned int b) { return a > b ? a : b; }
template <>
__device__ __inline__ long long _Max(long long a, long long b) { return a > b ? a : b; }
template <>
__device__ __inline__ unsigned long long _Max(unsigned long long a, unsigned long long b) { return a > b ? a : b; }
template <typename T>
__device__ __inline__ T _Abs(T a) { return a > (T)0 ? a : -a; }
@ -226,12 +250,6 @@ __device__ __inline__ T _Gelu(T a) {
#define CUDA_LONG int32_t
#endif
#define IDX2C(i, j, ld) (((j) * (ld)) + (i)) // 0 based indexing
// ---------------------------------------------------------------------------
// GridDim -- helper to choose the CUDA grid dimensions
// ---------------------------------------------------------------------------
template <class INT, class INT2>
static INT CeilDiv(INT a, INT2 b) // ceil(a/b)
{
@ -241,54 +259,13 @@ static INT CeilDiv(INT a, INT2 b) // ceil(a/b)
struct GridDim {
enum : CUDA_LONG {
maxThreadsPerBlock = 256, // max threads per block
maxWarpsPerBlock = 32, // max warps per block
maxElementsPerThread = 4, // max element processed per thread
};
// use these for launching
// GridDim grid(NN);
// kernel<<<grid.m_blocksPerGrid, grid.m_threadsPerBlock, ...>>>(...)
int blocks_per_grid_, threads_per_block_; // (these may in the future be extended to multi-dimensional ones)
CUDA_LONG N_;
GridDim(CUDA_LONG N) // linear grid
{
N_ = N;
if (N == 0) // CUDA will fail to launch with 0 blocks
N = 1;
// get device information
const auto& props = DeviceProp::GetDeviceProps();
CUDA_LONG numProcs = props.multiProcessorCount;
CUDA_LONG warpSize = props.warpSize;
// distribute warps evenly over processors
CUDA_LONG warpsPerProc = CeilDiv(N, numProcs * warpSize);
// if too many warps per block then reduce #warps
// This limits the number of threads to 512.
if (warpsPerProc > maxWarpsPerBlock) {
CUDA_LONG overBy = CeilDiv(warpsPerProc, maxWarpsPerBlock); // we are over by this factor
warpsPerProc = CeilDiv(warpsPerProc, overBy);
}
// put it back together
threads_per_block_ = warpsPerProc * warpSize; // =a multiple of 32 that is as close to 1024 as makes sense given NN
blocks_per_grid_ = CeilDiv(N, threads_per_block_);
if (blocks_per_grid_ == 1)
threads_per_block_ = N; // don't launch more than necessary
assert(blocks_per_grid_ * threads_per_block_ >= N);
}
// compute our location on the grid
static __device__ CUDA_LONG GetLinearThreadId() {
return blockDim.x * blockIdx.x + threadIdx.x;
}
};
#define CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, N) \
CUDA_LONG id = GridDim::GetLinearThreadId(); \
if (id >= N) \
#define CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, N) \
CUDA_LONG id = blockDim.x * blockIdx.x + threadIdx.x; \
if (id >= N) \
return;
// CUDA_KERNEL_ASSERT is a macro that wraps an assert() call inside cuda kernels.
@ -296,9 +273,9 @@ struct GridDim {
// See http://docs.nvidia.com/cuda/cuda-c-programming-guide/#assertion
#if defined(__APPLE__) || defined(__HIP_PLATFORM_HCC__)
#define CUDA_KERNEL_ASSERT(...)
#else // __APPLE__
#else // __APPLE__
#define CUDA_KERNEL_ASSERT(...) assert(__VA_ARGS__)
#endif // __APPLE__
#endif // __APPLE__
} // namespace cuda
} // namespace onnxruntime

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

@ -318,15 +318,29 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain,
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, float, Max);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, double, Max);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, MLFloat16, Max);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, float, Max);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, double, Max);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, MLFloat16, Max);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, float, Max);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, double, Max);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, MLFloat16, Max);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int32_t, Max);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int64_t, Max);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint32_t, Max);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint64_t, Max);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, Max);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, Max);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, Max);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, float, Min);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, double, Min);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, MLFloat16, Min);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, float, Min);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, double, Min);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, MLFloat16, Min);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, float, Min);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, double, Min);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, MLFloat16, Min);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int32_t, Min);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int64_t, Min);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint32_t, Min);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint64_t, Min);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, Min);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, Min);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, Min);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 8, float, Greater);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 8, double, Greater);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 8, MLFloat16, Greater);
@ -690,7 +704,7 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain,
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, Resize);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, int32_t, Resize);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, uint8_t, Resize);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, Clip);
class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, Clip);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, Pad);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, Pad);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, Pad);
@ -707,6 +721,8 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain,
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, CumSum);
// OpSet 12
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, Clip);
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);
@ -812,15 +828,29 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, float, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, double, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, MLFloat16, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, float, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, double, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, MLFloat16, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, float, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, double, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, MLFloat16, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int32_t, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int64_t, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint32_t, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint64_t, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, Max)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, float, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, double, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 6, 7, MLFloat16, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, float, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, double, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, MLFloat16, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, float, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, double, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 8, 11, MLFloat16, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int32_t, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int64_t, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint32_t, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint64_t, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, Min)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 8, float, Greater)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 8, double, Greater)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 8, MLFloat16, Greater)>,
@ -1187,7 +1217,7 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, Resize)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, int32_t, Resize)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, uint8_t, Resize)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, Clip)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, 11, Clip)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, float, Pad)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, double, Pad)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, MLFloat16, Pad)>,
@ -1200,6 +1230,8 @@ static void RegisterCudaKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 11, CumSum)>,
// OpSet 12
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, Clip)>,
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)>,
@ -1368,12 +1400,12 @@ CUDAExecutionProvider::GetCapability(const onnxruntime::GraphViewer& graph,
for (auto registry : kernel_registries) {
cuda_kernel_def = registry->TryFindKernel(node, Type());
// atleast one registry has a CUDA kernel for this node
// atleast one registry has a CUDA kernel for this node
if (cuda_kernel_def)
break;
break;
}
// none of the provided registries has a CUDA kernel for this node
// none of the provided registries has a CUDA kernel for this node
if (cuda_kernel_def == nullptr) {
// node is not in cuda exeuction provider if no kernel def found,
// or if other execution provider already assigned to it

6
onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc
View file

@ -384,9 +384,11 @@ BINARY_LOGICALOP_REGISTER_UZILHFD(Greater, 9)
BINARY_OP_REGISTER_VERSIONED_OIL(Equal, 7, 10)
BINARY_OP_REGISTER_OIL(Equal, 11)
BINARY_OP_REGISTER_VERSIONED_HFD(Greater, 7, 8)
BINARY_OP_REGISTER_HFD(Max, 8)
BINARY_OP_REGISTER_UZILHFD(Max, 12)
BINARY_OP_REGISTER_VERSIONED_HFD(Max, 8, 11)
BINARY_OP_REGISTER_VERSIONED_HFD(Max, 6, 7)
BINARY_OP_REGISTER_HFD(Min, 8)
BINARY_OP_REGISTER_UZILHFD(Min, 12)
BINARY_OP_REGISTER_VERSIONED_HFD(Min, 8, 11)
BINARY_OP_REGISTER_VERSIONED_HFD(Min, 6, 7)
BINARY_LOGICALOP_REGISTER_UZILHFD(Less, 9)
BINARY_OP_REGISTER_VERSIONED_HFD(Less, 7, 8)

4
onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu
View file

@ -92,8 +92,8 @@ SPECIALIZED_BINARY_ELEMENTWISE_IMPL(Xor, bool)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_HFD(PRelu)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Greater)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_OIL(Equal)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_HFD(Max)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_HFD(Min)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Max)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Min)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Less)
} // namespace cuda

139
onnxruntime/core/providers/cuda/math/clip.cc
View file

@ -8,61 +8,114 @@
namespace onnxruntime {
namespace cuda {
#define REGISTER_KERNEL_TYPED(T) \
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX( \
Clip, \
kOnnxDomain, \
6, \
10, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder() \
.TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Clip<T>); \
ONNX_OPERATOR_TYPED_KERNEL_EX( \
Clip, \
kOnnxDomain, \
11, \
T, \
kCudaExecutionProvider, \
KernelDefBuilder() \
.InputMemoryType<OrtMemTypeCPUInput>(1) \
.InputMemoryType<OrtMemTypeCPUInput>(2) \
.TypeConstraint("T", DataTypeImpl::GetTensorType<T>()), \
Clip<T>);
ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_EX(
Clip,
kOnnxDomain,
6,
10,
float,
kCudaExecutionProvider,
KernelDefBuilder()
.TypeConstraint("T", DataTypeImpl::GetTensorType<float>()),
Clip_6<float>);
ONNX_OPERATOR_VERSIONED_KERNEL_EX(
Clip,
kOnnxDomain,
11, 11,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(1)
.InputMemoryType<OrtMemTypeCPUInput>(2)
.TypeConstraint("T", DataTypeImpl::GetTensorType<float>()),
Clip);
ONNX_OPERATOR_KERNEL_EX(
Clip,
kOnnxDomain,
12,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(1)
.InputMemoryType<OrtMemTypeCPUInput>(2)
.TypeConstraint("T", BuildKernelDefConstraints<float, double, int8_t, uint8_t, int64_t, uint64_t>()),
Clip);
template <typename T>
Status Clip<T>::ComputeInternal(OpKernelContext* ctx) const {
T min_val = min_;
T max_val = max_;
if (is_min_max_input_) {
const auto* min_input = ctx->Input<Tensor>(1);
const auto* max_input = ctx->Input<Tensor>(2);
if (min_input) {
ORT_ENFORCE(min_input->Shape().NumDimensions() == 0, "min should be a scalar.");
min_val = *(min_input->template Data<T>());
}
if (max_input) {
ORT_ENFORCE(max_input->Shape().NumDimensions() == 0, "max should be a scalar.");
max_val = *(max_input->template Data<T>());
}
ORT_ENFORCE(min_val <= max_val);
}
Status Clip_6<T>::ComputeInternal(OpKernelContext* ctx) const {
const Tensor& X = *ctx->Input<Tensor>(0);
const TensorShape& input_shape{X.Shape()};
Tensor* Y = ctx->Output(0, input_shape);
size_t count = input_shape.Size();
const size_t count = input_shape.Size();
if (count > 0) {
auto* y_data = Y->template MutableData<T>();
const auto* x_data = X.template Data<T>();
ClipImpl<T>(x_data, y_data, min_val, max_val, count);
ClipImpl<T>(x_data, y_data, this->min_, this->max_, count);
}
return Status::OK();
}
namespace clip_internal {
template <typename T>
struct LowMax {
constexpr static T low() {
return std::numeric_limits<T>::lowest();
}
constexpr static T max() {
return std::numeric_limits<T>::max();
}
};
template <>
struct LowMax<MLFloat16> {
static MLFloat16 low() {
return MLFloat16(math::floatToHalf(std::numeric_limits<float>::lowest()));
}
static MLFloat16 max() {
return MLFloat16(math::floatToHalf(std::numeric_limits<float>::max()));
}
};
} // namespace clip_internal
template <typename T>
struct Clip::ComputeImpl {
void operator()(const Tensor* X, const Tensor* min, const Tensor* max, Tensor* Y) const {
auto min_val = clip_internal::LowMax<T>::low();
auto max_val = clip_internal::LowMax<T>::max();
// 1-2 Input on CPU
if (min) {
ORT_ENFORCE(min->Shape().NumDimensions() == 0, "min should be a scalar.");
min_val = *(min->template Data<T>());
}
if (max) {
ORT_ENFORCE(max->Shape().NumDimensions() == 0, "max should be a scalar.");
max_val = *(max->template Data<T>());
}
const size_t count = X->Shape().Size();
if (count > 0) {
auto* y_data = Y->template MutableData<T>();
const auto* x_data = X->template Data<T>();
ClipImpl<T>(x_data, y_data, min_val, max_val, count);
}
}
};
Status Clip::ComputeInternal(OpKernelContext* ctx) const {
const auto* X = ctx->Input<Tensor>(0);
const auto* min = ctx->Input<Tensor>(1);
const auto* max = ctx->Input<Tensor>(2);
Tensor* Y = ctx->Output(0, X->Shape());
utils::MLTypeCallDispatcher<ComputeImpl, float, double, int8_t, uint8_t, int64_t, uint64_t>
t_disp(X->GetElementType());
t_disp.Invoke(X, min, max, Y);
return Status::OK();
}
REGISTER_KERNEL_TYPED(float)
} // namespace cuda
} // namespace onnxruntime

32
onnxruntime/core/providers/cuda/math/clip.h
View file

@ -3,36 +3,32 @@
#pragma once
#include "core/providers/cuda/cuda_common.h"
#include "core/providers/cpu/math/clip.h"
namespace onnxruntime {
namespace cuda {
template <typename T>
class Clip_6 final : public onnxruntime::clip_internal::Clip_6Base<T>, public CudaKernel {
public:
explicit Clip_6(const OpKernelInfo& info) : onnxruntime::clip_internal::Clip_6Base<T>(info), CudaKernel{info} {
}
Status ComputeInternal(OpKernelContext* context) const override;
};
// Since version 11. Min and Max are inputs
// version 12 adds type support
class Clip final : public CudaKernel {
public:
Clip(const OpKernelInfo& info) : CudaKernel{info}, is_min_max_input_(false) {
int start_version;
int end_version;
info.GetKernelDef().SinceVersion(&start_version, &end_version);
if (start_version < 11) {
auto min_val = -std::numeric_limits<T>::infinity();
auto max_val = std::numeric_limits<T>::infinity();
info.GetAttrOrDefault("min", &min_, min_val);
info.GetAttrOrDefault("max", &max_, max_val);
ORT_ENFORCE(min_ <= max_);
} else {
min_ = -std::numeric_limits<T>::infinity();
max_ = std::numeric_limits<T>::infinity();
is_min_max_input_ = true;
}
explicit Clip(const OpKernelInfo& info) : CudaKernel{info} {
}
Status ComputeInternal(OpKernelContext* context) const override;
private:
T min_, max_;
bool is_min_max_input_;
template<typename T>
struct ComputeImpl;
};
} // namespace cuda

4
onnxruntime/core/providers/cuda/math/clip_impl.cu
View file

@ -27,6 +27,10 @@ void ClipImpl(const T* input_data, T* output_data, T min, T max, size_t count) {
template void ClipImpl<float>(const float* input_data, float* output_data, float min, float max, size_t count);
template void ClipImpl<double>(const double* input_data, double* output_data, double min, double max, size_t count);
template void ClipImpl<MLFloat16>(const MLFloat16* input_data, MLFloat16* output_data, MLFloat16 min, MLFloat16 max, size_t count);
template void ClipImpl<int8_t>(const int8_t* input_data, int8_t* output_data, int8_t min, int8_t max, size_t count);
template void ClipImpl<uint8_t>(const uint8_t* input_data, uint8_t* output_data, uint8_t min, uint8_t max, size_t count);
template void ClipImpl<int64_t>(const int64_t* input_data, int64_t* output_data, int64_t min, int64_t max, size_t count);
template void ClipImpl<uint64_t>(const uint64_t* input_data, uint64_t* output_data, uint64_t min, uint64_t max, size_t count);
} // namespace cuda
} // namespace onnxruntime

77
onnxruntime/test/providers/cpu/math/clip_test.cc
View file

@ -26,7 +26,7 @@ TEST(MathOpTest, Clip_6) {
}
TEST(MathOpTest, Clip_Default) {
OpTester test("Clip", 11);
OpTester test("Clip", -1);
std::vector<int64_t> dims{3, 3};
test.AddInput<float>("X", dims,
@ -38,10 +38,79 @@ TEST(MathOpTest, Clip_Default) {
-1.3f, 3.5f, 64.0f,
-5.4f, 9.3f, 82.4f});
// nGraph does not support Clip opset 11 yet.
// nGraph does not support Clip opset 12 yet.
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kNGraphExecutionProvider});
}
TEST(MathOpTest, Clip_Default_int8) {
OpTester test("Clip", -1);
std::vector<int64_t> dims{3, 3};
test.AddInput<int8_t>("X", dims,
{11, 4, 127,
-1, 3, 64,
-5, 9, 82});
test.AddOutput<int8_t>("Y", dims,
{11, 4, 127,
-1, 3, 64,
-5, 9, 82});
// TensorRT, nGraph does not support Clip opset 12 yet.
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kNGraphExecutionProvider});
}
TEST(MathOpTest, Clip_Default_uint8) {
OpTester test("Clip", -1);
std::vector<int64_t> dims{3, 3};
test.AddInput<uint8_t>("X", dims,
{11, 4, 255,
1, 3, 64,
5, 9, 82});
test.AddOutput<uint8_t>("Y", dims,
{11, 4, 255,
1, 3, 64,
5, 9, 82});
// TensorRT, nGraph does not support Clip opset 12 yet.
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kNGraphExecutionProvider});
}
TEST(MathOpTest, Clip_Default_int64) {
OpTester test("Clip", -1);
std::vector<int64_t> dims{3, 3};
test.AddInput<int64_t>("X", dims,
{11, 4, 432,
-1, 3, 64,
-5, 9, 82});
test.AddOutput<int64_t>("Y", dims,
{11, 4, 432,
-1, 3, 64,
-5, 9, 82});
// TensorRT, nGraph does not support Clip opset 12 yet.
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kNGraphExecutionProvider});
}
TEST(MathOpTest, Clip_Default_uint64) {
OpTester test("Clip", -1);
std::vector<int64_t> dims{3, 3};
test.AddInput<uint64_t>("X", dims,
{11, 4, 432,
1, 3, 64,
5, 9, 82});
test.AddOutput<uint64_t>("Y", dims,
{11, 4, 432,
1, 3, 64,
5, 9, 82});
// TensorRT, nGraph does not support Clip opset 12 yet.
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kNGraphExecutionProvider});
}
TEST(MathOpTest, Clip) {
OpTester test("Clip", 11);
@ -57,14 +126,14 @@ TEST(MathOpTest, Clip) {
-5.0f, 0.0f, 5.0f,
-5.0f, 2.0f, 5.0f});
// nGraph and Tensorrt does not support Clip opset 11 yet.
// TensorRT, nGraph and Tensorrt does not support Clip opset 11 yet.
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kNGraphExecutionProvider, kTensorrtExecutionProvider});
}
TEST(MathOpTest, ClipDimWithZero) {
std::vector<int64_t> dims{3, 0}; // dim with value of zero should be handled
OpTester test("Clip", 11); // latest opset
OpTester test("Clip", -1); // latest opset
test.AddInput<float>("X", dims, {});
test.AddInput<float>("min", {}, {-5});
test.AddInput<float>("max", {}, {5});

206
onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc
View file

@ -700,7 +700,7 @@ TEST(MathOpTest, Sum_8_Test1) {
//This test runs fine on CPU and GPU Plugins
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider, kOpenVINOExecutionProvider});
#else
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: Expected output shape [{3,3,3}] did not match run output shape [{3,1,1}] for sum
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: Expected output shape [{3,3,3}] did not match run output shape [{3,1,1}] for sum
#endif
}
@ -783,6 +783,108 @@ TEST(MathOpTest, Min_8) {
test.Run();
}
TEST(MathOpTest, Min_12_Float) {
OpTester test("Min", 12);
test.AddInput<float>("data_0", {1, 3},
{1.0f, 2.0f, 3.0f});
test.AddInput<float>("data_2", {3, 3},
{10.0f, 20.0f, 30.0f,
40.0f, 50.0f, 60.0f,
-70.0f, -80.0f, -90.0f});
test.AddInput<float>("data_1", {3, 1},
{-1.0f, 20.0f, 300.0f});
test.AddOutput<float>("min", {3, 3},
{-1.0f, -1.0f, -1.0f,
1.0f, 2.0f, 3.0f,
-70.0f, -80.0f, -90.0f});
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: Input batch size is inconsistent
}
TEST(MathOpTest, Min_12_Double) {
OpTester test("Min", 12);
test.AddInput<double>("data_0", {1, 3},
{1.0f, 2.0f, 3.0f});
test.AddInput<double>("data_2", {3, 3},
{10.0f, 20.0f, 30.0f,
40.0f, 50.0f, 60.0f,
-70.0f, -80.0f, -90.0f});
test.AddInput<double>("data_1", {3, 1},
{-1.0f, 20.0f, 300.0f});
test.AddOutput<double>("min", {3, 3},
{-1.0f, -1.0f, -1.0f,
1.0f, 2.0f, 3.0f,
-70.0f, -80.0f, -90.0f});
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: Input batch size is inconsistent
}
TEST(MathOpTest, Min_12_Int32) {
OpTester test("Min", 12);
test.AddInput<int32_t>("data_0", {1, 3},
{1, 2, 3});
test.AddInput<int32_t>("data_2", {3, 3},
{10, 20, 30,
40, 50, 60,
-70, -80, -90});
test.AddInput<int32_t>("data_1", {3, 1},
{-1, 20, 300});
test.AddOutput<int32_t>("min", {3, 3},
{-1, -1, -1,
1, 2, 3,
-70, -80, -90});
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: Input batch size is inconsistent
}
TEST(MathOpTest, Min_12_Int64) {
OpTester test("Min", 12);
test.AddInput<int64_t>("data_0", {1, 3},
{1, 2, 3});
test.AddInput<int64_t>("data_2", {3, 3},
{10, 20, 30,
40, 50, 60,
-70, -80, -90});
test.AddInput<int64_t>("data_1", {3, 1},
{-1, 20, 300});
test.AddOutput<int64_t>("min", {3, 3},
{-1, -1, -1,
1, 2, 3,
-70, -80, -90});
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: Input batch size is inconsistent
}
TEST(MathOpTest, Min_12_UInt32) {
OpTester test("Min", 12);
test.AddInput<uint32_t>("data_0", {1, 3},
{1, 20, 30});
test.AddInput<uint32_t>("data_2", {3, 3},
{10, 20, 30,
40, 50, 60,
70, 80, 90});
test.AddInput<uint32_t>("data_1", {3, 1},
{1, 20, 30});
test.AddOutput<uint32_t>("min", {3, 3},
{1, 1, 1,
1, 20, 20,
1, 20, 30});
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: Input batch size is inconsistent
}
TEST(MathOpTest, Min_12_UInt64) {
OpTester test("Min", 12);
test.AddInput<uint64_t>("data_0", {1, 3},
{1, 20, 30});
test.AddInput<uint64_t>("data_2", {3, 3},
{10, 20, 30,
40, 50, 60,
70, 80, 90});
test.AddInput<uint64_t>("data_1", {3, 1},
{1, 20, 30});
test.AddOutput<uint64_t>("min", {3, 3},
{1, 1, 1,
1, 20, 20,
1, 20, 30});
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: Input batch size is inconsistent
}
TEST(MathOpTest, Max_6) {
OpTester test("Max", 6);
std::vector<int64_t> dims{3, 3};
@ -854,6 +956,108 @@ TEST(MathOpTest, Max_8_2inputbroadcast) {
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: Input batch size is inconsistent
}
TEST(MathOpTest, Max_12_Float) {
OpTester test("Max", 12);
test.AddInput<float>("data_0", {1, 3},
{1.0f, 2.0f, 3.0f});
test.AddInput<float>("data_2", {3, 3},
{10.0f, 20.0f, 30.0f,
40.0f, 50.0f, 60.0f,
70.0f, 80.0f, 90.0f});
test.AddInput<float>("data_1", {3, 1},
{-1.0f, -2.0f, 300.0f});
test.AddOutput<float>("max", {3, 3},
{10.0f, 20.0f, 30.0f,
40.0f, 50.0f, 60.0f,
300.0f, 300.0f, 300.0f});
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: Input batch size is inconsistent
}
TEST(MathOpTest, Max_12_Double) {
OpTester test("Max", 12);
test.AddInput<double>("data_0", {1, 3},
{1.0, 2.0, 3.0});
test.AddInput<double>("data_2", {3, 3},
{10.0, 20.0, 30.0,
40.0, 50.0, 60.0,
70.0, 80.0, 90.0});
test.AddInput<double>("data_1", {3, 1},
{-1.0, -2.0, 300.0});
test.AddOutput<double>("max", {3, 3},
{10.0, 20.0, 30.0,
40.0, 50.0, 60.0,
300.0, 300.0, 300.0});
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: Input batch size is inconsistent
}
TEST(MathOpTest, Max_12_Int32) {
OpTester test("Max", 12);
test.AddInput<int32_t>("data_0", {1, 3},
{1, 2, 3});
test.AddInput<int32_t>("data_2", {3, 3},
{10, 20, 30,
40, 50, 60,
70, 80, 90});
test.AddInput<int32_t>("data_1", {3, 1},
{-1, -2, 300});
test.AddOutput<int32_t>("max", {3, 3},
{10, 20, 30,
40, 50, 60,
300, 300, 300});
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: Input batch size is inconsistent
}
TEST(MathOpTest, Max_12_Int64) {
OpTester test("Max", 12);
test.AddInput<int64_t>("data_0", {1, 3},
{1, 2, 3});
test.AddInput<int64_t>("data_2", {3, 3},
{10, 20, 30,
40, 50, 60,
70, 80, 90});
test.AddInput<int64_t>("data_1", {3, 1},
{-1, -2, 300});
test.AddOutput<int64_t>("max", {3, 3},
{10, 20, 30,
40, 50, 60,
300, 300, 300});
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: Input batch size is inconsistent
}
TEST(MathOpTest, Max_12_UInt32) {
OpTester test("Max", 12);
test.AddInput<uint32_t>("data_0", {1, 3},
{1, 2, 3});
test.AddInput<uint32_t>("data_2", {3, 3},
{10, 20, 30,
40, 50, 60,
70, 80, 90});
test.AddInput<uint32_t>("data_1", {3, 1},
{1, 2, 300});
test.AddOutput<uint32_t>("max", {3, 3},
{10, 20, 30,
40, 50, 60,
300, 300, 300});
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: Input batch size is inconsistent
}
TEST(MathOpTest, Max_12_UInt64) {
OpTester test("Max", 12);
test.AddInput<uint64_t>("data_0", {1, 3},
{1, 2, 3});
test.AddInput<uint64_t>("data_2", {3, 3},
{10, 20, 30,
40, 50, 60,
70, 80, 90});
test.AddInput<uint64_t>("data_1", {3, 1},
{1, 2, 300});
test.AddOutput<uint64_t>("max", {3, 3},
{10, 20, 30,
40, 50, 60,
300, 300, 300});
test.Run(OpTester::ExpectResult::kExpectSuccess, "", {kTensorrtExecutionProvider}); //TensorRT: Input batch size is inconsistent
}
TEST(MathOpTest, Not) {
OpTester test("Not");
std::vector<int64_t> dims{2};