[Mlir-commits] [mlir] [MLIR][NVVM] Enable import of nvvm.barrier0 (PR #119965)

[Ivan R. Ivanov]

https://github.com/ivanradanov updated https://github.com/llvm/llvm-project/pull/119965

>From ff70aa0100f263fb3330d87a1e7e908d67d74677 Mon Sep 17 00:00:00 2001
From: Ivan Radanov Ivanov <ivanov.i.aa at m.titech.ac.jp>
Date: Sat, 14 Dec 2024 21:19:01 +0900
Subject: [PATCH 1/3] [MLIR][NVVM] Enable import of nvvm.barrier0

---
 mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td |  7 +++----
 mlir/test/Target/LLVMIR/Import/nvvmir.ll    | 11 ++++++-----
 2 files changed, 9 insertions(+), 9 deletions(-)

diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 296a3c305e5bf4..dc8b61630df1cc 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -117,14 +117,13 @@ class NVVM_Attr<string attrName, string attrMnemonic, list<Trait> traits = []>
 // NVVM intrinsic operations
 //===----------------------------------------------------------------------===//
 
-class NVVM_IntrOp<string mnem, list<Trait> traits,
-                  int numResults>
+class NVVM_IntrOp<string mnem, list<Trait> traits = [],
+                  int numResults = 0>
   : LLVM_IntrOpBase<NVVM_Dialect, mnem, "nvvm_" # !subst(".", "_", mnem),
                     /*list<int> overloadedResults=*/[],
                     /*list<int> overloadedOperands=*/[],
                     traits, numResults>;
 
-
 //===----------------------------------------------------------------------===//
 // NVVM special register op definitions
 //===----------------------------------------------------------------------===//
@@ -431,7 +430,7 @@ def NVVM_MBarrierTestWaitSharedOp : NVVM_Op<"mbarrier.test.wait.shared">,
 // NVVM synchronization op definitions
 //===----------------------------------------------------------------------===//
 
-def NVVM_Barrier0Op : NVVM_Op<"barrier0"> {
+def NVVM_Barrier0Op : NVVM_IntrOp<"barrier0"> {
   string llvmBuilder = [{
       createIntrinsicCall(builder, llvm::Intrinsic::nvvm_barrier0);
   }];
diff --git a/mlir/test/Target/LLVMIR/Import/nvvmir.ll b/mlir/test/Target/LLVMIR/Import/nvvmir.ll
index 131e9065b2d883..f169b1cf79332b 100644
--- a/mlir/test/Target/LLVMIR/Import/nvvmir.ll
+++ b/mlir/test/Target/LLVMIR/Import/nvvmir.ll
@@ -71,12 +71,13 @@ define float @nvvm_rcp(float %0) {
   ret float %2
 }
 
-; TODO: Support the intrinsics below once they derive from NVVM_IntrOp rather than from NVVM_Op.
+define void @llvm_nvvm_barrier0() {
+  call void @llvm.nvvm.barrier0()
+  ret void
+}
 
-; define void @llvm_nvvm_barrier0() {
-;   call void @llvm.nvvm.barrier0()
-;   ret void
-; }
+
+; TODO: Support the intrinsics below once they derive from NVVM_IntrOp rather than from NVVM_Op.
 ;
 ; define i32 @nvvm_shfl(i32 %0, i32 %1, i32 %2, i32 %3, float %4) {
 ;   %6 = call i32 @llvm.nvvm.shfl.sync.bfly.i32(i32 %0, i32 %3, i32 %1, i32 %2)

>From 389db834d5b8552d91536b6d580a64b157cae5e5 Mon Sep 17 00:00:00 2001
From: "Ivan R. Ivanov" <ivanov.i.aa at m.titech.ac.jp>
Date: Tue, 17 Dec 2024 11:58:56 +0900
Subject: [PATCH 2/3] Update mlir/test/Target/LLVMIR/Import/nvvmir.ll

Co-authored-by: Tobias Gysi <tobias.gysi at nextsilicon.com>
---
 mlir/test/Target/LLVMIR/Import/nvvmir.ll | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/mlir/test/Target/LLVMIR/Import/nvvmir.ll b/mlir/test/Target/LLVMIR/Import/nvvmir.ll
index f169b1cf79332b..c8b7b82f47fd93 100644
--- a/mlir/test/Target/LLVMIR/Import/nvvmir.ll
+++ b/mlir/test/Target/LLVMIR/Import/nvvmir.ll
@@ -71,7 +71,9 @@ define float @nvvm_rcp(float %0) {
   ret float %2
 }
 
+; CHECK-LABEL: @llvm_nvvm_barrier0()
 define void @llvm_nvvm_barrier0() {
+  ; CHECK: nvvm.barrier0
   call void @llvm.nvvm.barrier0()
   ret void
 }

>From b4e65f54bca4c58a6dae9f123f437d0515a4fbfb Mon Sep 17 00:00:00 2001
From: Ivan Radanov Ivanov <ivanov.i.aa at m.titech.ac.jp>
Date: Tue, 17 Dec 2024 12:16:47 +0900
Subject: [PATCH 3/3] Remove llvm constructor

---
 mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td | 3 ---
 1 file changed, 3 deletions(-)

diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index dc8b61630df1cc..2dca576c344050 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -431,9 +431,6 @@ def NVVM_MBarrierTestWaitSharedOp : NVVM_Op<"mbarrier.test.wait.shared">,
 //===----------------------------------------------------------------------===//
 
 def NVVM_Barrier0Op : NVVM_IntrOp<"barrier0"> {
-  string llvmBuilder = [{
-      createIntrinsicCall(builder, llvm::Intrinsic::nvvm_barrier0);
-  }];
   let assemblyFormat = "attr-dict";
 }