Software Ecosystems & Lock-In

Chapter 7.1 through Chapter 7.8 compared accelerators on hardware: HBM capacity, FP4 PFLOPS, NVLink bandwidth, dollars per GPU. That comparison is incomplete in a way that costs real money. An accelerator also commits you to a runtime, a kernel language, a collective library, a compiler, a serving stack and a fabric. That stack 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. 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, or 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 holds captive.

CUDA is the deepest and broadest moat, and it is mostly a library and tooling moat rather than 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), the profilers and debuggers, and the fact that nearly every new model, attention variant and 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. As of 2026 it has gone 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 and rocBLAS for the math, RCCL for collectives. 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. The 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. The lock is not the absence of a kernel language: JAX's Pallas lets 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, so the lock-in is that your code is written in JAX against XLA's compilation model rather than in a portable imperative GPU style. The upside is that XLA is genuinely multi-target, since it compiles to GPU too, so the regime locks the programming model more than the silicon. The downside is that the ecosystem of pretrained checkpoints, third-party kernels and community tooling around JAX/TPU is narrower than CUDA's. Supply is also single-source. TPUs are mostly rented through Google Cloud; since Q2 2026 Google also delivers TPU systems to select customers' own data centers.

Neuron (AWS's SDK for Trainium and Inferentia) is the most vertically integrated of the four and the most tied to one supplier. Neuron compiles PyTorch or JAX down to the Trainium ISA. PyTorch reaches it through the native TorchNeuron backend, available on Trn2 and Trn3 since Neuron 2.27 (December 2025). The lock-in is the whole stack plus the fact that Trainium is available only from AWS, in its regions or as AWS-operated AI Factories in a customer's own data center. The trade is explicit: AWS reports customers cutting training and inference costs by up to 50% (a vendor figure with no stated baseline), 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. The gap reflects software, memory, fabric and the model. Model-FLOPS utilization is the useful floating-point work your model actually performs (its model FLOPs) divided by the chip's matching dense peak over the same wall-clock. Fix the model, precision, FLOP-count convention and device count before you compare two chips. Then profile the difference in the kernels, the fused attention path, the collective-communication overlap and the compiler's ability to keep the tensor cores fed.

This is where ROCm's story changed 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 inference benchmark, May 2025). By mid-2026 the gap was closed for some models and still open for others. In May 2026, MI355X running GLM-5 on ROCm and SGLang matched B200's recipe for FP8 KV cache and the same MLA path. FP4 with MTP was not yet at parity, and MI355X had no disaggregated or wide-EP recipe. Across roughly 10 to 77 tok/s per user, MI355X was cheaper per million tokens by up to about 29%. At the point where its lead was widest, about 18 tok/s per user, B200 cost 1.41 times as much: $0.30 against $0.22 per million tokens, rounded. B200 pulled ahead at about 90 tok/s per user (SemiAnalysis InferenceX, May 2026). SemiAnalysis's InferenceX dashboard in October 2026 set the best published run on each chip side by side, in an agentic scenario. On that view, Qwen3.5 397B ran 58% cheaper per token on MI355X than on B200, and DeepSeek V4 Pro ran 51% dearer. The lesson is not "AMD won" or "NVIDIA won." 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 only counts if the quotes are matched and useful output, porting and recurring support are measured on the same boundary. In the illustration, a chip priced 20% lower that delivers 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, in 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.

If lock-in is a contract, portability layers are the clause that lets you exit, and you pay for that 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 you make it years before you would ever switch. Three layers matter in 2026, at three different altitudes.

The framework layer (PyTorch / JAX). The 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 and ROCm. It reaches TPU today through PyTorch/XLA. In April 2026 Google introduced TorchTPU, a native, eager-first PyTorch backend with torch.compile and DTensor support; the PyTorch/XLA project says it will replace PyTorch/XLA once it is public, so check its release status and operator coverage for your model. Trainium is reached through TorchNeuron (Neuron 2.27, December 2025). Operator coverage, numerical behavior, compilation, performance and distributed-runtime support must still be validated per backend. That 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 its 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. Since PyTorch 2.0 it has been the default code generator inside torch.compile's TorchInductor. The project now lives in the community triton-lang repository, with major contributions from NVIDIA, AMD, Intel and Meta. NVIDIA and AMD GPUs are supported in tree, and Intel maintains its own GPU backend. A custom Triton kernel can reduce porting work, but 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 and SGLang (which run on both CUDA and ROCm) and NVIDIA's TensorRT-LLM all expose a near-identical OpenAI-compatible request API, so the application above never learns which silicon it is talking to. That 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. The number is the total cost of switching regimes, which is not one cost but four, paid at different times.

The porting cost is the 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. The 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 today's switching cost.

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. Between first light and tuned throughput you pay for silicon at a fraction of its peak while engineers chase the missing 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 window is measurable. On DeepSeek V4 (April 2026), MI355X ran at 1–2 tokens per second per user on day 0 and improved more than 100x within a month. Over the same period B200 gained about 1.7x in tokens per megawatt from software alone, at 50 tok/s per user (SemiAnalysis, June 2026). A benchmark taken in a model's first month measures the software date more than the silicon. Record the software date with every result, and re-test a new model within weeks.

The day-0 coverage cost is the opportunity cost of not being on the regime where new models land first. Most new architectures, attention variants and quantization schemes ship and are tuned on CUDA first. The exceptions are labs that co-develop with another stack: Google on TPU, and DeepSeek V4, which launched day-0 on Huawei's CANN as well as CUDA. On ROCm, Neuron and XLA, third-party models 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 and firmware matrices, two sets of failure modes, two on-call playbooks. It is paid every day, not once, and it is why many operators stay single-vendor despite a price advantage elsewhere. The second toolchain needs an owner, and that owner is a headcount.

The port takes 12 × 8 + 40 + 40 + 24 = 200 hours, or $30,000 at $150/h. Sustaining costs 20 h × 12 months × $150/h = $36,000 in the first year. Twelve months of $8,000 savings bring in $96,000, leaving a first-year surplus of $30,000 after both costs. The two costs total $66,000, so the decision flips at a gross saving of $5,500/month; at $5,000/month the first year loses $6,000. Benchmark the untouched baseline separately from the tuned manifest. The hours engineers spend closing the gap between the two are the realized-MFU recovery cost from earlier in this chapter. Count them as a switching cost on top of the port, not as a ROCm tax to be written off.

In this scenario, fund the port, and accept it only on the complete combination of hardware, software release and model. The release has to pass quality, p99 latency, collective scale, checkpoint recovery and rollback on that exact manifest; an API-compatible NCCL/RCCL call proves none of them. Below the $5,500 threshold, keep the incumbent unless a separately funded continuity requirement pays the difference.

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, because 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).

