Function assign_stalls_by_latency 

pub fn assign_stalls_by_latency(cubin_path: &Path) -> Option<Vec<u8>>

CONTROL-CODE ASSIGNER (correctness-first, stalls). Recompute every instruction’s STALL count from our own latency model — stall = fixed_latency(class) — which is provably >= the cycles any fixed-latency consumer needs (the immediate next instruction already issues stall cycles later, so later consumers have even more). We KEEP ptxas’s write/read barriers + wait masks (the variable-latency memory / cp.async tracking, the fragile part) untouched. If the kernel stays numerically correct with OUR stalls, the stall model is validated on-device — the first half of a working control-code assigner. Returns the patched cubin bytes.