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