Figure: GPU integration with CXL.mem — memory mapping, hint interface, driver-to-firmware translation and an unmeasured fault-path component budget. Four endpoints contribute 1 TB of tier-2 DRAM and 16 TB of tier-3 NVMe; the GPU virtual address space holds 192 GB of local HBM at 0x0000 to 0x4FFF plus CXL.mem regions of 256 GB, 256 GB and 512 GB, while the model's weights stay HBM-resident and never cross the tier boundary. Those capacities are arithmetic over stated device sizes and stand. Allocation hints — data type, access pattern, consistency, priority, expected size, plus scoring parameters including an LRFU alpha and a GQA head count of 8 — travel from the application through the driver into endpoint firmware in five stages over a CXL.io mailbox. The scoring rule is LRFU, Lee et al., ACM SIGMETRICS 1999: established prior art, with alpha as its documented LRU-to-LFU knob. Alpha = 0.2 was never validated and is withdrawn; the simulated optimum is about 0.05, and alpha is non-monotonic with no convergence onto LFU as alpha approaches zero. The head count is a scoring input only: every head of every layer is required for every token, so evicting a head produces wrong outputs, not a cache miss, and the placed, moved and evicted unit is a (sequence, token block) of 16 tokens, 5.00 mebibytes. There is no RoPE-aware prefetch hint, because RoPE is a rotation applied to Q and K at compute time and carries no information about which blocks are needed next. The fault-path panel decomposes a hypothetical load-store access into 15 nanoseconds for request issue, 20 for the PCIe Gen5 by 16 PHY, 25 for CXL switch traversal, 40 for endpoint controller decode, 80 for the DDR5-5600 array access and 20 for the return path, totalling 200 nanoseconds. That total is an unmeasured component budget, not a result: no CXL hardware was measured anywhere in this package, and the 200 nanosecond figure is demoted to an assumption inside a 32 to 121 gigabytes per second effective-bandwidth range. What the package models instead is the transfer of one 5.00 mebibyte block of 16 tokens, which takes 164, 82 or 43 microseconds at an assumed 32, 64 or 121 gigabytes per second. The 65 times ratio formerly claimed here — 13.0 microseconds of driver-mediated PCIe DMA divided by 200 nanoseconds of hypothetical load-store access — is retracted as a headline: the two paths are not the same operation and neither was measured. The comparison that survives is move versus regenerate: moving that block costs 164, 82 or 43 microseconds, while regenerating the same 16 tokens by forward compute costs 9, 11 or 18 milliseconds at 50, 40 or 25 percent model FLOPs utilisation — two to three orders of magnitude, with no precise multiple claimed, and regenerating an arbitrary interior block is at least that expensive because interior KV is not independently reproducible.

Section 9

GPU Integration

Memory mapping, hint interfaces, and fault handling for CXL.mem

9.1 Memory Mapping

CXL.mem regions appear in the GPU's unified virtual address space via PCIe BAR mapping with CXL.mem bridging. The GPU accesses remote memory using standard load/store semantics. Four endpoints contribute 1 TB of Tier-2 DRAM and 16 TB of Tier-3 flash. The model's 140 GB of weights stay HBM-resident and never cross CXL.

Target-platform capability, not a generic CXL guarantee. Whether a GPU can issue ordinary load/store directly against a CXL.mem region — rather than going through a driver-mediated copy — depends on the specific GPU, its MMU and BAR/HDM decoder support, the platform firmware and the CXL topology. It is assumed here as a property of the target platform. It is not a demonstrated capability of CXL in general, nor a published capability of any particular B200 system.

GPU Virtual Address Space
CXL.mem
Bridge
load/store
semantics
Endpoint Physical Memory
GPU Kernel Access Pattern — CONCEPTUAL PSEUDOCODE

