AlphaGitHub
Source/Packages

Accelerator.SM86.InstructionEncoding

packages/hardware/architectures/nvidia-sm86/src/Accelerator/SM86/InstructionEncoding.alpha

2,401 lines164 declarations94.5 KiBSHA-256 92e7bf9c2555
Download exact fileView at revisionf5034773ad4bb52133

def · lines 1488–1496

sm86EncodeBarrierSynchronize

Full file
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.
1488def sm86EncodeBarrierSynchronize =
1489  (lambda unrestricted guard : (family SM86InstructionGuard) .
1490    (lambda unrestricted control : (family SM86Control) .
1491      (sm86EncodeInstructionFields
1492        sm86OpcodeBarrierSynchronize
1493        guard
1494        control
1495        (sm86InstructionPrependField sm86InstructionNaturalEighty sm86InstructionNaturalOne 1
1496          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.