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

Add cuda kernels for GreaterOrEqual, LessOrEqual, Where; modify Clip to avoid memcpy (#7187)

Browse source 
* Where and Clip cuda kernel support

* GreaterOrEqual and LessOrEqual cuda kernels

* Clip input GPU mem

* review comments

* Add CPU kernel as well

* review comment

* Add kernel def hash for new op kernels

* Fix CI
This commit is contained in:
ashbhandare 2021-04-06 09:04:38 -07:00 committed by GitHub
parent c85657cfd7
commit e9ffcfa247
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
16 changed files with 376 additions and 33 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

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

@ -468,6 +468,15 @@ class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOn
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, 12, double_double, Dropout);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, Celu);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, float, GreaterOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, double, GreaterOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int32_t, GreaterOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int64_t, GreaterOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, float, LessOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, double, LessOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int32_t, LessOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int64_t, LessOrEqual);
// opset 13
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, float, Erf);
class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, Cast);
@ -1441,6 +1450,22 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, 12, double_float, Dropout)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, 12, double_double, Dropout)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, Celu)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, float,
GreaterOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, double,
GreaterOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int32_t,
GreaterOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int64_t,
GreaterOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, float,
LessOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, double,
LessOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int32_t,
LessOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 12, int64_t,
LessOrEqual)>,
// opset 13
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, Cast)>,

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

@ -89,11 +89,11 @@ struct Clip::ComputeImpl {
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.");
ORT_ENFORCE(min->Shape().IsScalar(), "min should be a scalar.");
min_val = *(min->template Data<T>());
}
if (max) {
ORT_ENFORCE(max->Shape().NumDimensions() == 0, "max should be a scalar.");
ORT_ENFORCE(max->Shape().IsScalar(), "max should be a scalar.");
max_val = *(max->template Data<T>());
}

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

