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.
1544def sm86EncodeBarrierSynchronize =
1545 (lambda unrestricted guard : (family SM86InstructionGuard) .
1546 (lambda unrestricted control : (family SM86Control) .
1547 (sm86EncodeInstructionFields
1548 sm86OpcodeBarrierSynchronize
1549 guard
1550 control
1551 (sm86InstructionPrependField sm86InstructionNaturalEighty sm86InstructionNaturalOne 1
1552 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.