Function tgemm_kernel 

pub fn tgemm_kernel(gpu: &dyn GpuBackend) -> KernelHandle

Resolve the N128/M64 tile GEMM, preferring the 3-deep weight-pipeline variant. ON by default; ATLAS_NO_TGEMM_PIPELINE3 (presence — =0 is NOT “off”) falls back to the 2-stage parent. Falls back automatically on any target that does not ship _p3.

Same mechanism as k64_kernel: the parent drains its cp.async group before the dequant phase, which only a co-resident CTA can cover. This kernel’s live shapes — ssm_qkvz (128 CTAs) and the fused QKV (112) — sit in the exposed band of the grid.x-vs-efficiency curve. Bit-identical.