pub fn cublas_scale_layout_kmajor() -> boolExpand description
ATLAS_CUBLAS_SCALE_LAYOUT — which VEC128 activation-scale layout the
cuBLASLt block-scaled arm feeds the library.
kmajor(DEFAULT) —[K/128, ceil16(M)], tokens contiguous. What the cuBLAS manual’s “Scaling factors layouts” specifies for the B operand (“N-major for B with shape N x L”); seespark_runtime::cublaslt::scale_layoutfor the full quotes.rowmajor— the quantizer’s[M, K/128]handed over untransposed, i.e. the pre-fix reading. KEPT ONLY as a measurement control: it is what the 2026-09-11 H100 run measured at rel_rms 7.7e-2 / cosine 0.996 vs the in-tree kernel, and an operator comparing the two arms on one box should not have to check out an old commit to reproduce it.
OnceLock-cached for the same reason ffn_w8a16_only() is: the selector
runs per projection per layer per prefill and env::var_os walks the
environment block every call.