Software Ecosystems & Lock-In

Chapter 7.1 through Chapter 7.8 compared accelerators on hardware — HBM capacity, FP4 PFLOPS, NVLink bandwidth, dollars per GPU — and that comparison is incomplete in a way that costs real money. Buying an accelerator also commits you to a runtime, a kernel language, a collective library, a compiler, a serving stack, and a fabric, and that stack is what decides how much of the silicon's peak you ever see, how fast a new model lands on the box, and how much it costs to leave. Four stacks matter in 2026: CUDA (NVIDIA), ROCm (AMD), XLA (Google TPU, via JAX/TensorFlow), and Neuron (AWS Trainium/Inferentia). Each is a distinct lock-in regime.

Three questions follow — which regime, which portability posture, which fleet composition — and each carries a downstream cost: the porting project that delays revenue, the realized-MFU gap that erases a price advantage, the kernel you write yourself because no library shipped it, a model whose NVIDIA path ships before a required operation is supported on the alternative release. Ignore the software stack at purchase and you do not avoid the bill; you defer it to the day the cluster is live and the delivered goodput misses the acceptance target.

These four ecosystems are not "CUDA, plus three challengers." Each locks a different layer of the stack, erodes at a different rate, and charges a different exit cost — so the useful way to sort them is by what each one actually holds captive.

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 subtler than 'you cannot leave': on NVIDIA the kernel you need already exists and is tuned, and everywhere else you write it 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. The lock is not the absence of a kernel language — JAX's Pallas does let you hand-write TPU kernels, as AWS's Neuron Kernel Interface does on Trainium. 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 reflects software, memory, fabric and the model. 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. For a fixed model, precision, FLOP-count convention and device count, compare useful model FLOPs with the matching dense peak over the same wall-clock. No universal 35–55% band classifies a stack as mature. Profile the difference in 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 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.3x the paper FLOPS of an H100 but, on real inference, showed workload-specific shortfalls on the tested kernels and serving stack at the reported model, precision and release (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 hardware-price advantage requires matched quotations, then useful output, porting and recurring support on the same boundary. In the stated illustration, a chip priced 20% lower but delivering 40% less matched useful output costs 0.8/0.6 ≈ 1.3 times as much per unit of output before other costs. It breaks even only when its useful output reaches 80% of the reference; below that, the cheaper chip buys a more expensive cluster. Compare the full cash flows in Chapter 1.8 — 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 complete service costs less, it 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: idiomatic PyTorch can target CUDA, ROCm and, through PyTorch/XLA, TPU, but operator coverage, numerical behavior, compilation, performance and distributed-runtime support must be validated per backend. 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 — offers multiple backends, but a kernel is portable only across the specific devices, operators, layouts, numerical behavior and performance targets validated in CI, 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 Triton kernel can reduce porting work, though it is not automatically portable or performance-equivalent across vendors; the same kernel written in CUDA C++ with wgmma and tensor-memory-accelerator intrinsics is not portable at all. 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 only actionable once 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.

Add operator, collective, precision and release-engineering hours; multiply by the loaded hourly rate for the initial port. Annual sustaining is monthly hours × billing months × rate. First-year surplus is gross monthly saving × months minus port and sustaining. Dividing those two costs by billing months gives the gross-saving flip threshold; keep the untouched baseline and tuned manifest as separate runs.

Fund the bounded port in this scenario, with acceptance at the complete tuple. Release requires the quality, p99 latency, collective scale, checkpoint/recovery and rollback gates to pass on that manifest; an API-compatible NCCL/RCCL call is not that evidence. Below the stated saving threshold, keep the incumbent unless a separately funded continuity requirement pays the difference.

Method: AMD ROCm documentation and compatibility guidance. Chapter 7.14 owns the next handoff.

Fleet composition is a strategy decision that reaches you disguised 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 at the inference layer, where an OpenAI-compatible API keeps the calling application unchanged across the swap and the porting, throughput-recovery and operating costs separated above are what you actually pay. It is least defensible at the frontier-training layer, where day-0 coverage and the deepest scale-up domain still favor a single vendor.

What forces this question onto every roadmap is the forecast rise of custom ASICs toward roughly a quarter of 2026 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. Whether everyone smaller can capture the same upside, or whether the operational-complexity cost eats it, comes down to their workload mix and the depth of their engineering bench. → the merchant-silicon disruption in Chapter 7.5; the qualified configuration from Chapter 7.11 enters the lifecycle cost model in Chapter 1.8.