[msprengholz]

The problem

Two things cost time on a Maxwell card (compute capability 5.x), and both are measured here on a GTX TITAN X:

1. The dense projections' FP16-operand GEMMs. They go through cuBLAS's Maxwell kernels
   (maxwell_sgemm_fp16_32x128, _64x64), and there is no tensor-core path for that call on such a card. They are
   the second biggest compute item of a prompt read: 1.93 s of the 9.8 s the GPU spends reading a 1,024-token
   prompt. As on the P40 (compute capability 6.1), the same work as cublasSgemm on FP32 copies of the FP16
   operands is faster - 1.18-1.65x for the big shapes on this card - and bit for bit the same.
2. The dense weights' MMVQ. The engine has an interleaved small-batch matvec (a warp takes 2 or 4 rows, so the
   weights and the activations are reused across rows) with a table per architecture - a P100, a P40, a V100 - but
   il_arch_ok stopped at compute capability 6, so a Maxwell card never used it. That is where the gain is on a
   card without __dp4a: the kernels are integer-bound, so the reuse pays.

What this changes

• src/prefill/gemm.cu: the route (STRATA_DENSE_F32=1, opt-in, unchanged as a switch) was gated to compute
  capability 6.x; it now also accepts 5.x, the cards whose cuBLAS has no tensor-core path for this call.
