}
def NVVM_MBarrierArriveOp : NVVM_Op<"mbarrier.arrive">,
+ Results<(outs LLVM_Type:$res)>,
Arguments<(ins LLVM_i64ptr_any:$addr)> {
string llvmBuilder = [{
- createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_arrive, {$addr});
+ $res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_arrive, {$addr});
}];
- let assemblyFormat = "$addr attr-dict `:` type(operands)";
+ let assemblyFormat = "$addr attr-dict `:` type($addr) `->` type($res)";
}
def NVVM_MBarrierArriveSharedOp : NVVM_Op<"mbarrier.arrive.shared">,
+ Results<(outs LLVM_Type:$res)>,
Arguments<(ins LLVM_i64ptr_shared:$addr)> {
string llvmBuilder = [{
- createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_arrive_shared, {$addr});
+ $res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_arrive_shared, {$addr});
}];
- let assemblyFormat = "$addr attr-dict `:` type(operands)";
+ let assemblyFormat = "$addr attr-dict `:` type($addr) `->` type($res)";
}
def NVVM_MBarrierArriveNocompleteOp : NVVM_Op<"mbarrier.arrive.nocomplete">,
+ Results<(outs LLVM_Type:$res)>,
Arguments<(ins LLVM_i64ptr_any:$addr, I32:$count)> {
string llvmBuilder = [{
- createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_arrive_noComplete, {$addr, $count});
+ $res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_arrive_noComplete, {$addr, $count});
}];
- let assemblyFormat = "$addr `,` $count attr-dict `:` type(operands)";
+ let assemblyFormat = "$addr `,` $count attr-dict `:` type(operands) `->` type($res)";
}
def NVVM_MBarrierArriveNocompleteSharedOp : NVVM_Op<"mbarrier.arrive.nocomplete.shared">,
+ Results<(outs LLVM_Type:$res)>,
Arguments<(ins LLVM_i64ptr_shared:$addr, I32:$count)> {
string llvmBuilder = [{
- createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_arrive_noComplete_shared, {$addr, $count});
+ $res = createIntrinsicCall(builder, llvm::Intrinsic::nvvm_mbarrier_arrive_noComplete_shared, {$addr, $count});
}];
- let assemblyFormat = "$addr `,` $count attr-dict `:` type(operands)";
+ let assemblyFormat = "$addr `,` $count attr-dict `:` type(operands) `->` type($res)";
}
//===----------------------------------------------------------------------===//
llvm.func private @mbarrier_arrive(%barrier: !llvm.ptr) {
// CHECK: nvvm.mbarrier.arrive %{{.*}} : !llvm.ptr
- nvvm.mbarrier.arrive %barrier : !llvm.ptr
+ %0 = nvvm.mbarrier.arrive %barrier : !llvm.ptr -> i64
llvm.return
}
llvm.func private @mbarrier_arrive_shared(%barrier: !llvm.ptr<3>) {
// CHECK: nvvm.mbarrier.arrive.shared %{{.*}} : !llvm.ptr<3>
- nvvm.mbarrier.arrive.shared %barrier : !llvm.ptr<3>
+ %0 = nvvm.mbarrier.arrive.shared %barrier : !llvm.ptr<3> -> i64
llvm.return
}
llvm.func private @mbarrier_arrive_nocomplete(%barrier: !llvm.ptr) {
%count = nvvm.read.ptx.sreg.ntid.x : i32
// CHECK: nvvm.mbarrier.arrive.nocomplete %{{.*}} : !llvm.ptr
- nvvm.mbarrier.arrive.nocomplete %barrier, %count : !llvm.ptr, i32
+ %0 = nvvm.mbarrier.arrive.nocomplete %barrier, %count : !llvm.ptr, i32 -> i64
llvm.return
}
llvm.func private @mbarrier_arrive_nocomplete_shared(%barrier: !llvm.ptr<3>) {
%count = nvvm.read.ptx.sreg.ntid.x : i32
// CHECK: nvvm.mbarrier.arrive.nocomplete.shared %{{.*}} : !llvm.ptr<3>
- nvvm.mbarrier.arrive.nocomplete.shared %barrier, %count : !llvm.ptr<3>, i32
+ %0 = nvvm.mbarrier.arrive.nocomplete.shared %barrier, %count : !llvm.ptr<3>, i32 -> i64
llvm.return
}