llvm/test/CodeGen/NVPTX/mbarrier_wait_sm90_ptx80.ll

; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80| FileCheck --check-prefixes=CHECK-PTX64 %s
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80 --nvptx-short-ptr| FileCheck --check-prefixes=CHECK-PTX-SHARED32 %s
; RUN: %if ptxas-sm_90 && ptxas-isa-8.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80| %ptxas-verify -arch=sm_90 %}
; 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 %}

; with sm-90 and ptx-80, we have support for cluster-scope

; --- test.wait ---
declare i1 @llvm.nvvm.mbarrier.test.wait.parity.scope.cluster.space.cta(ptr addrspace(3), i32)
declare i1 @llvm.nvvm.mbarrier.test.wait.scope.cluster.space.cta(ptr addrspace(3), i64)

; --- try.wait without timelimit ---
declare i1 @llvm.nvvm.mbarrier.try.wait.scope.cluster.space.cta(ptr addrspace(3), i64)
declare i1 @llvm.nvvm.mbarrier.try.wait.parity.scope.cluster.space.cta(ptr addrspace(3), i32)

; --- try.wait with timelimit ---
declare i1 @llvm.nvvm.mbarrier.try.wait.tl.scope.cluster.space.cta(ptr addrspace(3), i64, i32)
declare i1 @llvm.nvvm.mbarrier.try.wait.parity.tl.scope.cluster.space.cta(ptr addrspace(3), i32, i32)

