1; RUN: llc < %s -march=nvptx -mcpu=sm_80 -mattr=+ptx70 | FileCheck -check-prefixes=ALL,CHECK_PTX32 %s
2; RUN: llc < %s -march=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | FileCheck -check-prefixes=ALL,CHECK_PTX64 %s
3; RUN: %if ptxas-11.0 %{ llc < %s -march=nvptx -mcpu=sm_80 -mattr=+ptx70 | %ptxas-verify -arch=sm_80 %}
4; RUN: %if ptxas-11.0 %{ llc < %s -march=nvptx64 -mcpu=sm_80 -mattr=+ptx70 | %ptxas-verify -arch=sm_80 %}
5
6declare void @llvm.nvvm.cp.async.wait.group(i32)
7
8; ALL-LABEL: asyncwaitgroup
9define void @asyncwaitgroup() {
10  ; ALL: cp.async.wait_group 8;
11  tail call void @llvm.nvvm.cp.async.wait.group(i32 8)
12  ; ALL: cp.async.wait_group 0;
13  tail call void @llvm.nvvm.cp.async.wait.group(i32 0)
14  ; ALL: cp.async.wait_group 16;
15  tail call void @llvm.nvvm.cp.async.wait.group(i32 16)
16  ret void
17}
18
19declare void @llvm.nvvm.cp.async.wait.all()
20
21; ALL-LABEL: asyncwaitall
22define void @asyncwaitall() {
23; ALL: cp.async.wait_all
24  tail call void @llvm.nvvm.cp.async.wait.all()
25  ret void
26}
27
28declare void @llvm.nvvm.cp.async.commit.group()
29
30; ALL-LABEL: asynccommitgroup
31define void @asynccommitgroup() {
32; ALL: cp.async.commit_group
33  tail call void @llvm.nvvm.cp.async.commit.group()
34  ret void
35}
36
37declare void @llvm.nvvm.cp.async.mbarrier.arrive(i64* %a)
38declare void @llvm.nvvm.cp.async.mbarrier.arrive.shared(i64 addrspace(3)* %a)
39declare void @llvm.nvvm.cp.async.mbarrier.arrive.noinc(i64* %a)
40declare void @llvm.nvvm.cp.async.mbarrier.arrive.noinc.shared(i64 addrspace(3)* %a)
41
42; CHECK-LABEL: asyncmbarrier
43define void @asyncmbarrier(i64* %a) {
44; CHECK_PTX32: cp.async.mbarrier.arrive.b64 [%r{{[0-9]+}}];
45; CHECK_PTX64: cp.async.mbarrier.arrive.b64 [%rd{{[0-9]+}}];
46  tail call void @llvm.nvvm.cp.async.mbarrier.arrive(i64* %a)
47  ret void
48}
49
50; CHECK-LABEL: asyncmbarriershared
51define void @asyncmbarriershared(i64 addrspace(3)* %a) {
52; CHECK_PTX32: cp.async.mbarrier.arrive.shared.b64 [%r{{[0-9]+}}];
53; CHECK_PTX64: cp.async.mbarrier.arrive.shared.b64 [%rd{{[0-9]+}}];
54  tail call void @llvm.nvvm.cp.async.mbarrier.arrive.shared(i64 addrspace(3)* %a)
55  ret void
56}
57
58; CHECK-LABEL: asyncmbarriernoinc
59define void @asyncmbarriernoinc(i64* %a) {
60; CHECK_PTX32: cp.async.mbarrier.arrive.noinc.b64 [%r{{[0-9]+}}];
61; CHECK_PTX64: cp.async.mbarrier.arrive.noinc.b64 [%rd{{[0-9]+}}];
62  tail call void @llvm.nvvm.cp.async.mbarrier.arrive.noinc(i64* %a)
63  ret void
64}
65
66; CHECK-LABEL: asyncmbarriernoincshared
67define void @asyncmbarriernoincshared(i64 addrspace(3)* %a) {
68; CHECK_PTX32: cp.async.mbarrier.arrive.noinc.shared.b64 [%r{{[0-9]+}}];
69; CHECK_PTX64: cp.async.mbarrier.arrive.noinc.shared.b64 [%rd{{[0-9]+}}];
70  tail call void @llvm.nvvm.cp.async.mbarrier.arrive.noinc.shared(i64 addrspace(3)* %a)
71  ret void
72}
73
74declare void @llvm.nvvm.cp.async.ca.shared.global.4(i8 addrspace(3)* %a, i8 addrspace(1)* %b)
75
76; CHECK-LABEL: asynccasharedglobal4i8
77define void @asynccasharedglobal4i8(i8 addrspace(3)* %a, i8 addrspace(1)* %b) {
78; CHECK_PTX32: cp.async.ca.shared.global [%r{{[0-9]+}}], [%r{{[0-9]+}}], 4;
79; CHECK_PTX64: cp.async.ca.shared.global [%rd{{[0-9]+}}], [%rd{{[0-9]+}}], 4;
80  tail call void @llvm.nvvm.cp.async.ca.shared.global.4(i8 addrspace(3)* %a, i8 addrspace(1)* %b)
81  ret void
82}
83
84declare void @llvm.nvvm.cp.async.ca.shared.global.8(i8 addrspace(3)* %a, i8 addrspace(1)* %b)
85
86; CHECK-LABEL: asynccasharedglobal8i8
87define void @asynccasharedglobal8i8(i8 addrspace(3)* %a, i8 addrspace(1)* %b) {
88; CHECK_PTX32: cp.async.ca.shared.global [%r{{[0-9]+}}], [%r{{[0-9]+}}], 8;
89; CHECK_PTX64: cp.async.ca.shared.global [%rd{{[0-9]+}}], [%rd{{[0-9]+}}], 8;
90  tail call void @llvm.nvvm.cp.async.ca.shared.global.8(i8 addrspace(3)* %a, i8 addrspace(1)* %b)
91  ret void
92}
93
94declare void @llvm.nvvm.cp.async.ca.shared.global.16(i8 addrspace(3)* %a, i8 addrspace(1)* %b)
95
96; CHECK-LABEL: asynccasharedglobal16i8
97define void @asynccasharedglobal16i8(i8 addrspace(3)* %a, i8 addrspace(1)* %b) {
98; CHECK_PTX32: cp.async.ca.shared.global [%r{{[0-9]+}}], [%r{{[0-9]+}}], 16;
99; CHECK_PTX64: cp.async.ca.shared.global [%rd{{[0-9]+}}], [%rd{{[0-9]+}}], 16;
100  tail call void @llvm.nvvm.cp.async.ca.shared.global.16(i8 addrspace(3)* %a, i8 addrspace(1)* %b)
101  ret void
102}
103
104declare void @llvm.nvvm.cp.async.cg.shared.global.16(i8 addrspace(3)* %a, i8 addrspace(1)* %b)
105
106; CHECK-LABEL: asynccgsharedglobal16i8
107define void @asynccgsharedglobal16i8(i8 addrspace(3)* %a, i8 addrspace(1)* %b) {
108; CHECK_PTX32: cp.async.cg.shared.global [%r{{[0-9]+}}], [%r{{[0-9]+}}], 16;
109; CHECK_PTX64: cp.async.cg.shared.global [%rd{{[0-9]+}}], [%rd{{[0-9]+}}], 16;
110  tail call void @llvm.nvvm.cp.async.cg.shared.global.16(i8 addrspace(3)* %a, i8 addrspace(1)* %b)
111  ret void
112}
113