CUDA vs ROCm: The Gap Is Software, and HIP Moves the Cost

CUDA vs ROCm: the gap is mostly software maturity and tuning effort, HIPIFY moves switching cost rather than removing it, and one benchmark cannot decide.

CUDA vs ROCm: The Gap Is Software, and HIP Moves the Cost
Written by TechnoLynx Published on 13 May 2026

The CUDA vs ROCm choice comes down mostly to software maturity and tuning effort, not silicon. AMD’s ROCm stack has HIP as its CUDA-like programming interface, and the HipKittens authors (Stanford and AMD) write that AMD GPU hardware now offers state-of-the-art peak compute and memory bandwidth. In SemiAnalysis’s MI300X training tests, results on public AMD software releases still trailed NVIDIA because of weak software and testing. HIPIFY automates CUDA-to-HIP translation, but not all of it, and converted code still needs manual review and rework. No single benchmark settles the choice either, because the winner changes with the workload, the latency target and the software release. Profile your own workload on both stacks, then decide.

AMD’s equivalent of CUDA is a stack, not a single API

ROCm is the platform: compilers, libraries, tools and drivers, mostly open source. HIP is the C++ runtime API and kernel language inside ROCm, and it was designed to mirror CUDA closely. API calls are named almost the same (hipMalloc() where CUDA has cudaMalloc()), and the kernel syntax keywords are identical.

The rough mapping, component by component:

NVIDIA stack AMD counterpart What to know
CUDA platform ROCm Compilers, libraries, tools and drivers, mostly open source
CUDA C++ runtime and kernel language HIP Mirrors CUDA; HIP code can also target NVIDIA GPUs through a CUDA backend
nvcc hipcc Both compile C++ kernels down to vendor assembly
cuBLAS / cuBLASLt rocBLAS / hipBLASLt Two AMD GEMM libraries, with hipBLASLt the more optimised for MI300X in SemiAnalysis’s testing
cuDNN AITER SemiAnalysis describes AITER as AMD’s cuDNN equivalent
NCCL RCCL RCCL is a fork of NCCL
(source migration) HIPIFY: hipify-clang and hipify-perl Automates CUDA-to-HIP translation, with limits covered below

Most CUDA libraries have a corresponding HIP library. Where a library ships both a roc and a hip version (rocSparse and hipSparse, for example), AMD recommends the roc version on AMD GPUs because it is optimised for AMD architectures. HIP code is also compiled for a specific AMD GPU architecture, and the binaries carry code tailored to that target, which plays the same role on AMD as targeting a CUDA compute capability plays on NVIDIA.

Why is CUDA often faster than ROCm on the same workload?

In the benchmarks below, ROCm vs CUDA performance gaps trace to software: library defaults, kernel maturity, testing coverage and user tuning. SemiAnalysis’s MI300X vs H100/H200 training benchmark concluded that the MI300X’s on-paper advantage was not realised because of weaknesses in AMD’s public release software and testing. In its BF16 matrix-multiply tests, the H100 and H200 reached roughly 720 TFLOP/s against a marketed 989.5, while the MI300X reached about 620 against a marketed 1,307, making it 14% slower despite the higher headline figure. That MI300X result already used a custom Docker image built by an AMD principal engineer plus environment flags that are not set by default.

The mechanism behind part of that gap is unglamorous. According to SemiAnalysis, cuBLASLt’s heuristics pick the right GEMM algorithm for most shapes on H100/H200 out of the box, while hipBLASLt and rocBLAS heuristics picked the wrong one for most shapes, pushing the tuning onto the user. One bug they reported had F.Linear (the path most PyTorch users take to GEMMs) running on unoptimised rocBLAS while torch.matmul used hipBLASLt, so two wrappers around the same operation performed differently.

Maturity also shows up in correctness, not just speed. SemiAnalysis’s inference benchmark found 25% of tested models failing accuracy tests on AMD, and attributed models scoring worse on ROCm to missing CI and numeric-accuracy kernels. Its CI coverage for SGLang on AMD had improved from none but was still far from NVIDIA’s. On training, the same analysts measured a roughly six-month gap between FlexAttention working on NVIDIA and on AMD.

At the kernel level, the HipKittens authors (Stanford and AMD) report that peak-performance AMD kernels are written in raw assembly by a handful of experts, and that AITER and PyTorch reach just 30% and 24% of state-of-the-art performance respectively on Llama GQA backward passes on MI355X. Another source notes that FlashAttention kernels are typically written in NVIDIA CUDA assembly first, with ROCm versions lagging in features and peak optimisation. None of this is unique to AMD’s history: it took two years after the H100 launched for peak-performance open-source attention kernels to appear in low-level CUDA/CUTLASS.

