A practical mental model for CUDA programming

How to map work to threads, reason about memory, and write kernels that make sense

Many CUDA tutorials explain the syntax, but leave out the act of translation.

You learn threadIdx, blockIdx, blockDim, __syncthreads(), shared memory, streams, and a few example kernels. The pieces make sense one by one. But when you sit down to write a CUDA program, the question is still unclear:

What do I actually do first?

CUDA starts to make sense when you stop seeing it as a list of APIs and start seeing it as a way to translate a problem:

problem shape -> units of work -> thread layout -> memory access -> cooperation -> pipeline

This article is about that translation.

why CUDA feels confusing

Most CPU code gives you a loop:

for (int i = 0; i < n; i++) {
    c[i] = a[i] + b[i];
}

The loop gives you i. Each iteration has a clear piece of work.

CUDA is different. You launch many threads, and every thread runs the same kernel code. There is no single loop counter handed to each thread. Each thread must figure out which piece of the total work belongs to it.

That is why many CUDA kernels begin with this line:

int i = blockIdx.x * blockDim.x + threadIdx.x;

This line is not a trick. It answers a practical question:

Among all the GPU threads that were launched, which one am I?

Once a thread knows that, it can decide which data element to process.

CUDA gives coordinates, not work

CUDA does not assign data to threads automatically.

It gives each thread a coordinate. Your program decides what that coordinate means.

In a 1D launch, the useful values are:

blockIdx.x    // which block am I in?
blockDim.x    // how many threads are in each block?
threadIdx.x   // where am I inside this block?

The global index is:

int i = blockIdx.x * blockDim.x + threadIdx.x;

Read it as:

skip all threads in previous blocks,
then add my position inside this block

If each block has 4 threads:

block 0: thread 0 1 2 3 -> index 0 1 2 3
block 1: thread 0 1 2 3 -> index 4 5 6 7
block 2: thread 0 1 2 3 -> index 8 9 10 11

So the formula gives each thread a unique global index.

A simple kernel then looks like this:

__global__ void add(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];
    }
}

The if (i < n) matters because you often launch slightly more threads than data elements. For example, if n = 1000 and your block size is 256, you need 4 blocks, which gives you 1024 threads. The last 24 threads should do nothing.

This is the first checkpoint in CUDA correctness:

Does every thread compute the right data index?
Does every thread stay inside valid bounds?

what .x means

The .x is just the first dimension.

CUDA lets you organize threads in one, two, or three dimensions:

threadIdx.x, threadIdx.y, threadIdx.z
blockIdx.x,  blockIdx.y,  blockIdx.z
blockDim.x,  blockDim.y,  blockDim.z

This exists because many problems already have shapes.

For a 1D array, one dimension is enough:

x

For an image or matrix, two dimensions feel natural:

x = column
y = row

For a 3D volume or simulation grid, three dimensions may fit:

x = column
y = row
z = layer

The important point is that CUDA gives you a thread coordinate system. You choose how that coordinate system maps to your data.

memory is flat, even when data is not

A common source of confusion is this formula:

int i = y * width + x;

or, for 3D data:

int i = z * height * width + y * width + x;

The intuition is simple:

A linear index counts how many elements come before this position.

Memory is stored as a long line:

data[0], data[1], data[2], data[3], ...

But your data may have shape.

For a 2D image:

width = 4

row 0:  0  1  2  3
row 1:  4  5  6  7
row 2:  8  9 10 11

The coordinate (x = 2, y = 1) means:

row 1, column 2

To find its linear index, count what comes before it.

Before row 1, there is one full row:

y * width = 1 * 4 = 4

Inside that row, move 2 positions:

+ x = + 2

So:

i = y * width + x;
// i = 1 * 4 + 2 = 6

The same idea works for 3D data.

A 3D volume is like a stack of 2D layers. Each layer has:

height * width

elements.

So this:

z * height * width

means:

skip all complete layers before layer z

This:

y * width

means:

inside the current layer, skip all complete rows before row y

And this:

+x

means:

move x positions inside the current row

So:

int i = z * height * width + y * width + x;

means:

skipped layers + skipped rows + column offset

This is one of the most useful CUDA ideas:

coordinates describe the shape you think in
linear indices describe where values live in memory

CUDA programming constantly moves between these two views.

start from the unit of work

Before writing a kernel, ask:

What is one unit of work?

For array addition, one unit of work is one array element.

For image processing, one unit of work might be one pixel.

For matrix multiplication, one unit of work might be one output cell, or a tile of output cells.

For a simulation, one unit of work might be one grid point.

This question matters because CUDA code usually begins by deciding what one thread owns.

Simple case:

one thread -> one input item -> one output item

Example:

one thread processes one pixel

More complex case:

one block -> one tile
many threads inside the block cooperate

Example:

one block computes part of a matrix multiplication

If you cannot name the unit of work, the CUDA code will feel arbitrary. Once the unit is clear, the indexing formula has a purpose.

