Accelerator Architecture Field Guide

An 8-billion-parameter code assistant receives eight repository prompts, each 4,096 tokens long. It ingests the prompts in a flash, yet stumbles when generating its answers token by token. A chip advertised with double the peak arithmetic per second sounds tempting, but it might leave generation speed almost completely unchanged. What is the machine actually waiting for?

The CUDA foundation introduced kernels, thread blocks, warps, device memory, and synchronized timing. The model-parallelism lesson demonstrated that a logical shard count isn't a physical route. Keep both ideas: an accelerator runs work through a concrete execution hierarchy, and every tensor or collective travels through a concrete physical memory path.

One projection, two different machines inside it

Use an illustrative feed-forward projection for the assistant, not the specifications of a named 8B checkpoint. Each input row has K = 4,096 features, and the projection produces N = 11,008 features. Its simplified matrix multiplication is:

$$Y_{M \times N} = X_{M \times K} W_{K \times N}$$

M is the number of token rows processed together. BF16 stores each matrix element in 2 bytes. The weight matrix contains about 86 MiB, independent of phase.

During prefill, prompt positions can be processed together. Assume the scheduler puts all eight 4,096-token prompts into this projection at once, giving M = 32,768. Chunked prefill would use smaller groups. During ordinary one-token decode, each active sequence contributes one new token, so the same eight requests give M = 8. The weight shape hasn't changed. Reuse has.

One output element combines 4,096 products. Counting a multiply and an add as two floating-point operations (FLOPs), one row costs about 2 × 4,096 × 11,008 = 90,177,536 FLOPs. For M rows:

$$\text{FLOPs} = 2MKN$$

For an isolated operation with cold inputs and a materialized BF16 output, count each input and weight read once and each output write once:

$$\text{bytes} = 2(MK + KN + MN)$$

Dividing work by traffic gives arithmetic intensity, measured in floating-point operations per byte. The script computes both phases as a shape calculation, not a hardware benchmark.

Real kernels reread some tiles, reuse caches, fuse away intermediate writes, and may use quantized weights. The count isn't a measured HBM traffic total or a universal lower bound across fused operations. In the unfused model, prefill presents much more arithmetic per byte than low-concurrency decode. Whether either phase reaches a hardware limit still needs a second number: the machine's compute-to-bandwidth ratio.

Compare work and traffic before predicting time

