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.
1559def sm86EncodeBarrierSynchronizeCount =
1560 (lambda unrestricted guard : (family SM86InstructionGuard) .
1561 (lambda unrestricted barrier : Byte . (lambda unrestricted threads : Nat .
1562 (lambda unrestricted control : (family SM86Control) .
1563 (sm86EncodeInstructionFields
1564 sm86OpcodeBarrierSynchronize
1565 guard
1566 control
1567 (sm86InstructionPrependField (byte-to-nat (byte 42)) (byte-to-nat (byte 12)) threads
1568 (sm86InstructionPrependField sm86InstructionNaturalFiftyFour (byte-to-nat (byte 4)) (byte-to-nat barrier)
1569 (sm86InstructionPrependField sm86InstructionNaturalEighty sm86InstructionNaturalOne 1
1570 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.