diff options
| author | Guray Ozen <guray.ozen@gmail.com> | 2023-07-11 17:34:44 +0200 |
|---|---|---|
| committer | Guray Ozen <guray.ozen@gmail.com> | 2023-07-11 17:35:27 +0200 |
| commit | affcfccd3c1c51388c74b1083d0967332be3909d (patch) | |
| tree | c08eed1c5fa14e3b47d10b95b763311f5bfddab9 /llvm/tools/llvm-objdump/llvm-objdump.cpp | |
| parent | be29fe2f987b5bf58d7f6aa77c06e58d9402064a (diff) | |
| download | llvm-affcfccd3c1c51388c74b1083d0967332be3909d.zip llvm-affcfccd3c1c51388c74b1083d0967332be3909d.tar.gz llvm-affcfccd3c1c51388c74b1083d0967332be3909d.tar.bz2 | |
[mlir][nvgpu] Add initial support for `mbarrier`
`mbarrier` is a barrier created in shared memory that supports different flavors of synchronizing threads other than `__syncthreads`, for more information see below.
https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-mbarrier
This work adds initial Ops wrt `mbarrier` to nvgpu dialect.
First, it introduces to two types:
`mbarrier.barrier` that is barrier object in shared memory
`mbarrier.barrier.token` that is token
It introduces following Ops:
`mbarrier.create` creates `mbarrier.barrier`
`mbarrier.init` initializes `mbarrier.barrier`
`mbarrier.arrive` performs arrive-on `mbarrier.barrier` returns `mbarrier.barrier.token`
`mbarrier.arrive.nocomplete` performs arrive-on (non-blocking) `mbarrier.barrier` returns `mbarrier.barrier.token`
`mbarrier.test_wait` waits on `mbarrier.barrier` and `mbarrier.barrier.token`
Reviewed By: nicolasvasilache
Differential Revision: https://reviews.llvm.org/D154090
Diffstat (limited to 'llvm/tools/llvm-objdump/llvm-objdump.cpp')
0 files changed, 0 insertions, 0 deletions