@ -334,6 +334,16 @@ REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(Equal, 13, int64_t, Equal);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(Equal, 13, float, Equal);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(Equal, 13, double, Equal);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(LessOrEqual, 12, float, LessOrEqual);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(LessOrEqual, 12, double, LessOrEqual);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(LessOrEqual, 12, int32_t, LessOrEqual);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(LessOrEqual, 12, int64_t, LessOrEqual);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(GreaterOrEqual, 12, float, GreaterOrEqual);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(GreaterOrEqual, 12, double, GreaterOrEqual);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(GreaterOrEqual, 12, int32_t, GreaterOrEqual);
REG_ELEMENTWISE_LOGICALOP_TYPED_KERNEL(GreaterOrEqual, 12, int64_t, GreaterOrEqual);
REG_ELEMENTWISE_VERSIONED_TYPED_KERNEL(Mean, 6, 7, float, Mean_6);
REG_ELEMENTWISE_VERSIONED_TYPED_KERNEL(Mean, 8, 12, float, Mean_8);
// Supposed to add BFloat16 but we are not supporting now, however, separate registration
@ -950,6 +960,41 @@ Status Greater<T>::Compute(OpKernelContext* context) const {
return Status::OK();
}
template <typename T>
Status LessOrEqual<T>::Compute(OpKernelContext* context) const {
ProcessBroadcastSpanFuncs funcs{
[](BroadcastHelper& per_iter_bh) {
per_iter_bh.OutputEigen<bool>() = per_iter_bh.EigenInput1<T>().array() >= per_iter_bh.ScalarInput0<T>();
},
[](BroadcastHelper& per_iter_bh) {
per_iter_bh.OutputEigen<bool>() = per_iter_bh.EigenInput0<T>().array() <= per_iter_bh.ScalarInput1<T>();
},
[](BroadcastHelper& per_iter_bh) {
per_iter_bh.OutputEigen<bool>() = per_iter_bh.EigenInput0<T>().array() <= per_iter_bh.EigenInput1<T>().array();
}};
UntypedBroadcastTwo(*context, funcs, 1.0);
return Status::OK();
}
template <typename T>
Status GreaterOrEqual<T>::Compute(OpKernelContext* context) const {
ProcessBroadcastSpanFuncs funcs{
[](BroadcastHelper& per_iter_bh) {
per_iter_bh.OutputEigen<bool>() = per_iter_bh.EigenInput1<T>().array() <= per_iter_bh.ScalarInput0<T>();
},
[](BroadcastHelper& per_iter_bh) {
per_iter_bh.OutputEigen<bool>() = per_iter_bh.EigenInput0<T>().array() >= per_iter_bh.ScalarInput1<T>();
},
[](BroadcastHelper& per_iter_bh) {
per_iter_bh.OutputEigen<bool>() =
per_iter_bh.EigenInput0<T>().array() >= per_iter_bh.EigenInput1<T>().array();
}};
UntypedBroadcastTwo(*context, funcs, 1.0);
return Status::OK();
}
template <>
Status Mean_6<float>::Compute(OpKernelContext* ctx) const {
auto inputCount = Node().InputArgCount().front();

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

@ -374,6 +374,24 @@ class Greater final : public OpKernel {
Status Compute(OpKernelContext* context) const override;
};
template <typename T>
class LessOrEqual final : public OpKernel {
public:
LessOrEqual(const OpKernelInfo& info) : OpKernel(info) {
}
Status Compute(OpKernelContext* context) const override;
};
template <typename T>
class GreaterOrEqual final : public OpKernel {
public:
GreaterOrEqual(const OpKernelInfo& info) : OpKernel(info) {
}
Status Compute(OpKernelContext* context) const override;
};
template <typename T>
class Mean_6 final : public OpKernel {
public:

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

@ -412,6 +412,20 @@ class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kO
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, 12, float, Greater);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, 12, double, Greater);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, 12, MLFloat16, Greater);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int32_t, GreaterOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int64_t, GreaterOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint32_t, GreaterOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint64_t, GreaterOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, GreaterOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, GreaterOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, GreaterOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int32_t, LessOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int64_t, LessOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint32_t, LessOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint64_t, LessOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, LessOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, LessOrEqual);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, LessOrEqual);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 12, int32_t, Add);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 12, int64_t, Add);
class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 12, uint32_t, Add);
@ -644,6 +658,7 @@ class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, Ey
class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, 10, Scatter);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, MLFloat16, Where);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, float, Where);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, double_t, Where);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, int32_t, Where);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, int64_t, Where);
class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, uint8_t, Where);
@ -1159,6 +1174,20 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, 12, float, Greater)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, 12, double, Greater)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, 12, MLFloat16, Greater)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int32_t, GreaterOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int64_t, GreaterOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint32_t, GreaterOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint64_t, GreaterOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, GreaterOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, GreaterOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, GreaterOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int32_t, LessOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, int64_t, LessOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint32_t, LessOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, uint64_t, LessOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, float, LessOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, double, LessOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 12, MLFloat16, LessOrEqual)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 12, int32_t, Add)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 12, int64_t, Add)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 7, 12, uint32_t, Add)>,
@ -1390,6 +1419,7 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, 10, Scatter)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, MLFloat16, Where)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, float, Where)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, double_t, Where)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, int32_t, Where)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, int64_t, Where)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 9, uint8_t, Where)>,

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

@ -477,6 +477,24 @@ Status Less<T>::ComputeInternal(OpKernelContext* context) const {
return Status::OK();
}
//GreaterOrEqual op output tensor type is bool, so it cannot directly fit in the macros
//for other elementwise ops
template <typename T>
Status GreaterOrEqual<T>::ComputeInternal(OpKernelContext* context) const {
this->CompareMethod(context, &ImplT2_GreaterOrEqual);
return Status::OK();
}
//LessOrEqual op output tensor type is bool, so it cannot directly fit in the macros
//for other elementwise ops
template <typename T>
Status LessOrEqual<T>::ComputeInternal(OpKernelContext* context) const {
this->CompareMethod(context, &ImplT2_LessOrEqual);
return Status::OK();
}
BINARY_LOGICALOP_REGISTER_UZILHFD(Equal, 13)
BINARY_ELEMENTWISE_LOGICALOP_REGISTER_KERNEL_TYPED(Equal, 13, bool)
BINARY_OP_REGISTER_VERSIONED_UZILHFD(Equal, 11, 12)
@ -488,6 +506,9 @@ BINARY_OP_REGISTER_VERSIONED_HFD(Greater, 7, 8)
BINARY_LOGICALOP_REGISTER_UZILHFD(Less, 13)
BINARY_OP_REGISTER_VERSIONED_UZILHFD(Less, 9, 12)
BINARY_OP_REGISTER_VERSIONED_HFD(Less, 7, 8)
BINARY_LOGICALOP_REGISTER_UZILHFD(GreaterOrEqual, 12)
BINARY_LOGICALOP_REGISTER_UZILHFD(LessOrEqual, 12)
} // namespace cuda
} // namespace onnxruntime

