84 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 62; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78| FileCheck --check-prefixes=CHECK-PTX64 %s3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78 --nvptx-short-ptr| FileCheck --check-prefixes=CHECK-PTX-SHARED32 %s4; RUN: %if ptxas-sm_90 && ptxas-isa-7.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78| %ptxas-verify -arch=sm_90 %}5; RUN: %if ptxas-sm_90 && ptxas-isa-7.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx78 --nvptx-short-ptr| %ptxas-verify -arch=sm_90 %}6 7; --- try.wait without timelimit ---8declare i1 @llvm.nvvm.mbarrier.try.wait.scope.cta.space.cta(ptr addrspace(3), i64)9declare i1 @llvm.nvvm.mbarrier.try.wait.parity.scope.cta.space.cta(ptr addrspace(3), i32)10 11; --- try.wait with timelimit ---12declare i1 @llvm.nvvm.mbarrier.try.wait.tl.scope.cta.space.cta(ptr addrspace(3), i64, i32)13declare i1 @llvm.nvvm.mbarrier.try.wait.parity.tl.scope.cta.space.cta(ptr addrspace(3), i32, i32)14 15define void @mbar_try_wait(ptr addrspace(3) %mbar, i64 %state, i32 %parity) {16; CHECK-PTX64-LABEL: mbar_try_wait(17; CHECK-PTX64: {18; CHECK-PTX64-NEXT: .reg .pred %p<3>;19; CHECK-PTX64-NEXT: .reg .b32 %r<2>;20; CHECK-PTX64-NEXT: .reg .b64 %rd<3>;21; CHECK-PTX64-EMPTY:22; CHECK-PTX64-NEXT: // %bb.0:23; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [mbar_try_wait_param_0];24; CHECK-PTX64-NEXT: ld.param.b64 %rd2, [mbar_try_wait_param_1];25; CHECK-PTX64-NEXT: mbarrier.try_wait.shared.b64 %p1, [%rd1], %rd2;26; CHECK-PTX64-NEXT: ld.param.b32 %r1, [mbar_try_wait_param_2];27; CHECK-PTX64-NEXT: mbarrier.try_wait.parity.shared.b64 %p2, [%rd1], %r1;28; CHECK-PTX64-NEXT: ret;29;30; CHECK-PTX-SHARED32-LABEL: mbar_try_wait(31; CHECK-PTX-SHARED32: {32; CHECK-PTX-SHARED32-NEXT: .reg .pred %p<3>;33; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<3>;34; CHECK-PTX-SHARED32-NEXT: .reg .b64 %rd<2>;35; CHECK-PTX-SHARED32-EMPTY:36; CHECK-PTX-SHARED32-NEXT: // %bb.0:37; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r1, [mbar_try_wait_param_0];38; CHECK-PTX-SHARED32-NEXT: ld.param.b64 %rd1, [mbar_try_wait_param_1];39; CHECK-PTX-SHARED32-NEXT: mbarrier.try_wait.shared.b64 %p1, [%r1], %rd1;40; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r2, [mbar_try_wait_param_2];41; CHECK-PTX-SHARED32-NEXT: mbarrier.try_wait.parity.shared.b64 %p2, [%r1], %r2;42; CHECK-PTX-SHARED32-NEXT: ret;43 %ret0 = call i1 @llvm.nvvm.mbarrier.try.wait.scope.cta.space.cta(ptr addrspace(3) %mbar, i64 %state)44 %ret1 = call i1 @llvm.nvvm.mbarrier.try.wait.parity.scope.cta.space.cta(ptr addrspace(3) %mbar, i32 %parity)45 46 ret void47}48 49define void @mbar_try_wait_tl(ptr addrspace(3) %mbar, i64 %state, i32 %parity, i32 %tl) {50; CHECK-PTX64-LABEL: mbar_try_wait_tl(51; CHECK-PTX64: {52; CHECK-PTX64-NEXT: .reg .pred %p<3>;53; CHECK-PTX64-NEXT: .reg .b32 %r<3>;54; CHECK-PTX64-NEXT: .reg .b64 %rd<3>;55; CHECK-PTX64-EMPTY:56; CHECK-PTX64-NEXT: // %bb.0:57; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [mbar_try_wait_tl_param_0];58; CHECK-PTX64-NEXT: ld.param.b64 %rd2, [mbar_try_wait_tl_param_1];59; CHECK-PTX64-NEXT: ld.param.b32 %r1, [mbar_try_wait_tl_param_3];60; CHECK-PTX64-NEXT: mbarrier.try_wait.shared.b64 %p1, [%rd1], %rd2, %r1;61; CHECK-PTX64-NEXT: ld.param.b32 %r2, [mbar_try_wait_tl_param_2];62; CHECK-PTX64-NEXT: mbarrier.try_wait.parity.shared.b64 %p2, [%rd1], %r2, %r1;63; CHECK-PTX64-NEXT: ret;64;65; CHECK-PTX-SHARED32-LABEL: mbar_try_wait_tl(66; CHECK-PTX-SHARED32: {67; CHECK-PTX-SHARED32-NEXT: .reg .pred %p<3>;68; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<4>;69; CHECK-PTX-SHARED32-NEXT: .reg .b64 %rd<2>;70; CHECK-PTX-SHARED32-EMPTY:71; CHECK-PTX-SHARED32-NEXT: // %bb.0:72; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r1, [mbar_try_wait_tl_param_0];73; CHECK-PTX-SHARED32-NEXT: ld.param.b64 %rd1, [mbar_try_wait_tl_param_1];74; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r2, [mbar_try_wait_tl_param_3];75; CHECK-PTX-SHARED32-NEXT: mbarrier.try_wait.shared.b64 %p1, [%r1], %rd1, %r2;76; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r3, [mbar_try_wait_tl_param_2];77; CHECK-PTX-SHARED32-NEXT: mbarrier.try_wait.parity.shared.b64 %p2, [%r1], %r3, %r2;78; CHECK-PTX-SHARED32-NEXT: ret;79 %ret0 = call i1 @llvm.nvvm.mbarrier.try.wait.tl.scope.cta.space.cta(ptr addrspace(3) %mbar, i64 %state, i32 %tl)80 %ret1 = call i1 @llvm.nvvm.mbarrier.try.wait.parity.tl.scope.cta.space.cta(ptr addrspace(3) %mbar, i32 %parity, i32 %tl)81 82 ret void83}84