Module w4a16_gemv_tiers 

SSOT for the narrow w4a16_gemv_batch{M} tier family (M = 4..8).

Why this module exists

w4a16_gemv_batchm_impl<MAX_M> sizes acc[], s_vl[] and smem[] by MAX_M, and because its row loop is #pragma unrolled, MAX_M also sizes the CODE: at 80 static SASS instructions per row on sm_121f, the MAX_M=8 tier is 760 instructions against MAX_M=4’s 440. The t >= M guard skips a dead row’s WORK at run time but not its instructions, and the template is issue-bound (not DRAM-bound) at these M — so running M=5 on the MAX_M=8 tier pays for three rows that are not there.

Measured on the real 27B qkv/o shape (N=5120 K=5120, cold-cycled weights, 273 GB/s peak) before the exact-M tiers existed:

tier	M	time	eff BW	% peak
batch4	4	70.5 us	209.2 GB/s	76.6%
batch8	4	74.8 us	197.3 GB/s	72.3%
batch8	5	89.0 us	165.7 GB/s	60.7%
batch8	6	91.5 us	161.2 GB/s	59.0%
batch8	8	106.4 us	138.5 GB/s	50.7%

The batch8 @ M=4 row is the argument: same rows, same weight stream, +6.1% for nothing. It is NOT occupancy — batch4 lands on 48 registers / 5 CTA per SM with no __launch_bounds__ at all, which is exactly what the pragma pins batch8 to.

Why a shared table instead of a match per call site

Before this module, FIVE structs each carried a w4a16_gemv_batch4_k / w4a16_gemv_batch8_k pair and each re-derived 1..=4 => batch4, 5..=8 => batch8 inline (dense_ffn, qwen3_ssm x2, qwen3_attention, mtp_head, model). Adding three tiers would have meant five more copies of a widening decision. The decision now lives here once, as a PURE function over which tiers the loaded target actually resolved.

Kill switch

ATLAS_NO_GEMV_EXACT_M_TIERS=1 (presence-checked per the house convention; =0 is NOT off) hides widths 5/6/7 from the decision, restoring exactly the batch4/batch8 dispatch that shipped before them. It does not unload the kernels — it only removes them from selection, so an A/B needs no rebuild.

Structs

W4a16BatchmTiers

Resolved handles for the narrow w4a16_gemv_batch{M} family.

Constants

W4A16_BATCHM_WIDTHS

Tier widths in this family, narrowest first. Parallel to the handles field of W4a16BatchmTiers and to the present array of select_tier.

Functions

exact_m_tiers_enabled

Are the exact-M tiers (5/6/7) allowed in the dispatch decision?

select_tier

PURE tier decision: index into W4A16_BATCHM_WIDTHS of the narrowest tier that both COVERS m rows and is present in the loaded target, or None when this family cannot serve m (caller falls back to the tile GEMMs / the wide tiers).