4 ms·

Damn, wow. Good stuff. First of all, I assumed looping within kernels was "always bad" from a performance standpoint (not that I'd done any benchmarks on loopin

Damn, wow. Good stuff. First of all, I assumed looping within kernels was "always bad" from a performance standpoint (not that I'd done any benchmarks on looping vs repeated kernel invocation). I also assumed it was "always bad" to include conditionals. That certainly opens up a lot of possibilities on algorithms. It also explains why people think I'm crazy when I say I don't know how to multiply matrices that are not in power of 2 dimensions haha. Thank you, not sure why I was limiting myself like that. Not too many people to chat to about low level GPGPU stuff!

saagarjha 2y ago
Loops and conditions within kernels are much better than the alternative, since synchronizing with the CPU is quite expensive. If you unroll appropriately and hide latencies they are very effective.
- 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 ago
    Even 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 ago
      Automatic 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 ago
        Sure it is. Both loads and stores can be vectorized (and should, otherwise you’re leaving memory bandwidth on the table).
        ryao 2y ago
        That 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.
ryao 2y ago
It is an area where I am still learning, but as far as I know, the badness of things is relative. CPU-GPU communication is the major thing to avoid. In my CUDA llama3.c fork, making changes to reduce CPU-GPU communication generally increases performance. That means combining kernels and reducing cudaMemcpy and cudaMalloc calls. CUDA graphs likely could give further reductions, but I have yet to run out of ideas for more normal methods for getting them. To give a concrete example of CPU-GPU communication in a tight loop being bad (much like what you were planning to do), when I initially moved softmax to the GPU, I had been invoking the kernel in a host side loop (since attention calls softmax multiple times per pass), and it slowed down inference by a few tokens per second versus the CPU softmax. I eliminated the host side loop by using the grid dimension so that 1 CUDA call did all of the loop iterations implicitly and inference became a few tokens per second faster than the CPU softmax. As for branches, avoiding them on the GPU where possible can be beneficial, as long as you are not doing something more expensive like CPU-GPU communication in their place. For example, in my softmax kernel, rather than do a branch to calculate the maximum value inside my loop, I use fmaxf(). This is better since there is no branch, and fmaxf is incredibly cheap since it is a single PTX instruction (which presumably maps to a single SASS intruction). GPGPU grew out of work for having programmable shaders on GPUs and a big part of that was branching. As long as you are minimizing branch divergence and only using necessary branches, you are fine as far as branching on a GPU are concerned.

saagarjha 2y ago

Loops and conditions within kernels are much better than the alternative, since synchronizing with the CPU is quite expensive. If you unroll appropriately and hide latencies they are very effective.

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 ago
  Even 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 ago
    Automatic 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 ago
      Sure it is. Both loads and stores can be vectorized (and should, otherwise you’re leaving memory bandwidth on the table).
      - ryao 2y ago
        That 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.

ryao 2y ago

It is an area where I am still learning, but as far as I know, the badness of things is relative. CPU-GPU communication is the major thing to avoid. In my CUDA llama3.c fork, making changes to reduce CPU-GPU communication generally increases performance. That means combining kernels and reducing cudaMemcpy and cudaMalloc calls. CUDA graphs likely could give further reductions, but I have yet to run out of ideas for more normal methods for getting them. To give a concrete example of CPU-GPU communication in a tight loop being bad (much like what you were planning to do), when I initially moved softmax to the GPU, I had been invoking the kernel in a host side loop (since attention calls softmax multiple times per pass), and it slowed down inference by a few tokens per second versus the CPU softmax. I eliminated the host side loop by using the grid dimension so that 1 CUDA call did all of the loop iterations implicitly and inference became a few tokens per second faster than the CPU softmax. As for branches, avoiding them on the GPU where possible can be beneficial, as long as you are not doing something more expensive like CPU-GPU communication in their place. For example, in my softmax kernel, rather than do a branch to calculate the maximum value inside my loop, I use fmaxf(). This is better since there is no branch, and fmaxf is incredibly cheap since it is a single PTX instruction (which presumably maps to a single SASS intruction). GPGPU grew out of work for having programmable shaders on GPUs and a big part of that was branching. As long as you are minimizing branch divergence and only using necessary branches, you are fine as far as branching on a GPU are concerned.