5 ms·
Sometimes, the compiler will unroll for you. I have noticed that nvcc will unroll this: #define warpSize 32 float warpMax = partial_max; for (int o
by ryao 2y ago
Sometimes, the compiler will unroll for you. I have noticed that nvcc will unroll this:
#define warpSize 32
float warpMax = partial_max;
for (int offset = warpSize / 2; offset > 0; offset >>= 1) {
float otherValue = __shfl_down_sync(0xFFFFFFFF, warpMax, offset);
warpMax = fmaxf(warpMax, otherValue);
}
Into this:
bar.sync 0;
shl.b32 %r40, %r2, 2;
add.s32 %r8, %r39, %r40;
ld.shared.f32 %f24, [%r8];
mov.u32 %r43, 0;
mov.b32 %r44, %f24;
shfl.sync.down.b32 %r48|%p13, %r44, %r23, %r22, %r24;
mov.b32 %f25, %r48;
max.ftz.f32 %f26, %f24, %f25;
mov.b32 %r49, %f26;
shfl.sync.down.b32 %r51|%p14, %r49, %r27, %r22, %r24;
mov.b32 %f27, %r51;
max.ftz.f32 %f28, %f26, %f27;
mov.b32 %r52, %f28;
shfl.sync.down.b32 %r54|%p15, %r52, %r30, %r22, %r24;
mov.b32 %f29, %r54;
max.ftz.f32 %f30, %f28, %f29;
mov.b32 %r55, %f30;
shfl.sync.down.b32 %r56|%p16, %r55, %r21, %r22, %r24;
mov.b32 %f31, %r56;
max.ftz.f32 %f32, %f30, %f31;
mov.b32 %r57, %f32;
shfl.sync.down.b32 %r59|%p17, %r57, %r35, %r22, %r24;
mov.b32 %f33, %r59;
max.ftz.f32 %f34, %f32, %f33;
mov.b32 %r60, %f34;
This avoids loop branching, which is unnecessary in this case because we know in advance when the loop terminates. The use of ftz instructions is because I compiled using `nvcc -ftz true -ptx -arch=sm_86 rung.cu -o rung.ptx`.
I am not sure if unrolling loops makes sense on a GPU when we don't know when the loop will terminate in advance. If you can suggest any references explaining when that is helpful, I would be happy to read them. :)
- saagarjha 2y agoEven if you don't know the number of iterations it can be helpful to "partially unroll" the loop, for example going by 8 or 16 elements at a time and doing a single check for all elements to make sure you're not doing more work than you were asked to do. Not only does this amortize checks for the loop body it can also enable optimizations like automatic vectorization.
- ryao 2y agoAutomatic vectorization is not relevant to CUDA device code because the threading model is implicitly vector based. Every single instruction is a vector instruction that is executed on all threads in the CUDA block simultaneously (unless it is a no-op due to divergent branching). To be fair, there are actually 14 SIMD instructions intended for use by video operations, but I would be surprised if any compiler implemented optimization passes to use them since most code cannot use them: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#simd-video-instructions https://docs.nvidia.com/cuda/parallel-thread-execution/index... Reducing loop overhead does make sense, although I suspect this would be best informed by profiling the performance. That said, I found the following explanation for unrolling on a GPU helping: https://www.researchgate.net/post/CUDA-Is_it_worth_it_to_unwind_a_static_loop https://www.researchgate.net/post/CUDA-Is_it_worth_it_to_unw... It mentions what had me wondering if loop unrolling made sense on GPUs, which is: > Branch prediction does not even exists on GPUs. The GPU thread scheduler will just switch execution to a different warp until the outcome of the branch has been resolved. It then goes on to mention a few other benefits to loop unrolling on a GPU, including reducing loop overhead.
- saagarjha 2y agoSure it is. Both loads and stores can be vectorized (and should, otherwise you’re leaving memory bandwidth on the table).
- ryao 2y agoThat is on a CPU. A GPU works differently, such that threads on a GPU implicitly vectorize loads and stores as part of their warp/block. My question had concerned GPUs, where you cannot vectorize instructions by loop unrolling since the instructions are already vector instructions.
- saagarjha 2y agoI think you have a mistaken understanding of how GPUs work? There is some "vectorization" across threads in the form of coalescing but what I am talking about is literally a vectorized load/store, the same you would see on a CPU. Like, you can do a ld/ld.64/ld.128 to specify the width of the memory operation. If your loops load individual elements and it is possible to load them together then the compiler can fuse them together.