microsoft / microsoft/onnxruntime

CUDA Mod kernel has no zero divisor check, unlike CPU after #27833

Open
#32,494 0 comments 1 reaction 2 assignees View on GitHub

@tianleiwu is already working on this.

Since Sep 10, 2026.

  • #32546 by @copilot-swe-agent — open
ep:CUDA
Dominant language
C++
Stars
21.9k
Forks
4.2k
Avg merge
4d 11h
Merged PRs (30d)
184

Description

### Describe the issue

The CUDA `Mod` kernel has no zero divisor check, so an integer `Mod` with a zero in the divisor is undefined behaviour there. The CPU kernel got a check in #27833 and the two now disagree.

`_Mod` in `onnxruntime/core/providers/cuda/cu_inc/common.cuh`, line 503:

```cpp
template
__device__ __inline__ T _Mod(T a, T b) {
T r = a % b;
T zero = T(0);
if ((r > zero && b < zero) || (r < zero && b > zero)) {
r += b;
}
return r;
}
```

`a % b` runs with no guard on `b`. `_Fmod` seven lines below has the same shape.

On CPU the same input is rejected. `CheckZeroDivisorImpl` in `onnxruntime/core/providers/cpu/math/element_wise_ops.cc` does `ORT_RETURN_IF(b_data[i] == T{0}, "Integer modulo by zero")` before the kernel runs, which came from #27833.

So the same model errors cleanly on CPU and hits undefined behaviour on CUDA.

### To reproduce

Run an integer `Mod` node where the divisor tensor contains a zero, on the CUDA EP. For example `Mod(int64[3] {-3, 4, 7}, int64[3] {0, 2, 3})` with `fmod=0`.

I have not run this. I do not have NVIDIA hardware and cannot build the CUDA EP, so this is a read of the source rather than a measured crash, and I would not want it treated as a verified repro. Filing it because the asymmetry with #27833 looked worth recording. Happy to close if it turns out the divisor is already validated somewhere upstream of the kernel.

### Urgency

Low.

### Platform

Linux

### OS Version

n/a, reported from source

### ONNX Runtime Installation

Built from Source

### ONNX Runtime Version or Commit ID

main, `common.cuh` line 503 as of 2026-09-09

### ONNX Runtime API

Python

### Architecture

X64

### Execution Provider

CUDA

Contributor guide

Open the contributing guide

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.