diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td index 296a3c305e5bf..2dca576c34405 100644 --- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td +++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td @@ -117,14 +117,13 @@ class NVVM_Attr traits = []> // NVVM intrinsic operations //===----------------------------------------------------------------------===// -class NVVM_IntrOp traits, - int numResults> +class NVVM_IntrOp traits = [], + int numResults = 0> : LLVM_IntrOpBase overloadedResults=*/[], /*list overloadedOperands=*/[], traits, numResults>; - //===----------------------------------------------------------------------===// // NVVM special register op definitions //===----------------------------------------------------------------------===// @@ -431,10 +430,7 @@ def NVVM_MBarrierTestWaitSharedOp : NVVM_Op<"mbarrier.test.wait.shared">, // NVVM synchronization op definitions //===----------------------------------------------------------------------===// -def NVVM_Barrier0Op : NVVM_Op<"barrier0"> { - string llvmBuilder = [{ - createIntrinsicCall(builder, llvm::Intrinsic::nvvm_barrier0); - }]; +def NVVM_Barrier0Op : NVVM_IntrOp<"barrier0"> { let assemblyFormat = "attr-dict"; } diff --git a/mlir/test/Target/LLVMIR/Import/nvvmir.ll b/mlir/test/Target/LLVMIR/Import/nvvmir.ll index 131e9065b2d88..c8b7b82f47fd9 100644 --- a/mlir/test/Target/LLVMIR/Import/nvvmir.ll +++ b/mlir/test/Target/LLVMIR/Import/nvvmir.ll @@ -71,12 +71,15 @@ 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. +; CHECK-LABEL: @llvm_nvvm_barrier0() +define void @llvm_nvvm_barrier0() { + ; CHECK: 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)