203family SM121LowerContext : Type 0
204constructor SM121LowerContextValue
205-- the scoreboard the program names nowhere (5 when it names all six)
206field unrestricted sm121LowerContextBarrier : Nat
207-- the least register the program never names, below its count (the count
208-- when there is none)
209field unrestricted sm121LowerContextFree : Nat
210-- the least even register that and the next are never named
211field unrestricted sm121LowerContextFreePair : Nat
212field unrestricted sm121LowerContextCount : Nat
213field unrestricted sm121LowerContextAddressing : (family SM121SharedAddressing)
214-- the wait mask naming every scoreboard a register rides on (one an
215-- instruction reading or writing registers names) and the free one: a join
216-- waits on all of them. A scoreboard only register-free instructions name
217-- -- LDGDEPBAR's, cp.async's count of groups in flight -- holds no register
218-- the schedule tracks, and a join leaves it running: the copies stay in
219-- flight across a loop's back edge, as the program's own waits (DEPBAR)
220-- intend
221field unrestricted sm121LowerContextWaitAll : NatThe compiler supplied declaration spans and resolved links from this source snapshot. This page does not assert that this file belongs to a checked closure.