; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | FileCheck --check-prefixes=CHECK_PTX64 %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 --nvptx-short-ptr | FileCheck --check-prefixes=CHECK_PTX64_SHARED32 %s
; RUN: llc < %s -march=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | FileCheck --check-prefixes=CHECK_PTX64 %s
; RUN: %if ptxas-sm_100a && ptxas-isa-8.6 %{ llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | %ptxas-verify -arch=sm_100a %}
; RUN: %if ptxas-sm_100a && ptxas-isa-8.6 %{ llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 --nvptx-short-ptr | %ptxas-verify -arch=sm_100a %}
; RUN: %if ptxas-sm_103a && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | %ptxas-verify -arch=sm_103a %}


declare void @llvm.nvvm.tcgen05.alloc.cg1(ptr %addr, i32 %ncols)
declare void @llvm.nvvm.tcgen05.alloc.cg2(ptr %addr, i32 %ncols)
declare void @llvm.nvvm.tcgen05.alloc.shared.cg1(ptr addrspace(3) %addr, i32 %ncols)
declare void @llvm.nvvm.tcgen05.alloc.shared.cg2(ptr addrspace(3) %addr, i32 %ncols)

define void @test_tcgen05_alloc_cg1(ptr %addr, i32 %ncols) {
; CHECK_PTX64-LABEL: test_tcgen05_alloc_cg1(
; CHECK_PTX64:       {
; CHECK_PTX64-NEXT:    .reg .b32 %r<2>;
; CHECK_PTX64-NEXT:    .reg .b64 %rd<2>;
; CHECK_PTX64-EMPTY:
; CHECK_PTX64-NEXT:  // %bb.0:
; CHECK_PTX64-NEXT:    ld.param.b64 %rd1, [test_tcgen05_alloc_cg1_param_0];
; CHECK_PTX64-NEXT:    ld.param.b32 %r1, [test_tcgen05_alloc_cg1_param_1];
; CHECK_PTX64-NEXT:    tcgen05.alloc.cta_group::1.sync.aligned.b32 [%rd1], %r1;
; CHECK_PTX64-NEXT:    ret;
;
; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_alloc_cg1(
; CHECK_PTX64_SHARED32:       {
; CHECK_PTX64_SHARED32-NEXT:    .reg .b32 %r<2>;
; CHECK_PTX64_SHARED32-NEXT:    .reg .b64 %rd<2>;
; CHECK_PTX64_SHARED32-EMPTY:
; CHECK_PTX64_SHARED32-NEXT:  // %bb.0:
; CHECK_PTX64_SHARED32-NEXT:    ld.param.b64 %rd1, [test_tcgen05_alloc_cg1_param_0];
; CHECK_PTX64_SHARED32-NEXT:    ld.param.b32 %r1, [test_tcgen05_alloc_cg1_param_1];
; CHECK_PTX64_SHARED32-NEXT:    tcgen05.alloc.cta_group::1.sync.aligned.b32 [%rd1], %r1;
; CHECK_PTX64_SHARED32-NEXT:    ret;
  call void @llvm.nvvm.tcgen05.alloc.cg1(ptr %addr, i32 %ncols)
  ret void
}

define void @test_tcgen05_alloc_cg2(ptr %addr, i32 %ncols) {
; CHECK_PTX64-LABEL: test_tcgen05_alloc_cg2(
; CHECK_PTX64:       {
; CHECK_PTX64-NEXT:    .reg .b32 %r<2>;
; CHECK_PTX64-NEXT:    .reg .b64 %rd<2>;
; CHECK_PTX64-EMPTY:
; CHECK_PTX64-NEXT:  // %bb.0:
; CHECK_PTX64-NEXT:    ld.param.b64 %rd1, [test_tcgen05_alloc_cg2_param_0];
; CHECK_PTX64-NEXT:    ld.param.b32 %r1, [test_tcgen05_alloc_cg2_param_1];
; CHECK_PTX64-NEXT:    tcgen05.alloc.cta_group::2.sync.aligned.b32 [%rd1], %r1;
; CHECK_PTX64-NEXT:    ret;
;
; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_alloc_cg2(
; CHECK_PTX64_SHARED32:       {
; CHECK_PTX64_SHARED32-NEXT:    .reg .b32 %r<2>;
; CHECK_PTX64_SHARED32-NEXT:    .reg .b64 %rd<2>;
; CHECK_PTX64_SHARED32-EMPTY:
; CHECK_PTX64_SHARED32-NEXT:  // %bb.0:
; CHECK_PTX64_SHARED32-NEXT:    ld.param.b64 %rd1, [test_tcgen05_alloc_cg2_param_0];
; CHECK_PTX64_SHARED32-NEXT:    ld.param.b32 %r1, [test_tcgen05_alloc_cg2_param_1];
; CHECK_PTX64_SHARED32-NEXT:    tcgen05.alloc.cta_group::2.sync.aligned.b32 [%rd1], %r1;
; CHECK_PTX64_SHARED32-NEXT:    ret;
  call void @llvm.nvvm.tcgen05.alloc.cg2(ptr %addr, i32 %ncols)
  ret void
}

define void @test_tcgen05_alloc_shared_cg1(ptr addrspace(3) %addr, i32 %ncols) {
; CHECK_PTX64-LABEL: test_tcgen05_alloc_shared_cg1(
; CHECK_PTX64:       {
; CHECK_PTX64-NEXT:    .reg .b32 %r<2>;
; CHECK_PTX64-NEXT:    .reg .b64 %rd<2>;
; CHECK_PTX64-EMPTY:
; CHECK_PTX64-NEXT:  // %bb.0:
; CHECK_PTX64-NEXT:    ld.param.b64 %rd1, [test_tcgen05_alloc_shared_cg1_param_0];
; CHECK_PTX64-NEXT:    ld.param.b32 %r1, [test_tcgen05_alloc_shared_cg1_param_1];
; CHECK_PTX64-NEXT:    tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%rd1], %r1;
; CHECK_PTX64-NEXT:    ret;
;
; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_alloc_shared_cg1(
; CHECK_PTX64_SHARED32:       {
; CHECK_PTX64_SHARED32-NEXT:    .reg .b32 %r<3>;
; CHECK_PTX64_SHARED32-EMPTY:
; CHECK_PTX64_SHARED32-NEXT:  // %bb.0:
; CHECK_PTX64_SHARED32-NEXT:    ld.param.b32 %r1, [test_tcgen05_alloc_shared_cg1_param_0];
; CHECK_PTX64_SHARED32-NEXT:    ld.param.b32 %r2, [test_tcgen05_alloc_shared_cg1_param_1];
; CHECK_PTX64_SHARED32-NEXT:    tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%r1], %r2;
; CHECK_PTX64_SHARED32-NEXT:    ret;
  call void @llvm.nvvm.tcgen05.alloc.shared.cg1(ptr addrspace(3) %addr, i32 %ncols)
  ret void
}

define void @test_tcgen05_alloc_shared_cg2(ptr addrspace(3) %addr, i32 %ncols) {
; CHECK_PTX64-LABEL: test_tcgen05_alloc_shared_cg2(
; CHECK_PTX64:       {
; CHECK_PTX64-NEXT:    .reg .b32 %r<2>;
; CHECK_PTX64-NEXT:    .reg .b64 %rd<2>;
; CHECK_PTX64-EMPTY:
; CHECK_PTX64-NEXT:  // %bb.0:
; CHECK_PTX64-NEXT:    ld.param.b64 %rd1, [test_tcgen05_alloc_shared_cg2_param_0];
; CHECK_PTX64-NEXT:    ld.param.b32 %r1, [test_tcgen05_alloc_shared_cg2_param_1];
; CHECK_PTX64-NEXT:    tcgen05.alloc.cta_group::2.sync.aligned.shared::cta.b32 [%rd1], %r1;
; CHECK_PTX64-NEXT:    ret;
;
; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_alloc_shared_cg2(
; CHECK_PTX64_SHARED32:       {
; CHECK_PTX64_SHARED32-NEXT:    .reg .b32 %r<3>;
; CHECK_PTX64_SHARED32-EMPTY:
; CHECK_PTX64_SHARED32-NEXT:  // %bb.0:
; CHECK_PTX64_SHARED32-NEXT:    ld.param.b32 %r1, [test_tcgen05_alloc_shared_cg2_param_0];
; CHECK_PTX64_SHARED32-NEXT:    ld.param.b32 %r2, [test_tcgen05_alloc_shared_cg2_param_1];
; CHECK_PTX64_SHARED32-NEXT:    tcgen05.alloc.cta_group::2.sync.aligned.shared::cta.b32 [%r1], %r2;
; CHECK_PTX64_SHARED32-NEXT:    ret;
  call void @llvm.nvvm.tcgen05.alloc.shared.cg2(ptr addrspace(3) %addr, i32 %ncols)
  ret void
}

declare void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols)
declare void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols)

define void @test_tcgen05_dealloc_cg1(ptr addrspace(6) %tmem_addr, i32 %ncols) {
; CHECK_PTX64-LABEL: test_tcgen05_dealloc_cg1(
; CHECK_PTX64:       {
; CHECK_PTX64-NEXT:    .reg .b32 %r<3>;
; CHECK_PTX64-EMPTY:
; CHECK_PTX64-NEXT:  // %bb.0:
; CHECK_PTX64-NEXT:    ld.param.b32 %r1, [test_tcgen05_dealloc_cg1_param_0];
; CHECK_PTX64-NEXT:    ld.param.b32 %r2, [test_tcgen05_dealloc_cg1_param_1];
; CHECK_PTX64-NEXT:    tcgen05.dealloc.cta_group::1.sync.aligned.b32 %r1, %r2;
; CHECK_PTX64-NEXT:    ret;
;
; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_dealloc_cg1(
; CHECK_PTX64_SHARED32:       {
; CHECK_PTX64_SHARED32-NEXT:    .reg .b32 %r<3>;
; CHECK_PTX64_SHARED32-EMPTY:
; CHECK_PTX64_SHARED32-NEXT:  // %bb.0:
; CHECK_PTX64_SHARED32-NEXT:    ld.param.b32 %r1, [test_tcgen05_dealloc_cg1_param_0];
; CHECK_PTX64_SHARED32-NEXT:    ld.param.b32 %r2, [test_tcgen05_dealloc_cg1_param_1];
; CHECK_PTX64_SHARED32-NEXT:    tcgen05.dealloc.cta_group::1.sync.aligned.b32 %r1, %r2;
; CHECK_PTX64_SHARED32-NEXT:    ret;
  call void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols)
  ret void
}

define void @test_tcgen05_dealloc_cg2(ptr addrspace(6) %tmem_addr, i32 %ncols) {
; CHECK_PTX64-LABEL: test_tcgen05_dealloc_cg2(
; CHECK_PTX64:       {
; CHECK_PTX64-NEXT:    .reg .b32 %r<3>;
; CHECK_PTX64-EMPTY:
; CHECK_PTX64-NEXT:  // %bb.0:
; CHECK_PTX64-NEXT:    ld.param.b32 %r1, [test_tcgen05_dealloc_cg2_param_0];
; CHECK_PTX64-NEXT:    ld.param.b32 %r2, [test_tcgen05_dealloc_cg2_param_1];
; CHECK_PTX64-NEXT:    tcgen05.dealloc.cta_group::2.sync.aligned.b32 %r1, %r2;
; CHECK_PTX64-NEXT:    ret;
;
; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_dealloc_cg2(
; CHECK_PTX64_SHARED32:       {
; CHECK_PTX64_SHARED32-NEXT:    .reg .b32 %r<3>;
; CHECK_PTX64_SHARED32-EMPTY:
; CHECK_PTX64_SHARED32-NEXT:  // %bb.0:
; CHECK_PTX64_SHARED32-NEXT:    ld.param.b32 %r1, [test_tcgen05_dealloc_cg2_param_0];
; CHECK_PTX64_SHARED32-NEXT:    ld.param.b32 %r2, [test_tcgen05_dealloc_cg2_param_1];
; CHECK_PTX64_SHARED32-NEXT:    tcgen05.dealloc.cta_group::2.sync.aligned.b32 %r1, %r2;
; CHECK_PTX64_SHARED32-NEXT:    ret;
  call void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols)
  ret void
}

declare void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg1()
declare void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg2()

define void @test_tcgen05_relinquish_alloc_permit_cg1() {
; CHECK_PTX64-LABEL: test_tcgen05_relinquish_alloc_permit_cg1(
; CHECK_PTX64:       {
; CHECK_PTX64-EMPTY:
; CHECK_PTX64-EMPTY:
; CHECK_PTX64-NEXT:  // %bb.0:
; CHECK_PTX64-NEXT:    tcgen05.relinquish_alloc_permit.cta_group::1.sync.aligned;
; CHECK_PTX64-NEXT:    ret;
;
; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_relinquish_alloc_permit_cg1(
; CHECK_PTX64_SHARED32:       {
; CHECK_PTX64_SHARED32-EMPTY:
; CHECK_PTX64_SHARED32-EMPTY:
; CHECK_PTX64_SHARED32-NEXT:  // %bb.0:
; CHECK_PTX64_SHARED32-NEXT:    tcgen05.relinquish_alloc_permit.cta_group::1.sync.aligned;
; CHECK_PTX64_SHARED32-NEXT:    ret;
  call void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg1()
  ret void
}

define void @test_tcgen05_relinquish_alloc_permit_cg2() {
; CHECK_PTX64-LABEL: test_tcgen05_relinquish_alloc_permit_cg2(
; CHECK_PTX64:       {
; CHECK_PTX64-EMPTY:
; CHECK_PTX64-EMPTY:
; CHECK_PTX64-NEXT:  // %bb.0:
; CHECK_PTX64-NEXT:    tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;
; CHECK_PTX64-NEXT:    ret;
;
; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_relinquish_alloc_permit_cg2(
; CHECK_PTX64_SHARED32:       {
; CHECK_PTX64_SHARED32-EMPTY:
; CHECK_PTX64_SHARED32-EMPTY:
; CHECK_PTX64_SHARED32-NEXT:  // %bb.0:
; CHECK_PTX64_SHARED32-NEXT:    tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;
; CHECK_PTX64_SHARED32-NEXT:    ret;
  call void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg2()
  ret void
}