17
onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h
View file

@ -257,5 +257,22 @@ class Less final : public CompareFunction<T, typename ToCudaType<T>::MappedType>
Status ComputeInternal(OpKernelContext* context) const override;
};
template <typename T>
class GreaterOrEqual final : public CompareFunction<T, typename ToCudaType<T>::MappedType> {
public:
GreaterOrEqual(const OpKernelInfo& info) : CompareFunction<T, typename ToCudaType<T>::MappedType>(info) {}
Status ComputeInternal(OpKernelContext* context) const override;
};
template <typename T>
class LessOrEqual final : public CompareFunction<T, typename ToCudaType<T>::MappedType> {
public:
LessOrEqual(const OpKernelInfo& info) : CompareFunction<T, typename ToCudaType<T>::MappedType>(info) {}
Status ComputeInternal(OpKernelContext* context) const override;
};
} // namespace cuda
} // namespace onnxruntime

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

@ -180,6 +180,8 @@ SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD2(Greater)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD2(Equal)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_T2(Equal, bool, bool, bool)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD2(Less)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD2(GreaterOrEqual)
SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD2(LessOrEqual)
} // namespace cuda
} // namespace onnxruntime

5
onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.h
View file

@ -85,7 +85,10 @@ BINARY_ELEMENTWISE_IMPL_DECLARATION_T1(Pow);
#define BINARY_OPS2() \
BINARY_OP_NAME_EXPR2(Greater, (a > b)) \
BINARY_OP_NAME_EXPR2(Equal, (a == b)) \
BINARY_OP_NAME_EXPR2(Less, (a < b))
BINARY_OP_NAME_EXPR2(Less, (a < b)) \
BINARY_OP_NAME_EXPR2(GreaterOrEqual, (a >= b)) \
BINARY_OP_NAME_EXPR2(LessOrEqual, (a <= b))
#define BINARY_OP_NAME_EXPR2(name, expr) BINARY_ELEMENTWISE_IMPL_DECLARATION_T2(name);
BINARY_OPS2()

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

