Skip to content

[ONNX] Support Mod opset 28 #32369

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
Justin Chu (justinchuby) wants to merge 2 commits into
base: justinchuby-onnx-1-23-opset-28-base
Choose a base branch
from
Open

[ONNX] Support Mod opset 28 #32369

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
ec515d2
Support Mod opset 28 semantics
justinchuby Sep 2, 2026
31dc366
Fix CUDA integer fmod overflow
justinchuby Sep 2, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
9 changes: 6 additions & 3 deletions docs/OperatorKernels.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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**<br> *in* B:**T**<br> *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**<br> *in* B:**T**<br> *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**<br> *in* B:**T**<br> *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)|
Expand Down Expand Up @@ -858,7 +859,8 @@ The **OpSet Version** column uses the following notation:
|Min|*in* data_0:**T**<br> *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**<br> *in* B:**T**<br> *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**<br> *in* B:**T**<br> *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**<br> *in* B:**T**<br> *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)|
Expand Down Expand Up @@ -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**<br> *in* B:**T**<br> *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**<br> *in* B:**T**<br> *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**<br> *in* B:**T**<br> *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)|
Expand Down
1 change: 1 addition & 0 deletions onnxruntime/core/providers/cann/cann_graph.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -62,6 +62,7 @@ std::vector<NodeIndex> 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;
Expand Down
9 changes: 7 additions & 2 deletions onnxruntime/core/providers/cpu/cpu_execution_provider.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down Expand Up @@ -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! ***
Expand Down Expand Up @@ -2760,7 +2763,7 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, int16_t, Neg)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, int32_t, Neg)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, int64_t, Neg)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, Mod)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, 27, Mod)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, float, Abs)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, double, Abs)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 13, int8_t, Abs)>,
Expand Down Expand Up @@ -3742,6 +3745,8 @@ Status RegisterOnnxOperatorKernels(KernelRegistry& kernel_registry) {

// opset 27
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 27, Range)>,
// opset 28
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCpuExecutionProvider, kOnnxDomain, 28, Mod)>,
};
for (auto& function_table_entry : function_table) {
KernelCreateInfo info = function_table_entry();
Expand Down
133 changes: 124 additions & 9 deletions onnxruntime/core/providers/cpu/math/element_wise_ops.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,6 +12,7 @@
#include "core/mlas/inc/mlas.h"

#include <cmath>
#include <limits>

Check warning on line 15 in onnxruntime/core/providers/cpu/math/element_wise_ops.cc

View workflow job for this annotation

GitHub Actions / Optional Lint C++ 

[cpplint] reported by reviewdog 🐶
 Found C++ system header after other header. Should be: element_wise_ops.h, c system, c++ system, other.  [build/include_order] [4]

Raw Output:
onnxruntime/core/providers/cpu/math/element_wise_ops.cc:15:  Found C++ system header after other header. Should be: element_wise_ops.h, c system, c++ system, other.  [build/include_order] [4]

namespace onnxruntime {
// Supported types for operators that have type reduction enabled
Expand Down Expand Up @@ -2097,6 +2098,7 @@

private:
bool fmod_{false};
bool supports_float_floor_mod_{false};
bool divisor_is_validated_constant_{false};
};

Expand All @@ -2110,9 +2112,19 @@
BuildKernelDefConstraintsFromTypeList<EnabledModTypes>()),
Mod);

