1 //===- IntrinsicsNVVM.td - Defines NVVM intrinsics ---------*- tablegen -*-===//
3 // The LLVM Compiler Infrastructure
5 // This file is distributed under the University of Illinois Open Source
6 // License. See LICENSE.TXT for details.
8 //===----------------------------------------------------------------------===//
10 // This file defines all of the NVVM-specific intrinsics for use with NVPTX.
12 //===----------------------------------------------------------------------===//
14 def llvm_anyi64ptr_ty : LLVMAnyPointerType<llvm_i64_ty>; // (space)i64*
20 let TargetPrefix = "nvvm" in {
21 def int_nvvm_clz_i : GCCBuiltin<"__nvvm_clz_i">,
22 Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem]>;
23 def int_nvvm_clz_ll : GCCBuiltin<"__nvvm_clz_ll">,
24 Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem]>;
26 def int_nvvm_popc_i : GCCBuiltin<"__nvvm_popc_i">,
27 Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem]>;
28 def int_nvvm_popc_ll : GCCBuiltin<"__nvvm_popc_ll">,
29 Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem]>;
31 def int_nvvm_prmt : GCCBuiltin<"__nvvm_prmt">,
32 Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
33 [IntrNoMem, Commutative]>;
39 def int_nvvm_min_i : GCCBuiltin<"__nvvm_min_i">,
40 Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
41 [IntrNoMem, Commutative]>;
42 def int_nvvm_min_ui : GCCBuiltin<"__nvvm_min_ui">,
43 Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
44 [IntrNoMem, Commutative]>;
46 def int_nvvm_min_ll : GCCBuiltin<"__nvvm_min_ll">,
47 Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty],
48 [IntrNoMem, Commutative]>;
49 def int_nvvm_min_ull : GCCBuiltin<"__nvvm_min_ull">,
50 Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty],
51 [IntrNoMem, Commutative]>;
53 def int_nvvm_max_i : GCCBuiltin<"__nvvm_max_i">,
54 Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
55 [IntrNoMem, Commutative]>;
56 def int_nvvm_max_ui : GCCBuiltin<"__nvvm_max_ui">,
57 Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
58 [IntrNoMem, Commutative]>;
60 def int_nvvm_max_ll : GCCBuiltin<"__nvvm_max_ll">,
61 Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty],
62 [IntrNoMem, Commutative]>;
63 def int_nvvm_max_ull : GCCBuiltin<"__nvvm_max_ull">,
64 Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty],
65 [IntrNoMem, Commutative]>;
67 def int_nvvm_fmin_f : GCCBuiltin<"__nvvm_fmin_f">,
68 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
69 [IntrNoMem, Commutative]>;
70 def int_nvvm_fmin_ftz_f : GCCBuiltin<"__nvvm_fmin_ftz_f">,
71 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
72 [IntrNoMem, Commutative]>;
74 def int_nvvm_fmax_f : GCCBuiltin<"__nvvm_fmax_f">,
75 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty]
76 , [IntrNoMem, Commutative]>;
77 def int_nvvm_fmax_ftz_f : GCCBuiltin<"__nvvm_fmax_ftz_f">,
78 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
79 [IntrNoMem, Commutative]>;
81 def int_nvvm_fmin_d : GCCBuiltin<"__nvvm_fmin_d">,
82 Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
83 [IntrNoMem, Commutative]>;
84 def int_nvvm_fmax_d : GCCBuiltin<"__nvvm_fmax_d">,
85 Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
86 [IntrNoMem, Commutative]>;
92 def int_nvvm_mulhi_i : GCCBuiltin<"__nvvm_mulhi_i">,
93 Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
94 [IntrNoMem, Commutative]>;
95 def int_nvvm_mulhi_ui : GCCBuiltin<"__nvvm_mulhi_ui">,
96 Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
97 [IntrNoMem, Commutative]>;
99 def int_nvvm_mulhi_ll : GCCBuiltin<"__nvvm_mulhi_ll">,
100 Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty],
101 [IntrNoMem, Commutative]>;
102 def int_nvvm_mulhi_ull : GCCBuiltin<"__nvvm_mulhi_ull">,
103 Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i64_ty],
104 [IntrNoMem, Commutative]>;
106 def int_nvvm_mul_rn_ftz_f : GCCBuiltin<"__nvvm_mul_rn_ftz_f">,
107 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
108 [IntrNoMem, Commutative]>;
109 def int_nvvm_mul_rn_f : GCCBuiltin<"__nvvm_mul_rn_f">,
110 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
111 [IntrNoMem, Commutative]>;
112 def int_nvvm_mul_rz_ftz_f : GCCBuiltin<"__nvvm_mul_rz_ftz_f">,
113 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
114 [IntrNoMem, Commutative]>;
115 def int_nvvm_mul_rz_f : GCCBuiltin<"__nvvm_mul_rz_f">,
116 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
117 [IntrNoMem, Commutative]>;
118 def int_nvvm_mul_rm_ftz_f : GCCBuiltin<"__nvvm_mul_rm_ftz_f">,
119 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
120 [IntrNoMem, Commutative]>;
121 def int_nvvm_mul_rm_f : GCCBuiltin<"__nvvm_mul_rm_f">,
122 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
123 [IntrNoMem, Commutative]>;
124 def int_nvvm_mul_rp_ftz_f : GCCBuiltin<"__nvvm_mul_rp_ftz_f">,
125 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
126 [IntrNoMem, Commutative]>;
127 def int_nvvm_mul_rp_f : GCCBuiltin<"__nvvm_mul_rp_f">,
128 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
129 [IntrNoMem, Commutative]>;
131 def int_nvvm_mul_rn_d : GCCBuiltin<"__nvvm_mul_rn_d">,
132 Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
133 [IntrNoMem, Commutative]>;
134 def int_nvvm_mul_rz_d : GCCBuiltin<"__nvvm_mul_rz_d">,
135 Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
136 [IntrNoMem, Commutative]>;
137 def int_nvvm_mul_rm_d : GCCBuiltin<"__nvvm_mul_rm_d">,
138 Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
139 [IntrNoMem, Commutative]>;
140 def int_nvvm_mul_rp_d : GCCBuiltin<"__nvvm_mul_rp_d">,
141 Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
142 [IntrNoMem, Commutative]>;
144 def int_nvvm_mul24_i : GCCBuiltin<"__nvvm_mul24_i">,
145 Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
146 [IntrNoMem, Commutative]>;
147 def int_nvvm_mul24_ui : GCCBuiltin<"__nvvm_mul24_ui">,
148 Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
149 [IntrNoMem, Commutative]>;
155 def int_nvvm_div_approx_ftz_f : GCCBuiltin<"__nvvm_div_approx_ftz_f">,
156 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
157 [IntrNoMem, Commutative]>;
158 def int_nvvm_div_approx_f : GCCBuiltin<"__nvvm_div_approx_f">,
159 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
160 [IntrNoMem, Commutative]>;
162 def int_nvvm_div_rn_ftz_f : GCCBuiltin<"__nvvm_div_rn_ftz_f">,
163 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
164 [IntrNoMem, Commutative]>;
165 def int_nvvm_div_rn_f : GCCBuiltin<"__nvvm_div_rn_f">,
166 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
167 [IntrNoMem, Commutative]>;
169 def int_nvvm_div_rz_ftz_f : GCCBuiltin<"__nvvm_div_rz_ftz_f">,
170 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
171 [IntrNoMem, Commutative]>;
172 def int_nvvm_div_rz_f : GCCBuiltin<"__nvvm_div_rz_f">,
173 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
174 [IntrNoMem, Commutative]>;
176 def int_nvvm_div_rm_ftz_f : GCCBuiltin<"__nvvm_div_rm_ftz_f">,
177 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
178 [IntrNoMem, Commutative]>;
179 def int_nvvm_div_rm_f : GCCBuiltin<"__nvvm_div_rm_f">,
180 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
181 [IntrNoMem, Commutative]>;
183 def int_nvvm_div_rp_ftz_f : GCCBuiltin<"__nvvm_div_rp_ftz_f">,
184 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
185 [IntrNoMem, Commutative]>;
186 def int_nvvm_div_rp_f : GCCBuiltin<"__nvvm_div_rp_f">,
187 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
188 [IntrNoMem, Commutative]>;
190 def int_nvvm_div_rn_d : GCCBuiltin<"__nvvm_div_rn_d">,
191 Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
192 [IntrNoMem, Commutative]>;
193 def int_nvvm_div_rz_d : GCCBuiltin<"__nvvm_div_rz_d">,
194 Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
195 [IntrNoMem, Commutative]>;
196 def int_nvvm_div_rm_d : GCCBuiltin<"__nvvm_div_rm_d">,
197 Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
198 [IntrNoMem, Commutative]>;
199 def int_nvvm_div_rp_d : GCCBuiltin<"__nvvm_div_rp_d">,
200 Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
201 [IntrNoMem, Commutative]>;
207 def int_nvvm_brev32 : GCCBuiltin<"__nvvm_brev32">,
208 Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem]>;
209 def int_nvvm_brev64 : GCCBuiltin<"__nvvm_brev64">,
210 Intrinsic<[llvm_i64_ty], [llvm_i64_ty], [IntrNoMem]>;
216 def int_nvvm_sad_i : GCCBuiltin<"__nvvm_sad_i">,
217 Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
218 [IntrNoMem, Commutative]>;
219 def int_nvvm_sad_ui : GCCBuiltin<"__nvvm_sad_ui">,
220 Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
221 [IntrNoMem, Commutative]>;
227 def int_nvvm_floor_ftz_f : GCCBuiltin<"__nvvm_floor_ftz_f">,
228 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
229 def int_nvvm_floor_f : GCCBuiltin<"__nvvm_floor_f">,
230 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
231 def int_nvvm_floor_d : GCCBuiltin<"__nvvm_floor_d">,
232 Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>;
234 def int_nvvm_ceil_ftz_f : GCCBuiltin<"__nvvm_ceil_ftz_f">,
235 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
236 def int_nvvm_ceil_f : GCCBuiltin<"__nvvm_ceil_f">,
237 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
238 def int_nvvm_ceil_d : GCCBuiltin<"__nvvm_ceil_d">,
239 Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>;
245 def int_nvvm_abs_i : GCCBuiltin<"__nvvm_abs_i">,
246 Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem]>;
247 def int_nvvm_abs_ll : GCCBuiltin<"__nvvm_abs_ll">,
248 Intrinsic<[llvm_i64_ty], [llvm_i64_ty], [IntrNoMem]>;
250 def int_nvvm_fabs_ftz_f : GCCBuiltin<"__nvvm_fabs_ftz_f">,
251 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
252 def int_nvvm_fabs_f : GCCBuiltin<"__nvvm_fabs_f">,
253 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
255 def int_nvvm_fabs_d : GCCBuiltin<"__nvvm_fabs_d">,
256 Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>;
262 def int_nvvm_round_ftz_f : GCCBuiltin<"__nvvm_round_ftz_f">,
263 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
264 def int_nvvm_round_f : GCCBuiltin<"__nvvm_round_f">,
265 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
267 def int_nvvm_round_d : GCCBuiltin<"__nvvm_round_d">,
268 Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>;
274 def int_nvvm_trunc_ftz_f : GCCBuiltin<"__nvvm_trunc_ftz_f">,
275 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
276 def int_nvvm_trunc_f : GCCBuiltin<"__nvvm_trunc_f">,
277 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
279 def int_nvvm_trunc_d : GCCBuiltin<"__nvvm_trunc_d">,
280 Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>;
286 def int_nvvm_saturate_ftz_f : GCCBuiltin<"__nvvm_saturate_ftz_f">,
287 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
288 def int_nvvm_saturate_f : GCCBuiltin<"__nvvm_saturate_f">,
289 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
291 def int_nvvm_saturate_d : GCCBuiltin<"__nvvm_saturate_d">,
292 Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>;
298 def int_nvvm_ex2_approx_ftz_f : GCCBuiltin<"__nvvm_ex2_approx_ftz_f">,
299 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
300 def int_nvvm_ex2_approx_f : GCCBuiltin<"__nvvm_ex2_approx_f">,
301 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
302 def int_nvvm_ex2_approx_d : GCCBuiltin<"__nvvm_ex2_approx_d">,
303 Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>;
305 def int_nvvm_lg2_approx_ftz_f : GCCBuiltin<"__nvvm_lg2_approx_ftz_f">,
306 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
307 def int_nvvm_lg2_approx_f : GCCBuiltin<"__nvvm_lg2_approx_f">,
308 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
309 def int_nvvm_lg2_approx_d : GCCBuiltin<"__nvvm_lg2_approx_d">,
310 Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>;
316 def int_nvvm_sin_approx_ftz_f : GCCBuiltin<"__nvvm_sin_approx_ftz_f">,
317 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
318 def int_nvvm_sin_approx_f : GCCBuiltin<"__nvvm_sin_approx_f">,
319 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
321 def int_nvvm_cos_approx_ftz_f : GCCBuiltin<"__nvvm_cos_approx_ftz_f">,
322 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
323 def int_nvvm_cos_approx_f : GCCBuiltin<"__nvvm_cos_approx_f">,
324 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
330 def int_nvvm_fma_rn_ftz_f : GCCBuiltin<"__nvvm_fma_rn_ftz_f">,
331 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty],
332 [IntrNoMem, Commutative]>;
333 def int_nvvm_fma_rn_f : GCCBuiltin<"__nvvm_fma_rn_f">,
334 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty],
335 [IntrNoMem, Commutative]>;
336 def int_nvvm_fma_rz_ftz_f : GCCBuiltin<"__nvvm_fma_rz_ftz_f">,
337 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty],
338 [IntrNoMem, Commutative]>;
339 def int_nvvm_fma_rz_f : GCCBuiltin<"__nvvm_fma_rz_f">,
340 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty],
341 [IntrNoMem, Commutative]>;
342 def int_nvvm_fma_rm_ftz_f : GCCBuiltin<"__nvvm_fma_rm_ftz_f">,
343 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty],
344 [IntrNoMem, Commutative]>;
345 def int_nvvm_fma_rm_f : GCCBuiltin<"__nvvm_fma_rm_f">,
346 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty],
347 [IntrNoMem, Commutative]>;
348 def int_nvvm_fma_rp_ftz_f : GCCBuiltin<"__nvvm_fma_rp_ftz_f">,
349 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty],
350 [IntrNoMem, Commutative]>;
351 def int_nvvm_fma_rp_f : GCCBuiltin<"__nvvm_fma_rp_f">,
352 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty, llvm_float_ty],
353 [IntrNoMem, Commutative]>;
355 def int_nvvm_fma_rn_d : GCCBuiltin<"__nvvm_fma_rn_d">,
356 Intrinsic<[llvm_double_ty],
357 [llvm_double_ty, llvm_double_ty, llvm_double_ty],
358 [IntrNoMem, Commutative]>;
359 def int_nvvm_fma_rz_d : GCCBuiltin<"__nvvm_fma_rz_d">,
360 Intrinsic<[llvm_double_ty],
361 [llvm_double_ty, llvm_double_ty, llvm_double_ty],
362 [IntrNoMem, Commutative]>;
363 def int_nvvm_fma_rm_d : GCCBuiltin<"__nvvm_fma_rm_d">,
364 Intrinsic<[llvm_double_ty],
365 [llvm_double_ty, llvm_double_ty, llvm_double_ty],
366 [IntrNoMem, Commutative]>;
367 def int_nvvm_fma_rp_d : GCCBuiltin<"__nvvm_fma_rp_d">,
368 Intrinsic<[llvm_double_ty],
369 [llvm_double_ty, llvm_double_ty, llvm_double_ty],
370 [IntrNoMem, Commutative]>;
376 def int_nvvm_rcp_rn_ftz_f : GCCBuiltin<"__nvvm_rcp_rn_ftz_f">,
377 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
378 def int_nvvm_rcp_rn_f : GCCBuiltin<"__nvvm_rcp_rn_f">,
379 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
380 def int_nvvm_rcp_rz_ftz_f : GCCBuiltin<"__nvvm_rcp_rz_ftz_f">,
381 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
382 def int_nvvm_rcp_rz_f : GCCBuiltin<"__nvvm_rcp_rz_f">,
383 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
384 def int_nvvm_rcp_rm_ftz_f : GCCBuiltin<"__nvvm_rcp_rm_ftz_f">,
385 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
386 def int_nvvm_rcp_rm_f : GCCBuiltin<"__nvvm_rcp_rm_f">,
387 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
388 def int_nvvm_rcp_rp_ftz_f : GCCBuiltin<"__nvvm_rcp_rp_ftz_f">,
389 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
390 def int_nvvm_rcp_rp_f : GCCBuiltin<"__nvvm_rcp_rp_f">,
391 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
393 def int_nvvm_rcp_rn_d : GCCBuiltin<"__nvvm_rcp_rn_d">,
394 Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>;
395 def int_nvvm_rcp_rz_d : GCCBuiltin<"__nvvm_rcp_rz_d">,
396 Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>;
397 def int_nvvm_rcp_rm_d : GCCBuiltin<"__nvvm_rcp_rm_d">,
398 Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>;
399 def int_nvvm_rcp_rp_d : GCCBuiltin<"__nvvm_rcp_rp_d">,
400 Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>;
402 def int_nvvm_rcp_approx_ftz_d : GCCBuiltin<"__nvvm_rcp_approx_ftz_d">,
403 Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>;
409 def int_nvvm_sqrt_f : GCCBuiltin<"__nvvm_sqrt_f">,
410 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
411 def int_nvvm_sqrt_rn_ftz_f : GCCBuiltin<"__nvvm_sqrt_rn_ftz_f">,
412 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
413 def int_nvvm_sqrt_rn_f : GCCBuiltin<"__nvvm_sqrt_rn_f">,
414 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
415 def int_nvvm_sqrt_rz_ftz_f : GCCBuiltin<"__nvvm_sqrt_rz_ftz_f">,
416 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
417 def int_nvvm_sqrt_rz_f : GCCBuiltin<"__nvvm_sqrt_rz_f">,
418 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
419 def int_nvvm_sqrt_rm_ftz_f : GCCBuiltin<"__nvvm_sqrt_rm_ftz_f">,
420 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
421 def int_nvvm_sqrt_rm_f : GCCBuiltin<"__nvvm_sqrt_rm_f">,
422 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
423 def int_nvvm_sqrt_rp_ftz_f : GCCBuiltin<"__nvvm_sqrt_rp_ftz_f">,
424 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
425 def int_nvvm_sqrt_rp_f : GCCBuiltin<"__nvvm_sqrt_rp_f">,
426 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
427 def int_nvvm_sqrt_approx_ftz_f : GCCBuiltin<"__nvvm_sqrt_approx_ftz_f">,
428 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
429 def int_nvvm_sqrt_approx_f : GCCBuiltin<"__nvvm_sqrt_approx_f">,
430 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
432 def int_nvvm_sqrt_rn_d : GCCBuiltin<"__nvvm_sqrt_rn_d">,
433 Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>;
434 def int_nvvm_sqrt_rz_d : GCCBuiltin<"__nvvm_sqrt_rz_d">,
435 Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>;
436 def int_nvvm_sqrt_rm_d : GCCBuiltin<"__nvvm_sqrt_rm_d">,
437 Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>;
438 def int_nvvm_sqrt_rp_d : GCCBuiltin<"__nvvm_sqrt_rp_d">,
439 Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>;
445 def int_nvvm_rsqrt_approx_ftz_f : GCCBuiltin<"__nvvm_rsqrt_approx_ftz_f">,
446 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
447 def int_nvvm_rsqrt_approx_f : GCCBuiltin<"__nvvm_rsqrt_approx_f">,
448 Intrinsic<[llvm_float_ty], [llvm_float_ty], [IntrNoMem]>;
449 def int_nvvm_rsqrt_approx_d : GCCBuiltin<"__nvvm_rsqrt_approx_d">,
450 Intrinsic<[llvm_double_ty], [llvm_double_ty], [IntrNoMem]>;
456 def int_nvvm_add_rn_ftz_f : GCCBuiltin<"__nvvm_add_rn_ftz_f">,
457 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
458 [IntrNoMem, Commutative]>;
459 def int_nvvm_add_rn_f : GCCBuiltin<"__nvvm_add_rn_f">,
460 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
461 [IntrNoMem, Commutative]>;
462 def int_nvvm_add_rz_ftz_f : GCCBuiltin<"__nvvm_add_rz_ftz_f">,
463 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
464 [IntrNoMem, Commutative]>;
465 def int_nvvm_add_rz_f : GCCBuiltin<"__nvvm_add_rz_f">,
466 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
467 [IntrNoMem, Commutative]>;
468 def int_nvvm_add_rm_ftz_f : GCCBuiltin<"__nvvm_add_rm_ftz_f">,
469 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
470 [IntrNoMem, Commutative]>;
471 def int_nvvm_add_rm_f : GCCBuiltin<"__nvvm_add_rm_f">,
472 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
473 [IntrNoMem, Commutative]>;
474 def int_nvvm_add_rp_ftz_f : GCCBuiltin<"__nvvm_add_rp_ftz_f">,
475 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
476 [IntrNoMem, Commutative]>;
477 def int_nvvm_add_rp_f : GCCBuiltin<"__nvvm_add_rp_f">,
478 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_float_ty],
479 [IntrNoMem, Commutative]>;
481 def int_nvvm_add_rn_d : GCCBuiltin<"__nvvm_add_rn_d">,
482 Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
483 [IntrNoMem, Commutative]>;
484 def int_nvvm_add_rz_d : GCCBuiltin<"__nvvm_add_rz_d">,
485 Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
486 [IntrNoMem, Commutative]>;
487 def int_nvvm_add_rm_d : GCCBuiltin<"__nvvm_add_rm_d">,
488 Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
489 [IntrNoMem, Commutative]>;
490 def int_nvvm_add_rp_d : GCCBuiltin<"__nvvm_add_rp_d">,
491 Intrinsic<[llvm_double_ty], [llvm_double_ty, llvm_double_ty],
492 [IntrNoMem, Commutative]>;
498 def int_nvvm_d2f_rn_ftz : GCCBuiltin<"__nvvm_d2f_rn_ftz">,
499 Intrinsic<[llvm_float_ty], [llvm_double_ty], [IntrNoMem]>;
500 def int_nvvm_d2f_rn : GCCBuiltin<"__nvvm_d2f_rn">,
501 Intrinsic<[llvm_float_ty], [llvm_double_ty], [IntrNoMem]>;
502 def int_nvvm_d2f_rz_ftz : GCCBuiltin<"__nvvm_d2f_rz_ftz">,
503 Intrinsic<[llvm_float_ty], [llvm_double_ty], [IntrNoMem]>;
504 def int_nvvm_d2f_rz : GCCBuiltin<"__nvvm_d2f_rz">,
505 Intrinsic<[llvm_float_ty], [llvm_double_ty], [IntrNoMem]>;
506 def int_nvvm_d2f_rm_ftz : GCCBuiltin<"__nvvm_d2f_rm_ftz">,
507 Intrinsic<[llvm_float_ty], [llvm_double_ty], [IntrNoMem]>;
508 def int_nvvm_d2f_rm : GCCBuiltin<"__nvvm_d2f_rm">,
509 Intrinsic<[llvm_float_ty], [llvm_double_ty], [IntrNoMem]>;
510 def int_nvvm_d2f_rp_ftz : GCCBuiltin<"__nvvm_d2f_rp_ftz">,
511 Intrinsic<[llvm_float_ty], [llvm_double_ty], [IntrNoMem]>;
512 def int_nvvm_d2f_rp : GCCBuiltin<"__nvvm_d2f_rp">,
513 Intrinsic<[llvm_float_ty], [llvm_double_ty], [IntrNoMem]>;
515 def int_nvvm_d2i_rn : GCCBuiltin<"__nvvm_d2i_rn">,
516 Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>;
517 def int_nvvm_d2i_rz : GCCBuiltin<"__nvvm_d2i_rz">,
518 Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>;
519 def int_nvvm_d2i_rm : GCCBuiltin<"__nvvm_d2i_rm">,
520 Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>;
521 def int_nvvm_d2i_rp : GCCBuiltin<"__nvvm_d2i_rp">,
522 Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>;
524 def int_nvvm_d2ui_rn : GCCBuiltin<"__nvvm_d2ui_rn">,
525 Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>;
526 def int_nvvm_d2ui_rz : GCCBuiltin<"__nvvm_d2ui_rz">,
527 Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>;
528 def int_nvvm_d2ui_rm : GCCBuiltin<"__nvvm_d2ui_rm">,
529 Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>;
530 def int_nvvm_d2ui_rp : GCCBuiltin<"__nvvm_d2ui_rp">,
531 Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>;
533 def int_nvvm_i2d_rn : GCCBuiltin<"__nvvm_i2d_rn">,
534 Intrinsic<[llvm_double_ty], [llvm_i32_ty], [IntrNoMem]>;
535 def int_nvvm_i2d_rz : GCCBuiltin<"__nvvm_i2d_rz">,
536 Intrinsic<[llvm_double_ty], [llvm_i32_ty], [IntrNoMem]>;
537 def int_nvvm_i2d_rm : GCCBuiltin<"__nvvm_i2d_rm">,
538 Intrinsic<[llvm_double_ty], [llvm_i32_ty], [IntrNoMem]>;
539 def int_nvvm_i2d_rp : GCCBuiltin<"__nvvm_i2d_rp">,
540 Intrinsic<[llvm_double_ty], [llvm_i32_ty], [IntrNoMem]>;
542 def int_nvvm_ui2d_rn : GCCBuiltin<"__nvvm_ui2d_rn">,
543 Intrinsic<[llvm_double_ty], [llvm_i32_ty], [IntrNoMem]>;
544 def int_nvvm_ui2d_rz : GCCBuiltin<"__nvvm_ui2d_rz">,
545 Intrinsic<[llvm_double_ty], [llvm_i32_ty], [IntrNoMem]>;
546 def int_nvvm_ui2d_rm : GCCBuiltin<"__nvvm_ui2d_rm">,
547 Intrinsic<[llvm_double_ty], [llvm_i32_ty], [IntrNoMem]>;
548 def int_nvvm_ui2d_rp : GCCBuiltin<"__nvvm_ui2d_rp">,
549 Intrinsic<[llvm_double_ty], [llvm_i32_ty], [IntrNoMem]>;
551 def int_nvvm_f2i_rn_ftz : GCCBuiltin<"__nvvm_f2i_rn_ftz">,
552 Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>;
553 def int_nvvm_f2i_rn : GCCBuiltin<"__nvvm_f2i_rn">,
554 Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>;
555 def int_nvvm_f2i_rz_ftz : GCCBuiltin<"__nvvm_f2i_rz_ftz">,
556 Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>;
557 def int_nvvm_f2i_rz : GCCBuiltin<"__nvvm_f2i_rz">,
558 Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>;
559 def int_nvvm_f2i_rm_ftz : GCCBuiltin<"__nvvm_f2i_rm_ftz">,
560 Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>;
561 def int_nvvm_f2i_rm : GCCBuiltin<"__nvvm_f2i_rm">,
562 Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>;
563 def int_nvvm_f2i_rp_ftz : GCCBuiltin<"__nvvm_f2i_rp_ftz">,
564 Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>;
565 def int_nvvm_f2i_rp : GCCBuiltin<"__nvvm_f2i_rp">,
566 Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>;
568 def int_nvvm_f2ui_rn_ftz : GCCBuiltin<"__nvvm_f2ui_rn_ftz">,
569 Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>;
570 def int_nvvm_f2ui_rn : GCCBuiltin<"__nvvm_f2ui_rn">,
571 Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>;
572 def int_nvvm_f2ui_rz_ftz : GCCBuiltin<"__nvvm_f2ui_rz_ftz">,
573 Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>;
574 def int_nvvm_f2ui_rz : GCCBuiltin<"__nvvm_f2ui_rz">,
575 Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>;
576 def int_nvvm_f2ui_rm_ftz : GCCBuiltin<"__nvvm_f2ui_rm_ftz">,
577 Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>;
578 def int_nvvm_f2ui_rm : GCCBuiltin<"__nvvm_f2ui_rm">,
579 Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>;
580 def int_nvvm_f2ui_rp_ftz : GCCBuiltin<"__nvvm_f2ui_rp_ftz">,
581 Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>;
582 def int_nvvm_f2ui_rp : GCCBuiltin<"__nvvm_f2ui_rp">,
583 Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>;
585 def int_nvvm_i2f_rn : GCCBuiltin<"__nvvm_i2f_rn">,
586 Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>;
587 def int_nvvm_i2f_rz : GCCBuiltin<"__nvvm_i2f_rz">,
588 Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>;
589 def int_nvvm_i2f_rm : GCCBuiltin<"__nvvm_i2f_rm">,
590 Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>;
591 def int_nvvm_i2f_rp : GCCBuiltin<"__nvvm_i2f_rp">,
592 Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>;
594 def int_nvvm_ui2f_rn : GCCBuiltin<"__nvvm_ui2f_rn">,
595 Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>;
596 def int_nvvm_ui2f_rz : GCCBuiltin<"__nvvm_ui2f_rz">,
597 Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>;
598 def int_nvvm_ui2f_rm : GCCBuiltin<"__nvvm_ui2f_rm">,
599 Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>;
600 def int_nvvm_ui2f_rp : GCCBuiltin<"__nvvm_ui2f_rp">,
601 Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>;
603 def int_nvvm_lohi_i2d : GCCBuiltin<"__nvvm_lohi_i2d">,
604 Intrinsic<[llvm_double_ty], [llvm_i32_ty, llvm_i32_ty],
605 [IntrNoMem, Commutative]>;
607 def int_nvvm_d2i_lo : GCCBuiltin<"__nvvm_d2i_lo">,
608 Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>;
609 def int_nvvm_d2i_hi : GCCBuiltin<"__nvvm_d2i_hi">,
610 Intrinsic<[llvm_i32_ty], [llvm_double_ty], [IntrNoMem]>;
612 def int_nvvm_f2ll_rn_ftz : GCCBuiltin<"__nvvm_f2ll_rn_ftz">,
613 Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>;
614 def int_nvvm_f2ll_rn : GCCBuiltin<"__nvvm_f2ll_rn">,
615 Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>;
616 def int_nvvm_f2ll_rz_ftz : GCCBuiltin<"__nvvm_f2ll_rz_ftz">,
617 Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>;
618 def int_nvvm_f2ll_rz : GCCBuiltin<"__nvvm_f2ll_rz">,
619 Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>;
620 def int_nvvm_f2ll_rm_ftz : GCCBuiltin<"__nvvm_f2ll_rm_ftz">,
621 Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>;
622 def int_nvvm_f2ll_rm : GCCBuiltin<"__nvvm_f2ll_rm">,
623 Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>;
624 def int_nvvm_f2ll_rp_ftz : GCCBuiltin<"__nvvm_f2ll_rp_ftz">,
625 Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>;
626 def int_nvvm_f2ll_rp : GCCBuiltin<"__nvvm_f2ll_rp">,
627 Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>;
629 def int_nvvm_f2ull_rn_ftz : GCCBuiltin<"__nvvm_f2ull_rn_ftz">,
630 Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>;
631 def int_nvvm_f2ull_rn : GCCBuiltin<"__nvvm_f2ull_rn">,
632 Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>;
633 def int_nvvm_f2ull_rz_ftz : GCCBuiltin<"__nvvm_f2ull_rz_ftz">,
634 Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>;
635 def int_nvvm_f2ull_rz : GCCBuiltin<"__nvvm_f2ull_rz">,
636 Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>;
637 def int_nvvm_f2ull_rm_ftz : GCCBuiltin<"__nvvm_f2ull_rm_ftz">,
638 Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>;
639 def int_nvvm_f2ull_rm : GCCBuiltin<"__nvvm_f2ull_rm">,
640 Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>;
641 def int_nvvm_f2ull_rp_ftz : GCCBuiltin<"__nvvm_f2ull_rp_ftz">,
642 Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>;
643 def int_nvvm_f2ull_rp : GCCBuiltin<"__nvvm_f2ull_rp">,
644 Intrinsic<[llvm_i64_ty], [llvm_float_ty], [IntrNoMem]>;
646 def int_nvvm_d2ll_rn : GCCBuiltin<"__nvvm_d2ll_rn">,
647 Intrinsic<[llvm_i64_ty], [llvm_double_ty], [IntrNoMem]>;
648 def int_nvvm_d2ll_rz : GCCBuiltin<"__nvvm_d2ll_rz">,
649 Intrinsic<[llvm_i64_ty], [llvm_double_ty], [IntrNoMem]>;
650 def int_nvvm_d2ll_rm : GCCBuiltin<"__nvvm_d2ll_rm">,
651 Intrinsic<[llvm_i64_ty], [llvm_double_ty], [IntrNoMem]>;
652 def int_nvvm_d2ll_rp : GCCBuiltin<"__nvvm_d2ll_rp">,
653 Intrinsic<[llvm_i64_ty], [llvm_double_ty], [IntrNoMem]>;
655 def int_nvvm_d2ull_rn : GCCBuiltin<"__nvvm_d2ull_rn">,
656 Intrinsic<[llvm_i64_ty], [llvm_double_ty], [IntrNoMem]>;
657 def int_nvvm_d2ull_rz : GCCBuiltin<"__nvvm_d2ull_rz">,
658 Intrinsic<[llvm_i64_ty], [llvm_double_ty], [IntrNoMem]>;
659 def int_nvvm_d2ull_rm : GCCBuiltin<"__nvvm_d2ull_rm">,
660 Intrinsic<[llvm_i64_ty], [llvm_double_ty], [IntrNoMem]>;
661 def int_nvvm_d2ull_rp : GCCBuiltin<"__nvvm_d2ull_rp">,
662 Intrinsic<[llvm_i64_ty], [llvm_double_ty], [IntrNoMem]>;
664 def int_nvvm_ll2f_rn : GCCBuiltin<"__nvvm_ll2f_rn">,
665 Intrinsic<[llvm_float_ty], [llvm_i64_ty], [IntrNoMem]>;
666 def int_nvvm_ll2f_rz : GCCBuiltin<"__nvvm_ll2f_rz">,
667 Intrinsic<[llvm_float_ty], [llvm_i64_ty], [IntrNoMem]>;
668 def int_nvvm_ll2f_rm : GCCBuiltin<"__nvvm_ll2f_rm">,
669 Intrinsic<[llvm_float_ty], [llvm_i64_ty], [IntrNoMem]>;
670 def int_nvvm_ll2f_rp : GCCBuiltin<"__nvvm_ll2f_rp">,
671 Intrinsic<[llvm_float_ty], [llvm_i64_ty], [IntrNoMem]>;
672 def int_nvvm_ull2f_rn : GCCBuiltin<"__nvvm_ull2f_rn">,
673 Intrinsic<[llvm_float_ty], [llvm_i64_ty], [IntrNoMem]>;
674 def int_nvvm_ull2f_rz : GCCBuiltin<"__nvvm_ull2f_rz">,
675 Intrinsic<[llvm_float_ty], [llvm_i64_ty], [IntrNoMem]>;
676 def int_nvvm_ull2f_rm : GCCBuiltin<"__nvvm_ull2f_rm">,
677 Intrinsic<[llvm_float_ty], [llvm_i64_ty], [IntrNoMem]>;
678 def int_nvvm_ull2f_rp : GCCBuiltin<"__nvvm_ull2f_rp">,
679 Intrinsic<[llvm_float_ty], [llvm_i64_ty], [IntrNoMem]>;
681 def int_nvvm_ll2d_rn : GCCBuiltin<"__nvvm_ll2d_rn">,
682 Intrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem]>;
683 def int_nvvm_ll2d_rz : GCCBuiltin<"__nvvm_ll2d_rz">,
684 Intrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem]>;
685 def int_nvvm_ll2d_rm : GCCBuiltin<"__nvvm_ll2d_rm">,
686 Intrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem]>;
687 def int_nvvm_ll2d_rp : GCCBuiltin<"__nvvm_ll2d_rp">,
688 Intrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem]>;
689 def int_nvvm_ull2d_rn : GCCBuiltin<"__nvvm_ull2d_rn">,
690 Intrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem]>;
691 def int_nvvm_ull2d_rz : GCCBuiltin<"__nvvm_ull2d_rz">,
692 Intrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem]>;
693 def int_nvvm_ull2d_rm : GCCBuiltin<"__nvvm_ull2d_rm">,
694 Intrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem]>;
695 def int_nvvm_ull2d_rp : GCCBuiltin<"__nvvm_ull2d_rp">,
696 Intrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem]>;
698 def int_nvvm_f2h_rn_ftz : GCCBuiltin<"__nvvm_f2h_rn_ftz">,
699 Intrinsic<[llvm_i16_ty], [llvm_float_ty], [IntrNoMem]>;
700 def int_nvvm_f2h_rn : GCCBuiltin<"__nvvm_f2h_rn">,
701 Intrinsic<[llvm_i16_ty], [llvm_float_ty], [IntrNoMem]>;
703 def int_nvvm_h2f : GCCBuiltin<"__nvvm_h2f">,
704 Intrinsic<[llvm_float_ty], [llvm_i16_ty], [IntrNoMem]>;
710 def int_nvvm_bitcast_f2i : GCCBuiltin<"__nvvm_bitcast_f2i">,
711 Intrinsic<[llvm_i32_ty], [llvm_float_ty], [IntrNoMem]>;
712 def int_nvvm_bitcast_i2f : GCCBuiltin<"__nvvm_bitcast_i2f">,
713 Intrinsic<[llvm_float_ty], [llvm_i32_ty], [IntrNoMem]>;
715 def int_nvvm_bitcast_ll2d : GCCBuiltin<"__nvvm_bitcast_ll2d">,
716 Intrinsic<[llvm_double_ty], [llvm_i64_ty], [IntrNoMem]>;
717 def int_nvvm_bitcast_d2ll : GCCBuiltin<"__nvvm_bitcast_d2ll">,
718 Intrinsic<[llvm_i64_ty], [llvm_double_ty], [IntrNoMem]>;
721 // Atomic not available as an llvm intrinsic.
722 def int_nvvm_atomic_load_add_f32 : Intrinsic<[llvm_float_ty],
723 [LLVMAnyPointerType<llvm_float_ty>, llvm_float_ty],
724 [IntrArgMemOnly, NoCapture<0>]>;
725 def int_nvvm_atomic_load_inc_32 : Intrinsic<[llvm_i32_ty],
726 [LLVMAnyPointerType<llvm_i32_ty>, llvm_i32_ty],
727 [IntrArgMemOnly, NoCapture<0>]>;
728 def int_nvvm_atomic_load_dec_32 : Intrinsic<[llvm_i32_ty],
729 [LLVMAnyPointerType<llvm_i32_ty>, llvm_i32_ty],
730 [IntrArgMemOnly, NoCapture<0>]>;
732 class SCOPED_ATOMIC2_impl<LLVMType elty>
734 [LLVMAnyPointerType<LLVMMatchType<0>>, LLVMMatchType<0>],
735 [IntrArgMemOnly, NoCapture<0>]>;
736 class SCOPED_ATOMIC3_impl<LLVMType elty>
738 [LLVMAnyPointerType<LLVMMatchType<0>>, LLVMMatchType<0>,
740 [IntrArgMemOnly, NoCapture<0>]>;
742 multiclass PTXAtomicWithScope2<LLVMType elty> {
743 def _cta : SCOPED_ATOMIC2_impl<elty>;
744 def _sys : SCOPED_ATOMIC2_impl<elty>;
746 multiclass PTXAtomicWithScope3<LLVMType elty> {
747 def _cta : SCOPED_ATOMIC3_impl<elty>;
748 def _sys : SCOPED_ATOMIC3_impl<elty>;
750 multiclass PTXAtomicWithScope2_fi {
751 defm _f: PTXAtomicWithScope2<llvm_anyfloat_ty>;
752 defm _i: PTXAtomicWithScope2<llvm_anyint_ty>;
754 defm int_nvvm_atomic_add_gen : PTXAtomicWithScope2_fi;
755 defm int_nvvm_atomic_inc_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>;
756 defm int_nvvm_atomic_dec_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>;
757 defm int_nvvm_atomic_exch_gen_i: PTXAtomicWithScope2<llvm_anyint_ty>;
758 defm int_nvvm_atomic_xor_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>;
759 defm int_nvvm_atomic_max_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>;
760 defm int_nvvm_atomic_min_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>;
761 defm int_nvvm_atomic_or_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>;
762 defm int_nvvm_atomic_and_gen_i : PTXAtomicWithScope2<llvm_anyint_ty>;
763 defm int_nvvm_atomic_cas_gen_i : PTXAtomicWithScope3<llvm_anyint_ty>;
767 // The builtin for "bar.sync 0" is called __syncthreads. Unlike most of the
768 // intrinsics in this file, this one is a user-facing API.
769 def int_nvvm_barrier0 : GCCBuiltin<"__syncthreads">,
770 Intrinsic<[], [], [IntrConvergent]>;
771 def int_nvvm_barrier0_popc : GCCBuiltin<"__nvvm_bar0_popc">,
772 Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrConvergent]>;
773 def int_nvvm_barrier0_and : GCCBuiltin<"__nvvm_bar0_and">,
774 Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrConvergent]>;
775 def int_nvvm_barrier0_or : GCCBuiltin<"__nvvm_bar0_or">,
776 Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrConvergent]>;
778 def int_nvvm_bar_sync :
779 Intrinsic<[], [llvm_i32_ty], [IntrConvergent]>,
780 GCCBuiltin<"__nvvm_bar_sync">;
783 def int_nvvm_membar_cta : GCCBuiltin<"__nvvm_membar_cta">,
784 Intrinsic<[], [], []>;
785 def int_nvvm_membar_gl : GCCBuiltin<"__nvvm_membar_gl">,
786 Intrinsic<[], [], []>;
787 def int_nvvm_membar_sys : GCCBuiltin<"__nvvm_membar_sys">,
788 Intrinsic<[], [], []>;
790 // Generated within nvvm. Use for ldu on sm_20 or later. Second arg is the
791 // pointer's alignment.
792 def int_nvvm_ldu_global_i : Intrinsic<[llvm_anyint_ty],
793 [LLVMAnyPointerType<LLVMMatchType<0>>, llvm_i32_ty],
794 [IntrReadMem, IntrArgMemOnly, NoCapture<0>],
795 "llvm.nvvm.ldu.global.i">;
796 def int_nvvm_ldu_global_f : Intrinsic<[llvm_anyfloat_ty],
797 [LLVMAnyPointerType<LLVMMatchType<0>>, llvm_i32_ty],
798 [IntrReadMem, IntrArgMemOnly, NoCapture<0>],
799 "llvm.nvvm.ldu.global.f">;
800 def int_nvvm_ldu_global_p : Intrinsic<[llvm_anyptr_ty],
801 [LLVMAnyPointerType<LLVMMatchType<0>>, llvm_i32_ty],
802 [IntrReadMem, IntrArgMemOnly, NoCapture<0>],
803 "llvm.nvvm.ldu.global.p">;
805 // Generated within nvvm. Use for ldg on sm_35 or later. Second arg is the
806 // pointer's alignment.
807 def int_nvvm_ldg_global_i : Intrinsic<[llvm_anyint_ty],
808 [LLVMAnyPointerType<LLVMMatchType<0>>, llvm_i32_ty],
809 [IntrReadMem, IntrArgMemOnly, NoCapture<0>],
810 "llvm.nvvm.ldg.global.i">;
811 def int_nvvm_ldg_global_f : Intrinsic<[llvm_anyfloat_ty],
812 [LLVMAnyPointerType<LLVMMatchType<0>>, llvm_i32_ty],
813 [IntrReadMem, IntrArgMemOnly, NoCapture<0>],
814 "llvm.nvvm.ldg.global.f">;
815 def int_nvvm_ldg_global_p : Intrinsic<[llvm_anyptr_ty],
816 [LLVMAnyPointerType<LLVMMatchType<0>>, llvm_i32_ty],
817 [IntrReadMem, IntrArgMemOnly, NoCapture<0>],
818 "llvm.nvvm.ldg.global.p">;
820 // Use for generic pointers
821 // - These intrinsics are used to convert address spaces.
822 // - The input pointer and output pointer must have the same type, except for
823 // the address-space. (This restriction is not enforced here as there is
824 // currently no way to describe it).
825 // - This complements the llvm bitcast, which can be used to cast one type
826 // of pointer to another type of pointer, while the address space remains
828 def int_nvvm_ptr_local_to_gen: Intrinsic<[llvm_anyptr_ty],
829 [llvm_anyptr_ty], [IntrNoMem],
830 "llvm.nvvm.ptr.local.to.gen">;
831 def int_nvvm_ptr_shared_to_gen: Intrinsic<[llvm_anyptr_ty],
832 [llvm_anyptr_ty], [IntrNoMem],
833 "llvm.nvvm.ptr.shared.to.gen">;
834 def int_nvvm_ptr_global_to_gen: Intrinsic<[llvm_anyptr_ty],
835 [llvm_anyptr_ty], [IntrNoMem],
836 "llvm.nvvm.ptr.global.to.gen">;
837 def int_nvvm_ptr_constant_to_gen: Intrinsic<[llvm_anyptr_ty],
838 [llvm_anyptr_ty], [IntrNoMem],
839 "llvm.nvvm.ptr.constant.to.gen">;
841 def int_nvvm_ptr_gen_to_global: Intrinsic<[llvm_anyptr_ty],
842 [llvm_anyptr_ty], [IntrNoMem],
843 "llvm.nvvm.ptr.gen.to.global">;
844 def int_nvvm_ptr_gen_to_shared: Intrinsic<[llvm_anyptr_ty],
845 [llvm_anyptr_ty], [IntrNoMem],
846 "llvm.nvvm.ptr.gen.to.shared">;
847 def int_nvvm_ptr_gen_to_local: Intrinsic<[llvm_anyptr_ty],
848 [llvm_anyptr_ty], [IntrNoMem],
849 "llvm.nvvm.ptr.gen.to.local">;
850 def int_nvvm_ptr_gen_to_constant: Intrinsic<[llvm_anyptr_ty],
851 [llvm_anyptr_ty], [IntrNoMem],
852 "llvm.nvvm.ptr.gen.to.constant">;
854 // Used in nvvm internally to help address space opt and ptx code generation
855 // This is for params that are passed to kernel functions by pointer by-val.
856 def int_nvvm_ptr_gen_to_param: Intrinsic<[llvm_anyptr_ty],
859 "llvm.nvvm.ptr.gen.to.param">;
861 // Move intrinsics, used in nvvm internally
863 def int_nvvm_move_i16 : Intrinsic<[llvm_i16_ty], [llvm_i16_ty], [IntrNoMem],
864 "llvm.nvvm.move.i16">;
865 def int_nvvm_move_i32 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem],
866 "llvm.nvvm.move.i32">;
867 def int_nvvm_move_i64 : Intrinsic<[llvm_i64_ty], [llvm_i64_ty], [IntrNoMem],
868 "llvm.nvvm.move.i64">;
869 def int_nvvm_move_float : Intrinsic<[llvm_float_ty], [llvm_float_ty],
870 [IntrNoMem], "llvm.nvvm.move.float">;
871 def int_nvvm_move_double : Intrinsic<[llvm_double_ty], [llvm_double_ty],
872 [IntrNoMem], "llvm.nvvm.move.double">;
873 def int_nvvm_move_ptr : Intrinsic<[llvm_anyptr_ty], [llvm_anyptr_ty],
874 [IntrNoMem, NoCapture<0>], "llvm.nvvm.move.ptr">;
877 // For getting the handle from a texture or surface variable
878 def int_nvvm_texsurf_handle
879 : Intrinsic<[llvm_i64_ty], [llvm_metadata_ty, llvm_anyi64ptr_ty],
880 [IntrNoMem], "llvm.nvvm.texsurf.handle">;
881 def int_nvvm_texsurf_handle_internal
882 : Intrinsic<[llvm_i64_ty], [llvm_anyptr_ty],
883 [IntrNoMem], "llvm.nvvm.texsurf.handle.internal">;
886 def int_nvvm_compiler_error :
887 Intrinsic<[], [llvm_anyptr_ty], [], "llvm.nvvm.compiler.error">;
888 def int_nvvm_compiler_warn :
889 Intrinsic<[], [llvm_anyptr_ty], [], "llvm.nvvm.compiler.warn">;
891 def int_nvvm_reflect :
892 Intrinsic<[llvm_i32_ty], [llvm_anyptr_ty], [IntrNoMem], "llvm.nvvm.reflect">;
894 // isspacep.{const, global, local, shared}
895 def int_nvvm_isspacep_const
896 : Intrinsic<[llvm_i1_ty], [llvm_ptr_ty], [IntrNoMem],
897 "llvm.nvvm.isspacep.const">,
898 GCCBuiltin<"__nvvm_isspacep_const">;
899 def int_nvvm_isspacep_global
900 : Intrinsic<[llvm_i1_ty], [llvm_ptr_ty], [IntrNoMem],
901 "llvm.nvvm.isspacep.global">,
902 GCCBuiltin<"__nvvm_isspacep_global">;
903 def int_nvvm_isspacep_local
904 : Intrinsic<[llvm_i1_ty], [llvm_ptr_ty], [IntrNoMem],
905 "llvm.nvvm.isspacep.local">,
906 GCCBuiltin<"__nvvm_isspacep_local">;
907 def int_nvvm_isspacep_shared
908 : Intrinsic<[llvm_i1_ty], [llvm_ptr_ty], [IntrNoMem],
909 "llvm.nvvm.isspacep.shared">,
910 GCCBuiltin<"__nvvm_isspacep_shared">;
912 // Environment register read
913 def int_nvvm_read_ptx_sreg_envreg0
914 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
915 "llvm.nvvm.read.ptx.sreg.envreg0">,
916 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg0">;
917 def int_nvvm_read_ptx_sreg_envreg1
918 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
919 "llvm.nvvm.read.ptx.sreg.envreg1">,
920 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg1">;
921 def int_nvvm_read_ptx_sreg_envreg2
922 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
923 "llvm.nvvm.read.ptx.sreg.envreg2">,
924 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg2">;
925 def int_nvvm_read_ptx_sreg_envreg3
926 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
927 "llvm.nvvm.read.ptx.sreg.envreg3">,
928 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg3">;
929 def int_nvvm_read_ptx_sreg_envreg4
930 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
931 "llvm.nvvm.read.ptx.sreg.envreg4">,
932 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg4">;
933 def int_nvvm_read_ptx_sreg_envreg5
934 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
935 "llvm.nvvm.read.ptx.sreg.envreg5">,
936 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg5">;
937 def int_nvvm_read_ptx_sreg_envreg6
938 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
939 "llvm.nvvm.read.ptx.sreg.envreg6">,
940 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg6">;
941 def int_nvvm_read_ptx_sreg_envreg7
942 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
943 "llvm.nvvm.read.ptx.sreg.envreg7">,
944 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg7">;
945 def int_nvvm_read_ptx_sreg_envreg8
946 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
947 "llvm.nvvm.read.ptx.sreg.envreg8">,
948 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg8">;
949 def int_nvvm_read_ptx_sreg_envreg9
950 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
951 "llvm.nvvm.read.ptx.sreg.envreg9">,
952 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg9">;
953 def int_nvvm_read_ptx_sreg_envreg10
954 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
955 "llvm.nvvm.read.ptx.sreg.envreg10">,
956 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg10">;
957 def int_nvvm_read_ptx_sreg_envreg11
958 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
959 "llvm.nvvm.read.ptx.sreg.envreg11">,
960 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg11">;
961 def int_nvvm_read_ptx_sreg_envreg12
962 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
963 "llvm.nvvm.read.ptx.sreg.envreg12">,
964 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg12">;
965 def int_nvvm_read_ptx_sreg_envreg13
966 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
967 "llvm.nvvm.read.ptx.sreg.envreg13">,
968 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg13">;
969 def int_nvvm_read_ptx_sreg_envreg14
970 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
971 "llvm.nvvm.read.ptx.sreg.envreg14">,
972 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg14">;
973 def int_nvvm_read_ptx_sreg_envreg15
974 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
975 "llvm.nvvm.read.ptx.sreg.envreg15">,
976 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg15">;
977 def int_nvvm_read_ptx_sreg_envreg16
978 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
979 "llvm.nvvm.read.ptx.sreg.envreg16">,
980 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg16">;
981 def int_nvvm_read_ptx_sreg_envreg17
982 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
983 "llvm.nvvm.read.ptx.sreg.envreg17">,
984 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg17">;
985 def int_nvvm_read_ptx_sreg_envreg18
986 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
987 "llvm.nvvm.read.ptx.sreg.envreg18">,
988 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg18">;
989 def int_nvvm_read_ptx_sreg_envreg19
990 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
991 "llvm.nvvm.read.ptx.sreg.envreg19">,
992 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg19">;
993 def int_nvvm_read_ptx_sreg_envreg20
994 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
995 "llvm.nvvm.read.ptx.sreg.envreg20">,
996 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg20">;
997 def int_nvvm_read_ptx_sreg_envreg21
998 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
999 "llvm.nvvm.read.ptx.sreg.envreg21">,
1000 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg21">;
1001 def int_nvvm_read_ptx_sreg_envreg22
1002 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
1003 "llvm.nvvm.read.ptx.sreg.envreg22">,
1004 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg22">;
1005 def int_nvvm_read_ptx_sreg_envreg23
1006 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
1007 "llvm.nvvm.read.ptx.sreg.envreg23">,
1008 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg23">;
1009 def int_nvvm_read_ptx_sreg_envreg24
1010 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
1011 "llvm.nvvm.read.ptx.sreg.envreg24">,
1012 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg24">;
1013 def int_nvvm_read_ptx_sreg_envreg25
1014 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
1015 "llvm.nvvm.read.ptx.sreg.envreg25">,
1016 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg25">;
1017 def int_nvvm_read_ptx_sreg_envreg26
1018 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
1019 "llvm.nvvm.read.ptx.sreg.envreg26">,
1020 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg26">;
1021 def int_nvvm_read_ptx_sreg_envreg27
1022 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
1023 "llvm.nvvm.read.ptx.sreg.envreg27">,
1024 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg27">;
1025 def int_nvvm_read_ptx_sreg_envreg28
1026 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
1027 "llvm.nvvm.read.ptx.sreg.envreg28">,
1028 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg28">;
1029 def int_nvvm_read_ptx_sreg_envreg29
1030 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
1031 "llvm.nvvm.read.ptx.sreg.envreg29">,
1032 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg29">;
1033 def int_nvvm_read_ptx_sreg_envreg30
1034 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
1035 "llvm.nvvm.read.ptx.sreg.envreg30">,
1036 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg30">;
1037 def int_nvvm_read_ptx_sreg_envreg31
1038 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem],
1039 "llvm.nvvm.read.ptx.sreg.envreg31">,
1040 GCCBuiltin<"__nvvm_read_ptx_sreg_envreg31">;
1044 // texmode_independent
1045 def int_nvvm_tex_1d_v4f32_s32
1046 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1047 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty], [],
1048 "llvm.nvvm.tex.1d.v4f32.s32">;
1049 def int_nvvm_tex_1d_v4f32_f32
1050 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1051 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty], [],
1052 "llvm.nvvm.tex.1d.v4f32.f32">;
1053 def int_nvvm_tex_1d_level_v4f32_f32
1054 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1055 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1056 "llvm.nvvm.tex.1d.level.v4f32.f32">;
1057 def int_nvvm_tex_1d_grad_v4f32_f32
1058 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1059 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1061 "llvm.nvvm.tex.1d.grad.v4f32.f32">;
1062 def int_nvvm_tex_1d_v4s32_s32
1063 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1064 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty], [],
1065 "llvm.nvvm.tex.1d.v4s32.s32">;
1066 def int_nvvm_tex_1d_v4s32_f32
1067 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1068 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty], [],
1069 "llvm.nvvm.tex.1d.v4s32.f32">;
1070 def int_nvvm_tex_1d_level_v4s32_f32
1071 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1072 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1073 "llvm.nvvm.tex.1d.level.v4s32.f32">;
1074 def int_nvvm_tex_1d_grad_v4s32_f32
1075 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1076 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1078 "llvm.nvvm.tex.1d.grad.v4s32.f32">;
1079 def int_nvvm_tex_1d_v4u32_s32
1080 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1081 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty], [],
1082 "llvm.nvvm.tex.1d.v4u32.s32">;
1083 def int_nvvm_tex_1d_v4u32_f32
1084 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1085 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty], [],
1086 "llvm.nvvm.tex.1d.v4u32.f32">;
1087 def int_nvvm_tex_1d_level_v4u32_f32
1088 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1089 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1090 "llvm.nvvm.tex.1d.level.v4u32.f32">;
1091 def int_nvvm_tex_1d_grad_v4u32_f32
1092 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1093 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1095 "llvm.nvvm.tex.1d.grad.v4u32.f32">;
1097 def int_nvvm_tex_1d_array_v4f32_s32
1098 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1099 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1100 "llvm.nvvm.tex.1d.array.v4f32.s32">;
1101 def int_nvvm_tex_1d_array_v4f32_f32
1102 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1103 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty], [],
1104 "llvm.nvvm.tex.1d.array.v4f32.f32">;
1105 def int_nvvm_tex_1d_array_level_v4f32_f32
1106 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1107 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1109 "llvm.nvvm.tex.1d.array.level.v4f32.f32">;
1110 def int_nvvm_tex_1d_array_grad_v4f32_f32
1111 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1112 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1113 llvm_float_ty, llvm_float_ty], [],
1114 "llvm.nvvm.tex.1d.array.grad.v4f32.f32">;
1115 def int_nvvm_tex_1d_array_v4s32_s32
1116 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1117 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1118 "llvm.nvvm.tex.1d.array.v4s32.s32">;
1119 def int_nvvm_tex_1d_array_v4s32_f32
1120 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1121 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty], [],
1122 "llvm.nvvm.tex.1d.array.v4s32.f32">;
1123 def int_nvvm_tex_1d_array_level_v4s32_f32
1124 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1125 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1127 "llvm.nvvm.tex.1d.array.level.v4s32.f32">;
1128 def int_nvvm_tex_1d_array_grad_v4s32_f32
1129 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1130 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1131 llvm_float_ty, llvm_float_ty], [],
1132 "llvm.nvvm.tex.1d.array.grad.v4s32.f32">;
1133 def int_nvvm_tex_1d_array_v4u32_s32
1134 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1135 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1136 "llvm.nvvm.tex.1d.array.v4u32.s32">;
1137 def int_nvvm_tex_1d_array_v4u32_f32
1138 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1139 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty], [],
1140 "llvm.nvvm.tex.1d.array.v4u32.f32">;
1141 def int_nvvm_tex_1d_array_level_v4u32_f32
1142 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1143 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1145 "llvm.nvvm.tex.1d.array.level.v4u32.f32">;
1146 def int_nvvm_tex_1d_array_grad_v4u32_f32
1147 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1148 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1149 llvm_float_ty, llvm_float_ty], [],
1150 "llvm.nvvm.tex.1d.array.grad.v4u32.f32">;
1152 def int_nvvm_tex_2d_v4f32_s32
1153 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1154 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1155 "llvm.nvvm.tex.2d.v4f32.s32">;
1156 def int_nvvm_tex_2d_v4f32_f32
1157 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1158 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1159 "llvm.nvvm.tex.2d.v4f32.f32">;
1160 def int_nvvm_tex_2d_level_v4f32_f32
1161 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1162 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1164 "llvm.nvvm.tex.2d.level.v4f32.f32">;
1165 def int_nvvm_tex_2d_grad_v4f32_f32
1166 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1167 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1168 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1169 "llvm.nvvm.tex.2d.grad.v4f32.f32">;
1170 def int_nvvm_tex_2d_v4s32_s32
1171 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1172 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1173 "llvm.nvvm.tex.2d.v4s32.s32">;
1174 def int_nvvm_tex_2d_v4s32_f32
1175 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1176 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1177 "llvm.nvvm.tex.2d.v4s32.f32">;
1178 def int_nvvm_tex_2d_level_v4s32_f32
1179 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1180 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1182 "llvm.nvvm.tex.2d.level.v4s32.f32">;
1183 def int_nvvm_tex_2d_grad_v4s32_f32
1184 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1185 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1186 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1187 "llvm.nvvm.tex.2d.grad.v4s32.f32">;
1188 def int_nvvm_tex_2d_v4u32_s32
1189 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1190 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1191 "llvm.nvvm.tex.2d.v4u32.s32">;
1192 def int_nvvm_tex_2d_v4u32_f32
1193 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1194 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1195 "llvm.nvvm.tex.2d.v4u32.f32">;
1196 def int_nvvm_tex_2d_level_v4u32_f32
1197 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1198 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1200 "llvm.nvvm.tex.2d.level.v4u32.f32">;
1201 def int_nvvm_tex_2d_grad_v4u32_f32
1202 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1203 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1204 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1205 "llvm.nvvm.tex.2d.grad.v4u32.f32">;
1207 def int_nvvm_tex_2d_array_v4f32_s32
1208 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1209 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
1211 "llvm.nvvm.tex.2d.array.v4f32.s32">;
1212 def int_nvvm_tex_2d_array_v4f32_f32
1213 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1214 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1216 "llvm.nvvm.tex.2d.array.v4f32.f32">;
1217 def int_nvvm_tex_2d_array_level_v4f32_f32
1218 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1219 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1220 llvm_float_ty, llvm_float_ty], [],
1221 "llvm.nvvm.tex.2d.array.level.v4f32.f32">;
1222 def int_nvvm_tex_2d_array_grad_v4f32_f32
1223 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1224 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1225 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty,
1227 "llvm.nvvm.tex.2d.array.grad.v4f32.f32">;
1228 def int_nvvm_tex_2d_array_v4s32_s32
1229 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1230 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
1232 "llvm.nvvm.tex.2d.array.v4s32.s32">;
1233 def int_nvvm_tex_2d_array_v4s32_f32
1234 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1235 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1237 "llvm.nvvm.tex.2d.array.v4s32.f32">;
1238 def int_nvvm_tex_2d_array_level_v4s32_f32
1239 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1240 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1241 llvm_float_ty, llvm_float_ty], [],
1242 "llvm.nvvm.tex.2d.array.level.v4s32.f32">;
1243 def int_nvvm_tex_2d_array_grad_v4s32_f32
1244 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1245 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1246 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty,
1248 "llvm.nvvm.tex.2d.array.grad.v4s32.f32">;
1249 def int_nvvm_tex_2d_array_v4u32_s32
1250 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1251 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
1253 "llvm.nvvm.tex.2d.array.v4u32.s32">;
1254 def int_nvvm_tex_2d_array_v4u32_f32
1255 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1256 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1258 "llvm.nvvm.tex.2d.array.v4u32.f32">;
1259 def int_nvvm_tex_2d_array_level_v4u32_f32
1260 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1261 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1262 llvm_float_ty, llvm_float_ty], [],
1263 "llvm.nvvm.tex.2d.array.level.v4u32.f32">;
1264 def int_nvvm_tex_2d_array_grad_v4u32_f32
1265 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1266 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1267 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty,
1269 "llvm.nvvm.tex.2d.array.grad.v4u32.f32">;
1271 def int_nvvm_tex_3d_v4f32_s32
1272 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1273 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1274 [], "llvm.nvvm.tex.3d.v4f32.s32">;
1275 def int_nvvm_tex_3d_v4f32_f32
1276 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1277 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1279 "llvm.nvvm.tex.3d.v4f32.f32">;
1280 def int_nvvm_tex_3d_level_v4f32_f32
1281 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1282 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1283 llvm_float_ty, llvm_float_ty], [],
1284 "llvm.nvvm.tex.3d.level.v4f32.f32">;
1285 def int_nvvm_tex_3d_grad_v4f32_f32
1286 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1287 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1288 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty,
1289 llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1290 "llvm.nvvm.tex.3d.grad.v4f32.f32">;
1291 def int_nvvm_tex_3d_v4s32_s32
1292 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1293 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1294 [], "llvm.nvvm.tex.3d.v4s32.s32">;
1295 def int_nvvm_tex_3d_v4s32_f32
1296 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1297 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1299 "llvm.nvvm.tex.3d.v4s32.f32">;
1300 def int_nvvm_tex_3d_level_v4s32_f32
1301 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1302 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1303 llvm_float_ty, llvm_float_ty], [],
1304 "llvm.nvvm.tex.3d.level.v4s32.f32">;
1305 def int_nvvm_tex_3d_grad_v4s32_f32
1306 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1307 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1308 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty,
1309 llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1310 "llvm.nvvm.tex.3d.grad.v4s32.f32">;
1311 def int_nvvm_tex_3d_v4u32_s32
1312 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1313 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1314 [], "llvm.nvvm.tex.3d.v4u32.s32">;
1315 def int_nvvm_tex_3d_v4u32_f32
1316 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1317 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1319 "llvm.nvvm.tex.3d.v4u32.f32">;
1320 def int_nvvm_tex_3d_level_v4u32_f32
1321 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1322 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1323 llvm_float_ty, llvm_float_ty], [],
1324 "llvm.nvvm.tex.3d.level.v4u32.f32">;
1325 def int_nvvm_tex_3d_grad_v4u32_f32
1326 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1327 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1328 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty,
1329 llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1330 "llvm.nvvm.tex.3d.grad.v4u32.f32">;
1332 def int_nvvm_tex_cube_v4f32_f32
1333 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1334 [llvm_i64_ty, llvm_i64_ty,
1335 llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1336 "llvm.nvvm.tex.cube.v4f32.f32">;
1337 def int_nvvm_tex_cube_level_v4f32_f32
1338 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1339 [llvm_i64_ty, llvm_i64_ty,
1340 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1341 "llvm.nvvm.tex.cube.level.v4f32.f32">;
1342 def int_nvvm_tex_cube_v4s32_f32
1343 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1344 [llvm_i64_ty, llvm_i64_ty,
1345 llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1346 "llvm.nvvm.tex.cube.v4s32.f32">;
1347 def int_nvvm_tex_cube_level_v4s32_f32
1348 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1349 [llvm_i64_ty, llvm_i64_ty,
1350 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1351 "llvm.nvvm.tex.cube.level.v4s32.f32">;
1352 def int_nvvm_tex_cube_v4u32_f32
1353 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1354 [llvm_i64_ty, llvm_i64_ty,
1355 llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1356 "llvm.nvvm.tex.cube.v4u32.f32">;
1357 def int_nvvm_tex_cube_level_v4u32_f32
1358 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1359 [llvm_i64_ty, llvm_i64_ty,
1360 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1361 "llvm.nvvm.tex.cube.level.v4u32.f32">;
1363 def int_nvvm_tex_cube_array_v4f32_f32
1364 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1365 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty,
1366 llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1367 "llvm.nvvm.tex.cube.array.v4f32.f32">;
1368 def int_nvvm_tex_cube_array_level_v4f32_f32
1369 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1370 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty,
1371 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1372 "llvm.nvvm.tex.cube.array.level.v4f32.f32">;
1373 def int_nvvm_tex_cube_array_v4s32_f32
1374 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1375 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty,
1376 llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1377 "llvm.nvvm.tex.cube.array.v4s32.f32">;
1378 def int_nvvm_tex_cube_array_level_v4s32_f32
1379 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1380 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty,
1381 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1382 "llvm.nvvm.tex.cube.array.level.v4s32.f32">;
1383 def int_nvvm_tex_cube_array_v4u32_f32
1384 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1385 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty,
1386 llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1387 "llvm.nvvm.tex.cube.array.v4u32.f32">;
1388 def int_nvvm_tex_cube_array_level_v4u32_f32
1389 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1390 [llvm_i64_ty, llvm_i64_ty, llvm_i32_ty,
1391 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1392 "llvm.nvvm.tex.cube.array.level.v4u32.f32">;
1394 def int_nvvm_tld4_r_2d_v4f32_f32
1395 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1396 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1397 "llvm.nvvm.tld4.r.2d.v4f32.f32">;
1398 def int_nvvm_tld4_g_2d_v4f32_f32
1399 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1400 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1401 "llvm.nvvm.tld4.g.2d.v4f32.f32">;
1402 def int_nvvm_tld4_b_2d_v4f32_f32
1403 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1404 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1405 "llvm.nvvm.tld4.b.2d.v4f32.f32">;
1406 def int_nvvm_tld4_a_2d_v4f32_f32
1407 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1408 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1409 "llvm.nvvm.tld4.a.2d.v4f32.f32">;
1410 def int_nvvm_tld4_r_2d_v4s32_f32
1411 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1412 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1413 "llvm.nvvm.tld4.r.2d.v4s32.f32">;
1414 def int_nvvm_tld4_g_2d_v4s32_f32
1415 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1416 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1417 "llvm.nvvm.tld4.g.2d.v4s32.f32">;
1418 def int_nvvm_tld4_b_2d_v4s32_f32
1419 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1420 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1421 "llvm.nvvm.tld4.b.2d.v4s32.f32">;
1422 def int_nvvm_tld4_a_2d_v4s32_f32
1423 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1424 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1425 "llvm.nvvm.tld4.a.2d.v4s32.f32">;
1426 def int_nvvm_tld4_r_2d_v4u32_f32
1427 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1428 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1429 "llvm.nvvm.tld4.r.2d.v4u32.f32">;
1430 def int_nvvm_tld4_g_2d_v4u32_f32
1431 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1432 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1433 "llvm.nvvm.tld4.g.2d.v4u32.f32">;
1434 def int_nvvm_tld4_b_2d_v4u32_f32
1435 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1436 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1437 "llvm.nvvm.tld4.b.2d.v4u32.f32">;
1438 def int_nvvm_tld4_a_2d_v4u32_f32
1439 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1440 [llvm_i64_ty, llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1441 "llvm.nvvm.tld4.a.2d.v4u32.f32">;
1445 def int_nvvm_tex_unified_1d_v4f32_s32
1446 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1447 [llvm_i64_ty, llvm_i32_ty], [],
1448 "llvm.nvvm.tex.unified.1d.v4f32.s32">;
1449 def int_nvvm_tex_unified_1d_v4f32_f32
1450 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1451 [llvm_i64_ty, llvm_float_ty], [],
1452 "llvm.nvvm.tex.unified.1d.v4f32.f32">;
1453 def int_nvvm_tex_unified_1d_level_v4f32_f32
1454 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1455 [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1456 "llvm.nvvm.tex.unified.1d.level.v4f32.f32">;
1457 def int_nvvm_tex_unified_1d_grad_v4f32_f32
1458 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1459 [llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1461 "llvm.nvvm.tex.unified.1d.grad.v4f32.f32">;
1462 def int_nvvm_tex_unified_1d_v4s32_s32
1463 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1464 [llvm_i64_ty, llvm_i32_ty], [],
1465 "llvm.nvvm.tex.unified.1d.v4s32.s32">;
1466 def int_nvvm_tex_unified_1d_v4s32_f32
1467 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1468 [llvm_i64_ty, llvm_float_ty], [],
1469 "llvm.nvvm.tex.unified.1d.v4s32.f32">;
1470 def int_nvvm_tex_unified_1d_level_v4s32_f32
1471 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1472 [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1473 "llvm.nvvm.tex.unified.1d.level.v4s32.f32">;
1474 def int_nvvm_tex_unified_1d_grad_v4s32_f32
1475 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1476 [llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1478 "llvm.nvvm.tex.unified.1d.grad.v4s32.f32">;
1479 def int_nvvm_tex_unified_1d_v4u32_s32
1480 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1481 [llvm_i64_ty, llvm_i32_ty], [],
1482 "llvm.nvvm.tex.unified.1d.v4u32.s32">;
1483 def int_nvvm_tex_unified_1d_v4u32_f32
1484 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1485 [llvm_i64_ty, llvm_float_ty], [],
1486 "llvm.nvvm.tex.unified.1d.v4u32.f32">;
1487 def int_nvvm_tex_unified_1d_level_v4u32_f32
1488 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1489 [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1490 "llvm.nvvm.tex.unified.1d.level.v4u32.f32">;
1491 def int_nvvm_tex_unified_1d_grad_v4u32_f32
1492 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1493 [llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1495 "llvm.nvvm.tex.unified.1d.grad.v4u32.f32">;
1497 def int_nvvm_tex_unified_1d_array_v4f32_s32
1498 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1499 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1500 "llvm.nvvm.tex.unified.1d.array.v4f32.s32">;
1501 def int_nvvm_tex_unified_1d_array_v4f32_f32
1502 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1503 [llvm_i64_ty, llvm_i32_ty, llvm_float_ty], [],
1504 "llvm.nvvm.tex.unified.1d.array.v4f32.f32">;
1505 def int_nvvm_tex_unified_1d_array_level_v4f32_f32
1506 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1507 [llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1509 "llvm.nvvm.tex.unified.1d.array.level.v4f32.f32">;
1510 def int_nvvm_tex_unified_1d_array_grad_v4f32_f32
1511 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1512 [llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1513 llvm_float_ty, llvm_float_ty], [],
1514 "llvm.nvvm.tex.unified.1d.array.grad.v4f32.f32">;
1515 def int_nvvm_tex_unified_1d_array_v4s32_s32
1516 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1517 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1518 "llvm.nvvm.tex.unified.1d.array.v4s32.s32">;
1519 def int_nvvm_tex_unified_1d_array_v4s32_f32
1520 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1521 [llvm_i64_ty, llvm_i32_ty, llvm_float_ty], [],
1522 "llvm.nvvm.tex.unified.1d.array.v4s32.f32">;
1523 def int_nvvm_tex_unified_1d_array_level_v4s32_f32
1524 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1525 [llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1527 "llvm.nvvm.tex.unified.1d.array.level.v4s32.f32">;
1528 def int_nvvm_tex_unified_1d_array_grad_v4s32_f32
1529 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1530 [llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1531 llvm_float_ty, llvm_float_ty], [],
1532 "llvm.nvvm.tex.unified.1d.array.grad.v4s32.f32">;
1533 def int_nvvm_tex_unified_1d_array_v4u32_s32
1534 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1535 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1536 "llvm.nvvm.tex.unified.1d.array.v4u32.s32">;
1537 def int_nvvm_tex_unified_1d_array_v4u32_f32
1538 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1539 [llvm_i64_ty, llvm_i32_ty, llvm_float_ty], [],
1540 "llvm.nvvm.tex.unified.1d.array.v4u32.f32">;
1541 def int_nvvm_tex_unified_1d_array_level_v4u32_f32
1542 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1543 [llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1545 "llvm.nvvm.tex.unified.1d.array.level.v4u32.f32">;
1546 def int_nvvm_tex_unified_1d_array_grad_v4u32_f32
1547 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1548 [llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1549 llvm_float_ty, llvm_float_ty], [],
1550 "llvm.nvvm.tex.unified.1d.array.grad.v4u32.f32">;
1552 def int_nvvm_tex_unified_2d_v4f32_s32
1553 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1554 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1555 "llvm.nvvm.tex.unified.2d.v4f32.s32">;
1556 def int_nvvm_tex_unified_2d_v4f32_f32
1557 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1558 [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1559 "llvm.nvvm.tex.unified.2d.v4f32.f32">;
1560 def int_nvvm_tex_unified_2d_level_v4f32_f32
1561 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1562 [llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1564 "llvm.nvvm.tex.unified.2d.level.v4f32.f32">;
1565 def int_nvvm_tex_unified_2d_grad_v4f32_f32
1566 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1567 [llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1568 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1569 "llvm.nvvm.tex.unified.2d.grad.v4f32.f32">;
1570 def int_nvvm_tex_unified_2d_v4s32_s32
1571 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1572 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1573 "llvm.nvvm.tex.unified.2d.v4s32.s32">;
1574 def int_nvvm_tex_unified_2d_v4s32_f32
1575 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1576 [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1577 "llvm.nvvm.tex.unified.2d.v4s32.f32">;
1578 def int_nvvm_tex_unified_2d_level_v4s32_f32
1579 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1580 [llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1582 "llvm.nvvm.tex.unified.2d.level.v4s32.f32">;
1583 def int_nvvm_tex_unified_2d_grad_v4s32_f32
1584 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1585 [llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1586 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1587 "llvm.nvvm.tex.unified.2d.grad.v4s32.f32">;
1588 def int_nvvm_tex_unified_2d_v4u32_s32
1589 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1590 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1591 "llvm.nvvm.tex.unified.2d.v4u32.s32">;
1592 def int_nvvm_tex_unified_2d_v4u32_f32
1593 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1594 [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1595 "llvm.nvvm.tex.unified.2d.v4u32.f32">;
1596 def int_nvvm_tex_unified_2d_level_v4u32_f32
1597 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1598 [llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1600 "llvm.nvvm.tex.unified.2d.level.v4u32.f32">;
1601 def int_nvvm_tex_unified_2d_grad_v4u32_f32
1602 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1603 [llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1604 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1605 "llvm.nvvm.tex.unified.2d.grad.v4u32.f32">;
1607 def int_nvvm_tex_unified_2d_array_v4f32_s32
1608 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1609 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
1611 "llvm.nvvm.tex.unified.2d.array.v4f32.s32">;
1612 def int_nvvm_tex_unified_2d_array_v4f32_f32
1613 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1614 [llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1616 "llvm.nvvm.tex.unified.2d.array.v4f32.f32">;
1617 def int_nvvm_tex_unified_2d_array_level_v4f32_f32
1618 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1619 [llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1620 llvm_float_ty, llvm_float_ty], [],
1621 "llvm.nvvm.tex.unified.2d.array.level.v4f32.f32">;
1622 def int_nvvm_tex_unified_2d_array_grad_v4f32_f32
1623 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1624 [llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1625 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty,
1627 "llvm.nvvm.tex.unified.2d.array.grad.v4f32.f32">;
1628 def int_nvvm_tex_unified_2d_array_v4s32_s32
1629 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1630 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
1632 "llvm.nvvm.tex.unified.2d.array.v4s32.s32">;
1633 def int_nvvm_tex_unified_2d_array_v4s32_f32
1634 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1635 [llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1637 "llvm.nvvm.tex.unified.2d.array.v4s32.f32">;
1638 def int_nvvm_tex_unified_2d_array_level_v4s32_f32
1639 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1640 [llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1641 llvm_float_ty, llvm_float_ty], [],
1642 "llvm.nvvm.tex.unified.2d.array.level.v4s32.f32">;
1643 def int_nvvm_tex_unified_2d_array_grad_v4s32_f32
1644 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1645 [llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1646 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty,
1648 "llvm.nvvm.tex.unified.2d.array.grad.v4s32.f32">;
1649 def int_nvvm_tex_unified_2d_array_v4u32_s32
1650 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1651 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
1653 "llvm.nvvm.tex.unified.2d.array.v4u32.s32">;
1654 def int_nvvm_tex_unified_2d_array_v4u32_f32
1655 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1656 [llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1658 "llvm.nvvm.tex.unified.2d.array.v4u32.f32">;
1659 def int_nvvm_tex_unified_2d_array_level_v4u32_f32
1660 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1661 [llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1662 llvm_float_ty, llvm_float_ty], [],
1663 "llvm.nvvm.tex.unified.2d.array.level.v4u32.f32">;
1664 def int_nvvm_tex_unified_2d_array_grad_v4u32_f32
1665 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1666 [llvm_i64_ty, llvm_i32_ty, llvm_float_ty,
1667 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty,
1669 "llvm.nvvm.tex.unified.2d.array.grad.v4u32.f32">;
1671 def int_nvvm_tex_unified_3d_v4f32_s32
1672 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1673 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1674 [], "llvm.nvvm.tex.unified.3d.v4f32.s32">;
1675 def int_nvvm_tex_unified_3d_v4f32_f32
1676 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1677 [llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1679 "llvm.nvvm.tex.unified.3d.v4f32.f32">;
1680 def int_nvvm_tex_unified_3d_level_v4f32_f32
1681 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1682 [llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1683 llvm_float_ty, llvm_float_ty], [],
1684 "llvm.nvvm.tex.unified.3d.level.v4f32.f32">;
1685 def int_nvvm_tex_unified_3d_grad_v4f32_f32
1686 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1687 [llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1688 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty,
1689 llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1690 "llvm.nvvm.tex.unified.3d.grad.v4f32.f32">;
1691 def int_nvvm_tex_unified_3d_v4s32_s32
1692 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1693 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1694 [], "llvm.nvvm.tex.unified.3d.v4s32.s32">;
1695 def int_nvvm_tex_unified_3d_v4s32_f32
1696 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1697 [llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1699 "llvm.nvvm.tex.unified.3d.v4s32.f32">;
1700 def int_nvvm_tex_unified_3d_level_v4s32_f32
1701 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1702 [llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1703 llvm_float_ty, llvm_float_ty], [],
1704 "llvm.nvvm.tex.unified.3d.level.v4s32.f32">;
1705 def int_nvvm_tex_unified_3d_grad_v4s32_f32
1706 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1707 [llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1708 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty,
1709 llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1710 "llvm.nvvm.tex.unified.3d.grad.v4s32.f32">;
1711 def int_nvvm_tex_unified_3d_v4u32_s32
1712 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1713 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1714 [], "llvm.nvvm.tex.unified.3d.v4u32.s32">;
1715 def int_nvvm_tex_unified_3d_v4u32_f32
1716 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1717 [llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1719 "llvm.nvvm.tex.unified.3d.v4u32.f32">;
1720 def int_nvvm_tex_unified_3d_level_v4u32_f32
1721 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1722 [llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1723 llvm_float_ty, llvm_float_ty], [],
1724 "llvm.nvvm.tex.unified.3d.level.v4u32.f32">;
1725 def int_nvvm_tex_unified_3d_grad_v4u32_f32
1726 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1727 [llvm_i64_ty, llvm_float_ty, llvm_float_ty,
1728 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty,
1729 llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1730 "llvm.nvvm.tex.unified.3d.grad.v4u32.f32">;
1732 def int_nvvm_tex_unified_cube_v4f32_f32
1733 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1735 llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1736 "llvm.nvvm.tex.unified.cube.v4f32.f32">;
1737 def int_nvvm_tex_unified_cube_level_v4f32_f32
1738 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1740 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1741 "llvm.nvvm.tex.unified.cube.level.v4f32.f32">;
1742 def int_nvvm_tex_unified_cube_v4s32_f32
1743 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1745 llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1746 "llvm.nvvm.tex.unified.cube.v4s32.f32">;
1747 def int_nvvm_tex_unified_cube_level_v4s32_f32
1748 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1750 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1751 "llvm.nvvm.tex.unified.cube.level.v4s32.f32">;
1752 def int_nvvm_tex_unified_cube_v4u32_f32
1753 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1755 llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1756 "llvm.nvvm.tex.unified.cube.v4u32.f32">;
1757 def int_nvvm_tex_unified_cube_level_v4u32_f32
1758 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1760 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1761 "llvm.nvvm.tex.unified.cube.level.v4u32.f32">;
1763 def int_nvvm_tex_unified_cube_array_v4f32_f32
1764 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1765 [llvm_i64_ty, llvm_i32_ty,
1766 llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1767 "llvm.nvvm.tex.unified.cube.array.v4f32.f32">;
1768 def int_nvvm_tex_unified_cube_array_level_v4f32_f32
1769 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1770 [llvm_i64_ty, llvm_i32_ty,
1771 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1772 "llvm.nvvm.tex.unified.cube.array.level.v4f32.f32">;
1773 def int_nvvm_tex_unified_cube_array_v4s32_f32
1774 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1775 [llvm_i64_ty, llvm_i32_ty,
1776 llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1777 "llvm.nvvm.tex.unified.cube.array.v4s32.f32">;
1778 def int_nvvm_tex_unified_cube_array_level_v4s32_f32
1779 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1780 [llvm_i64_ty, llvm_i32_ty,
1781 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1782 "llvm.nvvm.tex.unified.cube.array.level.v4s32.f32">;
1783 def int_nvvm_tex_unified_cube_array_v4u32_f32
1784 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1785 [llvm_i64_ty, llvm_i32_ty,
1786 llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1787 "llvm.nvvm.tex.unified.cube.array.v4u32.f32">;
1788 def int_nvvm_tex_unified_cube_array_level_v4u32_f32
1789 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1790 [llvm_i64_ty, llvm_i32_ty,
1791 llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty], [],
1792 "llvm.nvvm.tex.unified.cube.array.level.v4u32.f32">;
1794 def int_nvvm_tld4_unified_r_2d_v4f32_f32
1795 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1796 [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1797 "llvm.nvvm.tld4.unified.r.2d.v4f32.f32">;
1798 def int_nvvm_tld4_unified_g_2d_v4f32_f32
1799 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1800 [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1801 "llvm.nvvm.tld4.unified.g.2d.v4f32.f32">;
1802 def int_nvvm_tld4_unified_b_2d_v4f32_f32
1803 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1804 [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1805 "llvm.nvvm.tld4.unified.b.2d.v4f32.f32">;
1806 def int_nvvm_tld4_unified_a_2d_v4f32_f32
1807 : Intrinsic<[llvm_float_ty, llvm_float_ty, llvm_float_ty, llvm_float_ty],
1808 [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1809 "llvm.nvvm.tld4.unified.a.2d.v4f32.f32">;
1810 def int_nvvm_tld4_unified_r_2d_v4s32_f32
1811 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1812 [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1813 "llvm.nvvm.tld4.unified.r.2d.v4s32.f32">;
1814 def int_nvvm_tld4_unified_g_2d_v4s32_f32
1815 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1816 [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1817 "llvm.nvvm.tld4.unified.g.2d.v4s32.f32">;
1818 def int_nvvm_tld4_unified_b_2d_v4s32_f32
1819 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1820 [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1821 "llvm.nvvm.tld4.unified.b.2d.v4s32.f32">;
1822 def int_nvvm_tld4_unified_a_2d_v4s32_f32
1823 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1824 [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1825 "llvm.nvvm.tld4.unified.a.2d.v4s32.f32">;
1826 def int_nvvm_tld4_unified_r_2d_v4u32_f32
1827 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1828 [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1829 "llvm.nvvm.tld4.unified.r.2d.v4u32.f32">;
1830 def int_nvvm_tld4_unified_g_2d_v4u32_f32
1831 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1832 [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1833 "llvm.nvvm.tld4.unified.g.2d.v4u32.f32">;
1834 def int_nvvm_tld4_unified_b_2d_v4u32_f32
1835 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1836 [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1837 "llvm.nvvm.tld4.unified.b.2d.v4u32.f32">;
1838 def int_nvvm_tld4_unified_a_2d_v4u32_f32
1839 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1840 [llvm_i64_ty, llvm_float_ty, llvm_float_ty], [],
1841 "llvm.nvvm.tld4.unified.a.2d.v4u32.f32">;
1846 def int_nvvm_suld_1d_i8_clamp
1847 : Intrinsic<[llvm_i16_ty],
1848 [llvm_i64_ty, llvm_i32_ty], [],
1849 "llvm.nvvm.suld.1d.i8.clamp">;
1850 def int_nvvm_suld_1d_i16_clamp
1851 : Intrinsic<[llvm_i16_ty],
1852 [llvm_i64_ty, llvm_i32_ty], [],
1853 "llvm.nvvm.suld.1d.i16.clamp">;
1854 def int_nvvm_suld_1d_i32_clamp
1855 : Intrinsic<[llvm_i32_ty],
1856 [llvm_i64_ty, llvm_i32_ty], [],
1857 "llvm.nvvm.suld.1d.i32.clamp">;
1858 def int_nvvm_suld_1d_i64_clamp
1859 : Intrinsic<[llvm_i64_ty],
1860 [llvm_i64_ty, llvm_i32_ty], [],
1861 "llvm.nvvm.suld.1d.i64.clamp">;
1862 def int_nvvm_suld_1d_v2i8_clamp
1863 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
1864 [llvm_i64_ty, llvm_i32_ty], [],
1865 "llvm.nvvm.suld.1d.v2i8.clamp">;
1866 def int_nvvm_suld_1d_v2i16_clamp
1867 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
1868 [llvm_i64_ty, llvm_i32_ty], [],
1869 "llvm.nvvm.suld.1d.v2i16.clamp">;
1870 def int_nvvm_suld_1d_v2i32_clamp
1871 : Intrinsic<[llvm_i32_ty, llvm_i32_ty],
1872 [llvm_i64_ty, llvm_i32_ty], [],
1873 "llvm.nvvm.suld.1d.v2i32.clamp">;
1874 def int_nvvm_suld_1d_v2i64_clamp
1875 : Intrinsic<[llvm_i64_ty, llvm_i64_ty],
1876 [llvm_i64_ty, llvm_i32_ty], [],
1877 "llvm.nvvm.suld.1d.v2i64.clamp">;
1878 def int_nvvm_suld_1d_v4i8_clamp
1879 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
1880 [llvm_i64_ty, llvm_i32_ty], [],
1881 "llvm.nvvm.suld.1d.v4i8.clamp">;
1882 def int_nvvm_suld_1d_v4i16_clamp
1883 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
1884 [llvm_i64_ty, llvm_i32_ty], [],
1885 "llvm.nvvm.suld.1d.v4i16.clamp">;
1886 def int_nvvm_suld_1d_v4i32_clamp
1887 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1888 [llvm_i64_ty, llvm_i32_ty], [],
1889 "llvm.nvvm.suld.1d.v4i32.clamp">;
1891 def int_nvvm_suld_1d_array_i8_clamp
1892 : Intrinsic<[llvm_i16_ty],
1893 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1894 "llvm.nvvm.suld.1d.array.i8.clamp">;
1895 def int_nvvm_suld_1d_array_i16_clamp
1896 : Intrinsic<[llvm_i16_ty],
1897 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1898 "llvm.nvvm.suld.1d.array.i16.clamp">;
1899 def int_nvvm_suld_1d_array_i32_clamp
1900 : Intrinsic<[llvm_i32_ty],
1901 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1902 "llvm.nvvm.suld.1d.array.i32.clamp">;
1903 def int_nvvm_suld_1d_array_i64_clamp
1904 : Intrinsic<[llvm_i64_ty],
1905 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1906 "llvm.nvvm.suld.1d.array.i64.clamp">;
1907 def int_nvvm_suld_1d_array_v2i8_clamp
1908 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
1909 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1910 "llvm.nvvm.suld.1d.array.v2i8.clamp">;
1911 def int_nvvm_suld_1d_array_v2i16_clamp
1912 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
1913 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1914 "llvm.nvvm.suld.1d.array.v2i16.clamp">;
1915 def int_nvvm_suld_1d_array_v2i32_clamp
1916 : Intrinsic<[llvm_i32_ty, llvm_i32_ty],
1917 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1918 "llvm.nvvm.suld.1d.array.v2i32.clamp">;
1919 def int_nvvm_suld_1d_array_v2i64_clamp
1920 : Intrinsic<[llvm_i64_ty, llvm_i64_ty],
1921 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1922 "llvm.nvvm.suld.1d.array.v2i64.clamp">;
1923 def int_nvvm_suld_1d_array_v4i8_clamp
1924 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
1925 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1926 "llvm.nvvm.suld.1d.array.v4i8.clamp">;
1927 def int_nvvm_suld_1d_array_v4i16_clamp
1928 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
1929 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1930 "llvm.nvvm.suld.1d.array.v4i16.clamp">;
1931 def int_nvvm_suld_1d_array_v4i32_clamp
1932 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1933 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1934 "llvm.nvvm.suld.1d.array.v4i32.clamp">;
1936 def int_nvvm_suld_2d_i8_clamp
1937 : Intrinsic<[llvm_i16_ty],
1938 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1939 "llvm.nvvm.suld.2d.i8.clamp">;
1940 def int_nvvm_suld_2d_i16_clamp
1941 : Intrinsic<[llvm_i16_ty],
1942 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1943 "llvm.nvvm.suld.2d.i16.clamp">;
1944 def int_nvvm_suld_2d_i32_clamp
1945 : Intrinsic<[llvm_i32_ty],
1946 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1947 "llvm.nvvm.suld.2d.i32.clamp">;
1948 def int_nvvm_suld_2d_i64_clamp
1949 : Intrinsic<[llvm_i64_ty],
1950 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1951 "llvm.nvvm.suld.2d.i64.clamp">;
1952 def int_nvvm_suld_2d_v2i8_clamp
1953 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
1954 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1955 "llvm.nvvm.suld.2d.v2i8.clamp">;
1956 def int_nvvm_suld_2d_v2i16_clamp
1957 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
1958 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1959 "llvm.nvvm.suld.2d.v2i16.clamp">;
1960 def int_nvvm_suld_2d_v2i32_clamp
1961 : Intrinsic<[llvm_i32_ty, llvm_i32_ty],
1962 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1963 "llvm.nvvm.suld.2d.v2i32.clamp">;
1964 def int_nvvm_suld_2d_v2i64_clamp
1965 : Intrinsic<[llvm_i64_ty, llvm_i64_ty],
1966 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1967 "llvm.nvvm.suld.2d.v2i64.clamp">;
1968 def int_nvvm_suld_2d_v4i8_clamp
1969 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
1970 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1971 "llvm.nvvm.suld.2d.v4i8.clamp">;
1972 def int_nvvm_suld_2d_v4i16_clamp
1973 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
1974 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1975 "llvm.nvvm.suld.2d.v4i16.clamp">;
1976 def int_nvvm_suld_2d_v4i32_clamp
1977 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
1978 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
1979 "llvm.nvvm.suld.2d.v4i32.clamp">;
1981 def int_nvvm_suld_2d_array_i8_clamp
1982 : Intrinsic<[llvm_i16_ty],
1983 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
1984 "llvm.nvvm.suld.2d.array.i8.clamp">;
1985 def int_nvvm_suld_2d_array_i16_clamp
1986 : Intrinsic<[llvm_i16_ty],
1987 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
1988 "llvm.nvvm.suld.2d.array.i16.clamp">;
1989 def int_nvvm_suld_2d_array_i32_clamp
1990 : Intrinsic<[llvm_i32_ty],
1991 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
1992 "llvm.nvvm.suld.2d.array.i32.clamp">;
1993 def int_nvvm_suld_2d_array_i64_clamp
1994 : Intrinsic<[llvm_i64_ty],
1995 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
1996 "llvm.nvvm.suld.2d.array.i64.clamp">;
1997 def int_nvvm_suld_2d_array_v2i8_clamp
1998 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
1999 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2000 "llvm.nvvm.suld.2d.array.v2i8.clamp">;
2001 def int_nvvm_suld_2d_array_v2i16_clamp
2002 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2003 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2004 "llvm.nvvm.suld.2d.array.v2i16.clamp">;
2005 def int_nvvm_suld_2d_array_v2i32_clamp
2006 : Intrinsic<[llvm_i32_ty, llvm_i32_ty],
2007 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2008 "llvm.nvvm.suld.2d.array.v2i32.clamp">;
2009 def int_nvvm_suld_2d_array_v2i64_clamp
2010 : Intrinsic<[llvm_i64_ty, llvm_i64_ty],
2011 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2012 "llvm.nvvm.suld.2d.array.v2i64.clamp">;
2013 def int_nvvm_suld_2d_array_v4i8_clamp
2014 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2015 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2016 "llvm.nvvm.suld.2d.array.v4i8.clamp">;
2017 def int_nvvm_suld_2d_array_v4i16_clamp
2018 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2019 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2020 "llvm.nvvm.suld.2d.array.v4i16.clamp">;
2021 def int_nvvm_suld_2d_array_v4i32_clamp
2022 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
2023 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2024 "llvm.nvvm.suld.2d.array.v4i32.clamp">;
2026 def int_nvvm_suld_3d_i8_clamp
2027 : Intrinsic<[llvm_i16_ty],
2028 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2029 "llvm.nvvm.suld.3d.i8.clamp">;
2030 def int_nvvm_suld_3d_i16_clamp
2031 : Intrinsic<[llvm_i16_ty],
2032 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2033 "llvm.nvvm.suld.3d.i16.clamp">;
2034 def int_nvvm_suld_3d_i32_clamp
2035 : Intrinsic<[llvm_i32_ty],
2036 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2037 "llvm.nvvm.suld.3d.i32.clamp">;
2038 def int_nvvm_suld_3d_i64_clamp
2039 : Intrinsic<[llvm_i64_ty],
2040 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2041 "llvm.nvvm.suld.3d.i64.clamp">;
2042 def int_nvvm_suld_3d_v2i8_clamp
2043 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2044 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2045 "llvm.nvvm.suld.3d.v2i8.clamp">;
2046 def int_nvvm_suld_3d_v2i16_clamp
2047 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2048 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2049 "llvm.nvvm.suld.3d.v2i16.clamp">;
2050 def int_nvvm_suld_3d_v2i32_clamp
2051 : Intrinsic<[llvm_i32_ty, llvm_i32_ty],
2052 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2053 "llvm.nvvm.suld.3d.v2i32.clamp">;
2054 def int_nvvm_suld_3d_v2i64_clamp
2055 : Intrinsic<[llvm_i64_ty, llvm_i64_ty],
2056 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2057 "llvm.nvvm.suld.3d.v2i64.clamp">;
2058 def int_nvvm_suld_3d_v4i8_clamp
2059 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2060 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2061 "llvm.nvvm.suld.3d.v4i8.clamp">;
2062 def int_nvvm_suld_3d_v4i16_clamp
2063 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2064 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2065 "llvm.nvvm.suld.3d.v4i16.clamp">;
2066 def int_nvvm_suld_3d_v4i32_clamp
2067 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
2068 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2069 "llvm.nvvm.suld.3d.v4i32.clamp">;
2072 def int_nvvm_suld_1d_i8_trap
2073 : Intrinsic<[llvm_i16_ty],
2074 [llvm_i64_ty, llvm_i32_ty], [],
2075 "llvm.nvvm.suld.1d.i8.trap">;
2076 def int_nvvm_suld_1d_i16_trap
2077 : Intrinsic<[llvm_i16_ty],
2078 [llvm_i64_ty, llvm_i32_ty], [],
2079 "llvm.nvvm.suld.1d.i16.trap">;
2080 def int_nvvm_suld_1d_i32_trap
2081 : Intrinsic<[llvm_i32_ty],
2082 [llvm_i64_ty, llvm_i32_ty], [],
2083 "llvm.nvvm.suld.1d.i32.trap">;
2084 def int_nvvm_suld_1d_i64_trap
2085 : Intrinsic<[llvm_i64_ty],
2086 [llvm_i64_ty, llvm_i32_ty], [],
2087 "llvm.nvvm.suld.1d.i64.trap">;
2088 def int_nvvm_suld_1d_v2i8_trap
2089 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2090 [llvm_i64_ty, llvm_i32_ty], [],
2091 "llvm.nvvm.suld.1d.v2i8.trap">;
2092 def int_nvvm_suld_1d_v2i16_trap
2093 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2094 [llvm_i64_ty, llvm_i32_ty], [],
2095 "llvm.nvvm.suld.1d.v2i16.trap">;
2096 def int_nvvm_suld_1d_v2i32_trap
2097 : Intrinsic<[llvm_i32_ty, llvm_i32_ty],
2098 [llvm_i64_ty, llvm_i32_ty], [],
2099 "llvm.nvvm.suld.1d.v2i32.trap">;
2100 def int_nvvm_suld_1d_v2i64_trap
2101 : Intrinsic<[llvm_i64_ty, llvm_i64_ty],
2102 [llvm_i64_ty, llvm_i32_ty], [],
2103 "llvm.nvvm.suld.1d.v2i64.trap">;
2104 def int_nvvm_suld_1d_v4i8_trap
2105 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2106 [llvm_i64_ty, llvm_i32_ty], [],
2107 "llvm.nvvm.suld.1d.v4i8.trap">;
2108 def int_nvvm_suld_1d_v4i16_trap
2109 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2110 [llvm_i64_ty, llvm_i32_ty], [],
2111 "llvm.nvvm.suld.1d.v4i16.trap">;
2112 def int_nvvm_suld_1d_v4i32_trap
2113 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
2114 [llvm_i64_ty, llvm_i32_ty], [],
2115 "llvm.nvvm.suld.1d.v4i32.trap">;
2117 def int_nvvm_suld_1d_array_i8_trap
2118 : Intrinsic<[llvm_i16_ty],
2119 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2120 "llvm.nvvm.suld.1d.array.i8.trap">;
2121 def int_nvvm_suld_1d_array_i16_trap
2122 : Intrinsic<[llvm_i16_ty],
2123 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2124 "llvm.nvvm.suld.1d.array.i16.trap">;
2125 def int_nvvm_suld_1d_array_i32_trap
2126 : Intrinsic<[llvm_i32_ty],
2127 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2128 "llvm.nvvm.suld.1d.array.i32.trap">;
2129 def int_nvvm_suld_1d_array_i64_trap
2130 : Intrinsic<[llvm_i64_ty],
2131 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2132 "llvm.nvvm.suld.1d.array.i64.trap">;
2133 def int_nvvm_suld_1d_array_v2i8_trap
2134 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2135 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2136 "llvm.nvvm.suld.1d.array.v2i8.trap">;
2137 def int_nvvm_suld_1d_array_v2i16_trap
2138 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2139 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2140 "llvm.nvvm.suld.1d.array.v2i16.trap">;
2141 def int_nvvm_suld_1d_array_v2i32_trap
2142 : Intrinsic<[llvm_i32_ty, llvm_i32_ty],
2143 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2144 "llvm.nvvm.suld.1d.array.v2i32.trap">;
2145 def int_nvvm_suld_1d_array_v2i64_trap
2146 : Intrinsic<[llvm_i64_ty, llvm_i64_ty],
2147 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2148 "llvm.nvvm.suld.1d.array.v2i64.trap">;
2149 def int_nvvm_suld_1d_array_v4i8_trap
2150 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2151 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2152 "llvm.nvvm.suld.1d.array.v4i8.trap">;
2153 def int_nvvm_suld_1d_array_v4i16_trap
2154 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2155 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2156 "llvm.nvvm.suld.1d.array.v4i16.trap">;
2157 def int_nvvm_suld_1d_array_v4i32_trap
2158 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
2159 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2160 "llvm.nvvm.suld.1d.array.v4i32.trap">;
2162 def int_nvvm_suld_2d_i8_trap
2163 : Intrinsic<[llvm_i16_ty],
2164 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2165 "llvm.nvvm.suld.2d.i8.trap">;
2166 def int_nvvm_suld_2d_i16_trap
2167 : Intrinsic<[llvm_i16_ty],
2168 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2169 "llvm.nvvm.suld.2d.i16.trap">;
2170 def int_nvvm_suld_2d_i32_trap
2171 : Intrinsic<[llvm_i32_ty],
2172 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2173 "llvm.nvvm.suld.2d.i32.trap">;
2174 def int_nvvm_suld_2d_i64_trap
2175 : Intrinsic<[llvm_i64_ty],
2176 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2177 "llvm.nvvm.suld.2d.i64.trap">;
2178 def int_nvvm_suld_2d_v2i8_trap
2179 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2180 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2181 "llvm.nvvm.suld.2d.v2i8.trap">;
2182 def int_nvvm_suld_2d_v2i16_trap
2183 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2184 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2185 "llvm.nvvm.suld.2d.v2i16.trap">;
2186 def int_nvvm_suld_2d_v2i32_trap
2187 : Intrinsic<[llvm_i32_ty, llvm_i32_ty],
2188 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2189 "llvm.nvvm.suld.2d.v2i32.trap">;
2190 def int_nvvm_suld_2d_v2i64_trap
2191 : Intrinsic<[llvm_i64_ty, llvm_i64_ty],
2192 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2193 "llvm.nvvm.suld.2d.v2i64.trap">;
2194 def int_nvvm_suld_2d_v4i8_trap
2195 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2196 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2197 "llvm.nvvm.suld.2d.v4i8.trap">;
2198 def int_nvvm_suld_2d_v4i16_trap
2199 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2200 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2201 "llvm.nvvm.suld.2d.v4i16.trap">;
2202 def int_nvvm_suld_2d_v4i32_trap
2203 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
2204 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2205 "llvm.nvvm.suld.2d.v4i32.trap">;
2207 def int_nvvm_suld_2d_array_i8_trap
2208 : Intrinsic<[llvm_i16_ty],
2209 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2210 "llvm.nvvm.suld.2d.array.i8.trap">;
2211 def int_nvvm_suld_2d_array_i16_trap
2212 : Intrinsic<[llvm_i16_ty],
2213 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2214 "llvm.nvvm.suld.2d.array.i16.trap">;
2215 def int_nvvm_suld_2d_array_i32_trap
2216 : Intrinsic<[llvm_i32_ty],
2217 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2218 "llvm.nvvm.suld.2d.array.i32.trap">;
2219 def int_nvvm_suld_2d_array_i64_trap
2220 : Intrinsic<[llvm_i64_ty],
2221 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2222 "llvm.nvvm.suld.2d.array.i64.trap">;
2223 def int_nvvm_suld_2d_array_v2i8_trap
2224 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2225 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2226 "llvm.nvvm.suld.2d.array.v2i8.trap">;
2227 def int_nvvm_suld_2d_array_v2i16_trap
2228 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2229 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2230 "llvm.nvvm.suld.2d.array.v2i16.trap">;
2231 def int_nvvm_suld_2d_array_v2i32_trap
2232 : Intrinsic<[llvm_i32_ty, llvm_i32_ty],
2233 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2234 "llvm.nvvm.suld.2d.array.v2i32.trap">;
2235 def int_nvvm_suld_2d_array_v2i64_trap
2236 : Intrinsic<[llvm_i64_ty, llvm_i64_ty],
2237 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2238 "llvm.nvvm.suld.2d.array.v2i64.trap">;
2239 def int_nvvm_suld_2d_array_v4i8_trap
2240 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2241 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2242 "llvm.nvvm.suld.2d.array.v4i8.trap">;
2243 def int_nvvm_suld_2d_array_v4i16_trap
2244 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2245 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2246 "llvm.nvvm.suld.2d.array.v4i16.trap">;
2247 def int_nvvm_suld_2d_array_v4i32_trap
2248 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
2249 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2250 "llvm.nvvm.suld.2d.array.v4i32.trap">;
2252 def int_nvvm_suld_3d_i8_trap
2253 : Intrinsic<[llvm_i16_ty],
2254 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2255 "llvm.nvvm.suld.3d.i8.trap">;
2256 def int_nvvm_suld_3d_i16_trap
2257 : Intrinsic<[llvm_i16_ty],
2258 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2259 "llvm.nvvm.suld.3d.i16.trap">;
2260 def int_nvvm_suld_3d_i32_trap
2261 : Intrinsic<[llvm_i32_ty],
2262 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2263 "llvm.nvvm.suld.3d.i32.trap">;
2264 def int_nvvm_suld_3d_i64_trap
2265 : Intrinsic<[llvm_i64_ty],
2266 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2267 "llvm.nvvm.suld.3d.i64.trap">;
2268 def int_nvvm_suld_3d_v2i8_trap
2269 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2270 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2271 "llvm.nvvm.suld.3d.v2i8.trap">;
2272 def int_nvvm_suld_3d_v2i16_trap
2273 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2274 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2275 "llvm.nvvm.suld.3d.v2i16.trap">;
2276 def int_nvvm_suld_3d_v2i32_trap
2277 : Intrinsic<[llvm_i32_ty, llvm_i32_ty],
2278 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2279 "llvm.nvvm.suld.3d.v2i32.trap">;
2280 def int_nvvm_suld_3d_v2i64_trap
2281 : Intrinsic<[llvm_i64_ty, llvm_i64_ty],
2282 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2283 "llvm.nvvm.suld.3d.v2i64.trap">;
2284 def int_nvvm_suld_3d_v4i8_trap
2285 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2286 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2287 "llvm.nvvm.suld.3d.v4i8.trap">;
2288 def int_nvvm_suld_3d_v4i16_trap
2289 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2290 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2291 "llvm.nvvm.suld.3d.v4i16.trap">;
2292 def int_nvvm_suld_3d_v4i32_trap
2293 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
2294 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2295 "llvm.nvvm.suld.3d.v4i32.trap">;
2298 def int_nvvm_suld_1d_i8_zero
2299 : Intrinsic<[llvm_i16_ty],
2300 [llvm_i64_ty, llvm_i32_ty], [],
2301 "llvm.nvvm.suld.1d.i8.zero">;
2302 def int_nvvm_suld_1d_i16_zero
2303 : Intrinsic<[llvm_i16_ty],
2304 [llvm_i64_ty, llvm_i32_ty], [],
2305 "llvm.nvvm.suld.1d.i16.zero">;
2306 def int_nvvm_suld_1d_i32_zero
2307 : Intrinsic<[llvm_i32_ty],
2308 [llvm_i64_ty, llvm_i32_ty], [],
2309 "llvm.nvvm.suld.1d.i32.zero">;
2310 def int_nvvm_suld_1d_i64_zero
2311 : Intrinsic<[llvm_i64_ty],
2312 [llvm_i64_ty, llvm_i32_ty], [],
2313 "llvm.nvvm.suld.1d.i64.zero">;
2314 def int_nvvm_suld_1d_v2i8_zero
2315 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2316 [llvm_i64_ty, llvm_i32_ty], [],
2317 "llvm.nvvm.suld.1d.v2i8.zero">;
2318 def int_nvvm_suld_1d_v2i16_zero
2319 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2320 [llvm_i64_ty, llvm_i32_ty], [],
2321 "llvm.nvvm.suld.1d.v2i16.zero">;
2322 def int_nvvm_suld_1d_v2i32_zero
2323 : Intrinsic<[llvm_i32_ty, llvm_i32_ty],
2324 [llvm_i64_ty, llvm_i32_ty], [],
2325 "llvm.nvvm.suld.1d.v2i32.zero">;
2326 def int_nvvm_suld_1d_v2i64_zero
2327 : Intrinsic<[llvm_i64_ty, llvm_i64_ty],
2328 [llvm_i64_ty, llvm_i32_ty], [],
2329 "llvm.nvvm.suld.1d.v2i64.zero">;
2330 def int_nvvm_suld_1d_v4i8_zero
2331 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2332 [llvm_i64_ty, llvm_i32_ty], [],
2333 "llvm.nvvm.suld.1d.v4i8.zero">;
2334 def int_nvvm_suld_1d_v4i16_zero
2335 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2336 [llvm_i64_ty, llvm_i32_ty], [],
2337 "llvm.nvvm.suld.1d.v4i16.zero">;
2338 def int_nvvm_suld_1d_v4i32_zero
2339 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
2340 [llvm_i64_ty, llvm_i32_ty], [],
2341 "llvm.nvvm.suld.1d.v4i32.zero">;
2343 def int_nvvm_suld_1d_array_i8_zero
2344 : Intrinsic<[llvm_i16_ty],
2345 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2346 "llvm.nvvm.suld.1d.array.i8.zero">;
2347 def int_nvvm_suld_1d_array_i16_zero
2348 : Intrinsic<[llvm_i16_ty],
2349 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2350 "llvm.nvvm.suld.1d.array.i16.zero">;
2351 def int_nvvm_suld_1d_array_i32_zero
2352 : Intrinsic<[llvm_i32_ty],
2353 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2354 "llvm.nvvm.suld.1d.array.i32.zero">;
2355 def int_nvvm_suld_1d_array_i64_zero
2356 : Intrinsic<[llvm_i64_ty],
2357 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2358 "llvm.nvvm.suld.1d.array.i64.zero">;
2359 def int_nvvm_suld_1d_array_v2i8_zero
2360 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2361 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2362 "llvm.nvvm.suld.1d.array.v2i8.zero">;
2363 def int_nvvm_suld_1d_array_v2i16_zero
2364 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2365 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2366 "llvm.nvvm.suld.1d.array.v2i16.zero">;
2367 def int_nvvm_suld_1d_array_v2i32_zero
2368 : Intrinsic<[llvm_i32_ty, llvm_i32_ty],
2369 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2370 "llvm.nvvm.suld.1d.array.v2i32.zero">;
2371 def int_nvvm_suld_1d_array_v2i64_zero
2372 : Intrinsic<[llvm_i64_ty, llvm_i64_ty],
2373 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2374 "llvm.nvvm.suld.1d.array.v2i64.zero">;
2375 def int_nvvm_suld_1d_array_v4i8_zero
2376 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2377 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2378 "llvm.nvvm.suld.1d.array.v4i8.zero">;
2379 def int_nvvm_suld_1d_array_v4i16_zero
2380 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2381 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2382 "llvm.nvvm.suld.1d.array.v4i16.zero">;
2383 def int_nvvm_suld_1d_array_v4i32_zero
2384 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
2385 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2386 "llvm.nvvm.suld.1d.array.v4i32.zero">;
2388 def int_nvvm_suld_2d_i8_zero
2389 : Intrinsic<[llvm_i16_ty],
2390 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2391 "llvm.nvvm.suld.2d.i8.zero">;
2392 def int_nvvm_suld_2d_i16_zero
2393 : Intrinsic<[llvm_i16_ty],
2394 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2395 "llvm.nvvm.suld.2d.i16.zero">;
2396 def int_nvvm_suld_2d_i32_zero
2397 : Intrinsic<[llvm_i32_ty],
2398 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2399 "llvm.nvvm.suld.2d.i32.zero">;
2400 def int_nvvm_suld_2d_i64_zero
2401 : Intrinsic<[llvm_i64_ty],
2402 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2403 "llvm.nvvm.suld.2d.i64.zero">;
2404 def int_nvvm_suld_2d_v2i8_zero
2405 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2406 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2407 "llvm.nvvm.suld.2d.v2i8.zero">;
2408 def int_nvvm_suld_2d_v2i16_zero
2409 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2410 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2411 "llvm.nvvm.suld.2d.v2i16.zero">;
2412 def int_nvvm_suld_2d_v2i32_zero
2413 : Intrinsic<[llvm_i32_ty, llvm_i32_ty],
2414 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2415 "llvm.nvvm.suld.2d.v2i32.zero">;
2416 def int_nvvm_suld_2d_v2i64_zero
2417 : Intrinsic<[llvm_i64_ty, llvm_i64_ty],
2418 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2419 "llvm.nvvm.suld.2d.v2i64.zero">;
2420 def int_nvvm_suld_2d_v4i8_zero
2421 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2422 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2423 "llvm.nvvm.suld.2d.v4i8.zero">;
2424 def int_nvvm_suld_2d_v4i16_zero
2425 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2426 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2427 "llvm.nvvm.suld.2d.v4i16.zero">;
2428 def int_nvvm_suld_2d_v4i32_zero
2429 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
2430 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2431 "llvm.nvvm.suld.2d.v4i32.zero">;
2433 def int_nvvm_suld_2d_array_i8_zero
2434 : Intrinsic<[llvm_i16_ty],
2435 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2436 "llvm.nvvm.suld.2d.array.i8.zero">;
2437 def int_nvvm_suld_2d_array_i16_zero
2438 : Intrinsic<[llvm_i16_ty],
2439 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2440 "llvm.nvvm.suld.2d.array.i16.zero">;
2441 def int_nvvm_suld_2d_array_i32_zero
2442 : Intrinsic<[llvm_i32_ty],
2443 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2444 "llvm.nvvm.suld.2d.array.i32.zero">;
2445 def int_nvvm_suld_2d_array_i64_zero
2446 : Intrinsic<[llvm_i64_ty],
2447 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2448 "llvm.nvvm.suld.2d.array.i64.zero">;
2449 def int_nvvm_suld_2d_array_v2i8_zero
2450 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2451 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2452 "llvm.nvvm.suld.2d.array.v2i8.zero">;
2453 def int_nvvm_suld_2d_array_v2i16_zero
2454 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2455 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2456 "llvm.nvvm.suld.2d.array.v2i16.zero">;
2457 def int_nvvm_suld_2d_array_v2i32_zero
2458 : Intrinsic<[llvm_i32_ty, llvm_i32_ty],
2459 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2460 "llvm.nvvm.suld.2d.array.v2i32.zero">;
2461 def int_nvvm_suld_2d_array_v2i64_zero
2462 : Intrinsic<[llvm_i64_ty, llvm_i64_ty],
2463 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2464 "llvm.nvvm.suld.2d.array.v2i64.zero">;
2465 def int_nvvm_suld_2d_array_v4i8_zero
2466 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2467 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2468 "llvm.nvvm.suld.2d.array.v4i8.zero">;
2469 def int_nvvm_suld_2d_array_v4i16_zero
2470 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2471 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2472 "llvm.nvvm.suld.2d.array.v4i16.zero">;
2473 def int_nvvm_suld_2d_array_v4i32_zero
2474 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
2475 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2476 "llvm.nvvm.suld.2d.array.v4i32.zero">;
2478 def int_nvvm_suld_3d_i8_zero
2479 : Intrinsic<[llvm_i16_ty],
2480 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2481 "llvm.nvvm.suld.3d.i8.zero">;
2482 def int_nvvm_suld_3d_i16_zero
2483 : Intrinsic<[llvm_i16_ty],
2484 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2485 "llvm.nvvm.suld.3d.i16.zero">;
2486 def int_nvvm_suld_3d_i32_zero
2487 : Intrinsic<[llvm_i32_ty],
2488 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2489 "llvm.nvvm.suld.3d.i32.zero">;
2490 def int_nvvm_suld_3d_i64_zero
2491 : Intrinsic<[llvm_i64_ty],
2492 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2493 "llvm.nvvm.suld.3d.i64.zero">;
2494 def int_nvvm_suld_3d_v2i8_zero
2495 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2496 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2497 "llvm.nvvm.suld.3d.v2i8.zero">;
2498 def int_nvvm_suld_3d_v2i16_zero
2499 : Intrinsic<[llvm_i16_ty, llvm_i16_ty],
2500 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2501 "llvm.nvvm.suld.3d.v2i16.zero">;
2502 def int_nvvm_suld_3d_v2i32_zero
2503 : Intrinsic<[llvm_i32_ty, llvm_i32_ty],
2504 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2505 "llvm.nvvm.suld.3d.v2i32.zero">;
2506 def int_nvvm_suld_3d_v2i64_zero
2507 : Intrinsic<[llvm_i64_ty, llvm_i64_ty],
2508 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2509 "llvm.nvvm.suld.3d.v2i64.zero">;
2510 def int_nvvm_suld_3d_v4i8_zero
2511 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2512 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2513 "llvm.nvvm.suld.3d.v4i8.zero">;
2514 def int_nvvm_suld_3d_v4i16_zero
2515 : Intrinsic<[llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty],
2516 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2517 "llvm.nvvm.suld.3d.v4i16.zero">;
2518 def int_nvvm_suld_3d_v4i32_zero
2519 : Intrinsic<[llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
2520 [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2521 "llvm.nvvm.suld.3d.v4i32.zero">;
2523 //===- Texture Query ------------------------------------------------------===//
2525 def int_nvvm_txq_channel_order
2526 : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem],
2527 "llvm.nvvm.txq.channel.order">,
2528 GCCBuiltin<"__nvvm_txq_channel_order">;
2529 def int_nvvm_txq_channel_data_type
2530 : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem],
2531 "llvm.nvvm.txq.channel.data.type">,
2532 GCCBuiltin<"__nvvm_txq_channel_data_type">;
2533 def int_nvvm_txq_width
2534 : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem],
2535 "llvm.nvvm.txq.width">,
2536 GCCBuiltin<"__nvvm_txq_width">;
2537 def int_nvvm_txq_height
2538 : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem],
2539 "llvm.nvvm.txq.height">,
2540 GCCBuiltin<"__nvvm_txq_height">;
2541 def int_nvvm_txq_depth
2542 : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem],
2543 "llvm.nvvm.txq.depth">,
2544 GCCBuiltin<"__nvvm_txq_depth">;
2545 def int_nvvm_txq_array_size
2546 : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem],
2547 "llvm.nvvm.txq.array.size">,
2548 GCCBuiltin<"__nvvm_txq_array_size">;
2549 def int_nvvm_txq_num_samples
2550 : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem],
2551 "llvm.nvvm.txq.num.samples">,
2552 GCCBuiltin<"__nvvm_txq_num_samples">;
2553 def int_nvvm_txq_num_mipmap_levels
2554 : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem],
2555 "llvm.nvvm.txq.num.mipmap.levels">,
2556 GCCBuiltin<"__nvvm_txq_num_mipmap_levels">;
2558 //===- Surface Query ------------------------------------------------------===//
2560 def int_nvvm_suq_channel_order
2561 : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem],
2562 "llvm.nvvm.suq.channel.order">,
2563 GCCBuiltin<"__nvvm_suq_channel_order">;
2564 def int_nvvm_suq_channel_data_type
2565 : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem],
2566 "llvm.nvvm.suq.channel.data.type">,
2567 GCCBuiltin<"__nvvm_suq_channel_data_type">;
2568 def int_nvvm_suq_width
2569 : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem],
2570 "llvm.nvvm.suq.width">,
2571 GCCBuiltin<"__nvvm_suq_width">;
2572 def int_nvvm_suq_height
2573 : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem],
2574 "llvm.nvvm.suq.height">,
2575 GCCBuiltin<"__nvvm_suq_height">;
2576 def int_nvvm_suq_depth
2577 : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem],
2578 "llvm.nvvm.suq.depth">,
2579 GCCBuiltin<"__nvvm_suq_depth">;
2580 def int_nvvm_suq_array_size
2581 : Intrinsic<[llvm_i32_ty], [llvm_i64_ty], [IntrNoMem],
2582 "llvm.nvvm.suq.array.size">,
2583 GCCBuiltin<"__nvvm_suq_array_size">;
2586 //===- Handle Query -------------------------------------------------------===//
2588 def int_nvvm_istypep_sampler
2589 : Intrinsic<[llvm_i1_ty], [llvm_i64_ty], [IntrNoMem],
2590 "llvm.nvvm.istypep.sampler">,
2591 GCCBuiltin<"__nvvm_istypep_sampler">;
2592 def int_nvvm_istypep_surface
2593 : Intrinsic<[llvm_i1_ty], [llvm_i64_ty], [IntrNoMem],
2594 "llvm.nvvm.istypep.surface">,
2595 GCCBuiltin<"__nvvm_istypep_surface">;
2596 def int_nvvm_istypep_texture
2597 : Intrinsic<[llvm_i1_ty], [llvm_i64_ty], [IntrNoMem],
2598 "llvm.nvvm.istypep.texture">,
2599 GCCBuiltin<"__nvvm_istypep_texture">;
2603 //===- Surface Stores -----------------------------------------------------===//
2607 def int_nvvm_sust_b_1d_i8_clamp
2608 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty], [],
2609 "llvm.nvvm.sust.b.1d.i8.clamp">,
2610 GCCBuiltin<"__nvvm_sust_b_1d_i8_clamp">;
2611 def int_nvvm_sust_b_1d_i16_clamp
2612 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty], [],
2613 "llvm.nvvm.sust.b.1d.i16.clamp">,
2614 GCCBuiltin<"__nvvm_sust_b_1d_i16_clamp">;
2615 def int_nvvm_sust_b_1d_i32_clamp
2616 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2617 "llvm.nvvm.sust.b.1d.i32.clamp">,
2618 GCCBuiltin<"__nvvm_sust_b_1d_i32_clamp">;
2619 def int_nvvm_sust_b_1d_i64_clamp
2620 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty], [],
2621 "llvm.nvvm.sust.b.1d.i64.clamp">,
2622 GCCBuiltin<"__nvvm_sust_b_1d_i64_clamp">;
2623 def int_nvvm_sust_b_1d_v2i8_clamp
2624 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [],
2625 "llvm.nvvm.sust.b.1d.v2i8.clamp">,
2626 GCCBuiltin<"__nvvm_sust_b_1d_v2i8_clamp">;
2627 def int_nvvm_sust_b_1d_v2i16_clamp
2628 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [],
2629 "llvm.nvvm.sust.b.1d.v2i16.clamp">,
2630 GCCBuiltin<"__nvvm_sust_b_1d_v2i16_clamp">;
2631 def int_nvvm_sust_b_1d_v2i32_clamp
2632 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2633 "llvm.nvvm.sust.b.1d.v2i32.clamp">,
2634 GCCBuiltin<"__nvvm_sust_b_1d_v2i32_clamp">;
2635 def int_nvvm_sust_b_1d_v2i64_clamp
2636 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty, llvm_i64_ty], [],
2637 "llvm.nvvm.sust.b.1d.v2i64.clamp">,
2638 GCCBuiltin<"__nvvm_sust_b_1d_v2i64_clamp">;
2639 def int_nvvm_sust_b_1d_v4i8_clamp
2640 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty,
2641 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
2642 "llvm.nvvm.sust.b.1d.v4i8.clamp">,
2643 GCCBuiltin<"__nvvm_sust_b_1d_v4i8_clamp">;
2644 def int_nvvm_sust_b_1d_v4i16_clamp
2645 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty,
2646 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
2647 "llvm.nvvm.sust.b.1d.v4i16.clamp">,
2648 GCCBuiltin<"__nvvm_sust_b_1d_v4i16_clamp">;
2649 def int_nvvm_sust_b_1d_v4i32_clamp
2650 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2651 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2652 "llvm.nvvm.sust.b.1d.v4i32.clamp">,
2653 GCCBuiltin<"__nvvm_sust_b_1d_v4i32_clamp">;
2656 def int_nvvm_sust_b_1d_array_i8_clamp
2657 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [],
2658 "llvm.nvvm.sust.b.1d.array.i8.clamp">,
2659 GCCBuiltin<"__nvvm_sust_b_1d_array_i8_clamp">;
2660 def int_nvvm_sust_b_1d_array_i16_clamp
2661 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [],
2662 "llvm.nvvm.sust.b.1d.array.i16.clamp">,
2663 GCCBuiltin<"__nvvm_sust_b_1d_array_i16_clamp">;
2664 def int_nvvm_sust_b_1d_array_i32_clamp
2665 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2666 "llvm.nvvm.sust.b.1d.array.i32.clamp">,
2667 GCCBuiltin<"__nvvm_sust_b_1d_array_i32_clamp">;
2668 def int_nvvm_sust_b_1d_array_i64_clamp
2669 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty], [],
2670 "llvm.nvvm.sust.b.1d.array.i64.clamp">,
2671 GCCBuiltin<"__nvvm_sust_b_1d_array_i64_clamp">;
2672 def int_nvvm_sust_b_1d_array_v2i8_clamp
2673 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2674 llvm_i16_ty, llvm_i16_ty], [],
2675 "llvm.nvvm.sust.b.1d.array.v2i8.clamp">,
2676 GCCBuiltin<"__nvvm_sust_b_1d_array_v2i8_clamp">;
2677 def int_nvvm_sust_b_1d_array_v2i16_clamp
2678 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2679 llvm_i16_ty, llvm_i16_ty], [],
2680 "llvm.nvvm.sust.b.1d.array.v2i16.clamp">,
2681 GCCBuiltin<"__nvvm_sust_b_1d_array_v2i16_clamp">;
2682 def int_nvvm_sust_b_1d_array_v2i32_clamp
2683 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2684 llvm_i32_ty, llvm_i32_ty], [],
2685 "llvm.nvvm.sust.b.1d.array.v2i32.clamp">,
2686 GCCBuiltin<"__nvvm_sust_b_1d_array_v2i32_clamp">;
2687 def int_nvvm_sust_b_1d_array_v2i64_clamp
2688 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2689 llvm_i64_ty, llvm_i64_ty], [],
2690 "llvm.nvvm.sust.b.1d.array.v2i64.clamp">,
2691 GCCBuiltin<"__nvvm_sust_b_1d_array_v2i64_clamp">;
2692 def int_nvvm_sust_b_1d_array_v4i8_clamp
2693 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty,
2694 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
2695 "llvm.nvvm.sust.b.1d.array.v4i8.clamp">,
2696 GCCBuiltin<"__nvvm_sust_b_1d_array_v4i8_clamp">;
2697 def int_nvvm_sust_b_1d_array_v4i16_clamp
2698 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty,
2699 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
2700 "llvm.nvvm.sust.b.1d.array.v4i16.clamp">,
2701 GCCBuiltin<"__nvvm_sust_b_1d_array_v4i16_clamp">;
2702 def int_nvvm_sust_b_1d_array_v4i32_clamp
2703 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
2704 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2705 "llvm.nvvm.sust.b.1d.array.v4i32.clamp">,
2706 GCCBuiltin<"__nvvm_sust_b_1d_array_v4i32_clamp">;
2709 def int_nvvm_sust_b_2d_i8_clamp
2710 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [],
2711 "llvm.nvvm.sust.b.2d.i8.clamp">,
2712 GCCBuiltin<"__nvvm_sust_b_2d_i8_clamp">;
2713 def int_nvvm_sust_b_2d_i16_clamp
2714 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [],
2715 "llvm.nvvm.sust.b.2d.i16.clamp">,
2716 GCCBuiltin<"__nvvm_sust_b_2d_i16_clamp">;
2717 def int_nvvm_sust_b_2d_i32_clamp
2718 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2719 "llvm.nvvm.sust.b.2d.i32.clamp">,
2720 GCCBuiltin<"__nvvm_sust_b_2d_i32_clamp">;
2721 def int_nvvm_sust_b_2d_i64_clamp
2722 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty], [],
2723 "llvm.nvvm.sust.b.2d.i64.clamp">,
2724 GCCBuiltin<"__nvvm_sust_b_2d_i64_clamp">;
2725 def int_nvvm_sust_b_2d_v2i8_clamp
2726 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2727 llvm_i16_ty, llvm_i16_ty], [],
2728 "llvm.nvvm.sust.b.2d.v2i8.clamp">,
2729 GCCBuiltin<"__nvvm_sust_b_2d_v2i8_clamp">;
2730 def int_nvvm_sust_b_2d_v2i16_clamp
2731 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2732 llvm_i16_ty, llvm_i16_ty], [],
2733 "llvm.nvvm.sust.b.2d.v2i16.clamp">,
2734 GCCBuiltin<"__nvvm_sust_b_2d_v2i16_clamp">;
2735 def int_nvvm_sust_b_2d_v2i32_clamp
2736 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2737 llvm_i32_ty, llvm_i32_ty], [],
2738 "llvm.nvvm.sust.b.2d.v2i32.clamp">,
2739 GCCBuiltin<"__nvvm_sust_b_2d_v2i32_clamp">;
2740 def int_nvvm_sust_b_2d_v2i64_clamp
2741 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2742 llvm_i64_ty, llvm_i64_ty], [],
2743 "llvm.nvvm.sust.b.2d.v2i64.clamp">,
2744 GCCBuiltin<"__nvvm_sust_b_2d_v2i64_clamp">;
2745 def int_nvvm_sust_b_2d_v4i8_clamp
2746 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty,
2747 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
2748 "llvm.nvvm.sust.b.2d.v4i8.clamp">,
2749 GCCBuiltin<"__nvvm_sust_b_2d_v4i8_clamp">;
2750 def int_nvvm_sust_b_2d_v4i16_clamp
2751 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty,
2752 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
2753 "llvm.nvvm.sust.b.2d.v4i16.clamp">,
2754 GCCBuiltin<"__nvvm_sust_b_2d_v4i16_clamp">;
2755 def int_nvvm_sust_b_2d_v4i32_clamp
2756 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
2757 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2758 "llvm.nvvm.sust.b.2d.v4i32.clamp">,
2759 GCCBuiltin<"__nvvm_sust_b_2d_v4i32_clamp">;
2762 def int_nvvm_sust_b_2d_array_i8_clamp
2763 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2764 llvm_i32_ty, llvm_i16_ty], [],
2765 "llvm.nvvm.sust.b.2d.array.i8.clamp">,
2766 GCCBuiltin<"__nvvm_sust_b_2d_array_i8_clamp">;
2767 def int_nvvm_sust_b_2d_array_i16_clamp
2768 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2769 llvm_i32_ty, llvm_i16_ty], [],
2770 "llvm.nvvm.sust.b.2d.array.i16.clamp">,
2771 GCCBuiltin<"__nvvm_sust_b_2d_array_i16_clamp">;
2772 def int_nvvm_sust_b_2d_array_i32_clamp
2773 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2774 llvm_i32_ty, llvm_i32_ty], [],
2775 "llvm.nvvm.sust.b.2d.array.i32.clamp">,
2776 GCCBuiltin<"__nvvm_sust_b_2d_array_i32_clamp">;
2777 def int_nvvm_sust_b_2d_array_i64_clamp
2778 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2779 llvm_i32_ty, llvm_i64_ty], [],
2780 "llvm.nvvm.sust.b.2d.array.i64.clamp">,
2781 GCCBuiltin<"__nvvm_sust_b_2d_array_i64_clamp">;
2782 def int_nvvm_sust_b_2d_array_v2i8_clamp
2783 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
2784 llvm_i16_ty, llvm_i16_ty], [],
2785 "llvm.nvvm.sust.b.2d.array.v2i8.clamp">,
2786 GCCBuiltin<"__nvvm_sust_b_2d_array_v2i8_clamp">;
2787 def int_nvvm_sust_b_2d_array_v2i16_clamp
2788 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
2789 llvm_i16_ty, llvm_i16_ty], [],
2790 "llvm.nvvm.sust.b.2d.array.v2i16.clamp">,
2791 GCCBuiltin<"__nvvm_sust_b_2d_array_v2i16_clamp">;
2792 def int_nvvm_sust_b_2d_array_v2i32_clamp
2793 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
2794 llvm_i32_ty, llvm_i32_ty], [],
2795 "llvm.nvvm.sust.b.2d.array.v2i32.clamp">,
2796 GCCBuiltin<"__nvvm_sust_b_2d_array_v2i32_clamp">;
2797 def int_nvvm_sust_b_2d_array_v2i64_clamp
2798 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
2799 llvm_i64_ty, llvm_i64_ty], [],
2800 "llvm.nvvm.sust.b.2d.array.v2i64.clamp">,
2801 GCCBuiltin<"__nvvm_sust_b_2d_array_v2i64_clamp">;
2802 def int_nvvm_sust_b_2d_array_v4i8_clamp
2803 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
2804 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
2805 "llvm.nvvm.sust.b.2d.array.v4i8.clamp">,
2806 GCCBuiltin<"__nvvm_sust_b_2d_array_v4i8_clamp">;
2807 def int_nvvm_sust_b_2d_array_v4i16_clamp
2808 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
2809 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
2810 "llvm.nvvm.sust.b.2d.array.v4i16.clamp">,
2811 GCCBuiltin<"__nvvm_sust_b_2d_array_v4i16_clamp">;
2812 def int_nvvm_sust_b_2d_array_v4i32_clamp
2813 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
2814 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2815 "llvm.nvvm.sust.b.2d.array.v4i32.clamp">,
2816 GCCBuiltin<"__nvvm_sust_b_2d_array_v4i32_clamp">;
2819 def int_nvvm_sust_b_3d_i8_clamp
2820 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2821 llvm_i32_ty, llvm_i16_ty], [],
2822 "llvm.nvvm.sust.b.3d.i8.clamp">,
2823 GCCBuiltin<"__nvvm_sust_b_3d_i8_clamp">;
2824 def int_nvvm_sust_b_3d_i16_clamp
2825 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2826 llvm_i32_ty, llvm_i16_ty], [],
2827 "llvm.nvvm.sust.b.3d.i16.clamp">,
2828 GCCBuiltin<"__nvvm_sust_b_3d_i16_clamp">;
2829 def int_nvvm_sust_b_3d_i32_clamp
2830 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2831 llvm_i32_ty, llvm_i32_ty], [],
2832 "llvm.nvvm.sust.b.3d.i32.clamp">,
2833 GCCBuiltin<"__nvvm_sust_b_3d_i32_clamp">;
2834 def int_nvvm_sust_b_3d_i64_clamp
2835 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2836 llvm_i32_ty, llvm_i64_ty], [],
2837 "llvm.nvvm.sust.b.3d.i64.clamp">,
2838 GCCBuiltin<"__nvvm_sust_b_3d_i64_clamp">;
2839 def int_nvvm_sust_b_3d_v2i8_clamp
2840 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
2841 llvm_i16_ty, llvm_i16_ty], [],
2842 "llvm.nvvm.sust.b.3d.v2i8.clamp">,
2843 GCCBuiltin<"__nvvm_sust_b_3d_v2i8_clamp">;
2844 def int_nvvm_sust_b_3d_v2i16_clamp
2845 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
2846 llvm_i16_ty, llvm_i16_ty], [],
2847 "llvm.nvvm.sust.b.3d.v2i16.clamp">,
2848 GCCBuiltin<"__nvvm_sust_b_3d_v2i16_clamp">;
2849 def int_nvvm_sust_b_3d_v2i32_clamp
2850 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
2851 llvm_i32_ty, llvm_i32_ty], [],
2852 "llvm.nvvm.sust.b.3d.v2i32.clamp">,
2853 GCCBuiltin<"__nvvm_sust_b_3d_v2i32_clamp">;
2854 def int_nvvm_sust_b_3d_v2i64_clamp
2855 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
2856 llvm_i64_ty, llvm_i64_ty], [],
2857 "llvm.nvvm.sust.b.3d.v2i64.clamp">,
2858 GCCBuiltin<"__nvvm_sust_b_3d_v2i64_clamp">;
2859 def int_nvvm_sust_b_3d_v4i8_clamp
2860 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
2861 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
2862 "llvm.nvvm.sust.b.3d.v4i8.clamp">,
2863 GCCBuiltin<"__nvvm_sust_b_3d_v4i8_clamp">;
2864 def int_nvvm_sust_b_3d_v4i16_clamp
2865 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
2866 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
2867 "llvm.nvvm.sust.b.3d.v4i16.clamp">,
2868 GCCBuiltin<"__nvvm_sust_b_3d_v4i16_clamp">;
2869 def int_nvvm_sust_b_3d_v4i32_clamp
2870 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
2871 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2872 "llvm.nvvm.sust.b.3d.v4i32.clamp">,
2873 GCCBuiltin<"__nvvm_sust_b_3d_v4i32_clamp">;
2877 def int_nvvm_sust_b_1d_i8_trap
2878 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty], [],
2879 "llvm.nvvm.sust.b.1d.i8.trap">,
2880 GCCBuiltin<"__nvvm_sust_b_1d_i8_trap">;
2881 def int_nvvm_sust_b_1d_i16_trap
2882 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty], [],
2883 "llvm.nvvm.sust.b.1d.i16.trap">,
2884 GCCBuiltin<"__nvvm_sust_b_1d_i16_trap">;
2885 def int_nvvm_sust_b_1d_i32_trap
2886 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
2887 "llvm.nvvm.sust.b.1d.i32.trap">,
2888 GCCBuiltin<"__nvvm_sust_b_1d_i32_trap">;
2889 def int_nvvm_sust_b_1d_i64_trap
2890 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty], [],
2891 "llvm.nvvm.sust.b.1d.i64.trap">,
2892 GCCBuiltin<"__nvvm_sust_b_1d_i64_trap">;
2893 def int_nvvm_sust_b_1d_v2i8_trap
2894 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [],
2895 "llvm.nvvm.sust.b.1d.v2i8.trap">,
2896 GCCBuiltin<"__nvvm_sust_b_1d_v2i8_trap">;
2897 def int_nvvm_sust_b_1d_v2i16_trap
2898 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [],
2899 "llvm.nvvm.sust.b.1d.v2i16.trap">,
2900 GCCBuiltin<"__nvvm_sust_b_1d_v2i16_trap">;
2901 def int_nvvm_sust_b_1d_v2i32_trap
2902 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2903 "llvm.nvvm.sust.b.1d.v2i32.trap">,
2904 GCCBuiltin<"__nvvm_sust_b_1d_v2i32_trap">;
2905 def int_nvvm_sust_b_1d_v2i64_trap
2906 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty, llvm_i64_ty], [],
2907 "llvm.nvvm.sust.b.1d.v2i64.trap">,
2908 GCCBuiltin<"__nvvm_sust_b_1d_v2i64_trap">;
2909 def int_nvvm_sust_b_1d_v4i8_trap
2910 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty,
2911 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
2912 "llvm.nvvm.sust.b.1d.v4i8.trap">,
2913 GCCBuiltin<"__nvvm_sust_b_1d_v4i8_trap">;
2914 def int_nvvm_sust_b_1d_v4i16_trap
2915 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty,
2916 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
2917 "llvm.nvvm.sust.b.1d.v4i16.trap">,
2918 GCCBuiltin<"__nvvm_sust_b_1d_v4i16_trap">;
2919 def int_nvvm_sust_b_1d_v4i32_trap
2920 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2921 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2922 "llvm.nvvm.sust.b.1d.v4i32.trap">,
2923 GCCBuiltin<"__nvvm_sust_b_1d_v4i32_trap">;
2926 def int_nvvm_sust_b_1d_array_i8_trap
2927 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [],
2928 "llvm.nvvm.sust.b.1d.array.i8.trap">,
2929 GCCBuiltin<"__nvvm_sust_b_1d_array_i8_trap">;
2930 def int_nvvm_sust_b_1d_array_i16_trap
2931 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [],
2932 "llvm.nvvm.sust.b.1d.array.i16.trap">,
2933 GCCBuiltin<"__nvvm_sust_b_1d_array_i16_trap">;
2934 def int_nvvm_sust_b_1d_array_i32_trap
2935 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2936 "llvm.nvvm.sust.b.1d.array.i32.trap">,
2937 GCCBuiltin<"__nvvm_sust_b_1d_array_i32_trap">;
2938 def int_nvvm_sust_b_1d_array_i64_trap
2939 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty], [],
2940 "llvm.nvvm.sust.b.1d.array.i64.trap">,
2941 GCCBuiltin<"__nvvm_sust_b_1d_array_i64_trap">;
2942 def int_nvvm_sust_b_1d_array_v2i8_trap
2943 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2944 llvm_i16_ty, llvm_i16_ty], [],
2945 "llvm.nvvm.sust.b.1d.array.v2i8.trap">,
2946 GCCBuiltin<"__nvvm_sust_b_1d_array_v2i8_trap">;
2947 def int_nvvm_sust_b_1d_array_v2i16_trap
2948 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2949 llvm_i16_ty, llvm_i16_ty], [],
2950 "llvm.nvvm.sust.b.1d.array.v2i16.trap">,
2951 GCCBuiltin<"__nvvm_sust_b_1d_array_v2i16_trap">;
2952 def int_nvvm_sust_b_1d_array_v2i32_trap
2953 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2954 llvm_i32_ty, llvm_i32_ty], [],
2955 "llvm.nvvm.sust.b.1d.array.v2i32.trap">,
2956 GCCBuiltin<"__nvvm_sust_b_1d_array_v2i32_trap">;
2957 def int_nvvm_sust_b_1d_array_v2i64_trap
2958 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2959 llvm_i64_ty, llvm_i64_ty], [],
2960 "llvm.nvvm.sust.b.1d.array.v2i64.trap">,
2961 GCCBuiltin<"__nvvm_sust_b_1d_array_v2i64_trap">;
2962 def int_nvvm_sust_b_1d_array_v4i8_trap
2963 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty,
2964 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
2965 "llvm.nvvm.sust.b.1d.array.v4i8.trap">,
2966 GCCBuiltin<"__nvvm_sust_b_1d_array_v4i8_trap">;
2967 def int_nvvm_sust_b_1d_array_v4i16_trap
2968 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty,
2969 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
2970 "llvm.nvvm.sust.b.1d.array.v4i16.trap">,
2971 GCCBuiltin<"__nvvm_sust_b_1d_array_v4i16_trap">;
2972 def int_nvvm_sust_b_1d_array_v4i32_trap
2973 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
2974 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2975 "llvm.nvvm.sust.b.1d.array.v4i32.trap">,
2976 GCCBuiltin<"__nvvm_sust_b_1d_array_v4i32_trap">;
2979 def int_nvvm_sust_b_2d_i8_trap
2980 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [],
2981 "llvm.nvvm.sust.b.2d.i8.trap">,
2982 GCCBuiltin<"__nvvm_sust_b_2d_i8_trap">;
2983 def int_nvvm_sust_b_2d_i16_trap
2984 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [],
2985 "llvm.nvvm.sust.b.2d.i16.trap">,
2986 GCCBuiltin<"__nvvm_sust_b_2d_i16_trap">;
2987 def int_nvvm_sust_b_2d_i32_trap
2988 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
2989 "llvm.nvvm.sust.b.2d.i32.trap">,
2990 GCCBuiltin<"__nvvm_sust_b_2d_i32_trap">;
2991 def int_nvvm_sust_b_2d_i64_trap
2992 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty], [],
2993 "llvm.nvvm.sust.b.2d.i64.trap">,
2994 GCCBuiltin<"__nvvm_sust_b_2d_i64_trap">;
2995 def int_nvvm_sust_b_2d_v2i8_trap
2996 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
2997 llvm_i16_ty, llvm_i16_ty], [],
2998 "llvm.nvvm.sust.b.2d.v2i8.trap">,
2999 GCCBuiltin<"__nvvm_sust_b_2d_v2i8_trap">;
3000 def int_nvvm_sust_b_2d_v2i16_trap
3001 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3002 llvm_i16_ty, llvm_i16_ty], [],
3003 "llvm.nvvm.sust.b.2d.v2i16.trap">,
3004 GCCBuiltin<"__nvvm_sust_b_2d_v2i16_trap">;
3005 def int_nvvm_sust_b_2d_v2i32_trap
3006 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3007 llvm_i32_ty, llvm_i32_ty], [],
3008 "llvm.nvvm.sust.b.2d.v2i32.trap">,
3009 GCCBuiltin<"__nvvm_sust_b_2d_v2i32_trap">;
3010 def int_nvvm_sust_b_2d_v2i64_trap
3011 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3012 llvm_i64_ty, llvm_i64_ty], [],
3013 "llvm.nvvm.sust.b.2d.v2i64.trap">,
3014 GCCBuiltin<"__nvvm_sust_b_2d_v2i64_trap">;
3015 def int_nvvm_sust_b_2d_v4i8_trap
3016 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty,
3017 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3018 "llvm.nvvm.sust.b.2d.v4i8.trap">,
3019 GCCBuiltin<"__nvvm_sust_b_2d_v4i8_trap">;
3020 def int_nvvm_sust_b_2d_v4i16_trap
3021 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty,
3022 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3023 "llvm.nvvm.sust.b.2d.v4i16.trap">,
3024 GCCBuiltin<"__nvvm_sust_b_2d_v4i16_trap">;
3025 def int_nvvm_sust_b_2d_v4i32_trap
3026 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3027 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3028 "llvm.nvvm.sust.b.2d.v4i32.trap">,
3029 GCCBuiltin<"__nvvm_sust_b_2d_v4i32_trap">;
3032 def int_nvvm_sust_b_2d_array_i8_trap
3033 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3034 llvm_i32_ty, llvm_i16_ty], [],
3035 "llvm.nvvm.sust.b.2d.array.i8.trap">,
3036 GCCBuiltin<"__nvvm_sust_b_2d_array_i8_trap">;
3037 def int_nvvm_sust_b_2d_array_i16_trap
3038 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3039 llvm_i32_ty, llvm_i16_ty], [],
3040 "llvm.nvvm.sust.b.2d.array.i16.trap">,
3041 GCCBuiltin<"__nvvm_sust_b_2d_array_i16_trap">;
3042 def int_nvvm_sust_b_2d_array_i32_trap
3043 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3044 llvm_i32_ty, llvm_i32_ty], [],
3045 "llvm.nvvm.sust.b.2d.array.i32.trap">,
3046 GCCBuiltin<"__nvvm_sust_b_2d_array_i32_trap">;
3047 def int_nvvm_sust_b_2d_array_i64_trap
3048 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3049 llvm_i32_ty, llvm_i64_ty], [],
3050 "llvm.nvvm.sust.b.2d.array.i64.trap">,
3051 GCCBuiltin<"__nvvm_sust_b_2d_array_i64_trap">;
3052 def int_nvvm_sust_b_2d_array_v2i8_trap
3053 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3054 llvm_i16_ty, llvm_i16_ty], [],
3055 "llvm.nvvm.sust.b.2d.array.v2i8.trap">,
3056 GCCBuiltin<"__nvvm_sust_b_2d_array_v2i8_trap">;
3057 def int_nvvm_sust_b_2d_array_v2i16_trap
3058 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3059 llvm_i16_ty, llvm_i16_ty], [],
3060 "llvm.nvvm.sust.b.2d.array.v2i16.trap">,
3061 GCCBuiltin<"__nvvm_sust_b_2d_array_v2i16_trap">;
3062 def int_nvvm_sust_b_2d_array_v2i32_trap
3063 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3064 llvm_i32_ty, llvm_i32_ty], [],
3065 "llvm.nvvm.sust.b.2d.array.v2i32.trap">,
3066 GCCBuiltin<"__nvvm_sust_b_2d_array_v2i32_trap">;
3067 def int_nvvm_sust_b_2d_array_v2i64_trap
3068 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3069 llvm_i64_ty, llvm_i64_ty], [],
3070 "llvm.nvvm.sust.b.2d.array.v2i64.trap">,
3071 GCCBuiltin<"__nvvm_sust_b_2d_array_v2i64_trap">;
3072 def int_nvvm_sust_b_2d_array_v4i8_trap
3073 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3074 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3075 "llvm.nvvm.sust.b.2d.array.v4i8.trap">,
3076 GCCBuiltin<"__nvvm_sust_b_2d_array_v4i8_trap">;
3077 def int_nvvm_sust_b_2d_array_v4i16_trap
3078 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3079 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3080 "llvm.nvvm.sust.b.2d.array.v4i16.trap">,
3081 GCCBuiltin<"__nvvm_sust_b_2d_array_v4i16_trap">;
3082 def int_nvvm_sust_b_2d_array_v4i32_trap
3083 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3084 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3085 "llvm.nvvm.sust.b.2d.array.v4i32.trap">,
3086 GCCBuiltin<"__nvvm_sust_b_2d_array_v4i32_trap">;
3089 def int_nvvm_sust_b_3d_i8_trap
3090 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3091 llvm_i32_ty, llvm_i16_ty], [],
3092 "llvm.nvvm.sust.b.3d.i8.trap">,
3093 GCCBuiltin<"__nvvm_sust_b_3d_i8_trap">;
3094 def int_nvvm_sust_b_3d_i16_trap
3095 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3096 llvm_i32_ty, llvm_i16_ty], [],
3097 "llvm.nvvm.sust.b.3d.i16.trap">,
3098 GCCBuiltin<"__nvvm_sust_b_3d_i16_trap">;
3099 def int_nvvm_sust_b_3d_i32_trap
3100 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3101 llvm_i32_ty, llvm_i32_ty], [],
3102 "llvm.nvvm.sust.b.3d.i32.trap">,
3103 GCCBuiltin<"__nvvm_sust_b_3d_i32_trap">;
3104 def int_nvvm_sust_b_3d_i64_trap
3105 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3106 llvm_i32_ty, llvm_i64_ty], [],
3107 "llvm.nvvm.sust.b.3d.i64.trap">,
3108 GCCBuiltin<"__nvvm_sust_b_3d_i64_trap">;
3109 def int_nvvm_sust_b_3d_v2i8_trap
3110 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3111 llvm_i16_ty, llvm_i16_ty], [],
3112 "llvm.nvvm.sust.b.3d.v2i8.trap">,
3113 GCCBuiltin<"__nvvm_sust_b_3d_v2i8_trap">;
3114 def int_nvvm_sust_b_3d_v2i16_trap
3115 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3116 llvm_i16_ty, llvm_i16_ty], [],
3117 "llvm.nvvm.sust.b.3d.v2i16.trap">,
3118 GCCBuiltin<"__nvvm_sust_b_3d_v2i16_trap">;
3119 def int_nvvm_sust_b_3d_v2i32_trap
3120 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3121 llvm_i32_ty, llvm_i32_ty], [],
3122 "llvm.nvvm.sust.b.3d.v2i32.trap">,
3123 GCCBuiltin<"__nvvm_sust_b_3d_v2i32_trap">;
3124 def int_nvvm_sust_b_3d_v2i64_trap
3125 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3126 llvm_i64_ty, llvm_i64_ty], [],
3127 "llvm.nvvm.sust.b.3d.v2i64.trap">,
3128 GCCBuiltin<"__nvvm_sust_b_3d_v2i64_trap">;
3129 def int_nvvm_sust_b_3d_v4i8_trap
3130 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3131 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3132 "llvm.nvvm.sust.b.3d.v4i8.trap">,
3133 GCCBuiltin<"__nvvm_sust_b_3d_v4i8_trap">;
3134 def int_nvvm_sust_b_3d_v4i16_trap
3135 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3136 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3137 "llvm.nvvm.sust.b.3d.v4i16.trap">,
3138 GCCBuiltin<"__nvvm_sust_b_3d_v4i16_trap">;
3139 def int_nvvm_sust_b_3d_v4i32_trap
3140 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3141 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3142 "llvm.nvvm.sust.b.3d.v4i32.trap">,
3143 GCCBuiltin<"__nvvm_sust_b_3d_v4i32_trap">;
3147 def int_nvvm_sust_b_1d_i8_zero
3148 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty], [],
3149 "llvm.nvvm.sust.b.1d.i8.zero">,
3150 GCCBuiltin<"__nvvm_sust_b_1d_i8_zero">;
3151 def int_nvvm_sust_b_1d_i16_zero
3152 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty], [],
3153 "llvm.nvvm.sust.b.1d.i16.zero">,
3154 GCCBuiltin<"__nvvm_sust_b_1d_i16_zero">;
3155 def int_nvvm_sust_b_1d_i32_zero
3156 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
3157 "llvm.nvvm.sust.b.1d.i32.zero">,
3158 GCCBuiltin<"__nvvm_sust_b_1d_i32_zero">;
3159 def int_nvvm_sust_b_1d_i64_zero
3160 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty], [],
3161 "llvm.nvvm.sust.b.1d.i64.zero">,
3162 GCCBuiltin<"__nvvm_sust_b_1d_i64_zero">;
3163 def int_nvvm_sust_b_1d_v2i8_zero
3164 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [],
3165 "llvm.nvvm.sust.b.1d.v2i8.zero">,
3166 GCCBuiltin<"__nvvm_sust_b_1d_v2i8_zero">;
3167 def int_nvvm_sust_b_1d_v2i16_zero
3168 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [],
3169 "llvm.nvvm.sust.b.1d.v2i16.zero">,
3170 GCCBuiltin<"__nvvm_sust_b_1d_v2i16_zero">;
3171 def int_nvvm_sust_b_1d_v2i32_zero
3172 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3173 "llvm.nvvm.sust.b.1d.v2i32.zero">,
3174 GCCBuiltin<"__nvvm_sust_b_1d_v2i32_zero">;
3175 def int_nvvm_sust_b_1d_v2i64_zero
3176 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i64_ty, llvm_i64_ty], [],
3177 "llvm.nvvm.sust.b.1d.v2i64.zero">,
3178 GCCBuiltin<"__nvvm_sust_b_1d_v2i64_zero">;
3179 def int_nvvm_sust_b_1d_v4i8_zero
3180 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty,
3181 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3182 "llvm.nvvm.sust.b.1d.v4i8.zero">,
3183 GCCBuiltin<"__nvvm_sust_b_1d_v4i8_zero">;
3184 def int_nvvm_sust_b_1d_v4i16_zero
3185 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty,
3186 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3187 "llvm.nvvm.sust.b.1d.v4i16.zero">,
3188 GCCBuiltin<"__nvvm_sust_b_1d_v4i16_zero">;
3189 def int_nvvm_sust_b_1d_v4i32_zero
3190 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3191 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3192 "llvm.nvvm.sust.b.1d.v4i32.zero">,
3193 GCCBuiltin<"__nvvm_sust_b_1d_v4i32_zero">;
3196 def int_nvvm_sust_b_1d_array_i8_zero
3197 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [],
3198 "llvm.nvvm.sust.b.1d.array.i8.zero">,
3199 GCCBuiltin<"__nvvm_sust_b_1d_array_i8_zero">;
3200 def int_nvvm_sust_b_1d_array_i16_zero
3201 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [],
3202 "llvm.nvvm.sust.b.1d.array.i16.zero">,
3203 GCCBuiltin<"__nvvm_sust_b_1d_array_i16_zero">;
3204 def int_nvvm_sust_b_1d_array_i32_zero
3205 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3206 "llvm.nvvm.sust.b.1d.array.i32.zero">,
3207 GCCBuiltin<"__nvvm_sust_b_1d_array_i32_zero">;
3208 def int_nvvm_sust_b_1d_array_i64_zero
3209 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty], [],
3210 "llvm.nvvm.sust.b.1d.array.i64.zero">,
3211 GCCBuiltin<"__nvvm_sust_b_1d_array_i64_zero">;
3212 def int_nvvm_sust_b_1d_array_v2i8_zero
3213 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3214 llvm_i16_ty, llvm_i16_ty], [],
3215 "llvm.nvvm.sust.b.1d.array.v2i8.zero">,
3216 GCCBuiltin<"__nvvm_sust_b_1d_array_v2i8_zero">;
3217 def int_nvvm_sust_b_1d_array_v2i16_zero
3218 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3219 llvm_i16_ty, llvm_i16_ty], [],
3220 "llvm.nvvm.sust.b.1d.array.v2i16.zero">,
3221 GCCBuiltin<"__nvvm_sust_b_1d_array_v2i16_zero">;
3222 def int_nvvm_sust_b_1d_array_v2i32_zero
3223 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3224 llvm_i32_ty, llvm_i32_ty], [],
3225 "llvm.nvvm.sust.b.1d.array.v2i32.zero">,
3226 GCCBuiltin<"__nvvm_sust_b_1d_array_v2i32_zero">;
3227 def int_nvvm_sust_b_1d_array_v2i64_zero
3228 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3229 llvm_i64_ty, llvm_i64_ty], [],
3230 "llvm.nvvm.sust.b.1d.array.v2i64.zero">,
3231 GCCBuiltin<"__nvvm_sust_b_1d_array_v2i64_zero">;
3232 def int_nvvm_sust_b_1d_array_v4i8_zero
3233 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty,
3234 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3235 "llvm.nvvm.sust.b.1d.array.v4i8.zero">,
3236 GCCBuiltin<"__nvvm_sust_b_1d_array_v4i8_zero">;
3237 def int_nvvm_sust_b_1d_array_v4i16_zero
3238 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty,
3239 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3240 "llvm.nvvm.sust.b.1d.array.v4i16.zero">,
3241 GCCBuiltin<"__nvvm_sust_b_1d_array_v4i16_zero">;
3242 def int_nvvm_sust_b_1d_array_v4i32_zero
3243 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3244 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3245 "llvm.nvvm.sust.b.1d.array.v4i32.zero">,
3246 GCCBuiltin<"__nvvm_sust_b_1d_array_v4i32_zero">;
3249 def int_nvvm_sust_b_2d_i8_zero
3250 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [],
3251 "llvm.nvvm.sust.b.2d.i8.zero">,
3252 GCCBuiltin<"__nvvm_sust_b_2d_i8_zero">;
3253 def int_nvvm_sust_b_2d_i16_zero
3254 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [],
3255 "llvm.nvvm.sust.b.2d.i16.zero">,
3256 GCCBuiltin<"__nvvm_sust_b_2d_i16_zero">;
3257 def int_nvvm_sust_b_2d_i32_zero
3258 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3259 "llvm.nvvm.sust.b.2d.i32.zero">,
3260 GCCBuiltin<"__nvvm_sust_b_2d_i32_zero">;
3261 def int_nvvm_sust_b_2d_i64_zero
3262 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i64_ty], [],
3263 "llvm.nvvm.sust.b.2d.i64.zero">,
3264 GCCBuiltin<"__nvvm_sust_b_2d_i64_zero">;
3265 def int_nvvm_sust_b_2d_v2i8_zero
3266 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3267 llvm_i16_ty, llvm_i16_ty], [],
3268 "llvm.nvvm.sust.b.2d.v2i8.zero">,
3269 GCCBuiltin<"__nvvm_sust_b_2d_v2i8_zero">;
3270 def int_nvvm_sust_b_2d_v2i16_zero
3271 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3272 llvm_i16_ty, llvm_i16_ty], [],
3273 "llvm.nvvm.sust.b.2d.v2i16.zero">,
3274 GCCBuiltin<"__nvvm_sust_b_2d_v2i16_zero">;
3275 def int_nvvm_sust_b_2d_v2i32_zero
3276 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3277 llvm_i32_ty, llvm_i32_ty], [],
3278 "llvm.nvvm.sust.b.2d.v2i32.zero">,
3279 GCCBuiltin<"__nvvm_sust_b_2d_v2i32_zero">;
3280 def int_nvvm_sust_b_2d_v2i64_zero
3281 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3282 llvm_i64_ty, llvm_i64_ty], [],
3283 "llvm.nvvm.sust.b.2d.v2i64.zero">,
3284 GCCBuiltin<"__nvvm_sust_b_2d_v2i64_zero">;
3285 def int_nvvm_sust_b_2d_v4i8_zero
3286 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty,
3287 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3288 "llvm.nvvm.sust.b.2d.v4i8.zero">,
3289 GCCBuiltin<"__nvvm_sust_b_2d_v4i8_zero">;
3290 def int_nvvm_sust_b_2d_v4i16_zero
3291 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty,
3292 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3293 "llvm.nvvm.sust.b.2d.v4i16.zero">,
3294 GCCBuiltin<"__nvvm_sust_b_2d_v4i16_zero">;
3295 def int_nvvm_sust_b_2d_v4i32_zero
3296 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3297 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3298 "llvm.nvvm.sust.b.2d.v4i32.zero">,
3299 GCCBuiltin<"__nvvm_sust_b_2d_v4i32_zero">;
3302 def int_nvvm_sust_b_2d_array_i8_zero
3303 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3304 llvm_i32_ty, llvm_i16_ty], [],
3305 "llvm.nvvm.sust.b.2d.array.i8.zero">,
3306 GCCBuiltin<"__nvvm_sust_b_2d_array_i8_zero">;
3307 def int_nvvm_sust_b_2d_array_i16_zero
3308 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3309 llvm_i32_ty, llvm_i16_ty], [],
3310 "llvm.nvvm.sust.b.2d.array.i16.zero">,
3311 GCCBuiltin<"__nvvm_sust_b_2d_array_i16_zero">;
3312 def int_nvvm_sust_b_2d_array_i32_zero
3313 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3314 llvm_i32_ty, llvm_i32_ty], [],
3315 "llvm.nvvm.sust.b.2d.array.i32.zero">,
3316 GCCBuiltin<"__nvvm_sust_b_2d_array_i32_zero">;
3317 def int_nvvm_sust_b_2d_array_i64_zero
3318 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3319 llvm_i32_ty, llvm_i64_ty], [],
3320 "llvm.nvvm.sust.b.2d.array.i64.zero">,
3321 GCCBuiltin<"__nvvm_sust_b_2d_array_i64_zero">;
3322 def int_nvvm_sust_b_2d_array_v2i8_zero
3323 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3324 llvm_i16_ty, llvm_i16_ty], [],
3325 "llvm.nvvm.sust.b.2d.array.v2i8.zero">,
3326 GCCBuiltin<"__nvvm_sust_b_2d_array_v2i8_zero">;
3327 def int_nvvm_sust_b_2d_array_v2i16_zero
3328 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3329 llvm_i16_ty, llvm_i16_ty], [],
3330 "llvm.nvvm.sust.b.2d.array.v2i16.zero">,
3331 GCCBuiltin<"__nvvm_sust_b_2d_array_v2i16_zero">;
3332 def int_nvvm_sust_b_2d_array_v2i32_zero
3333 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3334 llvm_i32_ty, llvm_i32_ty], [],
3335 "llvm.nvvm.sust.b.2d.array.v2i32.zero">,
3336 GCCBuiltin<"__nvvm_sust_b_2d_array_v2i32_zero">;
3337 def int_nvvm_sust_b_2d_array_v2i64_zero
3338 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3339 llvm_i64_ty, llvm_i64_ty], [],
3340 "llvm.nvvm.sust.b.2d.array.v2i64.zero">,
3341 GCCBuiltin<"__nvvm_sust_b_2d_array_v2i64_zero">;
3342 def int_nvvm_sust_b_2d_array_v4i8_zero
3343 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3344 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3345 "llvm.nvvm.sust.b.2d.array.v4i8.zero">,
3346 GCCBuiltin<"__nvvm_sust_b_2d_array_v4i8_zero">;
3347 def int_nvvm_sust_b_2d_array_v4i16_zero
3348 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3349 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3350 "llvm.nvvm.sust.b.2d.array.v4i16.zero">,
3351 GCCBuiltin<"__nvvm_sust_b_2d_array_v4i16_zero">;
3352 def int_nvvm_sust_b_2d_array_v4i32_zero
3353 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3354 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3355 "llvm.nvvm.sust.b.2d.array.v4i32.zero">,
3356 GCCBuiltin<"__nvvm_sust_b_2d_array_v4i32_zero">;
3359 def int_nvvm_sust_b_3d_i8_zero
3360 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3361 llvm_i32_ty, llvm_i16_ty], [],
3362 "llvm.nvvm.sust.b.3d.i8.zero">,
3363 GCCBuiltin<"__nvvm_sust_b_3d_i8_zero">;
3364 def int_nvvm_sust_b_3d_i16_zero
3365 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3366 llvm_i32_ty, llvm_i16_ty], [],
3367 "llvm.nvvm.sust.b.3d.i16.zero">,
3368 GCCBuiltin<"__nvvm_sust_b_3d_i16_zero">;
3369 def int_nvvm_sust_b_3d_i32_zero
3370 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3371 llvm_i32_ty, llvm_i32_ty], [],
3372 "llvm.nvvm.sust.b.3d.i32.zero">,
3373 GCCBuiltin<"__nvvm_sust_b_3d_i32_zero">;
3374 def int_nvvm_sust_b_3d_i64_zero
3375 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3376 llvm_i32_ty, llvm_i64_ty], [],
3377 "llvm.nvvm.sust.b.3d.i64.zero">,
3378 GCCBuiltin<"__nvvm_sust_b_3d_i64_zero">;
3379 def int_nvvm_sust_b_3d_v2i8_zero
3380 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3381 llvm_i16_ty, llvm_i16_ty], [],
3382 "llvm.nvvm.sust.b.3d.v2i8.zero">,
3383 GCCBuiltin<"__nvvm_sust_b_3d_v2i8_zero">;
3384 def int_nvvm_sust_b_3d_v2i16_zero
3385 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3386 llvm_i16_ty, llvm_i16_ty], [],
3387 "llvm.nvvm.sust.b.3d.v2i16.zero">,
3388 GCCBuiltin<"__nvvm_sust_b_3d_v2i16_zero">;
3389 def int_nvvm_sust_b_3d_v2i32_zero
3390 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3391 llvm_i32_ty, llvm_i32_ty], [],
3392 "llvm.nvvm.sust.b.3d.v2i32.zero">,
3393 GCCBuiltin<"__nvvm_sust_b_3d_v2i32_zero">;
3394 def int_nvvm_sust_b_3d_v2i64_zero
3395 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3396 llvm_i64_ty, llvm_i64_ty], [],
3397 "llvm.nvvm.sust.b.3d.v2i64.zero">,
3398 GCCBuiltin<"__nvvm_sust_b_3d_v2i64_zero">;
3399 def int_nvvm_sust_b_3d_v4i8_zero
3400 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3401 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3402 "llvm.nvvm.sust.b.3d.v4i8.zero">,
3403 GCCBuiltin<"__nvvm_sust_b_3d_v4i8_zero">;
3404 def int_nvvm_sust_b_3d_v4i16_zero
3405 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3406 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3407 "llvm.nvvm.sust.b.3d.v4i16.zero">,
3408 GCCBuiltin<"__nvvm_sust_b_3d_v4i16_zero">;
3409 def int_nvvm_sust_b_3d_v4i32_zero
3410 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3411 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3412 "llvm.nvvm.sust.b.3d.v4i32.zero">,
3413 GCCBuiltin<"__nvvm_sust_b_3d_v4i32_zero">;
3419 def int_nvvm_sust_p_1d_i8_trap
3420 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty], [],
3421 "llvm.nvvm.sust.p.1d.i8.trap">,
3422 GCCBuiltin<"__nvvm_sust_p_1d_i8_trap">;
3423 def int_nvvm_sust_p_1d_i16_trap
3424 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty], [],
3425 "llvm.nvvm.sust.p.1d.i16.trap">,
3426 GCCBuiltin<"__nvvm_sust_p_1d_i16_trap">;
3427 def int_nvvm_sust_p_1d_i32_trap
3428 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty], [],
3429 "llvm.nvvm.sust.p.1d.i32.trap">,
3430 GCCBuiltin<"__nvvm_sust_p_1d_i32_trap">;
3431 def int_nvvm_sust_p_1d_v2i8_trap
3432 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [],
3433 "llvm.nvvm.sust.p.1d.v2i8.trap">,
3434 GCCBuiltin<"__nvvm_sust_p_1d_v2i8_trap">;
3435 def int_nvvm_sust_p_1d_v2i16_trap
3436 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty, llvm_i16_ty], [],
3437 "llvm.nvvm.sust.p.1d.v2i16.trap">,
3438 GCCBuiltin<"__nvvm_sust_p_1d_v2i16_trap">;
3439 def int_nvvm_sust_p_1d_v2i32_trap
3440 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3441 "llvm.nvvm.sust.p.1d.v2i32.trap">,
3442 GCCBuiltin<"__nvvm_sust_p_1d_v2i32_trap">;
3443 def int_nvvm_sust_p_1d_v4i8_trap
3444 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty,
3445 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3446 "llvm.nvvm.sust.p.1d.v4i8.trap">,
3447 GCCBuiltin<"__nvvm_sust_p_1d_v4i8_trap">;
3448 def int_nvvm_sust_p_1d_v4i16_trap
3449 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i16_ty,
3450 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3451 "llvm.nvvm.sust.p.1d.v4i16.trap">,
3452 GCCBuiltin<"__nvvm_sust_p_1d_v4i16_trap">;
3453 def int_nvvm_sust_p_1d_v4i32_trap
3454 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3455 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3456 "llvm.nvvm.sust.p.1d.v4i32.trap">,
3457 GCCBuiltin<"__nvvm_sust_p_1d_v4i32_trap">;
3460 def int_nvvm_sust_p_1d_array_i8_trap
3461 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [],
3462 "llvm.nvvm.sust.p.1d.array.i8.trap">,
3463 GCCBuiltin<"__nvvm_sust_p_1d_array_i8_trap">;
3464 def int_nvvm_sust_p_1d_array_i16_trap
3465 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [],
3466 "llvm.nvvm.sust.p.1d.array.i16.trap">,
3467 GCCBuiltin<"__nvvm_sust_p_1d_array_i16_trap">;
3468 def int_nvvm_sust_p_1d_array_i32_trap
3469 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3470 "llvm.nvvm.sust.p.1d.array.i32.trap">,
3471 GCCBuiltin<"__nvvm_sust_p_1d_array_i32_trap">;
3472 def int_nvvm_sust_p_1d_array_v2i8_trap
3473 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3474 llvm_i16_ty, llvm_i16_ty], [],
3475 "llvm.nvvm.sust.p.1d.array.v2i8.trap">,
3476 GCCBuiltin<"__nvvm_sust_p_1d_array_v2i8_trap">;
3477 def int_nvvm_sust_p_1d_array_v2i16_trap
3478 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3479 llvm_i16_ty, llvm_i16_ty], [],
3480 "llvm.nvvm.sust.p.1d.array.v2i16.trap">,
3481 GCCBuiltin<"__nvvm_sust_p_1d_array_v2i16_trap">;
3482 def int_nvvm_sust_p_1d_array_v2i32_trap
3483 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3484 llvm_i32_ty, llvm_i32_ty], [],
3485 "llvm.nvvm.sust.p.1d.array.v2i32.trap">,
3486 GCCBuiltin<"__nvvm_sust_p_1d_array_v2i32_trap">;
3487 def int_nvvm_sust_p_1d_array_v4i8_trap
3488 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty,
3489 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3490 "llvm.nvvm.sust.p.1d.array.v4i8.trap">,
3491 GCCBuiltin<"__nvvm_sust_p_1d_array_v4i8_trap">;
3492 def int_nvvm_sust_p_1d_array_v4i16_trap
3493 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty,
3494 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3495 "llvm.nvvm.sust.p.1d.array.v4i16.trap">,
3496 GCCBuiltin<"__nvvm_sust_p_1d_array_v4i16_trap">;
3497 def int_nvvm_sust_p_1d_array_v4i32_trap
3498 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3499 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3500 "llvm.nvvm.sust.p.1d.array.v4i32.trap">,
3501 GCCBuiltin<"__nvvm_sust_p_1d_array_v4i32_trap">;
3504 def int_nvvm_sust_p_2d_i8_trap
3505 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [],
3506 "llvm.nvvm.sust.p.2d.i8.trap">,
3507 GCCBuiltin<"__nvvm_sust_p_2d_i8_trap">;
3508 def int_nvvm_sust_p_2d_i16_trap
3509 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty], [],
3510 "llvm.nvvm.sust.p.2d.i16.trap">,
3511 GCCBuiltin<"__nvvm_sust_p_2d_i16_trap">;
3512 def int_nvvm_sust_p_2d_i32_trap
3513 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3514 "llvm.nvvm.sust.p.2d.i32.trap">,
3515 GCCBuiltin<"__nvvm_sust_p_2d_i32_trap">;
3516 def int_nvvm_sust_p_2d_v2i8_trap
3517 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3518 llvm_i16_ty, llvm_i16_ty], [],
3519 "llvm.nvvm.sust.p.2d.v2i8.trap">,
3520 GCCBuiltin<"__nvvm_sust_p_2d_v2i8_trap">;
3521 def int_nvvm_sust_p_2d_v2i16_trap
3522 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3523 llvm_i16_ty, llvm_i16_ty], [],
3524 "llvm.nvvm.sust.p.2d.v2i16.trap">,
3525 GCCBuiltin<"__nvvm_sust_p_2d_v2i16_trap">;
3526 def int_nvvm_sust_p_2d_v2i32_trap
3527 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3528 llvm_i32_ty, llvm_i32_ty], [],
3529 "llvm.nvvm.sust.p.2d.v2i32.trap">,
3530 GCCBuiltin<"__nvvm_sust_p_2d_v2i32_trap">;
3531 def int_nvvm_sust_p_2d_v4i8_trap
3532 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty,
3533 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3534 "llvm.nvvm.sust.p.2d.v4i8.trap">,
3535 GCCBuiltin<"__nvvm_sust_p_2d_v4i8_trap">;
3536 def int_nvvm_sust_p_2d_v4i16_trap
3537 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i16_ty,
3538 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3539 "llvm.nvvm.sust.p.2d.v4i16.trap">,
3540 GCCBuiltin<"__nvvm_sust_p_2d_v4i16_trap">;
3541 def int_nvvm_sust_p_2d_v4i32_trap
3542 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3543 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3544 "llvm.nvvm.sust.p.2d.v4i32.trap">,
3545 GCCBuiltin<"__nvvm_sust_p_2d_v4i32_trap">;
3548 def int_nvvm_sust_p_2d_array_i8_trap
3549 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3550 llvm_i32_ty, llvm_i16_ty], [],
3551 "llvm.nvvm.sust.p.2d.array.i8.trap">,
3552 GCCBuiltin<"__nvvm_sust_p_2d_array_i8_trap">;
3553 def int_nvvm_sust_p_2d_array_i16_trap
3554 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3555 llvm_i32_ty, llvm_i16_ty], [],
3556 "llvm.nvvm.sust.p.2d.array.i16.trap">,
3557 GCCBuiltin<"__nvvm_sust_p_2d_array_i16_trap">;
3558 def int_nvvm_sust_p_2d_array_i32_trap
3559 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3560 llvm_i32_ty, llvm_i32_ty], [],
3561 "llvm.nvvm.sust.p.2d.array.i32.trap">,
3562 GCCBuiltin<"__nvvm_sust_p_2d_array_i32_trap">;
3563 def int_nvvm_sust_p_2d_array_v2i8_trap
3564 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3565 llvm_i16_ty, llvm_i16_ty], [],
3566 "llvm.nvvm.sust.p.2d.array.v2i8.trap">,
3567 GCCBuiltin<"__nvvm_sust_p_2d_array_v2i8_trap">;
3568 def int_nvvm_sust_p_2d_array_v2i16_trap
3569 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3570 llvm_i16_ty, llvm_i16_ty], [],
3571 "llvm.nvvm.sust.p.2d.array.v2i16.trap">,
3572 GCCBuiltin<"__nvvm_sust_p_2d_array_v2i16_trap">;
3573 def int_nvvm_sust_p_2d_array_v2i32_trap
3574 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3575 llvm_i32_ty, llvm_i32_ty], [],
3576 "llvm.nvvm.sust.p.2d.array.v2i32.trap">,
3577 GCCBuiltin<"__nvvm_sust_p_2d_array_v2i32_trap">;
3578 def int_nvvm_sust_p_2d_array_v4i8_trap
3579 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3580 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3581 "llvm.nvvm.sust.p.2d.array.v4i8.trap">,
3582 GCCBuiltin<"__nvvm_sust_p_2d_array_v4i8_trap">;
3583 def int_nvvm_sust_p_2d_array_v4i16_trap
3584 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3585 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3586 "llvm.nvvm.sust.p.2d.array.v4i16.trap">,
3587 GCCBuiltin<"__nvvm_sust_p_2d_array_v4i16_trap">;
3588 def int_nvvm_sust_p_2d_array_v4i32_trap
3589 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3590 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3591 "llvm.nvvm.sust.p.2d.array.v4i32.trap">,
3592 GCCBuiltin<"__nvvm_sust_p_2d_array_v4i32_trap">;
3595 def int_nvvm_sust_p_3d_i8_trap
3596 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3597 llvm_i32_ty, llvm_i16_ty], [],
3598 "llvm.nvvm.sust.p.3d.i8.trap">,
3599 GCCBuiltin<"__nvvm_sust_p_3d_i8_trap">;
3600 def int_nvvm_sust_p_3d_i16_trap
3601 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3602 llvm_i32_ty, llvm_i16_ty], [],
3603 "llvm.nvvm.sust.p.3d.i16.trap">,
3604 GCCBuiltin<"__nvvm_sust_p_3d_i16_trap">;
3605 def int_nvvm_sust_p_3d_i32_trap
3606 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty,
3607 llvm_i32_ty, llvm_i32_ty], [],
3608 "llvm.nvvm.sust.p.3d.i32.trap">,
3609 GCCBuiltin<"__nvvm_sust_p_3d_i32_trap">;
3610 def int_nvvm_sust_p_3d_v2i8_trap
3611 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3612 llvm_i16_ty, llvm_i16_ty], [],
3613 "llvm.nvvm.sust.p.3d.v2i8.trap">,
3614 GCCBuiltin<"__nvvm_sust_p_3d_v2i8_trap">;
3615 def int_nvvm_sust_p_3d_v2i16_trap
3616 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3617 llvm_i16_ty, llvm_i16_ty], [],
3618 "llvm.nvvm.sust.p.3d.v2i16.trap">,
3619 GCCBuiltin<"__nvvm_sust_p_3d_v2i16_trap">;
3620 def int_nvvm_sust_p_3d_v2i32_trap
3621 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3622 llvm_i32_ty, llvm_i32_ty], [],
3623 "llvm.nvvm.sust.p.3d.v2i32.trap">,
3624 GCCBuiltin<"__nvvm_sust_p_3d_v2i32_trap">;
3625 def int_nvvm_sust_p_3d_v4i8_trap
3626 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3627 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3628 "llvm.nvvm.sust.p.3d.v4i8.trap">,
3629 GCCBuiltin<"__nvvm_sust_p_3d_v4i8_trap">;
3630 def int_nvvm_sust_p_3d_v4i16_trap
3631 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3632 llvm_i16_ty, llvm_i16_ty, llvm_i16_ty, llvm_i16_ty], [],
3633 "llvm.nvvm.sust.p.3d.v4i16.trap">,
3634 GCCBuiltin<"__nvvm_sust_p_3d_v4i16_trap">;
3635 def int_nvvm_sust_p_3d_v4i32_trap
3636 : Intrinsic<[], [llvm_i64_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty,
3637 llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty], [],
3638 "llvm.nvvm.sust.p.3d.v4i32.trap">,
3639 GCCBuiltin<"__nvvm_sust_p_3d_v4i32_trap">;
3642 def int_nvvm_rotate_b32
3643 : Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty],
3644 [IntrNoMem], "llvm.nvvm.rotate.b32">,
3645 GCCBuiltin<"__nvvm_rotate_b32">;
3647 def int_nvvm_rotate_b64
3648 :Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty],
3649 [IntrNoMem], "llvm.nvvm.rotate.b64">,
3650 GCCBuiltin<"__nvvm_rotate_b64">;
3652 def int_nvvm_rotate_right_b64
3653 : Intrinsic<[llvm_i64_ty], [llvm_i64_ty, llvm_i32_ty],
3654 [IntrNoMem], "llvm.nvvm.rotate.right.b64">,
3655 GCCBuiltin<"__nvvm_rotate_right_b64">;
3657 def int_nvvm_swap_lo_hi_b64
3658 : Intrinsic<[llvm_i64_ty], [llvm_i64_ty],
3659 [IntrNoMem], "llvm.nvvm.swap.lo.hi.b64">,
3660 GCCBuiltin<"__nvvm_swap_lo_hi_b64">;
3663 // Accessing special registers.
3664 multiclass PTXReadSRegIntrinsic_v4i32<string regname> {
3665 // FIXME: Do we need the 128-bit integer type version?
3666 // def _r64 : Intrinsic<[llvm_i128_ty], [], [IntrNoMem]>;
3668 // FIXME: Enable this once v4i32 support is enabled in back-end.
3669 // def _v4i16 : Intrinsic<[llvm_v4i32_ty], [], [IntrNoMem]>;
3671 def _x : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
3672 GCCBuiltin<"__nvvm_read_ptx_sreg_" # regname # "_x">;
3673 def _y : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
3674 GCCBuiltin<"__nvvm_read_ptx_sreg_" # regname # "_y">;
3675 def _z : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
3676 GCCBuiltin<"__nvvm_read_ptx_sreg_" # regname # "_z">;
3677 def _w : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
3678 GCCBuiltin<"__nvvm_read_ptx_sreg_" # regname # "_w">;
3681 class PTXReadSRegIntrinsic_r32<string name>
3682 : Intrinsic<[llvm_i32_ty], [], [IntrNoMem]>,
3683 GCCBuiltin<"__nvvm_read_ptx_sreg_" # name>;
3685 class PTXReadSRegIntrinsic_r64<string name>
3686 : Intrinsic<[llvm_i64_ty], [], [IntrNoMem]>,
3687 GCCBuiltin<"__nvvm_read_ptx_sreg_" # name>;
3689 defm int_nvvm_read_ptx_sreg_tid : PTXReadSRegIntrinsic_v4i32<"tid">;
3690 defm int_nvvm_read_ptx_sreg_ntid : PTXReadSRegIntrinsic_v4i32<"ntid">;
3692 def int_nvvm_read_ptx_sreg_laneid : PTXReadSRegIntrinsic_r32<"laneid">;
3693 def int_nvvm_read_ptx_sreg_warpid : PTXReadSRegIntrinsic_r32<"warpid">;
3694 def int_nvvm_read_ptx_sreg_nwarpid : PTXReadSRegIntrinsic_r32<"nwarpid">;
3696 defm int_nvvm_read_ptx_sreg_ctaid : PTXReadSRegIntrinsic_v4i32<"ctaid">;
3697 defm int_nvvm_read_ptx_sreg_nctaid : PTXReadSRegIntrinsic_v4i32<"nctaid">;
3699 def int_nvvm_read_ptx_sreg_smid : PTXReadSRegIntrinsic_r32<"smid">;
3700 def int_nvvm_read_ptx_sreg_nsmid : PTXReadSRegIntrinsic_r32<"nsmid">;
3701 def int_nvvm_read_ptx_sreg_gridid : PTXReadSRegIntrinsic_r32<"gridid">;
3703 def int_nvvm_read_ptx_sreg_lanemask_eq :
3704 PTXReadSRegIntrinsic_r32<"lanemask_eq">;
3705 def int_nvvm_read_ptx_sreg_lanemask_le :
3706 PTXReadSRegIntrinsic_r32<"lanemask_le">;
3707 def int_nvvm_read_ptx_sreg_lanemask_lt :
3708 PTXReadSRegIntrinsic_r32<"lanemask_lt">;
3709 def int_nvvm_read_ptx_sreg_lanemask_ge :
3710 PTXReadSRegIntrinsic_r32<"lanemask_ge">;
3711 def int_nvvm_read_ptx_sreg_lanemask_gt :
3712 PTXReadSRegIntrinsic_r32<"lanemask_gt">;
3714 def int_nvvm_read_ptx_sreg_clock : PTXReadSRegIntrinsic_r32<"clock">;
3715 def int_nvvm_read_ptx_sreg_clock64 : PTXReadSRegIntrinsic_r64<"clock64">;
3717 def int_nvvm_read_ptx_sreg_pm0 : PTXReadSRegIntrinsic_r32<"pm0">;
3718 def int_nvvm_read_ptx_sreg_pm1 : PTXReadSRegIntrinsic_r32<"pm1">;
3719 def int_nvvm_read_ptx_sreg_pm2 : PTXReadSRegIntrinsic_r32<"pm2">;
3720 def int_nvvm_read_ptx_sreg_pm3 : PTXReadSRegIntrinsic_r32<"pm3">;
3722 def int_nvvm_read_ptx_sreg_warpsize : PTXReadSRegIntrinsic_r32<"warpsize">;
3728 // shfl.down.b32 dest, val, offset, mask_and_clamp
3729 def int_nvvm_shfl_down_i32 :
3730 Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
3731 [IntrNoMem, IntrConvergent], "llvm.nvvm.shfl.down.i32">,
3732 GCCBuiltin<"__nvvm_shfl_down_i32">;
3733 def int_nvvm_shfl_down_f32 :
3734 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty],
3735 [IntrNoMem, IntrConvergent], "llvm.nvvm.shfl.down.f32">,
3736 GCCBuiltin<"__nvvm_shfl_down_f32">;
3738 // shfl.up.b32 dest, val, offset, mask_and_clamp
3739 def int_nvvm_shfl_up_i32 :
3740 Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
3741 [IntrNoMem, IntrConvergent], "llvm.nvvm.shfl.up.i32">,
3742 GCCBuiltin<"__nvvm_shfl_up_i32">;
3743 def int_nvvm_shfl_up_f32 :
3744 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty],
3745 [IntrNoMem, IntrConvergent], "llvm.nvvm.shfl.up.f32">,
3746 GCCBuiltin<"__nvvm_shfl_up_f32">;
3748 // shfl.bfly.b32 dest, val, offset, mask_and_clamp
3749 def int_nvvm_shfl_bfly_i32 :
3750 Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
3751 [IntrNoMem, IntrConvergent], "llvm.nvvm.shfl.bfly.i32">,
3752 GCCBuiltin<"__nvvm_shfl_bfly_i32">;
3753 def int_nvvm_shfl_bfly_f32 :
3754 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty],
3755 [IntrNoMem, IntrConvergent], "llvm.nvvm.shfl.bfly.f32">,
3756 GCCBuiltin<"__nvvm_shfl_bfly_f32">;
3758 // shfl.idx.b32 dest, val, lane, mask_and_clamp
3759 def int_nvvm_shfl_idx_i32 :
3760 Intrinsic<[llvm_i32_ty], [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
3761 [IntrNoMem, IntrConvergent], "llvm.nvvm.shfl.idx.i32">,
3762 GCCBuiltin<"__nvvm_shfl_idx_i32">;
3763 def int_nvvm_shfl_idx_f32 :
3764 Intrinsic<[llvm_float_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty],
3765 [IntrNoMem, IntrConvergent], "llvm.nvvm.shfl.idx.f32">,
3766 GCCBuiltin<"__nvvm_shfl_idx_f32">;