diff options
author | Guray Ozen <guray.ozen@gmail.com> | 2023-12-06 12:03:20 +0100 |
---|---|---|
committer | GitHub <noreply@github.com> | 2023-12-06 12:03:20 +0100 |
commit | 84e01450a30eabab571103d2f1195565c0ab851a (patch) | |
tree | 26b84c710c2b8af018aaec9b4a110853e72f31dd /mlir | |
parent | 17544fa49df3f890ba22280d10cb28dedd789f9f (diff) | |
download | llvm-84e01450a30eabab571103d2f1195565c0ab851a.zip llvm-84e01450a30eabab571103d2f1195565c0ab851a.tar.gz llvm-84e01450a30eabab571103d2f1195565c0ab851a.tar.bz2 |
[mlir][nvvm] Introduce `fence.mbarrier.init` (#74058)
This PR introduce `fence.mbarrier.init` OP
Diffstat (limited to 'mlir')
-rw-r--r-- | mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 16 | ||||
-rw-r--r-- | mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir | 9 |
2 files changed, 25 insertions, 0 deletions
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td index 6670d94..57986f2 100644 --- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td +++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td @@ -471,6 +471,22 @@ def NVVM_SetMaxRegisterOp : NVVM_PTXBuilder_Op<"setmaxregister"> { let hasVerifier = 1; } +def NVVM_FenceMbarrierInitOp : NVVM_PTXBuilder_Op<"fence.mbarrier.init"> { + let arguments = (ins ); + let description = [{ + Fence operation that applies on the prior nvvm.mbarrier.init + [For more information, see PTX ISA] + (https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#parallel-synchronization-and-communication-instructions-membar) + }]; + + let assemblyFormat = "attr-dict"; + let extraClassDefinition = [{ + std::string $cppClass::getPtx() { + return std::string("fence.mbarrier_init.release.cluster;"); + } + }]; +} + def ShflKindBfly : I32EnumAttrCase<"bfly", 0>; def ShflKindUp : I32EnumAttrCase<"up", 1>; def ShflKindDown : I32EnumAttrCase<"down", 2>; diff --git a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir index 4d2d152..43de50f 100644 --- a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir +++ b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir @@ -642,6 +642,15 @@ func.func @cp_bulk_commit() { nvvm.cp.async.bulk.commit.group func.return } + + +// ----- + +func.func @fence_mbarrier_init() { + //CHECK: llvm.inline_asm has_side_effects asm_dialect = att "fence.mbarrier_init.release.cluster;" + nvvm.fence.mbarrier.init + func.return +} // ----- func.func @fence_proxy() { |