https://github.com/grypp updated https://github.com/llvm/llvm-project/pull/74058
>From 9f35504e81246f97a9d8c14a06043685660ae15e Mon Sep 17 00:00:00 2001
From: Guray Ozen <guray.o...@gmail.com>
Date: Fri, 1 Dec 2023 11:10:40 +0100
Subject: [PATCH 1/3] [mlir][nvvm] Introduce `fence.mbarrier.init`
This PR introduce `fence.mbarrier.init` OP
---
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 10 ++++++++++
mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir | 8 ++++++++
2 files changed, 18 insertions(+)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index ecad1a16eb6c5..f400c18b5f32c 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -422,6 +422,16 @@ def NVVM_SetMaxRegisterOp :
NVVM_PTXBuilder_Op<"setmaxregister"> {
let hasVerifier = 1;
}
+def NVVM_FenceMbarrierInitOp : NVVM_PTXBuilder_Op<"fence.mbarrier.init"> {
+ let arguments = (ins );
+ 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 5482cc194192d..8366f1d109b1c 100644
--- a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
+++ b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
@@ -629,3 +629,11 @@ 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
+}
>From 34e29b2bef58739dbcc2e34efcec644accd5c089 Mon Sep 17 00:00:00 2001
From: Guray Ozen <guray.o...@gmail.com>
Date: Fri, 1 Dec 2023 16:00:37 +0100
Subject: [PATCH 2/3] add descripton
---
mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 6 ++++++
1 file changed, 6 insertions(+)
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index f400c18b5f32c..adc60e72fdf82 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -424,6 +424,12 @@ def NVVM_SetMaxRegisterOp :
NVVM_PTXBuilder_Op<"setmaxregister"> {
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() {
>From c5d66888946d4397fb29cafa2555f13b9bec8e42 Mon Sep 17 00:00:00 2001
From: Guray Ozen <guray.o...@gmail.com>
Date: Wed, 6 Dec 2023 11:40:56 +0100
Subject: [PATCH 3/3] fix typo in test
---
mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
index ec79ad3e8c187..a4336a30999a1 100644
--- a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
+++ b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
@@ -636,7 +636,8 @@ func.func @cp_bulk_commit() {
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() {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
