AMD vs. NVIDIA: What Actually Differs — Why the Stack Is the Problem, Not the Hardware

• Introduction — The 10 Percent Left After HIPIFY Ports the 90 Percent
• Hardware — SM and CU, Tensor Core and Matrix Core
• Warp 32 vs. Wavefront 64 — The Single Most Expensive Line
• The Software Stack — CUDA, ROCm, and HIP
  • Library Mapping Table
• The Porting Path — What HIPIFY Handles and What It Doesn't
  • What Gets Ported Automatically
  • What You Have to Fix by Hand
  • Verification You Must Run After Porting
• What Actually Gets in the Way
  • 1. The Accumulated Kernel Ecosystem
  • 2. Lag in Framework Support and the Support Matrix
  • 3. The Volume of Information When Something Breaks
• When AMD Is the Reasonable Choice
• Closing — The Gap Isn't in the Silicon, It's in the Accumulated Kernels
• References

Introduction — The 10 Percent Left After HIPIFY Ports the 90 Percent

Porting a CUDA codebase to AMD tends to go like this. Run hipify-perl and most of the files convert in seconds. cudaMalloc becomes hipMalloc, __global__ stays the same, and kernel launch syntax is identical. The build passes, the small tests pass, and that's half a day.

The next two weeks is where the real time goes. Chasing down one kernel whose results are subtly wrong, you find a 32 hardcoded right after __shfl_xor. Stuck on a kernel running at half speed, you discover the shared-memory tile size was sized assuming a warp of 32. One file with inline PTX has to be rewritten from scratch.

This piece covers exactly what that 10 percent is, and — more importantly — what the ecosystem gap underneath it actually looks like. No cheerleading, no dunking, just structure. The reference points are ROCm 7.14.0 (released July 16, 2026), the vLLM v0.26.0 documentation, and AMD's official ROCm documentation.

Hardware — SM and CU, Tensor Core and Matrix Core

Start with the terminology mapping. The conceptual layers line up almost one to one.

You also have to specify the target architecture at build time on both sides. Just as NVIDIA takes -arch=sm_90, ROCm takes a value like gfx942. Check your own device's value with rocminfo.

rocminfo | grep gfx        # e.g., gfx942 (MI300 series), gfx950 (MI350 series)
rocm-smi                   # nvidia-smi equivalent. Utilization, temperature, power

Here's a rundown of the generations you'll run into in the market right now. The hardware figures below are compiled from vendor announcements and secondary reporting — I have not benchmarked them myself. For any purchasing decision, always check the vendor's official spec sheet.

What matters here isn't any individual number. It's the fact that the two companies' hardware lands in roughly the same class of memory capacity and bandwidth within a given generation. There have been generations where AMD led on capacity, and generations where NVIDIA caught up. There is no decisive gap on this axis.

So the gap that shows up in practice comes from somewhere else. That somewhere else is the rest of this piece.

Warp 32 vs. Wavefront 64 — The Single Most Expensive Line

Of all the hardware differences, this is the one that cuts directly into code. An NVIDIA warp is 32 threads. A wavefront on AMD's CDNA-family datacenter GPUs is 64. The RDNA family is 32.

The HIP documentation calls this out explicitly.

Code should not assume a warp size of 32 or 64, as AMD GPU architectures have different warp sizes. The warpSize built-in should be used in device code.

This one line creates three distinct classes of bugs in practice.

First, the iteration count in shuffle reductions. Let's revisit the warp reduction from an earlier post.

// Code written assuming NVIDIA. Reduces only half the values on AMD.
__inline__ __device__ float warpSum(float v) {
  for (int off = 16; off > 0; off >>= 1)          // 16 assumes a warp of 32
    v += __shfl_xor_sync(0xffffffffu, v, off);
  return v;
}

If the wavefront is 64, off needs to start at 32. Starting at 16 sums the first 32 lanes and the last 32 lanes separately, so only half the final result is correct. It compiles, it doesn't crash, it's just wrong. This is the kind of bug that eats the most time in a porting effort.

A portable version looks like this.

// HIP. Works correctly on both AMD and NVIDIA.
__device__ float waveSum(float v) {
  for (int off = warpSize / 2; off > 0; off >>= 1)
    v += __shfl_xor(v, off, warpSize);
  return v;
}

