161 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80| FileCheck --check-prefixes=CHECK-PTX64 %s
3; RUN: %if ptxas-sm_90 && ptxas-isa-8.0 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx80| %ptxas-verify -arch=sm_90 %}
4
5target triple = "nvptx64-nvidia-cuda"
6
7declare void @llvm.nvvm.prefetch.global.L1(ptr addrspace(1) %global_ptr)
8declare void @llvm.nvvm.prefetch.global.L2(ptr addrspace(1) %global_ptr)
9declare void @llvm.nvvm.prefetch.local.L1(ptr addrspace(5) %local_ptr)
10declare void @llvm.nvvm.prefetch.local.L2(ptr addrspace(5) %local_ptr)
11
12declare void @llvm.nvvm.prefetch.L1(ptr %ptr)
13declare void @llvm.nvvm.prefetch.L2(ptr %ptr)
14
15declare void @llvm.nvvm.prefetch.tensormap.p0(ptr %ptr)
16declare void @llvm.nvvm.prefetch.tensormap.p4(ptr addrspace(4) %const_ptr)
17declare void @llvm.nvvm.prefetch.tensormap.p101(ptr addrspace(101) %param_ptr)
18
19declare void @llvm.nvvm.prefetch.global.L2.evict.normal(ptr addrspace(1) %global_ptr)
20declare void @llvm.nvvm.prefetch.global.L2.evict.last(ptr addrspace(1) %global_ptr)
21
22declare void @llvm.nvvm.prefetchu.L1(ptr %ptr)
23
24define void @prefetch_local(ptr addrspace(5) %local_ptr) {
25; CHECK-PTX64-LABEL: prefetch_local(
26; CHECK-PTX64: {
27; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
28; CHECK-PTX64-EMPTY:
29; CHECK-PTX64-NEXT: // %bb.0:
30; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_local_param_0];
31; CHECK-PTX64-NEXT: prefetch.local.L1 [%rd1];
32; CHECK-PTX64-NEXT: prefetch.local.L2 [%rd1];
33; CHECK-PTX64-NEXT: ret;
34 tail call void @llvm.nvvm.prefetch.local.L1(ptr addrspace(5) %local_ptr)
35 tail call void @llvm.nvvm.prefetch.local.L2(ptr addrspace(5) %local_ptr)
36 ret void
37}
38
39define void @prefetch_global(ptr addrspace(1) %global_ptr) {
40; CHECK-PTX64-LABEL: prefetch_global(
41; CHECK-PTX64: {
42; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
43; CHECK-PTX64-EMPTY:
44; CHECK-PTX64-NEXT: // %bb.0:
45; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_global_param_0];
46; CHECK-PTX64-NEXT: prefetch.global.L1 [%rd1];
47; CHECK-PTX64-NEXT: prefetch.global.L2 [%rd1];
48; CHECK-PTX64-NEXT: prefetch.global.L2::evict_normal [%rd1];
49; CHECK-PTX64-NEXT: prefetch.global.L2::evict_last [%rd1];
50; CHECK-PTX64-NEXT: ret;
51 tail call void @llvm.nvvm.prefetch.global.L1(ptr addrspace(1) %global_ptr)
52 tail call void @llvm.nvvm.prefetch.global.L2(ptr addrspace(1) %global_ptr)
53 tail call void @llvm.nvvm.prefetch.global.L2.evict.normal(ptr addrspace(1) %global_ptr)
54 tail call void @llvm.nvvm.prefetch.global.L2.evict.last(ptr addrspace(1) %global_ptr)
55 ret void
56}
57
58
59define void @prefetch_(ptr %ptr) {
60; CHECK-PTX64-LABEL: prefetch_(
61; CHECK-PTX64: {
62; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
63; CHECK-PTX64-EMPTY:
64; CHECK-PTX64-NEXT: // %bb.0:
65; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch__param_0];
66; CHECK-PTX64-NEXT: prefetch.L1 [%rd1];
67; CHECK-PTX64-NEXT: prefetch.L2 [%rd1];
68; CHECK-PTX64-NEXT: ret;
69 tail call void @llvm.nvvm.prefetch.L1(ptr %ptr)
70 tail call void @llvm.nvvm.prefetch.L2(ptr %ptr)
71 ret void
72}
73
74define void @prefetchu_l1(ptr %ptr) {
75; CHECK-PTX64-LABEL: prefetchu_l1(
76; CHECK-PTX64: {
77; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
78; CHECK-PTX64-EMPTY:
79; CHECK-PTX64-NEXT: // %bb.0:
80; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetchu_l1_param_0];
81; CHECK-PTX64-NEXT: prefetchu.L1 [%rd1];
82; CHECK-PTX64-NEXT: ret;
83 tail call void @llvm.nvvm.prefetchu.L1(ptr %ptr)
84 ret void
85}
86
87define void @prefetch_tensormap(ptr %ptr) {
88; CHECK-PTX64-LABEL: prefetch_tensormap(
89; CHECK-PTX64: {
90; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
91; CHECK-PTX64-EMPTY:
92; CHECK-PTX64-NEXT: // %bb.0:
93; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_tensormap_param_0];
94; CHECK-PTX64-NEXT: prefetch.tensormap [%rd1];
95; CHECK-PTX64-NEXT: ret;
96 tail call void @llvm.nvvm.prefetch.tensormap.p0(ptr %ptr)
97 ret void
98}
99
100define void @prefetch_const_tensormap(ptr addrspace(4) %const_ptr) {
101; CHECK-PTX64-LABEL: prefetch_const_tensormap(
102; CHECK-PTX64: {
103; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
104; CHECK-PTX64-EMPTY:
105; CHECK-PTX64-NEXT: // %bb.0:
106; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_const_tensormap_param_0];
107; CHECK-PTX64-NEXT: prefetch.const.tensormap [%rd1];
108; CHECK-PTX64-NEXT: ret;
109 tail call void @llvm.nvvm.prefetch.tensormap.p4(ptr addrspace(4) %const_ptr)
110 ret void
111}
112
113define void @prefetch_param_tensormap(ptr addrspace(101) %param_ptr) {
114; CHECK-PTX64-LABEL: prefetch_param_tensormap(
115; CHECK-PTX64: {
116; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
117; CHECK-PTX64-EMPTY:
118; CHECK-PTX64-NEXT: // %bb.0:
119; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_param_tensormap_param_0];
120; CHECK-PTX64-NEXT: prefetch.param.tensormap [%rd1];
121; CHECK-PTX64-NEXT: ret;
122 tail call void @llvm.nvvm.prefetch.tensormap.p101(ptr addrspace(101) %param_ptr)
123 ret void
124}
125
126define ptx_kernel void @prefetch_generic_tensormap_kernel(ptr %ptr) {
127; CHECK-PTX64-LABEL: prefetch_generic_tensormap_kernel(
128; CHECK-PTX64: {
129; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
130; CHECK-PTX64-EMPTY:
131; CHECK-PTX64-NEXT: // %bb.0:
132; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_generic_tensormap_kernel_param_0];
133; CHECK-PTX64-NEXT: prefetch.tensormap [%rd1];
134; CHECK-PTX64-NEXT: ret;
135 tail call void @llvm.nvvm.prefetch.tensormap.p0(ptr %ptr)
136 ret void
137}
138
139define ptx_kernel void @prefetch_param_tensormap_kernel(ptr addrspace(101) %param_ptr) {
140; CHECK-PTX64-LABEL: prefetch_param_tensormap_kernel(
141; CHECK-PTX64: {
142; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
143; CHECK-PTX64-EMPTY:
144; CHECK-PTX64-NEXT: // %bb.0:
145; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_param_tensormap_kernel_param_0];
146; CHECK-PTX64-NEXT: prefetch.param.tensormap [%rd1];
147; CHECK-PTX64-NEXT: ret;
148 tail call void @llvm.nvvm.prefetch.tensormap.p101(ptr addrspace(101) %param_ptr)
149 ret void
150}
151
152define ptx_kernel void @prefetch_grid_const_tensormap(ptr byval([64 x i8]) align 64 "nvvm.grid_constant" %ptr) {
153; CHECK-PTX64-LABEL: .visible .entry prefetch_grid_const_tensormap(
154; CHECK-PTX64: prefetch.tensormap [%{{(SP|rd[0-9]+).*}}];
155; CHECK-PTX64: ret;
156
157entry:
158 call void @llvm.nvvm.prefetch.tensormap.p0(ptr addrspace(0) %ptr)
159 ret void
160}
161