77 lines · plain
1// RUN: mlir-translate -mlir-to-llvmir %s | FileCheck %s2 3llvm.func @cp_async_mbarrier_arrive(%bar_shared: !llvm.ptr<3>, %bar_gen: !llvm.ptr) {4 // CHECK-LABEL: define void @cp_async_mbarrier_arrive(ptr addrspace(3) %0, ptr %1) {5 // CHECK-NEXT: call void @llvm.nvvm.cp.async.mbarrier.arrive(ptr %1)6 // CHECK-NEXT: call void @llvm.nvvm.cp.async.mbarrier.arrive.noinc(ptr %1)7 // CHECK-NEXT: call void @llvm.nvvm.cp.async.mbarrier.arrive.shared(ptr addrspace(3) %0)8 // CHECK-NEXT: call void @llvm.nvvm.cp.async.mbarrier.arrive.noinc.shared(ptr addrspace(3) %0)9 // CHECK-NEXT: ret void10 // CHECK-NEXT: }11 nvvm.cp.async.mbarrier.arrive %bar_gen : !llvm.ptr12 nvvm.cp.async.mbarrier.arrive %bar_gen {noinc = true} : !llvm.ptr13 nvvm.cp.async.mbarrier.arrive %bar_shared : !llvm.ptr<3>14 nvvm.cp.async.mbarrier.arrive %bar_shared {noinc = true} : !llvm.ptr<3>15 llvm.return16}17 18llvm.func @mbarrier_init_generic(%barrier: !llvm.ptr) {19 // CHECK-LABEL: define void @mbarrier_init_generic(ptr %0) {20 // CHECK-NEXT: %2 = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()21 // CHECK-NEXT: call void @llvm.nvvm.mbarrier.init(ptr %0, i32 %2)22 // CHECK-NEXT: ret void23 // CHECK-NEXT: }24 %count = nvvm.read.ptx.sreg.ntid.x : i3225 nvvm.mbarrier.init %barrier, %count : !llvm.ptr, i3226 llvm.return27}28 29llvm.func @mbarrier_init_shared(%barrier: !llvm.ptr<3>) {30 // CHECK-LABEL: define void @mbarrier_init_shared(ptr addrspace(3) %0) {31 // CHECK-NEXT: %2 = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()32 // CHECK-NEXT: call void @llvm.nvvm.mbarrier.init.shared(ptr addrspace(3) %0, i32 %2)33 // CHECK-NEXT: ret void34 // CHECK-NEXT: }35 %count = nvvm.read.ptx.sreg.ntid.x : i3236 nvvm.mbarrier.init %barrier, %count : !llvm.ptr<3>, i3237 llvm.return38}39 40llvm.func @mbarrier_inval_generic(%barrier: !llvm.ptr) {41 // CHECK-LABEL: define void @mbarrier_inval_generic(ptr %0) {42 // CHECK-NEXT: call void @llvm.nvvm.mbarrier.inval(ptr %0)43 // CHECK-NEXT: ret void44 // CHECK-NEXT: }45 nvvm.mbarrier.inval %barrier : !llvm.ptr46 llvm.return47}48 49llvm.func @mbarrier_inval_shared(%barrier: !llvm.ptr<3>) {50 // CHECK-LABEL: define void @mbarrier_inval_shared(ptr addrspace(3) %0) {51 // CHECK-NEXT: call void @llvm.nvvm.mbarrier.inval.shared(ptr addrspace(3) %0)52 // CHECK-NEXT: ret void53 // CHECK-NEXT: }54 nvvm.mbarrier.inval %barrier : !llvm.ptr<3>55 llvm.return56}57 58llvm.func @mbarrier_test_wait(%barrier: !llvm.ptr, %token : i64) -> i1 {59 // CHECK-LABEL: define i1 @mbarrier_test_wait(ptr %0, i64 %1) {60 // CHECK-NEXT: %3 = call i1 @llvm.nvvm.mbarrier.test.wait(ptr %0, i64 %1)61 // CHECK-NEXT: ret i1 %362 // CHECK-NEXT: }63 %isComplete = nvvm.mbarrier.test.wait %barrier, %token : !llvm.ptr, i64 -> i164 llvm.return %isComplete : i165}66 67llvm.func @mbarrier_test_wait_shared(%barrier: !llvm.ptr<3>, %token : i64) {68 // CHECK-LABEL: define void @mbarrier_test_wait_shared(ptr addrspace(3) %0, i64 %1) {69 // CHECK-NEXT: %3 = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()70 // CHECK-NEXT: %4 = call i1 @llvm.nvvm.mbarrier.test.wait.shared(ptr addrspace(3) %0, i64 %1)71 // CHECK-NEXT: ret void72 // CHECK-NEXT: }73 %count = nvvm.read.ptx.sreg.ntid.x : i3274 %isComplete = nvvm.mbarrier.test.wait %barrier, %token : !llvm.ptr<3>, i64 -> i175 llvm.return76}77