Skip to content

Commit 384fb78

Browse files
durga4githubdurga4github
committed
[NVPTX] Add tcgen05 alloc/dealloc intrinsics
This patch adds intrinsics for the tcgen05 alloc/dealloc family of PTX instructions. This patch also adds addrspace 6 for tensor memory which is used by these intrinsics. lit tests are added and verified with a ptxas-12.8 executable. Documentation for these additions is also added in NVPTXUsage.rst. Signed-off-by: Durgadoss R <[email protected]>
1 parent 9534d27 commit 384fb78

File tree

8 files changed

+338
-0
lines changed

8 files changed

+338
-0
lines changed

‎llvm/docs/NVPTXUsage.rst

Lines changed: 103 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -962,6 +962,109 @@ The ``griddepcontrol`` intrinsics allows the dependent grids and prerequisite gr
962962
For more information, refer
963963
`PTX ISA <https://docs.nvidia.com/cuda/parallel-thread-execution/#parallel-synchronization-and-communication-instructions-griddepcontrol>`__.
964964

965+
TCGEN05 family of Intrinsics
966+
----------------------------
967+
968+
The llvm.nvvm.tcgen05.* intrinsics model the TCGEN05 family of instructions
969+
exposed by PTX. These intrinsics use 'Tensor Memory' (henceforth ``tmem``).
970+
NVPTX represents this memory using ``addrspace(6)`` and is always 32-bits.
971+
972+
For more information, refer to the PTX ISA
973+
`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory>`_.
974+
975+
The tensor-memory pointers may only be used with the tcgen05 intrinsics.
976+
There are specialized load/store instructions provided (tcgen05.ld/st) to
977+
work with tensor-memory.
978+
979+
See the PTX ISA for more information on tensor-memory load/store instructions
980+
`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-and-register-load-store-instructions>`_.
981+
982+
All tcgen05 intrinsics use a ``null`` pointer in tmem address
983+
space as their last operand. This helps to preserve ordering among the tcgen05
984+
operations especially when the intrinsic lacks any tmem operands. This
985+
last operand is dropped during Codegen.
986+
987+
'``llvm.nvvm.tcgen05.alloc``'
988+
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
989+
990+
Syntax:
991+
"""""""
992+
993+
.. code-block:: llvm
994+
995+
declare void @llvm.nvvm.tcgen05.alloc.cg1(ptr %dst, i32 %ncols, ptr addrspace(6) null)
996+
declare void @llvm.nvvm.tcgen05.alloc.cg2(ptr %dst, i32 %ncols, ptr addrspace(6) null)
997+
declare void @llvm.nvvm.tcgen05.alloc.shared.cg1(ptr addrspace(3) %dst, i32 %ncols, ptr addrspace(6) null)
998+
declare void @llvm.nvvm.tcgen05.alloc.shared.cg2(ptr addrspace(3) %dst, i32 %ncols, ptr addrspace(6) null)
999+
1000+
Overview:
1001+
"""""""""
1002+
1003+
The '``@llvm.nvvm.tcgen05.alloc.*``' intrinsics correspond to the
1004+
``tcgen05.alloc.cta_group*.sync.aligned.b32`` family of PTX instructions.
1005+
The ``tcgen05.alloc`` is a potentially blocking instruction which dynamically
1006+
allocates the specified number of columns in the Tensor Memory and writes
1007+
the address of the allocated Tensor Memory into shared memory at the
1008+
location specified by ``%dst``. The 32-bit operand ``%ncols`` specifies
1009+
the number of columns to be allocated and it must be a power-of-two.
1010+
The ``.shared`` variant explicitly uses shared memory address space for
1011+
the ``%dst`` operand. The ``.cg1`` and ``.cg2`` variants generate
1012+
``cta_group::1`` and ``cta_group::2`` variants of the instruction respectively.
1013+
1014+
For more information, refer to the PTX ISA
1015+
`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions>`_.
1016+
1017+
'``llvm.nvvm.tcgen05.dealloc``'
1018+
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
1019+
1020+
Syntax:
1021+
"""""""
1022+
1023+
.. code-block:: llvm
1024+
1025+
declare void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols, ptr addrspace(6) null)
1026+
declare void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols, ptr addrspace(6) null)
1027+
1028+
Overview:
1029+
"""""""""
1030+
1031+
The '``@llvm.nvvm.tcgen05.dealloc.*``' intrinsics correspond to the
1032+
``tcgen05.dealloc.*`` set of PTX instructions. The ``tcgen05.dealloc``
1033+
instructions deallocates the Tensor Memory specified by the Tensor Memory
1034+
address ``%tmem_addr``. The operand ``%tmem_addr`` must point to a previous
1035+
Tensor Memory allocation. The 32-bit operand ``%ncols`` specifies the number
1036+
of columns to be de-allocated. The ``.cg1`` and ``.cg2`` variants generate
1037+
``cta_group::1`` and ``cta_group::2`` variants of the instruction respectively.
1038+
1039+
For more information, refer to the PTX ISA
1040+
`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions>`_.
1041+
1042+
'``llvm.nvvm.tcgen05.relinq.alloc.permit``'
1043+
^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
1044+
1045+
Syntax:
1046+
"""""""
1047+
1048+
.. code-block:: llvm
1049+
1050+
declare void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg1(ptr addrspace(6) null)
1051+
declare void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg2(ptr addrspace(6) null)
1052+
1053+
Overview:
1054+
"""""""""
1055+
1056+
The '``@llvm.nvvm.tcgen05.relinq.alloc.permit.*``' intrinsics correspond
1057+
to the ``tcgen05.relinquish_alloc_permit.*`` set of PTX instructions.
1058+
This instruction specifies that the CTA of the executing thread is
1059+
relinquishing the right to allocate Tensor Memory. So, it is illegal
1060+
for a CTA to perform ``tcgen05.alloc`` after any of its constituent
1061+
threads execute ``tcgen05.relinquish_alloc_permit``. The ``.cg1``
1062+
and ``.cg2`` variants generate ``cta_group::1`` and ``cta_group::2``
1063+
flavors of the instruction respectively.
1064+
1065+
For more information, refer to the PTX ISA
1066+
`<https://docs.nvidia.com/cuda/parallel-thread-execution/#tensor-memory-allocation-and-management-instructions>`_.
1067+
9651068
Other Intrinsics
9661069
----------------
9671070

‎llvm/include/llvm/IR/IntrinsicsNVVM.td

Lines changed: 34 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -48,6 +48,7 @@
4848

4949
def llvm_global_ptr_ty : LLVMQualPointerType<1>; // (global)ptr
5050
def llvm_shared_ptr_ty : LLVMQualPointerType<3>; // (shared)ptr
51+
def llvm_tmem_ptr_ty : LLVMQualPointerType<6>; // (tensor memory)ptr
5152

5253
//
5354
// MISC
@@ -5055,4 +5056,37 @@ def int_nvvm_cp_async_bulk_prefetch_L2
50555056
def int_nvvm_griddepcontrol_launch_dependents: Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects]>;
50565057
def int_nvvm_griddepcontrol_wait: Intrinsic<[], [], [IntrNoMem, IntrHasSideEffects]>;
50575058

5059+
//
5060+
// Tcgen05 family of Intrinsics
5061+
//
5062+
5063+
// Tcgen05 alloc/dealloc related intrinsics
5064+
5065+
foreach cta_group = ["cg1", "cg2"] in {
5066+
def int_nvvm_tcgen05_alloc_ # cta_group : Intrinsic<[],
5067+
[llvm_ptr_ty, // dst_ptr
5068+
llvm_i32_ty, // num_columns
5069+
llvm_tmem_ptr_ty], // tmem_token
5070+
[IntrConvergent, IntrArgMemOnly, WriteOnly<ArgIndex<0>>,
5071+
NoCapture<ArgIndex<0>>, NoCapture<ArgIndex<2>>]>;
5072+
5073+
def int_nvvm_tcgen05_alloc_shared_ # cta_group : Intrinsic<[],
5074+
[llvm_shared_ptr_ty, // dst_ptr
5075+
llvm_i32_ty, // num_columns
5076+
llvm_tmem_ptr_ty], // tmem_token
5077+
[IntrConvergent, IntrArgMemOnly, WriteOnly<ArgIndex<0>>,
5078+
NoCapture<ArgIndex<0>>, NoCapture<ArgIndex<2>>]>;
5079+
5080+
def int_nvvm_tcgen05_dealloc_ # cta_group : Intrinsic<[],
5081+
[llvm_tmem_ptr_ty, // tmem_addr
5082+
llvm_i32_ty, // num_columns
5083+
llvm_tmem_ptr_ty], // tmem_token
5084+
[IntrConvergent, IntrArgMemOnly,
5085+
NoCapture<ArgIndex<0>>, NoCapture<ArgIndex<2>>]>;
5086+
5087+
def int_nvvm_tcgen05_relinq_alloc_permit_ # cta_group : Intrinsic<[],
5088+
[llvm_tmem_ptr_ty], // tmem_token
5089+
[IntrConvergent, IntrArgMemOnly, NoCapture<ArgIndex<0>>]>;
5090+
}
5091+
50585092
} // let TargetPrefix = "nvvm"

‎llvm/include/llvm/Support/NVPTXAddrSpace.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -23,6 +23,7 @@ enum AddressSpace : unsigned {
2323
ADDRESS_SPACE_SHARED = 3,
2424
ADDRESS_SPACE_CONST = 4,
2525
ADDRESS_SPACE_LOCAL = 5,
26+
ADDRESS_SPACE_TENSOR = 6,
2627

2728
ADDRESS_SPACE_PARAM = 101,
2829
};

‎llvm/lib/Target/NVPTX/NVPTXInstrInfo.td

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -163,6 +163,7 @@ def do_SQRTF32_RN : Predicate<"usePrecSqrtF32()">;
163163
def hasHWROT32 : Predicate<"Subtarget->hasHWROT32()">;
164164
def noHWROT32 : Predicate<"!Subtarget->hasHWROT32()">;
165165
def hasDotInstructions : Predicate<"Subtarget->hasDotInstructions()">;
166+
def hasTcgen05Instructions : Predicate<"Subtarget->hasTcgen05Instructions()">;
166167

167168
def True : Predicate<"true">;
168169
def False : Predicate<"false">;

‎llvm/lib/Target/NVPTX/NVPTXIntrinsics.td

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7582,3 +7582,45 @@ def GRIDDEPCONTROL_WAIT :
75827582
Requires<[hasSM<90>, hasPTX<78>]>;
75837583

75847584
def INT_EXIT : NVPTXInst<(outs), (ins), "exit;", [(int_nvvm_exit)]>;
7585+
7586+
// Tcgen05 intrinsics
7587+
let isConvergent = true in {
7588+
7589+
multiclass TCGEN05_ALLOC_INTR<NVPTXRegClass rc, string AS, string num, Intrinsic Intr> {
7590+
def NAME : NVPTXInst<(outs),
7591+
(ins rc:$dst, Int32Regs:$ncols, Int32Regs:$tmem_token),
7592+
!strconcat("tcgen05.alloc.cta_group::", num, ".sync.aligned", AS, ".b32 [$dst], $ncols;"),
7593+
[(Intr rc:$dst, Int32Regs:$ncols, Int32Regs:$tmem_token)]>,
7594+
Requires<[hasTcgen05Instructions]>;
7595+
}
7596+
7597+
defm TCGEN05_ALLOC_CG1 : TCGEN05_ALLOC_INTR<Int64Regs, "", "1", int_nvvm_tcgen05_alloc_cg1>;
7598+
defm TCGEN05_ALLOC_CG2 : TCGEN05_ALLOC_INTR<Int64Regs, "", "2", int_nvvm_tcgen05_alloc_cg2>;
7599+
7600+
defm TCGEN05_ALLOC_S64_CG1 : TCGEN05_ALLOC_INTR<Int64Regs, ".shared::cta", "1", int_nvvm_tcgen05_alloc_shared_cg1>;
7601+
defm TCGEN05_ALLOC_S64_CG2 : TCGEN05_ALLOC_INTR<Int64Regs, ".shared::cta", "2", int_nvvm_tcgen05_alloc_shared_cg2>;
7602+
7603+
defm TCGEN05_ALLOC_S32_CG1 : TCGEN05_ALLOC_INTR<Int32Regs, ".shared::cta", "1", int_nvvm_tcgen05_alloc_shared_cg1>;
7604+
defm TCGEN05_ALLOC_S32_CG2 : TCGEN05_ALLOC_INTR<Int32Regs, ".shared::cta", "2", int_nvvm_tcgen05_alloc_shared_cg2>;
7605+
7606+
multiclass TCGEN05_DEALLOC_INTR<string num, Intrinsic Intr> {
7607+
def NAME : NVPTXInst<(outs),
7608+
(ins Int32Regs:$tmem_addr, Int32Regs:$ncols, Int32Regs:$tmem_token),
7609+
!strconcat("tcgen05.dealloc.cta_group::", num, ".sync.aligned.b32 $tmem_addr, $ncols;"),
7610+
[(Intr Int32Regs:$tmem_addr, Int32Regs:$ncols, Int32Regs:$tmem_token)]>,
7611+
Requires<[hasTcgen05Instructions]>;
7612+
}
7613+
defm TCGEN05_DEALLOC_CG1: TCGEN05_DEALLOC_INTR<"1", int_nvvm_tcgen05_dealloc_cg1>;
7614+
defm TCGEN05_DEALLOC_CG2: TCGEN05_DEALLOC_INTR<"2", int_nvvm_tcgen05_dealloc_cg2>;
7615+
7616+
multiclass TCGEN05_RELINQ_PERMIT_INTR<string num, Intrinsic Intr> {
7617+
def NAME : NVPTXInst<(outs),
7618+
(ins Int32Regs:$tmem_token),
7619+
!strconcat("tcgen05.relinquish_alloc_permit.cta_group::", num, ".sync.aligned;"),
7620+
[(Intr Int32Regs:$tmem_token)]>,
7621+
Requires<[hasTcgen05Instructions]>;
7622+
}
7623+
defm TCGEN05_RELINQ_CG1: TCGEN05_RELINQ_PERMIT_INTR<"1", int_nvvm_tcgen05_relinq_alloc_permit_cg1>;
7624+
defm TCGEN05_RELINQ_CG2: TCGEN05_RELINQ_PERMIT_INTR<"2", int_nvvm_tcgen05_relinq_alloc_permit_cg2>;
7625+
7626+
} // isConvergent

‎llvm/lib/Target/NVPTX/NVPTXSubtarget.h

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -93,6 +93,21 @@ class NVPTXSubtarget : public NVPTXGenSubtargetInfo {
9393
bool hasDotInstructions() const {
9494
return SmVersion >= 61 && PTXVersion >= 50;
9595
}
96+
// Tcgen05 instructions in Blackwell family
97+
bool hasTcgen05Instructions() const {
98+
bool HasTcgen05 = false;
99+
switch (FullSmVersion) {
100+
default:
101+
break;
102+
case 1001: // sm_100a
103+
case 1011: // sm_101a
104+
HasTcgen05 = true;
105+
break;
106+
}
107+
108+
return HasTcgen05 && PTXVersion >= 86;
109+
}
110+
96111
// Prior to CUDA 12.3 ptxas did not recognize that the trap instruction
97112
// terminates a basic block. Instead, it would assume that control flow
98113
// continued to the next instruction. The next instruction could be in the

‎llvm/lib/Target/NVPTX/NVPTXTargetMachine.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -139,6 +139,9 @@ static std::string computeDataLayout(bool is64Bit, bool UseShortPointers) {
139139
else if (UseShortPointers)
140140
Ret += "-p3:32:32-p4:32:32-p5:32:32";
141141

142+
// Tensor Memory (addrspace:6) is always 32-bits.
143+
Ret += "-p6:32:32";
144+
142145
Ret += "-i64:64-i128:128-v16:16-v32:32-n16:32:64";
143146

144147
return Ret;

‎llvm/test/CodeGen/NVPTX/tcgen05-alloc.ll

Lines changed: 139 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,139 @@
1+
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2+
; RUN: llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | FileCheck --check-prefixes=CHECK_PTX64 %s
3+
; RUN: llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 --nvptx-short-ptr | FileCheck --check-prefixes=CHECK_PTX64_SHARED32 %s
4+
; RUN: %if ptxas-12.8 %{ llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 | %ptxas-verify -arch=sm_100a %}
5+
; RUN: %if ptxas-12.8 %{ llc < %s -march=nvptx64 -mcpu=sm_100a -mattr=+ptx86 --nvptx-short-ptr | %ptxas-verify -arch=sm_100a %}
6+
7+
declare void @llvm.nvvm.tcgen05.alloc.cg1(ptr %addr, i32 %ncols, ptr addrspace(6) %tmem_token)
8+
declare void @llvm.nvvm.tcgen05.alloc.cg2(ptr %addr, i32 %ncols, ptr addrspace(6) %tmem_token)
9+
declare void @llvm.nvvm.tcgen05.alloc.shared.cg1(ptr addrspace(3) %addr, i32 %ncols, ptr addrspace(6) %tmem_token)
10+
declare void @llvm.nvvm.tcgen05.alloc.shared.cg2(ptr addrspace(3) %addr, i32 %ncols, ptr addrspace(6) %tmem_token)
11+
12+
; CHECK-LABEL: test_tcgen05_alloc
13+
define void @test_tcgen05_alloc(ptr %addr, i32 %ncols) {
14+
; CHECK_PTX64-LABEL: test_tcgen05_alloc(
15+
; CHECK_PTX64: {
16+
; CHECK_PTX64-NEXT: .reg .b32 %r<3>;
17+
; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;
18+
; CHECK_PTX64-EMPTY:
19+
; CHECK_PTX64-NEXT: // %bb.0:
20+
; CHECK_PTX64-NEXT: ld.param.u64 %rd1, [test_tcgen05_alloc_param_0];
21+
; CHECK_PTX64-NEXT: ld.param.u32 %r1, [test_tcgen05_alloc_param_1];
22+
; CHECK_PTX64-NEXT: mov.b32 %r2, 0;
23+
; CHECK_PTX64-NEXT: tcgen05.alloc.cta_group::1.sync.aligned.b32 [%rd1], %r1;
24+
; CHECK_PTX64-NEXT: tcgen05.alloc.cta_group::2.sync.aligned.b32 [%rd1], %r1;
25+
; CHECK_PTX64-NEXT: ret;
26+
;
27+
; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_alloc(
28+
; CHECK_PTX64_SHARED32: {
29+
; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<3>;
30+
; CHECK_PTX64_SHARED32-NEXT: .reg .b64 %rd<2>;
31+
; CHECK_PTX64_SHARED32-EMPTY:
32+
; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
33+
; CHECK_PTX64_SHARED32-NEXT: ld.param.u64 %rd1, [test_tcgen05_alloc_param_0];
34+
; CHECK_PTX64_SHARED32-NEXT: ld.param.u32 %r1, [test_tcgen05_alloc_param_1];
35+
; CHECK_PTX64_SHARED32-NEXT: mov.b32 %r2, 0;
36+
; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.cta_group::1.sync.aligned.b32 [%rd1], %r1;
37+
; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.cta_group::2.sync.aligned.b32 [%rd1], %r1;
38+
; CHECK_PTX64_SHARED32-NEXT: ret;
39+
call void @llvm.nvvm.tcgen05.alloc.cg1(ptr %addr, i32 %ncols, ptr addrspace(6) null)
40+
call void @llvm.nvvm.tcgen05.alloc.cg2(ptr %addr, i32 %ncols, ptr addrspace(6) null)
41+
42+
ret void
43+
}
44+
45+
; CHECK-LABEL: test_tcgen05_alloc_shared
46+
define void @test_tcgen05_alloc_shared(ptr addrspace(3) %addr, i32 %ncols) {
47+
; CHECK_PTX64-LABEL: test_tcgen05_alloc_shared(
48+
; CHECK_PTX64: {
49+
; CHECK_PTX64-NEXT: .reg .b32 %r<3>;
50+
; CHECK_PTX64-NEXT: .reg .b64 %rd<2>;
51+
; CHECK_PTX64-EMPTY:
52+
; CHECK_PTX64-NEXT: // %bb.0:
53+
; CHECK_PTX64-NEXT: ld.param.u64 %rd1, [test_tcgen05_alloc_shared_param_0];
54+
; CHECK_PTX64-NEXT: ld.param.u32 %r1, [test_tcgen05_alloc_shared_param_1];
55+
; CHECK_PTX64-NEXT: mov.b32 %r2, 0;
56+
; CHECK_PTX64-NEXT: tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%rd1], %r1;
57+
; CHECK_PTX64-NEXT: tcgen05.alloc.cta_group::2.sync.aligned.shared::cta.b32 [%rd1], %r1;
58+
; CHECK_PTX64-NEXT: ret;
59+
;
60+
; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_alloc_shared(
61+
; CHECK_PTX64_SHARED32: {
62+
; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<4>;
63+
; CHECK_PTX64_SHARED32-EMPTY:
64+
; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
65+
; CHECK_PTX64_SHARED32-NEXT: ld.param.u32 %r1, [test_tcgen05_alloc_shared_param_0];
66+
; CHECK_PTX64_SHARED32-NEXT: ld.param.u32 %r2, [test_tcgen05_alloc_shared_param_1];
67+
; CHECK_PTX64_SHARED32-NEXT: mov.b32 %r3, 0;
68+
; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.cta_group::1.sync.aligned.shared::cta.b32 [%r1], %r2;
69+
; CHECK_PTX64_SHARED32-NEXT: tcgen05.alloc.cta_group::2.sync.aligned.shared::cta.b32 [%r1], %r2;
70+
; CHECK_PTX64_SHARED32-NEXT: ret;
71+
call void @llvm.nvvm.tcgen05.alloc.shared.cg1(ptr addrspace(3) %addr, i32 %ncols, ptr addrspace(6) null)
72+
73+
call void @llvm.nvvm.tcgen05.alloc.shared.cg2(ptr addrspace(3) %addr, i32 %ncols, ptr addrspace(6) null)
74+
ret void
75+
}
76+
77+
declare void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols, ptr addrspace(6) %tmem_token)
78+
declare void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols, ptr addrspace(6) %tmem_token)
79+
80+
; CHECK-LABEL: test_tcgen05_dealloc
81+
define void @test_tcgen05_dealloc(ptr addrspace(6) %tmem_addr, i32 %ncols) {
82+
; CHECK_PTX64-LABEL: test_tcgen05_dealloc(
83+
; CHECK_PTX64: {
84+
; CHECK_PTX64-NEXT: .reg .b32 %r<4>;
85+
; CHECK_PTX64-EMPTY:
86+
; CHECK_PTX64-NEXT: // %bb.0:
87+
; CHECK_PTX64-NEXT: ld.param.u32 %r1, [test_tcgen05_dealloc_param_0];
88+
; CHECK_PTX64-NEXT: ld.param.u32 %r2, [test_tcgen05_dealloc_param_1];
89+
; CHECK_PTX64-NEXT: mov.b32 %r3, 0;
90+
; CHECK_PTX64-NEXT: tcgen05.dealloc.cta_group::1.sync.aligned.b32 %r1, %r2;
91+
; CHECK_PTX64-NEXT: tcgen05.dealloc.cta_group::2.sync.aligned.b32 %r1, %r2;
92+
; CHECK_PTX64-NEXT: ret;
93+
;
94+
; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_dealloc(
95+
; CHECK_PTX64_SHARED32: {
96+
; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<4>;
97+
; CHECK_PTX64_SHARED32-EMPTY:
98+
; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
99+
; CHECK_PTX64_SHARED32-NEXT: ld.param.u32 %r1, [test_tcgen05_dealloc_param_0];
100+
; CHECK_PTX64_SHARED32-NEXT: ld.param.u32 %r2, [test_tcgen05_dealloc_param_1];
101+
; CHECK_PTX64_SHARED32-NEXT: mov.b32 %r3, 0;
102+
; CHECK_PTX64_SHARED32-NEXT: tcgen05.dealloc.cta_group::1.sync.aligned.b32 %r1, %r2;
103+
; CHECK_PTX64_SHARED32-NEXT: tcgen05.dealloc.cta_group::2.sync.aligned.b32 %r1, %r2;
104+
; CHECK_PTX64_SHARED32-NEXT: ret;
105+
call void @llvm.nvvm.tcgen05.dealloc.cg1(ptr addrspace(6) %tmem_addr, i32 %ncols, ptr addrspace(6) null)
106+
107+
call void @llvm.nvvm.tcgen05.dealloc.cg2(ptr addrspace(6) %tmem_addr, i32 %ncols, ptr addrspace(6) null)
108+
ret void
109+
}
110+
111+
declare void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg1(ptr addrspace(6) %tmem_token)
112+
declare void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg2(ptr addrspace(6) %tmem_token)
113+
114+
; CHECK-LABEL: test_tcgen05_relinquish_alloc_permit
115+
define void @test_tcgen05_relinquish_alloc_permit() {
116+
; CHECK_PTX64-LABEL: test_tcgen05_relinquish_alloc_permit(
117+
; CHECK_PTX64: {
118+
; CHECK_PTX64-NEXT: .reg .b32 %r<2>;
119+
; CHECK_PTX64-EMPTY:
120+
; CHECK_PTX64-NEXT: // %bb.0:
121+
; CHECK_PTX64-NEXT: mov.b32 %r1, 0;
122+
; CHECK_PTX64-NEXT: tcgen05.relinquish_alloc_permit.cta_group::1.sync.aligned;
123+
; CHECK_PTX64-NEXT: tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;
124+
; CHECK_PTX64-NEXT: ret;
125+
;
126+
; CHECK_PTX64_SHARED32-LABEL: test_tcgen05_relinquish_alloc_permit(
127+
; CHECK_PTX64_SHARED32: {
128+
; CHECK_PTX64_SHARED32-NEXT: .reg .b32 %r<2>;
129+
; CHECK_PTX64_SHARED32-EMPTY:
130+
; CHECK_PTX64_SHARED32-NEXT: // %bb.0:
131+
; CHECK_PTX64_SHARED32-NEXT: mov.b32 %r1, 0;
132+
; CHECK_PTX64_SHARED32-NEXT: tcgen05.relinquish_alloc_permit.cta_group::1.sync.aligned;
133+
; CHECK_PTX64_SHARED32-NEXT: tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;
134+
; CHECK_PTX64_SHARED32-NEXT: ret;
135+
call void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg1(ptr addrspace(6) null)
136+
137+
call void @llvm.nvvm.tcgen05.relinq.alloc.permit.cg2(ptr addrspace(6) null)
138+
ret void
139+
}

0 commit comments

Comments
 (0)