Reject Q2K repack candidates whose dims truncate to uint32 kernel params (OOB GPU write)

### Summary

`ds4_repack_q2k_candidate()` (`cuda/mmq/ds4_repack.cu:465-474`) bounds only
`dims[2]` (`> UINT32_MAX` rejected at line 467). `dims[0]` and `dims[1]` are
unbounded `uint64_t` values read straight from the GGUF tensor table
(`ds4_repack_collect_catalog`, the fully-checked parser).

The artifact is sized with full 64-bit math:

```c
// ds4_repack.cu:1030-1040
const uint64_t nb_row = t.dims[0] / 256u;
const uint64_t npair = nblk / 2u; // nblk = t.bytes/84, u64
const uint64_t art_bytes = dm_bytes + sc_bytes + npair * 128u; // cudaMalloc
```

but the kernel launch narrows both to `uint32_t` parameters:

```c
// ds4_repack.cu:1076-1078
repack_q2_k_aligned_kernel<<<...>>>(
dm2, sc4, qs2, scratch, blk0, cblk,
(uint32_t)nb_row, // truncates dims[0]/256
(uint32_t)t.dims[1]); // truncates dims[1]
```

The kernel (`ds4_repack.cu:368-400`) reconstructs the pair-block index from
the *truncated* values:

```c
const uint32_t b = (uint32_t)(g % nb_row);
const uint32_t r = (uint32_t)((g / nb_row) % nrows);
const uint64_t e = g / ((uint64_t)nb_row * nrows);
const uint64_t pblk = ((uint64_t)e * (nrows/2u) + r/2u) * nb_row + b;
...
qs2[pblk * 32ull + (uint64_t)p * 2ull + parity] = w;
sc4[pblk * 8ull + ...] = w;
dm2[pblk * 2ull + parity] = w;
```

Concrete example: `dims[1] = 2^32 + 2`, `dims[0] = 256` (`nb_row=1`),
`dims[2] = 1` — the file-bounds and geometry checks (`t.bytes ==
nb_row*dims[1]*dims[2]*84`) all pass, but the kernel sees `nrows = 2` while `e`
runs to `≈ dims[1]*dims[2]`, so `pblk` reaches `≈ dims[1]` against an artifact
holding only `npair ≈ dims[1]/2` pair-blocks. Result: `qs2`/`sc4`/`dm2` write
up to ~2x past the end of the `cudaMalloc`ed artifact — out-of-bounds **GPU
writes** with file-supplied 32-bit word values, corrupting adjacent device
allocations. The same effect is reachable via `nb_row` truncation
(`dims[0] ≥ 256 * 2^32`).

Reachability chain: engine CLI `-m model.gguf` → `ds4.c` load →
`ds4_gpu_build_derived_artifacts` (`ds4_cuda.cu:4349`) →
`ds4_repack_build_q2k_aligned`. Gated on an **integrated CUDA GPU**
(`ds4_cuda.cu:4360-4365`) and the q2k-aligned repack being enabled.

Triggering requires a tensor with `t.bytes = nb_row * dims[1] * dims[2] * 84 ≥
2^32 * 84 ≈ 361 GB` — large, but not implausible for the very large DeepSeek
checkpoints this engine targets (multi-hundred-GB Q2_K expert stacks), and a
sparse file satisfies every bounds check in the catalog parser without
occupying disk.

### Fix

Reject the truncating dims at candidate time (this PR):

```c
if (t.dims[0] / 256u > UINT32_MAX || t.dims[1] > UINT32_MAX) return false;
```

(alternatively, widen the kernel parameters to `uint64_t`; the candidate bound
is the smaller change and no real tensor needs `nb_row` or `nrows` above
2^32).

### Impact

Out-of-bounds device-memory write from a crafted model file on the
q2k-aligned repack path. No host ASan trace is possible for a CUDA-side write;
a `compute-sanitizer` run on a sparse 361 GB logical GGUF on an iGPU would
show the OOB store. Treat as hardening for the same trust boundary as #542 /
#543 / #823: model files people download and run.