diff options
| author | Guray Ozen <guray.ozen@gmail.com> | 2023-06-16 10:03:30 +0200 |
|---|---|---|
| committer | Guray Ozen <guray.ozen@gmail.com> | 2023-06-16 13:35:14 +0200 |
| commit | 58950d4addd6d1dd920801b32cc75ddc8b9f6c3a (patch) | |
| tree | fc320f936c190d643b9d200aafad7f9a0188fad7 /llvm/lib/Transforms/Utils/LoopVersioning.cpp | |
| parent | da7892f7295f31b46486418e2abf15334db96cbb (diff) | |
| download | llvm-58950d4addd6d1dd920801b32cc75ddc8b9f6c3a.zip llvm-58950d4addd6d1dd920801b32cc75ddc8b9f6c3a.tar.gz llvm-58950d4addd6d1dd920801b32cc75ddc8b9f6c3a.tar.bz2 | |
[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
Diffstat (limited to 'llvm/lib/Transforms/Utils/LoopVersioning.cpp')
0 files changed, 0 insertions, 0 deletions