correct code can still be slow

After you can map threads to data, you can write correct CUDA code.

That is only the first checkpoint.

A CUDA kernel can be correct and still be slow.

Correctness asks:

Does each thread access the right data?

Performance asks:

Do groups of threads access memory in a way the GPU can serve efficiently?

These are different questions.

For example, these two patterns may both be correct:

thread 0 reads data[0]
thread 1 reads data[1]
thread 2 reads data[2]
thread 3 reads data[3]

and:

thread 0 reads data[93021]
thread 1 reads data[17]
thread 2 reads data[600004]
thread 3 reads data[88]

The first pattern is usually better because nearby threads read nearby memory.

The GPU does not usually fetch one tiny value for one isolated thread. It moves memory in larger transactions. If neighboring threads ask for neighboring addresses, one memory transaction can serve many threads. If neighboring threads ask for scattered addresses, the GPU may need more memory transactions.

This is the reason behind coalesced memory access.

Coalescing is not a style preference. It reduces memory work.

you write threads, but the GPU runs groups

CUDA lets you write code for one thread. That programming model is useful.

The hardware usually executes threads in groups called warps. A warp is commonly 32 threads.

This has a practical consequence:

nearby threads should usually do similar work
nearby threads should usually read nearby memory

If threads in the same warp read consecutive addresses, the GPU can serve them efficiently.

If they read scattered addresses, memory access becomes more expensive.

Branching has a related issue.

If threads in the same warp take different paths:

if (x > 0) {
    doA();
} else {
    doB();
}

the GPU may need to execute both paths internally and mask off inactive threads.

This does not mean branches are forbidden. Real CUDA code has branches. The point is more practical:

performance depends on group behavior, not only individual thread behavior

A thread looks independent in your code. At runtime, it is often moving with a group.

memory movement explains much of CUDA performance

A beginner model is:

more threads -> more speed

A better model is:

more useful work with efficient memory access -> more speed

GPU memory has levels:

registers       private to one thread, very fast
shared memory   shared by threads in one block, fast
global memory   large, slower
CPU memory      outside the GPU, expensive to transfer

The practical rule is:

arithmetic is often cheaper than moving data

This is why many CUDA optimizations are really memory optimizations.

Ask these questions:

Do nearby threads read nearby memory?
Does the same data get loaded repeatedly?
Can a block reuse data through shared memory?
Is the program copying data between CPU and GPU too often?

Many slow CUDA programs are slow because the GPU is waiting for data, not because it lacks arithmetic power.

shared memory is explicit reuse

GPUs have caches, but shared memory is different.

Cache is automatic. Shared memory is intentional.

With shared memory, the programmer says:

this block of threads will reuse this small piece of data,
so keep it close and let these threads cooperate around it

A typical pattern is:

1. threads in a block load data from global memory into shared memory
2. the block waits until the loading is finished
3. threads reuse the shared data many times
4. the block moves on

The waiting step uses:

__syncthreads();

It means:

all threads in this block must reach this point before any continue

This matters because shared memory is often used for cooperation. If one thread writes shared memory and another thread reads it too early, the result may be wrong.

So shared memory and synchronization usually belong together.

Shared memory is not just “faster memory.” It is a way to make data reuse explicit inside a block.

reduction is the first serious pattern

Array addition is simple:

one input item -> one output item

Thread i owns output c[i].

Many real problems do not have that shape.

Suppose you want to sum an array:

sum = a[0] + a[1] + ... + a[n-1]

Now many inputs produce one result.

This is called reduction.

A reduction usually works in stages:

many threads compute small partial sums
threads inside a block combine those sums
each block produces one partial result
another step combines the block results

Reduction is worth studying because it changes the question.

Simple CUDA asks:

which thread owns this output?

Reduction asks:

how do many threads combine their work?

That is a different skill.

It introduces several ideas at once:

partial results
shared memory
synchronization
tree-like combining
multiple kernel launches or multi-stage computation

Reduction is often the first point where CUDA stops feeling like “parallel loops” and starts feeling like parallel algorithm design.

tiling is controlled data reuse

Matrix multiplication is the classic CUDA example because it makes data reuse obvious.

For each output value in matrix C, you use a row from A and a column from B.

A naive CUDA version may launch many threads, but each thread repeatedly reads from global memory. Many threads end up loading overlapping data again and again.

Tiling changes the pattern.

Instead of having every thread fetch everything directly from global memory, a block cooperatively loads a small tile of A and a small tile of B into shared memory.

Then threads reuse those tiles to compute part of the output.

The pattern is:

load a tile once
synchronize
reuse it many times
move to the next tile

The useful idea is:

pay the global memory cost once, then reuse the data nearby

Tiled matrix multiplication brings many CUDA ideas together:

2D thread indexing
linear memory layout
coalesced global memory access
shared memory
__syncthreads()
data reuse
block-level cooperation

You do not study tiled matrix multiplication only because matrices are important. You study it because it shows the CUDA way of thinking.