@ -25,8 +25,6 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX(
11, 11,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(1)
.InputMemoryType<OrtMemTypeCPUInput>(2)
.TypeConstraint("T", DataTypeImpl::GetTensorType<float>()),
Clip);
@ -36,9 +34,7 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX(
12, 12,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(1)
.InputMemoryType<OrtMemTypeCPUInput>(2)
.TypeConstraint("T", BuildKernelDefConstraints<float, double, int8_t, uint8_t, int64_t, uint64_t>()),
.TypeConstraint("T", BuildKernelDefConstraints<float, double, MLFloat16, int8_t, uint8_t, int64_t, uint64_t>()),
Clip);
ONNX_OPERATOR_KERNEL_EX(
@ -47,9 +43,7 @@ ONNX_OPERATOR_KERNEL_EX(
13,
kCudaExecutionProvider,
KernelDefBuilder()
.InputMemoryType<OrtMemTypeCPUInput>(1)
.InputMemoryType<OrtMemTypeCPUInput>(2)
.TypeConstraint("T", BuildKernelDefConstraints<float, double, int8_t, uint8_t, int64_t, uint64_t>()),
.TypeConstraint("T", BuildKernelDefConstraints<float, double, MLFloat16, int8_t, uint8_t, int64_t, uint64_t>()),
Clip);
@ -62,7 +56,7 @@ Status Clip_6<T>::ComputeInternal(OpKernelContext* ctx) const {
if (count > 0) {
auto* y_data = Y->template MutableData<T>();
const auto* x_data = X.template Data<T>();
ClipImpl<T>(Stream(), x_data, y_data, this->min_, this->max_, count);
ClipImpl<T>(Stream(), x_data, y_data, nullptr, nullptr, this->min_, this->max_, count);
}
return Status::OK();
}
@ -92,25 +86,27 @@ struct LowMax<MLFloat16> {
template <typename T>
struct Clip::ComputeImpl {
void operator()(cudaStream_t stream, 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();
auto min_default = clip_internal::LowMax<T>::low();
auto max_default = clip_internal::LowMax<T>::max();
const T* min_data = nullptr;
const T* max_data = nullptr;
// 1-2 Input on CPU
if (min) {
ORT_ENFORCE(min->Shape().NumDimensions() == 0, "min should be a scalar.");
min_val = *(min->template Data<T>());
ORT_ENFORCE(min->Shape().IsScalar(), "min should be a scalar.");
min_data = min->template Data<T>();
}
if (max) {
ORT_ENFORCE(max->Shape().NumDimensions() == 0, "max should be a scalar.");
max_val = *(max->template Data<T>());
ORT_ENFORCE(max->Shape().IsScalar(), "max should be a scalar.");
max_data = 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>(stream, x_data, y_data, min_val, max_val, count);
ClipImpl<T>(stream, x_data, y_data, min_data, max_data, min_default, max_default, count);
}
}
};
@ -121,7 +117,7 @@ Status Clip::ComputeInternal(OpKernelContext* ctx) const {
const auto* max = ctx->Input<Tensor>(2);
Tensor* Y = ctx->Output(0, X->Shape());
utils::MLTypeCallDispatcher<float, double, int8_t, uint8_t, int64_t, uint64_t>
utils::MLTypeCallDispatcher<float, double, MLFloat16, int8_t, uint8_t, int64_t, uint64_t>
t_disp(X->GetElementType());
t_disp.Invoke<ComputeImpl>(Stream(), X, min, max, Y);

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

@ -7,30 +7,34 @@
namespace onnxruntime {
namespace cuda {
template <typename T>
__global__ void _Clip(const T* input, T* output, T min, T max, size_t N) {
__global__ void _Clip(const T* input, T* output, const T* min, const T* max, T min_default, T max_default, size_t N) {
auto min_val = (min) ? *min : min_default;
auto max_val = (max) ? *max : max_default;
CALCULATE_ELEMENTWISE_INDEX_OR_EXIT(id, N);
output[id] = (input[id] < min) ? min : ((input[id] > max) ? max : input[id]);
output[id] = (input[id] < min_val) ? min_val : ((input[id] > max_val) ? max_val : input[id]);
}
template <typename T>
void ClipImpl(cudaStream_t stream, const T* input_data, T* output_data, T min, T max, size_t count) {
void ClipImpl(cudaStream_t stream, const T* input_data, T* output_data, const T* min, const T* max, T min_default, T max_default, size_t count) {
typedef typename ToCudaType<T>::MappedType CudaT;
int blocksPerGrid = (int)(ceil(static_cast<float>(count) / GridDim::maxThreadsPerBlock));
_Clip<CudaT><<<blocksPerGrid, GridDim::maxThreadsPerBlock, 0, stream>>>(reinterpret_cast<const CudaT*>(input_data),
reinterpret_cast<CudaT*>(output_data),
*reinterpret_cast<CudaT*>(&min),
*reinterpret_cast<CudaT*>(&max),
reinterpret_cast<const CudaT*>(min),
reinterpret_cast<const CudaT*>(max),
*reinterpret_cast<CudaT*>(&min_default),
*reinterpret_cast<CudaT*>(&max_default),
count);
}
template void ClipImpl<float>(cudaStream_t stream, const float* input_data, float* output_data, float min, float max, size_t count);
template void ClipImpl<double>(cudaStream_t stream, const double* input_data, double* output_data, double min, double max, size_t count);
template void ClipImpl<MLFloat16>(cudaStream_t stream, const MLFloat16* input_data, MLFloat16* output_data, MLFloat16 min, MLFloat16 max, size_t count);
template void ClipImpl<int8_t>(cudaStream_t stream, const int8_t* input_data, int8_t* output_data, int8_t min, int8_t max, size_t count);
template void ClipImpl<uint8_t>(cudaStream_t stream, const uint8_t* input_data, uint8_t* output_data, uint8_t min, uint8_t max, size_t count);
template void ClipImpl<int64_t>(cudaStream_t stream, const int64_t* input_data, int64_t* output_data, int64_t min, int64_t max, size_t count);
template void ClipImpl<uint64_t>(cudaStream_t stream, const uint64_t* input_data, uint64_t* output_data, uint64_t min, uint64_t max, size_t count);
template void ClipImpl<float>(cudaStream_t stream, const float* input_data, float* output_data, const float* min, const float* max, float min_default, float max_default, size_t count);
template void ClipImpl<double>(cudaStream_t stream, const double* input_data, double* output_data, const double* min, const double* max, double min_default, double max_default, size_t count);
template void ClipImpl<MLFloat16>(cudaStream_t stream, const MLFloat16* input_data, MLFloat16* output_data, const MLFloat16* min, const MLFloat16* max, MLFloat16 min_default, MLFloat16 max_default, size_t count);
template void ClipImpl<int8_t>(cudaStream_t stream, const int8_t* input_data, int8_t* output_data, const int8_t* min, const int8_t* max, int8_t min_default, int8_t max_default, size_t count);
template void ClipImpl<uint8_t>(cudaStream_t stream, const uint8_t* input_data, uint8_t* output_data, const uint8_t* min, const uint8_t* max, uint8_t min_default, uint8_t max_default, size_t count);
template void ClipImpl<int64_t>(cudaStream_t stream, const int64_t* input_data, int64_t* output_data, const int64_t* min, const int64_t* max, int64_t min_default, int64_t max_default, size_t count);
template void ClipImpl<uint64_t>(cudaStream_t stream, const uint64_t* input_data, uint64_t* output_data, const uint64_t* min, const uint64_t* max, uint64_t min_default, uint64_t max_default, size_t count);
} // namespace cuda
} // namespace onnxruntime

2
onnxruntime/core/providers/cuda/math/clip_impl.h
View file

@ -10,7 +10,7 @@
namespace onnxruntime {
namespace cuda {
template <typename T>
void ClipImpl(cudaStream_t stream, const T* input_data, T* output_data, T min, T max, size_t count);
void ClipImpl(cudaStream_t stream, const T* input_data, T* output_data, const T* min, const T* max, T min_default, T max_default, size_t count);
} // namespace cuda
} // namespace onnxruntime

1
onnxruntime/core/providers/cuda/tensor/where.cc
View file

@ -203,6 +203,7 @@ SPECIALIZED_COMPUTE(uint8_t)
SPECIALIZED_COMPUTE(int32_t)
SPECIALIZED_COMPUTE(int64_t)
SPECIALIZED_COMPUTE(float)
SPECIALIZED_COMPUTE(double_t)
SPECIALIZED_COMPUTE(MLFloat16)
} // namespace cuda
} // namespace onnxruntime

1
onnxruntime/core/providers/cuda/tensor/where_impl.cu
View file

@ -232,6 +232,7 @@ SPECIALIZED_IMPL(uint8_t)
SPECIALIZED_IMPL(int32_t)
SPECIALIZED_IMPL(int64_t)
SPECIALIZED_IMPL(float)
SPECIALIZED_IMPL(double_t)
SPECIALIZED_IMPL(half)
} // namespace cuda

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

@ -1628,6 +1628,78 @@ TEST(MathOpTest, Less_multidiretional_broadcastBA) {
test.Run();
}
TEST(MathOpTest, LessOrEqual) {
OpTester test("LessOrEqual", 12);
std::vector<int64_t> dims{4};
test.AddInput<float>("A", dims, {1.0f, 0.0f, -1.0f, -1.0f});
test.AddInput<float>("B", dims, {1.0f, 1.0f, 2.0f, -1.0f});
test.AddOutput<bool>("C", dims, {true, true, true, true});
test.Run(OpTester::ExpectResult::kExpectSuccess, "",
{kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider});
}
TEST(MathOpTest, LessOrEqual_Scalar0) {
OpTester test("LessOrEqual", 12);
test.AddInput<float>("A", {1}, {1.0f});
test.AddInput<float>("B", {4}, {1.0f, 1.5f, 2.0f, -1.0f});
test.AddOutput<bool>("C", {4}, {true, true, true, false});
test.Run(OpTester::ExpectResult::kExpectSuccess, "",
{kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider});
}
TEST(MathOpTest, LessOrEqual_Scalar1) {
OpTester test("LessOrEqual", 12);
test.AddInput<float>("A", {4}, {1.0f, 0.5f, 2.0f, -1.0f});
test.AddInput<float>("B", {1}, {1.0f});
test.AddOutput<bool>("C", {4}, {true, true, false, true});
test.Run(OpTester::ExpectResult::kExpectSuccess, "",
{kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider});
}
TEST(MathOpTest, LessOrEqual_int64_Scalar1) {
OpTester test("LessOrEqual", 12);
test.AddInput<int64_t>("A", {4}, {1, 0, 2, -1});
test.AddInput<int64_t>("B", {1}, {1});
test.AddOutput<bool>("C", {4}, {true, true, false, true});
test.Run(OpTester::ExpectResult::kExpectSuccess, "",
{kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider});
}
TEST(MathOpTest, LessOrEqual_broadcastAB) {
OpTester test("LessOrEqual", 12);
test.AddInput<int32_t>("A", {4, 2}, {10, 11, 12, 13, 14, 15, 16, 17});
test.AddInput<int32_t>("B", {2}, {15, 7});
test.AddOutput<bool>("C", {4, 2}, {true, false, true, false, true, false, false, false});
test.Run(OpTester::ExpectResult::kExpectSuccess, "",
{kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider});
}
TEST(MathOpTest, LessOrEqual_broadcastBA) {
OpTester test("LessOrEqual", 12);
test.AddInput<int32_t>("A", {2}, {15, 7});
test.AddInput<int32_t>("B", {4, 2}, {10, 11, 12, 13, 14, 15, 16, 17});
test.AddOutput<bool>("C", {4, 2}, {false, true, false, true, false, true, true, true});
test.Run(OpTester::ExpectResult::kExpectSuccess, "",
{kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider});
}
TEST(MathOpTest, LessOrEqual_multidiretional_broadcastAB) {
OpTester test("LessOrEqual", 12);
test.AddInput<int32_t>("A", {4, 1}, {10, 11, 12, 13});
test.AddInput<int32_t>("B", {2}, {15, 7});
test.AddOutput<bool>("C", {4, 2}, {true, false, true, false, true, false, true, false});
test.Run(OpTester::ExpectResult::kExpectSuccess, "",
{kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider});
}
TEST(MathOpTest, LessOrEqual_multidiretional_broadcastBA) {
OpTester test("LessOrEqual", 12);
test.AddInput<int32_t>("A", {2}, {15, 7});
test.AddInput<int32_t>("B", {4, 1}, {10, 11, 12, 13});
test.AddOutput<bool>("C", {4, 2}, {false, true, false, true, false, true, false, true});
test.Run(OpTester::ExpectResult::kExpectSuccess, "",
{kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider});
}
TEST(MathOpTest, Greater_7) {
OpTester test("Greater");
std::vector<int64_t> dims{4};
@ -1705,6 +1777,82 @@ TEST(MathOpTest, Greater_multidiretional_broadcastBA) {
test.Run();
}
TEST(MathOpTest, GreaterOrEqual_12_float) {
OpTester test("GreaterOrEqual", 12);
std::vector<int64_t> dims{4};
test.AddInput<float>("A", dims, {1.0f, 0.0f, -1.0f, -1.0f});
test.AddInput<float>("B", dims, {1.0f, 1.0f, 2.0f, -1.0f});
test.AddOutput<bool>("C", dims, {true, false, false, true});
test.Run(OpTester::ExpectResult::kExpectSuccess, "",
{kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider});
}
TEST(MathOpTest, GreaterOrEqual_12_double) {
OpTester test("GreaterOrEqual", 12);
std::vector<int64_t> dims{4};
test.AddInput<double>("A", dims, {1.0, 0.0, 3.0, -1.0});
test.AddInput<double>("B", dims, {1.0, 1.0, 2.0, -1.0});
test.AddOutput<bool>("C", dims, {true, false, true, true});
test.Run(OpTester::ExpectResult::kExpectSuccess, "",
{kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider});
}
TEST(MathOpTest, GreaterOrEqual_12_int32) {
OpTester test("GreaterOrEqual", 12);
std::vector<int64_t> dims{4};
test.AddInput<int32_t>("A", dims, {10, 11, 12, 13});
test.AddInput<int32_t>("B", dims, {15, 7, 12, 9});
test.AddOutput<bool>("C", dims, {false, true, true, true});
test.Run(OpTester::ExpectResult::kExpectSuccess, "",
{kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider});
}
TEST(MathOpTest, GreaterOrEqual_12_int64) {
OpTester test("GreaterOrEqual", 12);
std::vector<int64_t> dims{4};
test.AddInput<int64_t>("A", dims, {10, 11, 12, 13});
test.AddInput<int64_t>("B", dims, {15, 7, 12, 9});
test.AddOutput<bool>("C", dims, {false, true, true, true});
test.Run(OpTester::ExpectResult::kExpectSuccess, "",
{kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider});
}
TEST(MathOpTest, GreaterOrEqual_broadcastAB) {
OpTester test("GreaterOrEqual", 12);
test.AddInput<int32_t>("A", {4, 2}, {10, 11, 12, 13, 14, 15, 16, 17});
test.AddInput<int32_t>("B", {2}, {15, 7});
test.AddOutput<bool>("C", {4, 2}, {false, true, false, true, false, true, true, true});
test.Run(OpTester::ExpectResult::kExpectSuccess, "",
{kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider});
}
TEST(MathOpTest, GreaterOrEqual_broadcastBA) {
OpTester test("GreaterOrEqual", 12);
test.AddInput<int32_t>("A", {2}, {15, 7});
test.AddInput<int32_t>("B", {4, 2}, {10, 11, 12, 13, 14, 15, 16, 17});
test.AddOutput<bool>("C", {4, 2}, {true, false, true, false, true, false, false, false});
test.Run(OpTester::ExpectResult::kExpectSuccess, "",
{kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider});
}
TEST(MathOpTest, GreaterOrEqual_multidiretional_broadcastAB) {
OpTester test("GreaterOrEqual", 12);
test.AddInput<int32_t>("A", {4, 1}, {10, 11, 12, 13});
test.AddInput<int32_t>("B", {2}, {15, 7});
test.AddOutput<bool>("C", {4, 2}, {false, true, false, true, false, true, false, true});
test.Run(OpTester::ExpectResult::kExpectSuccess, "",
{kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider});
}
TEST(MathOpTest, GreaterOrEqual_multidiretional_broadcastBA) {
OpTester test("GreaterOrEqual", 12);
test.AddInput<int32_t>("A", {2}, {15, 7});
test.AddInput<int32_t>("B", {4, 1}, {10, 11, 12, 13});
test.AddOutput<bool>("C", {4, 2}, {true, false, true, false, true, false, true, false});
test.Run(OpTester::ExpectResult::kExpectSuccess, "",
{kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider});
}
TEST(MathOpTest, Equal_bool) {
OpTester test("Equal");
std::vector<int64_t> dims{4};

32
onnxruntime/test/testdata/kernel_def_hashes/onnx.cpu.json vendored
View file

@ -823,6 +823,22 @@
"Greater ai.onnx CPUExecutionProvider",
16852011221046024392
],
[
"GreaterOrEqual ai.onnx CPUExecutionProvider",
3999586969438630368
],
[
"GreaterOrEqual ai.onnx CPUExecutionProvider",
8317279776362716048
],
[
"GreaterOrEqual ai.onnx CPUExecutionProvider",
14896183015337647264
],
[
"GreaterOrEqual ai.onnx CPUExecutionProvider",
17416867432093505280
],
[
"GRU ai.onnx CPUExecutionProvider",
2706165712066264784
@ -931,6 +947,22 @@
"Less ai.onnx CPUExecutionProvider",
17960128831236491008
],
[
"LessOrEqual ai.onnx CPUExecutionProvider",
1261667279452953168
],
[
"LessOrEqual ai.onnx CPUExecutionProvider",
2051143717905239376
],
[
"LessOrEqual ai.onnx CPUExecutionProvider",
4697898477799165704
],
[
"LessOrEqual ai.onnx CPUExecutionProvider",
8848289292300988248
],
[
"Log ai.onnx CPUExecutionProvider",
268464912229648680