diff --git a/onnxruntime/core/providers/cpu/cpu_execution_provider.cc b/onnxruntime/core/providers/cpu/cpu_execution_provider.cc index f2c8653bdf..78778dc3ec 100644 --- a/onnxruntime/core/providers/cpu/cpu_execution_provider.cc +++ b/onnxruntime/core/providers/cpu/cpu_execution_provider.cc @@ -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, BuildKernelCreateInfo, BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, // opset 13 BuildKernelCreateInfo, diff --git a/onnxruntime/core/providers/cpu/math/clip.cc b/onnxruntime/core/providers/cpu/math/clip.cc index efd707eb8e..e4fbdae49a 100644 --- a/onnxruntime/core/providers/cpu/math/clip.cc +++ b/onnxruntime/core/providers/cpu/math/clip.cc @@ -89,11 +89,11 @@ struct Clip::ComputeImpl { auto min_val = std::numeric_limits::lowest(); auto max_val = std::numeric_limits::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()); } 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()); } diff --git a/onnxruntime/core/providers/cpu/math/element_wise_ops.cc b/onnxruntime/core/providers/cpu/math/element_wise_ops.cc index 30d6c7c011..4f3f6de569 100644 --- a/onnxruntime/core/providers/cpu/math/element_wise_ops.cc +++ b/onnxruntime/core/providers/cpu/math/element_wise_ops.cc @@ -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::Compute(OpKernelContext* context) const { return Status::OK(); } +template +Status LessOrEqual::Compute(OpKernelContext* context) const { + ProcessBroadcastSpanFuncs funcs{ + [](BroadcastHelper& per_iter_bh) { + per_iter_bh.OutputEigen() = per_iter_bh.EigenInput1().array() >= per_iter_bh.ScalarInput0(); + }, + [](BroadcastHelper& per_iter_bh) { + per_iter_bh.OutputEigen() = per_iter_bh.EigenInput0().array() <= per_iter_bh.ScalarInput1(); + }, + [](BroadcastHelper& per_iter_bh) { + per_iter_bh.OutputEigen() = per_iter_bh.EigenInput0().array() <= per_iter_bh.EigenInput1().array(); + }}; + + UntypedBroadcastTwo(*context, funcs, 1.0); + return Status::OK(); +} + +template +Status GreaterOrEqual::Compute(OpKernelContext* context) const { + ProcessBroadcastSpanFuncs funcs{ + [](BroadcastHelper& per_iter_bh) { + per_iter_bh.OutputEigen() = per_iter_bh.EigenInput1().array() <= per_iter_bh.ScalarInput0(); + }, + [](BroadcastHelper& per_iter_bh) { + per_iter_bh.OutputEigen() = per_iter_bh.EigenInput0().array() >= per_iter_bh.ScalarInput1(); + }, + [](BroadcastHelper& per_iter_bh) { + per_iter_bh.OutputEigen() = + per_iter_bh.EigenInput0().array() >= per_iter_bh.EigenInput1().array(); + }}; + + UntypedBroadcastTwo(*context, funcs, 1.0); + return Status::OK(); +} + template <> Status Mean_6::Compute(OpKernelContext* ctx) const { auto inputCount = Node().InputArgCount().front(); diff --git a/onnxruntime/core/providers/cpu/math/element_wise_ops.h b/onnxruntime/core/providers/cpu/math/element_wise_ops.h index e916c95e0d..ad34cc04ca 100644 --- a/onnxruntime/core/providers/cpu/math/element_wise_ops.h +++ b/onnxruntime/core/providers/cpu/math/element_wise_ops.h @@ -374,6 +374,24 @@ class Greater final : public OpKernel { Status Compute(OpKernelContext* context) const override; }; +template +class LessOrEqual final : public OpKernel { + public: + LessOrEqual(const OpKernelInfo& info) : OpKernel(info) { + } + + Status Compute(OpKernelContext* context) const override; +}; + +template +class GreaterOrEqual final : public OpKernel { + public: + GreaterOrEqual(const OpKernelInfo& info) : OpKernel(info) { + } + + Status Compute(OpKernelContext* context) const override; +}; + template class Mean_6 final : public OpKernel { public: diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc index 5e6b0a3be0..ef465a5362 100644 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc @@ -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, BuildKernelCreateInfo, BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -1390,6 +1419,7 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc index 52bf4c0a41..bfaf8af10c 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc @@ -477,6 +477,24 @@ Status Less::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 +Status GreaterOrEqual::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 +Status LessOrEqual::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 diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h index f4c1675aaf..5cd728189d 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h @@ -257,5 +257,22 @@ class Less final : public CompareFunction::MappedType> Status ComputeInternal(OpKernelContext* context) const override; }; + +template +class GreaterOrEqual final : public CompareFunction::MappedType> { + public: + GreaterOrEqual(const OpKernelInfo& info) : CompareFunction::MappedType>(info) {} + + Status ComputeInternal(OpKernelContext* context) const override; +}; + +template +class LessOrEqual final : public CompareFunction::MappedType> { + public: + LessOrEqual(const OpKernelInfo& info) : CompareFunction::MappedType>(info) {} + + Status ComputeInternal(OpKernelContext* context) const override; +}; + } // namespace cuda } // namespace onnxruntime diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu index 8dc09b7fbc..39ae0acfd8 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu @@ -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 diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.h b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.h index c9a8c0f1d3..4fd616dc7c 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.h +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.h @@ -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() diff --git a/onnxruntime/core/providers/cuda/math/clip.cc b/onnxruntime/core/providers/cuda/math/clip.cc index 6c17ecc523..02d986c393 100644 --- a/onnxruntime/core/providers/cuda/math/clip.cc +++ b/onnxruntime/core/providers/cuda/math/clip.cc @@ -25,8 +25,6 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX( 11, 11, kCudaExecutionProvider, KernelDefBuilder() - .InputMemoryType(1) - .InputMemoryType(2) .TypeConstraint("T", DataTypeImpl::GetTensorType()), Clip); @@ -36,9 +34,7 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX( 12, 12, kCudaExecutionProvider, KernelDefBuilder() - .InputMemoryType(1) - .InputMemoryType(2) - .TypeConstraint("T", BuildKernelDefConstraints()), + .TypeConstraint("T", BuildKernelDefConstraints()), Clip); ONNX_OPERATOR_KERNEL_EX( @@ -47,9 +43,7 @@ ONNX_OPERATOR_KERNEL_EX( 13, kCudaExecutionProvider, KernelDefBuilder() - .InputMemoryType(1) - .InputMemoryType(2) - .TypeConstraint("T", BuildKernelDefConstraints()), + .TypeConstraint("T", BuildKernelDefConstraints()), Clip); @@ -62,7 +56,7 @@ Status Clip_6::ComputeInternal(OpKernelContext* ctx) const { if (count > 0) { auto* y_data = Y->template MutableData(); const auto* x_data = X.template Data(); - ClipImpl(Stream(), x_data, y_data, this->min_, this->max_, count); + ClipImpl(Stream(), x_data, y_data, nullptr, nullptr, this->min_, this->max_, count); } return Status::OK(); } @@ -92,25 +86,27 @@ struct LowMax { template 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::low(); - auto max_val = clip_internal::LowMax::max(); + auto min_default = clip_internal::LowMax::low(); + auto max_default = clip_internal::LowMax::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()); + ORT_ENFORCE(min->Shape().IsScalar(), "min should be a scalar."); + min_data = min->template Data(); } if (max) { - ORT_ENFORCE(max->Shape().NumDimensions() == 0, "max should be a scalar."); - max_val = *(max->template Data()); + ORT_ENFORCE(max->Shape().IsScalar(), "max should be a scalar."); + max_data = max->template Data(); } const size_t count = X->Shape().Size(); if (count > 0) { auto* y_data = Y->template MutableData(); const auto* x_data = X->template Data(); - ClipImpl(stream, x_data, y_data, min_val, max_val, count); + ClipImpl(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(2); Tensor* Y = ctx->Output(0, X->Shape()); - utils::MLTypeCallDispatcher + utils::MLTypeCallDispatcher t_disp(X->GetElementType()); t_disp.Invoke(Stream(), X, min, max, Y); diff --git a/onnxruntime/core/providers/cuda/math/clip_impl.cu b/onnxruntime/core/providers/cuda/math/clip_impl.cu index 4164038972..901b002218 100644 --- a/onnxruntime/core/providers/cuda/math/clip_impl.cu +++ b/onnxruntime/core/providers/cuda/math/clip_impl.cu @@ -7,30 +7,34 @@ namespace onnxruntime { namespace cuda { template -__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 -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::MappedType CudaT; int blocksPerGrid = (int)(ceil(static_cast(count) / GridDim::maxThreadsPerBlock)); _Clip<<>>(reinterpret_cast(input_data), reinterpret_cast(output_data), - *reinterpret_cast(&min), - *reinterpret_cast(&max), + reinterpret_cast(min), + reinterpret_cast(max), + *reinterpret_cast(&min_default), + *reinterpret_cast(&max_default), count); } -template void ClipImpl(cudaStream_t stream, const float* input_data, float* output_data, float min, float max, size_t count); -template void ClipImpl(cudaStream_t stream, const double* input_data, double* output_data, double min, double max, size_t count); -template void ClipImpl(cudaStream_t stream, const MLFloat16* input_data, MLFloat16* output_data, MLFloat16 min, MLFloat16 max, size_t count); -template void ClipImpl(cudaStream_t stream, const int8_t* input_data, int8_t* output_data, int8_t min, int8_t max, size_t count); -template void ClipImpl(cudaStream_t stream, const uint8_t* input_data, uint8_t* output_data, uint8_t min, uint8_t max, size_t count); -template void ClipImpl(cudaStream_t stream, const int64_t* input_data, int64_t* output_data, int64_t min, int64_t max, size_t count); -template void ClipImpl(cudaStream_t stream, const uint64_t* input_data, uint64_t* output_data, uint64_t min, uint64_t max, size_t count); +template void ClipImpl(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(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(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(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(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(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(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 diff --git a/onnxruntime/core/providers/cuda/math/clip_impl.h b/onnxruntime/core/providers/cuda/math/clip_impl.h index 6aefa7b90b..9f537110aa 100644 --- a/onnxruntime/core/providers/cuda/math/clip_impl.h +++ b/onnxruntime/core/providers/cuda/math/clip_impl.h @@ -10,7 +10,7 @@ namespace onnxruntime { namespace cuda { template -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 diff --git a/onnxruntime/core/providers/cuda/tensor/where.cc b/onnxruntime/core/providers/cuda/tensor/where.cc index ba85c2cd4c..5dd3d6a098 100644 --- a/onnxruntime/core/providers/cuda/tensor/where.cc +++ b/onnxruntime/core/providers/cuda/tensor/where.cc @@ -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 diff --git a/onnxruntime/core/providers/cuda/tensor/where_impl.cu b/onnxruntime/core/providers/cuda/tensor/where_impl.cu index 0f2d4a4543..91d143ddc2 100644 --- a/onnxruntime/core/providers/cuda/tensor/where_impl.cu +++ b/onnxruntime/core/providers/cuda/tensor/where_impl.cu @@ -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 diff --git a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc index e923e16854..deef42e50d 100644 --- a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc +++ b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc @@ -1628,6 +1628,78 @@ TEST(MathOpTest, Less_multidiretional_broadcastBA) { test.Run(); } +TEST(MathOpTest, LessOrEqual) { + OpTester test("LessOrEqual", 12); + std::vector dims{4}; + test.AddInput("A", dims, {1.0f, 0.0f, -1.0f, -1.0f}); + test.AddInput("B", dims, {1.0f, 1.0f, 2.0f, -1.0f}); + test.AddOutput("C", dims, {true, true, true, true}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", + {kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider}); +} + +TEST(MathOpTest, LessOrEqual_Scalar0) { + OpTester test("LessOrEqual", 12); + test.AddInput("A", {1}, {1.0f}); + test.AddInput("B", {4}, {1.0f, 1.5f, 2.0f, -1.0f}); + test.AddOutput("C", {4}, {true, true, true, false}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", + {kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider}); +} + +TEST(MathOpTest, LessOrEqual_Scalar1) { + OpTester test("LessOrEqual", 12); + test.AddInput("A", {4}, {1.0f, 0.5f, 2.0f, -1.0f}); + test.AddInput("B", {1}, {1.0f}); + test.AddOutput("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("A", {4}, {1, 0, 2, -1}); + test.AddInput("B", {1}, {1}); + test.AddOutput("C", {4}, {true, true, false, true}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", + {kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider}); +} +TEST(MathOpTest, LessOrEqual_broadcastAB) { + OpTester test("LessOrEqual", 12); + test.AddInput("A", {4, 2}, {10, 11, 12, 13, 14, 15, 16, 17}); + test.AddInput("B", {2}, {15, 7}); + test.AddOutput("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("A", {2}, {15, 7}); + test.AddInput("B", {4, 2}, {10, 11, 12, 13, 14, 15, 16, 17}); + test.AddOutput("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("A", {4, 1}, {10, 11, 12, 13}); + test.AddInput("B", {2}, {15, 7}); + test.AddOutput("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("A", {2}, {15, 7}); + test.AddInput("B", {4, 1}, {10, 11, 12, 13}); + test.AddOutput("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 dims{4}; @@ -1705,6 +1777,82 @@ TEST(MathOpTest, Greater_multidiretional_broadcastBA) { test.Run(); } +TEST(MathOpTest, GreaterOrEqual_12_float) { + OpTester test("GreaterOrEqual", 12); + std::vector dims{4}; + test.AddInput("A", dims, {1.0f, 0.0f, -1.0f, -1.0f}); + test.AddInput("B", dims, {1.0f, 1.0f, 2.0f, -1.0f}); + test.AddOutput("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 dims{4}; + test.AddInput("A", dims, {1.0, 0.0, 3.0, -1.0}); + test.AddInput("B", dims, {1.0, 1.0, 2.0, -1.0}); + test.AddOutput("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 dims{4}; + test.AddInput("A", dims, {10, 11, 12, 13}); + test.AddInput("B", dims, {15, 7, 12, 9}); + test.AddOutput("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 dims{4}; + test.AddInput("A", dims, {10, 11, 12, 13}); + test.AddInput("B", dims, {15, 7, 12, 9}); + test.AddOutput("C", dims, {false, true, true, true}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", + {kTensorrtExecutionProvider, kNnapiExecutionProvider, kOpenVINOExecutionProvider}); +} + +TEST(MathOpTest, GreaterOrEqual_broadcastAB) { + OpTester test("GreaterOrEqual", 12); + test.AddInput("A", {4, 2}, {10, 11, 12, 13, 14, 15, 16, 17}); + test.AddInput("B", {2}, {15, 7}); + test.AddOutput("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("A", {2}, {15, 7}); + test.AddInput("B", {4, 2}, {10, 11, 12, 13, 14, 15, 16, 17}); + test.AddOutput("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("A", {4, 1}, {10, 11, 12, 13}); + test.AddInput("B", {2}, {15, 7}); + test.AddOutput("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("A", {2}, {15, 7}); + test.AddInput("B", {4, 1}, {10, 11, 12, 13}); + test.AddOutput("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 dims{4}; diff --git a/onnxruntime/test/testdata/kernel_def_hashes/onnx.cpu.json b/onnxruntime/test/testdata/kernel_def_hashes/onnx.cpu.json index 34fb50318e..6cc4bd1ec5 100644 --- a/onnxruntime/test/testdata/kernel_def_hashes/onnx.cpu.json +++ b/onnxruntime/test/testdata/kernel_def_hashes/onnx.cpu.json @@ -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