[CUDA:Bugfix] Fix int32 overflow in im2col packing for large inputs (NaN/garbage output on conv layers with e*l > INT32_MAX)

[baicai-1145]

Problem

On CUDA, conv layers whose (output pixels) x (packed input channels) exceeds INT32_MAX silently produce garbage output (NaN probability maps, scrambled detections). We hit this with PP-OCR server segmentation models on 1280x1920 inputs (e.g. a 256-channel 9x9 conv at 320x480 output: 153600 x 20736 = 3,185,049,600 > 2^31).

Root cause - two int32 multiplication overflows

1. callIm2ColPack() (ConvBaseKernel.cu): size_t maxCount = e * lp; - e and lp are int; the product overflows before the assignment to size_t. The garbage value becomes the grid dimension, so cudaLaunchKernel fails with cudaErrorInvalidValue. Since checkKernelErrors is compiled out in release builds, the failure is silent and downstream layers consume uninitialized memory.

2. Im2Col_FilterC_Vec4 / Im2Col_FilterC kernels: size_t dst_offset = eIndex * l_p + kI * ic + iz; - the RHS is computed entirely in int and overflows before promotion, corrupting the im2col buffer write addressing even when the launch succeeds.

Fix

Promote the multiplications to size_t (cast operands before the multiply). Same-pattern overflows fixed in ConvCutlassExecution.cu, DeconvSingleInputExecution.cu, MultiInputConvExecution.cu.

Verification (RTX A10G, sm_86, CUDA 13.0, MNN master @ cda4a6f (rebased from 3.6.1 d407447))

• PP-OCRv4 server det @ 1280x1920: maxdiff vs CPU output drops from 1.0 (garbage) to 0.06-0.15 (fp16-mix rounding); NaN count 96,170 -> 0.
• Overflow boundary confirmed by sweep: input height 1280 (e.lp = 2.12e9) passes; 1408 (2.34e9) fails before the patch and passes after - matching the INT32_MAX threshold exactly.
• End-to-end: detection boxes restored from 0 to expected count; small models (mobile/tiny, e.lp < 2^31) unaffected.

[CLAassistant]

All committers have signed the CLA.

[wangzhaode]

Thanks for the contribution, and for tracking down this overflow — the direction is right and it continues the size_t hardening that already exists on the allocation side (e.g. ConvCutlassExecution.cu:187). I reviewed all six changes and found that three of them are effective while three are unfortunately no-ops. Could you please take a look at the following before we merge?

1. Three of the six changes are truncated back to int

The maxCount widening in ConvCutlassExecution.cu:286, DeconvSingleInputExecution.cu:220 and MultiInputConvExecution.cu:160 all feed into callFloat2Half, whose parameter is still int:

// ConvBaseKernel.cuh:20
void callFloat2Half(const void* input, void* output, const int count, CUDARuntime* runtime);
//                                                   ^^^^^^^^^^^^^^^

The freshly computed size_t value is immediately narrowed at the call boundary, and inside the implementation int thread_count = count / 4; (ConvBaseKernel.cu:201) narrows it again before the kernel launch. So in the very scenario this PR targets, these three lines do not change behavior — they only make the code look fixed, which is arguably worse than leaving it untouched.

Could you either:

• widen count to size_t in callFloat2Half / callFloat2BFloat16 (declaration, definition, and thread_count), which would make these three edits meaningful; or
• drop these three edits from the PR so the change set reflects what is actually fixed.

2. The remaining three changes are correct, with a path-dependent caveat

ConvBaseKernel.cu:61, :130 and :266 are genuine fixes: the two dst_offset expressions really did multiply two ints before promotion, and the callIm2ColPack kernels do take const size_t maxCount, so widening there takes effect.

However, the fix is only complete for one of the two im2col paths, because DivModFast is int-based:

// MNNCUDAFunction.cuh:30
__device__ __inline__ void divmod(int idx, int &quo, int &rem)

• ic % 4 == 0 (the Im2Col_FilterC_Vec4 path): maxCount = e * lp / 4, so the grid-stride indexO stays below INT32_MAX while dst_offset (~e * lp) exceeds it. This is exactly the regime your patch repairs, and it is repaired correctly.
• ic % 4 != 0 (the Im2Col_FilterC path): maxCount = e * lp, so indexO itself reaches INT32_MAX and is truncated when passed into divmod. The magic-division in div() also stops being valid once idx >= 2^31, so eIndex / lpIndex are garbage regardless of how wide dst_offset is.

Widening dst_offset cannot rescue that second path. Would you mind either mentioning this limitation explicitly in the PR description, or adding a guard/diagnostic in callIm2ColPack when maxCount > INT32_MAX on the non-vectorized path, so users get a clear error instead of silent corruption?

3. Note on verification

Please be aware that our CI does not include a CUDA build or test job, so this change carries no automated coverage and rests entirely on code review. If you are able to share the reproducing shape (the e, l, ic values and precision mode) and a before/after result on real hardware, that would give us much more confidence in merging.

Regression risk of the change itself looks negligible — it is pure integer widening with no behavioral change at normal sizes, and the extra 64-bit arithmetic is immaterial in a memory-bound kernel. So once the points above are addressed this should be straightforward to land. Thanks again for the careful work.

[baicai-1145]

Thank you for the careful review — all points addressed in f338f39.

1. Widened callFloat2Half / callFloat2BFloat16 (declaration in ConvBaseKernel.cuh, both definitions in ConvBaseKernel.cu, and thread_count / block_num / block_size in the launch code) to size_t. The kernel-side Float22Half2 / Float22BFloat16 already take size_t maxCount and use grid-stride loops over size_t index, so the host-side widening now flows through end to end, and the three call-site edits are meaningful. CUDARuntime::blocks_num(const size_t) / threads_num() were already size_t, so no further narrowing exists on this path.

2. Added a guard in callIm2ColPack right before the non-vectorized Im2Col_FilterC dispatch: if maxCount > 2^31-1 it now prints Im2Col non-vectorized path does not support maxCount > 2^31-1 and returns, instead of silently corrupting output. Agreed that widening dst_offset cannot rescue the int-based DivModFast indexing on that path; the guard turns the silent corruption into a diagnosable error.

3. Verification data (A10G 24GB, SM86, CUDA 13, fp32 inference via BackendConfig::Precision_High):

Reproducing shape: PP-OCRv4_server_det, the 256ch 9x9 conv layer (p2o.pd_op.conv2d.50) at input 1280x1920 → that layer's e*l = 153600 * 20736 = 3,185,049,600 > 2^31.

input	layer e*l	before fix	after fix
1280x864	2.12e9 (< 2^31)	OK	OK
1280x1920	4.71e9 (> 2^31)	NaN prob map, 0 boxes	23/23 boxes match CPU
3552x1280	5.89e9 (> 2^31)	garbage	correct

Full PP-OCR detection+recognition matrix on this build: the previously failing server-det CUDA cells went from 0% to 191/224 passing against the paddle CPU baseline; the 33 remaining deltas are single-character recognition differences from fp32 reduction order (prob corr 1.000000), unrelated to this fix. Re-ran normal-size inputs (mobile/tiny det + rec across 18 languages) after the widening — no behavioral change, as expected for pure integer widening.

Worth noting for future work: this fix does not address the separate issue that the cutlass path materializes the full im2col buffer on the device (e * lp * sizeof(T) — ~12.7 GB for the shape above), which limits server-class models on large inputs regardless of the overflow. That would need either the (currently disabled) blocked-im2col path or an implicit-GEMM fallback; happy to file a separate issue if useful.