Source/Packages

Accelerator.SM121.Lowering

packages/hardware/architectures/nvidia-sm121/src/Accelerator/SM121/Lowering.alpha

2,621 lines365 declarations134.1 KiBSHA-256 b7b3bbc05e9c

def · lines 404–404

sm121LowerAsyncCopyOffsetPlace

Full file
LDGSTS (cp.async.cg ... 16): the shared offset at 44 (20 bits); the high word as ptxas -arch=sm_121a writes it against -arch=sm_86 (nvdisasm -hex on the DGX Spark, 2026-09-26, research p4-hmma/cp_async*.cu): bits 70 and 74 clear, 83 set -- desc[UR4] in place of sm_86's implicit descriptor (0x...0b901c44 becomes 0x...0b981804; the byte offsets and registers are the same)
404def sm121LowerAsyncCopyOffsetPlace : Nat = (compile-time (sm121LowerPlace 44))

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.