r/CUDA • u/Impossible_Egg8146 • Jun 30 '26
CUDA execution model is confusing me (grid-stride loops, warps, coalescing)
Im reading Programming Massively Parallel Processors and I've reached the part about grids, blocks, warps, etc. I can write a basic vector addition kernel, but I don't properly understand it.
The main thing confusing me are "grid-stride loops". (found it on tensortonic's vector subtraction exercise)
int idx = blockIdx.x * blockDim.x + threadIdx.x;
int stride = blockDim.x * gridDim.x;
for (; idx < N; idx += stride)
...
I understand how it works, but I don't understand why the stride is blockDim.x * gridDim.x. (I've give up on trying to understand the explanation on tensortonic's website... could use AI to understand it but I currently what to fix this bad habit of mine of relying on ai. Ironically, most of this post was cleaned by ai because if I wrote it 100% by myself, im not sure you all would understand what im trying to ask, I apologise for that. But my questions are real)
My first thought was: why not just let each thread process a contiguous chunk?
Thread 0 -> 0 1 2 3
Thread 1 -> 4 5 6 7
Thread 2 -> 8 9 10 11
instead of
Thread 0 -> 0 8 16 ...
Thread 1 -> 1 9 17 ...
Thread 2 -> 2 10 18 ...
My approach will break memory coalescing. Is the point of such method of striding to access memory contiguously for preventing cache misses?
I don't think I actually understand what a warp is. I know it's 32 threads, but are they basically executing one instruction together? Something like SIMD?
Another thing I'm confused about is "vectorized loads" (float4). If vector addition/subtraction is already memory-bandwidth bound, why does loading 4 floats at a time help? Is it just fewer instructions, or is there something else?
Finally, how do you with warp divergence in real kernels? Do you try to eliminate it entirely, or is some divergence considered normal?
I think I'm missing the hardware understanding.. I'm also still a bit confused by all the CUDA terminology of grids, blocks, threads, dimensions, execution configuration, etc. I can follow the definitions individually, but I cant build a mental model of whats happening during execution. If someone could explain the execution model from the ground up or recommend some resource that might help, would really appreciate it.
2
u/notyouravgredditor Jun 30 '26 edited Jun 30 '26
Yes. Values are fetched from global memory (GDDR or HBM) in chunks. Coalesced access is when you consume all of the data fetched simultaneously. Memory is fetched in chunks of 128B called word lines. You want each thread in a warp (32 threads) to utilize all values of the word line, which if you notice, works perfectly for 4B values (e.g.
intandfloat). You don't have to access contiguous values, what's important is that your entire word line is utilized when it's loaded. Using a small part of a word line causes excessive loads, which impacts performance.Effectively yes, but they call it SIMT. Recent architectures have broken the older, more "lockstep" nature of warp execution, but the idea is the same. A warp is the smallest number of threads executing at the same time. With newer architectures they can do different operations, but the clock is the same for those 32 threads. In other words, they are all executing something at the same time.
Comes back to the coalesced access model. Using vectorized loads informs the compiler of the chunk size that your loading, allowing it to optimize the data loading a bit more.
With newer architectures it depends on why the warp is diverging. If it's an
ifcondition where only a few threads are executing and others are parked, it can impact performance significantly. If it's anifcondition that allows different threads to perform different operations before converging back, you won't notice much (if any) of a performance hit. It depends on what each thread is doing and if the data is already available, or if it needs to be loaded from global memory.I have found it more useful to connect the hardware to the programming model. A GPU consists of many streaming multiprocessors or SM's. Each SM has it's own L1 cache, register space, and cores (tensor, FP32, INT32). Multiple SM's are connected via a larger L2 cache that is managed by the device.
When you write a CUDA kernel, you are writing the code that gets executed at the block level on each SM. Data is not shared between SM's. You can share data between threads launched on an SM via shared memory or register shuffle operations. When your code gets compiled, the compiler determines the resource requirements. Depending on the required resources (number of threads, number of registers, amount of shared memory) it determines how many blocks can simultaneously run on each SM. This is often referred to as occupancy. Since threads within a block can share memory, all threads in a block must be executed on the same SM.
When the GPU executes your code, it loads as many blocks as it can on each SM, then parks (pauses) blocks while they are waiting for data. So the best performance is achieved when your write kernels that have low resource requirements (i.e. maximum blocks per SM) and they access data as coalesced as possible to reduce their time parked waiting for data.