define void @mbar_test_wait(ptr addrspace(3) %mbar, i64 %state, i32 %parity) {
; CHECK-PTX64-LABEL: mbar_test_wait(
; CHECK-PTX64:       {
; CHECK-PTX64-NEXT:    .reg .pred %p<3>;
; CHECK-PTX64-NEXT:    .reg .b32 %r<2>;
; CHECK-PTX64-NEXT:    .reg .b64 %rd<3>;
; CHECK-PTX64-EMPTY:
; CHECK-PTX64-NEXT:  // %bb.0:
; CHECK-PTX64-NEXT:    ld.param.b64 %rd1, [mbar_test_wait_param_0];
; CHECK-PTX64-NEXT:    ld.param.b64 %rd2, [mbar_test_wait_param_1];
; CHECK-PTX64-NEXT:    mbarrier.test_wait.acquire.cluster.shared.b64 %p1, [%rd1], %rd2;
; CHECK-PTX64-NEXT:    ld.param.b32 %r1, [mbar_test_wait_param_2];
; CHECK-PTX64-NEXT:    mbarrier.test_wait.parity.acquire.cluster.shared.b64 %p2, [%rd1], %r1;
; CHECK-PTX64-NEXT:    ret;
;
; CHECK-PTX-SHARED32-LABEL: mbar_test_wait(
; CHECK-PTX-SHARED32:       {
; CHECK-PTX-SHARED32-NEXT:    .reg .pred %p<3>;
; CHECK-PTX-SHARED32-NEXT:    .reg .b32 %r<3>;
; CHECK-PTX-SHARED32-NEXT:    .reg .b64 %rd<2>;
; CHECK-PTX-SHARED32-EMPTY:
; CHECK-PTX-SHARED32-NEXT:  // %bb.0:
; CHECK-PTX-SHARED32-NEXT:    ld.param.b32 %r1, [mbar_test_wait_param_0];
; CHECK-PTX-SHARED32-NEXT:    ld.param.b64 %rd1, [mbar_test_wait_param_1];
; CHECK-PTX-SHARED32-NEXT:    mbarrier.test_wait.acquire.cluster.shared.b64 %p1, [%r1], %rd1;
; CHECK-PTX-SHARED32-NEXT:    ld.param.b32 %r2, [mbar_test_wait_param_2];
; CHECK-PTX-SHARED32-NEXT:    mbarrier.test_wait.parity.acquire.cluster.shared.b64 %p2, [%r1], %r2;
; CHECK-PTX-SHARED32-NEXT:    ret;
  %ret0 = call i1 @llvm.nvvm.mbarrier.test.wait.scope.cluster.space.cta(ptr addrspace(3) %mbar, i64 %state)
  %ret1 = call i1 @llvm.nvvm.mbarrier.test.wait.parity.scope.cluster.space.cta(ptr addrspace(3) %mbar, i32 %parity)

  ret void
}

define void @mbar_try_wait(ptr addrspace(3) %mbar, i64 %state, i32 %parity) {
; CHECK-PTX64-LABEL: mbar_try_wait(
; CHECK-PTX64:       {
; CHECK-PTX64-NEXT:    .reg .pred %p<3>;
; CHECK-PTX64-NEXT:    .reg .b32 %r<2>;
; CHECK-PTX64-NEXT:    .reg .b64 %rd<3>;
; CHECK-PTX64-EMPTY:
; CHECK-PTX64-NEXT:  // %bb.0:
; CHECK-PTX64-NEXT:    ld.param.b64 %rd1, [mbar_try_wait_param_0];
; CHECK-PTX64-NEXT:    ld.param.b64 %rd2, [mbar_try_wait_param_1];
; CHECK-PTX64-NEXT:    mbarrier.try_wait.acquire.cluster.shared.b64 %p1, [%rd1], %rd2;
; CHECK-PTX64-NEXT:    ld.param.b32 %r1, [mbar_try_wait_param_2];
; CHECK-PTX64-NEXT:    mbarrier.try_wait.parity.acquire.cluster.shared.b64 %p2, [%rd1], %r1;
; CHECK-PTX64-NEXT:    ret;
;
; CHECK-PTX-SHARED32-LABEL: mbar_try_wait(
; CHECK-PTX-SHARED32:       {
; CHECK-PTX-SHARED32-NEXT:    .reg .pred %p<3>;
; CHECK-PTX-SHARED32-NEXT:    .reg .b32 %r<3>;
; CHECK-PTX-SHARED32-NEXT:    .reg .b64 %rd<2>;
; CHECK-PTX-SHARED32-EMPTY:
; CHECK-PTX-SHARED32-NEXT:  // %bb.0:
; CHECK-PTX-SHARED32-NEXT:    ld.param.b32 %r1, [mbar_try_wait_param_0];
; CHECK-PTX-SHARED32-NEXT:    ld.param.b64 %rd1, [mbar_try_wait_param_1];
; CHECK-PTX-SHARED32-NEXT:    mbarrier.try_wait.acquire.cluster.shared.b64 %p1, [%r1], %rd1;
; CHECK-PTX-SHARED32-NEXT:    ld.param.b32 %r2, [mbar_try_wait_param_2];
; CHECK-PTX-SHARED32-NEXT:    mbarrier.try_wait.parity.acquire.cluster.shared.b64 %p2, [%r1], %r2;
; CHECK-PTX-SHARED32-NEXT:    ret;
  %ret0 = call i1 @llvm.nvvm.mbarrier.try.wait.scope.cluster.space.cta(ptr addrspace(3) %mbar, i64 %state)
  %ret1 = call i1 @llvm.nvvm.mbarrier.try.wait.parity.scope.cluster.space.cta(ptr addrspace(3) %mbar, i32 %parity)

  ret void
}

define void @mbar_try_wait_tl(ptr addrspace(3) %mbar, i64 %state, i32 %parity, i32 %tl) {
; CHECK-PTX64-LABEL: mbar_try_wait_tl(
; CHECK-PTX64:       {
; CHECK-PTX64-NEXT:    .reg .pred %p<3>;
; CHECK-PTX64-NEXT:    .reg .b32 %r<3>;
; CHECK-PTX64-NEXT:    .reg .b64 %rd<3>;
; CHECK-PTX64-EMPTY:
; CHECK-PTX64-NEXT:  // %bb.0:
; CHECK-PTX64-NEXT:    ld.param.b64 %rd1, [mbar_try_wait_tl_param_0];
; CHECK-PTX64-NEXT:    ld.param.b64 %rd2, [mbar_try_wait_tl_param_1];
; CHECK-PTX64-NEXT:    ld.param.b32 %r1, [mbar_try_wait_tl_param_3];
; CHECK-PTX64-NEXT:    mbarrier.try_wait.acquire.cluster.shared.b64 %p1, [%rd1], %rd2, %r1;
; CHECK-PTX64-NEXT:    ld.param.b32 %r2, [mbar_try_wait_tl_param_2];
; CHECK-PTX64-NEXT:    mbarrier.try_wait.parity.acquire.cluster.shared.b64 %p2, [%rd1], %r2, %r1;
; CHECK-PTX64-NEXT:    ret;
;
; CHECK-PTX-SHARED32-LABEL: mbar_try_wait_tl(
; CHECK-PTX-SHARED32:       {
; CHECK-PTX-SHARED32-NEXT:    .reg .pred %p<3>;
; CHECK-PTX-SHARED32-NEXT:    .reg .b32 %r<4>;
; CHECK-PTX-SHARED32-NEXT:    .reg .b64 %rd<2>;
; CHECK-PTX-SHARED32-EMPTY:
; CHECK-PTX-SHARED32-NEXT:  // %bb.0:
; CHECK-PTX-SHARED32-NEXT:    ld.param.b32 %r1, [mbar_try_wait_tl_param_0];
; CHECK-PTX-SHARED32-NEXT:    ld.param.b64 %rd1, [mbar_try_wait_tl_param_1];
; CHECK-PTX-SHARED32-NEXT:    ld.param.b32 %r2, [mbar_try_wait_tl_param_3];
; CHECK-PTX-SHARED32-NEXT:    mbarrier.try_wait.acquire.cluster.shared.b64 %p1, [%r1], %rd1, %r2;
; CHECK-PTX-SHARED32-NEXT:    ld.param.b32 %r3, [mbar_try_wait_tl_param_2];
; CHECK-PTX-SHARED32-NEXT:    mbarrier.try_wait.parity.acquire.cluster.shared.b64 %p2, [%r1], %r3, %r2;
; CHECK-PTX-SHARED32-NEXT:    ret;
  %ret0 = call i1 @llvm.nvvm.mbarrier.try.wait.tl.scope.cluster.space.cta(ptr addrspace(3) %mbar, i64 %state, i32 %tl)
  %ret1 = call i1 @llvm.nvvm.mbarrier.try.wait.parity.tl.scope.cluster.space.cta(ptr addrspace(3) %mbar, i32 %parity, i32 %tl)

  ret void
}