microsoft / microsoft/onnxruntime
CUDA Mod kernel has no zero divisor check, unlike CPU after #27833
- 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
First steps
- Read the whole issue, then the project's contributing guide.
- Comment on the issue to say you are picking it up — it saves two people doing the same work.
- Fork the repository and make your change on a branch.
- Open a pull request that references the issue number.
Assessment
This issue has not been assessed yet.