[mlir][nvvm] Implement `mbarrier.init`
authorGuray Ozen <guray.ozen@gmail.com>
Fri, 16 Jun 2023 08:03:30 +0000 (10:03 +0200)
committerGuray Ozen <guray.ozen@gmail.com>
Fri, 16 Jun 2023 11:35:14 +0000 (13:35 +0200)
commit58950d4addd6d1dd920801b32cc75ddc8b9f6c3a
treefc320f936c190d643b9d200aafad7f9a0188fad7
parentda7892f7295f31b46486418e2abf15334db96cbb
[mlir][nvvm] Implement `mbarrier.init`

NV GPUs provides split arrive/wait barriers that one can syncronize a subgroup of threads in CTA. It is particularly important for Hopper GPUs and allows tracking engines like TMA. See for more details:
https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-mbarrier

This initial implementation sets the foundation for future enhancements and additions.

Reviewed By: qcolombet

Differential Revision: https://reviews.llvm.org/D151334
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
mlir/test/Dialect/LLVMIR/nvvm.mlir