Illustrative only. This fragment is not compilable CUDA and does not demonstrate a capability of CXL, of CUDA, or of any shipping B200 platform. It sketches what a kernel would look like if the target platform exposes a CXL.mem region as GPU-addressable mapped memory with load/store semantics. Real systems require explicit region registration, a platform-specific mapping path, and driver and firmware support; a plain pointer dereference into a remote CXL region is not portable and is not guaranteed by the CXL specification.

// GPU kernel accessing CXL.mem KV-cache
__global__ void attention_kernel(float* kv_cache) {
  // Direct load from CXL.mem address
  float4 kv = *(float4*)(kv_cache + offset);
  
  // GPU MMU translates VA → CXL.mem address
  // CXL bridge handles coherent access
  // No per-access DMA setup and no per-access driver
  // involvement in steady state; the control plane, region
  // mapping, fault handling and firmware still run on the host.
}
9.2 Hint Interface

Extended allocation API communicates memory characteristics to the endpoint, so PLACE can score and site each block — including the identity key that lets a block's policy state survive eviction. Everything here except that persistence mechanism is Design / unvalidated.

struct cxl_alloc_hints
enum data_type KV_CACHE, WEIGHTS, ACTIVATIONS, SCRATCH
enum access_pattern SEQUENTIAL, RANDOM, STRIDED, BROADCAST
enum consistency RELAXED, ACQUIRE_RELEASE, SEQ_CST
u32 batch_window Lookahead distance in cache lines
u32 priority Eviction priority (0 = evict first)
u64 expected_size Allocation size hint for placement
struct placement_params
u32 window_size Blocks moved per batch (MOVE stage) — Design / unvalidated
u32 stride Access stride for strided patterns
u128 block_identity_key Survives eviction. Score is restored on re-admission, not reset — worth +34.15 pts to LFU and +24.85 pts to EMA α=0.01 at 32 GiB [Simulated (provisional)]
u32 block_tokens Tokens per block — 16 → 5.00 MiB at 320 KiB/token
float ema_alpha LRFU decay α for block scoring (Lee et al., SIGMETRICS 1999) — 0.05, half-life ln0.5/ln0.95 = 13.5 decode steps. α = 0.2 was never validated; α is non-monotonic, no convergence onto LFU as α→0
u32 head_count GQA head count — a scoring input only. Heads are never evicted independently; the unit is a (sequence, token block), 16 tokens = 5.00 MiB
DATA_TYPE_KV_CACHE
Block-granular placement (16 tokens = 5.00 MiB), LRFU scoring, persistent identity-keyed state. No per-head eviction: dropping a head gives wrong outputs, not a miss. No RoPE prefetch: RoPE carries no information about which blocks are needed next.
DATA_TYPE_WEIGHTS
Read-only, layer prefetch, broadcast-optimized
DATA_TYPE_ACTIVATIONS
High-bandwidth, sequential, short-lived
DATA_TYPE_SCRATCH
Temporary workspace, lowest eviction priority
9.3 Driver → Firmware Translation

The GPU driver translates allocation hints into endpoint firmware configuration via CXL.io mailbox commands.

1
Application Allocation Request
Runtime calls extended allocation API with hints
cxl_malloc(size, &hints)
2
Driver Hint Processing
Driver validates hints, selects target endpoint based on capacity and locality
select_endpoint(hints) → EP0
3
Mailbox Command Construction
Pack hints into CXL.io mailbox payload with vendor-specific extensions
build_mailbox_cmd(VENDOR_ALLOC, hints)
4
CXL.io Mailbox Transfer
Command sent over PCIe to endpoint controller
cxl_mailbox_send(ep, cmd)
5
Firmware Configuration
Endpoint firmware configures placement policy, batching, and eviction strategy — and registers the block identity key so its score survives the round trip
configure_region(addr, policy)
CXL.io Mailbox Protocol
GPU Driver
Host CPU
PCIe / CXL.io
Transport
Endpoint FW
ARM Cores
Mailbox Command Payload
opcode: 0xC0 (VENDOR_ALLOC) size: 4096 data_type: KV_CACHE pattern: RANDOM batch_window: 64 lrfu_alpha: 0.05 (half-life 13.5 steps) head_count: 8 (scoring only) block_tokens: 16 (5.00 MiB) persist_block_state: true
9.4 Fault Handling & Latency Breakdown