Second, the width of the lane mask. On NVIDIA, the active-lane mask is 32 bits, so 0xffffffff comes naturally. A 64-wide wavefront needs 64 bits. The HIP documentation flags one more specific trap here: on a 64-wide wavefront, shifting a 32-bit integer by more than 31 zeroes out the register. The fix is to use uint64_t for the lane mask.

Third, implicit assumptions baked into tile size and block size. If you thought of a block size of 256 as "8 warps" and sized a shared-memory partial-sum array to 8 slots accordingly, on AMD there are only 4 wavefronts, so only 4 slots get used and the logic breaks. Conversely, a block size that's only aligned to a multiple of 32 might not be a multiple of 64 on AMD, leaving half of the last wavefront idle.

The principle is singular. Never hardcode 32 or 64 as a constant in source. Use warpSize, or if you truly need a compile-time constant, keep it in one place and swap it per architecture.

The Software Stack — CUDA, ROCm, and HIP

Laid side by side, the layers look like this.

NVIDIA                          AMD
──────────────────────          ──────────────────────
PyTorch / JAX / vLLM            PyTorch / vLLM
   │                               │
cuBLAS, cuDNN, NCCL             rocBLAS, MIOpen, RCCL
   │                               │
CUDA Runtime API                HIP Runtime API
   │                               │
CUDA Driver                     ROCr runtime + ROCk kernel driver
   │                               │
NVCC → PTX → SASS               hipcc(LLVM) → AMDGCN ISA

HIP is deliberately shaped to look almost identical to CUDA. The cuda prefix in function names becomes hip, and kernel definitions and launch syntax are effectively the same. That's not an accident — it's a design goal. And HIP code compiles on NVIDIA GPUs too. That means HIP can be used as a thin layer on top of CUDA, which is exactly why libraries that need to support both vendors choose HIP as a single source.

There's something worth flagging about ROCm 7.14.0. The version number jumping from 7.2.4 to 7.14.0 looks like a typo, but it's real. AMD's release notes explicitly note a "version-number discontinuity starting with the 7.9.0 preview," and explain that ROCm is simultaneously moving to TheRock, a modular build-and-release system. The idea is a lighter core SDK with optional domain SDKs for AI, data science, and HPC installed on top.

That fact alone tells you something practical. ROCm's structure is still in motion. Installation steps and package names can change from one major release to the next, which means you have to check the version whenever you're reading the docs. CUDA is far more stable on this axis.

Library Mapping Table

This is where you need to understand AMD's particular dual-naming scheme. AMD's official documentation explains it this way: roc-prefixed libraries are native, high-performance implementations written in HIP and targeting AMD GPUs, while hip-prefixed libraries are portability wrappers that implement a CUDA-equivalent API. hipBLAS describes itself as a "marshalling library," and can sit in front of either rocBLAS or cuBLAS.

The choice is clear. If you're porting from CUDA, use the hip side. The call shapes match, so code changes are minimal. If you're writing new code targeting AMD primarily, use the roc side. There's one less layer, and you get direct access to AMD-specific features.

The bottom three rows of the table each have an empty cell. cuDNN, NCCL, and CUTLASS have no portability wrapper — only a differently named counterpart. That means the API shapes differ enough that you can't swap them without touching the source, though this isn't a problem for most users since deep learning frameworks absorb this layer on their behalf. If your code calls these libraries directly, though, that's where the real cost of porting concentrates.

Scoped to LLM inference specifically, the recently introduced AITER (AI Tensor Engine for ROCm) matters a lot. It's AMD's repository of kernels for LLM inference, and vLLM's ROCm installation docs include it in the build procedure. It occupies roughly the spot that FlashInfer or FlashAttention occupies on the NVIDIA side.

The Porting Path — What HIPIFY Handles and What It Doesn't

There are two tools.

In practice, teams often start with hipify-perl, because getting every CUDA header lined up across a large codebase is a hassle.

# Preview a single file (leaves the original untouched, prints the converted result only)
hipify-perl kernel.cu

# Convert an entire directory in place, leaving backups behind
find . -name "*.cu" -o -name "*.cuh" | xargs hipify-perl -inplace -print-stats

