From Fragment Shaders to Tensor Cores: How CUDA Actually Works ?

Your system has a CPU. It probably also has a GPU. They look similar on a spec sheet - both have "cores," both run at some clock speed, but they were built to solve completely different problems. Understanding why is the key that unlocks everything else about CUDA.

Why does this hardware even exist?

A CPU is latency-optimized. Its job is to finish one complex task as fast as possible, so a large fraction of its transistor budget goes toward things that have nothing to do with raw arithmetic: branch predictors that guess which way an if statement will go before it's even evaluated, out-of-order execution engines that reorder instructions to avoid stalls, and large multi-level caches (L1, L2, sometimes L3) that keep frequently used data close to the core. A CPU is built to handle unpredictable, branching, sequential logic.

A GPU is totally opposite of that of CPU. It assumes the workload is predictable, regular, and mathematically dense. The same instruction applied to enormous amounts of data, like every pixel on a screen or every element in a large matrix. Because of that assumption, a GPU strips away most of the branch prediction and out-of-order machinery a CPU relies on, and instead spends its transistor budget on raw arithmetic units -> thousands of small, simple cores instead of a handful of large, complex ones.

CPUs are optimized to finish one task fast. GPUs are optimized to finish millions of similar tasks at once.

Machine learning workloads like matrix multiplications, convolutions, attention mechanisms are, at the arithmetic level, the same handful of operations repeated across enormous tensors. That's exactly the kind of regular, parallel workload GPUs were built for, which is why they became the default hardware for training and running Neural Networks.

Before CUDA: computing by pretending to render graphics

It's easy to forget that GPUs weren't originally designed to run arbitrary programs at all. Before 2007, a GPU's entire reason for existing was the graphics pipeline: take a scene made of triangles, run a vertex shader on each vertex to figure out where it lands on screen, then run a fragment (pixel) shader on each resulting pixel to decide its final colour, and write that colour to a framebuffer.

Researchers noticed something important that a fragment shader is really just a small program that runs once per pixel, in parallel, across the whole screen. If you could trick the GPU into treating your data as if it were a picture, you could get that same massive parallelism applied to ordinary numerical computation. This workaround became known as GPGPU —> General Purpose computing on Graphics Processing Units and it required contorting real math problems into a form the graphics pipeline would accept:

• Input data (say, two matrices you wanted to multiply) had to be encoded as textures, because textures were the only large, structured input the GPU knew how to read.
• To actually trigger the computation, you had to render a full-screen quad two triangles covering the entire viewport, a purely as a pretext to force the GPU to invoke your fragment shader once per output element. The "image" being rendered had no visual meaning; it was a disguise to make the hardware execute your math.
• The fragment shader, written in a shading language meant for lighting and colour, had to be repurposed for arithmetic — reading inputs via texture lookups instead of normal memory reads, and outputting a numeric result as if it were a pixel colour.

A wildly simplified GLSL fragment shader "computing" a[i] + b[i] by disguising it as pixel colour math looked something like this:

// Old-style GPGPU: addition disguised as a fragment shader
uniform sampler2D textureA; // matrix A, encoded as a texture
uniform sampler2D textureB; // matrix B, encoded as a texture

void main() {
    vec4 a = texture2D(textureA, gl_TexCoord[0].xy);
    vec4 b = texture2D(textureB, gl_TexCoord[0].xy);
    gl_FragColour = a + b; // the "pixel colour" is actually your result
}

The result didn't come back as an array, rather it came back as an image sitting in a framebuffer, which then had to be read back out and reinterpreted as numbers. Control flow was extremely limited, loops and conditionals were weak or unsupported in early shader models, and there was no way to write to an arbitrary memory location only to the specific pixel a shader invocation happened to own.

NVIDIA released Compute Unified Device Architecture (CUDA) in 2007 specifically to remove this barrier, exposing the GPU's parallel execution units through a C-like programming model with real memory reads and writes, real control flow, and no requirement to pretend any of it was a picture. This is where GPU became a general purpose tool.

How a GPU organizes its execution

CUDA uses a model called SIMT — Single Instruction, Multiple Threads. One instruction is issued, and many threads execute it simultaneously, each on its own piece of data. This is organized into a hierarchy that maps directly onto the physical hardware:

• Thread is the smallest unit of execution, with its own program counter and private registers.
• Warp a group of exactly 32 threads, the actual unit the hardware schedules. All 32 threads in a warp are issued the same instruction on the same clock cycle.
• Block up to 1024 threads scheduled onto one Streaming Multiprocessor (SM). Threads in a block can synchronize with __syncthreads() and share fast on-chip memory.
• Grid every block needed to run the kernel, spread across all available SMs.

Here's what a minimal CUDA kernel actually looks like — adding two vectors, with one thread handling one element:

