Software Ecosystems & Lock-In

Every accelerator decision in Chapter 7.1 through Chapter 7.8 was framed as a hardware comparison: HBM capacity, FP4 PFLOPS, NVLink bandwidth, dollars per GPU. That framing is incomplete in a way that costs real money. When you buy an accelerator you are also signing a software contract — a runtime, a kernel language, a collective library, a compiler, a serving stack, and a fabric — and that contract is what actually determines how much of the silicon's peak you will ever see, how fast a new model lands on the box, and how much it costs to leave. The four regimes that matter in 2026 are CUDA (NVIDIA), ROCm (AMD), XLA (Google TPU, via JAX/TensorFlow), and Neuron (AWS Trainium/Inferentia). They are not interchangeable. Choosing one is choosing a lock-in regime.

This chapter works through three forks — which regime, which portability posture, which fleet composition — and their downstream costs: the porting project that delays revenue, the realized-MFU gap that erases a price advantage, the kernel you have to write yourself because no library shipped it, the model that lands on NVIDIA on day zero and on everyone else on day forty. The hardware buyer who ignores the software contract does not save money. They defer the bill to the day the cluster is live and the MFU is half of what the datasheet promised.

It is tempting to model the four ecosystems as "CUDA, plus three things trying to be CUDA." That is wrong in a way that hides the decision. Each regime locks a different layer of the stack, erodes at a different rate, and exacts a different exit cost. Sorting them by what they lock is the first analytical move.

CUDA is the deepest and broadest moat, and it is mostly a library and tooling moat, not merely a language one. The language (CUDA C++) is replaceable; what is hard to replace is two decades of accreted libraries — cuDNN, cuBLAS, CUTLASS, NCCL, TensorRT-LLM — plus the profilers, debuggers, and the simple fact that every new model, every new attention variant, every new quantization scheme ships and is tuned on NVIDIA first. The lock-in is not that you cannot leave; it is that on NVIDIA the kernel you need already exists and is tuned, and everywhere else you may have to write or wait for it.

ROCm is the open challenger, and as of 2026 it has closed the gap from "unusable for production" to "feature-matched on the workloads AMD chooses to fund." ROCm mirrors CUDA layer-for-layer — HIP for the language, MIOpen/rocBLAS for the math, RCCL for collectives — and the porting tax is genuinely low for code written against PyTorch rather than raw CUDA. The residual lock-in is coverage: the long tail of custom kernels and bleeding-edge model architectures that are tuned on CUDA months before ROCm catches up. ROCm's exit cost is asymmetric — coming from CUDA-portable PyTorch is cheap; coming from hand-written CUDA C++ is not.

XLA (the compiler behind JAX and TensorFlow on TPU) inverts the model entirely. There is no kernel language to lock; XLA is a graph compiler that takes a whole-program traced computation and emits TPU machine code, and the lock-in is that your code is written in JAX against XLA's compilation model, not in a portable imperative GPU style. The upside is that XLA is genuinely multi-target — it compiles to GPU too — so the regime locks the programming model more than the silicon. The downside: the ecosystem of pretrained checkpoints, third-party kernels, and community tooling around JAX/TPU is narrower than CUDA's, and TPUs are rentable only from Google.