# Use the clang-based tool when you need accuracy
hipify-clang kernel.cu -- -I/usr/local/cuda/include

What Gets Ported Automatically

• Runtime API call names (cudaMalloc, cudaMemcpy, cudaStreamCreate, etc.)
• Kernel qualifiers and launch syntax (__global__, __device__, triple-angle-bracket launches)
• Built-in variables (threadIdx, blockIdx, blockDim)
• Most math built-ins
• Header includes

That covers roughly 90 percent of the code, by volume.

What You Have to Fix by Hand

• Inline PTX assembly. PTX is NVIDIA's virtual ISA. AMD has no counterpart. These sections have to be rewritten using HIP intrinsics or rewritten from scratch as AMDGCN inline assembly. This is the single most reliable time sink in a porting effort.
• Warp-size assumptions. Everything from the previous section. Automated conversion has no way to know whether a given 32 means "warp size" or something else entirely, so it leaves it untouched.
• Semantic differences in warp-level primitives. The mask argument on the __shfl_sync family, the return width of __ballot, and explicit synchronization semantics all differ. The names get ported, but the meaning needs verification.
• Calls into CUDA-only libraries. Things like cuDNN, CUTLASS, and cuBLASLt that have no portability wrapper.
• Code that uses the driver API. Anything calling low-level driver APIs directly, like cuModuleLoad.
• Every performance-tuning constant. Tile size, block size, unroll factor, pipeline stage count. These convert, but the optimal values differ. Porting and optimization are separate jobs — skip this step and you end up with "it runs, but at half speed."

Verification You Must Run After Porting

# 1. Numerical verification first. Performance comes after.
#    Focus especially on kernels that involve reductions.
pytest tests/ -k "reduction or attention or norm"

# 2. Re-measure the actual bottlenecks with AMD's profiler.
#    Don't just trust tuning results carried over from NVIDIA.
rocprofv3 --stats -- ./my_app
rocprofv3 --kernel-trace -- ./my_app

The AMD tools that correspond to Nsight Compute are rocprofv3 and the ROCm Compute Profiler. The concepts are the same: per-kernel time, memory throughput, cache hit rate, occupancy.

What Actually Gets in the Way

Here are the three reasons the real-world gap persists even though the hardware is comparable.

1. The Accumulated Kernel Ecosystem

This is the single biggest item. Whenever a new model architecture or a new quantization format shows up, the kernels for it come out in CUDA first, almost without exception. New FlashAttention variants, new MoE routing kernels, new low-precision GEMMs — all of it. AMD support follows weeks to months later.

Triton is meaningfully closing this gap. The Triton repository's third_party directory holds nvidia and amd backends side by side, so a kernel written in Triton compiles for both. As covered in an earlier post, recent custom attention kernels have tended to appear in Triton first, and if that trend continues, the gap keeps shrinking. Anything that still depends on hand-written CUDA kernels, though, still lags.

2. Lag in Framework Support and the Support Matrix

Here's one concrete example. The vLLM v0.26.0 ROCm installation docs state the following.

• It supports ROCm 6.3 and above, and prebuilt wheels are provided for ROCm 7.0 and ROCm 7.2.1.
• Supported GPUs are the MI200 series (gfx90a), MI300 (gfx942), MI350 (gfx950), the Radeon RX 7900 series, the RX 9000 series, and the Ryzen AI series.
• MI350 requires ROCm 7.0 or higher.

The lag is visible right here. The latest ROCm is 7.14.0, but vLLM's prebuilt wheels only go up to 7.2.1. To use the latest ROCm, you have to switch to a source build, which brings a procedure of building against the verified branches of Triton, FlashAttention, and AITER individually. The fact that the vLLM docs point you to a Dockerfile to check those branch values is itself evidence of how fragile this combination is.

On the NVIDIA side, the same task is usually one line: pip install vllm. That difference is pure software-engineering friction, and it eats a team's time.

The practical conclusion is clear. On AMD, using a verified container image is effectively the default path. vLLM publishes the official vllm/vllm-openai-rocm image on Docker Hub, and the rocm/vllm-family images AMD used to distribute have now folded into the official image. Deciding to build it yourself is, on its own, a decision to take on substantial ongoing maintenance cost.

