BAR.SYNC.DEFER_BLOCKING b, n: BAR.SYNC 0x0's word with the barrier in
bits 54..57 and the thread count in bits 42..53 (read off nvdisasm -hex of
ptxas sm_86 code on the RTX 3070 pod, 2026-10-09: bar.sync 1, 128 is
0x0042000000007b1d, bar.sync 15, 96 0x03c1800000007b1d; count 0 is the
whole CTA, bar.sync 0's word). RuntimeMachine's native encoder is its twin.
1639def sm86EncodeBarrierSynchronizeCount =
1640 (lambda unrestricted guard : (family SM86InstructionGuard) .
1641 (lambda unrestricted barrier : Byte . (lambda unrestricted threads : Nat .
1642 (lambda unrestricted control : (family SM86Control) .
1643 (sm86EncodeInstructionFields
1644 sm86OpcodeBarrierSynchronize
1645 guard
1646 control
1647 (sm86InstructionPrependField (byte-to-nat (byte 42)) (byte-to-nat (byte 12)) threads
1648 (sm86InstructionPrependField sm86InstructionNaturalFiftyFour (byte-to-nat (byte 4)) (byte-to-nat barrier)
1649 (sm86InstructionPrependField sm86InstructionNaturalEighty sm86InstructionNaturalOne 1
1650 sm86InstructionFieldsEmpty))))))))The compiler supplied declaration spans and resolved links from this source snapshot. This page does not assert that this file belongs to a checked closure.