LDGSTS.E.BYPASS.128 [address + offset], [source.64 + sourceOffset], as
ptxas encodes cp.async.cg ... 16 for sm_86 (nvdisasm -hex on the DGX
Spark, 2026-09-26; research p4-hmma/cp_async*.cu): the shared address
register at 16, the global pair at 24, the global offset at 32 (signed 12
bits in the silicon; this encoder admits 0..0x7ff, bit 43 clear), the
shared offset at 44 (20 bits), and bits 64..95 0x0b901c44: BYPASS (bit 81
clear), the global-memory descriptor UR4 with its enable (0x44) -- the
descriptor LDG's bits 32..39 name too. The compiler's native encoder
(RuntimeMachine) is its twin.
1577def sm86EncodeLoadGlobalToShared =
1578 (lambda unrestricted guard : (family SM86InstructionGuard) .
1579 (lambda unrestricted address : (family SM86Register) .
1580 (lambda unrestricted offset : (family SM86Unsigned32) .
1581 (lambda unrestricted source : (family SM86Register) .
1582 (lambda unrestricted sourceOffset : (family SM86Unsigned32) .
1583 (lambda unrestricted control : (family SM86Control) .
1584 (sm86EncodeInstructionFields
1585 sm86OpcodeLoadGlobalToShared
1586 guard
1587 control
1588 (sm86InstructionPrependField sm86InstructionNaturalSixteen sm86InstructionNaturalEight
1589 (sm86RegisterNatural address)
1590 (sm86InstructionPrependField sm86InstructionNaturalTwentyFour sm86InstructionNaturalEight
1591 (sm86RegisterNatural source)
1592 (sm86InstructionPrependField sm86InstructionNaturalThirtyTwo (byte-to-nat (byte 11))
1593 (sm86Unsigned32Natural sourceOffset)
1594 (sm86InstructionPrependField (byte-to-nat (byte 44)) (byte-to-nat (byte 20))
1595 (sm86Unsigned32Natural offset)
1596 (sm86InstructionOneWord32Field
1597 sm86InstructionNaturalSixtyFour
1598 (constructor SM86Unsigned32 SM86Unsigned32Value (byte 68) (byte 28) (byte 144) (byte 11))))))))))))))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.