Neuron (AWS's SDK for Trainium and Inferentia) is the most vertically integrated and the most single-tenant. The Neuron compiler ingests PyTorch/JAX and targets the Trainium ISA; the lock-in is the whole stack plus the fact that Trainium exists only inside AWS. The trade is explicit: external adopters cite roughly 50% cost savings versus comparable NVIDIA capacity, paid for with a Neuron porting project and the loss of any exit to another cloud. → Chapter 7.4.

A costly misconception in accelerator procurement is that two chips with similar peak FLOPS deliver similar work. They do not, and the gap is a software phenomenon. Model-FLOPS-utilization is the ratio of the useful floating-point work your model actually performs to the chip's theoretical peak over the same wall-clock. On a mature stack running a well-tuned dense transformer, training MFU lands in the 35–55% range; on an immature stack, or an exotic architecture, or a workload nobody has optimized yet, it can be half that — and the silicon is identical. The difference is entirely the kernels, the fused attention path, the collective-communication overlap, the compiler's ability to keep the tensor cores fed.

This is where ROCm's story turned in 2026, and it is worth stating precisely because the direction of travel matters more than any single snapshot. The historical knock on AMD was a severe realized gap: the MI300X carried roughly 1.5x the paper FLOPS of an H100 but, on real inference, delivered something like 37–66% of H100/H200 throughput because the kernels and serving stack were immature (SemiAnalysis AMD-vs-NVIDIA benchmarking, 2025). By mid-2026 that gap had narrowed dramatically on the workloads AMD and the community chose to fund. On specific recent models, the MI355X running ROCm on SGLang reached feature parity with B200 running CUDA — same FP8 KV-cache path, same MLA kernels — and undercut the B200 on cost per million tokens across much of the single-node Pareto frontier (SemiAnalysis InferenceX, 2026). The lesson is not "AMD won" or "NVIDIA won." The lesson is that the realized-MFU gap is a moving target set by software investment, workload by workload, and a procurement decision made on last quarter's gap may be wrong this quarter.

The consequence for the buyer is structural. A 15–30% hardware-price advantage (AMD's typical discount to NVIDIA) is real money only if the realized-MFU gap on your workload is smaller than that discount. If the chip is 20% cheaper but delivers 40% less goodput on your model, you have bought a more expensive cluster measured the only way that matters — cost per token or cost per useful FLOP. Flip it around and the same logic favors the challenger: on a workload where ROCm has reached parity and the chip is cheaper, the cheaper silicon now wins on cost-per-token. The number you must produce before signing is not on any datasheet; it is the realized cost-per-token of your workload on the software you will run.

If lock-in is a contract, portability layers are the clause that lets you exit — but you pay for the clause whether or not you exercise it. The strategic decision is how much of your stack to write against a portable abstraction versus a vendor-native one, and it is made years before you would ever switch. Three layers matter in 2026, at three different altitudes.

The framework layer (PyTorch / JAX). The single most important portability decision most teams make is simply to write models in PyTorch (or JAX) rather than in raw CUDA. PyTorch is the great equalizer: a model written in idiomatic PyTorch runs on CUDA, ROCm, and — via the PyTorch/XLA bridge — on TPU, with the framework absorbing most of the vendor difference. This is why ROCm's porting tax is low for PyTorch code and high for hand-written CUDA: the abstraction already exists, and AMD invests heavily in keeping the AMD backend upstreamed. The lock-in you create by dropping into vendor-native kernels for the last 10% of performance is the lock-in you will pay to undo.

The kernel layer (Triton / MLIR). Below the framework sits the kernel, and this is where CUDA's deepest moat is being abstracted away. Triton — the kernel language built on MLIR — lets you write a fused attention or quantization kernel once and compile it to NVIDIA PTX, AMD AMDGCN, or Intel GPU code, and since PyTorch 2.0 it has been the default code generator inside torch.compile's TorchInductor. Its governance moved from OpenAI to a community triton-lang project with contributions from NVIDIA, AMD, Intel, Meta, IBM and Red Hat — a deliberately multi-vendor structure. A custom kernel written in Triton is portable; the same kernel written in CUDA C++ with wgmma and tensor-memory-accelerator intrinsics is not. This is the layer where the CUDA moat is most actively eroding, and where your kernel-authoring policy directly sets your future switching cost.

The serving layer (OpenAI-compatible APIs). For inference, the cheapest portability you can buy is an OpenAI-compatible HTTP surface in front of whatever engine you run. vLLM, SGLang, TensorRT-LLM, and their ROCm equivalents all expose a near-identical request API, so the application above never learns which silicon it is talking to. This makes the inference fleet the easiest place to be heterogeneous: route traffic to whichever accelerator delivers the best cost-per-token this quarter, behind a stable API. The lock-in at this layer is shallow by design.

"Lock-in" is a slogan until you put a number on it, and the number is the total cost of switching regimes — which is not one cost but four, and they are paid at different times.

The porting cost is engineering time to make your stack run at all on the new regime. For PyTorch-native training and inference it is genuinely modest — days to weeks, mostly dependency and container work. For a codebase laced with hand-written CUDA kernels, custom NCCL tuning, or TensorRT-LLM-specific serving paths, it is a quarters-long project staffed by scarce specialists. This cost scales with how much vendor-native code you wrote in the first place, which is why the kernel-authoring policy of two years ago is the switching cost of today.

The realized-MFU recovery cost is the larger and more often ignored term. Getting your model to run on the new regime is not getting it to run well. The gap between first-light and tuned-throughput is the period where you are paying for silicon at a fraction of its peak while engineers chase the missing 20–40% of MFU — re-tuning kernels, fixing collective overlap, matching the serving framework's batching to the new memory system. On an immature stack this recovery can take longer than the port itself, and during it your cost-per-token is underwater.

The day-0 coverage cost is the opportunity cost of not being on the regime where new models land first. New architectures, attention variants, and quantization schemes ship and are tuned on CUDA first; on other regimes they arrive days to weeks later, after community or vendor porting. For a frontier lab that must serve the newest model the hour it drops, that lag is lost revenue and a competitive gap — and it is a recurring tax, not a one-time switching cost.

The operational-complexity cost is the standing overhead of running a heterogeneous fleet: two toolchains, two driver/firmware matrices, two sets of failure modes, two on-call playbooks. It is paid every day, not once, and it is the reason many operators stay single-vendor despite a price advantage elsewhere — the second toolchain needs an owner, and that owner is a headcount.

The final fork is fleet composition, a strategy decision that arrives as a procurement one. A single-vendor fleet buys operational simplicity: one toolchain, one driver matrix, one set of failure modes, one on-call playbook, and the deepest day-0 model coverage if that vendor is NVIDIA. It pays for that simplicity with price exposure — you are a captive buyer with no credible alternative to walk to — and with supply exposure, because a single allocation queue gates your entire build (→ Chapter 7.6 on the HBM/CoWoS upstream gate).

A heterogeneous fleet inverts both terms. Running NVIDIA for the bleeding edge and AMD, TPU, or Trainium for the steady-state workloads where they have reached parity buys price discipline — a real outside option in every negotiation — and supply diversity across multiple allocation queues. It pays for that leverage with the operational-complexity cost above: two toolchains that each need an owner, and a routing layer (best built at the serving API) that sends each workload to the silicon with the best current cost-per-token. The heterogeneous strategy is most defensible exactly at the inference layer, where the OpenAI-compatible API makes the silicon swap nearly free, and least defensible at the frontier-training layer, where day-0 coverage and the deepest scale-up domain still favor a single vendor.

The 2026 reality that forces this question onto every roadmap is the rise of custom ASICs to roughly a quarter of AI-server shipments. Every Maia, MTIA, TPU, and Trainium is, by construction, a non-CUDA software regime — and the hyperscalers building them have already paid the heterogeneity tax internally because the cost-per-token and supply-control upside justified it at their scale. The strategic question for everyone smaller is whether their workload mix and engineering bench can capture the same upside, or whether the operational-complexity cost eats it. → the merchant-silicon disruption in Chapter 7.5; TCO that scores all of this in Chapter 7.11 and Chapter 1.8.