fix(gguf): chunk the moe_vec z grid past the 65535 cap

[avlp12]

The bug

moe_vec_q puts the routed-row count (tokens * top_k) on gridDim.z, which the
hardware caps at 65535 — only gridDim.x reaches 2^31-1. Any grouped MoE GEMV with
more than 65535 routed rows therefore fails to launch outright with
cudaErrorInvalidValue.

At top_k = 6 the largest chunk that fits is 10922 tokens (10923 * 6 = 65538). A
DeepSeek-V4 deployment at --max-prefill-length 16384 (z = 98304) dies on its very
first prefill chunk, while 8192 (z = 49152) works. The limit is inherited from the
vendored kernel and was replicated across all 19 per-type launchers.

The fix

The row count has to stay on z. blockIdx.x is the fastest-varying axis and
consecutive x-blocks are consecutive weight rows of the same (token, expert)
pair, which is what keeps one expert's rows co-resident in L2 for this
bandwidth-bound GEMV. Moving the count to x would scatter the weight streaming.

So slice z instead, and pass each launch the base of its chunk in a new
z_offset argument — the same shape as quantize_row_q8_1_cuda's existing
MAX_BLOCK_SIZE loop. An explicit offset rather than a bumped base pointer,
because the kernel derives token = z / topk from the absolute z, and a
pointer scheme would also force every chunk stride to be a multiple of topk.

Also promoted to 64-bit, so this does not merely turn a hard crash into silent
corruption:

• dst[z * nrows + row], which wrapped around T ~ 174762
• token * token_stride (token is now int64, so the product widens)
• tokens / top_k in the launcher signatures, so the product cannot overflow at
  the call boundary
• &x[off * kx] in quantize_row_q8_1_cuda, which wraps for off > 524288 at
  kx = 4096 and genuinely sees large ky (= tokens * top_k on the down GEMV)

The 19 per-type launchers differed only in their template arguments and all 19 had
copied the capped launch, so they are now generated from one shared
moe_vec_launch helper via a macro instead of hand-maintained. The generated
launchers are argument-for-argument identical to the ones they replace.

Test

tests/kernels/test_moe_vec_large_rows.py covers 12288 and 16384 tokens at
top_k = 6 (z = 73728 / 98304) plus the top_k = 1 down-projection shape, and
anchors the sub-cap case against a dense ggml_dequantize reference.

The assertion has teeth. "Does not throw" is the cheap half — a wrong z_offset
still launches fine, it just reads the wrong token, the wrong expert id, or writes
the wrong dst row. So the tests compare one over-cap call against the
concatenation of sub-cap slices and assert bit-identical results (every routed
row is an independent dot product with an identical reduction order in both paths).

Mutation-tested: with the kernel reverted to z = blockIdx.z, i.e. ignoring
z_offset, all four value comparisons fail while all four "does not throw" cases
stay green.

FAILED test_over_cap_matches_sub_cap_slices[12288]
FAILED test_over_cap_matches_sub_cap_slices[16384]
FAILED test_rows_straddling_the_old_cap_are_individually_correct
FAILED test_down_projection_over_cap_matches_slices
4 failed, 4 passed

With the fix in place: 8 passed. tests/kernels goes from 204 to 212 passed with
an unchanged failure set (5 pre-existing test_pinned_tensor.py failures in this
environment); tests/moe is unchanged.

Verified on RTX 5090 (sm_120), CUDA 13.3 / nvcc 13.3.73, torch 2.11.0+cu130.

Note: on a host without clang++, building this branch also needs #159 (the
-std=c++20 flag) — the two are independent fixes but #159 is what lets the
extension compile with gcc at all.

[salekseev]

Measured this against #185 on sm_89, since the two PRs fix the same bug with opposite
reasoning about the axis. Your locality argument holds on this hardware.

Same image, same flags, same model, only the kernel patch swapped; separate JIT cache
directories so neither variant can load the other's .so. JIT warmed before measuring.

prompt tokens	#185 (index on grid.x)	#211 (z chunked)	Δ
35415	770.7 tok/s	808.6	+4.9%
51543	1953.1	2019.1	+3.4%
77303	1567.9	1627.1	+3.8%
103063	1787.0	1844.8	+3.2%

Decode is a dead heat, as expected — a decode step routes only batch * top_k (~8-32) rows,
so grid geometry barely matters there:

	#185	#211
decode @400 tok	149.11 tok/s	148.79
decode @1200 tok	145.56	145.34

Small but consistent: the same sign at all four prompt sizes, and it appears exactly where
your reasoning predicts (prefill, thousands of routed rows) rather than in decode. I would not
read much into any single row, but four out of four in the same direction is not noise.

Setup: FreeToken 0.1.2, RTX 4080 SUPER (sm_89, 16 GB), driver 610.57.04 / CUDA 13.3,
unsloth/gemma-4-26B-A4B-it-qat-GGUF (Q4_0, top_k = 8, 30 layers x 128 experts),
--kv-reserve-tokens 131072 --kv-cache-dtype q8_0 --max-running-requests 4, chunked prefill
at the default 8192.

One thing worth keeping from #185 if this is the version that lands: it adds
FT_MOE_VEC_LAUNCH_CHECK() after each of the 19 launches. Without it a rejected launch only
sets the error flag and gets blamed on the next unrelated CUDA call — I independently chased
this same grid.z bug into flashinfer's activation kernel before finding #186, and the
misattribution cost hours. The check is orthogonal to which axis you choose.