From 54e8f3b5e33ac2fb552909a771f6282929fddf78 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Thu, 10 Sep 2026 15:47:14 +0000 Subject: [PATCH 1/5] Initial plan From b158bbeeb01cfd23e0b01597b4ae4b79ad15164a Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Thu, 10 Sep 2026 15:55:42 +0000 Subject: [PATCH 2/5] Fix CUDA Mod zero divisor handling Co-authored-by: tianleiwu <30328909+tianleiwu@users.noreply.github.com> --- .../core/providers/cuda/cu_inc/common.cuh | 6 ++++ .../cuda/math/binary_elementwise_ops.cc | 28 +++++++++++++++++++ .../cuda/math/binary_elementwise_ops_impl.cu | 20 +++++++++++++ .../cuda/math/binary_elementwise_ops_impl.h | 3 ++ .../cpu/math/element_wise_ops_test.cc | 14 ++++++++++ 5 files changed, 71 insertions(+) diff --git a/onnxruntime/core/providers/cuda/cu_inc/common.cuh b/onnxruntime/core/providers/cuda/cu_inc/common.cuh index f53da982538bd..dbe84ada5488a 100644 --- a/onnxruntime/core/providers/cuda/cu_inc/common.cuh +++ b/onnxruntime/core/providers/cuda/cu_inc/common.cuh @@ -501,6 +501,9 @@ __device__ __inline__ half _Gelu(half a) { template __device__ __inline__ T _Mod(T a, T b) { + if (b == T(0)) { + return T(0); + } T r = a % b; T zero = T(0); if ((r > zero && b < zero) || (r < zero && b > zero)) { @@ -511,6 +514,9 @@ __device__ __inline__ T _Mod(T a, T b) { template __device__ __inline__ T _Fmod(T a, T b) { + if (b == T(0)) { + return T(0); + } return a % b; } diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc index babbb4b3ba672..dd05b88b88146 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc @@ -482,6 +482,34 @@ Status Mod::ComputeInternal(OpKernelContext* context) const { 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."); + if (element_type == on::TensorProto_DataType_INT32 || element_type == on::TensorProto_DataType_INT64 || + element_type == on::TensorProto_DataType_UINT32 || element_type == on::TensorProto_DataType_UINT64) { + const size_t divisor_count = prepare.rhs_tensor->Shape().Size(); + if (divisor_count > 0) { + auto has_zero_buffer = GetScratchBuffer(1, GetComputeStream(context)); + CUDA_RETURN_IF_ERROR(cudaMemsetAsync(has_zero_buffer.get(), 0, sizeof(int), Stream(context))); + switch (element_type) { + case on::TensorProto_DataType_INT32: + CheckZeroDivisor(Stream(context), prepare.rhs_tensor->Data(), divisor_count, has_zero_buffer.get()); + break; + case on::TensorProto_DataType_INT64: + CheckZeroDivisor(Stream(context), prepare.rhs_tensor->Data(), divisor_count, has_zero_buffer.get()); + break; + case on::TensorProto_DataType_UINT32: + CheckZeroDivisor(Stream(context), prepare.rhs_tensor->Data(), divisor_count, has_zero_buffer.get()); + break; + case on::TensorProto_DataType_UINT64: + CheckZeroDivisor(Stream(context), prepare.rhs_tensor->Data(), divisor_count, has_zero_buffer.get()); + break; + } + + int has_zero = 0; + // A device-to-pageable-host copy completes before cudaMemcpyAsync returns. + CUDA_RETURN_IF_ERROR( + cudaMemcpyAsync(&has_zero, has_zero_buffer.get(), sizeof(int), cudaMemcpyDeviceToHost, Stream(context))); + ORT_RETURN_IF(has_zero != 0, "Integer modulo by zero"); + } + } #define CASE_MOD_ELEMENT_TYPE(name, onnx_type, data_type) \ case onnx_type: { \ Impl_##name::MappedType>( \ 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..af02df2dd2c66 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.cu @@ -10,6 +10,26 @@ namespace onnxruntime { namespace cuda { +template +__global__ void CheckZeroDivisorKernel(const T* divisor_data, size_t count, int* has_zero) { + const size_t idx = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + if (idx < count && divisor_data[idx] == T{0}) { + atomicExch(has_zero, 1); + } +} + +template +void CheckZeroDivisor(cudaStream_t stream, const T* divisor_data, size_t count, int* has_zero) { + constexpr int threads_per_block = GridDim::maxThreadsPerBlock; + const int blocks_per_grid = static_cast(CeilDiv(count, threads_per_block)); + CheckZeroDivisorKernel<<>>(divisor_data, count, has_zero); +} + +template void CheckZeroDivisor(cudaStream_t, const int32_t*, size_t, int*); +template void CheckZeroDivisor(cudaStream_t, const int64_t*, size_t, int*); +template void CheckZeroDivisor(cudaStream_t, const uint32_t*, size_t, int*); +template void CheckZeroDivisor(cudaStream_t, const uint64_t*, size_t, int*); + #define BINARY_ELEMENTWISE_IMPL(name) \ BINARY_ELEMENTWISE_IMPL_DECLARATION(name) { \ BinaryElementWiseImpl(stream, \ 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 5311f73f01a78..089d5b741e906 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.h +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops_impl.h @@ -8,6 +8,9 @@ namespace onnxruntime { namespace cuda { +template +void CheckZeroDivisor(cudaStream_t stream, const T* divisor_data, size_t count, int* has_zero); + // These macros simplifies coding. To add a new op with following steps: // 1. Add a new entry in BINARY_OPS() list // 2. (optional) Define templated single element operator in binary_elementwise_ops_impl.cu 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..35cfa54758066 100644 --- a/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc +++ b/onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc @@ -4831,6 +4831,20 @@ TEST(ModOpTest, Mod_int32_by_zero_constant_initializer) { {}, nullptr, &execution_providers); } +#ifdef USE_CUDA +TEST(ModOpTest, Mod_int64_by_zero_CUDA) { + OpTester test("Mod", ModOp_ver); + test.AddInput("X", {3}, {-3, 4, 7}); + test.AddInput("Y", {3}, {0, 2, 3}); + test.AddOutput("Z", {3}, {0, 0, 0}); + std::vector> execution_providers; + execution_providers.push_back(DefaultCudaExecutionProvider()); + test.Run(OpTester::ExpectResult::kExpectFailure, + "Integer modulo by zero", + {}, nullptr, &execution_providers); +} +#endif + TEST(BitShiftOpTest, SimpleLeft) { OpTester test("BitShift", 11); test.AddAttribute("direction", "LEFT"); From b72682706a64de71da88270b202a703a813b724e Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Thu, 10 Sep 2026 15:56:31 +0000 Subject: [PATCH 3/5] Synchronize CUDA divisor validation Co-authored-by: tianleiwu <30328909+tianleiwu@users.noreply.github.com> --- onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc index dd05b88b88146..93280526a3fdc 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc @@ -504,9 +504,9 @@ Status Mod::ComputeInternal(OpKernelContext* context) const { } int has_zero = 0; - // A device-to-pageable-host copy completes before cudaMemcpyAsync returns. CUDA_RETURN_IF_ERROR( cudaMemcpyAsync(&has_zero, has_zero_buffer.get(), sizeof(int), cudaMemcpyDeviceToHost, Stream(context))); + CUDA_RETURN_IF_ERROR(cudaStreamSynchronize(Stream(context))); ORT_RETURN_IF(has_zero != 0, "Integer modulo by zero"); } } From d7f08645a1ee576999ac11976546863d3bc54848 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Thu, 10 Sep 2026 15:57:21 +0000 Subject: [PATCH 4/5] Separate CUDA validation dispatch Co-authored-by: tianleiwu <30328909+tianleiwu@users.noreply.github.com> --- onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc | 1 + 1 file changed, 1 insertion(+) diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc index 93280526a3fdc..ebdc390a8fd66 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc @@ -510,6 +510,7 @@ Status Mod::ComputeInternal(OpKernelContext* context) const { ORT_RETURN_IF(has_zero != 0, "Integer modulo by zero"); } } + #define CASE_MOD_ELEMENT_TYPE(name, onnx_type, data_type) \ case onnx_type: { \ Impl_##name::MappedType>( \ From 1d8e693eee9ce60d23d932110d395834049c8b76 Mon Sep 17 00:00:00 2001 From: "copilot-swe-agent[bot]" <198982749+Copilot@users.noreply.github.com> Date: Thu, 10 Sep 2026 15:58:35 +0000 Subject: [PATCH 5/5] Use pinned host validation flag Co-authored-by: tianleiwu <30328909+tianleiwu@users.noreply.github.com> --- .../core/providers/cuda/math/binary_elementwise_ops.cc | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc index ebdc390a8fd66..3f8a6e48711d5 100644 --- a/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc +++ b/onnxruntime/core/providers/cuda/math/binary_elementwise_ops.cc @@ -503,11 +503,11 @@ Status Mod::ComputeInternal(OpKernelContext* context) const { break; } - int has_zero = 0; + auto has_zero = AllocateBufferOnCPUPinned(1); CUDA_RETURN_IF_ERROR( - cudaMemcpyAsync(&has_zero, has_zero_buffer.get(), sizeof(int), cudaMemcpyDeviceToHost, Stream(context))); + cudaMemcpyAsync(has_zero.get(), has_zero_buffer.get(), sizeof(int), cudaMemcpyDeviceToHost, Stream(context))); CUDA_RETURN_IF_ERROR(cudaStreamSynchronize(Stream(context))); - ORT_RETURN_IF(has_zero != 0, "Integer modulo by zero"); + ORT_RETURN_IF(*has_zero != 0, "Integer modulo by zero"); } }