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