Function ssm_ba_gates_hopper_reject 

pub fn ssm_ba_gates_hopper_reject(
    requested: bool,
    kernel_present: bool,
    m: u32,
    n: u32,
    k: u32,
    k_stride: u32,
    sm_count: u32,
) -> Option<&'static str>

Why the twin is NOT running — None means it is.

Pure, so the grammar is testable without a GPU or the process environment, and it NAMES THE GUARD THAT REFUSED. A perf path that asks to be enabled and silently is not measures as “no effect”, which is how PR #296 shipped an ldmatrix GEMM that fell back with both gates green.

The shape guards are not defensive padding:

• K % 8 != 0 would leave a tail the uint4 sweep never reads — the parent has the same limit and the same silence about it, so it is stated here;
• K_stride < K would read past the row;
• N == 0 has nothing to compute;
• the token-count floor is the decode guard described in this module’s header.