node-8gpu.vx

//===- node-8gpu.vx - 8-GPU node (NVLink / PCIe / IB) -----------*- Vx -*-===//
//
// Part of the Vx Project, under the Apache License v2.0 with LLVM Exceptions.
// See LICENSE for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
// An 8-GPU HGX-class node: one GPU's local hierarchy plus the interconnect a
// multi-GPU placement crosses. Use as `vxc --machine fleet/node-8gpu.vx prog.vx`
// for a program whose tensors move *between* GPUs or off-node; the per-SKU
// files model a single GPU and are the right choice otherwise.
//
// NOTE: figures transcribed from vendor specs, NOT re-verified against the
// datasheet PDF. Confirm before publication; see fleet/README.md.
//
// This file re-declares the H100 SXM local hierarchy rather than composing with
// fleet/h100-sxm.vx, because `--machine` takes a single file (#281) and two
// files declaring `Memory HBM` is E6012, by design. Declaration import (#224)
// would let a node model compose with a SKU model; until then, a node variant
// per SKU means duplicating three declarations. Deliberate: silent shadowing
// would be much worse than duplication.

// --- One GPU's local hierarchy (H100 SXM; see fleet/h100-sxm.vx) ------------

// spec: NVIDIA H100 product page, SXM column (verified 2026-08-03) -- 80 GB
// HBM3 at 3.35 TB/s. Keep in sync with fleet/h100-sxm.vx.
Memory HBM {
  capacity: 80 GiB, bandwidth: 3.35 TB/s, managed: explicit, scope: device
}

// spec: NVIDIA Hopper Architecture Whitepaper -- 50 MB L2 cache.
Memory L2 {
  within: Memory::HBM, capacity: 50 MiB, bandwidth: 12 TB/s, managed: cached, scope: device
}

// spec: NVIDIA Hopper Architecture Whitepaper -- 228 KB shared memory per SM.
Memory SMEM {
  // replicas: measured: probe_edges reported SMs=132 on two rented H100 SXM boxes
  // clock: measured: nvidia-smi clocks.max.sm on two rented H100 SXM boxes, 2026-08-08
  within: Memory::L2, capacity: 228 KiB, bandwidth: 128 B/cyc, clock: 1.98 GHz, granule: 1 KiB, managed: explicit, scope: sm, replicas: 132
}

// --- The rest of the node ---------------------------------------------------

// The aggregate HBM of the other seven GPUs, reachable over NVLink. Modelled as
// one space rather than seven: Vx has no per-instance memory accounting (out of
// scope per #290), so a tensor sharded across peers is checked against the total
// peer budget. That is the honest bound the declarations can express -- it will
// admit a placement that no *single* peer could hold, which is why per-instance
// accounting is future work rather than something quietly approximated here.
//
// spec: derived -- 7 peers x 80 GB (verified per-GPU figure above).
Memory PEER_HBM {
  capacity: 560 GiB, bandwidth: 3.35 TB/s, managed: explicit, scope: device
}

// The node's transfer topology. Three physical tiers, and the whole point of
// declaring them: a placement that crosses IB costs ~40x one that stays on
// NVLink, and the admission verdict should say so rather than leaving it to a
// runtime surprise.
//
// spec: NVIDIA H100 Datasheet -- NVLink 4: 900 GB/s bidirectional per GPU,
//       modelled as 450 GB/s unidirectional (see fleet/README.md on halving).
// spec: PCIe Gen5 x16 -- 64 GB/s unidirectional (128 GB/s bidirectional).
// spec: NVIDIA ConnectX-7 / InfiniBand NDR -- 400 Gb/s = 50 GB/s per port.
Topology Node {
  // spec: Eight NVIDIA parts; the node's compute ISA is theirs.
  arch: nvptx64,
  memory: Memory::HBM,
  visible: [Memory::HBM, Memory::L2, Memory::SMEM, Memory::PEER_HBM],
  // spec: PCIe Gen5 x16 = 63 GB/s per direction (32 GT/s x 16 lanes, 128b/130b).
  // UNVERIFIED on this SKU; a Grace/C2C-attached part reaches the host at ~450 GB/s instead,
  // a 7x difference in every host-seam prediction. Settling it is an open experiment (utils/memalg).
  // The on-die hops below declare no cost -- their `bandwidth:` figures derive it (E6013).
  transfer Memory::CPU_DRAM -> Memory::HBM : 63 GB/s,
  transfer Memory::HBM -> Memory::L2,
  // spec: Ampere cp.async (PTX ISA 8.3 section 9.7.8.24) -- same engine the SKU file
  // fleet/a100-80.vx declares; node files repeat the SKU hierarchy verbatim (#224).
  transfer Memory::L2 -> Memory::SMEM copy_engine,
  // The three inter-device edges. Each carries the bandwidth of the physical link it
  // represents -- the same figures the `spec:` block above already cites, moved out of
  // prose and into the declaration so the compiler can actually price them.
  //
  // They used to carry unitless weights (60, 200, 2400). Those are relative latencies for
  // route selection and are NOT times, so nothing derived a cost from them: every one of
  // these edges reported `derived_cost: null`. The compiler would route a KV cache from
  // GPU0 to GPU7, or out through the NIC to another machine, and predict nothing about how
  // long it took -- which left the between-machine measurement with no predicted column
  // against. An unpriced edge cannot be validated by any amount of machine time.
  //
  // Worth recording: the old weights did not even encode the right ORDER. 450/63/50 GB/s
  // implies per-byte costs in the ratio 1 : 7.1 : 9, while the weights said 1 : 3.3 : 40.
  //
  // spec: NVIDIA H100 Datasheet -- NVLink 4 at 900 GB/s bidirectional per GPU, modelled as
  // 450 GB/s unidirectional (see fleet/README.md on halving). UNVERIFIED on any rented
  // pod: `nvidia-smi topo -m` reporting NV18 between a pair confirms full NVLink, and
  // anything less means this figure is optimistic in the direction a bound must not be.
  transfer Memory::HBM -> Memory::PEER_HBM : 450 GB/s,
  transfer Memory::PEER_HBM -> Memory::HBM : 450 GB/s,
  // spec: PCIe Gen5 x16 = 63 GB/s per direction. The GPU reaches the NIC across the same
  // root complex it reaches the host across, so this is the host-seam link rate again, not
  // the fabric's. UNVERIFIED, and it shares the Grace/C2C uncertainty noted above.
  transfer Memory::HBM -> Memory::NIC_RAM : 63 GB/s,
  // spec: NVIDIA ConnectX-7 / InfiniBand NDR -- 400 Gb/s = 50 GB/s per port. This is the
  // between-machine edge the fleet campaign exists to measure, and it is deliberately the *port* rate
  // rather than any aggregate: a placement crossing it is one flow on one port.
  transfer Memory::NIC_RAM -> Memory::Remote_HBM : 50 GB/s
}