1 //===- NVPTXIntrinsics.td - PTX Intrinsics Instructions -------*- tblgen -*-==//
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 def immFloat0 : PatLeaf<(fpimm), [{
11 float f = (float)N->getValueAPF().convertToFloat();
15 def immFloat1 : PatLeaf<(fpimm), [{
16 float f = (float)N->getValueAPF().convertToFloat();
20 def immDouble0 : PatLeaf<(fpimm), [{
21 double d = (double)N->getValueAPF().convertToDouble();
25 def immDouble1 : PatLeaf<(fpimm), [{
26 double d = (double)N->getValueAPF().convertToDouble();
32 //-----------------------------------
33 // Synchronization Functions
34 //-----------------------------------
35 def INT_CUDA_SYNCTHREADS : NVPTXInst<(outs), (ins),
37 [(int_cuda_syncthreads)]>;
38 def INT_BARRIER0 : NVPTXInst<(outs), (ins),
40 [(int_nvvm_barrier0)]>;
41 def INT_BARRIER0_POPC : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred),
43 !strconcat(".reg .pred \t%p1; \n\t",
44 !strconcat("setp.ne.u32 \t%p1, $pred, 0; \n\t",
45 !strconcat("bar.red.popc.u32 \t$dst, 0, %p1; \n\t",
46 !strconcat("}}", ""))))),
47 [(set Int32Regs:$dst, (int_nvvm_barrier0_popc Int32Regs:$pred))]>;
48 def INT_BARRIER0_AND : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred),
50 !strconcat(".reg .pred \t%p1; \n\t",
51 !strconcat(".reg .pred \t%p2; \n\t",
52 !strconcat("setp.ne.u32 \t%p1, $pred, 0; \n\t",
53 !strconcat("bar.red.and.pred \t%p2, 0, %p1; \n\t",
54 !strconcat("selp.u32 \t$dst, 1, 0, %p2; \n\t",
55 !strconcat("}}", ""))))))),
56 [(set Int32Regs:$dst, (int_nvvm_barrier0_and Int32Regs:$pred))]>;
57 def INT_BARRIER0_OR : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred),
59 !strconcat(".reg .pred \t%p1; \n\t",
60 !strconcat(".reg .pred \t%p2; \n\t",
61 !strconcat("setp.ne.u32 \t%p1, $pred, 0; \n\t",
62 !strconcat("bar.red.or.pred \t%p2, 0, %p1; \n\t",
63 !strconcat("selp.u32 \t$dst, 1, 0, %p2; \n\t",
64 !strconcat("}}", ""))))))),
65 [(set Int32Regs:$dst, (int_nvvm_barrier0_or Int32Regs:$pred))]>;
68 //-----------------------------------
69 // Explicit Memory Fence Functions
70 //-----------------------------------
71 class MEMBAR<string StrOp, Intrinsic IntOP> :
72 NVPTXInst<(outs), (ins),
75 def INT_MEMBAR_CTA : MEMBAR<"membar.cta;", int_nvvm_membar_cta>;
76 def INT_MEMBAR_GL : MEMBAR<"membar.gl;", int_nvvm_membar_gl>;
77 def INT_MEMBAR_SYS : MEMBAR<"membar.sys;", int_nvvm_membar_sys>;
80 //-----------------------------------
82 //-----------------------------------
84 // Map min(1.0, max(0.0, x)) to sat(x)
85 // Note that max(0.0, min(x, 1.0)) cannot be mapped to sat(x) because when x is
87 // max(0.0, min(x, 1.0)) is 1.0 while sat(x) is 0.
88 // Same story for fmax, fmin.
90 def : Pat<(int_nvvm_fmin_f immFloat1,
91 (int_nvvm_fmax_f immFloat0, Float32Regs:$a)),
92 (CVT_f32_f32 Float32Regs:$a, CvtSAT)>;
93 def : Pat<(int_nvvm_fmin_f immFloat1,
94 (int_nvvm_fmax_f Float32Regs:$a, immFloat0)),
95 (CVT_f32_f32 Float32Regs:$a, CvtSAT)>;
96 def : Pat<(int_nvvm_fmin_f
97 (int_nvvm_fmax_f immFloat0, Float32Regs:$a), immFloat1),
98 (CVT_f32_f32 Float32Regs:$a, CvtSAT)>;
99 def : Pat<(int_nvvm_fmin_f
100 (int_nvvm_fmax_f Float32Regs:$a, immFloat0), immFloat1),
101 (CVT_f32_f32 Float32Regs:$a, CvtSAT)>;
103 def : Pat<(int_nvvm_fmin_d immDouble1,
104 (int_nvvm_fmax_d immDouble0, Float64Regs:$a)),
105 (CVT_f64_f64 Float64Regs:$a, CvtSAT)>;
106 def : Pat<(int_nvvm_fmin_d immDouble1,
107 (int_nvvm_fmax_d Float64Regs:$a, immDouble0)),
108 (CVT_f64_f64 Float64Regs:$a, CvtSAT)>;
109 def : Pat<(int_nvvm_fmin_d
110 (int_nvvm_fmax_d immDouble0, Float64Regs:$a), immDouble1),
111 (CVT_f64_f64 Float64Regs:$a, CvtSAT)>;
112 def : Pat<(int_nvvm_fmin_d
113 (int_nvvm_fmax_d Float64Regs:$a, immDouble0), immDouble1),
114 (CVT_f64_f64 Float64Regs:$a, CvtSAT)>;
117 // We need a full string for OpcStr here because we need to deal with case like
119 class F_MATH_1<string OpcStr, NVPTXRegClass target_regclass,
120 NVPTXRegClass src_regclass, Intrinsic IntOP>
121 : NVPTXInst<(outs target_regclass:$dst), (ins src_regclass:$src0),
123 [(set target_regclass:$dst, (IntOP src_regclass:$src0))]>;
125 // We need a full string for OpcStr here because we need to deal with the case
126 // like INT_PTX_NATIVE_POWR_F.
127 class F_MATH_2<string OpcStr, NVPTXRegClass t_regclass,
128 NVPTXRegClass s0_regclass, NVPTXRegClass s1_regclass, Intrinsic IntOP>
129 : NVPTXInst<(outs t_regclass:$dst),
130 (ins s0_regclass:$src0, s1_regclass:$src1),
132 [(set t_regclass:$dst, (IntOP s0_regclass:$src0, s1_regclass:$src1))]>;
134 class F_MATH_3<string OpcStr, NVPTXRegClass t_regclass,
135 NVPTXRegClass s0_regclass, NVPTXRegClass s1_regclass,
136 NVPTXRegClass s2_regclass, Intrinsic IntOP>
137 : NVPTXInst<(outs t_regclass:$dst),
138 (ins s0_regclass:$src0, s1_regclass:$src1, s2_regclass:$src2),
140 [(set t_regclass:$dst,
141 (IntOP s0_regclass:$src0, s1_regclass:$src1, s2_regclass:$src2))]>;
147 def INT_NVVM_CLZ_I : F_MATH_1<"clz.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
149 def INT_NVVM_CLZ_LL : F_MATH_1<"clz.b64 \t$dst, $src0;", Int32Regs, Int64Regs,
152 def INT_NVVM_POPC_I : F_MATH_1<"popc.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
154 def INT_NVVM_POPC_LL : F_MATH_1<"popc.b64 \t$dst, $src0;", Int32Regs, Int64Regs,
157 def INT_NVVM_PRMT : F_MATH_3<"prmt.b32 \t$dst, $src0, $src1, $src2;", Int32Regs,
158 Int32Regs, Int32Regs, Int32Regs, int_nvvm_prmt>;
164 def INT_NVVM_MIN_I : F_MATH_2<"min.s32 \t$dst, $src0, $src1;", Int32Regs,
165 Int32Regs, Int32Regs, int_nvvm_min_i>;
166 def INT_NVVM_MIN_UI : F_MATH_2<"min.u32 \t$dst, $src0, $src1;", Int32Regs,
167 Int32Regs, Int32Regs, int_nvvm_min_ui>;
169 def INT_NVVM_MIN_LL : F_MATH_2<"min.s64 \t$dst, $src0, $src1;", Int64Regs,
170 Int64Regs, Int64Regs, int_nvvm_min_ll>;
171 def INT_NVVM_MIN_ULL : F_MATH_2<"min.u64 \t$dst, $src0, $src1;", Int64Regs,
172 Int64Regs, Int64Regs, int_nvvm_min_ull>;
174 def INT_NVVM_MAX_I : F_MATH_2<"max.s32 \t$dst, $src0, $src1;", Int32Regs,
175 Int32Regs, Int32Regs, int_nvvm_max_i>;
176 def INT_NVVM_MAX_UI : F_MATH_2<"max.u32 \t$dst, $src0, $src1;", Int32Regs,
177 Int32Regs, Int32Regs, int_nvvm_max_ui>;
179 def INT_NVVM_MAX_LL : F_MATH_2<"max.s64 \t$dst, $src0, $src1;", Int64Regs,
180 Int64Regs, Int64Regs, int_nvvm_max_ll>;
181 def INT_NVVM_MAX_ULL : F_MATH_2<"max.u64 \t$dst, $src0, $src1;", Int64Regs,
182 Int64Regs, Int64Regs, int_nvvm_max_ull>;
184 def INT_NVVM_FMIN_F : F_MATH_2<"min.f32 \t$dst, $src0, $src1;", Float32Regs,
185 Float32Regs, Float32Regs, int_nvvm_fmin_f>;
186 def INT_NVVM_FMIN_FTZ_F : F_MATH_2<"min.ftz.f32 \t$dst, $src0, $src1;",
187 Float32Regs, Float32Regs, Float32Regs, int_nvvm_fmin_ftz_f>;
189 def INT_NVVM_FMAX_F : F_MATH_2<"max.f32 \t$dst, $src0, $src1;", Float32Regs,
190 Float32Regs, Float32Regs, int_nvvm_fmax_f>;
191 def INT_NVVM_FMAX_FTZ_F : F_MATH_2<"max.ftz.f32 \t$dst, $src0, $src1;",
192 Float32Regs, Float32Regs, Float32Regs, int_nvvm_fmax_ftz_f>;
194 def INT_NVVM_FMIN_D : F_MATH_2<"min.f64 \t$dst, $src0, $src1;", Float64Regs,
195 Float64Regs, Float64Regs, int_nvvm_fmin_d>;
196 def INT_NVVM_FMAX_D : F_MATH_2<"max.f64 \t$dst, $src0, $src1;", Float64Regs,
197 Float64Regs, Float64Regs, int_nvvm_fmax_d>;
203 def INT_NVVM_MULHI_I : F_MATH_2<"mul.hi.s32 \t$dst, $src0, $src1;", Int32Regs,
204 Int32Regs, Int32Regs, int_nvvm_mulhi_i>;
205 def INT_NVVM_MULHI_UI : F_MATH_2<"mul.hi.u32 \t$dst, $src0, $src1;", Int32Regs,
206 Int32Regs, Int32Regs, int_nvvm_mulhi_ui>;
208 def INT_NVVM_MULHI_LL : F_MATH_2<"mul.hi.s64 \t$dst, $src0, $src1;", Int64Regs,
209 Int64Regs, Int64Regs, int_nvvm_mulhi_ll>;
210 def INT_NVVM_MULHI_ULL : F_MATH_2<"mul.hi.u64 \t$dst, $src0, $src1;", Int64Regs,
211 Int64Regs, Int64Regs, int_nvvm_mulhi_ull>;
213 def INT_NVVM_MUL_RN_FTZ_F : F_MATH_2<"mul.rn.ftz.f32 \t$dst, $src0, $src1;",
214 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rn_ftz_f>;
215 def INT_NVVM_MUL_RN_F : F_MATH_2<"mul.rn.f32 \t$dst, $src0, $src1;",
216 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rn_f>;
217 def INT_NVVM_MUL_RZ_FTZ_F : F_MATH_2<"mul.rz.ftz.f32 \t$dst, $src0, $src1;",
218 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rz_ftz_f>;
219 def INT_NVVM_MUL_RZ_F : F_MATH_2<"mul.rz.f32 \t$dst, $src0, $src1;",
220 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rz_f>;
221 def INT_NVVM_MUL_RM_FTZ_F : F_MATH_2<"mul.rm.ftz.f32 \t$dst, $src0, $src1;",
222 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rm_ftz_f>;
223 def INT_NVVM_MUL_RM_F : F_MATH_2<"mul.rm.f32 \t$dst, $src0, $src1;",
224 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rm_f>;
225 def INT_NVVM_MUL_RP_FTZ_F : F_MATH_2<"mul.rp.ftz.f32 \t$dst, $src0, $src1;",
226 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rp_ftz_f>;
227 def INT_NVVM_MUL_RP_F : F_MATH_2<"mul.rp.f32 \t$dst, $src0, $src1;",
228 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rp_f>;
230 def INT_NVVM_MUL_RN_D : F_MATH_2<"mul.rn.f64 \t$dst, $src0, $src1;",
231 Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rn_d>;
232 def INT_NVVM_MUL_RZ_D : F_MATH_2<"mul.rz.f64 \t$dst, $src0, $src1;",
233 Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rz_d>;
234 def INT_NVVM_MUL_RM_D : F_MATH_2<"mul.rm.f64 \t$dst, $src0, $src1;",
235 Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rm_d>;
236 def INT_NVVM_MUL_RP_D : F_MATH_2<"mul.rp.f64 \t$dst, $src0, $src1;",
237 Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rp_d>;
239 def INT_NVVM_MUL24_I : F_MATH_2<"mul24.lo.s32 \t$dst, $src0, $src1;",
240 Int32Regs, Int32Regs, Int32Regs, int_nvvm_mul24_i>;
241 def INT_NVVM_MUL24_UI : F_MATH_2<"mul24.lo.u32 \t$dst, $src0, $src1;",
242 Int32Regs, Int32Regs, Int32Regs, int_nvvm_mul24_ui>;
248 def INT_NVVM_DIV_APPROX_FTZ_F
249 : F_MATH_2<"div.approx.ftz.f32 \t$dst, $src0, $src1;", Float32Regs,
250 Float32Regs, Float32Regs, int_nvvm_div_approx_ftz_f>;
251 def INT_NVVM_DIV_APPROX_F : F_MATH_2<"div.approx.f32 \t$dst, $src0, $src1;",
252 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_approx_f>;
254 def INT_NVVM_DIV_RN_FTZ_F : F_MATH_2<"div.rn.ftz.f32 \t$dst, $src0, $src1;",
255 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rn_ftz_f>;
256 def INT_NVVM_DIV_RN_F : F_MATH_2<"div.rn.f32 \t$dst, $src0, $src1;",
257 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rn_f>;
258 def INT_NVVM_DIV_RZ_FTZ_F : F_MATH_2<"div.rz.ftz.f32 \t$dst, $src0, $src1;",
259 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rz_ftz_f>;
260 def INT_NVVM_DIV_RZ_F : F_MATH_2<"div.rz.f32 \t$dst, $src0, $src1;",
261 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rz_f>;
262 def INT_NVVM_DIV_RM_FTZ_F : F_MATH_2<"div.rm.ftz.f32 \t$dst, $src0, $src1;",
263 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rm_ftz_f>;
264 def INT_NVVM_DIV_RM_F : F_MATH_2<"div.rm.f32 \t$dst, $src0, $src1;",
265 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rm_f>;
266 def INT_NVVM_DIV_RP_FTZ_F : F_MATH_2<"div.rp.ftz.f32 \t$dst, $src0, $src1;",
267 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rp_ftz_f>;
268 def INT_NVVM_DIV_RP_F : F_MATH_2<"div.rp.f32 \t$dst, $src0, $src1;",
269 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rp_f>;
271 def INT_NVVM_DIV_RN_D : F_MATH_2<"div.rn.f64 \t$dst, $src0, $src1;",
272 Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rn_d>;
273 def INT_NVVM_DIV_RZ_D : F_MATH_2<"div.rz.f64 \t$dst, $src0, $src1;",
274 Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rz_d>;
275 def INT_NVVM_DIV_RM_D : F_MATH_2<"div.rm.f64 \t$dst, $src0, $src1;",
276 Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rm_d>;
277 def INT_NVVM_DIV_RP_D : F_MATH_2<"div.rp.f64 \t$dst, $src0, $src1;",
278 Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rp_d>;
284 def INT_NVVM_BREV32 : F_MATH_1<"brev.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
286 def INT_NVVM_BREV64 : F_MATH_1<"brev.b64 \t$dst, $src0;", Int64Regs, Int64Regs,
293 def INT_NVVM_SAD_I : F_MATH_3<"sad.s32 \t$dst, $src0, $src1, $src2;",
294 Int32Regs, Int32Regs, Int32Regs, Int32Regs, int_nvvm_sad_i>;
295 def INT_NVVM_SAD_UI : F_MATH_3<"sad.u32 \t$dst, $src0, $src1, $src2;",
296 Int32Regs, Int32Regs, Int32Regs, Int32Regs, int_nvvm_sad_ui>;
302 def : Pat<(int_nvvm_floor_ftz_f Float32Regs:$a),
303 (CVT_f32_f32 Float32Regs:$a, CvtRMI_FTZ)>;
304 def : Pat<(int_nvvm_floor_f Float32Regs:$a),
305 (CVT_f32_f32 Float32Regs:$a, CvtRMI)>;
306 def : Pat<(int_nvvm_floor_d Float64Regs:$a),
307 (CVT_f64_f64 Float64Regs:$a, CvtRMI)>;
309 def : Pat<(int_nvvm_ceil_ftz_f Float32Regs:$a),
310 (CVT_f32_f32 Float32Regs:$a, CvtRPI_FTZ)>;
311 def : Pat<(int_nvvm_ceil_f Float32Regs:$a),
312 (CVT_f32_f32 Float32Regs:$a, CvtRPI)>;
313 def : Pat<(int_nvvm_ceil_d Float64Regs:$a),
314 (CVT_f64_f64 Float64Regs:$a, CvtRPI)>;
320 def INT_NVVM_ABS_I : F_MATH_1<"abs.s32 \t$dst, $src0;", Int32Regs, Int32Regs,
322 def INT_NVVM_ABS_LL : F_MATH_1<"abs.s64 \t$dst, $src0;", Int64Regs, Int64Regs,
325 def INT_NVVM_FABS_FTZ_F : F_MATH_1<"abs.ftz.f32 \t$dst, $src0;", Float32Regs,
326 Float32Regs, int_nvvm_fabs_ftz_f>;
327 def INT_NVVM_FABS_F : F_MATH_1<"abs.f32 \t$dst, $src0;", Float32Regs,
328 Float32Regs, int_nvvm_fabs_f>;
330 def INT_NVVM_FABS_D : F_MATH_1<"abs.f64 \t$dst, $src0;", Float64Regs,
331 Float64Regs, int_nvvm_fabs_d>;
337 def : Pat<(int_nvvm_round_ftz_f Float32Regs:$a),
338 (CVT_f32_f32 Float32Regs:$a, CvtRNI_FTZ)>;
339 def : Pat<(int_nvvm_round_f Float32Regs:$a),
340 (CVT_f32_f32 Float32Regs:$a, CvtRNI)>;
341 def : Pat<(int_nvvm_round_d Float64Regs:$a),
342 (CVT_f64_f64 Float64Regs:$a, CvtRNI)>;
348 def : Pat<(int_nvvm_trunc_ftz_f Float32Regs:$a),
349 (CVT_f32_f32 Float32Regs:$a, CvtRZI_FTZ)>;
350 def : Pat<(int_nvvm_trunc_f Float32Regs:$a),
351 (CVT_f32_f32 Float32Regs:$a, CvtRZI)>;
352 def : Pat<(int_nvvm_trunc_d Float64Regs:$a),
353 (CVT_f64_f64 Float64Regs:$a, CvtRZI)>;
359 def : Pat<(int_nvvm_saturate_ftz_f Float32Regs:$a),
360 (CVT_f32_f32 Float32Regs:$a, CvtSAT_FTZ)>;
361 def : Pat<(int_nvvm_saturate_f Float32Regs:$a),
362 (CVT_f32_f32 Float32Regs:$a, CvtSAT)>;
363 def : Pat<(int_nvvm_saturate_d Float64Regs:$a),
364 (CVT_f64_f64 Float64Regs:$a, CvtSAT)>;
370 def INT_NVVM_EX2_APPROX_FTZ_F : F_MATH_1<"ex2.approx.ftz.f32 \t$dst, $src0;",
371 Float32Regs, Float32Regs, int_nvvm_ex2_approx_ftz_f>;
372 def INT_NVVM_EX2_APPROX_F : F_MATH_1<"ex2.approx.f32 \t$dst, $src0;",
373 Float32Regs, Float32Regs, int_nvvm_ex2_approx_f>;
374 def INT_NVVM_EX2_APPROX_D : F_MATH_1<"ex2.approx.f64 \t$dst, $src0;",
375 Float64Regs, Float64Regs, int_nvvm_ex2_approx_d>;
377 def INT_NVVM_LG2_APPROX_FTZ_F : F_MATH_1<"lg2.approx.ftz.f32 \t$dst, $src0;",
378 Float32Regs, Float32Regs, int_nvvm_lg2_approx_ftz_f>;
379 def INT_NVVM_LG2_APPROX_F : F_MATH_1<"lg2.approx.f32 \t$dst, $src0;",
380 Float32Regs, Float32Regs, int_nvvm_lg2_approx_f>;
381 def INT_NVVM_LG2_APPROX_D : F_MATH_1<"lg2.approx.f64 \t$dst, $src0;",
382 Float64Regs, Float64Regs, int_nvvm_lg2_approx_d>;
388 def INT_NVVM_SIN_APPROX_FTZ_F : F_MATH_1<"sin.approx.ftz.f32 \t$dst, $src0;",
389 Float32Regs, Float32Regs, int_nvvm_sin_approx_ftz_f>;
390 def INT_NVVM_SIN_APPROX_F : F_MATH_1<"sin.approx.f32 \t$dst, $src0;",
391 Float32Regs, Float32Regs, int_nvvm_sin_approx_f>;
393 def INT_NVVM_COS_APPROX_FTZ_F : F_MATH_1<"cos.approx.ftz.f32 \t$dst, $src0;",
394 Float32Regs, Float32Regs, int_nvvm_cos_approx_ftz_f>;
395 def INT_NVVM_COS_APPROX_F : F_MATH_1<"cos.approx.f32 \t$dst, $src0;",
396 Float32Regs, Float32Regs, int_nvvm_cos_approx_f>;
402 def INT_NVVM_FMA_RN_FTZ_F
403 : F_MATH_3<"fma.rn.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
404 Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rn_ftz_f>;
405 def INT_NVVM_FMA_RN_F : F_MATH_3<"fma.rn.f32 \t$dst, $src0, $src1, $src2;",
406 Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rn_f>;
407 def INT_NVVM_FMA_RZ_FTZ_F
408 : F_MATH_3<"fma.rz.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
409 Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rz_ftz_f>;
410 def INT_NVVM_FMA_RZ_F : F_MATH_3<"fma.rz.f32 \t$dst, $src0, $src1, $src2;",
411 Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rz_f>;
412 def INT_NVVM_FMA_RM_FTZ_F
413 : F_MATH_3<"fma.rm.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
414 Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rm_ftz_f>;
415 def INT_NVVM_FMA_RM_F : F_MATH_3<"fma.rm.f32 \t$dst, $src0, $src1, $src2;",
416 Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rm_f>;
417 def INT_NVVM_FMA_RP_FTZ_F
418 : F_MATH_3<"fma.rp.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
419 Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rp_ftz_f>;
420 def INT_NVVM_FMA_RP_F : F_MATH_3<"fma.rp.f32 \t$dst, $src0, $src1, $src2;",
421 Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rp_f>;
423 def INT_NVVM_FMA_RN_D : F_MATH_3<"fma.rn.f64 \t$dst, $src0, $src1, $src2;",
424 Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rn_d>;
425 def INT_NVVM_FMA_RZ_D : F_MATH_3<"fma.rz.f64 \t$dst, $src0, $src1, $src2;",
426 Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rz_d>;
427 def INT_NVVM_FMA_RM_D : F_MATH_3<"fma.rm.f64 \t$dst, $src0, $src1, $src2;",
428 Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rm_d>;
429 def INT_NVVM_FMA_RP_D : F_MATH_3<"fma.rp.f64 \t$dst, $src0, $src1, $src2;",
430 Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rp_d>;
436 def INT_NVVM_RCP_RN_FTZ_F : F_MATH_1<"rcp.rn.ftz.f32 \t$dst, $src0;",
437 Float32Regs, Float32Regs, int_nvvm_rcp_rn_ftz_f>;
438 def INT_NVVM_RCP_RN_F : F_MATH_1<"rcp.rn.f32 \t$dst, $src0;",
439 Float32Regs, Float32Regs, int_nvvm_rcp_rn_f>;
440 def INT_NVVM_RCP_RZ_FTZ_F : F_MATH_1<"rcp.rz.ftz.f32 \t$dst, $src0;",
441 Float32Regs, Float32Regs, int_nvvm_rcp_rz_ftz_f>;
442 def INT_NVVM_RCP_RZ_F : F_MATH_1<"rcp.rz.f32 \t$dst, $src0;",
443 Float32Regs, Float32Regs, int_nvvm_rcp_rz_f>;
444 def INT_NVVM_RCP_RM_FTZ_F : F_MATH_1<"rcp.rm.ftz.f32 \t$dst, $src0;",
445 Float32Regs, Float32Regs, int_nvvm_rcp_rm_ftz_f>;
446 def INT_NVVM_RCP_RM_F : F_MATH_1<"rcp.rm.f32 \t$dst, $src0;",
447 Float32Regs, Float32Regs, int_nvvm_rcp_rm_f>;
448 def INT_NVVM_RCP_RP_FTZ_F : F_MATH_1<"rcp.rp.ftz.f32 \t$dst, $src0;",
449 Float32Regs, Float32Regs, int_nvvm_rcp_rp_ftz_f>;
450 def INT_NVVM_RCP_RP_F : F_MATH_1<"rcp.rp.f32 \t$dst, $src0;",
451 Float32Regs, Float32Regs, int_nvvm_rcp_rp_f>;
453 def INT_NVVM_RCP_RN_D : F_MATH_1<"rcp.rn.f64 \t$dst, $src0;", Float64Regs,
454 Float64Regs, int_nvvm_rcp_rn_d>;
455 def INT_NVVM_RCP_RZ_D : F_MATH_1<"rcp.rz.f64 \t$dst, $src0;", Float64Regs,
456 Float64Regs, int_nvvm_rcp_rz_d>;
457 def INT_NVVM_RCP_RM_D : F_MATH_1<"rcp.rm.f64 \t$dst, $src0;", Float64Regs,
458 Float64Regs, int_nvvm_rcp_rm_d>;
459 def INT_NVVM_RCP_RP_D : F_MATH_1<"rcp.rp.f64 \t$dst, $src0;", Float64Regs,
460 Float64Regs, int_nvvm_rcp_rp_d>;
462 def INT_NVVM_RCP_APPROX_FTZ_D : F_MATH_1<"rcp.approx.ftz.f64 \t$dst, $src0;",
463 Float64Regs, Float64Regs, int_nvvm_rcp_approx_ftz_d>;
469 def INT_NVVM_SQRT_RN_FTZ_F : F_MATH_1<"sqrt.rn.ftz.f32 \t$dst, $src0;",
470 Float32Regs, Float32Regs, int_nvvm_sqrt_rn_ftz_f>;
471 def INT_NVVM_SQRT_RN_F : F_MATH_1<"sqrt.rn.f32 \t$dst, $src0;", Float32Regs,
472 Float32Regs, int_nvvm_sqrt_rn_f>;
473 def INT_NVVM_SQRT_RZ_FTZ_F : F_MATH_1<"sqrt.rz.ftz.f32 \t$dst, $src0;",
474 Float32Regs, Float32Regs, int_nvvm_sqrt_rz_ftz_f>;
475 def INT_NVVM_SQRT_RZ_F : F_MATH_1<"sqrt.rz.f32 \t$dst, $src0;", Float32Regs,
476 Float32Regs, int_nvvm_sqrt_rz_f>;
477 def INT_NVVM_SQRT_RM_FTZ_F : F_MATH_1<"sqrt.rm.ftz.f32 \t$dst, $src0;",
478 Float32Regs, Float32Regs, int_nvvm_sqrt_rm_ftz_f>;
479 def INT_NVVM_SQRT_RM_F : F_MATH_1<"sqrt.rm.f32 \t$dst, $src0;", Float32Regs,
480 Float32Regs, int_nvvm_sqrt_rm_f>;
481 def INT_NVVM_SQRT_RP_FTZ_F : F_MATH_1<"sqrt.rp.ftz.f32 \t$dst, $src0;",
482 Float32Regs, Float32Regs, int_nvvm_sqrt_rp_ftz_f>;
483 def INT_NVVM_SQRT_RP_F : F_MATH_1<"sqrt.rp.f32 \t$dst, $src0;", Float32Regs,
484 Float32Regs, int_nvvm_sqrt_rp_f>;
485 def INT_NVVM_SQRT_APPROX_FTZ_F : F_MATH_1<"sqrt.approx.ftz.f32 \t$dst, $src0;",
486 Float32Regs, Float32Regs, int_nvvm_sqrt_approx_ftz_f>;
487 def INT_NVVM_SQRT_APPROX_F : F_MATH_1<"sqrt.approx.f32 \t$dst, $src0;",
488 Float32Regs, Float32Regs, int_nvvm_sqrt_approx_f>;
490 def INT_NVVM_SQRT_RN_D : F_MATH_1<"sqrt.rn.f64 \t$dst, $src0;", Float64Regs,
491 Float64Regs, int_nvvm_sqrt_rn_d>;
492 def INT_NVVM_SQRT_RZ_D : F_MATH_1<"sqrt.rz.f64 \t$dst, $src0;", Float64Regs,
493 Float64Regs, int_nvvm_sqrt_rz_d>;
494 def INT_NVVM_SQRT_RM_D : F_MATH_1<"sqrt.rm.f64 \t$dst, $src0;", Float64Regs,
495 Float64Regs, int_nvvm_sqrt_rm_d>;
496 def INT_NVVM_SQRT_RP_D : F_MATH_1<"sqrt.rp.f64 \t$dst, $src0;", Float64Regs,
497 Float64Regs, int_nvvm_sqrt_rp_d>;
499 // nvvm_sqrt intrinsic
500 def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
501 (INT_NVVM_SQRT_RN_FTZ_F Float32Regs:$a)>, Requires<[doF32FTZ, do_SQRTF32_RN]>;
502 def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
503 (INT_NVVM_SQRT_RN_F Float32Regs:$a)>, Requires<[do_SQRTF32_RN]>;
504 def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
505 (INT_NVVM_SQRT_APPROX_FTZ_F Float32Regs:$a)>, Requires<[doF32FTZ]>;
506 def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
507 (INT_NVVM_SQRT_APPROX_F Float32Regs:$a)>;
513 def INT_NVVM_RSQRT_APPROX_FTZ_F
514 : F_MATH_1<"rsqrt.approx.ftz.f32 \t$dst, $src0;", Float32Regs, Float32Regs,
515 int_nvvm_rsqrt_approx_ftz_f>;
516 def INT_NVVM_RSQRT_APPROX_F : F_MATH_1<"rsqrt.approx.f32 \t$dst, $src0;",
517 Float32Regs, Float32Regs, int_nvvm_rsqrt_approx_f>;
518 def INT_NVVM_RSQRT_APPROX_D : F_MATH_1<"rsqrt.approx.f64 \t$dst, $src0;",
519 Float64Regs, Float64Regs, int_nvvm_rsqrt_approx_d>;
525 def INT_NVVM_ADD_RN_FTZ_F : F_MATH_2<"add.rn.ftz.f32 \t$dst, $src0, $src1;",
526 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rn_ftz_f>;
527 def INT_NVVM_ADD_RN_F : F_MATH_2<"add.rn.f32 \t$dst, $src0, $src1;",
528 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rn_f>;
529 def INT_NVVM_ADD_RZ_FTZ_F : F_MATH_2<"add.rz.ftz.f32 \t$dst, $src0, $src1;",
530 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rz_ftz_f>;
531 def INT_NVVM_ADD_RZ_F : F_MATH_2<"add.rz.f32 \t$dst, $src0, $src1;",
532 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rz_f>;
533 def INT_NVVM_ADD_RM_FTZ_F : F_MATH_2<"add.rm.ftz.f32 \t$dst, $src0, $src1;",
534 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rm_ftz_f>;
535 def INT_NVVM_ADD_RM_F : F_MATH_2<"add.rm.f32 \t$dst, $src0, $src1;",
536 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rm_f>;
537 def INT_NVVM_ADD_RP_FTZ_F : F_MATH_2<"add.rp.ftz.f32 \t$dst, $src0, $src1;",
538 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rp_ftz_f>;
539 def INT_NVVM_ADD_RP_F : F_MATH_2<"add.rp.f32 \t$dst, $src0, $src1;",
540 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rp_f>;
542 def INT_NVVM_ADD_RN_D : F_MATH_2<"add.rn.f64 \t$dst, $src0, $src1;",
543 Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rn_d>;
544 def INT_NVVM_ADD_RZ_D : F_MATH_2<"add.rz.f64 \t$dst, $src0, $src1;",
545 Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rz_d>;
546 def INT_NVVM_ADD_RM_D : F_MATH_2<"add.rm.f64 \t$dst, $src0, $src1;",
547 Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rm_d>;
548 def INT_NVVM_ADD_RP_D : F_MATH_2<"add.rp.f64 \t$dst, $src0, $src1;",
549 Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rp_d>;
555 def : Pat<(int_nvvm_d2f_rn_ftz Float64Regs:$a),
556 (CVT_f32_f64 Float64Regs:$a, CvtRN_FTZ)>;
557 def : Pat<(int_nvvm_d2f_rn Float64Regs:$a),
558 (CVT_f32_f64 Float64Regs:$a, CvtRN)>;
559 def : Pat<(int_nvvm_d2f_rz_ftz Float64Regs:$a),
560 (CVT_f32_f64 Float64Regs:$a, CvtRZ_FTZ)>;
561 def : Pat<(int_nvvm_d2f_rz Float64Regs:$a),
562 (CVT_f32_f64 Float64Regs:$a, CvtRZ)>;
563 def : Pat<(int_nvvm_d2f_rm_ftz Float64Regs:$a),
564 (CVT_f32_f64 Float64Regs:$a, CvtRM_FTZ)>;
565 def : Pat<(int_nvvm_d2f_rm Float64Regs:$a),
566 (CVT_f32_f64 Float64Regs:$a, CvtRM)>;
567 def : Pat<(int_nvvm_d2f_rp_ftz Float64Regs:$a),
568 (CVT_f32_f64 Float64Regs:$a, CvtRP_FTZ)>;
569 def : Pat<(int_nvvm_d2f_rp Float64Regs:$a),
570 (CVT_f32_f64 Float64Regs:$a, CvtRP)>;
572 def : Pat<(int_nvvm_d2i_rn Float64Regs:$a),
573 (CVT_s32_f64 Float64Regs:$a, CvtRNI)>;
574 def : Pat<(int_nvvm_d2i_rz Float64Regs:$a),
575 (CVT_s32_f64 Float64Regs:$a, CvtRZI)>;
576 def : Pat<(int_nvvm_d2i_rm Float64Regs:$a),
577 (CVT_s32_f64 Float64Regs:$a, CvtRMI)>;
578 def : Pat<(int_nvvm_d2i_rp Float64Regs:$a),
579 (CVT_s32_f64 Float64Regs:$a, CvtRPI)>;
581 def : Pat<(int_nvvm_d2ui_rn Float64Regs:$a),
582 (CVT_u32_f64 Float64Regs:$a, CvtRNI)>;
583 def : Pat<(int_nvvm_d2ui_rz Float64Regs:$a),
584 (CVT_u32_f64 Float64Regs:$a, CvtRZI)>;
585 def : Pat<(int_nvvm_d2ui_rm Float64Regs:$a),
586 (CVT_u32_f64 Float64Regs:$a, CvtRMI)>;
587 def : Pat<(int_nvvm_d2ui_rp Float64Regs:$a),
588 (CVT_u32_f64 Float64Regs:$a, CvtRPI)>;
590 def : Pat<(int_nvvm_i2d_rn Int32Regs:$a),
591 (CVT_f64_s32 Int32Regs:$a, CvtRN)>;
592 def : Pat<(int_nvvm_i2d_rz Int32Regs:$a),
593 (CVT_f64_s32 Int32Regs:$a, CvtRZ)>;
594 def : Pat<(int_nvvm_i2d_rm Int32Regs:$a),
595 (CVT_f64_s32 Int32Regs:$a, CvtRM)>;
596 def : Pat<(int_nvvm_i2d_rp Int32Regs:$a),
597 (CVT_f64_s32 Int32Regs:$a, CvtRP)>;
599 def : Pat<(int_nvvm_ui2d_rn Int32Regs:$a),
600 (CVT_f64_u32 Int32Regs:$a, CvtRN)>;
601 def : Pat<(int_nvvm_ui2d_rz Int32Regs:$a),
602 (CVT_f64_u32 Int32Regs:$a, CvtRZ)>;
603 def : Pat<(int_nvvm_ui2d_rm Int32Regs:$a),
604 (CVT_f64_u32 Int32Regs:$a, CvtRM)>;
605 def : Pat<(int_nvvm_ui2d_rp Int32Regs:$a),
606 (CVT_f64_u32 Int32Regs:$a, CvtRP)>;
608 def : Pat<(int_nvvm_f2i_rn_ftz Float32Regs:$a),
609 (CVT_s32_f32 Float32Regs:$a, CvtRNI_FTZ)>;
610 def : Pat<(int_nvvm_f2i_rn Float32Regs:$a),
611 (CVT_s32_f32 Float32Regs:$a, CvtRNI)>;
612 def : Pat<(int_nvvm_f2i_rz_ftz Float32Regs:$a),
613 (CVT_s32_f32 Float32Regs:$a, CvtRZI_FTZ)>;
614 def : Pat<(int_nvvm_f2i_rz Float32Regs:$a),
615 (CVT_s32_f32 Float32Regs:$a, CvtRZI)>;
616 def : Pat<(int_nvvm_f2i_rm_ftz Float32Regs:$a),
617 (CVT_s32_f32 Float32Regs:$a, CvtRMI_FTZ)>;
618 def : Pat<(int_nvvm_f2i_rm Float32Regs:$a),
619 (CVT_s32_f32 Float32Regs:$a, CvtRMI)>;
620 def : Pat<(int_nvvm_f2i_rp_ftz Float32Regs:$a),
621 (CVT_s32_f32 Float32Regs:$a, CvtRPI_FTZ)>;
622 def : Pat<(int_nvvm_f2i_rp Float32Regs:$a),
623 (CVT_s32_f32 Float32Regs:$a, CvtRPI)>;
625 def : Pat<(int_nvvm_f2ui_rn_ftz Float32Regs:$a),
626 (CVT_u32_f32 Float32Regs:$a, CvtRNI_FTZ)>;
627 def : Pat<(int_nvvm_f2ui_rn Float32Regs:$a),
628 (CVT_u32_f32 Float32Regs:$a, CvtRNI)>;
629 def : Pat<(int_nvvm_f2ui_rz_ftz Float32Regs:$a),
630 (CVT_u32_f32 Float32Regs:$a, CvtRZI_FTZ)>;
631 def : Pat<(int_nvvm_f2ui_rz Float32Regs:$a),
632 (CVT_u32_f32 Float32Regs:$a, CvtRZI)>;
633 def : Pat<(int_nvvm_f2ui_rm_ftz Float32Regs:$a),
634 (CVT_u32_f32 Float32Regs:$a, CvtRMI_FTZ)>;
635 def : Pat<(int_nvvm_f2ui_rm Float32Regs:$a),
636 (CVT_u32_f32 Float32Regs:$a, CvtRMI)>;
637 def : Pat<(int_nvvm_f2ui_rp_ftz Float32Regs:$a),
638 (CVT_u32_f32 Float32Regs:$a, CvtRPI_FTZ)>;
639 def : Pat<(int_nvvm_f2ui_rp Float32Regs:$a),
640 (CVT_u32_f32 Float32Regs:$a, CvtRPI)>;
642 def : Pat<(int_nvvm_i2f_rn Int32Regs:$a),
643 (CVT_f32_s32 Int32Regs:$a, CvtRN)>;
644 def : Pat<(int_nvvm_i2f_rz Int32Regs:$a),
645 (CVT_f32_s32 Int32Regs:$a, CvtRZ)>;
646 def : Pat<(int_nvvm_i2f_rm Int32Regs:$a),
647 (CVT_f32_s32 Int32Regs:$a, CvtRM)>;
648 def : Pat<(int_nvvm_i2f_rp Int32Regs:$a),
649 (CVT_f32_s32 Int32Regs:$a, CvtRP)>;
651 def : Pat<(int_nvvm_ui2f_rn Int32Regs:$a),
652 (CVT_f32_u32 Int32Regs:$a, CvtRN)>;
653 def : Pat<(int_nvvm_ui2f_rz Int32Regs:$a),
654 (CVT_f32_u32 Int32Regs:$a, CvtRZ)>;
655 def : Pat<(int_nvvm_ui2f_rm Int32Regs:$a),
656 (CVT_f32_u32 Int32Regs:$a, CvtRM)>;
657 def : Pat<(int_nvvm_ui2f_rp Int32Regs:$a),
658 (CVT_f32_u32 Int32Regs:$a, CvtRP)>;
660 def INT_NVVM_LOHI_I2D : F_MATH_2<"mov.b64 \t$dst, {{$src0, $src1}};",
661 Float64Regs, Int32Regs, Int32Regs, int_nvvm_lohi_i2d>;
663 def INT_NVVM_D2I_LO : F_MATH_1<!strconcat("{{\n\t",
664 !strconcat(".reg .b32 %temp; \n\t",
665 !strconcat("mov.b64 \t{$dst, %temp}, $src0;\n\t",
667 Int32Regs, Float64Regs, int_nvvm_d2i_lo>;
668 def INT_NVVM_D2I_HI : F_MATH_1<!strconcat("{{\n\t",
669 !strconcat(".reg .b32 %temp; \n\t",
670 !strconcat("mov.b64 \t{%temp, $dst}, $src0;\n\t",
672 Int32Regs, Float64Regs, int_nvvm_d2i_hi>;
674 def : Pat<(int_nvvm_f2ll_rn_ftz Float32Regs:$a),
675 (CVT_s64_f32 Float32Regs:$a, CvtRNI_FTZ)>;
676 def : Pat<(int_nvvm_f2ll_rn Float32Regs:$a),
677 (CVT_s64_f32 Float32Regs:$a, CvtRNI)>;
678 def : Pat<(int_nvvm_f2ll_rz_ftz Float32Regs:$a),
679 (CVT_s64_f32 Float32Regs:$a, CvtRZI_FTZ)>;
680 def : Pat<(int_nvvm_f2ll_rz Float32Regs:$a),
681 (CVT_s64_f32 Float32Regs:$a, CvtRZI)>;
682 def : Pat<(int_nvvm_f2ll_rm_ftz Float32Regs:$a),
683 (CVT_s64_f32 Float32Regs:$a, CvtRMI_FTZ)>;
684 def : Pat<(int_nvvm_f2ll_rm Float32Regs:$a),
685 (CVT_s64_f32 Float32Regs:$a, CvtRMI)>;
686 def : Pat<(int_nvvm_f2ll_rp_ftz Float32Regs:$a),
687 (CVT_s64_f32 Float32Regs:$a, CvtRPI_FTZ)>;
688 def : Pat<(int_nvvm_f2ll_rp Float32Regs:$a),
689 (CVT_s64_f32 Float32Regs:$a, CvtRPI)>;
691 def : Pat<(int_nvvm_f2ull_rn_ftz Float32Regs:$a),
692 (CVT_u64_f32 Float32Regs:$a, CvtRNI_FTZ)>;
693 def : Pat<(int_nvvm_f2ull_rn Float32Regs:$a),
694 (CVT_u64_f32 Float32Regs:$a, CvtRNI)>;
695 def : Pat<(int_nvvm_f2ull_rz_ftz Float32Regs:$a),
696 (CVT_u64_f32 Float32Regs:$a, CvtRZI_FTZ)>;
697 def : Pat<(int_nvvm_f2ull_rz Float32Regs:$a),
698 (CVT_u64_f32 Float32Regs:$a, CvtRZI)>;
699 def : Pat<(int_nvvm_f2ull_rm_ftz Float32Regs:$a),
700 (CVT_u64_f32 Float32Regs:$a, CvtRMI_FTZ)>;
701 def : Pat<(int_nvvm_f2ull_rm Float32Regs:$a),
702 (CVT_u64_f32 Float32Regs:$a, CvtRMI)>;
703 def : Pat<(int_nvvm_f2ull_rp_ftz Float32Regs:$a),
704 (CVT_u64_f32 Float32Regs:$a, CvtRPI_FTZ)>;
705 def : Pat<(int_nvvm_f2ull_rp Float32Regs:$a),
706 (CVT_u64_f32 Float32Regs:$a, CvtRPI)>;
708 def : Pat<(int_nvvm_d2ll_rn Float64Regs:$a),
709 (CVT_s64_f64 Float64Regs:$a, CvtRNI)>;
710 def : Pat<(int_nvvm_d2ll_rz Float64Regs:$a),
711 (CVT_s64_f64 Float64Regs:$a, CvtRZI)>;
712 def : Pat<(int_nvvm_d2ll_rm Float64Regs:$a),
713 (CVT_s64_f64 Float64Regs:$a, CvtRMI)>;
714 def : Pat<(int_nvvm_d2ll_rp Float64Regs:$a),
715 (CVT_s64_f64 Float64Regs:$a, CvtRPI)>;
717 def : Pat<(int_nvvm_d2ull_rn Float64Regs:$a),
718 (CVT_u64_f64 Float64Regs:$a, CvtRNI)>;
719 def : Pat<(int_nvvm_d2ull_rz Float64Regs:$a),
720 (CVT_u64_f64 Float64Regs:$a, CvtRZI)>;
721 def : Pat<(int_nvvm_d2ull_rm Float64Regs:$a),
722 (CVT_u64_f64 Float64Regs:$a, CvtRMI)>;
723 def : Pat<(int_nvvm_d2ull_rp Float64Regs:$a),
724 (CVT_u64_f64 Float64Regs:$a, CvtRPI)>;
726 def : Pat<(int_nvvm_ll2f_rn Int64Regs:$a),
727 (CVT_f32_s64 Int64Regs:$a, CvtRN)>;
728 def : Pat<(int_nvvm_ll2f_rz Int64Regs:$a),
729 (CVT_f32_s64 Int64Regs:$a, CvtRZ)>;
730 def : Pat<(int_nvvm_ll2f_rm Int64Regs:$a),
731 (CVT_f32_s64 Int64Regs:$a, CvtRM)>;
732 def : Pat<(int_nvvm_ll2f_rp Int64Regs:$a),
733 (CVT_f32_s64 Int64Regs:$a, CvtRP)>;
735 def : Pat<(int_nvvm_ull2f_rn Int64Regs:$a),
736 (CVT_f32_u64 Int64Regs:$a, CvtRN)>;
737 def : Pat<(int_nvvm_ull2f_rz Int64Regs:$a),
738 (CVT_f32_u64 Int64Regs:$a, CvtRZ)>;
739 def : Pat<(int_nvvm_ull2f_rm Int64Regs:$a),
740 (CVT_f32_u64 Int64Regs:$a, CvtRM)>;
741 def : Pat<(int_nvvm_ull2f_rp Int64Regs:$a),
742 (CVT_f32_u64 Int64Regs:$a, CvtRP)>;
744 def : Pat<(int_nvvm_ll2d_rn Int64Regs:$a),
745 (CVT_f64_s64 Int64Regs:$a, CvtRN)>;
746 def : Pat<(int_nvvm_ll2d_rz Int64Regs:$a),
747 (CVT_f64_s64 Int64Regs:$a, CvtRZ)>;
748 def : Pat<(int_nvvm_ll2d_rm Int64Regs:$a),
749 (CVT_f64_s64 Int64Regs:$a, CvtRM)>;
750 def : Pat<(int_nvvm_ll2d_rp Int64Regs:$a),
751 (CVT_f64_s64 Int64Regs:$a, CvtRP)>;
753 def : Pat<(int_nvvm_ull2d_rn Int64Regs:$a),
754 (CVT_f64_u64 Int64Regs:$a, CvtRN)>;
755 def : Pat<(int_nvvm_ull2d_rz Int64Regs:$a),
756 (CVT_f64_u64 Int64Regs:$a, CvtRZ)>;
757 def : Pat<(int_nvvm_ull2d_rm Int64Regs:$a),
758 (CVT_f64_u64 Int64Regs:$a, CvtRM)>;
759 def : Pat<(int_nvvm_ull2d_rp Int64Regs:$a),
760 (CVT_f64_u64 Int64Regs:$a, CvtRP)>;
763 // FIXME: Ideally, we could use these patterns instead of the scope-creating
764 // patterns, but ptxas does not like these since .s16 is not compatible with
765 // .f16. The solution is to use .bXX for all integer register types, but we
766 // are not there yet.
767 //def : Pat<(int_nvvm_f2h_rn_ftz Float32Regs:$a),
768 // (CVT_f16_f32 Float32Regs:$a, CvtRN_FTZ)>;
769 //def : Pat<(int_nvvm_f2h_rn Float32Regs:$a),
770 // (CVT_f16_f32 Float32Regs:$a, CvtRN)>;
772 //def : Pat<(int_nvvm_h2f Int16Regs:$a),
773 // (CVT_f32_f16 Int16Regs:$a, CvtNONE)>;
775 def INT_NVVM_F2H_RN_FTZ : F_MATH_1<!strconcat("{{\n\t",
776 !strconcat(".reg .b16 %temp;\n\t",
777 !strconcat("cvt.rn.ftz.f16.f32 \t%temp, $src0;\n\t",
778 !strconcat("mov.b16 \t$dst, %temp;\n",
780 Int16Regs, Float32Regs, int_nvvm_f2h_rn_ftz>;
781 def INT_NVVM_F2H_RN : F_MATH_1<!strconcat("{{\n\t",
782 !strconcat(".reg .b16 %temp;\n\t",
783 !strconcat("cvt.rn.f16.f32 \t%temp, $src0;\n\t",
784 !strconcat("mov.b16 \t$dst, %temp;\n",
786 Int16Regs, Float32Regs, int_nvvm_f2h_rn>;
788 def INT_NVVM_H2F : F_MATH_1<!strconcat("{{\n\t",
789 !strconcat(".reg .b16 %temp;\n\t",
790 !strconcat("mov.b16 \t%temp, $src0;\n\t",
791 !strconcat("cvt.f32.f16 \t$dst, %temp;\n\t",
793 Float32Regs, Int16Regs, int_nvvm_h2f>;
795 def : Pat<(f32 (f16_to_f32 Int16Regs:$a)),
796 (CVT_f32_f16 Int16Regs:$a, CvtNONE)>;
797 def : Pat<(i16 (f32_to_f16 Float32Regs:$a)),
798 (CVT_f16_f32 Float32Regs:$a, CvtRN_FTZ)>, Requires<[doF32FTZ]>;
799 def : Pat<(i16 (f32_to_f16 Float32Regs:$a)),
800 (CVT_f16_f32 Float32Regs:$a, CvtRN)>;
806 def INT_NVVM_BITCAST_F2I : F_MATH_1<"mov.b32 \t$dst, $src0;", Int32Regs,
807 Float32Regs, int_nvvm_bitcast_f2i>;
808 def INT_NVVM_BITCAST_I2F : F_MATH_1<"mov.b32 \t$dst, $src0;", Float32Regs,
809 Int32Regs, int_nvvm_bitcast_i2f>;
811 def INT_NVVM_BITCAST_LL2D : F_MATH_1<"mov.b64 \t$dst, $src0;", Float64Regs,
812 Int64Regs, int_nvvm_bitcast_ll2d>;
813 def INT_NVVM_BITCAST_D2LL : F_MATH_1<"mov.b64 \t$dst, $src0;", Int64Regs,
814 Float64Regs, int_nvvm_bitcast_d2ll>;
816 //-----------------------------------
818 //-----------------------------------
820 class ATOMIC_GLOBAL_CHK <dag ops, dag frag>
821 : PatFrag<ops, frag, [{
822 return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_GLOBAL);
824 class ATOMIC_SHARED_CHK <dag ops, dag frag>
825 : PatFrag<ops, frag, [{
826 return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_SHARED);
828 class ATOMIC_GENERIC_CHK <dag ops, dag frag>
829 : PatFrag<ops, frag, [{
830 return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_GENERIC);
833 multiclass F_ATOMIC_2_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
834 string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
835 Operand IMMType, SDNode IMM, Predicate Pred> {
836 def reg : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, regclass:$b),
841 !strconcat(" \t$dst, [$addr], $b;", ""))))),
842 [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b))]>,
844 def imm : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, IMMType:$b),
849 !strconcat(" \t$dst, [$addr], $b;", ""))))),
850 [(set regclass:$dst, (IntOp ptrclass:$addr, IMM:$b))]>,
853 multiclass F_ATOMIC_2<NVPTXRegClass regclass, string SpaceStr, string TypeStr,
854 string OpcStr, PatFrag IntOp, Operand IMMType, SDNode IMM, Predicate Pred> {
855 defm p32 : F_ATOMIC_2_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
856 IntOp, IMMType, IMM, Pred>;
857 defm p64 : F_ATOMIC_2_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
858 IntOp, IMMType, IMM, Pred>;
861 // has 2 operands, neg the second one
862 multiclass F_ATOMIC_2_NEG_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
863 string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
864 Operand IMMType, Predicate Pred> {
865 def reg : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, regclass:$b),
866 !strconcat("{{ \n\t",
867 !strconcat(".reg \t.s",
869 !strconcat(" temp; \n\t",
872 !strconcat(" \ttemp, $b; \n\t",
878 !strconcat(" \t$dst, [$addr], temp; \n\t",
879 !strconcat("}}", "")))))))))))))),
880 [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b))]>,
883 multiclass F_ATOMIC_2_NEG<NVPTXRegClass regclass, string SpaceStr,
884 string TypeStr, string OpcStr, PatFrag IntOp, Operand IMMType,
886 defm p32: F_ATOMIC_2_NEG_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
887 IntOp, IMMType, Pred> ;
888 defm p64: F_ATOMIC_2_NEG_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
889 IntOp, IMMType, Pred> ;
893 multiclass F_ATOMIC_3_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
894 string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
895 Operand IMMType, Predicate Pred> {
896 def reg : NVPTXInst<(outs regclass:$dst),
897 (ins ptrclass:$addr, regclass:$b, regclass:$c),
902 !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
904 (IntOp ptrclass:$addr, regclass:$b, regclass:$c))]>,
906 def imm1 : NVPTXInst<(outs regclass:$dst),
907 (ins ptrclass:$addr, IMMType:$b, regclass:$c),
912 !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
913 [(set regclass:$dst, (IntOp ptrclass:$addr, imm:$b, regclass:$c))]>,
915 def imm2 : NVPTXInst<(outs regclass:$dst),
916 (ins ptrclass:$addr, regclass:$b, IMMType:$c),
921 !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
922 [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b, imm:$c))]>,
924 def imm3 : NVPTXInst<(outs regclass:$dst),
925 (ins ptrclass:$addr, IMMType:$b, IMMType:$c),
930 !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
931 [(set regclass:$dst, (IntOp ptrclass:$addr, imm:$b, imm:$c))]>,
934 multiclass F_ATOMIC_3<NVPTXRegClass regclass, string SpaceStr, string TypeStr,
935 string OpcStr, PatFrag IntOp, Operand IMMType, Predicate Pred> {
936 defm p32 : F_ATOMIC_3_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
937 IntOp, IMMType, Pred>;
938 defm p64 : F_ATOMIC_3_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
939 IntOp, IMMType, Pred>;
944 def atomic_load_add_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
945 (atomic_load_add_32 node:$a, node:$b)>;
946 def atomic_load_add_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
947 (atomic_load_add_32 node:$a, node:$b)>;
948 def atomic_load_add_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
949 (atomic_load_add_32 node:$a, node:$b)>;
950 def atomic_load_add_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
951 (atomic_load_add_64 node:$a, node:$b)>;
952 def atomic_load_add_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
953 (atomic_load_add_64 node:$a, node:$b)>;
954 def atomic_load_add_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
955 (atomic_load_add_64 node:$a, node:$b)>;
956 def atomic_load_add_f32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
957 (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
958 def atomic_load_add_f32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
959 (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
960 def atomic_load_add_f32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
961 (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
963 defm INT_PTX_ATOM_ADD_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".add",
964 atomic_load_add_32_g, i32imm, imm, hasAtomRedG32>;
965 defm INT_PTX_ATOM_ADD_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".add",
966 atomic_load_add_32_s, i32imm, imm, hasAtomRedS32>;
967 defm INT_PTX_ATOM_ADD_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".add",
968 atomic_load_add_32_gen, i32imm, imm, hasAtomRedGen32>;
969 defm INT_PTX_ATOM_ADD_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
970 ".add", atomic_load_add_32_gen, i32imm, imm, useAtomRedG32forGen32>;
972 defm INT_PTX_ATOM_ADD_G_64 : F_ATOMIC_2<Int64Regs, ".global", ".u64", ".add",
973 atomic_load_add_64_g, i64imm, imm, hasAtomRedG64>;
974 defm INT_PTX_ATOM_ADD_S_64 : F_ATOMIC_2<Int64Regs, ".shared", ".u64", ".add",
975 atomic_load_add_64_s, i64imm, imm, hasAtomRedS64>;
976 defm INT_PTX_ATOM_ADD_GEN_64 : F_ATOMIC_2<Int64Regs, "", ".u64", ".add",
977 atomic_load_add_64_gen, i64imm, imm, hasAtomRedGen64>;
978 defm INT_PTX_ATOM_ADD_GEN_64_USE_G : F_ATOMIC_2<Int64Regs, ".global", ".u64",
979 ".add", atomic_load_add_64_gen, i64imm, imm, useAtomRedG64forGen64>;
981 defm INT_PTX_ATOM_ADD_G_F32 : F_ATOMIC_2<Float32Regs, ".global", ".f32", ".add",
982 atomic_load_add_f32_g, f32imm, fpimm, hasAtomAddF32>;
983 defm INT_PTX_ATOM_ADD_S_F32 : F_ATOMIC_2<Float32Regs, ".shared", ".f32", ".add",
984 atomic_load_add_f32_s, f32imm, fpimm, hasAtomAddF32>;
985 defm INT_PTX_ATOM_ADD_GEN_F32 : F_ATOMIC_2<Float32Regs, "", ".f32", ".add",
986 atomic_load_add_f32_gen, f32imm, fpimm, hasAtomAddF32>;
990 def atomic_load_sub_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
991 (atomic_load_sub_32 node:$a, node:$b)>;
992 def atomic_load_sub_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
993 (atomic_load_sub_32 node:$a, node:$b)>;
994 def atomic_load_sub_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
995 (atomic_load_sub_32 node:$a, node:$b)>;
996 def atomic_load_sub_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
997 (atomic_load_sub_64 node:$a, node:$b)>;
998 def atomic_load_sub_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
999 (atomic_load_sub_64 node:$a, node:$b)>;
1000 def atomic_load_sub_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1001 (atomic_load_sub_64 node:$a, node:$b)>;
1003 defm INT_PTX_ATOM_SUB_G_32 : F_ATOMIC_2_NEG<Int32Regs, ".global", "32", ".add",
1004 atomic_load_sub_32_g, i32imm, hasAtomRedG32>;
1005 defm INT_PTX_ATOM_SUB_G_64 : F_ATOMIC_2_NEG<Int64Regs, ".global", "64", ".add",
1006 atomic_load_sub_64_g, i64imm, hasAtomRedG64>;
1007 defm INT_PTX_ATOM_SUB_GEN_32 : F_ATOMIC_2_NEG<Int32Regs, "", "32", ".add",
1008 atomic_load_sub_32_gen, i32imm, hasAtomRedGen32>;
1009 defm INT_PTX_ATOM_SUB_GEN_32_USE_G : F_ATOMIC_2_NEG<Int32Regs, ".global", "32",
1010 ".add", atomic_load_sub_32_gen, i32imm, useAtomRedG32forGen32>;
1011 defm INT_PTX_ATOM_SUB_S_32 : F_ATOMIC_2_NEG<Int32Regs, ".shared", "32", ".add",
1012 atomic_load_sub_32_s, i32imm, hasAtomRedS32>;
1013 defm INT_PTX_ATOM_SUB_S_64 : F_ATOMIC_2_NEG<Int64Regs, ".shared", "64", ".add",
1014 atomic_load_sub_64_s, i64imm, hasAtomRedS64>;
1015 defm INT_PTX_ATOM_SUB_GEN_64 : F_ATOMIC_2_NEG<Int64Regs, "", "64", ".add",
1016 atomic_load_sub_64_gen, i64imm, hasAtomRedGen64>;
1017 defm INT_PTX_ATOM_SUB_GEN_64_USE_G : F_ATOMIC_2_NEG<Int64Regs, ".global", "64",
1018 ".add", atomic_load_sub_64_gen, i64imm, useAtomRedG64forGen64>;
1022 def atomic_swap_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1023 (atomic_swap_32 node:$a, node:$b)>;
1024 def atomic_swap_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1025 (atomic_swap_32 node:$a, node:$b)>;
1026 def atomic_swap_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1027 (atomic_swap_32 node:$a, node:$b)>;
1028 def atomic_swap_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1029 (atomic_swap_64 node:$a, node:$b)>;
1030 def atomic_swap_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1031 (atomic_swap_64 node:$a, node:$b)>;
1032 def atomic_swap_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1033 (atomic_swap_64 node:$a, node:$b)>;
1035 defm INT_PTX_ATOM_SWAP_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".exch",
1036 atomic_swap_32_g, i32imm, imm, hasAtomRedG32>;
1037 defm INT_PTX_ATOM_SWAP_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".exch",
1038 atomic_swap_32_s, i32imm, imm, hasAtomRedS32>;
1039 defm INT_PTX_ATOM_SWAP_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".exch",
1040 atomic_swap_32_gen, i32imm, imm, hasAtomRedGen32>;
1041 defm INT_PTX_ATOM_SWAP_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1042 ".exch", atomic_swap_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1043 defm INT_PTX_ATOM_SWAP_G_64 : F_ATOMIC_2<Int64Regs, ".global", ".b64", ".exch",
1044 atomic_swap_64_g, i64imm, imm, hasAtomRedG64>;
1045 defm INT_PTX_ATOM_SWAP_S_64 : F_ATOMIC_2<Int64Regs, ".shared", ".b64", ".exch",
1046 atomic_swap_64_s, i64imm, imm, hasAtomRedS64>;
1047 defm INT_PTX_ATOM_SWAP_GEN_64 : F_ATOMIC_2<Int64Regs, "", ".b64", ".exch",
1048 atomic_swap_64_gen, i64imm, imm, hasAtomRedGen64>;
1049 defm INT_PTX_ATOM_SWAP_GEN_64_USE_G : F_ATOMIC_2<Int64Regs, ".global", ".b64",
1050 ".exch", atomic_swap_64_gen, i64imm, imm, useAtomRedG64forGen64>;
1054 def atomic_load_max_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b)
1055 , (atomic_load_max_32 node:$a, node:$b)>;
1056 def atomic_load_max_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1057 (atomic_load_max_32 node:$a, node:$b)>;
1058 def atomic_load_max_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1059 (atomic_load_max_32 node:$a, node:$b)>;
1060 def atomic_load_umax_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1061 (atomic_load_umax_32 node:$a, node:$b)>;
1062 def atomic_load_umax_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1063 (atomic_load_umax_32 node:$a, node:$b)>;
1064 def atomic_load_umax_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1065 (atomic_load_umax_32 node:$a, node:$b)>;
1067 defm INT_PTX_ATOM_LOAD_MAX_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".s32",
1068 ".max", atomic_load_max_32_g, i32imm, imm, hasAtomRedG32>;
1069 defm INT_PTX_ATOM_LOAD_MAX_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".s32",
1070 ".max", atomic_load_max_32_s, i32imm, imm, hasAtomRedS32>;
1071 defm INT_PTX_ATOM_LOAD_MAX_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".s32", ".max",
1072 atomic_load_max_32_gen, i32imm, imm, hasAtomRedGen32>;
1073 defm INT_PTX_ATOM_LOAD_MAX_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1074 ".s32", ".max", atomic_load_max_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1075 defm INT_PTX_ATOM_LOAD_UMAX_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1076 ".max", atomic_load_umax_32_g, i32imm, imm, hasAtomRedG32>;
1077 defm INT_PTX_ATOM_LOAD_UMAX_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32",
1078 ".max", atomic_load_umax_32_s, i32imm, imm, hasAtomRedS32>;
1079 defm INT_PTX_ATOM_LOAD_UMAX_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".max",
1080 atomic_load_umax_32_gen, i32imm, imm, hasAtomRedGen32>;
1081 defm INT_PTX_ATOM_LOAD_UMAX_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1082 ".u32", ".max", atomic_load_umax_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1086 def atomic_load_min_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1087 (atomic_load_min_32 node:$a, node:$b)>;
1088 def atomic_load_min_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1089 (atomic_load_min_32 node:$a, node:$b)>;
1090 def atomic_load_min_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1091 (atomic_load_min_32 node:$a, node:$b)>;
1092 def atomic_load_umin_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1093 (atomic_load_umin_32 node:$a, node:$b)>;
1094 def atomic_load_umin_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1095 (atomic_load_umin_32 node:$a, node:$b)>;
1096 def atomic_load_umin_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1097 (atomic_load_umin_32 node:$a, node:$b)>;
1099 defm INT_PTX_ATOM_LOAD_MIN_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".s32",
1100 ".min", atomic_load_min_32_g, i32imm, imm, hasAtomRedG32>;
1101 defm INT_PTX_ATOM_LOAD_MIN_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".s32",
1102 ".min", atomic_load_min_32_s, i32imm, imm, hasAtomRedS32>;
1103 defm INT_PTX_ATOM_LOAD_MIN_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".s32", ".min",
1104 atomic_load_min_32_gen, i32imm, imm, hasAtomRedGen32>;
1105 defm INT_PTX_ATOM_LOAD_MIN_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1106 ".s32", ".min", atomic_load_min_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1107 defm INT_PTX_ATOM_LOAD_UMIN_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1108 ".min", atomic_load_umin_32_g, i32imm, imm, hasAtomRedG32>;
1109 defm INT_PTX_ATOM_LOAD_UMIN_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32",
1110 ".min", atomic_load_umin_32_s, i32imm, imm, hasAtomRedS32>;
1111 defm INT_PTX_ATOM_LOAD_UMIN_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".min",
1112 atomic_load_umin_32_gen, i32imm, imm, hasAtomRedGen32>;
1113 defm INT_PTX_ATOM_LOAD_UMIN_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1114 ".u32", ".min", atomic_load_umin_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1116 // atom_inc atom_dec
1118 def atomic_load_inc_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1119 (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
1120 def atomic_load_inc_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1121 (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
1122 def atomic_load_inc_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1123 (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
1124 def atomic_load_dec_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1125 (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
1126 def atomic_load_dec_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1127 (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
1128 def atomic_load_dec_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1129 (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
1131 defm INT_PTX_ATOM_INC_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".inc",
1132 atomic_load_inc_32_g, i32imm, imm, hasAtomRedG32>;
1133 defm INT_PTX_ATOM_INC_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".inc",
1134 atomic_load_inc_32_s, i32imm, imm, hasAtomRedS32>;
1135 defm INT_PTX_ATOM_INC_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".inc",
1136 atomic_load_inc_32_gen, i32imm, imm, hasAtomRedGen32>;
1137 defm INT_PTX_ATOM_INC_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1138 ".inc", atomic_load_inc_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1139 defm INT_PTX_ATOM_DEC_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".dec",
1140 atomic_load_dec_32_g, i32imm, imm, hasAtomRedG32>;
1141 defm INT_PTX_ATOM_DEC_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".dec",
1142 atomic_load_dec_32_s, i32imm, imm, hasAtomRedS32>;
1143 defm INT_PTX_ATOM_DEC_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".dec",
1144 atomic_load_dec_32_gen, i32imm, imm, hasAtomRedGen32>;
1145 defm INT_PTX_ATOM_DEC_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1146 ".dec", atomic_load_dec_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1150 def atomic_load_and_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1151 (atomic_load_and_32 node:$a, node:$b)>;
1152 def atomic_load_and_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1153 (atomic_load_and_32 node:$a, node:$b)>;
1154 def atomic_load_and_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1155 (atomic_load_and_32 node:$a, node:$b)>;
1157 defm INT_PTX_ATOM_AND_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".and",
1158 atomic_load_and_32_g, i32imm, imm, hasAtomRedG32>;
1159 defm INT_PTX_ATOM_AND_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".and",
1160 atomic_load_and_32_s, i32imm, imm, hasAtomRedS32>;
1161 defm INT_PTX_ATOM_AND_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".and",
1162 atomic_load_and_32_gen, i32imm, imm, hasAtomRedGen32>;
1163 defm INT_PTX_ATOM_AND_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1164 ".and", atomic_load_and_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1168 def atomic_load_or_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1169 (atomic_load_or_32 node:$a, node:$b)>;
1170 def atomic_load_or_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1171 (atomic_load_or_32 node:$a, node:$b)>;
1172 def atomic_load_or_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1173 (atomic_load_or_32 node:$a, node:$b)>;
1175 defm INT_PTX_ATOM_OR_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".or",
1176 atomic_load_or_32_g, i32imm, imm, hasAtomRedG32>;
1177 defm INT_PTX_ATOM_OR_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".or",
1178 atomic_load_or_32_gen, i32imm, imm, hasAtomRedGen32>;
1179 defm INT_PTX_ATOM_OR_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1180 ".or", atomic_load_or_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1181 defm INT_PTX_ATOM_OR_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".or",
1182 atomic_load_or_32_s, i32imm, imm, hasAtomRedS32>;
1186 def atomic_load_xor_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1187 (atomic_load_xor_32 node:$a, node:$b)>;
1188 def atomic_load_xor_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1189 (atomic_load_xor_32 node:$a, node:$b)>;
1190 def atomic_load_xor_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1191 (atomic_load_xor_32 node:$a, node:$b)>;
1193 defm INT_PTX_ATOM_XOR_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".xor",
1194 atomic_load_xor_32_g, i32imm, imm, hasAtomRedG32>;
1195 defm INT_PTX_ATOM_XOR_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".xor",
1196 atomic_load_xor_32_s, i32imm, imm, hasAtomRedS32>;
1197 defm INT_PTX_ATOM_XOR_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".xor",
1198 atomic_load_xor_32_gen, i32imm, imm, hasAtomRedGen32>;
1199 defm INT_PTX_ATOM_XOR_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1200 ".xor", atomic_load_xor_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1204 def atomic_cmp_swap_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b, node:$c),
1205 (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
1206 def atomic_cmp_swap_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b, node:$c),
1207 (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
1208 def atomic_cmp_swap_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b, node:$c),
1209 (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
1210 def atomic_cmp_swap_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b, node:$c),
1211 (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
1212 def atomic_cmp_swap_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b, node:$c),
1213 (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
1214 def atomic_cmp_swap_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b, node:$c),
1215 (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
1217 defm INT_PTX_ATOM_CAS_G_32 : F_ATOMIC_3<Int32Regs, ".global", ".b32", ".cas",
1218 atomic_cmp_swap_32_g, i32imm, hasAtomRedG32>;
1219 defm INT_PTX_ATOM_CAS_S_32 : F_ATOMIC_3<Int32Regs, ".shared", ".b32", ".cas",
1220 atomic_cmp_swap_32_s, i32imm, hasAtomRedS32>;
1221 defm INT_PTX_ATOM_CAS_GEN_32 : F_ATOMIC_3<Int32Regs, "", ".b32", ".cas",
1222 atomic_cmp_swap_32_gen, i32imm, hasAtomRedGen32>;
1223 defm INT_PTX_ATOM_CAS_GEN_32_USE_G : F_ATOMIC_3<Int32Regs, ".global", ".b32",
1224 ".cas", atomic_cmp_swap_32_gen, i32imm, useAtomRedG32forGen32>;
1225 defm INT_PTX_ATOM_CAS_G_64 : F_ATOMIC_3<Int64Regs, ".global", ".b64", ".cas",
1226 atomic_cmp_swap_64_g, i64imm, hasAtomRedG64>;
1227 defm INT_PTX_ATOM_CAS_S_64 : F_ATOMIC_3<Int64Regs, ".shared", ".b64", ".cas",
1228 atomic_cmp_swap_64_s, i64imm, hasAtomRedS64>;
1229 defm INT_PTX_ATOM_CAS_GEN_64 : F_ATOMIC_3<Int64Regs, "", ".b64", ".cas",
1230 atomic_cmp_swap_64_gen, i64imm, hasAtomRedGen64>;
1231 defm INT_PTX_ATOM_CAS_GEN_64_USE_G : F_ATOMIC_3<Int64Regs, ".global", ".b64",
1232 ".cas", atomic_cmp_swap_64_gen, i64imm, useAtomRedG64forGen64>;
1235 //-----------------------------------
1236 // Read Special Registers
1237 //-----------------------------------
1238 class F_SREG<string OpStr, NVPTXRegClass regclassOut, Intrinsic IntOp> :
1239 NVPTXInst<(outs regclassOut:$dst), (ins),
1241 [(set regclassOut:$dst, (IntOp))]>;
1243 def INT_PTX_SREG_TID_X : F_SREG<"mov.u32 \t$dst, %tid.x;", Int32Regs,
1244 int_nvvm_read_ptx_sreg_tid_x>;
1245 def INT_PTX_SREG_TID_Y : F_SREG<"mov.u32 \t$dst, %tid.y;", Int32Regs,
1246 int_nvvm_read_ptx_sreg_tid_y>;
1247 def INT_PTX_SREG_TID_Z : F_SREG<"mov.u32 \t$dst, %tid.z;", Int32Regs,
1248 int_nvvm_read_ptx_sreg_tid_z>;
1250 def INT_PTX_SREG_NTID_X : F_SREG<"mov.u32 \t$dst, %ntid.x;", Int32Regs,
1251 int_nvvm_read_ptx_sreg_ntid_x>;
1252 def INT_PTX_SREG_NTID_Y : F_SREG<"mov.u32 \t$dst, %ntid.y;", Int32Regs,
1253 int_nvvm_read_ptx_sreg_ntid_y>;
1254 def INT_PTX_SREG_NTID_Z : F_SREG<"mov.u32 \t$dst, %ntid.z;", Int32Regs,
1255 int_nvvm_read_ptx_sreg_ntid_z>;
1257 def INT_PTX_SREG_CTAID_X : F_SREG<"mov.u32 \t$dst, %ctaid.x;", Int32Regs,
1258 int_nvvm_read_ptx_sreg_ctaid_x>;
1259 def INT_PTX_SREG_CTAID_Y : F_SREG<"mov.u32 \t$dst, %ctaid.y;", Int32Regs,
1260 int_nvvm_read_ptx_sreg_ctaid_y>;
1261 def INT_PTX_SREG_CTAID_Z : F_SREG<"mov.u32 \t$dst, %ctaid.z;", Int32Regs,
1262 int_nvvm_read_ptx_sreg_ctaid_z>;
1264 def INT_PTX_SREG_NCTAID_X : F_SREG<"mov.u32 \t$dst, %nctaid.x;", Int32Regs,
1265 int_nvvm_read_ptx_sreg_nctaid_x>;
1266 def INT_PTX_SREG_NCTAID_Y : F_SREG<"mov.u32 \t$dst, %nctaid.y;", Int32Regs,
1267 int_nvvm_read_ptx_sreg_nctaid_y>;
1268 def INT_PTX_SREG_NCTAID_Z : F_SREG<"mov.u32 \t$dst, %nctaid.z;", Int32Regs,
1269 int_nvvm_read_ptx_sreg_nctaid_z>;
1271 def INT_PTX_SREG_WARPSIZE : F_SREG<"mov.u32 \t$dst, WARP_SZ;", Int32Regs,
1272 int_nvvm_read_ptx_sreg_warpsize>;
1275 //-----------------------------------
1276 // Support for ldu on sm_20 or later
1277 //-----------------------------------
1279 def ldu_i8 : PatFrag<(ops node:$ptr), (int_nvvm_ldu_global_i node:$ptr), [{
1280 MemIntrinsicSDNode *M = cast<MemIntrinsicSDNode>(N);
1281 return M->getMemoryVT() == MVT::i8;
1285 // @TODO: Revisit this, Changed imemAny to imem
1286 multiclass LDU_G<string TyStr, NVPTXRegClass regclass, Intrinsic IntOp> {
1287 def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
1288 !strconcat("ldu.global.", TyStr),
1289 [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDU]>;
1290 def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
1291 !strconcat("ldu.global.", TyStr),
1292 [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDU]>;
1293 def avar: NVPTXInst<(outs regclass:$result), (ins imem:$src),
1294 !strconcat("ldu.global.", TyStr),
1295 [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
1297 def ari : NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
1298 !strconcat("ldu.global.", TyStr),
1299 [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDU]>;
1300 def ari64 : NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
1301 !strconcat("ldu.global.", TyStr),
1302 [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDU]>;
1305 multiclass LDU_G_NOINTRIN<string TyStr, NVPTXRegClass regclass, PatFrag IntOp> {
1306 def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
1307 !strconcat("ldu.global.", TyStr),
1308 [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDU]>;
1309 def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
1310 !strconcat("ldu.global.", TyStr),
1311 [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDU]>;
1312 def avar: NVPTXInst<(outs regclass:$result), (ins imem:$src),
1313 !strconcat("ldu.global.", TyStr),
1314 [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
1316 def ari : NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
1317 !strconcat("ldu.global.", TyStr),
1318 [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDU]>;
1319 def ari64 : NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
1320 !strconcat("ldu.global.", TyStr),
1321 [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDU]>;
1324 defm INT_PTX_LDU_GLOBAL_i8 : LDU_G_NOINTRIN<"u8 \t$result, [$src];", Int16Regs,
1326 defm INT_PTX_LDU_GLOBAL_i16 : LDU_G<"u16 \t$result, [$src];", Int16Regs,
1327 int_nvvm_ldu_global_i>;
1328 defm INT_PTX_LDU_GLOBAL_i32 : LDU_G<"u32 \t$result, [$src];", Int32Regs,
1329 int_nvvm_ldu_global_i>;
1330 defm INT_PTX_LDU_GLOBAL_i64 : LDU_G<"u64 \t$result, [$src];", Int64Regs,
1331 int_nvvm_ldu_global_i>;
1332 defm INT_PTX_LDU_GLOBAL_f32 : LDU_G<"f32 \t$result, [$src];", Float32Regs,
1333 int_nvvm_ldu_global_f>;
1334 defm INT_PTX_LDU_GLOBAL_f64 : LDU_G<"f64 \t$result, [$src];", Float64Regs,
1335 int_nvvm_ldu_global_f>;
1336 defm INT_PTX_LDU_GLOBAL_p32 : LDU_G<"u32 \t$result, [$src];", Int32Regs,
1337 int_nvvm_ldu_global_p>;
1338 defm INT_PTX_LDU_GLOBAL_p64 : LDU_G<"u64 \t$result, [$src];", Int64Regs,
1339 int_nvvm_ldu_global_p>;
1343 // Elementized vector ldu
1344 multiclass VLDU_G_ELE_V2<string TyStr, NVPTXRegClass regclass> {
1345 def _areg32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1346 (ins Int32Regs:$src),
1347 !strconcat("ldu.global.", TyStr), []>;
1348 def _areg64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1349 (ins Int64Regs:$src),
1350 !strconcat("ldu.global.", TyStr), []>;
1351 def _ari32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1353 !strconcat("ldu.global.", TyStr), []>;
1354 def _ari64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1356 !strconcat("ldu.global.", TyStr), []>;
1357 def _avar: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1359 !strconcat("ldu.global.", TyStr), []>;
1362 multiclass VLDU_G_ELE_V4<string TyStr, NVPTXRegClass regclass> {
1363 def _areg32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1364 regclass:$dst4), (ins Int32Regs:$src),
1365 !strconcat("ldu.global.", TyStr), []>;
1366 def _areg64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1367 regclass:$dst4), (ins Int64Regs:$src),
1368 !strconcat("ldu.global.", TyStr), []>;
1369 def _ari32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1370 regclass:$dst4), (ins MEMri:$src),
1371 !strconcat("ldu.global.", TyStr), []>;
1372 def _ari64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1373 regclass:$dst4), (ins MEMri64:$src),
1374 !strconcat("ldu.global.", TyStr), []>;
1375 def _avar: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1376 regclass:$dst4), (ins imemAny:$src),
1377 !strconcat("ldu.global.", TyStr), []>;
1380 defm INT_PTX_LDU_G_v2i8_ELE
1381 : VLDU_G_ELE_V2<"v2.u8 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
1382 defm INT_PTX_LDU_G_v2i16_ELE
1383 : VLDU_G_ELE_V2<"v2.u16 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
1384 defm INT_PTX_LDU_G_v2i32_ELE
1385 : VLDU_G_ELE_V2<"v2.u32 \t{{$dst1, $dst2}}, [$src];", Int32Regs>;
1386 defm INT_PTX_LDU_G_v2f32_ELE
1387 : VLDU_G_ELE_V2<"v2.f32 \t{{$dst1, $dst2}}, [$src];", Float32Regs>;
1388 defm INT_PTX_LDU_G_v2i64_ELE
1389 : VLDU_G_ELE_V2<"v2.u64 \t{{$dst1, $dst2}}, [$src];", Int64Regs>;
1390 defm INT_PTX_LDU_G_v2f64_ELE
1391 : VLDU_G_ELE_V2<"v2.f64 \t{{$dst1, $dst2}}, [$src];", Float64Regs>;
1392 defm INT_PTX_LDU_G_v4i8_ELE
1393 : VLDU_G_ELE_V4<"v4.u8 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
1394 defm INT_PTX_LDU_G_v4i16_ELE
1395 : VLDU_G_ELE_V4<"v4.u16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
1397 defm INT_PTX_LDU_G_v4i32_ELE
1398 : VLDU_G_ELE_V4<"v4.u32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
1400 defm INT_PTX_LDU_G_v4f32_ELE
1401 : VLDU_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
1405 //-----------------------------------
1406 // Support for ldg on sm_35 or later
1407 //-----------------------------------
1409 def ldg_i8 : PatFrag<(ops node:$ptr), (int_nvvm_ldg_global_i node:$ptr), [{
1410 MemIntrinsicSDNode *M = cast<MemIntrinsicSDNode>(N);
1411 return M->getMemoryVT() == MVT::i8;
1414 multiclass LDG_G<string TyStr, NVPTXRegClass regclass, Intrinsic IntOp> {
1415 def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
1416 !strconcat("ld.global.nc.", TyStr),
1417 [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDG]>;
1418 def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
1419 !strconcat("ld.global.nc.", TyStr),
1420 [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDG]>;
1421 def avar: NVPTXInst<(outs regclass:$result), (ins imem:$src),
1422 !strconcat("ld.global.nc.", TyStr),
1423 [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
1425 def ari : NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
1426 !strconcat("ld.global.nc.", TyStr),
1427 [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDG]>;
1428 def ari64 : NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
1429 !strconcat("ld.global.nc.", TyStr),
1430 [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDG]>;
1433 multiclass LDG_G_NOINTRIN<string TyStr, NVPTXRegClass regclass, PatFrag IntOp> {
1434 def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
1435 !strconcat("ld.global.nc.", TyStr),
1436 [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDG]>;
1437 def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
1438 !strconcat("ld.global.nc.", TyStr),
1439 [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDG]>;
1440 def avar: NVPTXInst<(outs regclass:$result), (ins imem:$src),
1441 !strconcat("ld.global.nc.", TyStr),
1442 [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
1444 def ari : NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
1445 !strconcat("ld.global.nc.", TyStr),
1446 [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDG]>;
1447 def ari64 : NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
1448 !strconcat("ld.global.nc.", TyStr),
1449 [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDG]>;
1452 defm INT_PTX_LDG_GLOBAL_i8
1453 : LDG_G_NOINTRIN<"u8 \t$result, [$src];", Int16Regs, ldg_i8>;
1454 defm INT_PTX_LDG_GLOBAL_i16
1455 : LDG_G<"u16 \t$result, [$src];", Int16Regs, int_nvvm_ldg_global_i>;
1456 defm INT_PTX_LDG_GLOBAL_i32
1457 : LDG_G<"u32 \t$result, [$src];", Int32Regs, int_nvvm_ldg_global_i>;
1458 defm INT_PTX_LDG_GLOBAL_i64
1459 : LDG_G<"u64 \t$result, [$src];", Int64Regs, int_nvvm_ldg_global_i>;
1460 defm INT_PTX_LDG_GLOBAL_f32
1461 : LDG_G<"f32 \t$result, [$src];", Float32Regs, int_nvvm_ldg_global_f>;
1462 defm INT_PTX_LDG_GLOBAL_f64
1463 : LDG_G<"f64 \t$result, [$src];", Float64Regs, int_nvvm_ldg_global_f>;
1464 defm INT_PTX_LDG_GLOBAL_p32
1465 : LDG_G<"u32 \t$result, [$src];", Int32Regs, int_nvvm_ldg_global_p>;
1466 defm INT_PTX_LDG_GLOBAL_p64
1467 : LDG_G<"u64 \t$result, [$src];", Int64Regs, int_nvvm_ldg_global_p>;
1471 // Elementized vector ldg
1472 multiclass VLDG_G_ELE_V2<string TyStr, NVPTXRegClass regclass> {
1473 def _areg32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1474 (ins Int32Regs:$src),
1475 !strconcat("ld.global.nc.", TyStr), []>;
1476 def _areg64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1477 (ins Int64Regs:$src),
1478 !strconcat("ld.global.nc.", TyStr), []>;
1479 def _ari32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1481 !strconcat("ld.global.nc.", TyStr), []>;
1482 def _ari64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1484 !strconcat("ld.global.nc.", TyStr), []>;
1485 def _avar: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1487 !strconcat("ld.global.nc.", TyStr), []>;
1490 multiclass VLDG_G_ELE_V4<string TyStr, NVPTXRegClass regclass> {
1491 def _areg32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1492 regclass:$dst4), (ins Int32Regs:$src),
1493 !strconcat("ld.global.nc.", TyStr), []>;
1494 def _areg64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1495 regclass:$dst4), (ins Int64Regs:$src),
1496 !strconcat("ld.global.nc.", TyStr), []>;
1497 def _ari32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1498 regclass:$dst4), (ins MEMri:$src),
1499 !strconcat("ld.global.nc.", TyStr), []>;
1500 def _ari64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1501 regclass:$dst4), (ins MEMri64:$src),
1502 !strconcat("ld.global.nc.", TyStr), []>;
1503 def _avar: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1504 regclass:$dst4), (ins imemAny:$src),
1505 !strconcat("ld.global.nc.", TyStr), []>;
1508 // FIXME: 8-bit LDG should be fixed once LDG/LDU nodes are made into proper loads.
1509 defm INT_PTX_LDG_G_v2i8_ELE
1510 : VLDG_G_ELE_V2<"v2.u8 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
1511 defm INT_PTX_LDG_G_v2i16_ELE
1512 : VLDG_G_ELE_V2<"v2.u16 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
1513 defm INT_PTX_LDG_G_v2i32_ELE
1514 : VLDG_G_ELE_V2<"v2.u32 \t{{$dst1, $dst2}}, [$src];", Int32Regs>;
1515 defm INT_PTX_LDG_G_v2f32_ELE
1516 : VLDG_G_ELE_V2<"v2.f32 \t{{$dst1, $dst2}}, [$src];", Float32Regs>;
1517 defm INT_PTX_LDG_G_v2i64_ELE
1518 : VLDG_G_ELE_V2<"v2.u64 \t{{$dst1, $dst2}}, [$src];", Int64Regs>;
1519 defm INT_PTX_LDG_G_v2f64_ELE
1520 : VLDG_G_ELE_V2<"v2.f64 \t{{$dst1, $dst2}}, [$src];", Float64Regs>;
1521 defm INT_PTX_LDG_G_v4i8_ELE
1522 : VLDG_G_ELE_V4<"v4.u8 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
1523 defm INT_PTX_LDG_G_v4i16_ELE
1524 : VLDG_G_ELE_V4<"v4.u16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
1525 defm INT_PTX_LDG_G_v4i32_ELE
1526 : VLDG_G_ELE_V4<"v4.u32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int32Regs>;
1527 defm INT_PTX_LDG_G_v4f32_ELE
1528 : VLDG_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Float32Regs>;
1531 multiclass NG_TO_G<string Str, Intrinsic Intrin> {
1532 def _yes : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1533 !strconcat("cvta.", !strconcat(Str, ".u32 \t$result, $src;")),
1534 [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>,
1535 Requires<[hasGenericLdSt]>;
1536 def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1537 !strconcat("cvta.", !strconcat(Str, ".u64 \t$result, $src;")),
1538 [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>,
1539 Requires<[hasGenericLdSt]>;
1541 // @TODO: Are these actually needed? I believe global addresses will be copied
1542 // to register values anyway.
1543 /*def __addr_yes : NVPTXInst<(outs Int32Regs:$result), (ins imemAny:$src),
1544 !strconcat("cvta.", !strconcat(Str, ".u32 \t$result, $src;")),
1545 [(set Int32Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>,
1546 Requires<[hasGenericLdSt]>;
1547 def __addr_yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins imemAny:$src),
1548 !strconcat("cvta.", !strconcat(Str, ".u64 \t$result, $src;")),
1549 [(set Int64Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>,
1550 Requires<[hasGenericLdSt]>;*/
1552 def _no : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1553 "mov.u32 \t$result, $src;",
1554 [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>;
1555 def _no_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1556 "mov.u64 \t$result, $src;",
1557 [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>;
1559 // @TODO: Are these actually needed? I believe global addresses will be copied
1560 // to register values anyway.
1561 /*def _addr_no : NVPTXInst<(outs Int32Regs:$result), (ins imem:$src),
1562 "mov.u32 \t$result, $src;",
1563 [(set Int32Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>;
1564 def _addr_no_64 : NVPTXInst<(outs Int64Regs:$result), (ins imem:$src),
1565 "mov.u64 \t$result, $src;",
1566 [(set Int64Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>;*/
1569 multiclass G_TO_NG<string Str, Intrinsic Intrin> {
1570 def _yes : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1571 !strconcat("cvta.to.", !strconcat(Str, ".u32 \t$result, $src;")),
1572 [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>,
1573 Requires<[hasGenericLdSt]>;
1574 def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1575 !strconcat("cvta.to.", !strconcat(Str, ".u64 \t$result, $src;")),
1576 [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>,
1577 Requires<[hasGenericLdSt]>;
1578 def _no : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1579 "mov.u32 \t$result, $src;",
1580 [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>;
1581 def _no_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1582 "mov.u64 \t$result, $src;",
1583 [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>;
1586 defm cvta_local : NG_TO_G<"local", int_nvvm_ptr_local_to_gen>;
1587 defm cvta_shared : NG_TO_G<"shared", int_nvvm_ptr_shared_to_gen>;
1588 defm cvta_global : NG_TO_G<"global", int_nvvm_ptr_global_to_gen>;
1589 defm cvta_const : NG_TO_G<"const", int_nvvm_ptr_constant_to_gen>;
1591 defm cvta_to_local : G_TO_NG<"local", int_nvvm_ptr_gen_to_local>;
1592 defm cvta_to_shared : G_TO_NG<"shared", int_nvvm_ptr_gen_to_shared>;
1593 defm cvta_to_global : G_TO_NG<"global", int_nvvm_ptr_gen_to_global>;
1594 defm cvta_to_const : G_TO_NG<"const", int_nvvm_ptr_gen_to_constant>;
1597 // nvvm.ptr.gen.to.param
1598 def nvvm_ptr_gen_to_param : NVPTXInst<(outs Int32Regs:$result),
1599 (ins Int32Regs:$src),
1600 "mov.u32 \t$result, $src;",
1601 [(set Int32Regs:$result,
1602 (int_nvvm_ptr_gen_to_param Int32Regs:$src))]>;
1603 def nvvm_ptr_gen_to_param_64 : NVPTXInst<(outs Int64Regs:$result),
1604 (ins Int64Regs:$src),
1605 "mov.u64 \t$result, $src;",
1606 [(set Int64Regs:$result,
1607 (int_nvvm_ptr_gen_to_param Int64Regs:$src))]>;
1610 // nvvm.move intrinsicc
1611 def nvvm_move_i16 : NVPTXInst<(outs Int16Regs:$r), (ins Int16Regs:$s),
1612 "mov.b16 \t$r, $s;",
1614 (int_nvvm_move_i16 Int16Regs:$s))]>;
1615 def nvvm_move_i32 : NVPTXInst<(outs Int32Regs:$r), (ins Int32Regs:$s),
1616 "mov.b32 \t$r, $s;",
1618 (int_nvvm_move_i32 Int32Regs:$s))]>;
1619 def nvvm_move_i64 : NVPTXInst<(outs Int64Regs:$r), (ins Int64Regs:$s),
1620 "mov.b64 \t$r, $s;",
1622 (int_nvvm_move_i64 Int64Regs:$s))]>;
1623 def nvvm_move_float : NVPTXInst<(outs Float32Regs:$r), (ins Float32Regs:$s),
1624 "mov.f32 \t$r, $s;",
1625 [(set Float32Regs:$r,
1626 (int_nvvm_move_float Float32Regs:$s))]>;
1627 def nvvm_move_double : NVPTXInst<(outs Float64Regs:$r), (ins Float64Regs:$s),
1628 "mov.f64 \t$r, $s;",
1629 [(set Float64Regs:$r,
1630 (int_nvvm_move_double Float64Regs:$s))]>;
1631 def nvvm_move_ptr32 : NVPTXInst<(outs Int32Regs:$r), (ins Int32Regs:$s),
1632 "mov.u32 \t$r, $s;",
1634 (int_nvvm_move_ptr Int32Regs:$s))]>;
1635 def nvvm_move_ptr64 : NVPTXInst<(outs Int64Regs:$r), (ins Int64Regs:$s),
1636 "mov.u64 \t$r, $s;",
1638 (int_nvvm_move_ptr Int64Regs:$s))]>;
1640 // @TODO: Are these actually needed, or will we always just see symbols
1641 // copied to registers first?
1642 /*def nvvm_move_sym32 : NVPTXInst<(outs Int32Regs:$r), (ins imem:$s),
1643 "mov.u32 \t$r, $s;",
1645 (int_nvvm_move_ptr texternalsym:$s))]>;
1646 def nvvm_move_sym64 : NVPTXInst<(outs Int64Regs:$r), (ins imem:$s),
1647 "mov.u64 \t$r, $s;",
1649 (int_nvvm_move_ptr texternalsym:$s))]>;*/
1652 // MoveParam %r1, param
1653 // ptr_local_to_gen %r2, %r1
1654 // ptr_gen_to_local %r3, %r2
1658 // @TODO: Revisit this. There is a type
1659 // contradiction between iPTRAny and iPTR for the addr defs, so the move_sym
1660 // instructions are not currently defined. However, we can use the ptr
1661 // variants and the asm printer will do the right thing.
1662 def : Pat<(i64 (int_nvvm_ptr_gen_to_local (int_nvvm_ptr_local_to_gen
1663 (MoveParam texternalsym:$src)))),
1664 (nvvm_move_ptr64 texternalsym:$src)>;
1665 def : Pat<(i32 (int_nvvm_ptr_gen_to_local (int_nvvm_ptr_local_to_gen
1666 (MoveParam texternalsym:$src)))),
1667 (nvvm_move_ptr32 texternalsym:$src)>;
1670 : NVPTXInst<(outs Int64Regs:$result), (ins imem:$src),
1671 "mov.u64 \t$result, $src;", []>;
1673 //-----------------------------------
1674 // Compiler Error Warn
1675 // - Just ignore them in codegen
1676 //-----------------------------------
1678 def INT_NVVM_COMPILER_WARN_32 : NVPTXInst<(outs), (ins Int32Regs:$a),
1679 "// llvm.nvvm.compiler.warn()",
1680 [(int_nvvm_compiler_warn Int32Regs:$a)]>;
1681 def INT_NVVM_COMPILER_WARN_64 : NVPTXInst<(outs), (ins Int64Regs:$a),
1682 "// llvm.nvvm.compiler.warn()",
1683 [(int_nvvm_compiler_warn Int64Regs:$a)]>;
1684 def INT_NVVM_COMPILER_ERROR_32 : NVPTXInst<(outs), (ins Int32Regs:$a),
1685 "// llvm.nvvm.compiler.error()",
1686 [(int_nvvm_compiler_error Int32Regs:$a)]>;
1687 def INT_NVVM_COMPILER_ERROR_64 : NVPTXInst<(outs), (ins Int64Regs:$a),
1688 "// llvm.nvvm.compiler.error()",
1689 [(int_nvvm_compiler_error Int64Regs:$a)]>;
1694 def ISSPACEP_CONST_32
1695 : NVPTXInst<(outs Int1Regs:$d), (ins Int32Regs:$a),
1696 "isspacep.const \t$d, $a;",
1697 [(set Int1Regs:$d, (int_nvvm_isspacep_const Int32Regs:$a))]>,
1698 Requires<[hasPTX31]>;
1699 def ISSPACEP_CONST_64
1700 : NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
1701 "isspacep.const \t$d, $a;",
1702 [(set Int1Regs:$d, (int_nvvm_isspacep_const Int64Regs:$a))]>,
1703 Requires<[hasPTX31]>;
1704 def ISSPACEP_GLOBAL_32
1705 : NVPTXInst<(outs Int1Regs:$d), (ins Int32Regs:$a),
1706 "isspacep.global \t$d, $a;",
1707 [(set Int1Regs:$d, (int_nvvm_isspacep_global Int32Regs:$a))]>;
1708 def ISSPACEP_GLOBAL_64
1709 : NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
1710 "isspacep.global \t$d, $a;",
1711 [(set Int1Regs:$d, (int_nvvm_isspacep_global Int64Regs:$a))]>;
1712 def ISSPACEP_LOCAL_32
1713 : NVPTXInst<(outs Int1Regs:$d), (ins Int32Regs:$a),
1714 "isspacep.local \t$d, $a;",
1715 [(set Int1Regs:$d, (int_nvvm_isspacep_local Int32Regs:$a))]>;
1716 def ISSPACEP_LOCAL_64
1717 : NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
1718 "isspacep.local \t$d, $a;",
1719 [(set Int1Regs:$d, (int_nvvm_isspacep_local Int64Regs:$a))]>;
1720 def ISSPACEP_SHARED_32
1721 : NVPTXInst<(outs Int1Regs:$d), (ins Int32Regs:$a),
1722 "isspacep.shared \t$d, $a;",
1723 [(set Int1Regs:$d, (int_nvvm_isspacep_shared Int32Regs:$a))]>;
1724 def ISSPACEP_SHARED_64
1725 : NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
1726 "isspacep.shared \t$d, $a;",
1727 [(set Int1Regs:$d, (int_nvvm_isspacep_shared Int64Regs:$a))]>;
1730 // Special register reads
1731 def MOV_SPECIAL : NVPTXInst<(outs Int32Regs:$d),
1732 (ins SpecialRegs:$r),
1733 "mov.b32\t$d, $r;", []>;
1735 def : Pat<(int_nvvm_read_ptx_sreg_envreg0), (MOV_SPECIAL ENVREG0)>;
1736 def : Pat<(int_nvvm_read_ptx_sreg_envreg1), (MOV_SPECIAL ENVREG1)>;
1737 def : Pat<(int_nvvm_read_ptx_sreg_envreg2), (MOV_SPECIAL ENVREG2)>;
1738 def : Pat<(int_nvvm_read_ptx_sreg_envreg3), (MOV_SPECIAL ENVREG3)>;
1739 def : Pat<(int_nvvm_read_ptx_sreg_envreg4), (MOV_SPECIAL ENVREG4)>;
1740 def : Pat<(int_nvvm_read_ptx_sreg_envreg5), (MOV_SPECIAL ENVREG5)>;
1741 def : Pat<(int_nvvm_read_ptx_sreg_envreg6), (MOV_SPECIAL ENVREG6)>;
1742 def : Pat<(int_nvvm_read_ptx_sreg_envreg7), (MOV_SPECIAL ENVREG7)>;
1743 def : Pat<(int_nvvm_read_ptx_sreg_envreg8), (MOV_SPECIAL ENVREG8)>;
1744 def : Pat<(int_nvvm_read_ptx_sreg_envreg9), (MOV_SPECIAL ENVREG9)>;
1745 def : Pat<(int_nvvm_read_ptx_sreg_envreg10), (MOV_SPECIAL ENVREG10)>;
1746 def : Pat<(int_nvvm_read_ptx_sreg_envreg11), (MOV_SPECIAL ENVREG11)>;
1747 def : Pat<(int_nvvm_read_ptx_sreg_envreg12), (MOV_SPECIAL ENVREG12)>;
1748 def : Pat<(int_nvvm_read_ptx_sreg_envreg13), (MOV_SPECIAL ENVREG13)>;
1749 def : Pat<(int_nvvm_read_ptx_sreg_envreg14), (MOV_SPECIAL ENVREG14)>;
1750 def : Pat<(int_nvvm_read_ptx_sreg_envreg15), (MOV_SPECIAL ENVREG15)>;
1751 def : Pat<(int_nvvm_read_ptx_sreg_envreg16), (MOV_SPECIAL ENVREG16)>;
1752 def : Pat<(int_nvvm_read_ptx_sreg_envreg17), (MOV_SPECIAL ENVREG17)>;
1753 def : Pat<(int_nvvm_read_ptx_sreg_envreg18), (MOV_SPECIAL ENVREG18)>;
1754 def : Pat<(int_nvvm_read_ptx_sreg_envreg19), (MOV_SPECIAL ENVREG19)>;
1755 def : Pat<(int_nvvm_read_ptx_sreg_envreg20), (MOV_SPECIAL ENVREG20)>;
1756 def : Pat<(int_nvvm_read_ptx_sreg_envreg21), (MOV_SPECIAL ENVREG21)>;
1757 def : Pat<(int_nvvm_read_ptx_sreg_envreg22), (MOV_SPECIAL ENVREG22)>;
1758 def : Pat<(int_nvvm_read_ptx_sreg_envreg23), (MOV_SPECIAL ENVREG23)>;
1759 def : Pat<(int_nvvm_read_ptx_sreg_envreg24), (MOV_SPECIAL ENVREG24)>;
1760 def : Pat<(int_nvvm_read_ptx_sreg_envreg25), (MOV_SPECIAL ENVREG25)>;
1761 def : Pat<(int_nvvm_read_ptx_sreg_envreg26), (MOV_SPECIAL ENVREG26)>;
1762 def : Pat<(int_nvvm_read_ptx_sreg_envreg27), (MOV_SPECIAL ENVREG27)>;
1763 def : Pat<(int_nvvm_read_ptx_sreg_envreg28), (MOV_SPECIAL ENVREG28)>;
1764 def : Pat<(int_nvvm_read_ptx_sreg_envreg29), (MOV_SPECIAL ENVREG29)>;
1765 def : Pat<(int_nvvm_read_ptx_sreg_envreg30), (MOV_SPECIAL ENVREG30)>;
1766 def : Pat<(int_nvvm_read_ptx_sreg_envreg31), (MOV_SPECIAL ENVREG31)>;
1769 //-----------------------------------
1770 // Texture Intrinsics
1771 //-----------------------------------
1773 // NOTE: For Fermi support, any new texture/surface/sampler intrinsics must be
1774 // also defined in NVPTXReplaceImageHandles.cpp
1777 // Texture fetch instructions using handles
1779 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1780 Float32Regs:$b, Float32Regs:$a),
1781 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x),
1782 "tex.1d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
1785 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1786 Float32Regs:$b, Float32Regs:$a),
1787 (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x),
1788 "tex.1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
1790 def TEX_1D_F32_F32_LEVEL
1791 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1792 Float32Regs:$b, Float32Regs:$a),
1793 (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$lod),
1794 "tex.level.1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1795 "[$t, $s, \\{$x\\}], $lod;",
1797 def TEX_1D_F32_F32_GRAD
1798 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1799 Float32Regs:$b, Float32Regs:$a),
1800 (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x,
1801 Float32Regs:$gradx, Float32Regs:$grady),
1802 "tex.grad.1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1803 "[$t, $s, \\{$x\\}], \\{$gradx\\}, \\{$grady\\};",
1806 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1807 Int32Regs:$b, Int32Regs:$a),
1808 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x),
1809 "tex.1d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
1812 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1813 Int32Regs:$b, Int32Regs:$a),
1814 (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x),
1815 "tex.1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, [$t, $s, \\{$x\\}];",
1817 def TEX_1D_I32_F32_LEVEL
1818 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1819 Int32Regs:$b, Int32Regs:$a),
1820 (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x,
1822 "tex.level.1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
1823 "[$t, $s, \\{$x\\}], $lod;",
1825 def TEX_1D_I32_F32_GRAD
1826 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1827 Int32Regs:$b, Int32Regs:$a),
1828 (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x,
1829 Float32Regs:$gradx, Float32Regs:$grady),
1830 "tex.grad.1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
1831 "[$t, $s, \\{$x\\}], \\{$gradx\\}, \\{$grady\\};",
1834 def TEX_1D_ARRAY_F32_I32
1835 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1836 Float32Regs:$b, Float32Regs:$a),
1837 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
1838 "tex.a1d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
1839 "[$t, $s, \\{$l, $x\\}];",
1841 def TEX_1D_ARRAY_F32_F32
1842 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1843 Float32Regs:$b, Float32Regs:$a),
1844 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x),
1845 "tex.a1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1846 "[$t, $s, \\{$l, $x\\}];",
1848 def TEX_1D_ARRAY_F32_F32_LEVEL
1849 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1850 Float32Regs:$b, Float32Regs:$a),
1851 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
1853 "tex.level.a1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1854 "[$t, $s, \\{$l, $x\\}], $lod;",
1856 def TEX_1D_ARRAY_F32_F32_GRAD
1857 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1858 Float32Regs:$b, Float32Regs:$a),
1859 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
1860 Float32Regs:$gradx, Float32Regs:$grady),
1861 "tex.grad.a1d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1862 "[$t, $s, \\{$l, $x\\}], \\{$gradx\\}, \\{$grady\\};",
1864 def TEX_1D_ARRAY_I32_I32
1865 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1866 Int32Regs:$b, Int32Regs:$a),
1867 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
1868 "tex.a1d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
1869 "[$t, $s, \\{$l, $x\\}];",
1871 def TEX_1D_ARRAY_I32_F32
1872 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1873 Int32Regs:$b, Int32Regs:$a),
1874 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x),
1875 "tex.a1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
1876 "[$t, $s, \\{$l, $x\\}];",
1878 def TEX_1D_ARRAY_I32_F32_LEVEL
1879 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1880 Int32Regs:$b, Int32Regs:$a),
1881 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
1883 "tex.level.a1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
1884 "[$t, $s, \\{$l, $x\\}], $lod;",
1886 def TEX_1D_ARRAY_I32_F32_GRAD
1887 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1888 Int32Regs:$b, Int32Regs:$a),
1889 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
1890 Float32Regs:$gradx, Float32Regs:$grady),
1891 "tex.grad.a1d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
1892 "[$t, $s, \\{$l, $x\\}], \\{$gradx\\}, \\{$grady\\};",
1896 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1897 Float32Regs:$b, Float32Regs:$a),
1898 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
1899 "tex.2d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
1900 "[$t, $s, \\{$x, $y\\}];",
1903 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1904 Float32Regs:$b, Float32Regs:$a),
1905 (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
1906 "tex.2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1907 "[$t, $s, \\{$x, $y\\}];",
1909 def TEX_2D_F32_F32_LEVEL
1910 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1911 Float32Regs:$b, Float32Regs:$a),
1912 (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
1914 "tex.level.2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1915 "[$t, $s, \\{$x, $y\\}], $lod;",
1917 def TEX_2D_F32_F32_GRAD
1918 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1919 Float32Regs:$b, Float32Regs:$a),
1920 (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
1921 Float32Regs:$gradx0, Float32Regs:$gradx1,
1922 Float32Regs:$grady0, Float32Regs:$grady1),
1923 "tex.grad.2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1924 "[$t, $s, \\{$x, $y\\}], \\{$gradx0, $gradx1\\}, "
1925 "\\{$grady0, $grady1\\};",
1928 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1929 Int32Regs:$b, Int32Regs:$a),
1930 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
1931 "tex.2d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
1932 "[$t, $s, \\{$x, $y\\}];",
1935 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1936 Int32Regs:$b, Int32Regs:$a),
1937 (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y),
1938 "tex.2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
1939 "[$t, $s, \\{$x, $y\\}];",
1941 def TEX_2D_I32_F32_LEVEL
1942 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1943 Int32Regs:$b, Int32Regs:$a),
1944 (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
1946 "tex.level.2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
1947 "[$t, $s, \\{$x, $y\\}], $lod;",
1949 def TEX_2D_I32_F32_GRAD
1950 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1951 Int32Regs:$b, Int32Regs:$a),
1952 (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
1953 Float32Regs:$gradx0, Float32Regs:$gradx1,
1954 Float32Regs:$grady0, Float32Regs:$grady1),
1955 "tex.grad.2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
1956 "[$t, $s, \\{$x, $y\\}], \\{$gradx0, $gradx1\\}, "
1957 "\\{$grady0, $grady1\\};",
1960 def TEX_2D_ARRAY_F32_I32
1961 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1962 Float32Regs:$b, Float32Regs:$a),
1963 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
1965 "tex.a2d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
1966 "[$t, $s, \\{$l, $x, $y, $y\\}];",
1968 def TEX_2D_ARRAY_F32_F32
1969 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1970 Float32Regs:$b, Float32Regs:$a),
1971 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
1973 "tex.a2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1974 "[$t, $s, \\{$l, $x, $y, $y\\}];",
1976 def TEX_2D_ARRAY_F32_F32_LEVEL
1977 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1978 Float32Regs:$b, Float32Regs:$a),
1979 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
1980 Float32Regs:$y, Float32Regs:$lod),
1981 "tex.level.a2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1982 "[$t, $s, \\{$l, $x, $y, $y\\}], $lod;",
1984 def TEX_2D_ARRAY_F32_F32_GRAD
1985 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
1986 Float32Regs:$b, Float32Regs:$a),
1987 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
1988 Float32Regs:$y, Float32Regs:$gradx0, Float32Regs:$gradx1,
1989 Float32Regs:$grady0, Float32Regs:$grady1),
1990 "tex.grad.a2d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
1991 "[$t, $s, \\{$l, $x, $y, $y\\}], \\{$gradx0, $gradx1\\}, "
1992 "\\{$grady0, $grady1\\};",
1994 def TEX_2D_ARRAY_I32_I32
1995 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
1996 Int32Regs:$b, Int32Regs:$a),
1997 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
1999 "tex.a2d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
2000 "[$t, $s, \\{$l, $x, $y, $y\\}];",
2002 def TEX_2D_ARRAY_I32_F32
2003 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
2004 Int32Regs:$b, Int32Regs:$a),
2005 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
2007 "tex.a2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
2008 "[$t, $s, \\{$l, $x, $y, $y\\}];",
2010 def TEX_2D_ARRAY_I32_F32_LEVEL
2011 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
2012 Int32Regs:$b, Int32Regs:$a),
2013 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
2014 Float32Regs:$y, Float32Regs:$lod),
2015 "tex.level.a2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
2016 "[$t, $s, \\{$l, $x, $y, $y\\}], $lod;",
2018 def TEX_2D_ARRAY_I32_F32_GRAD
2019 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
2020 Int32Regs:$b, Int32Regs:$a),
2021 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$l, Float32Regs:$x,
2023 Float32Regs:$gradx0, Float32Regs:$gradx1,
2024 Float32Regs:$grady0, Float32Regs:$grady1),
2025 "tex.grad.a2d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
2026 "[$t, $s, \\{$l, $x, $y, $y\\}], \\{$gradx0, $gradx1\\}, "
2027 "\\{$grady0, $grady1\\};",
2031 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
2032 Float32Regs:$b, Float32Regs:$a),
2033 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
2035 "tex.3d.v4.f32.s32\t\\{$r, $g, $b, $a\\}, "
2036 "[$t, $s, \\{$x, $y, $z, $z\\}];",
2039 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
2040 Float32Regs:$b, Float32Regs:$a),
2041 (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
2043 "tex.3d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
2044 "[$t, $s, \\{$x, $y, $z, $z\\}];",
2046 def TEX_3D_F32_F32_LEVEL
2047 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
2048 Float32Regs:$b, Float32Regs:$a),
2049 (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
2050 Float32Regs:$z, Float32Regs:$lod),
2051 "tex.level.3d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
2052 "[$t, $s, \\{$x, $y, $z, $z\\}], $lod;",
2054 def TEX_3D_F32_F32_GRAD
2055 : NVPTXInst<(outs Float32Regs:$r, Float32Regs:$g,
2056 Float32Regs:$b, Float32Regs:$a),
2057 (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
2059 Float32Regs:$gradx0, Float32Regs:$gradx1,
2060 Float32Regs:$gradx2, Float32Regs:$grady0,
2061 Float32Regs:$grady1, Float32Regs:$grady2),
2062 "tex.grad.3d.v4.f32.f32\t\\{$r, $g, $b, $a\\}, "
2063 "[$t, $s, \\{$x, $y, $z, $z\\}], "
2064 "\\{$gradx0, $gradx1, $gradx2, $gradx2\\}, "
2065 "\\{$grady0, $grady1, $grady2, $grady2\\};",
2068 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
2069 Int32Regs:$b, Int32Regs:$a),
2070 (ins Int64Regs:$t, Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
2072 "tex.3d.v4.s32.s32\t\\{$r, $g, $b, $a\\}, "
2073 "[$t, $s, \\{$x, $y, $z, $z\\}];",
2076 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
2077 Int32Regs:$b, Int32Regs:$a),
2078 (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
2080 "tex.3d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
2081 "[$t, $s, \\{$x, $y, $z, $z\\}];",
2083 def TEX_3D_I32_F32_LEVEL
2084 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
2085 Int32Regs:$b, Int32Regs:$a),
2086 (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
2087 Float32Regs:$z, Float32Regs:$lod),
2088 "tex.level.3d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
2089 "[$t, $s, \\{$x, $y, $z, $z\\}], $lod;",
2091 def TEX_3D_I32_F32_GRAD
2092 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g,
2093 Int32Regs:$b, Int32Regs:$a),
2094 (ins Int64Regs:$t, Int64Regs:$s, Float32Regs:$x, Float32Regs:$y,
2096 Float32Regs:$gradx0, Float32Regs:$gradx1,
2097 Float32Regs:$gradx2, Float32Regs:$grady0,
2098 Float32Regs:$grady1, Float32Regs:$grady2),
2099 "tex.grad.3d.v4.s32.f32\t\\{$r, $g, $b, $a\\}, "
2100 "[$t, $s, \\{$x, $y, $z, $z\\}], "
2101 "\\{$gradx0, $gradx1, $gradx2, $gradx2\\}, "
2102 "\\{$grady0, $grady1, $grady2, $grady2\\};",
2106 // Surface load instructions
2108 : NVPTXInst<(outs Int16Regs:$r),
2109 (ins Int64Regs:$s, Int32Regs:$x),
2110 "suld.b.1d.b8.trap \\{$r\\}, [$s, \\{$x\\}];",
2112 def SULD_1D_I16_TRAP
2113 : NVPTXInst<(outs Int16Regs:$r),
2114 (ins Int64Regs:$s, Int32Regs:$x),
2115 "suld.b.1d.b16.trap \\{$r\\}, [$s, \\{$x\\}];",
2117 def SULD_1D_I32_TRAP
2118 : NVPTXInst<(outs Int32Regs:$r),
2119 (ins Int64Regs:$s, Int32Regs:$x),
2120 "suld.b.1d.b32.trap \\{$r\\}, [$s, \\{$x\\}];",
2122 def SULD_1D_V2I8_TRAP
2123 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2124 (ins Int64Regs:$s, Int32Regs:$x),
2125 "suld.b.1d.v2.b8.trap \\{$r, $g\\}, [$s, \\{$x\\}];",
2127 def SULD_1D_V2I16_TRAP
2128 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2129 (ins Int64Regs:$s, Int32Regs:$x),
2130 "suld.b.1d.v2.b16.trap \\{$r, $g\\}, [$s, \\{$x\\}];",
2132 def SULD_1D_V2I32_TRAP
2133 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g),
2134 (ins Int64Regs:$s, Int32Regs:$x),
2135 "suld.b.1d.v2.b32.trap \\{$r, $g\\}, [$s, \\{$x\\}];",
2137 def SULD_1D_V4I8_TRAP
2138 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2139 (ins Int64Regs:$s, Int32Regs:$x),
2140 "suld.b.1d.v4.b8.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x\\}];",
2142 def SULD_1D_V4I16_TRAP
2143 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2144 (ins Int64Regs:$s, Int32Regs:$x),
2145 "suld.b.1d.v4.b16.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x\\}];",
2147 def SULD_1D_V4I32_TRAP
2148 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2149 (ins Int64Regs:$s, Int32Regs:$x),
2150 "suld.b.1d.v4.b32.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x\\}];",
2153 def SULD_1D_ARRAY_I8_TRAP
2154 : NVPTXInst<(outs Int16Regs:$r),
2155 (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
2156 "suld.b.a1d.b8.trap \\{$r\\}, [$s, \\{$l, $x\\}];",
2158 def SULD_1D_ARRAY_I16_TRAP
2159 : NVPTXInst<(outs Int16Regs:$r),
2160 (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
2161 "suld.b.a1d.b16.trap \\{$r\\}, [$s, \\{$l, $x\\}];",
2163 def SULD_1D_ARRAY_I32_TRAP
2164 : NVPTXInst<(outs Int32Regs:$r),
2165 (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
2166 "suld.b.a1d.b32.trap \\{$r\\}, [$s, \\{$l, $x\\}];",
2168 def SULD_1D_ARRAY_V2I8_TRAP
2169 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2170 (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
2171 "suld.b.a1d.v2.b8.trap \\{$r, $g\\}, [$s, \\{$l, $x\\}];",
2173 def SULD_1D_ARRAY_V2I16_TRAP
2174 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2175 (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
2176 "suld.b.a1d.v2.b16.trap \\{$r, $g\\}, [$s, \\{$l, $x\\}];",
2178 def SULD_1D_ARRAY_V2I32_TRAP
2179 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g),
2180 (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
2181 "suld.b.a1d.v2.b32.trap \\{$r, $g\\}, [$s, \\{$l, $x\\}];",
2183 def SULD_1D_ARRAY_V4I8_TRAP
2184 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2185 (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
2186 "suld.b.a1d.v4.b8.trap \\{$r, $g, $b, $a\\}, "
2187 "[$s, \\{$l, $x\\}];",
2189 def SULD_1D_ARRAY_V4I16_TRAP
2190 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2191 (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
2192 "suld.b.a1d.v4.b16.trap \\{$r, $g, $b, $a\\}, "
2193 "[$s, \\{$l, $x\\}];",
2195 def SULD_1D_ARRAY_V4I32_TRAP
2196 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2197 (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x),
2198 "suld.b.a1d.v4.b32.trap \\{$r, $g, $b, $a\\}, "
2199 "[$s, \\{$l, $x\\}];",
2203 : NVPTXInst<(outs Int16Regs:$r),
2204 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
2205 "suld.b.2d.b8.trap \\{$r\\}, [$s, \\{$x, $y\\}];",
2207 def SULD_2D_I16_TRAP
2208 : NVPTXInst<(outs Int16Regs:$r),
2209 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
2210 "suld.b.2d.b16.trap \\{$r\\}, [$s, \\{$x, $y\\}];",
2212 def SULD_2D_I32_TRAP
2213 : NVPTXInst<(outs Int32Regs:$r),
2214 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
2215 "suld.b.2d.b32.trap \\{$r\\}, [$s, \\{$x, $y\\}];",
2217 def SULD_2D_V2I8_TRAP
2218 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2219 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
2220 "suld.b.2d.v2.b8.trap \\{$r, $g\\}, [$s, \\{$x, $y\\}];",
2222 def SULD_2D_V2I16_TRAP
2223 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2224 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
2225 "suld.b.2d.v2.b16.trap \\{$r, $g\\}, [$s, \\{$x, $y\\}];",
2227 def SULD_2D_V2I32_TRAP
2228 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g),
2229 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
2230 "suld.b.2d.v2.b32.trap \\{$r, $g\\}, [$s, \\{$x, $y\\}];",
2232 def SULD_2D_V4I8_TRAP
2233 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2234 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
2235 "suld.b.2d.v4.b8.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x, $y\\}];",
2237 def SULD_2D_V4I16_TRAP
2238 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2239 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
2240 "suld.b.2d.v4.b16.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x, $y\\}];",
2242 def SULD_2D_V4I32_TRAP
2243 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2244 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y),
2245 "suld.b.2d.v4.b32.trap \\{$r, $g, $b, $a\\}, [$s, \\{$x, $y\\}];",
2248 def SULD_2D_ARRAY_I8_TRAP
2249 : NVPTXInst<(outs Int16Regs:$r),
2250 (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
2251 "suld.b.a2d.b8.trap \\{$r\\}, [$s, \\{$l, $x, $y, $y\\}];",
2253 def SULD_2D_ARRAY_I16_TRAP
2254 : NVPTXInst<(outs Int16Regs:$r),
2255 (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
2256 "suld.b.a2d.b16.trap \\{$r\\}, [$s, \\{$l, $x, $y, $y\\}];",
2258 def SULD_2D_ARRAY_I32_TRAP
2259 : NVPTXInst<(outs Int32Regs:$r),
2260 (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
2261 "suld.b.a2d.b32.trap \\{$r\\}, [$s, \\{$l, $x, $y, $y\\}];",
2263 def SULD_2D_ARRAY_V2I8_TRAP
2264 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2265 (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
2266 "suld.b.a2d.v2.b8.trap \\{$r, $g\\}, "
2267 "[$s, \\{$l, $x, $y, $y\\}];",
2269 def SULD_2D_ARRAY_V2I16_TRAP
2270 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2271 (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
2272 "suld.b.a2d.v2.b16.trap \\{$r, $g\\}, "
2273 "[$s, \\{$l, $x, $y, $y\\}];",
2275 def SULD_2D_ARRAY_V2I32_TRAP
2276 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g),
2277 (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
2278 "suld.b.a2d.v2.b32.trap \\{$r, $g\\}, "
2279 "[$s, \\{$l, $x, $y, $y\\}];",
2281 def SULD_2D_ARRAY_V4I8_TRAP
2282 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2283 (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
2284 "suld.b.a2d.v4.b8.trap \\{$r, $g, $b, $a\\}, "
2285 "[$s, \\{$l, $x, $y, $y\\}];",
2287 def SULD_2D_ARRAY_V4I16_TRAP
2288 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2289 (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
2290 "suld.b.a2d.v4.b16.trap \\{$r, $g, $b, $a\\}, "
2291 "[$s, \\{$l, $x, $y, $y\\}];",
2293 def SULD_2D_ARRAY_V4I32_TRAP
2294 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2295 (ins Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y),
2296 "suld.b.a2d.v4.b32.trap \\{$r, $g, $b, $a\\}, "
2297 "[$s, \\{$l, $x, $y, $y\\}];",
2301 : NVPTXInst<(outs Int16Regs:$r),
2302 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
2303 "suld.b.3d.b8.trap \\{$r\\}, [$s, \\{$x, $y, $z, $z\\}];",
2305 def SULD_3D_I16_TRAP
2306 : NVPTXInst<(outs Int16Regs:$r),
2307 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
2308 "suld.b.3d.b16.trap \\{$r\\}, [$s, \\{$x, $y, $z, $z\\}];",
2310 def SULD_3D_I32_TRAP
2311 : NVPTXInst<(outs Int32Regs:$r),
2312 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
2313 "suld.b.3d.b32.trap \\{$r\\}, [$s, \\{$x, $y, $z, $z\\}];",
2315 def SULD_3D_V2I8_TRAP
2316 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2317 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
2318 "suld.b.3d.v2.b8.trap \\{$r, $g\\}, [$s, \\{$x, $y, $z, $z\\}];",
2320 def SULD_3D_V2I16_TRAP
2321 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g),
2322 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
2323 "suld.b.3d.v2.b16.trap \\{$r, $g\\}, [$s, \\{$x, $y, $z, $z\\}];",
2325 def SULD_3D_V2I32_TRAP
2326 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g),
2327 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
2328 "suld.b.3d.v2.b32.trap \\{$r, $g\\}, [$s, \\{$x, $y, $z, $z\\}];",
2330 def SULD_3D_V4I8_TRAP
2331 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2332 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
2333 "suld.b.3d.v4.b8.trap \\{$r, $g, $b, $a\\}, "
2334 "[$s, \\{$x, $y, $z, $z\\}];",
2336 def SULD_3D_V4I16_TRAP
2337 : NVPTXInst<(outs Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2338 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
2339 "suld.b.3d.v4.b16.trap \\{$r, $g, $b, $a\\}, "
2340 "[$s, \\{$x, $y, $z, $z\\}];",
2342 def SULD_3D_V4I32_TRAP
2343 : NVPTXInst<(outs Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2344 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z),
2345 "suld.b.3d.v4.b32.trap \\{$r, $g, $b, $a\\}, "
2346 "[$s, \\{$x, $y, $z, $z\\}];",
2350 //-----------------------------------
2351 // Texture Query Intrinsics
2352 //-----------------------------------
2353 def TXQ_CHANNEL_ORDER
2354 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2355 "txq.channel_order.b32 \t$d, [$a];",
2357 def TXQ_CHANNEL_DATA_TYPE
2358 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2359 "txq.channel_data_type.b32 \t$d, [$a];",
2362 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2363 "txq.width.b32 \t$d, [$a];",
2366 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2367 "txq.height.b32 \t$d, [$a];",
2370 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2371 "txq.depth.b32 \t$d, [$a];",
2374 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2375 "txq.array_size.b32 \t$d, [$a];",
2378 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2379 "txq.num_samples.b32 \t$d, [$a];",
2381 def TXQ_NUM_MIPMAP_LEVELS
2382 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2383 "txq.num_mipmap_levels.b32 \t$d, [$a];",
2386 def : Pat<(int_nvvm_txq_channel_order Int64Regs:$a),
2387 (TXQ_CHANNEL_ORDER Int64Regs:$a)>;
2388 def : Pat<(int_nvvm_txq_channel_data_type Int64Regs:$a),
2389 (TXQ_CHANNEL_DATA_TYPE Int64Regs:$a)>;
2390 def : Pat<(int_nvvm_txq_width Int64Regs:$a),
2391 (TXQ_WIDTH Int64Regs:$a)>;
2392 def : Pat<(int_nvvm_txq_height Int64Regs:$a),
2393 (TXQ_HEIGHT Int64Regs:$a)>;
2394 def : Pat<(int_nvvm_txq_depth Int64Regs:$a),
2395 (TXQ_DEPTH Int64Regs:$a)>;
2396 def : Pat<(int_nvvm_txq_array_size Int64Regs:$a),
2397 (TXQ_ARRAY_SIZE Int64Regs:$a)>;
2398 def : Pat<(int_nvvm_txq_num_samples Int64Regs:$a),
2399 (TXQ_NUM_SAMPLES Int64Regs:$a)>;
2400 def : Pat<(int_nvvm_txq_num_mipmap_levels Int64Regs:$a),
2401 (TXQ_NUM_MIPMAP_LEVELS Int64Regs:$a)>;
2404 //-----------------------------------
2405 // Surface Query Intrinsics
2406 //-----------------------------------
2407 def SUQ_CHANNEL_ORDER
2408 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2409 "suq.channel_order.b32 \t$d, [$a];",
2411 def SUQ_CHANNEL_DATA_TYPE
2412 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2413 "suq.channel_data_type.b32 \t$d, [$a];",
2416 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2417 "suq.width.b32 \t$d, [$a];",
2420 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2421 "suq.height.b32 \t$d, [$a];",
2424 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2425 "suq.depth.b32 \t$d, [$a];",
2428 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
2429 "suq.array_size.b32 \t$d, [$a];",
2432 def : Pat<(int_nvvm_suq_channel_order Int64Regs:$a),
2433 (SUQ_CHANNEL_ORDER Int64Regs:$a)>;
2434 def : Pat<(int_nvvm_suq_channel_data_type Int64Regs:$a),
2435 (SUQ_CHANNEL_DATA_TYPE Int64Regs:$a)>;
2436 def : Pat<(int_nvvm_suq_width Int64Regs:$a),
2437 (SUQ_WIDTH Int64Regs:$a)>;
2438 def : Pat<(int_nvvm_suq_height Int64Regs:$a),
2439 (SUQ_HEIGHT Int64Regs:$a)>;
2440 def : Pat<(int_nvvm_suq_depth Int64Regs:$a),
2441 (SUQ_DEPTH Int64Regs:$a)>;
2442 def : Pat<(int_nvvm_suq_array_size Int64Regs:$a),
2443 (SUQ_ARRAY_SIZE Int64Regs:$a)>;
2446 //===- Handle Query -------------------------------------------------------===//
2448 // TODO: These intrinsics are not yet finalized, pending PTX ISA design work
2450 : NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
2451 "istypep.samplerref \t$d, $a;",
2452 [(set Int1Regs:$d, (int_nvvm_istypep_sampler Int64Regs:$a))]>;
2454 : NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
2455 "istypep.surfref \t$d, $a;",
2456 [(set Int1Regs:$d, (int_nvvm_istypep_surface Int64Regs:$a))]>;
2458 : NVPTXInst<(outs Int1Regs:$d), (ins Int64Regs:$a),
2459 "istypep.texref \t$d, $a;",
2460 [(set Int1Regs:$d, (int_nvvm_istypep_texture Int64Regs:$a))]>;
2462 //===- Surface Stores -----------------------------------------------------===//
2466 def SUST_B_1D_B8_TRAP
2468 (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
2469 "sust.b.1d.b8.trap \t[$s, \\{$x\\}], \\{$r\\};",
2471 def SUST_B_1D_B16_TRAP
2473 (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
2474 "sust.b.1d.b16.trap \t[$s, \\{$x\\}], \\{$r\\};",
2476 def SUST_B_1D_B32_TRAP
2478 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$r),
2479 "sust.b.1d.b32.trap \t[$s, \\{$x\\}], \\{$r\\};",
2481 def SUST_B_1D_V2B8_TRAP
2483 (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
2484 "sust.b.1d.v2.b8.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
2486 def SUST_B_1D_V2B16_TRAP
2488 (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
2489 "sust.b.1d.v2.b16.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
2491 def SUST_B_1D_V2B32_TRAP
2493 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g),
2494 "sust.b.1d.v2.b32.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
2496 def SUST_B_1D_V4B8_TRAP
2498 (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g,
2499 Int16Regs:$b, Int16Regs:$a),
2500 "sust.b.1d.v4.b8.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
2502 def SUST_B_1D_V4B16_TRAP
2504 (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g,
2505 Int16Regs:$b, Int16Regs:$a),
2506 "sust.b.1d.v4.b16.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
2508 def SUST_B_1D_V4B32_TRAP
2510 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g,
2511 Int32Regs:$b, Int32Regs:$a),
2512 "sust.b.1d.v4.b32.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
2516 def SUST_B_1D_ARRAY_B8_TRAP
2518 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r),
2519 "sust.b.a1d.b8.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
2521 def SUST_B_1D_ARRAY_B16_TRAP
2523 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r),
2524 "sust.b.a1d.b16.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
2526 def SUST_B_1D_ARRAY_B32_TRAP
2528 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$r),
2529 "sust.b.a1d.b32.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
2531 def SUST_B_1D_ARRAY_V2B8_TRAP
2533 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
2535 "sust.b.a1d.v2.b8.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
2537 def SUST_B_1D_ARRAY_V2B16_TRAP
2539 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
2541 "sust.b.a1d.v2.b16.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
2543 def SUST_B_1D_ARRAY_V2B32_TRAP
2545 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$r,
2547 "sust.b.a1d.v2.b32.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
2549 def SUST_B_1D_ARRAY_V4B8_TRAP
2551 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
2552 Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2553 "sust.b.a1d.v4.b8.trap \t[$s, \\{$idx, $x\\}], "
2554 "\\{$r, $g, $b, $a\\};",
2556 def SUST_B_1D_ARRAY_V4B16_TRAP
2558 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
2559 Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2560 "sust.b.a1d.v4.b16.trap \t[$s, \\{$idx, $x\\}], "
2561 "\\{$r, $g, $b, $a\\};",
2563 def SUST_B_1D_ARRAY_V4B32_TRAP
2565 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$r,
2566 Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2567 "sust.b.a1d.v4.b32.trap \t[$s, \\{$idx, $x\\}], "
2568 "\\{$r, $g, $b, $a\\};",
2572 def SUST_B_2D_B8_TRAP
2574 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
2575 "sust.b.2d.b8.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
2577 def SUST_B_2D_B16_TRAP
2579 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
2580 "sust.b.2d.b16.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
2582 def SUST_B_2D_B32_TRAP
2584 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r),
2585 "sust.b.2d.b32.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
2587 def SUST_B_2D_V2B8_TRAP
2589 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
2591 "sust.b.2d.v2.b8.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
2593 def SUST_B_2D_V2B16_TRAP
2595 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
2597 "sust.b.2d.v2.b16.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
2599 def SUST_B_2D_V2B32_TRAP
2601 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r,
2603 "sust.b.2d.v2.b32.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
2605 def SUST_B_2D_V4B8_TRAP
2607 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
2608 Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2609 "sust.b.2d.v4.b8.trap \t[$s, \\{$x, $y\\}], "
2610 "\\{$r, $g, $b, $a\\};",
2612 def SUST_B_2D_V4B16_TRAP
2614 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
2615 Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2616 "sust.b.2d.v4.b16.trap \t[$s, \\{$x, $y\\}], "
2617 "\\{$r, $g, $b, $a\\};",
2619 def SUST_B_2D_V4B32_TRAP
2621 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r,
2622 Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2623 "sust.b.2d.v4.b32.trap \t[$s, \\{$x, $y\\}], "
2624 "\\{$r, $g, $b, $a\\};",
2628 def SUST_B_2D_ARRAY_B8_TRAP
2630 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2632 "sust.b.a2d.b8.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
2634 def SUST_B_2D_ARRAY_B16_TRAP
2636 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2638 "sust.b.a2d.b16.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
2640 def SUST_B_2D_ARRAY_B32_TRAP
2642 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2644 "sust.b.a2d.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
2646 def SUST_B_2D_ARRAY_V2B8_TRAP
2648 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2649 Int16Regs:$r, Int16Regs:$g),
2650 "sust.b.a2d.v2.b8.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2653 def SUST_B_2D_ARRAY_V2B16_TRAP
2655 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2656 Int16Regs:$r, Int16Regs:$g),
2657 "sust.b.a2d.v2.b16.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2660 def SUST_B_2D_ARRAY_V2B32_TRAP
2662 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2663 Int32Regs:$r, Int32Regs:$g),
2664 "sust.b.a2d.v2.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2667 def SUST_B_2D_ARRAY_V4B8_TRAP
2669 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2670 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2671 "sust.b.a2d.v4.b8.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2672 "\\{$r, $g, $b, $a\\};",
2674 def SUST_B_2D_ARRAY_V4B16_TRAP
2676 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2677 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2678 "sust.b.a2d.v4.b16.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2679 "\\{$r, $g, $b, $a\\};",
2681 def SUST_B_2D_ARRAY_V4B32_TRAP
2683 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2684 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2685 "sust.b.a2d.v4.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2686 "\\{$r, $g, $b, $a\\};",
2690 def SUST_B_3D_B8_TRAP
2692 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2694 "sust.b.3d.b8.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
2696 def SUST_B_3D_B16_TRAP
2698 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2700 "sust.b.3d.b16.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
2702 def SUST_B_3D_B32_TRAP
2704 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2706 "sust.b.3d.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
2708 def SUST_B_3D_V2B8_TRAP
2710 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2711 Int16Regs:$r, Int16Regs:$g),
2712 "sust.b.3d.v2.b8.trap \t[$s, \\{$x, $y, $z, $z\\}], "
2715 def SUST_B_3D_V2B16_TRAP
2717 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2718 Int16Regs:$r, Int16Regs:$g),
2719 "sust.b.3d.v2.b16.trap \t[$s, \\{$x, $y, $z, $z\\}], "
2722 def SUST_B_3D_V2B32_TRAP
2724 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2725 Int32Regs:$r, Int32Regs:$g),
2726 "sust.b.3d.v2.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], "
2729 def SUST_B_3D_V4B8_TRAP
2731 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2732 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2733 "sust.b.3d.v4.b8.trap \t[$s, \\{$x, $y, $z, $z\\}], "
2734 "\\{$r, $g, $b, $a\\};",
2736 def SUST_B_3D_V4B16_TRAP
2738 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2739 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2740 "sust.b.3d.v4.b16.trap \t[$s, \\{$x, $y, $z, $z\\}], "
2741 "\\{$r, $g, $b, $a\\};",
2743 def SUST_B_3D_V4B32_TRAP
2745 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2746 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2747 "sust.b.3d.v4.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], "
2748 "\\{$r, $g, $b, $a\\};",
2753 def SUST_P_1D_B8_TRAP
2755 (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
2756 "sust.p.1d.b8.trap \t[$s, \\{$x\\}], \\{$r\\};",
2758 def SUST_P_1D_B16_TRAP
2760 (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
2761 "sust.p.1d.b16.trap \t[$s, \\{$x\\}], \\{$r\\};",
2763 def SUST_P_1D_B32_TRAP
2765 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$r),
2766 "sust.p.1d.b32.trap \t[$s, \\{$x\\}], \\{$r\\};",
2768 def SUST_P_1D_V2B8_TRAP
2770 (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
2771 "sust.p.1d.v2.b8.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
2773 def SUST_P_1D_V2B16_TRAP
2775 (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
2776 "sust.p.1d.v2.b16.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
2778 def SUST_P_1D_V2B32_TRAP
2780 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g),
2781 "sust.p.1d.v2.b32.trap \t[$s, \\{$x\\}], \\{$r, $g\\};",
2783 def SUST_P_1D_V4B8_TRAP
2785 (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g,
2786 Int16Regs:$b, Int16Regs:$a),
2787 "sust.p.1d.v4.b8.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
2789 def SUST_P_1D_V4B16_TRAP
2791 (ins Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g,
2792 Int16Regs:$b, Int16Regs:$a),
2793 "sust.p.1d.v4.b16.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
2795 def SUST_P_1D_V4B32_TRAP
2797 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g,
2798 Int32Regs:$b, Int32Regs:$a),
2799 "sust.p.1d.v4.b32.trap \t[$s, \\{$x\\}], \\{$r, $g, $b, $a\\};",
2803 def SUST_P_1D_ARRAY_B8_TRAP
2805 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r),
2806 "sust.p.a1d.b8.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
2808 def SUST_P_1D_ARRAY_B16_TRAP
2810 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r),
2811 "sust.p.a1d.b16.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
2813 def SUST_P_1D_ARRAY_B32_TRAP
2815 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$r),
2816 "sust.p.a1d.b32.trap \t[$s, \\{$idx, $x\\}], \\{$r\\};",
2818 def SUST_P_1D_ARRAY_V2B8_TRAP
2820 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
2822 "sust.p.a1d.v2.b8.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
2824 def SUST_P_1D_ARRAY_V2B16_TRAP
2826 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
2828 "sust.p.a1d.v2.b16.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
2830 def SUST_P_1D_ARRAY_V2B32_TRAP
2832 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$r,
2834 "sust.p.a1d.v2.b32.trap \t[$s, \\{$idx, $x\\}], \\{$r, $g\\};",
2836 def SUST_P_1D_ARRAY_V4B8_TRAP
2838 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
2839 Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2840 "sust.p.a1d.v4.b8.trap \t[$s, \\{$idx, $x\\}], "
2841 "\\{$r, $g, $b, $a\\};",
2843 def SUST_P_1D_ARRAY_V4B16_TRAP
2845 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int16Regs:$r,
2846 Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2847 "sust.p.a1d.v4.b16.trap \t[$s, \\{$idx, $x\\}], "
2848 "\\{$r, $g, $b, $a\\};",
2850 def SUST_P_1D_ARRAY_V4B32_TRAP
2852 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$r,
2853 Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2854 "sust.p.a1d.v4.b32.trap \t[$s, \\{$idx, $x\\}], "
2855 "\\{$r, $g, $b, $a\\};",
2859 def SUST_P_2D_B8_TRAP
2861 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
2862 "sust.p.2d.b8.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
2864 def SUST_P_2D_B16_TRAP
2866 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
2867 "sust.p.2d.b16.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
2869 def SUST_P_2D_B32_TRAP
2871 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r),
2872 "sust.p.2d.b32.trap \t[$s, \\{$x, $y\\}], \\{$r\\};",
2874 def SUST_P_2D_V2B8_TRAP
2876 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
2878 "sust.p.2d.v2.b8.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
2880 def SUST_P_2D_V2B16_TRAP
2882 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
2884 "sust.p.2d.v2.b16.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
2886 def SUST_P_2D_V2B32_TRAP
2888 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r,
2890 "sust.p.2d.v2.b32.trap \t[$s, \\{$x, $y\\}], \\{$r, $g\\};",
2892 def SUST_P_2D_V4B8_TRAP
2894 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
2895 Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2896 "sust.p.2d.v4.b8.trap \t[$s, \\{$x, $y\\}], "
2897 "\\{$r, $g, $b, $a\\};",
2899 def SUST_P_2D_V4B16_TRAP
2901 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r,
2902 Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2903 "sust.p.2d.v4.b16.trap \t[$s, \\{$x, $y\\}], "
2904 "\\{$r, $g, $b, $a\\};",
2906 def SUST_P_2D_V4B32_TRAP
2908 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r,
2909 Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2910 "sust.p.2d.v4.b32.trap \t[$s, \\{$x, $y\\}], "
2911 "\\{$r, $g, $b, $a\\};",
2915 def SUST_P_2D_ARRAY_B8_TRAP
2917 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2919 "sust.p.a2d.b8.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
2921 def SUST_P_2D_ARRAY_B16_TRAP
2923 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2925 "sust.p.a2d.b16.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
2927 def SUST_P_2D_ARRAY_B32_TRAP
2929 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2931 "sust.p.a2d.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], \\{$r\\};",
2933 def SUST_P_2D_ARRAY_V2B8_TRAP
2935 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2936 Int16Regs:$r, Int16Regs:$g),
2937 "sust.p.a2d.v2.b8.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2940 def SUST_P_2D_ARRAY_V2B16_TRAP
2942 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2943 Int16Regs:$r, Int16Regs:$g),
2944 "sust.p.a2d.v2.b16.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2947 def SUST_P_2D_ARRAY_V2B32_TRAP
2949 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2950 Int32Regs:$r, Int32Regs:$g),
2951 "sust.p.a2d.v2.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2954 def SUST_P_2D_ARRAY_V4B8_TRAP
2956 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2957 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2958 "sust.p.a2d.v4.b8.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2959 "\\{$r, $g, $b, $a\\};",
2961 def SUST_P_2D_ARRAY_V4B16_TRAP
2963 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2964 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
2965 "sust.p.a2d.v4.b16.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2966 "\\{$r, $g, $b, $a\\};",
2968 def SUST_P_2D_ARRAY_V4B32_TRAP
2970 (ins Int64Regs:$s, Int32Regs:$idx, Int32Regs:$x, Int32Regs:$y,
2971 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
2972 "sust.p.a2d.v4.b32.trap \t[$s, \\{$idx, $x, $y, $y\\}], "
2973 "\\{$r, $g, $b, $a\\};",
2977 def SUST_P_3D_B8_TRAP
2979 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2981 "sust.p.3d.b8.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
2983 def SUST_P_3D_B16_TRAP
2985 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2987 "sust.p.3d.b16.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
2989 def SUST_P_3D_B32_TRAP
2991 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2993 "sust.p.3d.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], \\{$r\\};",
2995 def SUST_P_3D_V2B8_TRAP
2997 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
2998 Int16Regs:$r, Int16Regs:$g),
2999 "sust.p.3d.v2.b8.trap \t[$s, \\{$x, $y, $z, $z\\}], "
3002 def SUST_P_3D_V2B16_TRAP
3004 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3005 Int16Regs:$r, Int16Regs:$g),
3006 "sust.p.3d.v2.b16.trap \t[$s, \\{$x, $y, $z, $z\\}], "
3009 def SUST_P_3D_V2B32_TRAP
3011 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3012 Int32Regs:$r, Int32Regs:$g),
3013 "sust.p.3d.v2.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], "
3016 def SUST_P_3D_V4B8_TRAP
3018 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3019 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3020 "sust.p.3d.v4.b8.trap \t[$s, \\{$x, $y, $z, $z\\}], "
3021 "\\{$r, $g, $b, $a\\};",
3023 def SUST_P_3D_V4B16_TRAP
3025 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3026 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3027 "sust.p.3d.v4.b16.trap \t[$s, \\{$x, $y, $z, $z\\}], "
3028 "\\{$r, $g, $b, $a\\};",
3030 def SUST_P_3D_V4B32_TRAP
3032 (ins Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3033 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3034 "sust.p.3d.v4.b32.trap \t[$s, \\{$x, $y, $z, $z\\}], "
3035 "\\{$r, $g, $b, $a\\};",
3039 // Surface store instruction patterns
3040 // I'm not sure why we can't just include these in the instruction definitions,
3041 // but TableGen complains of type errors :(
3043 def : Pat<(int_nvvm_sust_b_1d_i8_trap
3044 Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
3045 (SUST_B_1D_B8_TRAP Int64Regs:$s, Int32Regs:$x, Int16Regs:$r)>;
3047 def : Pat<(int_nvvm_sust_b_1d_i16_trap
3048 Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
3049 (SUST_B_1D_B16_TRAP Int64Regs:$s, Int32Regs:$x, Int16Regs:$r)>;
3051 def : Pat<(int_nvvm_sust_b_1d_i32_trap
3052 Int64Regs:$s, Int32Regs:$x, Int32Regs:$r),
3053 (SUST_B_1D_B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$r)>;
3055 def : Pat<(int_nvvm_sust_b_1d_v2i8_trap
3056 Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
3057 (SUST_B_1D_V2B8_TRAP Int64Regs:$s, Int32Regs:$x,
3058 Int16Regs:$r, Int16Regs:$g)>;
3060 def : Pat<(int_nvvm_sust_b_1d_v2i16_trap
3061 Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
3062 (SUST_B_1D_V2B16_TRAP Int64Regs:$s, Int32Regs:$x,
3063 Int16Regs:$r, Int16Regs:$g)>;
3065 def : Pat<(int_nvvm_sust_b_1d_v2i32_trap
3066 Int64Regs:$s, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g),
3067 (SUST_B_1D_V2B32_TRAP Int64Regs:$s, Int32Regs:$x,
3068 Int32Regs:$r, Int32Regs:$g)>;
3070 def : Pat<(int_nvvm_sust_b_1d_v4i8_trap
3071 Int64Regs:$s, Int32Regs:$x,
3072 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3073 (SUST_B_1D_V4B8_TRAP Int64Regs:$s, Int32Regs:$x,
3074 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3076 def : Pat<(int_nvvm_sust_b_1d_v4i16_trap
3077 Int64Regs:$s, Int32Regs:$x,
3078 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3079 (SUST_B_1D_V4B16_TRAP Int64Regs:$s, Int32Regs:$x,
3080 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3082 def : Pat<(int_nvvm_sust_b_1d_v4i32_trap
3083 Int64Regs:$s, Int32Regs:$x,
3084 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3085 (SUST_B_1D_V4B32_TRAP Int64Regs:$s, Int32Regs:$x,
3086 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3090 def : Pat<(int_nvvm_sust_b_1d_array_i8_trap
3091 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r),
3092 (SUST_B_1D_ARRAY_B8_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3095 def : Pat<(int_nvvm_sust_b_1d_array_i16_trap
3096 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r),
3097 (SUST_B_1D_ARRAY_B16_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3100 def : Pat<(int_nvvm_sust_b_1d_array_i32_trap
3101 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$r),
3102 (SUST_B_1D_ARRAY_B32_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3105 def : Pat<(int_nvvm_sust_b_1d_array_v2i8_trap
3106 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
3107 (SUST_B_1D_ARRAY_V2B8_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3108 Int16Regs:$r, Int16Regs:$g)>;
3110 def : Pat<(int_nvvm_sust_b_1d_array_v2i16_trap
3111 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
3112 (SUST_B_1D_ARRAY_V2B16_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3113 Int16Regs:$r, Int16Regs:$g)>;
3115 def : Pat<(int_nvvm_sust_b_1d_array_v2i32_trap
3116 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g),
3117 (SUST_B_1D_ARRAY_V2B32_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3118 Int32Regs:$r, Int32Regs:$g)>;
3120 def : Pat<(int_nvvm_sust_b_1d_array_v4i8_trap
3121 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3122 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3123 (SUST_B_1D_ARRAY_V4B8_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3124 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3126 def : Pat<(int_nvvm_sust_b_1d_array_v4i16_trap
3127 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3128 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3129 (SUST_B_1D_ARRAY_V4B16_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3130 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3132 def : Pat<(int_nvvm_sust_b_1d_array_v4i32_trap
3133 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3134 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3135 (SUST_B_1D_ARRAY_V4B32_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3136 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3140 def : Pat<(int_nvvm_sust_b_2d_i8_trap
3141 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
3142 (SUST_B_2D_B8_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3145 def : Pat<(int_nvvm_sust_b_2d_i16_trap
3146 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
3147 (SUST_B_2D_B16_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3150 def : Pat<(int_nvvm_sust_b_2d_i32_trap
3151 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r),
3152 (SUST_B_2D_B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3155 def : Pat<(int_nvvm_sust_b_2d_v2i8_trap
3156 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r, Int16Regs:$g),
3157 (SUST_B_2D_V2B8_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3158 Int16Regs:$r, Int16Regs:$g)>;
3160 def : Pat<(int_nvvm_sust_b_2d_v2i16_trap
3161 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r, Int16Regs:$g),
3162 (SUST_B_2D_V2B16_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3163 Int16Regs:$r, Int16Regs:$g)>;
3165 def : Pat<(int_nvvm_sust_b_2d_v2i32_trap
3166 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r, Int32Regs:$g),
3167 (SUST_B_2D_V2B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3168 Int32Regs:$r, Int32Regs:$g)>;
3170 def : Pat<(int_nvvm_sust_b_2d_v4i8_trap
3171 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3172 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3173 (SUST_B_2D_V4B8_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3174 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3176 def : Pat<(int_nvvm_sust_b_2d_v4i16_trap
3177 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3178 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3179 (SUST_B_2D_V4B16_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3180 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3182 def : Pat<(int_nvvm_sust_b_2d_v4i32_trap
3183 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3184 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3185 (SUST_B_2D_V4B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3186 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3190 def : Pat<(int_nvvm_sust_b_2d_array_i8_trap
3191 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
3192 (SUST_B_2D_ARRAY_B8_TRAP Int64Regs:$s,
3193 Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3196 def : Pat<(int_nvvm_sust_b_2d_array_i16_trap
3197 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
3198 (SUST_B_2D_ARRAY_B16_TRAP Int64Regs:$s,
3199 Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3202 def : Pat<(int_nvvm_sust_b_2d_array_i32_trap
3203 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r),
3204 (SUST_B_2D_ARRAY_B32_TRAP Int64Regs:$s,
3205 Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3208 def : Pat<(int_nvvm_sust_b_2d_array_v2i8_trap
3209 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3210 Int16Regs:$r, Int16Regs:$g),
3211 (SUST_B_2D_ARRAY_V2B8_TRAP Int64Regs:$s, Int32Regs:$l,
3212 Int32Regs:$x, Int32Regs:$y,
3213 Int16Regs:$r, Int16Regs:$g)>;
3215 def : Pat<(int_nvvm_sust_b_2d_array_v2i16_trap
3216 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3217 Int16Regs:$r, Int16Regs:$g),
3218 (SUST_B_2D_ARRAY_V2B16_TRAP Int64Regs:$s, Int32Regs:$l,
3219 Int32Regs:$x, Int32Regs:$y,
3220 Int16Regs:$r, Int16Regs:$g)>;
3222 def : Pat<(int_nvvm_sust_b_2d_array_v2i32_trap
3223 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r,
3225 (SUST_B_2D_ARRAY_V2B32_TRAP Int64Regs:$s, Int32Regs:$l,
3226 Int32Regs:$x, Int32Regs:$y, Int32Regs:$r, Int32Regs:$g)>;
3228 def : Pat<(int_nvvm_sust_b_2d_array_v4i8_trap
3229 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3230 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3231 (SUST_B_2D_ARRAY_V4B8_TRAP Int64Regs:$s,
3232 Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3233 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3235 def : Pat<(int_nvvm_sust_b_2d_array_v4i16_trap
3236 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3237 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3238 (SUST_B_2D_ARRAY_V4B16_TRAP Int64Regs:$s,
3239 Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3240 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3242 def : Pat<(int_nvvm_sust_b_2d_array_v4i32_trap
3243 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3244 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3245 (SUST_B_2D_ARRAY_V4B32_TRAP Int64Regs:$s, Int32Regs:$l,
3246 Int32Regs:$x, Int32Regs:$y,
3247 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3251 def : Pat<(int_nvvm_sust_b_3d_i8_trap
3252 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3254 (SUST_B_3D_B8_TRAP Int64Regs:$s,
3255 Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3258 def : Pat<(int_nvvm_sust_b_3d_i16_trap
3259 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3261 (SUST_B_3D_B16_TRAP Int64Regs:$s,
3262 Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3265 def : Pat<(int_nvvm_sust_b_3d_i32_trap
3266 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3268 (SUST_B_3D_B32_TRAP Int64Regs:$s,
3269 Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3272 def : Pat<(int_nvvm_sust_b_3d_v2i8_trap
3273 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3274 Int16Regs:$r, Int16Regs:$g),
3275 (SUST_B_3D_V2B8_TRAP Int64Regs:$s,
3276 Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3277 Int16Regs:$r, Int16Regs:$g)>;
3279 def : Pat<(int_nvvm_sust_b_3d_v2i16_trap
3280 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3281 Int16Regs:$r, Int16Regs:$g),
3282 (SUST_B_3D_V2B16_TRAP Int64Regs:$s,
3283 Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3284 Int16Regs:$r, Int16Regs:$g)>;
3286 def : Pat<(int_nvvm_sust_b_3d_v2i32_trap
3287 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3288 Int32Regs:$r, Int32Regs:$g),
3289 (SUST_B_3D_V2B32_TRAP Int64Regs:$s,
3290 Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3291 Int32Regs:$r, Int32Regs:$g)>;
3293 def : Pat<(int_nvvm_sust_b_3d_v4i8_trap
3294 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3295 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3296 (SUST_B_3D_V4B8_TRAP Int64Regs:$s,
3297 Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3298 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3300 def : Pat<(int_nvvm_sust_b_3d_v4i16_trap
3301 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3302 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3303 (SUST_B_3D_V4B16_TRAP Int64Regs:$s,
3304 Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3305 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3307 def : Pat<(int_nvvm_sust_b_3d_v4i32_trap
3308 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3309 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3310 (SUST_B_3D_V4B32_TRAP Int64Regs:$s,
3311 Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3312 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3317 def : Pat<(int_nvvm_sust_p_1d_i8_trap
3318 Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
3319 (SUST_P_1D_B8_TRAP Int64Regs:$s, Int32Regs:$x, Int16Regs:$r)>;
3321 def : Pat<(int_nvvm_sust_p_1d_i16_trap
3322 Int64Regs:$s, Int32Regs:$x, Int16Regs:$r),
3323 (SUST_P_1D_B16_TRAP Int64Regs:$s, Int32Regs:$x, Int16Regs:$r)>;
3325 def : Pat<(int_nvvm_sust_p_1d_i32_trap
3326 Int64Regs:$s, Int32Regs:$x, Int32Regs:$r),
3327 (SUST_P_1D_B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$r)>;
3329 def : Pat<(int_nvvm_sust_p_1d_v2i8_trap
3330 Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
3331 (SUST_P_1D_V2B8_TRAP Int64Regs:$s, Int32Regs:$x,
3332 Int16Regs:$r, Int16Regs:$g)>;
3334 def : Pat<(int_nvvm_sust_p_1d_v2i16_trap
3335 Int64Regs:$s, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
3336 (SUST_P_1D_V2B16_TRAP Int64Regs:$s, Int32Regs:$x,
3337 Int16Regs:$r, Int16Regs:$g)>;
3339 def : Pat<(int_nvvm_sust_p_1d_v2i32_trap
3340 Int64Regs:$s, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g),
3341 (SUST_P_1D_V2B32_TRAP Int64Regs:$s, Int32Regs:$x,
3342 Int32Regs:$r, Int32Regs:$g)>;
3344 def : Pat<(int_nvvm_sust_p_1d_v4i8_trap
3345 Int64Regs:$s, Int32Regs:$x,
3346 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3347 (SUST_P_1D_V4B8_TRAP Int64Regs:$s, Int32Regs:$x,
3348 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3350 def : Pat<(int_nvvm_sust_p_1d_v4i16_trap
3351 Int64Regs:$s, Int32Regs:$x,
3352 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3353 (SUST_P_1D_V4B16_TRAP Int64Regs:$s, Int32Regs:$x,
3354 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3356 def : Pat<(int_nvvm_sust_p_1d_v4i32_trap
3357 Int64Regs:$s, Int32Regs:$x,
3358 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3359 (SUST_P_1D_V4B32_TRAP Int64Regs:$s, Int32Regs:$x,
3360 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3364 def : Pat<(int_nvvm_sust_p_1d_array_i8_trap
3365 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r),
3366 (SUST_P_1D_ARRAY_B8_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3369 def : Pat<(int_nvvm_sust_p_1d_array_i16_trap
3370 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r),
3371 (SUST_P_1D_ARRAY_B16_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3374 def : Pat<(int_nvvm_sust_p_1d_array_i32_trap
3375 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$r),
3376 (SUST_P_1D_ARRAY_B32_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3379 def : Pat<(int_nvvm_sust_p_1d_array_v2i8_trap
3380 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
3381 (SUST_P_1D_ARRAY_V2B8_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3382 Int16Regs:$r, Int16Regs:$g)>;
3384 def : Pat<(int_nvvm_sust_p_1d_array_v2i16_trap
3385 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int16Regs:$r, Int16Regs:$g),
3386 (SUST_P_1D_ARRAY_V2B16_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3387 Int16Regs:$r, Int16Regs:$g)>;
3389 def : Pat<(int_nvvm_sust_p_1d_array_v2i32_trap
3390 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$r, Int32Regs:$g),
3391 (SUST_P_1D_ARRAY_V2B32_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3392 Int32Regs:$r, Int32Regs:$g)>;
3394 def : Pat<(int_nvvm_sust_p_1d_array_v4i8_trap
3395 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3396 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3397 (SUST_P_1D_ARRAY_V4B8_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3398 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3400 def : Pat<(int_nvvm_sust_p_1d_array_v4i16_trap
3401 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3402 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3403 (SUST_P_1D_ARRAY_V4B16_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3404 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3406 def : Pat<(int_nvvm_sust_p_1d_array_v4i32_trap
3407 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3408 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3409 (SUST_P_1D_ARRAY_V4B32_TRAP Int64Regs:$s, Int32Regs:$l, Int32Regs:$x,
3410 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3414 def : Pat<(int_nvvm_sust_p_2d_i8_trap
3415 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
3416 (SUST_P_2D_B8_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3419 def : Pat<(int_nvvm_sust_p_2d_i16_trap
3420 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
3421 (SUST_P_2D_B16_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3424 def : Pat<(int_nvvm_sust_p_2d_i32_trap
3425 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r),
3426 (SUST_P_2D_B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3429 def : Pat<(int_nvvm_sust_p_2d_v2i8_trap
3430 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r, Int16Regs:$g),
3431 (SUST_P_2D_V2B8_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3432 Int16Regs:$r, Int16Regs:$g)>;
3434 def : Pat<(int_nvvm_sust_p_2d_v2i16_trap
3435 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r, Int16Regs:$g),
3436 (SUST_P_2D_V2B16_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3437 Int16Regs:$r, Int16Regs:$g)>;
3439 def : Pat<(int_nvvm_sust_p_2d_v2i32_trap
3440 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r, Int32Regs:$g),
3441 (SUST_P_2D_V2B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3442 Int32Regs:$r, Int32Regs:$g)>;
3444 def : Pat<(int_nvvm_sust_p_2d_v4i8_trap
3445 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3446 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3447 (SUST_P_2D_V4B8_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3448 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3450 def : Pat<(int_nvvm_sust_p_2d_v4i16_trap
3451 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3452 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3453 (SUST_P_2D_V4B16_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3454 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3456 def : Pat<(int_nvvm_sust_p_2d_v4i32_trap
3457 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3458 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3459 (SUST_P_2D_V4B32_TRAP Int64Regs:$s, Int32Regs:$x, Int32Regs:$y,
3460 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3464 def : Pat<(int_nvvm_sust_p_2d_array_i8_trap
3465 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
3466 (SUST_P_2D_ARRAY_B8_TRAP Int64Regs:$s,
3467 Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3470 def : Pat<(int_nvvm_sust_p_2d_array_i16_trap
3471 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int16Regs:$r),
3472 (SUST_P_2D_ARRAY_B16_TRAP Int64Regs:$s,
3473 Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3476 def : Pat<(int_nvvm_sust_p_2d_array_i32_trap
3477 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r),
3478 (SUST_P_2D_ARRAY_B32_TRAP Int64Regs:$s,
3479 Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3482 def : Pat<(int_nvvm_sust_p_2d_array_v2i8_trap
3483 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3484 Int16Regs:$r, Int16Regs:$g),
3485 (SUST_P_2D_ARRAY_V2B8_TRAP Int64Regs:$s, Int32Regs:$l,
3486 Int32Regs:$x, Int32Regs:$y,
3487 Int16Regs:$r, Int16Regs:$g)>;
3489 def : Pat<(int_nvvm_sust_p_2d_array_v2i16_trap
3490 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3491 Int16Regs:$r, Int16Regs:$g),
3492 (SUST_P_2D_ARRAY_V2B16_TRAP Int64Regs:$s, Int32Regs:$l,
3493 Int32Regs:$x, Int32Regs:$y,
3494 Int16Regs:$r, Int16Regs:$g)>;
3496 def : Pat<(int_nvvm_sust_p_2d_array_v2i32_trap
3497 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y, Int32Regs:$r,
3499 (SUST_P_2D_ARRAY_V2B32_TRAP Int64Regs:$s, Int32Regs:$l,
3500 Int32Regs:$x, Int32Regs:$y, Int32Regs:$r, Int32Regs:$g)>;
3502 def : Pat<(int_nvvm_sust_p_2d_array_v4i8_trap
3503 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3504 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3505 (SUST_P_2D_ARRAY_V4B8_TRAP Int64Regs:$s,
3506 Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3507 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3509 def : Pat<(int_nvvm_sust_p_2d_array_v4i16_trap
3510 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3511 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3512 (SUST_P_2D_ARRAY_V4B16_TRAP Int64Regs:$s,
3513 Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3514 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3516 def : Pat<(int_nvvm_sust_p_2d_array_v4i32_trap
3517 Int64Regs:$s, Int32Regs:$l, Int32Regs:$x, Int32Regs:$y,
3518 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3519 (SUST_P_2D_ARRAY_V4B32_TRAP Int64Regs:$s, Int32Regs:$l,
3520 Int32Regs:$x, Int32Regs:$y,
3521 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3525 def : Pat<(int_nvvm_sust_p_3d_i8_trap
3526 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3528 (SUST_P_3D_B8_TRAP Int64Regs:$s,
3529 Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3532 def : Pat<(int_nvvm_sust_p_3d_i16_trap
3533 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3535 (SUST_P_3D_B16_TRAP Int64Regs:$s,
3536 Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3539 def : Pat<(int_nvvm_sust_p_3d_i32_trap
3540 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3542 (SUST_P_3D_B32_TRAP Int64Regs:$s,
3543 Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3546 def : Pat<(int_nvvm_sust_p_3d_v2i8_trap
3547 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3548 Int16Regs:$r, Int16Regs:$g),
3549 (SUST_P_3D_V2B8_TRAP Int64Regs:$s,
3550 Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3551 Int16Regs:$r, Int16Regs:$g)>;
3553 def : Pat<(int_nvvm_sust_p_3d_v2i16_trap
3554 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3555 Int16Regs:$r, Int16Regs:$g),
3556 (SUST_P_3D_V2B16_TRAP Int64Regs:$s,
3557 Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3558 Int16Regs:$r, Int16Regs:$g)>;
3560 def : Pat<(int_nvvm_sust_p_3d_v2i32_trap
3561 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3562 Int32Regs:$r, Int32Regs:$g),
3563 (SUST_P_3D_V2B32_TRAP Int64Regs:$s,
3564 Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3565 Int32Regs:$r, Int32Regs:$g)>;
3567 def : Pat<(int_nvvm_sust_p_3d_v4i8_trap
3568 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3569 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3570 (SUST_P_3D_V4B8_TRAP Int64Regs:$s,
3571 Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3572 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3574 def : Pat<(int_nvvm_sust_p_3d_v4i16_trap
3575 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3576 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a),
3577 (SUST_P_3D_V4B16_TRAP Int64Regs:$s,
3578 Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3579 Int16Regs:$r, Int16Regs:$g, Int16Regs:$b, Int16Regs:$a)>;
3581 def : Pat<(int_nvvm_sust_p_3d_v4i32_trap
3582 Int64Regs:$s, Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3583 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a),
3584 (SUST_P_3D_V4B32_TRAP Int64Regs:$s,
3585 Int32Regs:$x, Int32Regs:$y, Int32Regs:$z,
3586 Int32Regs:$r, Int32Regs:$g, Int32Regs:$b, Int32Regs:$a)>;
3590 //===-- Old PTX Back-end Intrinsics ---------------------------------------===//
3592 // These intrinsics are handled to retain compatibility with the old backend.
3594 // PTX Special Purpose Register Accessor Intrinsics
3596 class PTX_READ_SPECIAL_REGISTER_R64<string regname, Intrinsic intop>
3597 : NVPTXInst<(outs Int64Regs:$d), (ins),
3598 !strconcat(!strconcat("mov.u64\t$d, %", regname), ";"),
3599 [(set Int64Regs:$d, (intop))]>;
3601 class PTX_READ_SPECIAL_REGISTER_R32<string regname, Intrinsic intop>
3602 : NVPTXInst<(outs Int32Regs:$d), (ins),
3603 !strconcat(!strconcat("mov.u32\t$d, %", regname), ";"),
3604 [(set Int32Regs:$d, (intop))]>;
3606 // TODO Add read vector-version of special registers
3608 def PTX_READ_TID_X : PTX_READ_SPECIAL_REGISTER_R32<"tid.x",
3609 int_ptx_read_tid_x>;
3610 def PTX_READ_TID_Y : PTX_READ_SPECIAL_REGISTER_R32<"tid.y",
3611 int_ptx_read_tid_y>;
3612 def PTX_READ_TID_Z : PTX_READ_SPECIAL_REGISTER_R32<"tid.z",
3613 int_ptx_read_tid_z>;
3614 def PTX_READ_TID_W : PTX_READ_SPECIAL_REGISTER_R32<"tid.w",
3615 int_ptx_read_tid_w>;
3617 def PTX_READ_NTID_X : PTX_READ_SPECIAL_REGISTER_R32<"ntid.x",
3618 int_ptx_read_ntid_x>;
3619 def PTX_READ_NTID_Y : PTX_READ_SPECIAL_REGISTER_R32<"ntid.y",
3620 int_ptx_read_ntid_y>;
3621 def PTX_READ_NTID_Z : PTX_READ_SPECIAL_REGISTER_R32<"ntid.z",
3622 int_ptx_read_ntid_z>;
3623 def PTX_READ_NTID_W : PTX_READ_SPECIAL_REGISTER_R32<"ntid.w",
3624 int_ptx_read_ntid_w>;
3626 def PTX_READ_LANEID : PTX_READ_SPECIAL_REGISTER_R32<"laneid",
3627 int_ptx_read_laneid>;
3628 def PTX_READ_WARPID : PTX_READ_SPECIAL_REGISTER_R32<"warpid",
3629 int_ptx_read_warpid>;
3630 def PTX_READ_NWARPID : PTX_READ_SPECIAL_REGISTER_R32<"nwarpid",
3631 int_ptx_read_nwarpid>;
3633 def PTX_READ_CTAID_X : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.x",
3634 int_ptx_read_ctaid_x>;
3635 def PTX_READ_CTAID_Y : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.y",
3636 int_ptx_read_ctaid_y>;
3637 def PTX_READ_CTAID_Z : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.z",
3638 int_ptx_read_ctaid_z>;
3639 def PTX_READ_CTAID_W : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.w",
3640 int_ptx_read_ctaid_w>;
3642 def PTX_READ_NCTAID_X : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.x",
3643 int_ptx_read_nctaid_x>;
3644 def PTX_READ_NCTAID_Y : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.y",
3645 int_ptx_read_nctaid_y>;
3646 def PTX_READ_NCTAID_Z : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.z",
3647 int_ptx_read_nctaid_z>;
3648 def PTX_READ_NCTAID_W : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.w",
3649 int_ptx_read_nctaid_w>;
3651 def PTX_READ_SMID : PTX_READ_SPECIAL_REGISTER_R32<"smid",
3653 def PTX_READ_NSMID : PTX_READ_SPECIAL_REGISTER_R32<"nsmid",
3654 int_ptx_read_nsmid>;
3655 def PTX_READ_GRIDID : PTX_READ_SPECIAL_REGISTER_R32<"gridid",
3656 int_ptx_read_gridid>;
3658 def PTX_READ_LANEMASK_EQ
3659 : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_eq", int_ptx_read_lanemask_eq>;
3660 def PTX_READ_LANEMASK_LE
3661 : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_le", int_ptx_read_lanemask_le>;
3662 def PTX_READ_LANEMASK_LT
3663 : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_lt", int_ptx_read_lanemask_lt>;
3664 def PTX_READ_LANEMASK_GE
3665 : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_ge", int_ptx_read_lanemask_ge>;
3666 def PTX_READ_LANEMASK_GT
3667 : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_gt", int_ptx_read_lanemask_gt>;
3670 : PTX_READ_SPECIAL_REGISTER_R32<"clock", int_ptx_read_clock>;
3671 def PTX_READ_CLOCK64
3672 : PTX_READ_SPECIAL_REGISTER_R64<"clock64", int_ptx_read_clock64>;
3674 def PTX_READ_PM0 : PTX_READ_SPECIAL_REGISTER_R32<"pm0", int_ptx_read_pm0>;
3675 def PTX_READ_PM1 : PTX_READ_SPECIAL_REGISTER_R32<"pm1", int_ptx_read_pm1>;
3676 def PTX_READ_PM2 : PTX_READ_SPECIAL_REGISTER_R32<"pm2", int_ptx_read_pm2>;
3677 def PTX_READ_PM3 : PTX_READ_SPECIAL_REGISTER_R32<"pm3", int_ptx_read_pm3>;
3679 // PTX Parallel Synchronization and Communication Intrinsics
3681 def PTX_BAR_SYNC : NVPTXInst<(outs), (ins i32imm:$i), "bar.sync\t$i;",
3682 [(int_ptx_bar_sync imm:$i)]>;