124 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=+ptx80| FileCheck --check-prefixes=CHECK-PTX64 %s3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80 --nvptx-short-ptr| FileCheck --check-prefixes=CHECK-PTX-SHARED32 %s4; RUN: %if ptxas-sm_90 && ptxas-isa-8.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80| %ptxas-verify -arch=sm_90 %}5; RUN: %if ptxas-sm_90 && ptxas-isa-8.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80 --nvptx-short-ptr| %ptxas-verify -arch=sm_90 %}6 7; with sm-90 and ptx-80, we have support for cluster-scope8 9; --- test.wait ---10declare i1 @llvm.nvvm.mbarrier.test.wait.parity.scope.cluster.space.cta(ptr addrspace(3), i32)11declare i1 @llvm.nvvm.mbarrier.test.wait.scope.cluster.space.cta(ptr addrspace(3), i64)12 13; --- try.wait without timelimit ---14declare i1 @llvm.nvvm.mbarrier.try.wait.scope.cluster.space.cta(ptr addrspace(3), i64)15declare i1 @llvm.nvvm.mbarrier.try.wait.parity.scope.cluster.space.cta(ptr addrspace(3), i32)16 17; --- try.wait with timelimit ---18declare i1 @llvm.nvvm.mbarrier.try.wait.tl.scope.cluster.space.cta(ptr addrspace(3), i64, i32)19declare i1 @llvm.nvvm.mbarrier.try.wait.parity.tl.scope.cluster.space.cta(ptr addrspace(3), i32, i32)20 21define void @mbar_test_wait(ptr addrspace(3) %mbar, i64 %state, i32 %parity) {22; CHECK-PTX64-LABEL: mbar_test_wait(23; CHECK-PTX64: {24; CHECK-PTX64-NEXT: .reg .pred %p<3>;25; CHECK-PTX64-NEXT: .reg .b32 %r<2>;26; CHECK-PTX64-NEXT: .reg .b64 %rd<3>;27; CHECK-PTX64-EMPTY:28; CHECK-PTX64-NEXT: // %bb.0:29; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [mbar_test_wait_param_0];30; CHECK-PTX64-NEXT: ld.param.b64 %rd2, [mbar_test_wait_param_1];31; CHECK-PTX64-NEXT: mbarrier.test_wait.acquire.cluster.shared.b64 %p1, [%rd1], %rd2;32; CHECK-PTX64-NEXT: ld.param.b32 %r1, [mbar_test_wait_param_2];33; CHECK-PTX64-NEXT: mbarrier.test_wait.parity.acquire.cluster.shared.b64 %p2, [%rd1], %r1;34; CHECK-PTX64-NEXT: ret;35;36; CHECK-PTX-SHARED32-LABEL: mbar_test_wait(37; CHECK-PTX-SHARED32: {38; CHECK-PTX-SHARED32-NEXT: .reg .pred %p<3>;39; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<3>;40; CHECK-PTX-SHARED32-NEXT: .reg .b64 %rd<2>;41; CHECK-PTX-SHARED32-EMPTY:42; CHECK-PTX-SHARED32-NEXT: // %bb.0:43; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r1, [mbar_test_wait_param_0];44; CHECK-PTX-SHARED32-NEXT: ld.param.b64 %rd1, [mbar_test_wait_param_1];45; CHECK-PTX-SHARED32-NEXT: mbarrier.test_wait.acquire.cluster.shared.b64 %p1, [%r1], %rd1;46; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r2, [mbar_test_wait_param_2];47; CHECK-PTX-SHARED32-NEXT: mbarrier.test_wait.parity.acquire.cluster.shared.b64 %p2, [%r1], %r2;48; CHECK-PTX-SHARED32-NEXT: ret;49 %ret0 = call i1 @llvm.nvvm.mbarrier.test.wait.scope.cluster.space.cta(ptr addrspace(3) %mbar, i64 %state)50 %ret1 = call i1 @llvm.nvvm.mbarrier.test.wait.parity.scope.cluster.space.cta(ptr addrspace(3) %mbar, i32 %parity)51 52 ret void53}54 55define void @mbar_try_wait(ptr addrspace(3) %mbar, i64 %state, i32 %parity) {56; CHECK-PTX64-LABEL: mbar_try_wait(57; CHECK-PTX64: {58; CHECK-PTX64-NEXT: .reg .pred %p<3>;59; CHECK-PTX64-NEXT: .reg .b32 %r<2>;60; CHECK-PTX64-NEXT: .reg .b64 %rd<3>;61; CHECK-PTX64-EMPTY:62; CHECK-PTX64-NEXT: // %bb.0:63; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [mbar_try_wait_param_0];64; CHECK-PTX64-NEXT: ld.param.b64 %rd2, [mbar_try_wait_param_1];65; CHECK-PTX64-NEXT: mbarrier.try_wait.acquire.cluster.shared.b64 %p1, [%rd1], %rd2;66; CHECK-PTX64-NEXT: ld.param.b32 %r1, [mbar_try_wait_param_2];67; CHECK-PTX64-NEXT: mbarrier.try_wait.parity.acquire.cluster.shared.b64 %p2, [%rd1], %r1;68; CHECK-PTX64-NEXT: ret;69;70; CHECK-PTX-SHARED32-LABEL: mbar_try_wait(71; CHECK-PTX-SHARED32: {72; CHECK-PTX-SHARED32-NEXT: .reg .pred %p<3>;73; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<3>;74; CHECK-PTX-SHARED32-NEXT: .reg .b64 %rd<2>;75; CHECK-PTX-SHARED32-EMPTY:76; CHECK-PTX-SHARED32-NEXT: // %bb.0:77; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r1, [mbar_try_wait_param_0];78; CHECK-PTX-SHARED32-NEXT: ld.param.b64 %rd1, [mbar_try_wait_param_1];79; CHECK-PTX-SHARED32-NEXT: mbarrier.try_wait.acquire.cluster.shared.b64 %p1, [%r1], %rd1;80; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r2, [mbar_try_wait_param_2];81; CHECK-PTX-SHARED32-NEXT: mbarrier.try_wait.parity.acquire.cluster.shared.b64 %p2, [%r1], %r2;82; CHECK-PTX-SHARED32-NEXT: ret;83 %ret0 = call i1 @llvm.nvvm.mbarrier.try.wait.scope.cluster.space.cta(ptr addrspace(3) %mbar, i64 %state)84 %ret1 = call i1 @llvm.nvvm.mbarrier.try.wait.parity.scope.cluster.space.cta(ptr addrspace(3) %mbar, i32 %parity)85 86 ret void87}88 89define void @mbar_try_wait_tl(ptr addrspace(3) %mbar, i64 %state, i32 %parity, i32 %tl) {90; CHECK-PTX64-LABEL: mbar_try_wait_tl(91; CHECK-PTX64: {92; CHECK-PTX64-NEXT: .reg .pred %p<3>;93; CHECK-PTX64-NEXT: .reg .b32 %r<3>;94; CHECK-PTX64-NEXT: .reg .b64 %rd<3>;95; CHECK-PTX64-EMPTY:96; CHECK-PTX64-NEXT: // %bb.0:97; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [mbar_try_wait_tl_param_0];98; CHECK-PTX64-NEXT: ld.param.b64 %rd2, [mbar_try_wait_tl_param_1];99; CHECK-PTX64-NEXT: ld.param.b32 %r1, [mbar_try_wait_tl_param_3];100; CHECK-PTX64-NEXT: mbarrier.try_wait.acquire.cluster.shared.b64 %p1, [%rd1], %rd2, %r1;101; CHECK-PTX64-NEXT: ld.param.b32 %r2, [mbar_try_wait_tl_param_2];102; CHECK-PTX64-NEXT: mbarrier.try_wait.parity.acquire.cluster.shared.b64 %p2, [%rd1], %r2, %r1;103; CHECK-PTX64-NEXT: ret;104;105; CHECK-PTX-SHARED32-LABEL: mbar_try_wait_tl(106; CHECK-PTX-SHARED32: {107; CHECK-PTX-SHARED32-NEXT: .reg .pred %p<3>;108; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<4>;109; CHECK-PTX-SHARED32-NEXT: .reg .b64 %rd<2>;110; CHECK-PTX-SHARED32-EMPTY:111; CHECK-PTX-SHARED32-NEXT: // %bb.0:112; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r1, [mbar_try_wait_tl_param_0];113; CHECK-PTX-SHARED32-NEXT: ld.param.b64 %rd1, [mbar_try_wait_tl_param_1];114; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r2, [mbar_try_wait_tl_param_3];115; CHECK-PTX-SHARED32-NEXT: mbarrier.try_wait.acquire.cluster.shared.b64 %p1, [%r1], %rd1, %r2;116; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r3, [mbar_try_wait_tl_param_2];117; CHECK-PTX-SHARED32-NEXT: mbarrier.try_wait.parity.acquire.cluster.shared.b64 %p2, [%r1], %r3, %r2;118; CHECK-PTX-SHARED32-NEXT: ret;119 %ret0 = call i1 @llvm.nvvm.mbarrier.try.wait.tl.scope.cluster.space.cta(ptr addrspace(3) %mbar, i64 %state, i32 %tl)120 %ret1 = call i1 @llvm.nvvm.mbarrier.try.wait.parity.tl.scope.cluster.space.cta(ptr addrspace(3) %mbar, i32 %parity, i32 %tl)121 122 ret void123}124