[all-commits] [llvm/llvm-project] 633893: [mlir][nvvm] Support predicates in `BasicPtxBuilde...

[Guray Ozen via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: 63389326f529fd3e3019f8f8afae662e765a3b72
      https://github.com/llvm/llvm-project/commit/63389326f529fd3e3019f8f8afae662e765a3b72
  Author: Guray Ozen <guray.ozen at gmail.com>
  Date:   2023-10-17 (Tue, 17 Oct 2023)

  Changed paths:
    M mlir/include/mlir/Dialect/LLVMIR/BasicPtxBuilderInterface.td
    M mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
    M mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp
    M mlir/lib/Conversion/NVVMToLLVM/NVVMToLLVM.cpp
    M mlir/lib/Dialect/LLVMIR/IR/BasicPtxBuilderInterface.cpp
    M mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir

  Log Message:
  -----------
  [mlir][nvvm] Support predicates in `BasicPtxBuilder` (#67102)

This PR enhances `BasicPtxBuilder` to support predicates in PTX code
generation. The `BasicPtxBuilder` interface was initially introduced for
generating PTX code automatically for Ops that aren't supported by LLVM
core. Predicates, which are typically not supported in LLVM core, are
now supported using the same mechanism.

In PTX programming, instructions can be guarded by predicates as shown
below:. Here `@p` is a predicate register and guard the execution of the
instruction.

```
@p ptx.code op1, op2, op3
```

This PR introduces the `getPredicate` function in the `BasicPtxBuilder`
interface to set an optional predicate. When a predicate is provided,
the instruction is generated with predicate and guarded, otherwise,
predicate is not genearted. Note that the predicate value must always
appear as the last argument on the Op definition.

Additionally, this PR implements predicate usage for the following ops:

- mbarrier.init
- mbarrier.init.shared
- mbarrier.arrive.expect_tx
- mbarrier.arrive.expect_tx.shared
- cp.async.bulk.tensor.shared.cluster.global
- cp.async.bulk.tensor.global.shared.cta

See for more detail in PTX programing model

https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#ptx-instructions