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

[edenfunf]

Overview

Fixes #27901, refs #27911.

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

Affected kernels:

• norm_f32
• rms_norm_f32 (plain and fused mul/add)
• l2_norm_f32
• rms_norm_mul_rope_f32

As in #25103 and #22944, clamp grid.y/z at launch and loop over the remaining channels/samples inside the kernel. Since the kernels previously derived nchannels from gridDim.y, the actual dimensions are now passed explicitly.

#27941 avoids one instance on the model side, but the op still fails on master, e.g. RMS_NORM with ne = [4,1,65536,1].

Additional information

• Multi-warp variants need a trailing __syncthreads() because block_reduce reuses the shared buffer, same as softmax.cu (CUDA: Fix data-races when reusing SMEM in block_reduce #26385).
• RMS_NORM_MUL_ROPE uses a slightly relaxed NMSE threshold for large positions due to the existing GPU/CPU trig difference (~5e-6 at ne2 = 60000).

Tested on RTX 5070. The new cases reproduce #27901 before the fix and pass against the CPU backend after it. The full test-backend-ops suite also passes.

A/B performance:

• RMS_NORM 4096x512: 15.07 -> 14.63 us
• RMS_NORM_MUL_ROPE 128x32x512: 16.66 -> 16.69 us

Requirements

• I have read and agree with the [contributing guidelines](https://github.com/ggml-org/llama.cpp/blob/master/CONTRIBUTING.md)
• 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.