From ec515d2f80c25d3904f7a954caacdda9f8c29e05 Mon Sep 17 00:00:00 2001 From: Justin Chu Date: Tue, 1 Sep 2026 18:34:01 -0700 Subject: [PATCH 1/2] Support Mod opset 28 semantics Add native CPU and CUDA floor-mod behavior, register DirectML support, and cap uncertified EPs at opset 27 fallback. Cover mixed signs, signed zero, infinities, NaNs, zero divisors, and signed overflow. Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com> --- docs/OperatorKernels.md | 9 +- onnxruntime/core/providers/cann/cann_graph.cc | 1 + .../providers/cpu/cpu_execution_provider.cc | 9 +- .../providers/cpu/math/element_wise_ops.cc | 133 ++++++++++++++++-- .../core/providers/cuda/cu_inc/common.cuh | 42 ++++++ .../providers/cuda/cuda_execution_provider.cc | 9 +- .../cuda/math/binary_elementwise_ops.cc | 26 +++- .../cuda/math/binary_elementwise_ops.h | 2 + .../cuda/math/binary_elementwise_ops_impl.cu | 2 +- .../src/Operators/OperatorRegistration.cpp | 1 + .../OperatorAuthorHelper/OperatorVersions.h | 5 + .../migraphx/migraphx_execution_provider.cc | 4 + .../openvino/ov_versions/data_ops.cc | 4 + .../qnn/builder/opbuilder/mod_op_builder.cc | 3 + .../cpu/math/element_wise_ops_test.cc | 72 ++++++++++ 15 files changed, 298 insertions(+), 24 deletions(-) diff --git a/docs/OperatorKernels.md b/docs/OperatorKernels.md index 96a6c76f63c09..c868c9f719c6b 100644 --- a/docs/OperatorKernels.md +++ b/docs/OperatorKernels.md @@ -287,7 +287,8 @@ The **OpSet Version** column uses the following notation: |||12|**T** = tensor(double), tensor(float), tensor(float16), tensor(int32), tensor(int64), tensor(int8), tensor(uint32), tensor(uint64), tensor(uint8)| |||[8, 11]|**T** = tensor(double), tensor(float)| |||[6, 7]|**T** = tensor(float)| -|Mod|*in* A:**T**
*in* B:**T**
*out* C:**T**|13+|**T** = tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| +|Mod|*in* A:**T**
*in* B:**T**
*out* C:**T**|28+|**T** = tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| +|||[13, 27]|**T** = tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |||[10, 12]|**T** = tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |Mul|*in* A:**T**
*in* B:**T**
*out* C:**T**|14+|**T** = tensor(double), tensor(float), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |||13|**T** = tensor(double), tensor(float), tensor(int32), tensor(int64), tensor(uint32), tensor(uint64)| @@ -858,7 +859,8 @@ The **OpSet Version** column uses the following notation: |Min|*in* data_0:**T**
*out* min:**T**|13+|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16), tensor(int32), tensor(int64), tensor(uint32), tensor(uint64)| |||12|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16), tensor(int32), tensor(int64), tensor(uint32), tensor(uint64)| |||[6, 11]|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16)| -|Mod|*in* A:**T**
*in* B:**T**
*out* C:**T**|13+|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16), tensor(int32), tensor(int64), tensor(uint32), tensor(uint64)| +|Mod|*in* A:**T**
*in* B:**T**
*out* C:**T**|28+|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16), tensor(int32), tensor(int64), tensor(uint32), tensor(uint64)| +|||[13, 27]|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16), tensor(int32), tensor(int64), tensor(uint32), tensor(uint64)| |||[10, 12]|**T** = tensor(double), tensor(float), tensor(float16), tensor(int32), tensor(int64), tensor(uint32), tensor(uint64)| |Mul|*in* A:**T**
*in* B:**T**
*out* C:**T**|14+|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |||13|**T** = tensor(bfloat16), tensor(double), tensor(float), tensor(float16), tensor(int32), tensor(int64), tensor(uint32), tensor(uint64)| @@ -1388,7 +1390,8 @@ The **OpSet Version** column uses the following notation: |||12+|**T** = tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |||8+|**T** = tensor(float), tensor(float16)| |||6+|**T** = tensor(float), tensor(float16)| -|Mod|*in* A:**T**
*in* B:**T**
*out* C:**T**|13+|**T** = tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint8)| +|Mod|*in* A:**T**
*in* B:**T**
*out* C:**T**|28+|**T** = tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint8)| +|||13+|**T** = tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint8)| |||10+|**T** = tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint8)| |Mul|*in* A:**T**
*in* B:**T**
*out* C:**T**|14+|**T** = tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| |||13+|**T** = tensor(float), tensor(float16), tensor(int16), tensor(int32), tensor(int64), tensor(int8), tensor(uint16), tensor(uint32), tensor(uint64), tensor(uint8)| diff --git a/onnxruntime/core/providers/cann/cann_graph.cc b/onnxruntime/core/providers/cann/cann_graph.cc index 6e7844293b1ed..8caf9e4bf4724 100644 --- a/onnxruntime/core/providers/cann/cann_graph.cc +++ b/onnxruntime/core/providers/cann/cann_graph.cc @@ -62,6 +62,7 @@ std::vector SupportONNXModel(const GraphViewer& graph_viewer) { const auto& node = graph_viewer.GetNode(index); if (node->Domain() != kOnnxDomain || domain_version < lower_bound || + (node->OpType() == "Mod" && node->SinceVersion() >= 28) || !cann_supported_ops.count(node->OpType())) { unsupported_nodes.push_back(index); continue; diff --git a/onnxruntime/core/providers/cpu/cpu_execution_provider.cc b/onnxruntime/core/providers/cpu/cpu_execution_provider.cc index a79d9c3d24cde..26ccfb879b349 100644 --- a/onnxruntime/core/providers/cpu/cpu_execution_provider.cc +++ b/onnxruntime/core/providers/cpu/cpu_execution_provider.cc @@ -834,7 +834,7 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, int16_t, Neg); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, int32_t, Neg); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, int64_t, Neg); -class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, Mod); +class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, 27, Mod); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, float, Abs); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, double, Abs); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, int8_t, Abs); @@ -1549,6 +1549,9 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, // Opset 27 class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 27, Range); +// Opset 28 +class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, Mod); + // !!PLEASE READ BELOW!! Following that, add new entries above this comment /* *** IMPORTANT! *** @@ -2760,7 +2763,7 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, - BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -3742,6 +3745,8 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) { // opset 27 BuildKernelCreateInfo, + // opset 28 + BuildKernelCreateInfo, }; for (auto& function_table_entry : function_table) { KernelCreateInfo info = function_table_entry(); diff --git a/onnxruntime/core/providers/cpu/math/element_wise_ops.cc b/onnxruntime/core/providers/cpu/math/element_wise_ops.cc index 4a9c4646a2c34..35ebe2ab146ad 100644 --- a/onnxruntime/core/providers/cpu/math/element_wise_ops.cc +++ b/onnxruntime/core/providers/cpu/math/element_wise_ops.cc @@ -12,6 +12,7 @@ #include "core/mlas/inc/mlas.h" #include +#include namespace onnxruntime { // Supported types for operators that have type reduction enabled @@ -2097,6 +2098,7 @@ class Mod final : public OpKernel { private: bool fmod_{false}; + bool supports_float_floor_mod_{false}; bool divisor_is_validated_constant_{false}; }; @@ -2110,9 +2112,19 @@ ONNX_CPU_OPERATOR_VERSIONED_KERNEL( BuildKernelDefConstraintsFromTypeList()), Mod); -ONNX_CPU_OPERATOR_KERNEL( +ONNX_CPU_OPERATOR_VERSIONED_KERNEL( Mod, 13, + 27, + KernelDefBuilder() + .TypeConstraint( + "T", + BuildKernelDefConstraintsFromTypeList()), + Mod); + +ONNX_CPU_OPERATOR_KERNEL( + Mod, + 28, KernelDefBuilder() .TypeConstraint( "T", @@ -2160,6 +2172,12 @@ void BroadCastFMod(OpKernelContext* context) { template inline T Modulus(T x, T y) { + if constexpr (std::is_signed_v) { + if (x == std::numeric_limits::min() && y == T{-1}) { + return T{0}; + } + } + auto res = x % y; if ((res < 0 && y > 0) || (res > 0 && y < 0)) { res += y; @@ -2204,6 +2222,56 @@ void BroadCastMod(OpKernelContext* context) { UntypedBroadcastTwo(*context, funcs); } +template +inline T FloorMod(T x, T y) { + auto res = std::fmod(x, y); + if (res == T{0}) { + return std::copysign(T{0}, y); + } + + if ((res < T{0} && y > T{0}) || (res > T{0} && y < T{0})) { + res += y; + } + return res; +} + +template +void BroadCastFloorMod(OpKernelContext* context) { + ProcessBroadcastSpanFuncs funcs{ + [](BroadcastHelper& per_iter_bh) { + const T& X = per_iter_bh.ScalarInput0(); + auto Y = per_iter_bh.SpanInput1(); + auto output = per_iter_bh.OutputSpan(); + + std::transform(Y.begin(), Y.end(), output.begin(), + [X](T y) { + return FloorMod(X, y); + }); + }, + [](BroadcastHelper& per_iter_bh) { + auto X = per_iter_bh.SpanInput0(); + const T& Y = per_iter_bh.ScalarInput1(); + auto output = per_iter_bh.OutputSpan(); + + std::transform(X.begin(), X.end(), output.begin(), + [Y](T x) { + return FloorMod(x, Y); + }); + }, + [](BroadcastHelper& per_iter_bh) { + auto X = per_iter_bh.SpanInput0(); + auto Y = per_iter_bh.SpanInput1(); + auto output = per_iter_bh.OutputSpan(); + + std::transform(X.begin(), X.end(), Y.begin(), output.begin(), + [](T x, T y) { + return FloorMod(x, y); + }); + }}; + + UntypedBroadcastTwo(*context, funcs); +} + void BroadCastMLFloat16FMod(OpKernelContext* context) { ProcessBroadcastSpanFuncs funcs{ [](BroadcastHelper& per_iter_bh) { @@ -2242,6 +2310,42 @@ void BroadCastMLFloat16FMod(OpKernelContext* context) { UntypedBroadcastTwo(*context, funcs); } +void BroadCastMLFloat16FloorMod(OpKernelContext* context) { + ProcessBroadcastSpanFuncs funcs{ + [](BroadcastHelper& per_iter_bh) { + const auto X = per_iter_bh.ScalarInput0(); + auto Y = per_iter_bh.SpanInput1(); + auto output = per_iter_bh.OutputSpan(); + + std::transform(Y.begin(), Y.end(), output.begin(), + [X_fl = X.ToFloat()](const MLFloat16& y) { + return MLFloat16(FloorMod(X_fl, y.ToFloat())); + }); + }, + [](BroadcastHelper& per_iter_bh) { + auto X = per_iter_bh.SpanInput0(); + const MLFloat16 Y = per_iter_bh.ScalarInput1(); + auto output = per_iter_bh.OutputSpan(); + + std::transform(X.begin(), X.end(), output.begin(), + [Y_fl = Y.ToFloat()](const MLFloat16& x) { + return MLFloat16(FloorMod(x.ToFloat(), Y_fl)); + }); + }, + [](BroadcastHelper& per_iter_bh) { + auto X = per_iter_bh.SpanInput0(); + auto Y = per_iter_bh.SpanInput1(); + auto output = per_iter_bh.OutputSpan(); + + std::transform(X.begin(), X.end(), Y.begin(), output.begin(), + [](const MLFloat16& x, const MLFloat16& y) { + return MLFloat16(FloorMod(x.ToFloat(), y.ToFloat())); + }); + }}; + + UntypedBroadcastTwo(*context, funcs); +} + template struct CallModImpl; @@ -2263,7 +2367,7 @@ struct CheckZeroDivisorImpl { // Generic implementation of Mod kernel, non-floating point types template struct CallModImpl::value>::type> { - void operator()(bool fmod, OpKernelContext* ctx) const { + void operator()(bool fmod, bool /*supports_float_floor_mod*/, OpKernelContext* ctx) const { if (fmod) { BroadCastFMod(ctx); } else { @@ -2275,24 +2379,35 @@ struct CallModImpl::value> // Generic implementation of Mod kernel, floating point types template struct CallModImpl::value, void>::type> { - void operator()(bool fmod, OpKernelContext* ctx) const { - ORT_ENFORCE(fmod, "fmod attribute must be true for floating point types"); - BroadCastFMod(ctx); + void operator()(bool fmod, bool supports_float_floor_mod, OpKernelContext* ctx) const { + ORT_ENFORCE(fmod || supports_float_floor_mod, + "fmod attribute must be true for floating point types before opset 28"); + if (fmod) { + BroadCastFMod(ctx); + } else { + BroadCastFloorMod(ctx); + } } }; // MLFloat16 implementation of Mod kernel template <> struct CallModImpl { - void operator()(bool fmod, OpKernelContext* ctx) const { - ORT_ENFORCE(fmod, "fmod attribute must be true for floating point types"); - BroadCastMLFloat16FMod(ctx); + void operator()(bool fmod, bool supports_float_floor_mod, OpKernelContext* ctx) const { + ORT_ENFORCE(fmod || supports_float_floor_mod, + "fmod attribute must be true for floating point types before opset 28"); + if (fmod) { + BroadCastMLFloat16FMod(ctx); + } else { + BroadCastMLFloat16FloorMod(ctx); + } } }; } // namespace mod_internal Mod::Mod(const OpKernelInfo& info) : OpKernel(info) { + supports_float_floor_mod_ = info.node().SinceVersion() >= 28; int64_t fmod = 0; Status s = info.GetAttr("fmod", &fmod); if (s.IsOK()) { @@ -2327,7 +2442,7 @@ Status Mod::Compute(OpKernelContext* context) const { } utils::MLTypeCallDispatcherFromTypeList t_disp(dt_type); - t_disp.Invoke(fmod_, context); + t_disp.Invoke(fmod_, supports_float_floor_mod_, context); return Status::OK(); } diff --git a/onnxruntime/core/providers/cuda/cu_inc/common.cuh b/onnxruntime/core/providers/cuda/cu_inc/common.cuh index f53da982538bd..767ae875d4db4 100644 --- a/onnxruntime/core/providers/cuda/cu_inc/common.cuh +++ b/onnxruntime/core/providers/cuda/cu_inc/common.cuh @@ -501,6 +501,10 @@ __device__ __inline__ half _Gelu(half a) { template __device__ __inline__ T _Mod(T a, T b) { + if (a == std::numeric_limits::min() && b == T(-1)) { + return T(0); + } + T r = a % b; T zero = T(0); if ((r > zero && b < zero) || (r < zero && b > zero)) { @@ -509,6 +513,44 @@ __device__ __inline__ T _Mod(T a, T b) { return r; } +__device__ __inline__ float _FloorMod(float a, float b) { + float r = fmodf(a, b); + if (r == 0.0f) { + return copysignf(0.0f, b); + } + if ((r > 0.0f && b < 0.0f) || (r < 0.0f && b > 0.0f)) { + r += b; + } + return r; +} + +template <> +__device__ __inline__ float _Mod(float a, float b) { + return _FloorMod(a, b); +} + +template <> +__device__ __inline__ double _Mod(double a, double b) { + double r = fmod(a, b); + if (r == 0.0) { + return copysign(0.0, b); + } + if ((r > 0.0 && b < 0.0) || (r < 0.0 && b > 0.0)) { + r += b; + } + return r; +} + +template <> +__device__ __inline__ half _Mod(half a, half b) { + return _FloorMod(static_cast(a), static_cast(b)); +} + +template <> +__device__ __inline__ BFloat16 _Mod(BFloat16 a, BFloat16 b) { + return _FloorMod(static_cast(a), static_cast(b)); +} + template __device__ __inline__ T _Fmod(T a, T b) { return a % b; diff --git a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc index 9729b495dc1f8..288b2f7282da3 100755 --- a/onnxruntime/core/providers/cuda/cuda_execution_provider.cc +++ b/onnxruntime/core/providers/cuda/cuda_execution_provider.cc @@ -1342,7 +1342,7 @@ class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, BFloat16, Tanh); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, BFloat16, Gemm); class ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, BFloat16, ReduceSum); -class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, Mod); +class ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 27, Mod); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 18, int8_t, QuantizeLinear); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 18, uint8_t, QuantizeLinear); class ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 18, int8_t, DequantizeLinear); @@ -1897,6 +1897,9 @@ class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, U // Opset 27. class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 27, Range); + +// Opset 28. +class ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, Mod); #endif static Status RegisterCudaKernels(KernelRegistry& kernel_registry) { @@ -2633,7 +2636,7 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, - BuildKernelCreateInfo, + BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, BuildKernelCreateInfo, @@ -3188,6 +3191,8 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) { BuildKernelCreateInfo, // Opset 27 BuildKernelCreateInfo, + // Opset 28 + BuildKernelCreateInfo, #endif }; diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc index babbb4b3ba672..e4c398938a26f 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc @@ -467,21 +467,29 @@ ONNX_OPERATOR_VERSIONED_KERNEL_EX( BuildKernelDefConstraints()), Mod); -ONNX_OPERATOR_KERNEL_EX(Mod, kOnnxDomain, 13, kCudaExecutionProvider, - (*KernelDefBuilder::Create()) - .TypeConstraint("T", BuildKernelDefConstraints()), - Mod); +ONNX_OPERATOR_VERSIONED_KERNEL_EX( + Mod, kOnnxDomain, 13, 27, kCudaExecutionProvider, + (*KernelDefBuilder::Create()) + .TypeConstraint("T", BuildKernelDefConstraints()), + Mod); + +ONNX_OPERATOR_KERNEL_EX( + Mod, kOnnxDomain, 28, kCudaExecutionProvider, + (*KernelDefBuilder::Create()) + .TypeConstraint("T", BuildKernelDefConstraints()), + Mod); Status Mod::ComputeInternal(OpKernelContext* context) const { namespace on = ONNX_NAMESPACE; BinaryElementwisePreparation prepare; ORT_RETURN_IF_ERROR(Prepare(context, &prepare)); auto element_type = prepare.lhs_tensor->GetElementType(); - ORT_ENFORCE(fmod_ || element_type == on::TensorProto_DataType_INT32 || + ORT_ENFORCE(fmod_ || supports_float_floor_mod_ || element_type == on::TensorProto_DataType_INT32 || element_type == on::TensorProto_DataType_INT64 || element_type == on::TensorProto_DataType_UINT32 || element_type == on::TensorProto_DataType_UINT64, - "Non-fmod can support integer types only."); + "Non-fmod can support floating point types starting in opset 28."); #define CASE_MOD_ELEMENT_TYPE(name, onnx_type, data_type) \ case onnx_type: { \ Impl_##name::MappedType>( \ @@ -514,6 +522,10 @@ Status Mod::ComputeInternal(OpKernelContext* context) const { CASE_MOD_ELEMENT_TYPE(Mod, on::TensorProto_DataType_INT64, int64_t); CASE_MOD_ELEMENT_TYPE(Mod, on::TensorProto_DataType_UINT32, uint32_t); CASE_MOD_ELEMENT_TYPE(Mod, on::TensorProto_DataType_UINT64, uint64_t); + CASE_MOD_ELEMENT_TYPE(Mod, on::TensorProto_DataType_FLOAT, float); + CASE_MOD_ELEMENT_TYPE(Mod, on::TensorProto_DataType_DOUBLE, double); + CASE_MOD_ELEMENT_TYPE(Mod, on::TensorProto_DataType_FLOAT16, MLFloat16); + CASE_MOD_ELEMENT_TYPE(Mod, on::TensorProto_DataType_BFLOAT16, BFloat16); default: return ORT_MAKE_STATUS(ONNXRUNTIME, INVALID_ARGUMENT, "Unsupported element type: ", DataTypeImpl::ToString(prepare.lhs_tensor->DataType())); diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h index 8bc40ae7d2ce7..337c7909f46e9 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.h @@ -221,11 +221,13 @@ class Mod final : public BinaryElementwise { Mod(const OpKernelInfo& info) : BinaryElementwise(info) { int64_t fmod = info.GetAttrOrDefault("fmod", 0LL); fmod_ = fmod != 0; + supports_float_floor_mod_ = info.node().SinceVersion() >= 28; } Status ComputeInternal(OpKernelContext* context) const override; private: bool fmod_{false}; + bool supports_float_floor_mod_{false}; }; template 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 4e71a9f7c5090..2b08614b1e9c4 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu @@ -162,7 +162,7 @@ SPECIALIZED_BINARY_ELEMENTWISE_IMPL(Xor, bool) SPECIALIZED_BINARY_ELEMENTWISE_IMPL_HFD(PRelu) SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Max) SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Min) -SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZIL(Mod) +SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Mod) SPECIALIZED_BINARY_ELEMENTWISE_IMPL_UZILHFD(Fmod) // create declarations for impl for Pow diff --git a/onnxruntime/core/providers/dml/DmlExecutionProvider/src/Operators/OperatorRegistration.cpp b/onnxruntime/core/providers/dml/DmlExecutionProvider/src/Operators/OperatorRegistration.cpp index 26f998c7521a2..9536e7b505e7c 100644 --- a/onnxruntime/core/providers/dml/DmlExecutionProvider/src/Operators/OperatorRegistration.cpp +++ b/onnxruntime/core/providers/dml/DmlExecutionProvider/src/Operators/OperatorRegistration.cpp @@ -1143,6 +1143,7 @@ constexpr static OperatorRegistrationInformation operatorRegistrationInformation {REG_INFO( 20, IsInf, typeNameListTwo, supportedTypeListIsInf, DmlGraphSupport::Supported)}, {REG_INFO( 10, Mod, typeNameListDefault, supportedTypeListNumericDefault, DmlGraphSupport::Supported)}, {REG_INFO( 13, Mod, typeNameListDefault, supportedTypeListNumericDefault, DmlGraphSupport::Supported)}, + {REG_INFO( 28, Mod, typeNameListDefault, supportedTypeListNumericDefault, DmlGraphSupport::Supported)}, {REG_INFO( 11, BitShift, typeNameListDefault, supportedTypeListUInt8to64, DmlGraphSupport::Supported)}, {REG_INFO( 11, Round, typeNameListDefault, supportedTypeListFloat16to32, DmlGraphSupport::Supported)}, {REG_INFO( 10, ReverseSequence, typeNameListDefault, supportedTypeListAllScalars, DmlGraphSupport::Supported)}, diff --git a/onnxruntime/core/providers/dml/OperatorAuthorHelper/OperatorVersions.h b/onnxruntime/core/providers/dml/OperatorAuthorHelper/OperatorVersions.h index b4d402a1d9e77..2fea7ad80effe 100644 --- a/onnxruntime/core/providers/dml/OperatorAuthorHelper/OperatorVersions.h +++ b/onnxruntime/core/providers/dml/OperatorAuthorHelper/OperatorVersions.h @@ -457,6 +457,11 @@ namespace OperatorHelper static const int sc_sinceVer_GroupNorm = 21; } + namespace OnnxOperatorSet28 + { + static const int sc_sinceVer_Mod = 28; + } + namespace MsftOperatorSet1 { static const int sc_sinceVer_DmlFusedConv = 1; diff --git a/onnxruntime/core/providers/migraphx/migraphx_execution_provider.cc b/onnxruntime/core/providers/migraphx/migraphx_execution_provider.cc index 8908013a52291..dbaac59e6ce15 100644 --- a/onnxruntime/core/providers/migraphx/migraphx_execution_provider.cc +++ b/onnxruntime/core/providers/migraphx/migraphx_execution_provider.cc @@ -666,6 +666,10 @@ static bool IsNodeSupported(const std::set& op_set, const auto& optype = node->OpType(); const auto& domain = node->Domain(); + if (domain == kOnnxDomain && optype == "Mod" && node->SinceVersion() >= 28) { + return false; + } + // Three types of checking: // 1. Check input and output data types are supported. // 2. Check op_type is implemented in migraphx diff --git a/onnxruntime/core/providers/openvino/ov_versions/data_ops.cc b/onnxruntime/core/providers/openvino/ov_versions/data_ops.cc index 7aae81b987c0d..2b66dc07c547f 100644 --- a/onnxruntime/core/providers/openvino/ov_versions/data_ops.cc +++ b/onnxruntime/core/providers/openvino/ov_versions/data_ops.cc @@ -736,6 +736,10 @@ bool DataOps::node_is_supported(const NodeIndex node_idx, bool& has_external_wei const auto& domain = node->Domain(); + if (domain == kOnnxDomain && optype == "Mod" && node->SinceVersion() >= 28) { + return false; + } + /* 0. Check if node is in the unsupported list 1. Check input and output data types are supported. diff --git a/onnxruntime/core/providers/qnn/builder/opbuilder/mod_op_builder.cc b/onnxruntime/core/providers/qnn/builder/opbuilder/mod_op_builder.cc index 1d61ac9b9599d..8fae4a3f9f9dd 100644 --- a/onnxruntime/core/providers/qnn/builder/opbuilder/mod_op_builder.cc +++ b/onnxruntime/core/providers/qnn/builder/opbuilder/mod_op_builder.cc @@ -46,6 +46,9 @@ Status ModOpBuilder::ProcessInputs(QnnModelWrapper& qnn_model_wrapper, std::vector& input_names, bool do_op_validation) const { ORT_UNUSED_PARAMETER(do_op_validation); + ORT_RETURN_IF(node_unit.SinceVersion() >= 28, + "QNN Mod does not support opset 28 floor-mod edge semantics."); + NodeAttrHelper node_helper(node_unit); int64_t fmod = node_helper.Get("fmod", static_cast(0)); // 0=integer mod. 1=float mod. if (1 == fmod) { 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 278372bcb09f4..4127fbf402b35 100644 --- a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc +++ b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc @@ -4584,6 +4584,7 @@ TEST(MathOpTest, ErfCheckMultiThreadDataChunking) { } constexpr int ModOp_ver = 10; +constexpr int ModOp_ver28 = 28; TEST(ModOpTest, Fmod_float_mixed_sign) { OpTester test("Mod", ModOp_ver); @@ -4636,6 +4637,77 @@ TEST(ModOpTest, Fmod_bfloat16_mixed_sign) { } #endif +TEST(ModOpTest, FloorMod_float_mixed_sign) { + OpTester test("Mod", ModOp_ver28); + test.AddInput("X", {6}, {-4.3f, 7.2f, 5.0f, 4.3f, -7.2f, 8.0f}); + test.AddInput("Y", {6}, {2.1f, -3.4f, 8.0f, -2.1f, 3.4f, 5.0f}); + test.AddOutput("Z", {6}, {2.0f, -3.0f, 5.0f, -2.0f, 3.0f, 3.0f}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", + {kTensorrtExecutionProvider, kQnnExecutionProvider}); +} + +TEST(ModOpTest, FloorMod_double_mixed_sign) { + OpTester test("Mod", ModOp_ver28); + test.AddInput("X", {6}, {-4.3, 7.2, 5.0, 4.3, -7.2, 8.0}); + test.AddInput("Y", {6}, {2.1, -3.4, 8.0, -2.1, 3.4, 5.0}); + test.AddOutput("Z", {6}, {2.0, -3.0, 5.0, -2.0, 3.0, 3.0}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", + {kTensorrtExecutionProvider, kQnnExecutionProvider}); +} + +TEST(ModOpTest, FloorMod_float16_mixed_sign) { + OpTester test("Mod", ModOp_ver28); + test.AddInput("X", {6}, MakeMLFloat16({-4.3f, 7.2f, 5.0f, 4.3f, -7.2f, 8.0f})); + test.AddInput("Y", {6}, MakeMLFloat16({2.1f, -3.4f, 8.0f, -2.1f, 3.4f, 5.0f})); + test.AddOutput("Z", {6}, MakeMLFloat16({2.0f, -3.0f, 5.0f, -2.0f, 3.0f, 3.0f})); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", + {kTensorrtExecutionProvider, kQnnExecutionProvider}); +} + +TEST(ModOpTest, FloorMod_float_edge_cases) { + const float nan = std::numeric_limits::quiet_NaN(); + const float inf = std::numeric_limits::infinity(); + OpTester test("Mod", ModOp_ver28); + test.AddInput("X", {14}, + {0.0f, -0.0f, 0.0f, -0.0f, -3.0f, 3.0f, -1.0f, 1.0f, + inf, -inf, 1.0f, 1.0f, nan, 1.0f}); + test.AddInput("Y", {14}, + {-2.0f, 2.0f, 2.0f, -2.0f, inf, inf, -inf, -inf, + 2.0f, 2.0f, 0.0f, -0.0f, 2.0f, nan}); + test.AddOutput("Z", {14}, + {-0.0f, 0.0f, 0.0f, -0.0f, inf, 3.0f, -1.0f, -inf, + nan, nan, nan, nan, nan, nan}); + test.SetCustomOutputVerifier([](const std::vector& fetches, + const std::string& /*provider_type*/) { + ASSERT_EQ(fetches.size(), 1u); + ASSERT_TRUE(fetches[0].IsTensor()); + const float* output = fetches[0].Get().Data(); + + EXPECT_TRUE(std::signbit(output[0])); + EXPECT_FALSE(std::signbit(output[1])); + EXPECT_FALSE(std::signbit(output[2])); + EXPECT_TRUE(std::signbit(output[3])); + EXPECT_EQ(output[4], std::numeric_limits::infinity()); + EXPECT_EQ(output[5], 3.0f); + EXPECT_EQ(output[6], -1.0f); + EXPECT_EQ(output[7], -std::numeric_limits::infinity()); + for (size_t i = 8; i < 14; ++i) { + EXPECT_TRUE(std::isnan(output[i])) << "output index " << i; + } + }); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", + {kTensorrtExecutionProvider, kQnnExecutionProvider}); +} + +TEST(ModOpTest, Signed_integer_overflow_case) { + OpTester test("Mod", ModOp_ver28); + test.AddInput("X", {2}, {std::numeric_limits::min(), 7}); + test.AddInput("Y", {2}, {-1, -3}); + test.AddOutput("Z", {2}, {0, -2}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", + {kTensorrtExecutionProvider, kQnnExecutionProvider}); +} + TEST(ModOpTest, Int8_mixed_sign) { OpTester test("Mod", ModOp_ver); test.AddInput("X", {6}, {-4, 7, 5, 4, -7, 8}); From 31dc366f662bd8c0bf173e565d6e38a09cb77491 Mon Sep 17 00:00:00 2001 From: Justin Chu Date: Wed, 2 Sep 2026 07:15:27 -0700 Subject: [PATCH 2/2] Fix CUDA integer fmod overflow Co-authored-by: Copilot App <223556219+Copilot@users.noreply.github.com> --- onnxruntime/core/providers/cuda/cu_inc/common.cuh | 4 ++++ .../test/providers/cpu/math/element_wise_ops_test.cc | 10 ++++++++++ 2 files changed, 14 insertions(+) diff --git a/onnxruntime/core/providers/cuda/cu_inc/common.cuh b/onnxruntime/core/providers/cuda/cu_inc/common.cuh index 767ae875d4db4..a1962a4248d05 100644 --- a/onnxruntime/core/providers/cuda/cu_inc/common.cuh +++ b/onnxruntime/core/providers/cuda/cu_inc/common.cuh @@ -553,6 +553,10 @@ __device__ __inline__ BFloat16 _Mod(BFloat16 a, BFloat16 b) { template __device__ __inline__ T _Fmod(T a, T b) { + if (a == std::numeric_limits::min() && b == T(-1)) { + return T(0); + } + return a % b; } 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 4127fbf402b35..aaf16ad80344b 100644 --- a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc +++ b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc @@ -4708,6 +4708,16 @@ TEST(ModOpTest, Signed_integer_overflow_case) { {kTensorrtExecutionProvider, kQnnExecutionProvider}); } +TEST(ModOpTest, Signed_integer_overflow_case_fmod) { + OpTester test("Mod", ModOp_ver28); + test.AddAttribute("fmod", 1); + test.AddInput("X", {2}, {std::numeric_limits::min(), 7}); + test.AddInput("Y", {2}, {-1, -3}); + test.AddOutput("Z", {2}, {0, 1}); + test.Run(OpTester::ExpectResult::kExpectSuccess, "", + {kTensorrtExecutionProvider, kQnnExecutionProvider}); +} + TEST(ModOpTest, Int8_mixed_sign) { OpTester test("Mod", ModOp_ver); test.AddInput("X", {6}, {-4, 7, 5, 4, -7, 8});