3. The Volume of Information When Something Breaks

This is hard to measure but very real. The volume of results when you search an error message, the number of Stack Overflow answers, blog posts from people who hit the same problem, GitHub issues that already have an answer — CUDA overwhelmingly wins on every one of these.

Debugging time is a cost that goes straight into the project schedule. It doesn't show up on the hardware price sheet, but it's in the total cost of ownership.

When AMD Is the Reasonable Choice

Even after acknowledging everything above, there are clear cases where AMD is the reasonable choice. The conditions are specific.

First, serving a widely used model for inference. Serving mainstream models from the Llama, Qwen, or DeepSeek families with vLLM is already a well-worn path. As long as you use a verified container and deploy a standard model, most of the friction described above has already been hit and resolved by someone else.

Second, when memory capacity is decisive. Fitting a large model onto fewer GPUs reduces tensor-parallel communication, which buys both performance and simplicity on its own. AMD has had generations where it led on capacity, and in those windows, this is a real advantage.

Third, when you need procurement and pricing leverage. In large-scale deployments, simply having a second supplier is valuable in itself. Having an alternative you can put on the negotiating table is different from not having one, even if you never actually use it.

Fourth, teams that work only at the top of the stack. If you work purely on top of PyTorch and don't maintain any custom CUDA kernels, most of the porting cost simply doesn't exist in the first place.

Conversely, the cases where you should avoid AMD are just as clear.

• A codebase whose performance depends on hand-written CUDA kernels or inline PTX
• A research organization that needs to be first to use new model architectures and new kernels
• A small team that can't dedicate headcount to GPU infrastructure
• Code that calls cuDNN or CUTLASS directly sitting at the core of the system

One practical piece of advice on top of this: if you're writing a new kernel in CUDA right now, checking whether it can be written in Triton instead is the cheapest insurance you can buy. When you eventually do evaluate AMD, a Triton kernel just needs a recompile — a CUDA kernel becomes a porting project.

Closing — The Gap Isn't in the Silicon, It's in the Accumulated Kernels

Within a given generation, both companies' GPUs land at comparable compute throughput and comparable memory bandwidth. SM and CU are conceptually the same, and Tensor Cores and Matrix Cores do the same job. HIP offers an API nearly identical to CUDA's, and HIPIFY ports 90 percent of the code automatically. Looking only this far, there's seemingly no difference.

The difference lies in what's stacked up underneath. The total volume of kernels written and tuned for CUDA over the past 15 years, the libraries built assuming those kernels exist, the frameworks written assuming those libraries exist, and the accumulated debugging knowledge on top of all of it. That's an asset that doesn't get replicated overnight, and it explains most of the real-world performance gap.

And this is also where the direction for closing the gap comes from. The more kernels get written in a portable layer like Triton rather than a vendor-specific language, the less that accumulation piles up on only one side. The thicker the layer where the compiler absorbs vendor lock-in becomes, the more genuinely free the choice gets. That's the subject of the next post in this series.

If you have to decide right now, the question collapses down to one thing. Does our performance hinge on CUDA code we wrote ourselves, or on a library someone else wrote? If it's the latter, AMD is worth evaluating. If it's the former, you need to honestly estimate the porting cost first.

References

• ROCm official documentation: https://rocm.docs.amd.com/
• ROCm releases (check version and date): https://github.com/ROCm/ROCm/releases
• HIP porting guide (includes the warp-size warning): https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_porting_guide.html
• ROCm API library list (CUDA equivalence mapping): https://rocm.docs.amd.com/en/latest/reference/api-libraries.html
• HIPIFY repository and documentation: https://github.com/ROCm/HIPIFY
• AITER (AI Tensor Engine for ROCm): https://github.com/ROCm/aiter
• vLLM ROCm installation docs: https://docs.vllm.ai/en/latest/getting_started/installation/gpu.html
• AMD MI300X tuning guide: https://rocm.docs.amd.com/en/latest/how-to/tuning-guides/mi300x/index.html
• AMD Instinct product page (official specs): https://www.amd.com/en/products/accelerators/instinct.html
• NVIDIA data center GPU product page: https://www.nvidia.com/en-us/data-center/