BAR.SYNC.DEFER_BLOCKING 0x0: bit 80 set, as ptxas encodes bar.sync for
sm_86 and sm_121 (read off nvdisasm -hex on the DGX Spark, 2026-09-26:
high word 0x...0001_0000); without it nvdisasm reads plain BAR.SYNC. The
compiler's native encoder (RuntimeMachine) is its twin.
1624def sm86EncodeBarrierSynchronize =
1625 (lambda unrestricted guard : (family SM86InstructionGuard) .
1626 (lambda unrestricted control : (family SM86Control) .
1627 (sm86EncodeInstructionFields
1628 sm86OpcodeBarrierSynchronize
1629 guard
1630 control
1631 (sm86InstructionPrependField sm86InstructionNaturalEighty sm86InstructionNaturalOne 1
1632 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.