CUDA

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, 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; 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. 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):

Bandwidth tells the same story. Shared memory sustains about 128 bytes per clock per SM. An H100 SXM has 132 SMs, each with 128 FP32 cores, 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 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 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 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. 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, which at the H100's 67 TFLOPS of vector FP32 is 2.1 ms.

The kernel is correct, and it spends roughly 80 times longer waiting for memory than computing. 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. A chip can carry more FLOPS and more HBM bandwidth than an H100 on paper and still lose by a large factor on a real workload, because the numbers on the datasheet are ceilings and the kernels that reach them have not been written for that chip. 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, 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, and low-precision kernels that make 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; 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 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 rather than for the graphics workloads CUDA was originally shaped around.

Future Trends

AMD's ROCm 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, 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) 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.

#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<<<grid, block>>>(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.