When the GPU accesses an uncached CXL.mem address, a fault triggers the full access path. The stages below are a component budget, summed from vendor specifications — not a measurement.

This 200 ns total is an assumption. No CXL hardware was measured.

The six-line budget below is demoted from a headline figure to an unmeasured assumption inside a 32–121 GB/s effective-bandwidth range. The 65× ratio once claimed here (13.0 µs ÷ 200 ns) is retracted: it divided a component sum for a hypothetical load/store path by a component sum for a driver-mediated DMA path — two different operations, neither of them measured.

What the package models instead is a 5.00 MiB block (16 tokens at 320 KiB/token) crossing the tier: 164 / 82 / 43 µs at an assumed 32 / 64 / 121 GB/s. The tier's value is avoiding regeneration — forward compute for the same 16 tokens is 9 / 11 / 18 ms at 50 / 40 / 25% MFU. Two to three orders of magnitude, no precise multiple claimed. Analytical model

Latency
GPU/host request issue
Attention kernel issues the access; GPU MMU resolves the VA and identifies the CXL.mem region from the PTE flags.
Latency
PCIe Gen5 ×16 PHY
Serialization and PHY processing onto the Gen5 ×16 link.
Latency
CXL switch traversal
CXL 3.0 switch decodes the address and routes the request to the owning endpoint port.
Latency
Endpoint CXL controller decode
Endpoint controller decodes the CXL.mem request and issues the local memory transaction.
Latency
DDR5-5600 array access
Endpoint DRAM row activate, column access and data return.
Latency
Return path
Response travels back through the switch and PHY to the requesting GPU.
Component-budget total — assumed, never measured
200 ns (assumption)
Access Type Latency Bandwidth CPU Involved
HBM (Local) ~100 ns 8 TB/s No
CXL.mem (Direct) — assumed, not measured 200 ns (component budget) 32 / 64 / 121 GB/s assumed range No
PCIe DMA (Traditional) — different operation 13.0 μs 64 GB/s Yes (driver)
NVMe Read (Tier 3) latency assumed 14 GB/s Yes (filesystem)
The comparison that survives: move vs. regenerate, per 5.00 MiB block (16 tokens, Llama-3 70B GQA)
Operation Cost Basis Evidence
Move the block @ 121 GB/s 43 µs CXL 3.0 ×16 theoretical, one direction Analytical model — assumed
Move the block @ 64 GB/s 82 µs mid-range assumption Analytical model — no CXL hardware measured
Move the block @ 32 GB/s 164 µs pessimistic Analytical model — assumed
Regenerate: forward compute, 16 tokens 9 / 11 / 18 ms 50% / 40% / 25% MFU, linear term Analytical model
Regenerate an arbitrary interior block ≥ the above interior KV is not independently reproducible — hidden states depend recursively on all preceding context Analytical model

Analytical model — capacities are arithmetic over vendor specifications and stand. No CXL hardware was measured. Retracted: the 65× single-access ratio (13,000 ÷ 200), which divided a component sum for a hypothetical load/store path by a component sum for a driver-mediated DMA path — two different operations, neither measured. The 200 ns CXL access is demoted to an unmeasured assumption inside a 32–121 GB/s range, under which a 5.00 MiB block is 164 / 82 / 43 µs. Also retracted here: α = 0.2, per-head eviction, and RoPE-aware prefetch. Scoring is LRFU (Lee et al., ACM SIGMETRICS 1999) External literature. Persistence deltas are Simulated (provisional): one data-structure defect in Revision 1 inverted four of five findings, so Revision 2 is provisional, not settled. Canonical numbers v6.0.