--- source: 'https://howaiworks.ai/glossary/cuda' section: glossary title: CUDA description: >- NVIDIA's software layer for GPUs — the programming model that turned graphics cards into AI hardware, and the real reason rivals with comparable silicon lose. tags: - CUDA - NVIDIA - GPU - GPU Computing - parallel processing - AI Infrastructure - Hardware category: Infrastructure & Hardware datePublished: '2026-07-13' lastUpdated: '2026-07-24' --- # CUDA > NVIDIA's software layer for GPUs — the programming model that turned graphics cards into AI hardware, and the real reason rivals with comparable silicon lose. ## Definition **CUDA is NVIDIA's software layer for its GPUs — and it is the reason competitors keep shipping chips with comparable specifications and losing anyway.** The hardware is only half of an accelerator. The other half is a kernel for every operation your model performs, tuned for the specific chip, and CUDA is where twenty years of those kernels already live. The programming model itself is simple enough to state in a sentence. On a CPU you write a loop that runs a million times. On a GPU you write **the body of the loop once** — a function called a *kernel* — and tell the hardware to run a million copies of it simultaneously, each one working out from its own thread index which slice of the data it owns. NVIDIA shipped CUDA 1.0 in [2007](https://insidehpc.com/2007/07/nvidia-releases-cuda-10/), when the point was scientific computing and nobody was training neural networks. Why a GPU suits AI at all is the subject of [GPU computing](https://howaiworks.ai/glossary/gpu-computing); this page is about the layer on top, and about why it has proved so much harder to replace than the silicon underneath it. ## How It Works ### Threads, blocks and warps A kernel launch creates a **grid** of thread **blocks**, and each block contains up to [1,024 threads](https://docs.nvidia.com/cuda/cuda-programming-guide/01-introduction/programming-model.html). A block is scheduled onto a single streaming multiprocessor (SM) and stays there: its threads can synchronise with each other and share a scratchpad of on-chip memory. Threads in *different* blocks cannot, which is the price of being able to schedule blocks in any order across any number of SMs. Underneath the block, the hardware imposes a unit the programmer never asked for: the **warp**, a group of **32 threads that execute the same instruction at the same time**. This is why a branch is expensive. If sixteen threads of a warp take the `if` and sixteen take the `else`, the warp does not split — it runs *both* paths in sequence, with half the lanes masked off each time. A branch on `threadIdx` can therefore double the cost of a region of code while the profiler still reports the GPU as busy. ### The memory hierarchy is the entire craft Here is the fact that everything else follows from. Measured on Hopper-generation hardware, the latency of a memory access depends on where the data is by more than an order of magnitude ([Luo et al., 2024](https://arxiv.org/abs/2402.13499)): | Where the data is | Latency | |---|---| | Shared memory (on-chip scratchpad) | **29 cycles** | | L1 cache | 41 cycles | | L2 cache | 263 cycles | | Global memory (HBM) | **479 cycles** | Bandwidth tells the same story. Shared memory sustains about [128 bytes per clock per SM](https://arxiv.org/abs/2402.13499). An H100 SXM has [132 SMs, each with 128 FP32 cores](https://developer.nvidia.com/blog/nvidia-hopper-architecture-in-depth/), and its shipping rate of 67 TFLOPS of vector FP32 therefore implies a clock just under 2 GHz — which puts aggregate on-chip bandwidth around **33 TB/s**, against [3.35 TB/s](https://www.nvidia.com/en-us/data-center/h100/) of HBM3 on the same card. **Ten times the bandwidth, one-sixteenth the latency, if you can get the data on-chip and keep it there.** That whole derivation is pinned to one chip — an **H100 SXM**, the Hopper-generation part (2022) that was the default AI accelerator while most of these kernels were being tuned, which is why it anchors this page. The absolute has since moved: NVIDIA's Blackwell **B200** (2024) and AMD's Instinct **MI355X** (2025) both raise HBM bandwidth to roughly **8 TB/s** ([NVIDIA](https://www.nvidia.com/en-us/data-center/dgx-b200/); [AMD](https://www.amd.com/en/products/accelerators/instinct/mi350/mi355x.html)) — about 2.4× the H100. The number worth carrying forward is not the 33 TB/s or the 3.35: it is the **ratio**. On-chip memory stays roughly an order of magnitude faster, and far lower in latency, than HBM in *every* generation, because it sits physically closer to the compute — and that gap, not any one datasheet figure, is what the craft exploits. That is why kernel authors think in **bytes moved, not operations performed**. A fast kernel is one that moves few bytes per unit of arithmetic — the framing the [memory wall](https://howaiworks.ai/glossary/memory-wall) develops properly — and every technique in the CUDA toolbox (tiling, coalescing, fusion, register blocking) is a different way of not going back to HBM. ### Worked example: a 4096 × 4096 matrix multiply Why this one operation dominates AI workloads — and the tensor cores built to accelerate it — belongs to [matrix multiplication](https://howaiworks.ai/glossary/matrix-multiplication). Here it is simply the cleanest place to watch bytes beat operations. Take the naive kernel that everyone writes first: one thread per output element, each thread reading a full row of A and a full column of B straight from global memory. - Global loads: `2 × N³` = 2 × 4096³ ≈ **137 billion** four-byte elements ≈ **550 GB**. - Time floor at 3.35 TB/s of HBM3: **164 ms**. - Useful arithmetic: also `2 × N³` ≈ 137 [GFLOP](https://howaiworks.ai/glossary/flops), which at the H100's [67 TFLOPS](https://www.nvidia.com/en-us/data-center/h100/) of vector FP32 is **2.1 ms**. The kernel is correct, and it spends roughly **80 times longer waiting for memory than computing**. Both floors here are Hopper-generation figures; a Blackwell or an AMD Instinct card with roughly 8 TB/s of HBM lowers each of them, but a naive kernel stays memory-bound by a large factor on any of them — the shape is a property of the algorithm, not the datasheet. Now tile it: each block cooperatively loads a 32 × 32 tile of A and of B into shared memory, and every thread reads each staged element **32 times** from on-chip memory instead of from HBM. Traffic drops by a factor of 32 — 550 GB becomes **17 GB**, and the memory floor becomes **5.1 ms**. Note where that leaves us. 5.1 ms of memory time still exceeds 2.1 ms of compute, so even the textbook tiled kernel is memory-bound. Real GEMM kernels tile a *second* time, from shared memory into registers, and a third time across warps — which is precisely what cuBLAS and CUTLASS have been doing, re-tuned for every architecture, for over a decade. That is the honest reason nobody hand-writes a matrix multiply. ### What breaks if you get this wrong A kernel that ignores the hierarchy is not broken. It is *correct and thirty times too slow*, which is far worse, because nothing alerts you. GPU "utilisation" counts a warp stalled on a memory load as active, so a profiler can report 95% utilisation on a kernel that is idling in front of HBM. The same mistake is made at the purchasing level with more zeros attached, and it is not hypothetical. AMD's Instinct accelerators lead the contemporaneous NVIDIA part on the datasheet — the [MI355X](https://www.amd.com/en/products/accelerators/instinct/mi350/mi355x.html) carries 288 GB of HBM against the H100's 80 GB, at comparable-or-higher bandwidth — yet a chip can top every [FLOPS](https://howaiworks.ai/glossary/flops) and bandwidth ceiling and still lose by a large factor on a real workload, because those ceilings are reached only by kernels tuned for that specific chip, and on AMD's ROCm stack the kernels for the newest operations arrive later, if at all. **You buy silicon; you rent performance from whoever wrote the kernels.** ## Real-World Applications **Almost nobody writes CUDA.** This is the most important practical fact about it, and it sounds like a contradiction until you look at what actually runs. Your `model.to("cuda")` in PyTorch calls into **cuBLAS** for matrix multiplies, **cuDNN** — shipped in [September 2014](https://www.infoq.com/news/2014/09/cudnn/), when NVIDIA first noticed deep learning was going to matter — for convolution and attention primitives, and **NCCL** for the all-reduce that synchronises gradients across a training cluster. NCCL is the unglamorous one and the hardest to replace: at 10,000 GPUs, a training run is bounded by how fast the interconnect can sum gradients, not by the arithmetic. The exceptions are where the money is. **FlashAttention** was, at bottom, a hand-written CUDA kernel that fused the attention computation so the intermediate score matrix never touched HBM — and it moved the economics of long context windows for the whole industry. Inference stacks like TensorRT-LLM and vLLM are, similarly, collections of CUDA kernels: paged-attention kernels that manage the [KV cache](https://howaiworks.ai/glossary/kv-cache), and low-precision kernels that make [quantization](https://howaiworks.ai/glossary/quantization) pay off in wall-clock latency rather than just in memory footprint. A quantized model with no fast kernel for its data type is smaller *and slower*. The pattern is consistent: a few hundred people write kernels, and everyone else in AI stands on them without ever seeing them. The kernels are the moat. The language is not. ## Key Concepts - **Coalescing**: when the 32 threads of a warp read 32 *adjacent* addresses, the hardware fuses them into a handful of memory transactions; when they read scattered addresses, it issues up to 32 separate ones. Same instruction, same data volume, an order of magnitude apart in time. - **Occupancy**: how many warps an SM has resident at once. Warps are how the GPU hides memory latency — while one waits on HBM, another computes — so a kernel that grabs too many registers or too much shared memory throttles its own ability to hide the 479-cycle stall. - **Fusion**: two separate passes that each read a tensor from HBM and write it back move it twice; merged into one pass, it moves once. Adding arithmetic to save a round trip is usually a win, and this — not clever maths — is most of what `torch.compile` does. - **A virtual instruction set**: CUDA compiles to PTX, an abstract ISA, which is lowered to the real per-architecture assembly (SASS) only at install or load time. It is why a binary built years ago still runs on a new GPU, and forward compatibility is a large, invisible part of the lock-in. ## Challenges **The moat is not the transistors, and it is not the language either.** CUDA is a modest extension of C++; a competent team could clone the syntax in a year. What cannot be cloned in a year is the accumulated stack behind it: cuBLAS, cuDNN, NCCL, CUTLASS and hundreds of CUDA-X libraries, each re-tuned for every architecture since 2007; the profilers and debuggers; and — the piece people underrate — **the fact that the reference implementation of every ML paper is a CUDA repository**. A new technique arrives as a CUDA kernel. It gets ported elsewhere months later, if anyone bothers. That is a treadmill a competitor has to run on forever, always one release behind. The reinforcing loop is the real problem. Fewer users means fewer kernel authors, which means slower kernels, which means fewer users. NVIDIA says CUDA now underpins [six million developers](https://forums.developer.nvidia.com/t/20-years-of-cuda-honoring-the-architects-of-the-accelerated-age/364654); the number is NVIDIA's own and should be read as marketing, but the direction it points is not in dispute. Two things cut the other way, and honesty requires them. First, a software moat is expensive but not physical: unlike a fab or a patent, it can be eroded by anyone willing to spend enough engineering years, and Google's [TPU](https://howaiworks.ai/glossary/tensor-processing-unit) stack is the standing proof that it can be done — as well as a demonstration of how much it costs. Second, the lock-in binds NVIDIA too. It must carry compatibility forward across every generation, and a competitor starting today gets to design for [transformers](https://howaiworks.ai/glossary/transformer) rather than for the graphics workloads CUDA was originally shaped around. ## Future Trends **AMD's ROCm**, the open stack behind its Instinct accelerators, is the most direct challenge, and it is further along than the discourse suggests. Its HIP layer is a source-level CUDA-compatibility interface — PyTorch's own build system mechanically translates its CUDA backend into HIP, so ROCm PyTorch exposes the same `torch.cuda` API and most user code runs unmodified. The gap is not the common path; it is the frontier. The newest attention kernels and inference libraries land on CUDA first and arrive on ROCm later, and a compatibility layer chasing a moving API is structurally always second. **OpenAI's Triton** attacks from a different angle. Released in [July 2021](https://openai.com/index/triton/), it lets you write a kernel in Python at the level of *blocks and tiles*, leaving coalescing, shared-memory staging and thread scheduling to the compiler. Triton does not beat CUDA on NVIDIA hardware; that is not the point. It removes the reason to care which vendor you are on, because the same Triton source can target multiple backends. **And that is the mechanism that would actually erode the moat.** PyTorch 2.0 ([March 2023](https://docs.pytorch.org/get-started/pytorch-2.0/)) made TorchInductor the default compiler behind `torch.compile`, and TorchInductor generates **Triton** kernels rather than calling hand-written CUDA ones. If the kernel your model runs is generated at compile time from framework IR, then the vendor-specific layer becomes a *backend* — and backends can be swapped. That decouples the framework from the vendor, which is the only thing that has ever dissolved a software moat. So rather than predict a winner, state the conditions. A rival stack wins when: **(1)** its compiler path reaches parity on the operations that dominate real workloads — attention, GEMM, collectives — not on average, but on the *newest* operation, within weeks of it being invented; **(2)** it has a collectives library as good as NCCL, because at cluster scale training is bound by interconnect, not arithmetic; and **(3)** its hardware is buyable in volume, since a stack nobody can get hold of attracts no kernel authors, and the loop closes again. None of those three is impossible. None of them is close to free, and they have to arrive together. ## Code Example The tiled matrix multiply from the worked example above. The thread/block structure and the `__shared__` staging buffer are the whole of CUDA in one screen. ```cuda #define TILE 32 // 32 x 32 = 1,024 threads = 32 warps = the maximum block size // C = A * B, all N x N and row-major. One thread computes one element of C. __global__ void matmulTiled(const float* A, const float* B, float* C, int N) { // On-chip scratchpad, ~29 cycles away, shared by all 1,024 threads of this // block. This declaration is the only reason the kernel is fast. __shared__ float As[TILE][TILE]; __shared__ float Bs[TILE][TILE]; int row = blockIdx.y * TILE + threadIdx.y; // which slice of the data I own int col = blockIdx.x * TILE + threadIdx.x; float acc = 0.0f; for (int t = 0; t < N / TILE; ++t) { // Each thread pulls exactly ONE element of each matrix from HBM... As[threadIdx.y][threadIdx.x] = A[row * N + t * TILE + threadIdx.x]; Bs[threadIdx.y][threadIdx.x] = B[(t * TILE + threadIdx.y) * N + col]; __syncthreads(); // ...and waits for the other 1,023 to do the same. // ...then reads it back 32 times from shared memory, not from HBM. // That factor of 32 is the entire optimisation. for (int k = 0; k < TILE; ++k) { acc += As[threadIdx.y][k] * Bs[k][threadIdx.x]; } __syncthreads(); // don't overwrite a tile others are still reading. } C[row * N + col] = acc; } // Host side: 128 x 128 blocks of 1,024 threads each = 16.8 million threads. dim3 block(TILE, TILE); dim3 grid(N / TILE, N / TILE); matmulTiled<<>>(dA, dB, dC, N); ``` Note what the kernel never says: it does not say *how many* GPUs, *which* GPU, or how the 16.8 million threads are scheduled. That is the abstraction CUDA sells, and the reason it has been so difficult to displace. ## Frequently Asked Questions ### What is CUDA in simple terms? CUDA is NVIDIA's programming model for its GPUs. Instead of writing a loop that runs a million times, you write the body of the loop once and tell the GPU to run a million copies of it at once. Everything above it — PyTorch, cuBLAS, cuDNN, every AI framework you have heard of — is built on that one idea. ### Do I need to learn CUDA to do machine learning? Almost certainly not. Practically all AI work goes through PyTorch or JAX, which call pre-tuned CUDA libraries for you; a handful of specialists write the kernels that everyone else calls. Learning CUDA matters when you are optimising inference or implementing an operation the libraries do not have. ### Why can't AMD just copy CUDA? It largely has: AMD's ROCm includes HIP, a source-level CUDA-compatibility layer, and PyTorch ships ROCm builds that run most code unmodified. The gap is not the API, it is the twenty years of tuned libraries behind it and the fact that every new technique appears as a CUDA kernel first — so a compatibility layer is always chasing a moving target. ### Is CUDA the same thing as a GPU? No. The GPU is the chip; CUDA is the software layer that lets general-purpose programs use it. A competitor can build a chip with comparable FLOPS and HBM bandwidth and still lose badly, because performance comes from kernels that keep data close to the compute — and those kernels have to be written for each stack. ### What is a warp? A group of 32 threads that the hardware executes in lockstep on the same instruction. It is the real unit of GPU execution, and it is why a branch that splits a warp costs you: both sides of the branch run, one after the other, with half the lanes idle each time. ## Related ### Related terms - [GPU Computing](https://howaiworks.ai/glossary/gpu-computing) - [NVIDIA GPU for AI](https://howaiworks.ai/glossary/nvidia-gpu-ai) - [Parallel Processing](https://howaiworks.ai/glossary/parallel-processing) - [Deep Learning](https://howaiworks.ai/glossary/deep-learning) - [Tensor Operations](https://howaiworks.ai/glossary/tensor-operations) - [Inference](https://howaiworks.ai/glossary/inference) --- Source: https://howaiworks.ai/glossary/cuda — HowAIWorks.ai