[MLIR][NVVM] Enable import of nvvm.barrier0

[ivanradanov]

No description provided.

[llvmbot]

@llvm/pr-subscribers-mlir-llvm

@llvm/pr-subscribers-mlir

Author: Ivan R. Ivanov (ivanradanov)

Changes

---

Full diff: https://github.com/llvm/llvm-project/pull/119965.diff

2 Files Affected:

• (modified) mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td (+1-1)
• (modified) mlir/test/Target/LLVMIR/Import/nvvmir.ll (+6-5)

diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 296a3c305e5bf4..8841897842c3d8 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
@@ -431,7 +431,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)

[schwarzschild-radius]

Given the barrier0 is an NVVM_IntrOp, Can the llvmBuilder code be removed?

[gysit]

Thanks, LGTM modulo comment.

[grypp]

The PR looks good, but we might want to deprecate this operation (nvvm.barrier0) and use nvvm.barrier OP. It is more capable and aligns better with the PTX instruction.

[ivanradanov]

Thank you for the comments.