The kernel is no longer just assigning work to threads. It is arranging how data moves.

streams matter when the kernel is not the whole program

So far, most of the discussion has been about one kernel.

A real CUDA program has a larger shape:

copy input from CPU to GPU
run kernel
copy output from GPU to CPU

If you process one batch at a time, the program may spend a lot of time waiting:

copy batch 1
compute batch 1
copy result 1

copy batch 2
compute batch 2
copy result 2

Streams and async copy address this outer problem.

They let the program overlap work:

copy batch 2 while computing batch 1
compute batch 2 while copying result 1

This does not change what a kernel computes. It changes when copies and kernels can happen relative to each other.

Kernel optimization asks:

how fast can this computation run once data is already on the GPU?

Stream optimization asks:

can the program avoid waiting while data moves between CPU and GPU?

For this part of CUDA, the important concepts are:

cudaMemcpyAsync
pinned memory
streams
events
double buffering

Streams are especially useful when transfer time and compute time are both significant. They are less useful if your program is dominated by one slow kernel or if the workload is too small.

A good way to think about streams is:

shared memory helps inside a kernel
streams help around kernels

profiling tells you what kind of problem you have

CUDA performance is hard to guess by intuition alone.

A kernel may be slow because:

memory access is scattered
global memory bandwidth is saturated
there is too much branching
occupancy is low
shared memory is used poorly
CPU-GPU transfers dominate runtime
the workload is too small

Before optimizing, try to identify the bottleneck.

A useful first distinction is:

memory-bound
compute-bound
transfer-bound

A memory-bound kernel spends most of its time waiting on memory.

A compute-bound kernel spends most of its time doing arithmetic.

A transfer-bound program spends too much time moving data between CPU and GPU.

These require different fixes.

For a memory-bound kernel, think about:

coalescing
data layout
shared memory
tiling
less redundant loading

For a compute-bound kernel, think about:

less arithmetic
better instruction mix
using specialized hardware if appropriate

For a transfer-bound program, think about:

fewer CPU-GPU copies
larger batches
pinned memory
streams
async copy
keeping data on the GPU longer

Without profiling, CUDA optimization easily becomes guessing.

a practical checklist for writing CUDA

When you face a new CUDA problem, use this checklist.

First, find the unit of work:

What is one independent piece of work?
Is it one array element, one pixel, one output cell, one grid point, or one tile?

Then map work to threads:

What does one thread own?
How do I compute its global index?
Do I need 1D, 2D, or 3D indexing?
Do I need a bounds check?

Then reason about memory:

Do nearby threads access nearby memory?
Is my data layout helping or hurting coalescing?
Is the same data loaded many times?
Would shared memory help reuse it?

Then check cooperation:

Do threads need to combine partial results?
Do they need shared memory?
Where is synchronization required?
Is this a map, reduction, stencil, histogram, or tiled computation?

Then check the whole program:

How much time is spent copying data?
Can data stay on the GPU longer?
Can copy and compute overlap?
Should I use streams?

This checklist is often more useful than memorizing more APIs.

a compact learning path

A reasonable CUDA learning order is:

1. 1D indexing
2. 2D and 3D flattening
3. bounds checks
4. coalesced memory access
5. shared memory
6. __syncthreads()
7. reduction
8. stencil or neighborhood access
9. tiled matrix multiplication
10. profiling memory-bound vs compute-bound kernels
11. pinned memory, streams, and async copy

This order moves through three levels.

First, correctness:

which thread owns which data?

Second, kernel performance:

how do memory access and cooperation affect speed?

Third, program throughput:

how do copy and compute overlap across the whole program?

Later topics like atomics, warp-level primitives, bank conflicts, occupancy, Tensor Cores, CUDA graphs, and multi-GPU programming are important. But they make more sense after the basic model is stable.

the practical model

CUDA programming becomes much clearer when you organize it into three layers.

Layer 1: mapping

thread identity -> data coordinate -> memory index

This answers:

which data does this thread own?

Layer 2: memory and cooperation

coalesced access -> shared memory -> synchronization -> reuse

This answers:

can the GPU feed and coordinate these threads efficiently?

Layer 3: pipeline

CPU-GPU copy -> kernel execution -> GPU-CPU copy
streams -> async copy -> overlap

This answers:

can the whole program avoid unnecessary waiting?

That is the model most beginners are missing.

CUDA is not only about launching many threads. That is the visible part.

The useful way to think is:

choose a unit of work
map it to threads
map coordinates to memory
make memory access regular
reuse data when possible
coordinate threads when needed
overlap copy and compute when the pipeline matters
measure the bottleneck

The quiet surprise is that CUDA programming is often less about writing computation and more about arranging work.

The GPU is powerful when the work is regular, the memory access is predictable, and enough independent work is available to keep the hardware busy.

If you can look at a problem and ask these questions naturally, CUDA stops being a collection of strange keywords. It becomes a way to design the path from data to parallel work to memory movement.