Constant STALL 

pub const STALL: &str = r#"
__global__ void lk_stall(unsigned long long ns){
  unsigned long long t0, t;
  asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(t0));
  do { __nanosleep(256); asm volatile("mov.u64 %0, %%globaltimer;" : "=l"(t)); } while(t - t0 < ns);
}
// The stall before launch `i` of the form of `rows` rows under `seed`: under 128 µs, drawn by splitmix64.
static unsigned long long lk_stall_ns(unsigned long long seed, unsigned long long rows, unsigned long long i){
  unsigned long long z = seed * 0x9e3779b97f4a7c15ULL + rows * 0xbf58476d1ce4e5b9ULL + i * 0x94d049bb133111ebULL;
  z = (z ^ (z >> 30)) * 0xbf58476d1ce4e5b9ULL;
  z = (z ^ (z >> 27)) * 0x94d049bb133111ebULL;
  return (z ^ (z >> 31)) % 128000ULL;
}
"#;

A stream held back for a span of the global timer. Under --stagger <seed> one runs before each launch of a step, its span drawn from the seed, the form and the launch, so launches its streams leave unordered run in another order under each seed.