CUDA : fix PAD 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 in this PR:

• For shapes that fit the original configuration: same kernel with only change being that threads past the end of the row exit early.
• Larger shapes: cap gridDim.y & gridDim.z at 65535 and each block loops over the remaining rows and slices.

Additional information

How to reproduce

Environment: RTX A5000 (sm_86), driver 580.173, CUDA 12.8, at upstream 926862e.

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:2428

How we hit this bug

Nemotron speech 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

Shape	master	fix	change
512x512, pad 1	3.91 us	3.72 us	-4.7%
4096x4096, pad 1	205.5 us	205.0 us	-0.3%
33x65x256x16 (Nemotron subsampling, batch 16)	401 us	364 us	-9.1%
33x65x256x255 (same, batch 255)	6365 us	5784 us	-9.1%

Rows narrower than a block get 5-10% faster because out of range threads now exit before computing the slice index.

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

[JohannesGaessler]

This is overengineered, just always do the looped version for ops that aren't performance critical.

[pskrunner14]

@JohannesGaessler understood. I'll submit the looped version