Single vendor no longer guarantees a single toolchain. Alongside the GPU racks, NVIDIA's Vera Rubin platform now includes Groq 3 LPX racks for latency-sensitive decode, each holding 256 SRAM-only LPUs (a vendor spec). NVIDIA took a non-exclusive licence to Groq's technology in December 2025, announced full production of LPX in August 2026, and uses its Dynamo software to split the work between GPUs and LPUs. The LPUs do not run CUDA natively. NVIDIA positions them as an accelerator attached to the CUDA stack on NVL72. Put LPX racks in an all-NVIDIA fleet and you run a second silicon family with its own compiler and its own failure domain, all bought from the same vendor, so it adds work without adding price leverage. Before sizing decode around LPX, find out which models and serving frameworks that path supports.

A heterogeneous fleet inverts both exposures, price and supply. 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. The cost is the operational complexity 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 strategy is most defensible at the inference layer, where an OpenAI-compatible API keeps the calling application unchanged across the swap, and where the porting, throughput-recovery and operating costs separated above are what you actually pay. It is least defensible at frontier training, where day-0 coverage and the deepest scale-up domain still favor a single vendor.

The forecast rise of ASIC-based systems to about 27% of 2026 AI-server shipments (TrendForce, April 2026) puts this question on every roadmap. 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, depends on 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.