ONNX_CPU_OPERATOR_KERNEL(
ONNX_CPU_OPERATOR_VERSIONED_KERNEL(
Mod,
13,
27,
KernelDefBuilder()
.TypeConstraint(
"T",
BuildKernelDefConstraintsFromTypeList<EnabledModTypes>()),
Mod);

ONNX_CPU_OPERATOR_KERNEL(
Mod,
28,
KernelDefBuilder()
.TypeConstraint(
"T",
Expand Down Expand Up @@ -2160,6 +2172,12 @@

template <class T>
inline T Modulus(T x, T y) {
if constexpr (std::is_signed_v<T>) {
if (x == std::numeric_limits<T>::min() && y == T{-1}) {
return T{0};
}
}

auto res = x % y;
if ((res < 0 && y > 0) || (res > 0 && y < 0)) {
res += y;
Expand Down Expand Up @@ -2204,6 +2222,56 @@
UntypedBroadcastTwo(*context, funcs);
}

template <class T>
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 <class T>
void BroadCastFloorMod(OpKernelContext* context) {
ProcessBroadcastSpanFuncs funcs{
[](BroadcastHelper& per_iter_bh) {
const T& X = per_iter_bh.ScalarInput0<T>();
auto Y = per_iter_bh.SpanInput1<T>();
auto output = per_iter_bh.OutputSpan<T>();

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<T>();
const T& Y = per_iter_bh.ScalarInput1<T>();
auto output = per_iter_bh.OutputSpan<T>();

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<T>();
auto Y = per_iter_bh.SpanInput1<T>();
auto output = per_iter_bh.OutputSpan<T>();

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) {
Expand Down Expand Up @@ -2242,6 +2310,42 @@
UntypedBroadcastTwo(*context, funcs);
}

void BroadCastMLFloat16FloorMod(OpKernelContext* context) {
ProcessBroadcastSpanFuncs funcs{
[](BroadcastHelper& per_iter_bh) {
const auto X = per_iter_bh.ScalarInput0<MLFloat16>();
auto Y = per_iter_bh.SpanInput1<MLFloat16>();
auto output = per_iter_bh.OutputSpan<MLFloat16>();

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<MLFloat16>();
const MLFloat16 Y = per_iter_bh.ScalarInput1<MLFloat16>();
auto output = per_iter_bh.OutputSpan<MLFloat16>();

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<MLFloat16>();
auto Y = per_iter_bh.SpanInput1<MLFloat16>();
auto output = per_iter_bh.OutputSpan<MLFloat16>();

std::transform(X.begin(), X.end(), Y.begin(), output.begin(),

Check warning on line 2340 in onnxruntime/core/providers/cpu/math/element_wise_ops.cc

View workflow job for this annotation

GitHub Actions / Optional Lint C++ 

[cpplint] reported by reviewdog 🐶
 Add #include <algorithm> for transform  [build/include_what_you_use] [4]

Raw Output:
onnxruntime/core/providers/cpu/math/element_wise_ops.cc:2340:  Add #include <algorithm> for transform  [build/include_what_you_use] [4]
[](const MLFloat16& x, const MLFloat16& y) {
return MLFloat16(FloorMod(x.ToFloat(), y.ToFloat()));
});
}};

UntypedBroadcastTwo(*context, funcs);
}

template <class T, typename Enable = void>
struct CallModImpl;

Expand All @@ -2263,7 +2367,7 @@
// Generic implementation of Mod kernel, non-floating point types
template <class T>
struct CallModImpl<T, typename std::enable_if<!std::is_floating_point<T>::value>::type> {
void operator()(bool fmod, OpKernelContext* ctx) const {
void operator()(bool fmod, bool /*supports_float_floor_mod*/, OpKernelContext* ctx) const {
if (fmod) {
BroadCastFMod<T>(ctx);
} else {
Expand All @@ -2275,24 +2379,35 @@
// Generic implementation of Mod kernel, floating point types
template <class T>
struct CallModImpl<T, typename std::enable_if<std::is_floating_point<T>::value, void>::type> {
void operator()(bool fmod, OpKernelContext* ctx) const {
ORT_ENFORCE(fmod, "fmod attribute must be true for floating point types");
BroadCastFMod<T>(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<T>(ctx);
} else {
BroadCastFloorMod<T>(ctx);
}
}
};

// MLFloat16 implementation of Mod kernel
template <>
struct CallModImpl<MLFloat16> {
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<int64_t>("fmod", &fmod);
if (s.IsOK()) {
Expand Down Expand Up @@ -2327,7 +2442,7 @@
}

utils::MLTypeCallDispatcherFromTypeList<EnabledModTypes> t_disp(dt_type);
t_disp.Invoke<mod_internal::CallModImpl>(fmod_, context);
t_disp.Invoke<mod_internal::CallModImpl>(fmod_, supports_float_floor_mod_, context);

return Status::OK();
}
Expand Down
46 changes: 46 additions & 0 deletions onnxruntime/core/providers/cuda/cu_inc/common.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -501,6 +501,10 @@ __device__ __inline__ half _Gelu(half a) {

template <typename T>
__device__ __inline__ T _Mod(T a, T b) {
if (a == std::numeric_limits<T>::min() && b == T(-1)) {
return T(0);
}
Comment thread
justinchuby marked this conversation as resolved.

T r = a % b;
T zero = T(0);
if ((r > zero && b < zero) || (r < zero && b > zero)) {
Expand All @@ -509,8 +513,50 @@ __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<float>(a), static_cast<float>(b));
}

template <>
__device__ __inline__ BFloat16 _Mod(BFloat16 a, BFloat16 b) {
return _FloorMod(static_cast<float>(a), static_cast<float>(b));
}

template <typename T>
__device__ __inline__ T _Fmod(T a, T b) {
if (a == std::numeric_limits<T>::min() && b == T(-1)) {
return T(0);
}

return a % b;
}

Expand Down
9 changes: 7 additions & 2 deletions onnxruntime/core/providers/cuda/cuda_execution_provider.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down Expand Up @@ -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) {
Expand Down Expand Up @@ -2633,7 +2636,7 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, BFloat16, Tanh)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, BFloat16, Gemm)>,
BuildKernelCreateInfo<ONNX_OPERATOR_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, BFloat16, ReduceSum)>,
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, Mod)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 27, Mod)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 18, int8_t, QuantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 18, uint8_t, QuantizeLinear)>,
BuildKernelCreateInfo<ONNX_OPERATOR_VERSIONED_TYPED_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 13, 18, int8_t, DequantizeLinear)>,
Expand Down Expand Up @@ -3188,6 +3191,8 @@ static Status RegisterCudaKernels(KernelRegistry& kernel_registry) {
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 25, Unsqueeze)>,
// Opset 27
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 27, Range)>,
// Opset 28
BuildKernelCreateInfo<ONNX_OPERATOR_KERNEL_CLASS_NAME(kCudaExecutionProvider, kOnnxDomain, 28, Mod)>,
#endif
};

Expand Down
Loading
Loading