• src/kernels/cuda/native_mmvq.cu: the interleaved matvec's table for this architecture, read off
  mmvq_il_parity --bench --emit-table on the card (the repo's documented procedure), with its own opt-in switch
  (STRATA_SM50_TABLE=1).
• The default path is untouched in both cases: without the environment variable, nothing changes.
• The commit also carries the profile that found them and the tools that made it: an Nsight Systems trace of a real
  run (every kernel with its time and count, the copies, the idle gaps), the tool that turns a window of the trace
  into the tables, the A/B harness that measures a candidate, and the ranked candidate list with every candidate's
  number - the nine that measured slower and were dropped included.
• docs/OLDER_GPUS.md and docs/GYRO.md: what the profile says, with the numbers, the limits of the card, what was
  not measured, and the ranked next steps in the order the profile gives them.

The first commit of this branch is the previous pull request (the Maxwell support the experimental CUDA 12 build
needs); the changes themselves are the two commits after it.

The profile, in one table (the TITAN X, a real 1,024-token prompt and 201 decoded tokens)

window	length	GPU busy	GPU idle	real compute	PCIe copies
prompt	27.0 s	9.77 s (36 %)	17.2 s (64 %)	9.8 s	0.3 s
decode	53.0 s	49.2 s (93 %)	3.8 s (7 %)	9.0 s (17 %)	0.29 s

The decode's 40.8 s of "busy" is one thread spinning on the doorbell (wait_flag_ge_kernel) while the host waits for
the expert data off the SSD: the experts are 18 GiB, the card has 12 GiB of VRAM and the PC 15.7 GiB of RAM. The
kernels are 9.0 s: the trellis expert kernels 3.30 s, the dense weights' MMVQ 1.72 s, the GDN hyper-connection
1.01 s, the head and the embedding 0.43 s, the attention 0.11 s, the MTP draft head 0.02 s. Every one of them is
named with its source file and its job in the report, and the four things this card does not have are measured on it
(no __dp4a and no vdp4a; no cp.async - the binary has zero CP.ASYNC, so the staged hyper-connection read is
the "split" variant only, although the start-up message claims otherwise; no tensor cores; 48 KB of shared memory per
block, which the GDN's staged down projection uses 40 KB of, so an sm_50 SM holds two of them where the registers
would allow three).

Validation

Environment: GeForce GTX TITAN X (sm_52, 12,288 MiB, 24 SMs), driver 580.178.04, Linux 6.2 (CachyOS), CUDA 12.9 at
/opt/cuda, host compiler g++-14. Engine: the experimental CUDA 12 build
(-DCMAKE_CUDA_ARCHITECTURES=50 -DSTRATA_EXPERIMENTAL_SM60=ON), engine 0.1.46.1. Model: Gyro-S (TQ1_0) with the
repo's strata-gyro-s.json (MTP on, --spec 4, --expert-cache auto, --resident-budget-gib 6, --kv int8).

Passed

• gemm_f32_parity --quick and --bench (the repo's own tool, built with this CUDA 12.9): 5,201 chunk sizes bitwise
  equal, and the route 1.18-1.65x faster for the big shapes (qkv 10,240x2,560 at T 7,040: 166.5 -> 120.0 ms; query
  12,288x2,560 at T 7,041: 203.4 -> 149.8; gate 6,144x2,560: 101.0 -> 67.8; out 2,560x6,144: 127.5 -> 77.1), neutral
  from T 560 to 2,048 (0.97-1.06x).
• mmvq_il_parity --bench --emit-table (the repo's own tool): every case, T 1-8, rows 1/2/4 and the table, bitwise.
• Through the server, the prompt read: 91.5 -> 95.7 tok/s with STRATA_DENSE_F32=1 (the medians of three repeats, the
  page cache evicted before each run; the two base runs of the same build agree to 0.2 % on that number).
• With STRATA_SM50_TABLE=1, the engine's own stage timers (STRATA_VERIFY_PROFILE=1): the dense projections'
  stage 8.92 -> 7.76 ms per window (-13 %, better in 3 of 3 repeats), the VRAM expert hits 36.7 -> 34.5 ms (-6 %). The
  request's tok/s does not move, because the GPU is idle 80 % of the time on this card (it waits 190-310 ms per
  window for the experts while it computes about 40 ms) - the change is for the cards where the experts fit.
• The answers are the same, character for character, in both cases (parity-check.py, the same prompt through the
  server with and without the switch, greedy; parity-answers-dense_f32.json, parity-answers-sm50-table.json).
• cmake --build build-cuda12 -j 8 - clean.

Failed

• ctest: ple_parity and platform_memory_test fail in this environment, and they fail the same way without these
  changes: ple_parity wants a Q2_0 model file that is not on this PC, and platform_memory_test wants to lock more
  memory than this shell's ulimit -l (8 MB) allows. The other 66 tests pass.

Not run

• The HIP and the SYCL engines. The changed lines in src/prefill/gemm.cu are inside
  #if defined(STRATA_EXPERIMENTAL_SM60) (lines 738-794), and that macro is defined only in a CUDA build
  (CMakeLists.txt:216-219); the SYCL build (sycl/CMakeLists.txt:265) compiles the same file without it. The
  interleaved matvec's table in src/kernels/cuda/native_mmvq.cu is CUDA-only source, but that file is in the
  kernel library the HIP build compiles too, and this PC has no ROCm and no Intel toolchain, so I could not build
  them to show it. This pull request is a draft for that reason: if you can build the HIP or the SYCL engine on a
  machine that has the toolchain, that is the check I am missing.
• The CUDA 13 release engine (only CUDA 12.9 is installed here). The changed lines in gemm.cu are not compiled
  there (the macro is off by default), and the table's switch is off by default too.
• Any other card. Every number here is this one TITAN X on this PC.

The limits, said plainly

• On this card the GPU computes 9.0 s of a 53 s decode and waits 40.8 s for the expert data. The kernel work matters
  where the model fits in the card's memory; here the data path is the limit, and the profile says so.
• Both switches are opt-in. The GEMM route's equality is a property of the cuBLAS build, not of the code: run
  gemm_f32_parity with your toolkit's library first. The matvec table is one card's measurement, so it is opt-in
  too.

[msprengholz]

Closing this one: the kernel optimizations (the FP32 dense-GEMM route for cc 5.x and the interleaved-matvec table for sm_52) were not asked for. The Maxwell support itself stays in #3595, which is the only part that is needed for the engine to run on a sm_50 card.