How much of the gap closes with optimisation effort

Part of it closes, but the effort is real and recurring. SemiAnalysis got MI300X training to within 75% of H100/H200 performance only with multiple AMD teams fixing bugs and a hand-built Docker image of about 60 commands that took around five hours to build. Training performance per TCO was worse on MI300X with public stable releases but changed with custom development builds.

PyTorch TunableOp automates one slice of that tuning work. It benchmarks every registered implementation of an operator such as a GEMM for the given inputs, across rocBLAS and hipBLASLt on ROCm, and caches the fastest. Only a ROCm TunableGemm is implemented; CUDA builds just route to the cuBLAS default. In SemiAnalysis’s MI300X tests, the prototype PYTORCH_TUNABLE_OPS flag took one to two hours per modern LLM, repeated whenever the code changed; they called it buggy and urged AMD to tune its GEMM libraries instead.

A study of performance portability across CUDA, HIP, SYCL, Kokkos, RAJA, OpenMP and OpenACC found CUDA best or within 3% of the best in eleven of fifteen cases on NVIDIA systems, and HIP best in most AMD cases, though SYCL beat HIP in five of ten AMD cases. Those were proxy scientific applications run “out of the box” without kernel grid or block tuning.

Does HIPIFY remove the switching cost, or just move it?

HIPIFY does not remove the switching cost, because some code still needs manual intervention. AMD’s own documentation says HIPIFY does not convert all CUDA code seamlessly, that CUDA or third-party libraries with no HIP equivalent cannot be translated, and that code optimised for NVIDIA GPUs may need rework to perform well on AMD. The recommended steps after conversion are a code review for functional correctness, replacing unsupported libraries or constructs, debugging, testing and optimising on the target AMD GPU.

The two tools behave differently. hipify-clang parses the code into an abstract syntax tree, so it needs correct CUDA code and a CUDA install with all headers; hipify-perl uses pattern matching, needs no CUDA install, and is easier to set up but less powerful. Small semantic differences survive translation too: hipMemoryType enum values differ from cudaMemoryType, and cudaMemoryTypeUnregistered has no HIP counterpart, so memory-type code must be mapped by hand.

When the CASS authors ran HIPIFY over their CUDA corpus, about 43.9% of files failed conversion and were discarded. They trace part of this to HIPIFY lacking mappings for WMMA, cuBLAS TensorOp functions and other low-level primitives, so kernels relying on Tensor Core or warp-level matrix operations fail to translate or compile. That figure describes one research corpus, and the missing mappings show where the tool runs out.

Converted code that compiles is not converted code that performs

XSBench in the performance-portability study illustrates this. Its HIP port was produced with HIPIFY and is identical to the CUDA version apart from syntax substitution. The kernel is longer and more complex in control flow than the others studied, and the authors conclude that HIP kernels translated directly from CUDA without further optimisation may not perform optimally on AMD GPUs, which have significantly smaller cache capacity per thread workgroup. In that case portable models such as SYCL matched or beat the vendor model with similar effort. The port was a syntax substitution; it did not guarantee optimal performance on AMD.

The Ginkgo linear algebra team’s account of adding a HIP backend makes the same point from production code: they had to extend HIP with missing features such as cooperative groups, paid a performance price compiling HIP for AMD architectures, and designed a shared code base to limit duplication. SemiAnalysis goes further and argues that building AMD libraries by HIPIFY-translating NVIDIA’s cannot match NVIDIA’s user experience. That is an analyst’s argument about strategy rather than a measurement.

PyTorch on ROCm hides most of this, until it does not

For teams working inside PyTorch, the porting effort is much smaller. PyTorch for HIP intentionally reuses the torch.cuda interfaces, so existing code and models need very few changes, if any. PyTorch generates HIP source from its CUDA backend with its own tooling, and ROCm support is upstreamed, though it runs on its own release cycle. ROCm’s documentation also lists GPU-accelerated PyTorch features not supported for given versions. Because both builds reuse the torch.cuda interfaces, see which CUDA version a PyTorch build reports before trusting an environment check.

The cost resurfaces at the edges of PyTorch. Bleeding-edge model repositories often bypass standard PyTorch APIs and hardcode custom CUDA C++ extensions that must be ported by hand. In InferenceMAX results, an H200 on NVIDIA’s TensorRT-LLM beat an MI325X that won against the same H200 when both ran vLLM. Multi-node features are another gap: SemiAnalysis found AMD lacking disaggregated prefill, smart routing and NVMe KV cache tiering in inference, and measured MI300X collectives at two to four times slower than InfiniBand for all-gather, all-to-all and reduce-scatter.