__global__ void vectorAdd(float *a, float *b, float *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

Compare that to the GLSL disguise above — real array indexing, real addition, no textures, no framebuffers.

Occupancy and latency hiding

An SM typically holds far more warps resident at once than it can execute in a single cycle. When a warp stalls, most often waiting on a slow global memory reads the SM's scheduler switches to a different warp that's ready to go, and the stalled warp resumes once its data arrives. This is called latency hiding, and it's the main reason GPUs don't need CPU-style caches and speculative execution to stay busy. The fraction of an SM's maximum warp capacity actually kept resident is called occupancy.

Warp divergence

Because all 32 threads in a warp share a single instruction fetch, they must step through the code together:

if (data[i] > threshold) {
    // path A
} else {
    // path B
}

If threads within the same warp evaluate this condition differently, the warp diverges. The hardware masks off the threads taking path B, executes path A for the active threads, flips the mask, executes path B, and reconverges once both paths finish. Work that could have run in parallel now runs sequentially . A simple 2 way branch can roughly half effective throughput for that warp.

Takeaway: uniform, branch-free workloads run efficiently on a GPU. Conditional, data-dependent logic that splits a warp's behaviour is where you lose parallel performance.

The real bottleneck: memory, not compute

A modern GPU can typically perform far more arithmetic operations per second than it can pull bytes from memory per second. In most real workloads, memory bandwidth is the limiting factor. This is formalized in the roofline model: a kernel's achievable performance is capped either by peak compute throughput or by peak memory bandwidth, whichever constraint that kernel's ratio of math-to-bytes hits first.

Efficient CUDA code minimizes round trips to global memory. A common pattern used heavily in matrix multiplication is tiling. Load a chunk of input data from global memory into shared memory once, let every thread in the block reuse that tile for all its arithmetic, and only write the final result back to global memory at the end.

__global__ void matMulTiled(float *A, float *B, float *C, int n) {
    __shared__ float tileA[16][16];
    __shared__ float tileB[16][16];

    int row = blockIdx.y * 16 + threadIdx.y;
    int col = blockIdx.x * 16 + threadIdx.x;
    float sum = 0.0f;

    for (int t = 0; t < n / 16; t++) {
        tileA[threadIdx.y][threadIdx.x] = A[row * n + t * 16 + threadIdx.x];
        tileB[threadIdx.y][threadIdx.x] = B[(t * 16 + threadIdx.y) * n + col];
        __syncthreads();

        for (int k = 0; k < 16; k++)
            sum += tileA[threadIdx.y][k] * tileB[k][threadIdx.x];
        __syncthreads();
    }
    C[row * n + col] = sum;
}

Each value is fetched from global memory once per tile and reused 16 times from shared memory instead of being re-fetched for every multiplication. This is often the single biggest optimization available for compute-heavy kernels like this one.

Memory coalescing

When a warp requests data from global memory, the hardware tries to satisfy that request in as few large, contiguous transactions as possible. This works efficiently if thread 0 requests index 0, thread 1 requests index 1, and so on — coalesced access, where one wide transaction can satisfy the entire warp.

// Coalesced: consecutive threads read consecutive memory
c[i] = a[i] + b[i];              // i = threadIdx.x + blockIdx.x * blockDim.x

// Uncoalesced: large stride scatters each thread's access far apart
c[i] = a[i * stride] + b[i * stride];

If the data is scattered like - a linked list, a fragmented hashmap, or simply a large stride, the hardware can't fetch it in one transaction and instead issues a separate transaction per thread (or small group). This uncoalesced access can cut effective memory bandwidth by an order of magnitude, even though the computation itself hasn't changed — only the layout of the data being read.

Takeaway: the layout of your data in memory can matter as much as the algorithm itself.

From CUDA cores to Tensor Cores

Everything above involves a standard CUDA core which is a simple ALU capable of one scalar operation per clock cycle, per thread. That's already effective for general-purpose parallel work. But deep learning workloads are dominated by one specific operation: matrix multiplication. A fully connected layer is a matrix multiply; a convolution can be reformulated as one; the core operation inside transformer attention is, again, matrix multiplication.

To accelerate this specifically, NVIDIA introduced the Tensor Core with the Volta architecture in 2017.

Where a standard CUDA core computes one scalar operation per cycle, a Tensor Core computes a full mixed-precision multiply-accumulate — D = A × B + C — on small matrix tiles (originally 4×4) in a single clock cycle. Later architectures expanded this substantially: larger effective tile sizes, and support for more numeric precisions — FP16, BF16, TF32, INT8, and structured sparsity, which skips known-zero values to save additional cycles. "Mixed precision" here means the multiplication happens in a lower-precision format for speed, while accumulation happens in a higher-precision format to preserve accuracy.

This is the central hardware reason GPUs shifted from being primarily graphics chips that happened to be useful for AI, to being AI accelerators that also happen to still render graphics.

Putting it together

1. GPUs trade single-task speed for massive parallel throughput — well suited to the repetitive, regular math in ML workloads.
2. Before CUDA, that parallelism could only be reached by disguising computation as graphics rendering — data encoded as textures, math run inside fragment shaders, results read back out of a framebuffer.
3. Threads execute in warps of 32, in lockstep, and SMs hide memory latency by switching between resident warps (occupancy) — branching that splits a warp's execution path (divergence) reduces performance.
4. Memory bandwidth, not raw compute, is usually the actual bottleneck — structure data so neighbouring threads access neighbouring memory (coalescing), and reuse shared memory via tiling instead of repeatedly hitting global memory.
5. Tensor Cores exist specifically to accelerate the matrix multiplications neural networks are built from — a core hardware reason large-scale AI training is feasible at all.

Most CUDA optimization techniques you'll encounter fall into one of these five ideas: tiling, shared memory caching, avoiding divergent branches, maximizing occupancy, and choosing the right numeric precision. It's a long way from disguising math as pixel colours to letting a single instruction multiply entire matrix tiles. But the underlying goal never changed: keep thousands of simple cores fed, and keep them all doing the same thing at once.