Most CUDA introductions start with the hardware: streaming multiprocessors, register files, memory hierarchies. That's all true and none of it stuck for me until I translated it into things I already deal with every day as a backend engineer. It turns out most of the rules for fast GPU code have a backend twin. Here's the translation table I wish I'd started with.

It's a thread pool, just absurdly wide

A CUDA kernel is a function that runs once per thread, and you launch thousands of threads at a time. They're grouped into blocks (up to 1024 threads each), and blocks form a grid. Each thread works out which piece of data is its own from where it sits:

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

int threads = 256;
int blocks = (n + threads - 1) / threads;   // one thread per element
add<<<blocks, threads>>>(d_a, d_b, d_out, n);

If you've ever fanned work out to a worker pool by chunking a list, you already understand the launch. The difference is scale: a CPU pool is tens of workers, each one smart. A GPU is tens of thousands, each one simple. You don't optimise one worker. You keep all of them busy.

The warp scheduler is an event loop

Reading from GPU memory takes hundreds of clock cycles. Threads don't sit and wait for it. Threads run in groups of 32 called warps, and when a warp stalls on memory, the scheduler switches to another warp that's ready, at essentially no cost.

That's an event loop. A request that awaits the database gives the loop to someone else, and throughput comes from having enough requests in flight that someone is always ready. Same here: the GPU hides memory latency by having enough warps to swap between. Launch too little work and there's nothing to switch to, so the hardware just waits. A GPU with ten thousand idle cores is a server with one request per minute.

Uncoalesced memory is an N+1 query

When the 32 threads of a warp read memory, the hardware tries to serve them together. If neighbouring threads read neighbouring addresses, the whole warp's request becomes a few wide reads. If they're scattered, it becomes up to 32 separate ones:

float x = data[i];        // coalesced: a warp reads 32 neighbours in one go
float y = data[i * 32];   // strided: every thread lands somewhere different

Both lines are correct. One is dramatically slower, and nothing in the code looks wrong. That's exactly an N+1 query: one row per round trip when a single batched query would do. The fix has the same shape too. Lay the data out the way it will be read, usually structure-of-arrays rather than array-of-structs, so that thread i wants element i.

The PCIe bus is the network

The GPU has its own memory. Data starts on the CPU side and has to be copied across the PCIe bus with cudaMemcpy, which moves data an order of magnitude or more slower than the GPU reads its own memory. A kernel that runs in a millisecond can easily spend longer than that waiting for its inputs to arrive.

Backend engineers know this one in their bones: the slow part is rarely the computation, it's the round trip. So treat device memory like a cache you fought to fill. Copy once, run many kernels on the data while it's there, and copy back only the answer. Shipping data back and forth between every step is the GPU version of calling a microservice inside a loop.

Divergence: why it hates if-statements

The 32 threads of a warp execute together, in step. When they disagree at a branch, the warp can't split up. It runs both sides, one after the other, with half its threads switched off each time:

if (threadIdx.x % 2 == 0) {
    doA();   // half the warp works, half waits
} else {
    doB();   // then they swap
}

To be fair to the title, branching itself is fine. A condition that every thread in a warp answers the same way costs almost nothing. What hurts is a branch that splits a warp down the middle. The backend analogy is a batch job where one slow item holds the whole batch: the batch is only as fast as its slowest path. The fix is the same as it would be on a server: sort or group the work so similar items travel together.

Launch overhead is per-request overhead

Each kernel launch costs a few microseconds before any work happens. For a big kernel that's nothing. For a pipeline of a hundred tiny elementwise steps, it's most of the runtime. That's the case for fusing kernels, doing several operations in one pass while the data is already in registers, and it's why libraries make such a fuss about it. It's request batching: per-call overhead is fixed, so make each call count.

The rule

All of this collapses into one sentence I can actually remember:

Give the GPU one simple thing to do to a lot of data that's already there.

One simple thing, so warps don't diverge. A lot of data, so there are enough warps to hide latency. Already there, so the bus isn't the bottleneck. Lay it out so neighbours read neighbours, and you've covered most of what makes CUDA fast before writing anything clever.

None of these ideas are new. They're the same ones that make a backend fast: keep everyone busy, batch the round trips, don't let one slow path hold the others. The GPU just punishes you for forgetting them a thousand times faster.