37 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_80 -mattr=+ptx71| FileCheck --check-prefixes=CHECK-PTX64 %s3; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx71 --nvptx-short-ptr| FileCheck --check-prefixes=CHECK-PTX-SHARED32 %s4; RUN: %if ptxas-sm_80 && ptxas-isa-7.1 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx71| %ptxas-verify -arch=sm_80 %}5; RUN: %if ptxas-sm_80 && ptxas-isa-7.1 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx71 --nvptx-short-ptr| %ptxas-verify -arch=sm_80 %}6 7; --- test.wait.parity ---8declare i1 @llvm.nvvm.mbarrier.test.wait.parity.scope.cta.space.cta(ptr addrspace(3), i32)9 10define void @mbar_test_wait(ptr addrspace(3) %mbar, i32 %parity) {11; CHECK-PTX64-LABEL: mbar_test_wait(12; CHECK-PTX64: {13; CHECK-PTX64-NEXT: .reg .pred %p<2>;14; CHECK-PTX64-NEXT: .reg .b32 %r<2>;15; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;16; CHECK-PTX64-EMPTY:17; CHECK-PTX64-NEXT: // %bb.0:18; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [mbar_test_wait_param_0];19; CHECK-PTX64-NEXT: ld.param.b32 %r1, [mbar_test_wait_param_1];20; CHECK-PTX64-NEXT: mbarrier.test_wait.parity.shared.b64 %p1, [%rd1], %r1;21; CHECK-PTX64-NEXT: ret;22;23; CHECK-PTX-SHARED32-LABEL: mbar_test_wait(24; CHECK-PTX-SHARED32: {25; CHECK-PTX-SHARED32-NEXT: .reg .pred %p<2>;26; CHECK-PTX-SHARED32-NEXT: .reg .b32 %r<3>;27; CHECK-PTX-SHARED32-EMPTY:28; CHECK-PTX-SHARED32-NEXT: // %bb.0:29; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r1, [mbar_test_wait_param_0];30; CHECK-PTX-SHARED32-NEXT: ld.param.b32 %r2, [mbar_test_wait_param_1];31; CHECK-PTX-SHARED32-NEXT: mbarrier.test_wait.parity.shared.b64 %p1, [%r1], %r2;32; CHECK-PTX-SHARED32-NEXT: ret;33 %ret0 = call i1 @llvm.nvvm.mbarrier.test.wait.parity.scope.cta.space.cta(ptr addrspace(3) %mbar, i32 %parity)34 35 ret void36}37