Consider a hypothetical device with a 200 TFLOP/s compute ceiling and a 1 TB/s memory ceiling for this operation. These are exercise inputs, not a vendor specification or measured rates. Its balance point is 200 FLOP/byte. Below that intensity, the traffic term is larger; above it, the arithmetic term is larger. This is the basic roofline model.^{[1]Reference 1Roofline: An Insightful Visual Performance Model for Multicore Architectureshttps://doi.org/10.1145/1498765.1498785}

For work $F$, bytes $Q$, compute rate $P$, and bandwidth $B$, the optimistic time is:

$$t_{\mathrm{floor}} = \max(F/P, Q/B)$$

The maximum assumes perfect overlap. It leaves out launch overhead, synchronization, contention, and work outside this projection. Predict which term wins before running this CPU-only calculation.

Doubling only the compute ceiling halves the prefill floor here but leaves the decode floor unchanged. That isn't a measured speedup. A real kernel can sit far below either ceiling, and a full decode step also reads KV cache, runs other layers, and may exchange collective messages. Long context, higher concurrency, speculative decoding, and cache residency can change which resource dominates.

Why can the same projection be compute-heavy during prefill and memory-heavy during decode?

Answer

Prefill applies each loaded weight tile to 32,768 token rows in this workload, while decode applies it to only 8. The large prefill matrix gets far more arithmetic from each byte moved. Decode has much lower arithmetic intensity, so weight traffic and fixed dispatch costs can dominate.

The memory hierarchy and the 52x bandwidth cliff

Every accelerator bridges a massive physical gap between fast, tiny execution storage and large, slow capacity. When data moves between compute units and memory, it steps through a strict hierarchy of bandwidth tiers:

Notice the staggering drop across boundaries. Stepping from on-chip registers down to HBM3 represents an order-of-magnitude reduction in throughput. But stepping from device HBM3 down to the host PCIe 5.0 bus drops available bandwidth from 3,350 GB/s to 64 GB/s, a brutal 52x bandwidth cliff.

During decode, the model reads its entire weight set to generate a single token per stream. Offloading even a fraction of those weights to host memory over PCIe means the GPU spends over 98% of its time waiting for the bus. Fast inference requires keeping weights and active KV caches pinned inside device HBM.

Compute engines: SIMT warps vs systolic arrays

Accelerators calculate matrix products through two primary hardware philosophies: Single Instruction, Multiple Threads (SIMT) with dedicated matrix accelerators, and hardwired 2D systolic arrays. Both evaluate the fundamental matrix multiply-accumulate primitive:

$$D = A \cdot B + C$$

Their physical data paths could hardly be more different.

SIMT warps and matrix cores

In an NVIDIA SM or AMD Compute Unit, threads execute in lockstep groups: 32 threads in a CUDA warp, 64 threads in an AMD CDNA wavefront. For general arithmetic, each thread accesses its own registers and executes independent operations.

Matrix multiplication uses specialized hardware units: Tensor Cores on NVIDIA, Matrix Cores on AMD. Rather than issuing independent scalar instructions, threads in a warp coordinate to execute a collective instruction (such as mma.sync or Hopper's asynchronous wgmma). In Hopper, a warp group of 128 threads issues matrix operations directly from shared memory into Tensor Cores, bypassing register-file pressure.

SIMT's greatest strength is flexibility. When the model finishes its projection and moves to layer normalization, rotary position embeddings, or token sampling, those exact same SM cores immediately run scalar and vector kernels without leaving silicon idle.

Systolic arrays

Google TPUs and AWS Trainium chips route matrix math through 2D systolic arrays. A TPU v6e Matrix Multiply Unit (MXU) contains a grid of Processing Elements (PEs) wired directly to their immediate horizontal and vertical neighbors.

In a weight-stationary systolic array, weights from matrix $B$ are preloaded into PE registers and remain stationary throughout computation. Activations from matrix $A$ stream in from the left boundary, with rows staggered by one clock cycle so products align correctly. As activations flow horizontally, each PE computes its product, adds it to the incoming partial sum from above, and passes the updated sum downward.

Systolic arrays avoid repeatedly reading and writing large multiported register files because operands travel directly between adjacent processing elements over short internal wires. That design delivers exceptional energy efficiency and silicon density for large, dense GEMMs.

The trade-off is rigid geometry. A $128 \times 128$ systolic array achieves peak efficiency only when both dimensions are populated. When executing a skinny decode step with $M = 8$, over 93% of the systolic array sits idle, waiting for work that isn't there.

How does the execution of a matrix multiply-accumulate differ between a SIMT warp and a systolic array?

Answer

A SIMT warp executes cooperatively through software-scheduled instructions (such as mma or wgmma), feeding matrix fragments into Tensor Cores while retaining the flexibility to run general scalar or reduction code. A systolic array routes data directly through a hardwired 2D mesh of processing elements where weights remain stationary and activations stream across neighbors, maximizing energy efficiency for large dense GEMMs at the expense of utilization on skinny decode shapes.

The invariant behind every accelerator

Fast accelerator code keeps reused data close to compute, performs enough work before eviction, and overlaps the next transfer with current arithmetic. Vendor names differ, but four ownership questions don't:

1. Which group executes one instruction together?
2. Which memory can that group share at low latency?
3. Who schedules movement between large memory and local memory?
4. Which physical link carries bytes when work spans devices?

Time to first token (TTFT) measures from request submission to the first returned token, including queueing and any transport within the measurement boundary. Time per output token (TPOT) commonly averages the subsequent generation interval over output_tokens - 1; inter-token latency measures individual gaps. Report the convention, since p95 of request-average TPOT can hide isolated long pauses.

For a tiled projection, large memory supplies input and weight tiles, on-chip storage holds the active pieces, and compute consumes them. Results travel back in the opposite direction. The arrows below describe data movement, not a promise that all platforms expose identical memory instructions.

A platform port is incomplete until latency and output checks use the same prompt distribution. Holding the graph constant doesn't hold its execution path constant.

Documentation boundary, checked September 2, 2026: The examples target NVIDIA H100 (Hopper) and B200-class Blackwell, AMD MI300/CDNA 3 and MI350/CDNA 4 using pinned ROCm 7.14 docs, Google TPU v6e, AWS Trainium2 using Neuron 2.32 memory docs, and Apple silicon through PyTorch MPS. These are selected architectures, not a list of each vendor's newest products. Check exact hardware, OS, driver, and library compatibility before choosing a deployment image.^{[2]Reference 2NVIDIA Hopper Tuning Guidehttps://docs.nvidia.com/cuda/archive/13.0.0/hopper-tuning-guide/index.html[3]Reference 3Blackwell Tuning Guidehttps://docs.nvidia.com/cuda/archive/13.0.2/pdf/Blackwell_Tuning_Guide.pdf[4]Reference 4ROCm Core SDK 7.14.0 Release Noteshttps://rocm.docs.amd.com/en/docs-7.14.0/about/release-notes.html[5]Reference 5TPU v6ehttps://cloud.google.com/tpu/docs/v6e[6]Reference 6The Trainium Memory Hierarchyhttps://awsdocs-neuron.readthedocs-hosted.com/en/v2.32.0/nki/get-started/about/memory-hierarchy-overview.html[7]Reference 7Accelerated PyTorch training on Mac.https://developer.apple.com/metal/pytorch/}

The local-memory vocabulary is different, but the job is recognizable:

The table compares ownership, not speed. After a symptom appears, inspect the platform-specific owner in its row.

NVIDIA Hopper and Blackwell: preserve mapping, retune pressure

CUDA presents grids of thread blocks, blocks of threads, 32-thread warps, and streaming multiprocessors (SMs). Threads in one block share on-chip shared memory and can synchronize. Global device memory holds large tensors. Those boundaries from the CUDA prerequisite remain valid on Hopper and Blackwell.^{[8]Reference 8CUDA Programming Guide.https://docs.nvidia.com/cuda/cuda-programming-guide/}

Hopper added the Tensor Memory Accelerator (TMA), which can move multidimensional tensor tiles between global and shared memory while thread blocks continue independent work. Hopper also added thread-block clusters and distributed shared memory across blocks in a cluster.^{[2]Reference 2NVIDIA Hopper Tuning Guidehttps://docs.nvidia.com/cuda/archive/13.0.0/hopper-tuning-guide/index.html} Those features are useful only when a kernel has a tile worth reusing. Copying the decode projection's 86 MiB weight matrix through shared memory without enough rows to reuse each tile adds staging without changing the bandwidth problem.

Blackwell retains and extends the CUDA programming model. NVIDIA's Blackwell tuning guide still starts with coalesced global access, reduced redundant traffic, suitable launch configuration, and limited warp divergence. It also warns that occupancy limits differ across Blackwell compute capabilities, so “Blackwell” isn't one register-and-shared-memory budget.^{[3]Reference 3Blackwell Tuning Guidehttps://docs.nvidia.com/cuda/archive/13.0.2/pdf/Blackwell_Tuning_Guide.pdf}

For the running workload, use the same algorithmic split on both generations:

• Prefill: choose a matrix path whose tile shape keeps Tensor Core work dense, then overlap global-to-shared movement with current computation.
• Decode: batch enough active rows to reuse weights, use a kernel specialized for small M, or reduce bytes through a verified quantization path.
• Port from Hopper to Blackwell: ship native code for the target or a compatible PTX intermediate representation, then retune. PTX using architecture-conditional features is an exception to general forward compatibility: Hopper compute_90a PTX doesn't run on Blackwell. A binary that launches is only a compatibility result.^{[9]Reference 9Blackwell Architecture Compatibilityhttps://docs.nvidia.com/cuda/blackwell-compatibility-guide/}

A CUDA failure that looks like architecture progress

Suppose a Hopper kernel uses more shared memory to stage a larger weight tile. One block now reuses more bytes, but fewer blocks remain resident on each SM. TPOT gets worse.

The symptom isn't proof that shared-memory tiling failed. Larger tiles improved reuse and reduced concurrency at the same time. Compare achieved occupancy, memory traffic, and eligible warps, then test smaller tiles. If decode still moves almost the same weight bytes per token, more staging can't create missing reuse.

A Hopper decode kernel gets slower after its shared-memory tile doubles. Which two effects must be separated?

Answer

The larger tile may reduce global-memory traffic per block, but its shared-memory and register footprint may lower occupancy. Measure both data movement and resident or eligible warps before deciding whether the tile helped.

AMD CDNA and ROCm: familiar syntax, different lane contract

Hopper-to-Blackwell tuning preserves the warp width. Crossing to AMD's selected CDNA generations changes that assumption. HIP, AMD's C++ GPU programming interface, deliberately resembles CUDA: kernels launch grids of blocks, threads use threadIdx and blockIdx, and a work-group shares Local Data Share (LDS). That similarity helps source portability, but not every lane-level algorithm survives it.

The CDNA 3 and CDNA 4 targets here use 64-thread wavefronts. NVIDIA code commonly assumes a 32-thread warp. A reduction that hardcodes masks, lane counts, or “four warps per 128-thread block” can return wrong values or waste half a wave after a mechanical HIP conversion. Don't generalize this width to every AMD architecture.^{[10]Reference 10Introduction to the HIP Programming Modelhttps://rocm.docs.amd.com/projects/HIP/en/develop/understand/programming_model.html}

Repair the execution model before changing syntax:

• derive lane behavior from warpSize or use library primitives with documented semantics;
• recalculate work-group size as a number of waves on the selected CDNA target;
• remeasure vector-register and LDS pressure, because both limit resident waves;
• compile for the exact LLVM gfx target in the pinned environment.

As of the snapshot date, ROCm 7.14 documentation lists MI300-series accelerators as CDNA 3 with gfx942 and MI350-series accelerators as CDNA 4 with gfx950.^{[4]Reference 4ROCm Core SDK 7.14.0 Release Noteshttps://rocm.docs.amd.com/en/docs-7.14.0/about/release-notes.html} Keep those identifiers in build and benchmark records. “ROCm passed” without hardware target, ROCm version, and library versions isn't reproducible.

The ROCm Communication Collectives Library (RCCL) provides operations such as all-reduce across GPUs. Its MI300X guidance describes eight-GPU systems where every accelerator pair has dedicated Infinity Fabric links. Using only part of that topology changes available collective routes, so a tensor-parallel result needs the selected ranks, not only TP=4.^{[11]Reference 11RCCL Documentationhttps://rocm.docs.amd.com/projects/rccl/en/docs-7.2.4/}

For the projection, CDNA's decision remains phase-specific. Dense prefill can keep matrix units busy through large tiles. Decode needs enough concurrent rows, a small-M kernel, or fewer weight bytes. Replacing CUDA API names with HIP API names can't change M = 8.

A 128-thread CUDA reduction is ported to CDNA 3 and keeps a hardcoded 32-lane shuffle mask. Why is this a correctness risk rather than only a tuning issue?

Answer

CDNA 3 uses 64-thread wavefronts. A hardcoded 32-lane algorithm may combine only half of each wave or apply invalid lane assumptions, so output can be wrong. Replace the assumption with documented cross-lane primitives or a wave-size-aware algorithm, then test against a reference.

TPU and Pallas: map tiles, not CUDA threads

CUDA and HIP both expose thread hierarchies. A Tensor Processing Unit (TPU) asks for a different mapping. Each v6e chip has one TensorCore containing two matrix-multiply units (MXUs), a vector unit, and a scalar unit. The MXUs handle dense matrix work, while vector and scalar operations have separate resources. Google's TensorCore names this larger assembly, not NVIDIA's matrix unit.^{[5]Reference 5TPU v6ehttps://cloud.google.com/tpu/docs/v6e}

Pallas is JAX's custom-kernel layer for GPU and TPU. On TPU, kernel inputs usually reside in high-bandwidth memory (HBM), while kernel-body references point into faster vector memory (VMEM) or scalar memory (SMEM). BlockSpec describes which input and output tile each program sees; the compiler can overlap HBM transfers with computation.^{[12]Reference 12Writing TPU Kernels with Pallashttps://docs.jax.dev/en/latest/pallas/tpu/details.html}

Grid order carries extra meaning. Pallas TPU programs normally advance sequentially in lexicographic grid order. Consecutive programs that use the same input slice can reuse data already in VMEM and skip another HBM transfer.^{[12]Reference 12Writing TPU Kernels with Pallashttps://docs.jax.dev/en/latest/pallas/tpu/details.html} Reordering grid axes can therefore change traffic without changing the mathematical output.

Map the running projection this way:

1. Put the 86 MiB weights in HBM.
2. Use BlockSpec to select weight and activation tiles that fit VMEM.
3. Arrange the prefill grid so consecutive output tiles reuse one weight or input slice where possible.
4. Accumulate matrix work on MXUs and keep softmax, normalization, or element-wise work on suitable vector paths.
5. For M = 8 decode, test padded or batched shapes against real TPOT. A large MXU doesn't guarantee high use for a skinny matrix.

Pallas block shapes have backend constraints, including divisibility requirements on the trailing dimensions for TPU paths.^{[12]Reference 12Writing TPU Kernels with Pallashttps://docs.jax.dev/en/latest/pallas/tpu/details.html} A shape rejection at compile time isn't an availability incident. The tile contract failed. Pad with correctness masking, choose a legal block, or keep the operation in compiled JAX when a custom kernel doesn't earn its maintenance cost.

Which operand does grid order reuse?

For one fixed reduction tile, write an output tile as Y[i,j] = X[i] @ W[j]. Visiting (0,0), (0,1), (1,0), (1,1) keeps X[0] for the first two programs, then X[1]. Visiting (0,0), (1,0), (0,1), (1,1) instead keeps each weight tile for two programs. Neither order saves both operands automatically.

Count loads under a deliberately small model: one currently resident tile per operand, no reuse after eviction. Each activation tile is 1 KiB and each weight tile is 4 KiB. These sizes teach the trade-off; they aren't a legal Pallas kernel configuration or a TPU performance simulation.

Columns-first saves activation loads, but rows-first moves fewer bytes in this fixture because weights are larger. If both tiles have equal size, the byte totals tie. For a real matrix multiplication with multiple reduction tiles, also preserve Pallas's requirement that updates to the same output slice be consecutive; the reduction axis normally varies last. Inspect the complete grid and transfer trace, not just one reused input.^{[12]Reference 12Writing TPU Kernels with Pallashttps://docs.jax.dev/en/latest/pallas/tpu/details.html}

Trainium and NKI: assign work to engines and memories

Pallas exposes tile movement through a compiler-managed grid. AWS's Neuron Kernel Interface (NKI) exposes another arrangement of engines and buffers. Trainium2's NeuronCore-v3 contains tensor, vector, scalar, and general-purpose single-instruction, multiple-data (GPSIMD) engines plus software-managed on-chip SRAM.^{[13]Reference 13NeuronCore-v3 Architecturehttps://awsdocs-neuron.readthedocs-hosted.com/en/v2.29.1/about-neuron/arch/neuron-hardware/neuron-core-v3.html}

NKI's memory path has three named levels for our projection:

• HBM holds kernel inputs and outputs.
• SBUF is the main software-managed on-chip buffer shared by compute engines.
• PSUM holds partial matrix-multiply accumulations near the tensor engine.

An NKI kernel loads HBM tiles into SBUF, performs work from internal memory, accumulates matrix results in PSUM when appropriate, and stores completed outputs back through SBUF to HBM. If live tiles exceed a buffer's capacity, the compiler can spill to the next memory level: PSUM pressure can add SBUF traffic, while SBUF pressure can add HBM traffic. Check which level overflowed before attributing every spill to HBM.^{[6]Reference 6The Trainium Memory Hierarchyhttps://awsdocs-neuron.readthedocs-hosted.com/en/v2.32.0/nki/get-started/about/memory-hierarchy-overview.html}

The running projection suggests an engine schedule: direct-memory-access engines prefetch the next X and W tiles, the tensor engine performs matrix multiplication, vector or scalar engines handle fused follow-up operations they support, and the current output tile accumulates in PSUM. NKI tile dimensions also distinguish partition and free dimensions, so a valid NumPy shape isn't yet a valid physical layout.

Trainium2 systems connect chips through NeuronLink-v3 in a 4-by-4 two-dimensional torus within a 16-chip instance.^{[14]Reference 14Amazon EC2 Trn2 Architecturehttps://awsdocs-neuron.readthedocs-hosted.com/en/v2.32.0/about-neuron/arch/neuron-hardware/trn2-arch.html} Keep tensor-parallel ranks aligned with the actual topology. A rank count alone hides routing, exactly as it did for CUDA and ROCm.

An NKI failure that keeps answers correct

The prefill kernel grows its fused region to remove one HBM round trip. Live intermediates no longer fit SBUF, so the compiler adds spills. Latency increases and the trace shows extra DMA traffic.

Split the fusion or shrink tiles until the live set stays on chip. “More fusion” isn't a monotonic optimization when software-managed local memory is the limiting resource.

Why can a larger fused NKI kernel run slower even though it launches fewer kernels?

Answer

Fusion lengthens the lifetime of intermediate tiles. If that live set overflows a buffer, spills add traffic to the next memory level. SBUF spills can reach HBM; PSUM spills can first increase SBUF pressure. Fewer launches can lose to extra movement.

Apple MPS: unified memory changes transfer, not capacity

The previous stacks distinguish host memory from large accelerator memory. Apple silicon changes that physical boundary. PyTorch's Metal Performance Shaders (MPS) backend maps tensor operations to MPS Graph and tuned MPS kernels on Apple platforms.^{[7]Reference 7Accelerated PyTorch training on Mac.https://developer.apple.com/metal/pytorch/[15]Reference 15MPS backend.https://docs.pytorch.org/docs/stable/notes/mps} MPS Graph represents operations and tensors as a symbolic compute graph that can be compiled into an executable.^{[16]Reference 16Metal Performance Shaders Graphhttps://developer.apple.com/documentation/metalperformanceshadersgraph} This is a GPU backend, not automatic execution on the Apple Neural Engine.

Apple silicon uses a unified physical memory pool. The CPU and GPU don't have the discrete host-RAM-to-VRAM boundary used by a typical data-center GPU. PyTorch still treats cpu and mps as different device targets, and unified memory is still finite.^{[17]Reference 17MLX: An array framework for Apple siliconhttps://github.com/ml-explore/mlx[15]Reference 15MPS backend.https://docs.pytorch.org/docs/stable/notes/mps}

For the 8B BF16 model, weights alone need about 16 GB in decimal units. The machine also needs KV cache, temporary activations, allocator headroom, macOS, and other applications. A Mac advertised with enough total memory can still enter pressure or fail allocation once the full serving ledger is counted.

For this comparison, use MPS as the single-Mac path: validate graph behavior, test product logic, and measure an on-device latency envelope. Keep three questions separate:

1. Is the operation implemented on MPS?
2. Does the whole model fit within safe system-memory headroom?
3. Does the graph execute fast enough after warmup and synchronization?

When PYTORCH_ENABLE_MPS_FALLBACK=1 is enabled, an unsupported MPS operation can fall back to the CPU and create a latency cliff.^{[18]Reference 18MPS Environment Variables.https://docs.pytorch.org/docs/stable/mps_environment_variables.html} Unified memory makes the detour less visually obvious than a discrete device copy, but it doesn't make CPU and GPU execution equally fast. Profile for CPU operations and synchronization gaps before blaming GPU arithmetic.

Interconnect topologies: why tensor parallelism is strictly intra-node

When a model exceeds the memory or compute capacity of a single chip, workloads shard across multiple accelerators. The physical interconnect determines which parallelization strategies are viable:

Notice the sharp boundary between intra-node fabrics (NVLink, Infinity Fabric) and cross-node network links (InfiniBand, RoCE). This boundary explains why tensor parallelism (TP) is strictly intra-node.

In Megatron-style tensor parallelism, every transformer layer contains two communication collectives: one All-Reduce after the multi-head attention projection, and another All-Reduce after the MLP down-projection. For an 80-layer model such as Llama 3 70B, generating a single decode token requires:

$$\text{Collectives per token} = 2 \times 80 = 160\ \text{All-Reduces}$$

During decode with batch size 8 and hidden dimension 8,192 in BF16, each All-Reduce moves only 8 × 8,192 × 2 = 128 KB. At this tiny payload, transfer duration is completely dominated by network latency, packet serialization, and kernel launch overhead, not raw link bandwidth.

On an NVLink or NVSwitch crossbar with hardware reduction engines, an All-Reduce finishes in 1 to 2 microseconds. Across the full 80 layers, 160 collectives consume under 0.3 ms total.

Across an inter-node network, traversing PCIe to the network interface card (NIC), passing through leaf-spine switches, and handling network protocol stacks adds 5 to 15 microseconds per collective. Multiplying 160 collectives by an optimistic 10 microseconds yields 1.6 ms of pure idle flight time per token, before counting network contention or actual byte transfer. For an interactive target of 25 ms TPOT, burning several milliseconds on network latency destroys the user experience.

Tensor parallelism belongs strictly inside the high-speed intra-node fabric. Scaling across cluster nodes relies on Pipeline Parallelism (where communication occurs only at pipeline boundaries) or Data and Expert Parallelism (where collectives execute per forward pass rather than twice per layer per token).

Why is tensor parallelism strictly restricted to intra-node interconnects rather than crossing InfiniBand or Ethernet?

Answer

Tensor parallelism requires two All-Reduce collectives per transformer layer. An 80-layer model executes 160 All-Reduces per generated token. At decode batch sizes, the payload is tiny (tens of kilobytes), so communication is completely latency-dominated. While NVLink completes each collective in 1 to 2 microseconds, inter-node network stacks add 5 to 15 microseconds per hop, accumulating several milliseconds of pure idle flight time per token and destroying token-generation latency.

Choose programming depth after the bottleneck

Most ports should stop at the highest layer that meets correctness and service-level objectives (SLOs). Each step downward gains control and creates a new maintenance surface:

Portability is a test matrix, not a source-language property. A Triton kernel may support NVIDIA and AMD while using different legal tile sizes. A Pallas kernel may share BlockSpec vocabulary across GPU and TPU while requiring backend-specific memory and shape rules. NKI exposes Trainium's memory hierarchy directly. MPS Graph may compile a complete graph without offering the same custom-kernel surface as CUDA.

For our assistant, start with graph-level BF16 inference on each candidate. Preserve tokenization, weights, prompts, decoding settings, and output checks. Move down only after a trace identifies a stable projection, attention, normalization, or collective hotspot.

Read failures as ownership mistakes

The same symptom can point to a different owner on each platform. Use evidence in this order: output, phase, memory path, execution mapping, then topology.

Don't change precision, batching, kernel code, and topology in one experiment. That destroys attribution. Keep one baseline and change one ownership boundary at a time.

The four-gate hardware evaluation funnel

Suppose the assistant must meet p95 TTFT below 450 ms and p95 TPOT below 45 ms for eight 4,096-token prompts. Those thresholds are exercise inputs, not published accelerator results.

Selecting hardware requires passing four non-negotiable evaluation gates in strict sequence:

Gate 1: Correctness and numerical precision

Evaluate whether the candidate runs the model accurately across supported precisions:

• Dtype mechanics: Compare baseline BF16 against 8-bit floating point (such as FP8 E4M3 for weights and activations, or E5M2 for wider dynamic range) and microscaling block formats (such as MXFP4 or NVFP4).
• Tolerance verification: Check layer outputs against reference implementations using absolute and relative error bounds. Inspect attention logit distributions for underflow and verify that downstream task benchmarks (such as MMLU or coding benchmarks) show no degradation.
• Decision: Any candidate that produces numerical drift, NaN/Inf values, or degraded task quality is immediately disqualified.

Gate 2: Memory fit and residency

Calculate the full memory ledger before launching serving benchmarks:

$$\text{Memory required} = \text{Weights} + \text{KV cache}(B, L) + \text{Activation buffers} + \text{Runtime workspace}$$

• For the 8B BF16 model, weights consume about 16 GB. At batch 64 and context 4,096, the KV cache adds substantial memory.
• Verify that total memory leaves at least 10% to 15% headroom to prevent allocator thrashing or out-of-memory crashes during burst traffic.
• Decision: If weights and target concurrency exceed safe physical capacity without aggressive eviction, the configuration fails.

Gate 3: Latency SLOs under production concurrency

Measure response times under the target prompt arrival distribution:

• TTFT (Time to First Token): Tests prefill throughput, prompt queueing, and scheduling efficiency.
• TPOT (Time per Output Token): Tests decode memory bandwidth, continuous-batch coordination, and collective communication.
• Decision: Disqualify any candidate that fails either p95 TTFT (<450 ms) or p95 TPOT (<45 ms), regardless of its peak batch throughput.

Gate 4: Goodput and cost-efficiency

Only configurations that clear Gates 1, 2, and 3 compete on efficiency:

• Goodput: Measures tokens per second that strictly meet the latency SLOs (dropping or penalizing requests that breached tail limits).
• Cost calculation: Compute goodput divided by total cost of ownership (hardware capital expenses, cloud hourly rates, and power).
• Decision: Rank the remaining candidates by valid tokens per dollar ($T/$$).

Two accelerators pass numerical checks. Candidate A has higher total tokens per second but misses the p95 TPOT limit. Candidate B meets TTFT and TPOT with lower total throughput. Which candidate is eligible?

Answer

Candidate B. Correctness, fit, and latency SLOs are non-negotiable gates. Throughput ranks only the candidates that pass all prior gates.

The field guide leaves one unresolved problem: each diagnosis above depends on trustworthy timing, traces, counters, and output comparisons. The next lesson turns those words into a repeatable benchmark and correctness workflow.