What about running CUDA binaries on AMD directly?

On the CASS assembly-translation benchmark, the compatibility layer ZLUDA scored 7.86% accuracy, which the authors attribute to it translating at the LLVM IR level and so missing backend optimisations like NVIDIA’s instruction scheduling and register allocation that exist only in SASS. We would treat a compatibility layer as a way to get code running for evaluation, not as evidence of production performance, because it swaps porting cost for runtime performance risk and a dependency outside both vendors’ support.

Why one CUDA-vs-ROCm benchmark is a weak basis for choosing

A single benchmark captures one model, one latency target and one software snapshot, and all three move the result. SemiAnalysis set out to find a simple answer in its inference comparison and instead found NVIDIA ahead on performance per dollar in some workloads and AMD ahead in others. For reasoning tasks, the MI325X beat the H200 by up to 20% in performance per dollar beyond 100 seconds of latency, while the H200 won handily below it. For chat and translation, the small perf-per-dollar gains of the MI325X were, in their judgement, not worth the pain of adopting ROCm.

The software stack flips results as well. In InferenceMAX’s comparisons, the MI355X on vLLM had a lower cost per million tokens than the B200 on vLLM for GPT-OSS 120B summarisation, and even beat the B200 on TRT-LLM below 225 tok/s/user, but lost above that interactivity level. The analysts behind it note that hardware changes arrive yearly while software improves in releases days apart; their own runs took 60 to 120 hours per GPU type while serving frameworks updated almost weekly.

Tuning access matters too. Serving frameworks expose so many under-documented flags that SemiAnalysis could not guarantee optimal configurations and relied on NVIDIA and AMD engineers for their best settings. They recommend reading relative performance across GPU types and scenarios rather than any single absolute datapoint, since micro-optimising one point tells you nothing about the whole curve. The performance-portability researchers add that a single small benchmark is unlikely to represent complex production applications. The same logic applies within one vendor, where a GPU model number is not a performance contract either.

Reading results within a matched software release, and how a published benchmark should be structured, is a measurement question in its own right.

Profile your workload, then choose the stack

Make the decision on your own workload’s profile, measured the way it will run. For serving, that means throughput against end-to-end latency per user as concurrency rises, not offline throughput under ideal conditions. For kernels, the portability study’s authors found that their ability to find bottlenecks depended heavily on profiling tools. Best-practice rules only make sense in the context profiling gives them.

How the evidence tends to lean, by situation (judgement, to be confirmed by your own profile):

Your situation Lean Why
Heavy use of custom CUDA extensions, day-1 research repos or Tensor Core intrinsics CUDA, unless a measured AMD advantage pays for the port Hand porting and missing HIPIFY mappings sit here
Standard PyTorch ops, single-node inference, latency-tolerant or throughput-bound ROCm is a serious candidate; measure on your model torch.cuda reuse keeps porting small, and AMD wins some perf-per-dollar regimes
Fast research iteration with frequent code changes CUDA, or budget recurring tuning time Prototype GEMM tuning on MI300X took one to two hours per code change in SemiAnalysis’s tests
Multi-node training or disaggregated serving CUDA today, re-check per release Collective performance and missing serving features were measured gaps
Codebase that must run on more than one vendor HIP or a portable model such as SYCL or Kokkos HIP targets both vendors; SYCL was the most consistent performance-portable model in one study

Whatever you pick, weigh how the API decision interacts with your three-year hardware roadmap. We would also put the cost of the port and its ongoing tuning into the same comparison as the hardware price; a perf-per-dollar win that needs a dedicated engineer to maintain is a different number.

Who can profile and port our workload between CUDA and ROCm?

TechnoLynx profiles GPU workloads on the stacks a team is weighing, ports CUDA code to HIP where the numbers justify it, and tunes the result on the target hardware, through its GPU Performance Engineering work. The first step is profiling your own model or kernels, so the CUDA vs ROCm decision rests on your workload rather than on someone else’s benchmark.

FAQ

Will ROCm ever catch up to CUDA?

The available evidence shows progress without parity: SemiAnalysis credits AMD with meaningful CI and developer-experience changes while still seeing lots of room for improvement. InferenceMAX already shows AMD ahead for some workloads and interactivity levels and NVIDIA ahead for others.

Are there any competitors to CUDA?

Yes: HIP mimics CUDA and can target both NVIDIA and AMD, SYCL is a different model supported on NVIDIA, AMD and Intel GPUs, and OpenMP, Kokkos and Alpaka also span all three. One survey counted more than 50 routes for programming a GPU.

Back See Blogs
arrow icon