#124 Pascal (sm_60): bring the engine up on compute capability 6.0

Beschreibung

## What this does

Runs the engine on compute capability 6.0 (Pascal, GP100-class) cards. Two build gates stand in the way, plus
one attribute the card does not honour - and none of them needs the kernels rewritten:

1. `CMakeLists.txt` refuses to configure below sm_75;
2. `__dp4a` (6.1+) appears 29 times in the i-quant kernels (`native_mmvq.cu`, `s2_expert_grouped.cu`,
   `iq_kernels.cu`), and `__nanosleep` (7.0+) in the three doorbell waits (`elementwise.cu`,
   `verify_kernels.cu`);
3. the Turing port sizes `fused_gr_read_multi`'s slice from `cudaDevAttrMaxSharedMemoryPerBlockOptin`, which on
   an sm_60 answers 65536 - while that card's launch validator enforces the 49152 B **per-block** limit. The
   capacity comes out 6 tokens and every 6- or 8-token window fails the launch with `invalid argument`.

## The changes

* **`include/strata/kernels/dp4a.hpp` (new)** - `STRATA_DP4A` and `strata_spin_pause()`. The dp4a fallback is
  llama.cpp's own (`ggml/src/ggml-cuda/common.cuh`): both operands read as SIGNED bytes, accumulated into int32,
  which is what `__dp4a` does for exactly these call sites - the kernels here are transcribed from
  `vecdotq.cuh`/`mmvq.cu`, so the fallback is **bit-exact**, not approximate. The nanosleep sites only pace
  single-thread doorbell waits, so on sm_6x the loop simply spins.
* **`CMakeLists.txt`** - `-DSTRATA_EXPERIMENTAL_SM60=ON` relaxes the architecture gate, and the refusal message
  names it.
* **`fused_gr.cu`** - below sm_70 the slice capacity comes from `cudaDevAttrMaxSharedMemoryPerBlock`
  (49152 B -> 4 tokens) instead of the opt-in, which is the limit the launch validator actually enforces.

## What the HIP review asked for

* **the HIP tile stays.** `fused_gr.cu`'s `#if defined(__HIPCC__)` tile (1280 against gfx1100's 64 KiB LDS) is
  upstream's, untouched. The 0.1.21 form of this patch carried a `kFusedGrTile` constant and a window clamp for
  the same shared-memory problem; both are gone, so no tile value is overridden anywhere.
* **the HIP sleep stays.** `strata_spin_pause()` is a no-op only under `__CUDA_ARCH__ < 700`; every other
  target - HIP included - calls `__nanosleep`, so AMD keeps `hip_compat`'s
  `__builtin_amdgcn_s_sleep(1)` and does not busy-spin. `STRATA_DP4A` on HIP is likewise `hip_compat`'s
  `__dp4a` (the gfx1100 `sudot4` path), because `__CUDA_ARCH__` is undefined there.
* **`fused_gr_max_window` is gone**, and what remains of the shared-memory work in `fused_gr.cu` is HIP-guarded:
  the HIP branch keeps the exact expression it had (`optin > 0 ? optin : 48 * 1024`).
* **the UTF-8 BOM stays** in `CMakeLists.txt`.

## Rebased onto 0.1.27

0.1.21 -> 0.1.24 -> 0.1.27. Two things came out of the last rebase:

* `STRATA_EXPERIMENTAL_SM75` is gone (Turing is supported now), so the gate reads `_base LESS 75 AND NOT
  STRATA_EXPERIMENTAL_SM60` and keeps the current wording.
* `fused_gr_read_multi` now slices an oversized batch itself. That is a better fix than the window clamp this
  patch used to carry, so the clamp, its `fused_gr_max_window` helper and the two `generate.cpp` call sites were
  dropped in favour of correcting the one number the slicing is computed from. The full window now runs here:

  ```
  strata verify: window up to 6 tokens, 69.1 MiB of device buffers
  strata verify: captured the 6-token window (upload no error, sync no error)
  ```

## Behaviour on supported cards

A no-op, by construction: from sm_70 the opt-in is honoured, so the slice capacity is the value it already was
(and the clamp never applied there anyway); `STRATA_DP4A` and `strata_spin_pause()` compile to `__dp4a` and
`__nanosleep` exactly as before; a HIP build takes neither branch. Only the 6.0 path was exercised on hardware,
so the sm_70+/HIP paths here are reasoned rather than measured.

## Notes for a Pascal build

* Needs a CUDA release that still targets Pascal: **CUDA 13 removed it, 12.x has `compute_60`**.
* `CMAKE_CUDA_STANDARD` has to stay **17** for that combination: nvcc 12.9 refuses `-std=c++20` when the host
  compiler is VS2019's `cl` and silently falls back to C++14, which rejects the nested
  `namespace strata::kernels {` declarations everywhere.
* `CMAKE_CUDA_RUNTIME_LIBRARY=Shared` if the local CUDA keeps `cudart.lib` as a DLL import library (CUDA 13 made
  it static); the default `-cudart static` then collides with the project's `CUDA::cudart` (LNK2005).
* A Pascal build also loads `cublas64_12.dll` at runtime (the prefill GEMM). CUDA 12.9's cuBLAS does serve the
  BF16 `cublasGemmEx` those kernels ask for on sm_60 - verified by reading prompts on hardware, not assumed.
* The tensor-core paths stay off by their own guards: `qsa_select.cu` and `native_qsa_score.cu` compile their
  `mma`/`ldmatrix` blocks only at `__CUDA_ARCH__ >= 800` (with the portable fp32-FMA fallback the Turing port
  added) and `qsa_prompt_attn.cu` keeps the older kernel. Nothing extra was needed for them.

## Measured (6.0 card, GP100-class, four dies through the layer split, IQ3_XXS)

883-token prompt, 256 generated tokens, the shipped `--spec 4` and `--expert-cache auto`. Both rows were
measured in the same session on the same machine with the same load, so the difference between them is the
engine:

| | prefill (883 tok) | decode | expert tier |
|---|---:|---:|---|
| four sm_60 dies, **this tree (0.1.27)** | 110 tok/s | **21.3 tok/s** (warm median) | 24576/24576 resident, 100% hit |
| four sm_60 dies, 0.1.21 base, same prompt | 110.6 tok/s | 17.9 tok/s (warm median) | 24576/24576 resident, 100% hit |
| one RTX 3090 Ti, release engine, same model | 305 tok/s | 79 tok/s | 8000 slots |

The 0.1.21 row is this port's original base, whose fused GR read could not launch a 6-token window at all and
therefore ran with the MTP window at 2 and the verify window at 4. With the slicing corrected, the windows are
the shipped ones and decode is **~19% faster** (warm prefill ~7% faster); cold prefill is unchanged. Earlier in
this PR I reported a decode **drop** between the 0.1.21 and 0.1.24 bases - that was measured with the clamp
still in place, and it no longer reproduces.

So this is about **fitting, not speed**: with every expert resident the wall is Pascal's per-layer arithmetic
(no `__dp4a`, 140 W caps), not capacity and not memory. If the maintainer would rather not carry the extra
`#if` branches for a configuration this slow, that is a fair call - the alternative is that these cards cannot
run the engine at all.