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.