[llvm] [NVPTX] Fold full-warp add reductions to redux.sync.add (PR #226345)

[via llvm-commits]

llvmorg-github-actions[bot] wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-backend-nvptx

Author: frost-miner

<details>
<summary>Changes</summary>

  ## Summary

  Fold full-warp integer butterfly add reductions into `redux.sync.add` in the NVPTX IR peephole pass.

  Compiler-generated full-warp reductions can be represented as a chain of five `shfl.sync.bfly.i32` operations and five integer additions. When the target supports `redux.sync.add`, this sequence can be replaced with a single reduction intrinsic, reducing the amount of generated code.

  ## Implementation

  - Add warp reduction folding to `NVPTXIRPeephole`.
  - Enable the transformation only when the subtarget supports both SM80 and PTX 7.0 or newer.
  - Match five butterfly steps covering each lane-id bit exactly once.
  - Require a full member mask, a clamp of `31`, and matching shuffle source operands.
  - Accept different lane-mask orders and either operand ordering of the add.
  - Replace the matched chain with `llvm.nvvm.redux.sync.add`.

  ## Tests

  Added `nvptx-fold-warp-reduce.ll` to test supported and rejected warp-reduction folding cases on SM80/PTX70 and SM75. Extended `redux-sync.ll` to verify `redux.sync.add.s32` emission.

  Both tests passed.

---
Full diff: https://github.com/llvm/llvm-project/pull/226345.diff


7 Files Affected:

- (modified) llvm/lib/Target/NVPTX/NVPTX.h (+3) 
- (modified) llvm/lib/Target/NVPTX/NVPTXCodeGenPassBuilder.cpp (+1-1) 
- (modified) llvm/lib/Target/NVPTX/NVPTXIRPeephole.cpp (+114-4) 
- (modified) llvm/lib/Target/NVPTX/NVPTXPassRegistry.def (+1-1) 
- (modified) llvm/lib/Target/NVPTX/NVPTXSubtarget.h (+3) 
- (added) llvm/test/CodeGen/NVPTX/nvptx-fold-warp-reduce.ll (+310) 
- (modified) llvm/test/CodeGen/NVPTX/redux-sync.ll (+20) 


``````````diff
The server is unavailable at this time. Please wait a few minutes before you try again.
``````````

</details>


https://github.com/llvm/llvm-project/pull/226345