CUDA: fix NORM/RMS_NORM/L2_NORM for more than 65535 channels or samples

[edenfunf]

Overview

Fixes #27901, refs #27911.

The CUDA norm kernels use (nrows, nchannels, nsamples) as the launch grid, but gridDim.y/z is limited to 65535. When ne[2] or ne[3] exceeds that, the launch fails with invalid argument.

Affected kernels:

• norm_f32
• rms_norm_f32
• l2_norm_f32
• rms_norm_mul_rope_f32

The fix follows #25103 and #22944: clamp grid.y/z to 65535 and loop over the remaining channels/samples in the kernel.

Since gridDim.y can no longer be used as the real channel count, nchannels and nsamples are now passed explicitly.

Multi-warp variants also need a trailing __syncthreads() because block_reduce reuses the same shared buffer across loop iterations.

Since the previous version

• Rebased onto master with CUDA: fuse RMS_NORM + SCALE into one kernel #29393 and integrated the new do_scale path.
• Removed the relocated ggml_cuda_pdl_lc() in l2_norm_f32 based on review feedback.
• Merged the rms_norm_mul_rope NMSE override with the existing one on master.
• Added [[maybe_unused]] for mulc / addc in non-fused instantiations.

Testing

Tested on RTX 5070, CUDA 13.3.

Added 7 cases with ne[2] or ne[3] = 65536. They fail on master and pass against the CPU backend with this patch.

Full test-backend-ops:

• 17689 / 17689 passed

Also checked additional boundary cases for 1024-thread paths, nrows > 1, sample-axis overflow, 65535 boundary, multiple loop iterations, non-contiguous views, and fused variants.

Performance

Shape	master	patched
RMS_NORM 4096x512	14.8 us	14.3 us
RMS_NORM 8192x1	2.5 us	2.0 us
NORM 4096x512	17.1 us	17.0 us
RMS_NORM [4,1,65535,1]	136 us	159 us
NORM [4,1,65535,1]	52 us	57 us

Normal shapes are flat or slightly faster.

The slowdown is only on the degenerate ~65k-block cases with very few active threads per block. I tried a few loop variants and none improved it, so this looks like loop overhead rather than a specific implementation detail.

A split fast/slow path could avoid that regression, but would add another kernel instantiation. For now this keeps a single loop-based path, consistent with the existing getrows.cu approach.

Requirements

• I have read and agree with the [contributing guidelines](...)
• AI usage disclosure: YES - AI was used for issue investigation and part of the implementation. I reviewed and verified the changes.

[ggml-gh-bot]

Hi @edenfunf, thanks for your contribution!

Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:

• PR Template not respected: Please respect the template when creating a new pull request. Make sure to fill out all required sections.

---

Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below.

[edenfunf]

@ggerganov If you have a chance, could you take a look at this PR? Thanks!

[JohannesGaessler]

Sorry for the long radio silence. I took a break from llama.cpp and am currently working through my backlog.

[JohannesGaessler]

Suggested change

ggml_cuda_pdl_lc();

The PDL launch completion does not affect correctness has to be placed essentially experimentally. If should not just be moved around like this. So I think the safest bet is to just remove it. cc @aendk

[edenfunf]

Removed ggml_cuda_pdl_lc() from l2_norm_f32 as suggested, so it now matches master.

Also rebased onto master with #29393 and integrated the do_scale path. Full test-backend-ops passes.