CUDA : looped PAD kernel for more than 65535 rows or slices

[pskrunner14]

Overview

CUDA GGML_OP_PAD kernel puts the output's row count in gridDim.y and its slice count (ne2*ne3) in gridDim.z. CUDA limits both to 65535. When a padded tensor exceeds either limit, the launch fails with invalid configuration argument and ggml aborts. The CPU backend handles the same shapes correctly.

ggml/src/ggml-cuda/pad.cu, pad_f32_cuda():

int  num_blocks = (ne0 + CUDA_PAD_BLOCK_SIZE - 1) / CUDA_PAD_BLOCK_SIZE;
dim3 gridDim(num_blocks, ne1, ne2 * ne3);   // y and z must each be <= 65535

The kernel reads i1 = blockIdx.y and i2/i3 from blockIdx.z, so the grid has to cover the tensor exactly. Any ne1 > 65535 or ne2*ne3 > 65535 (counted on the padded output) cannot be launched.

The fix is to do a looped version of the kernel. The grid is capped at 65535 in y and z and blocks loop over rows or slices beyond that. As @JohannesGaessler suggested in last PR this is preferable over the two kernel variant.

Additional information

How to reproduce

New test-backend-ops eval cases:

// more than 65535 rows or slices, beyond the CUDA grid.y/grid.z limit
test_pad(GGML_TYPE_F32, {4, 70000, 1, 1}, 1, 1, false);
test_pad(GGML_TYPE_F32, {4, 70000, 1, 1}, 1, 1, true);
test_pad_ext(GGML_TYPE_F32, {4, 2, 300, 300}, 1, 1, 0, 0, 0, 0, 0, 0, 0, false);

On unmodified master, build/bin/test-backend-ops -o PAD -b CUDA0 passes the existing cases, then aborts on the first new one:

CUDA error: invalid configuration argument
  current device: 0, in function ggml_cuda_compute_forward at .../ggml-cuda.cu

How we hit this bug

The cache-aware streaming Nemotron ASR encoder's causal conv subsampling pads a [T, F, 256, B] tensor (256 conv channels, batch B). Batch 255 is the largest that launches on master, batch 256 gives ne2 * ne3 = 65536 and aborts.

Perf Impact

Environment: RTX A5000 (sm_86), driver 580.173, CUDA 12.8, at upstream 03aa006.

Shape	master	fix	delta
512x512, pad 1	3.88 us	4.63 us	+19.3%
4096x4096, pad 1	205.0 us	228.3 us	+11.4%
128x64x32x4, rp1=63	25.4 us	29.3 us	+15.5%
33x65x256x16, pad 3	395.3 us	408.1 us	+3.3%
33x65x256x255, pad 3	6283 us	6482 us	+3.2%

Note: there is 3-20% perf hit with this change

Correctness: all 37 PAD cases pass on CUDA0 (compared against CPU), including the 3 new ones.

Requirements

• I have read and agree with the contributing guidelines
• AI usage disclosure: Yes, AI was used to speed up the original version of the fix here