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 multiclass SAT<NVPTXRegClass regclass, Operand fimm, Intrinsic IntMinOp,
86 Intrinsic IntMaxOp, PatLeaf f0, PatLeaf f1, string OpStr> {
88 // fmin(1.0, fmax(0.0, x)) => sat(x)
89 def SAT11 : NVPTXInst<(outs regclass:$dst),
90 (ins fimm:$srcf0, fimm:$srcf1, regclass:$src),
92 [(set regclass:$dst, (IntMinOp f1:$srcf0 ,
93 (IntMaxOp f0:$srcf1, regclass:$src)))]>;
95 // fmin(1.0, fmax(x, 0.0)) => sat(x)
96 def SAT12 : NVPTXInst<(outs regclass:$dst),
97 (ins fimm:$srcf0, fimm:$srcf1, regclass:$src),
99 [(set regclass:$dst, (IntMinOp f1:$srcf0 ,
100 (IntMaxOp regclass:$src, f0:$srcf1)))]>;
102 // fmin(fmax(0.0, x), 1.0) => sat(x)
103 def SAT13 : NVPTXInst<(outs regclass:$dst),
104 (ins fimm:$srcf0, fimm:$srcf1, regclass:$src),
106 [(set regclass:$dst, (IntMinOp
107 (IntMaxOp f0:$srcf0, regclass:$src), f1:$srcf1))]>;
109 // fmin(fmax(x, 0.0), 1.0) => sat(x)
110 def SAT14 : NVPTXInst<(outs regclass:$dst),
111 (ins fimm:$srcf0, fimm:$srcf1, regclass:$src),
113 [(set regclass:$dst, (IntMinOp
114 (IntMaxOp regclass:$src, f0:$srcf0), f1:$srcf1))]>;
117 // Note that max(0.0, min(x, 1.0)) cannot be mapped to sat(x) because when x
119 // max(0.0, min(x, 1.0)) is 1.0 while sat(x) is 0.
120 // Same story for fmax, fmin.
122 defm SAT_fmin_fmax_f : SAT<Float32Regs, f32imm, int_nvvm_fmin_f,
123 int_nvvm_fmax_f, immFloat0, immFloat1,
124 "cvt.sat.f32.f32 \t$dst, $src; \n">;
125 defm SAT_fmin_fmax_d : SAT<Float64Regs, f64imm, int_nvvm_fmin_d,
126 int_nvvm_fmax_d, immDouble0, immDouble1,
127 "cvt.sat.f64.f64 \t$dst, $src; \n">;
130 // We need a full string for OpcStr here because we need to deal with case like
132 class F_MATH_1<string OpcStr, NVPTXRegClass target_regclass,
133 NVPTXRegClass src_regclass, Intrinsic IntOP>
134 : NVPTXInst<(outs target_regclass:$dst), (ins src_regclass:$src0),
136 [(set target_regclass:$dst, (IntOP src_regclass:$src0))]>;
138 // We need a full string for OpcStr here because we need to deal with the case
139 // like INT_PTX_NATIVE_POWR_F.
140 class F_MATH_2<string OpcStr, NVPTXRegClass t_regclass,
141 NVPTXRegClass s0_regclass, NVPTXRegClass s1_regclass, Intrinsic IntOP>
142 : NVPTXInst<(outs t_regclass:$dst),
143 (ins s0_regclass:$src0, s1_regclass:$src1),
145 [(set t_regclass:$dst, (IntOP s0_regclass:$src0, s1_regclass:$src1))]>;
147 class F_MATH_3<string OpcStr, NVPTXRegClass t_regclass,
148 NVPTXRegClass s0_regclass, NVPTXRegClass s1_regclass,
149 NVPTXRegClass s2_regclass, Intrinsic IntOP>
150 : NVPTXInst<(outs t_regclass:$dst),
151 (ins s0_regclass:$src0, s1_regclass:$src1, s2_regclass:$src2),
153 [(set t_regclass:$dst,
154 (IntOP s0_regclass:$src0, s1_regclass:$src1, s2_regclass:$src2))]>;
160 def INT_NVVM_CLZ_I : F_MATH_1<"clz.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
162 def INT_NVVM_CLZ_LL : F_MATH_1<"clz.b64 \t$dst, $src0;", Int32Regs, Int64Regs,
165 def INT_NVVM_POPC_I : F_MATH_1<"popc.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
167 def INT_NVVM_POPC_LL : F_MATH_1<"popc.b64 \t$dst, $src0;", Int32Regs, Int64Regs,
170 def INT_NVVM_PRMT : F_MATH_3<"prmt.b32 \t$dst, $src0, $src1, $src2;", Int32Regs,
171 Int32Regs, Int32Regs, Int32Regs, int_nvvm_prmt>;
177 def INT_NVVM_MIN_I : F_MATH_2<"min.s32 \t$dst, $src0, $src1;", Int32Regs,
178 Int32Regs, Int32Regs, int_nvvm_min_i>;
179 def INT_NVVM_MIN_UI : F_MATH_2<"min.u32 \t$dst, $src0, $src1;", Int32Regs,
180 Int32Regs, Int32Regs, int_nvvm_min_ui>;
182 def INT_NVVM_MIN_LL : F_MATH_2<"min.s64 \t$dst, $src0, $src1;", Int64Regs,
183 Int64Regs, Int64Regs, int_nvvm_min_ll>;
184 def INT_NVVM_MIN_ULL : F_MATH_2<"min.u64 \t$dst, $src0, $src1;", Int64Regs,
185 Int64Regs, Int64Regs, int_nvvm_min_ull>;
187 def INT_NVVM_MAX_I : F_MATH_2<"max.s32 \t$dst, $src0, $src1;", Int32Regs,
188 Int32Regs, Int32Regs, int_nvvm_max_i>;
189 def INT_NVVM_MAX_UI : F_MATH_2<"max.u32 \t$dst, $src0, $src1;", Int32Regs,
190 Int32Regs, Int32Regs, int_nvvm_max_ui>;
192 def INT_NVVM_MAX_LL : F_MATH_2<"max.s64 \t$dst, $src0, $src1;", Int64Regs,
193 Int64Regs, Int64Regs, int_nvvm_max_ll>;
194 def INT_NVVM_MAX_ULL : F_MATH_2<"max.u64 \t$dst, $src0, $src1;", Int64Regs,
195 Int64Regs, Int64Regs, int_nvvm_max_ull>;
197 def INT_NVVM_FMIN_F : F_MATH_2<"min.f32 \t$dst, $src0, $src1;", Float32Regs,
198 Float32Regs, Float32Regs, int_nvvm_fmin_f>;
199 def INT_NVVM_FMIN_FTZ_F : F_MATH_2<"min.ftz.f32 \t$dst, $src0, $src1;",
200 Float32Regs, Float32Regs, Float32Regs, int_nvvm_fmin_ftz_f>;
202 def INT_NVVM_FMAX_F : F_MATH_2<"max.f32 \t$dst, $src0, $src1;", Float32Regs,
203 Float32Regs, Float32Regs, int_nvvm_fmax_f>;
204 def INT_NVVM_FMAX_FTZ_F : F_MATH_2<"max.ftz.f32 \t$dst, $src0, $src1;",
205 Float32Regs, Float32Regs, Float32Regs, int_nvvm_fmax_ftz_f>;
207 def INT_NVVM_FMIN_D : F_MATH_2<"min.f64 \t$dst, $src0, $src1;", Float64Regs,
208 Float64Regs, Float64Regs, int_nvvm_fmin_d>;
209 def INT_NVVM_FMAX_D : F_MATH_2<"max.f64 \t$dst, $src0, $src1;", Float64Regs,
210 Float64Regs, Float64Regs, int_nvvm_fmax_d>;
216 def INT_NVVM_MULHI_I : F_MATH_2<"mul.hi.s32 \t$dst, $src0, $src1;", Int32Regs,
217 Int32Regs, Int32Regs, int_nvvm_mulhi_i>;
218 def INT_NVVM_MULHI_UI : F_MATH_2<"mul.hi.u32 \t$dst, $src0, $src1;", Int32Regs,
219 Int32Regs, Int32Regs, int_nvvm_mulhi_ui>;
221 def INT_NVVM_MULHI_LL : F_MATH_2<"mul.hi.s64 \t$dst, $src0, $src1;", Int64Regs,
222 Int64Regs, Int64Regs, int_nvvm_mulhi_ll>;
223 def INT_NVVM_MULHI_ULL : F_MATH_2<"mul.hi.u64 \t$dst, $src0, $src1;", Int64Regs,
224 Int64Regs, Int64Regs, int_nvvm_mulhi_ull>;
226 def INT_NVVM_MUL_RN_FTZ_F : F_MATH_2<"mul.rn.ftz.f32 \t$dst, $src0, $src1;",
227 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rn_ftz_f>;
228 def INT_NVVM_MUL_RN_F : F_MATH_2<"mul.rn.f32 \t$dst, $src0, $src1;",
229 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rn_f>;
230 def INT_NVVM_MUL_RZ_FTZ_F : F_MATH_2<"mul.rz.ftz.f32 \t$dst, $src0, $src1;",
231 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rz_ftz_f>;
232 def INT_NVVM_MUL_RZ_F : F_MATH_2<"mul.rz.f32 \t$dst, $src0, $src1;",
233 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rz_f>;
234 def INT_NVVM_MUL_RM_FTZ_F : F_MATH_2<"mul.rm.ftz.f32 \t$dst, $src0, $src1;",
235 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rm_ftz_f>;
236 def INT_NVVM_MUL_RM_F : F_MATH_2<"mul.rm.f32 \t$dst, $src0, $src1;",
237 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rm_f>;
238 def INT_NVVM_MUL_RP_FTZ_F : F_MATH_2<"mul.rp.ftz.f32 \t$dst, $src0, $src1;",
239 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rp_ftz_f>;
240 def INT_NVVM_MUL_RP_F : F_MATH_2<"mul.rp.f32 \t$dst, $src0, $src1;",
241 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rp_f>;
243 def INT_NVVM_MUL_RN_D : F_MATH_2<"mul.rn.f64 \t$dst, $src0, $src1;",
244 Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rn_d>;
245 def INT_NVVM_MUL_RZ_D : F_MATH_2<"mul.rz.f64 \t$dst, $src0, $src1;",
246 Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rz_d>;
247 def INT_NVVM_MUL_RM_D : F_MATH_2<"mul.rm.f64 \t$dst, $src0, $src1;",
248 Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rm_d>;
249 def INT_NVVM_MUL_RP_D : F_MATH_2<"mul.rp.f64 \t$dst, $src0, $src1;",
250 Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rp_d>;
252 def INT_NVVM_MUL24_I : F_MATH_2<"mul24.lo.s32 \t$dst, $src0, $src1;",
253 Int32Regs, Int32Regs, Int32Regs, int_nvvm_mul24_i>;
254 def INT_NVVM_MUL24_UI : F_MATH_2<"mul24.lo.u32 \t$dst, $src0, $src1;",
255 Int32Regs, Int32Regs, Int32Regs, int_nvvm_mul24_ui>;
261 def INT_NVVM_DIV_APPROX_FTZ_F
262 : F_MATH_2<"div.approx.ftz.f32 \t$dst, $src0, $src1;", Float32Regs,
263 Float32Regs, Float32Regs, int_nvvm_div_approx_ftz_f>;
264 def INT_NVVM_DIV_APPROX_F : F_MATH_2<"div.approx.f32 \t$dst, $src0, $src1;",
265 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_approx_f>;
267 def INT_NVVM_DIV_RN_FTZ_F : F_MATH_2<"div.rn.ftz.f32 \t$dst, $src0, $src1;",
268 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rn_ftz_f>;
269 def INT_NVVM_DIV_RN_F : F_MATH_2<"div.rn.f32 \t$dst, $src0, $src1;",
270 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rn_f>;
271 def INT_NVVM_DIV_RZ_FTZ_F : F_MATH_2<"div.rz.ftz.f32 \t$dst, $src0, $src1;",
272 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rz_ftz_f>;
273 def INT_NVVM_DIV_RZ_F : F_MATH_2<"div.rz.f32 \t$dst, $src0, $src1;",
274 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rz_f>;
275 def INT_NVVM_DIV_RM_FTZ_F : F_MATH_2<"div.rm.ftz.f32 \t$dst, $src0, $src1;",
276 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rm_ftz_f>;
277 def INT_NVVM_DIV_RM_F : F_MATH_2<"div.rm.f32 \t$dst, $src0, $src1;",
278 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rm_f>;
279 def INT_NVVM_DIV_RP_FTZ_F : F_MATH_2<"div.rp.ftz.f32 \t$dst, $src0, $src1;",
280 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rp_ftz_f>;
281 def INT_NVVM_DIV_RP_F : F_MATH_2<"div.rp.f32 \t$dst, $src0, $src1;",
282 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rp_f>;
284 def INT_NVVM_DIV_RN_D : F_MATH_2<"div.rn.f64 \t$dst, $src0, $src1;",
285 Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rn_d>;
286 def INT_NVVM_DIV_RZ_D : F_MATH_2<"div.rz.f64 \t$dst, $src0, $src1;",
287 Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rz_d>;
288 def INT_NVVM_DIV_RM_D : F_MATH_2<"div.rm.f64 \t$dst, $src0, $src1;",
289 Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rm_d>;
290 def INT_NVVM_DIV_RP_D : F_MATH_2<"div.rp.f64 \t$dst, $src0, $src1;",
291 Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rp_d>;
297 def INT_NVVM_BREV32 : F_MATH_1<"brev.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
299 def INT_NVVM_BREV64 : F_MATH_1<"brev.b64 \t$dst, $src0;", Int64Regs, Int64Regs,
306 def INT_NVVM_SAD_I : F_MATH_3<"sad.s32 \t$dst, $src0, $src1, $src2;",
307 Int32Regs, Int32Regs, Int32Regs, Int32Regs, int_nvvm_sad_i>;
308 def INT_NVVM_SAD_UI : F_MATH_3<"sad.u32 \t$dst, $src0, $src1, $src2;",
309 Int32Regs, Int32Regs, Int32Regs, Int32Regs, int_nvvm_sad_ui>;
315 def INT_NVVM_FLOOR_FTZ_F : F_MATH_1<"cvt.rmi.ftz.f32.f32 \t$dst, $src0;",
316 Float32Regs, Float32Regs, int_nvvm_floor_ftz_f>;
317 def INT_NVVM_FLOOR_F : F_MATH_1<"cvt.rmi.f32.f32 \t$dst, $src0;",
318 Float32Regs, Float32Regs, int_nvvm_floor_f>;
319 def INT_NVVM_FLOOR_D : F_MATH_1<"cvt.rmi.f64.f64 \t$dst, $src0;",
320 Float64Regs, Float64Regs, int_nvvm_floor_d>;
322 def INT_NVVM_CEIL_FTZ_F : F_MATH_1<"cvt.rpi.ftz.f32.f32 \t$dst, $src0;",
323 Float32Regs, Float32Regs, int_nvvm_ceil_ftz_f>;
324 def INT_NVVM_CEIL_F : F_MATH_1<"cvt.rpi.f32.f32 \t$dst, $src0;",
325 Float32Regs, Float32Regs, int_nvvm_ceil_f>;
326 def INT_NVVM_CEIL_D : F_MATH_1<"cvt.rpi.f64.f64 \t$dst, $src0;",
327 Float64Regs, Float64Regs, int_nvvm_ceil_d>;
333 def INT_NVVM_ABS_I : F_MATH_1<"abs.s32 \t$dst, $src0;", Int32Regs, Int32Regs,
335 def INT_NVVM_ABS_LL : F_MATH_1<"abs.s64 \t$dst, $src0;", Int64Regs, Int64Regs,
338 def INT_NVVM_FABS_FTZ_F : F_MATH_1<"abs.ftz.f32 \t$dst, $src0;", Float32Regs,
339 Float32Regs, int_nvvm_fabs_ftz_f>;
340 def INT_NVVM_FABS_F : F_MATH_1<"abs.f32 \t$dst, $src0;", Float32Regs,
341 Float32Regs, int_nvvm_fabs_f>;
343 def INT_NVVM_FABS_D : F_MATH_1<"abs.f64 \t$dst, $src0;", Float64Regs,
344 Float64Regs, int_nvvm_fabs_d>;
350 def INT_NVVM_ROUND_FTZ_F : F_MATH_1<"cvt.rni.ftz.f32.f32 \t$dst, $src0;",
351 Float32Regs, Float32Regs, int_nvvm_round_ftz_f>;
352 def INT_NVVM_ROUND_F : F_MATH_1<"cvt.rni.f32.f32 \t$dst, $src0;", Float32Regs,
353 Float32Regs, int_nvvm_round_f>;
355 def INT_NVVM_ROUND_D : F_MATH_1<"cvt.rni.f64.f64 \t$dst, $src0;", Float64Regs,
356 Float64Regs, int_nvvm_round_d>;
362 def INT_NVVM_TRUNC_FTZ_F : F_MATH_1<"cvt.rzi.ftz.f32.f32 \t$dst, $src0;",
363 Float32Regs, Float32Regs, int_nvvm_trunc_ftz_f>;
364 def INT_NVVM_TRUNC_F : F_MATH_1<"cvt.rzi.f32.f32 \t$dst, $src0;", Float32Regs,
365 Float32Regs, int_nvvm_trunc_f>;
367 def INT_NVVM_TRUNC_D : F_MATH_1<"cvt.rzi.f64.f64 \t$dst, $src0;", Float64Regs,
368 Float64Regs, int_nvvm_trunc_d>;
374 def INT_NVVM_SATURATE_FTZ_F : F_MATH_1<"cvt.sat.ftz.f32.f32 \t$dst, $src0;",
375 Float32Regs, Float32Regs, int_nvvm_saturate_ftz_f>;
376 def INT_NVVM_SATURATE_F : F_MATH_1<"cvt.sat.f32.f32 \t$dst, $src0;",
377 Float32Regs, Float32Regs, int_nvvm_saturate_f>;
379 def INT_NVVM_SATURATE_D : F_MATH_1<"cvt.sat.f64.f64 \t$dst, $src0;",
380 Float64Regs, Float64Regs, int_nvvm_saturate_d>;
386 def INT_NVVM_EX2_APPROX_FTZ_F : F_MATH_1<"ex2.approx.ftz.f32 \t$dst, $src0;",
387 Float32Regs, Float32Regs, int_nvvm_ex2_approx_ftz_f>;
388 def INT_NVVM_EX2_APPROX_F : F_MATH_1<"ex2.approx.f32 \t$dst, $src0;",
389 Float32Regs, Float32Regs, int_nvvm_ex2_approx_f>;
390 def INT_NVVM_EX2_APPROX_D : F_MATH_1<"ex2.approx.f64 \t$dst, $src0;",
391 Float64Regs, Float64Regs, int_nvvm_ex2_approx_d>;
393 def INT_NVVM_LG2_APPROX_FTZ_F : F_MATH_1<"lg2.approx.ftz.f32 \t$dst, $src0;",
394 Float32Regs, Float32Regs, int_nvvm_lg2_approx_ftz_f>;
395 def INT_NVVM_LG2_APPROX_F : F_MATH_1<"lg2.approx.f32 \t$dst, $src0;",
396 Float32Regs, Float32Regs, int_nvvm_lg2_approx_f>;
397 def INT_NVVM_LG2_APPROX_D : F_MATH_1<"lg2.approx.f64 \t$dst, $src0;",
398 Float64Regs, Float64Regs, int_nvvm_lg2_approx_d>;
404 def INT_NVVM_SIN_APPROX_FTZ_F : F_MATH_1<"sin.approx.ftz.f32 \t$dst, $src0;",
405 Float32Regs, Float32Regs, int_nvvm_sin_approx_ftz_f>;
406 def INT_NVVM_SIN_APPROX_F : F_MATH_1<"sin.approx.f32 \t$dst, $src0;",
407 Float32Regs, Float32Regs, int_nvvm_sin_approx_f>;
409 def INT_NVVM_COS_APPROX_FTZ_F : F_MATH_1<"cos.approx.ftz.f32 \t$dst, $src0;",
410 Float32Regs, Float32Regs, int_nvvm_cos_approx_ftz_f>;
411 def INT_NVVM_COS_APPROX_F : F_MATH_1<"cos.approx.f32 \t$dst, $src0;",
412 Float32Regs, Float32Regs, int_nvvm_cos_approx_f>;
418 def INT_NVVM_FMA_RN_FTZ_F
419 : F_MATH_3<"fma.rn.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
420 Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rn_ftz_f>;
421 def INT_NVVM_FMA_RN_F : F_MATH_3<"fma.rn.f32 \t$dst, $src0, $src1, $src2;",
422 Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rn_f>;
423 def INT_NVVM_FMA_RZ_FTZ_F
424 : F_MATH_3<"fma.rz.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
425 Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rz_ftz_f>;
426 def INT_NVVM_FMA_RZ_F : F_MATH_3<"fma.rz.f32 \t$dst, $src0, $src1, $src2;",
427 Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rz_f>;
428 def INT_NVVM_FMA_RM_FTZ_F
429 : F_MATH_3<"fma.rm.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
430 Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rm_ftz_f>;
431 def INT_NVVM_FMA_RM_F : F_MATH_3<"fma.rm.f32 \t$dst, $src0, $src1, $src2;",
432 Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rm_f>;
433 def INT_NVVM_FMA_RP_FTZ_F
434 : F_MATH_3<"fma.rp.ftz.f32 \t$dst, $src0, $src1, $src2;", Float32Regs,
435 Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rp_ftz_f>;
436 def INT_NVVM_FMA_RP_F : F_MATH_3<"fma.rp.f32 \t$dst, $src0, $src1, $src2;",
437 Float32Regs, Float32Regs, Float32Regs, Float32Regs, int_nvvm_fma_rp_f>;
439 def INT_NVVM_FMA_RN_D : F_MATH_3<"fma.rn.f64 \t$dst, $src0, $src1, $src2;",
440 Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rn_d>;
441 def INT_NVVM_FMA_RZ_D : F_MATH_3<"fma.rz.f64 \t$dst, $src0, $src1, $src2;",
442 Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rz_d>;
443 def INT_NVVM_FMA_RM_D : F_MATH_3<"fma.rm.f64 \t$dst, $src0, $src1, $src2;",
444 Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rm_d>;
445 def INT_NVVM_FMA_RP_D : F_MATH_3<"fma.rp.f64 \t$dst, $src0, $src1, $src2;",
446 Float64Regs, Float64Regs, Float64Regs, Float64Regs, int_nvvm_fma_rp_d>;
452 def INT_NVVM_RCP_RN_FTZ_F : F_MATH_1<"rcp.rn.ftz.f32 \t$dst, $src0;",
453 Float32Regs, Float32Regs, int_nvvm_rcp_rn_ftz_f>;
454 def INT_NVVM_RCP_RN_F : F_MATH_1<"rcp.rn.f32 \t$dst, $src0;",
455 Float32Regs, Float32Regs, int_nvvm_rcp_rn_f>;
456 def INT_NVVM_RCP_RZ_FTZ_F : F_MATH_1<"rcp.rz.ftz.f32 \t$dst, $src0;",
457 Float32Regs, Float32Regs, int_nvvm_rcp_rz_ftz_f>;
458 def INT_NVVM_RCP_RZ_F : F_MATH_1<"rcp.rz.f32 \t$dst, $src0;",
459 Float32Regs, Float32Regs, int_nvvm_rcp_rz_f>;
460 def INT_NVVM_RCP_RM_FTZ_F : F_MATH_1<"rcp.rm.ftz.f32 \t$dst, $src0;",
461 Float32Regs, Float32Regs, int_nvvm_rcp_rm_ftz_f>;
462 def INT_NVVM_RCP_RM_F : F_MATH_1<"rcp.rm.f32 \t$dst, $src0;",
463 Float32Regs, Float32Regs, int_nvvm_rcp_rm_f>;
464 def INT_NVVM_RCP_RP_FTZ_F : F_MATH_1<"rcp.rp.ftz.f32 \t$dst, $src0;",
465 Float32Regs, Float32Regs, int_nvvm_rcp_rp_ftz_f>;
466 def INT_NVVM_RCP_RP_F : F_MATH_1<"rcp.rp.f32 \t$dst, $src0;",
467 Float32Regs, Float32Regs, int_nvvm_rcp_rp_f>;
469 def INT_NVVM_RCP_RN_D : F_MATH_1<"rcp.rn.f64 \t$dst, $src0;", Float64Regs,
470 Float64Regs, int_nvvm_rcp_rn_d>;
471 def INT_NVVM_RCP_RZ_D : F_MATH_1<"rcp.rz.f64 \t$dst, $src0;", Float64Regs,
472 Float64Regs, int_nvvm_rcp_rz_d>;
473 def INT_NVVM_RCP_RM_D : F_MATH_1<"rcp.rm.f64 \t$dst, $src0;", Float64Regs,
474 Float64Regs, int_nvvm_rcp_rm_d>;
475 def INT_NVVM_RCP_RP_D : F_MATH_1<"rcp.rp.f64 \t$dst, $src0;", Float64Regs,
476 Float64Regs, int_nvvm_rcp_rp_d>;
478 def INT_NVVM_RCP_APPROX_FTZ_D : F_MATH_1<"rcp.approx.ftz.f64 \t$dst, $src0;",
479 Float64Regs, Float64Regs, int_nvvm_rcp_approx_ftz_d>;
485 def INT_NVVM_SQRT_RN_FTZ_F : F_MATH_1<"sqrt.rn.ftz.f32 \t$dst, $src0;",
486 Float32Regs, Float32Regs, int_nvvm_sqrt_rn_ftz_f>;
487 def INT_NVVM_SQRT_RN_F : F_MATH_1<"sqrt.rn.f32 \t$dst, $src0;", Float32Regs,
488 Float32Regs, int_nvvm_sqrt_rn_f>;
489 def INT_NVVM_SQRT_RZ_FTZ_F : F_MATH_1<"sqrt.rz.ftz.f32 \t$dst, $src0;",
490 Float32Regs, Float32Regs, int_nvvm_sqrt_rz_ftz_f>;
491 def INT_NVVM_SQRT_RZ_F : F_MATH_1<"sqrt.rz.f32 \t$dst, $src0;", Float32Regs,
492 Float32Regs, int_nvvm_sqrt_rz_f>;
493 def INT_NVVM_SQRT_RM_FTZ_F : F_MATH_1<"sqrt.rm.ftz.f32 \t$dst, $src0;",
494 Float32Regs, Float32Regs, int_nvvm_sqrt_rm_ftz_f>;
495 def INT_NVVM_SQRT_RM_F : F_MATH_1<"sqrt.rm.f32 \t$dst, $src0;", Float32Regs,
496 Float32Regs, int_nvvm_sqrt_rm_f>;
497 def INT_NVVM_SQRT_RP_FTZ_F : F_MATH_1<"sqrt.rp.ftz.f32 \t$dst, $src0;",
498 Float32Regs, Float32Regs, int_nvvm_sqrt_rp_ftz_f>;
499 def INT_NVVM_SQRT_RP_F : F_MATH_1<"sqrt.rp.f32 \t$dst, $src0;", Float32Regs,
500 Float32Regs, int_nvvm_sqrt_rp_f>;
501 def INT_NVVM_SQRT_APPROX_FTZ_F : F_MATH_1<"sqrt.approx.ftz.f32 \t$dst, $src0;",
502 Float32Regs, Float32Regs, int_nvvm_sqrt_approx_ftz_f>;
503 def INT_NVVM_SQRT_APPROX_F : F_MATH_1<"sqrt.approx.f32 \t$dst, $src0;",
504 Float32Regs, Float32Regs, int_nvvm_sqrt_approx_f>;
506 def INT_NVVM_SQRT_RN_D : F_MATH_1<"sqrt.rn.f64 \t$dst, $src0;", Float64Regs,
507 Float64Regs, int_nvvm_sqrt_rn_d>;
508 def INT_NVVM_SQRT_RZ_D : F_MATH_1<"sqrt.rz.f64 \t$dst, $src0;", Float64Regs,
509 Float64Regs, int_nvvm_sqrt_rz_d>;
510 def INT_NVVM_SQRT_RM_D : F_MATH_1<"sqrt.rm.f64 \t$dst, $src0;", Float64Regs,
511 Float64Regs, int_nvvm_sqrt_rm_d>;
512 def INT_NVVM_SQRT_RP_D : F_MATH_1<"sqrt.rp.f64 \t$dst, $src0;", Float64Regs,
513 Float64Regs, int_nvvm_sqrt_rp_d>;
515 // nvvm_sqrt intrinsic
516 def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
517 (INT_NVVM_SQRT_RN_FTZ_F Float32Regs:$a)>, Requires<[doF32FTZ, do_SQRTF32_RN]>;
518 def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
519 (INT_NVVM_SQRT_RN_F Float32Regs:$a)>, Requires<[do_SQRTF32_RN]>;
520 def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
521 (INT_NVVM_SQRT_APPROX_FTZ_F Float32Regs:$a)>, Requires<[doF32FTZ]>;
522 def : Pat<(int_nvvm_sqrt_f Float32Regs:$a),
523 (INT_NVVM_SQRT_APPROX_F Float32Regs:$a)>;
529 def INT_NVVM_RSQRT_APPROX_FTZ_F
530 : F_MATH_1<"rsqrt.approx.ftz.f32 \t$dst, $src0;", Float32Regs, Float32Regs,
531 int_nvvm_rsqrt_approx_ftz_f>;
532 def INT_NVVM_RSQRT_APPROX_F : F_MATH_1<"rsqrt.approx.f32 \t$dst, $src0;",
533 Float32Regs, Float32Regs, int_nvvm_rsqrt_approx_f>;
534 def INT_NVVM_RSQRT_APPROX_D : F_MATH_1<"rsqrt.approx.f64 \t$dst, $src0;",
535 Float64Regs, Float64Regs, int_nvvm_rsqrt_approx_d>;
541 def INT_NVVM_ADD_RN_FTZ_F : F_MATH_2<"add.rn.ftz.f32 \t$dst, $src0, $src1;",
542 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rn_ftz_f>;
543 def INT_NVVM_ADD_RN_F : F_MATH_2<"add.rn.f32 \t$dst, $src0, $src1;",
544 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rn_f>;
545 def INT_NVVM_ADD_RZ_FTZ_F : F_MATH_2<"add.rz.ftz.f32 \t$dst, $src0, $src1;",
546 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rz_ftz_f>;
547 def INT_NVVM_ADD_RZ_F : F_MATH_2<"add.rz.f32 \t$dst, $src0, $src1;",
548 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rz_f>;
549 def INT_NVVM_ADD_RM_FTZ_F : F_MATH_2<"add.rm.ftz.f32 \t$dst, $src0, $src1;",
550 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rm_ftz_f>;
551 def INT_NVVM_ADD_RM_F : F_MATH_2<"add.rm.f32 \t$dst, $src0, $src1;",
552 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rm_f>;
553 def INT_NVVM_ADD_RP_FTZ_F : F_MATH_2<"add.rp.ftz.f32 \t$dst, $src0, $src1;",
554 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rp_ftz_f>;
555 def INT_NVVM_ADD_RP_F : F_MATH_2<"add.rp.f32 \t$dst, $src0, $src1;",
556 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rp_f>;
558 def INT_NVVM_ADD_RN_D : F_MATH_2<"add.rn.f64 \t$dst, $src0, $src1;",
559 Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rn_d>;
560 def INT_NVVM_ADD_RZ_D : F_MATH_2<"add.rz.f64 \t$dst, $src0, $src1;",
561 Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rz_d>;
562 def INT_NVVM_ADD_RM_D : F_MATH_2<"add.rm.f64 \t$dst, $src0, $src1;",
563 Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rm_d>;
564 def INT_NVVM_ADD_RP_D : F_MATH_2<"add.rp.f64 \t$dst, $src0, $src1;",
565 Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rp_d>;
571 def INT_NVVM_D2F_RN_FTZ : F_MATH_1<"cvt.rn.ftz.f32.f64 \t$dst, $src0;",
572 Float32Regs, Float64Regs, int_nvvm_d2f_rn_ftz>;
573 def INT_NVVM_D2F_RN : F_MATH_1<"cvt.rn.f32.f64 \t$dst, $src0;",
574 Float32Regs, Float64Regs, int_nvvm_d2f_rn>;
575 def INT_NVVM_D2F_RZ_FTZ : F_MATH_1<"cvt.rz.ftz.f32.f64 \t$dst, $src0;",
576 Float32Regs, Float64Regs, int_nvvm_d2f_rz_ftz>;
577 def INT_NVVM_D2F_RZ : F_MATH_1<"cvt.rz.f32.f64 \t$dst, $src0;",
578 Float32Regs, Float64Regs, int_nvvm_d2f_rz>;
579 def INT_NVVM_D2F_RM_FTZ : F_MATH_1<"cvt.rm.ftz.f32.f64 \t$dst, $src0;",
580 Float32Regs, Float64Regs, int_nvvm_d2f_rm_ftz>;
581 def INT_NVVM_D2F_RM : F_MATH_1<"cvt.rm.f32.f64 \t$dst, $src0;",
582 Float32Regs, Float64Regs, int_nvvm_d2f_rm>;
583 def INT_NVVM_D2F_RP_FTZ : F_MATH_1<"cvt.rp.ftz.f32.f64 \t$dst, $src0;",
584 Float32Regs, Float64Regs, int_nvvm_d2f_rp_ftz>;
585 def INT_NVVM_D2F_RP : F_MATH_1<"cvt.rp.f32.f64 \t$dst, $src0;",
586 Float32Regs, Float64Regs, int_nvvm_d2f_rp>;
588 def INT_NVVM_D2I_RN : F_MATH_1<"cvt.rni.s32.f64 \t$dst, $src0;",
589 Int32Regs, Float64Regs, int_nvvm_d2i_rn>;
590 def INT_NVVM_D2I_RZ : F_MATH_1<"cvt.rzi.s32.f64 \t$dst, $src0;",
591 Int32Regs, Float64Regs, int_nvvm_d2i_rz>;
592 def INT_NVVM_D2I_RM : F_MATH_1<"cvt.rmi.s32.f64 \t$dst, $src0;",
593 Int32Regs, Float64Regs, int_nvvm_d2i_rm>;
594 def INT_NVVM_D2I_RP : F_MATH_1<"cvt.rpi.s32.f64 \t$dst, $src0;",
595 Int32Regs, Float64Regs, int_nvvm_d2i_rp>;
597 def INT_NVVM_D2UI_RN : F_MATH_1<"cvt.rni.u32.f64 \t$dst, $src0;",
598 Int32Regs, Float64Regs, int_nvvm_d2ui_rn>;
599 def INT_NVVM_D2UI_RZ : F_MATH_1<"cvt.rzi.u32.f64 \t$dst, $src0;",
600 Int32Regs, Float64Regs, int_nvvm_d2ui_rz>;
601 def INT_NVVM_D2UI_RM : F_MATH_1<"cvt.rmi.u32.f64 \t$dst, $src0;",
602 Int32Regs, Float64Regs, int_nvvm_d2ui_rm>;
603 def INT_NVVM_D2UI_RP : F_MATH_1<"cvt.rpi.u32.f64 \t$dst, $src0;",
604 Int32Regs, Float64Regs, int_nvvm_d2ui_rp>;
606 def INT_NVVM_I2D_RN : F_MATH_1<"cvt.rn.f64.s32 \t$dst, $src0;",
607 Float64Regs, Int32Regs, int_nvvm_i2d_rn>;
608 def INT_NVVM_I2D_RZ : F_MATH_1<"cvt.rz.f64.s32 \t$dst, $src0;",
609 Float64Regs, Int32Regs, int_nvvm_i2d_rz>;
610 def INT_NVVM_I2D_RM : F_MATH_1<"cvt.rm.f64.s32 \t$dst, $src0;",
611 Float64Regs, Int32Regs, int_nvvm_i2d_rm>;
612 def INT_NVVM_I2D_RP : F_MATH_1<"cvt.rp.f64.s32 \t$dst, $src0;",
613 Float64Regs, Int32Regs, int_nvvm_i2d_rp>;
615 def INT_NVVM_UI2D_RN : F_MATH_1<"cvt.rn.f64.u32 \t$dst, $src0;",
616 Float64Regs, Int32Regs, int_nvvm_ui2d_rn>;
617 def INT_NVVM_UI2D_RZ : F_MATH_1<"cvt.rz.f64.u32 \t$dst, $src0;",
618 Float64Regs, Int32Regs, int_nvvm_ui2d_rz>;
619 def INT_NVVM_UI2D_RM : F_MATH_1<"cvt.rm.f64.u32 \t$dst, $src0;",
620 Float64Regs, Int32Regs, int_nvvm_ui2d_rm>;
621 def INT_NVVM_UI2D_RP : F_MATH_1<"cvt.rp.f64.u32 \t$dst, $src0;",
622 Float64Regs, Int32Regs, int_nvvm_ui2d_rp>;
624 def INT_NVVM_F2I_RN_FTZ : F_MATH_1<"cvt.rni.ftz.s32.f32 \t$dst, $src0;",
625 Int32Regs, Float32Regs, int_nvvm_f2i_rn_ftz>;
626 def INT_NVVM_F2I_RN : F_MATH_1<"cvt.rni.s32.f32 \t$dst, $src0;", Int32Regs,
627 Float32Regs, int_nvvm_f2i_rn>;
628 def INT_NVVM_F2I_RZ_FTZ : F_MATH_1<"cvt.rzi.ftz.s32.f32 \t$dst, $src0;",
629 Int32Regs, Float32Regs, int_nvvm_f2i_rz_ftz>;
630 def INT_NVVM_F2I_RZ : F_MATH_1<"cvt.rzi.s32.f32 \t$dst, $src0;", Int32Regs,
631 Float32Regs, int_nvvm_f2i_rz>;
632 def INT_NVVM_F2I_RM_FTZ : F_MATH_1<"cvt.rmi.ftz.s32.f32 \t$dst, $src0;",
633 Int32Regs, Float32Regs, int_nvvm_f2i_rm_ftz>;
634 def INT_NVVM_F2I_RM : F_MATH_1<"cvt.rmi.s32.f32 \t$dst, $src0;", Int32Regs,
635 Float32Regs, int_nvvm_f2i_rm>;
636 def INT_NVVM_F2I_RP_FTZ : F_MATH_1<"cvt.rpi.ftz.s32.f32 \t$dst, $src0;",
637 Int32Regs, Float32Regs, int_nvvm_f2i_rp_ftz>;
638 def INT_NVVM_F2I_RP : F_MATH_1<"cvt.rpi.s32.f32 \t$dst, $src0;", Int32Regs,
639 Float32Regs, int_nvvm_f2i_rp>;
641 def INT_NVVM_F2UI_RN_FTZ : F_MATH_1<"cvt.rni.ftz.u32.f32 \t$dst, $src0;",
642 Int32Regs, Float32Regs, int_nvvm_f2ui_rn_ftz>;
643 def INT_NVVM_F2UI_RN : F_MATH_1<"cvt.rni.u32.f32 \t$dst, $src0;", Int32Regs,
644 Float32Regs, int_nvvm_f2ui_rn>;
645 def INT_NVVM_F2UI_RZ_FTZ : F_MATH_1<"cvt.rzi.ftz.u32.f32 \t$dst, $src0;",
646 Int32Regs, Float32Regs, int_nvvm_f2ui_rz_ftz>;
647 def INT_NVVM_F2UI_RZ : F_MATH_1<"cvt.rzi.u32.f32 \t$dst, $src0;", Int32Regs,
648 Float32Regs, int_nvvm_f2ui_rz>;
649 def INT_NVVM_F2UI_RM_FTZ : F_MATH_1<"cvt.rmi.ftz.u32.f32 \t$dst, $src0;",
650 Int32Regs, Float32Regs, int_nvvm_f2ui_rm_ftz>;
651 def INT_NVVM_F2UI_RM : F_MATH_1<"cvt.rmi.u32.f32 \t$dst, $src0;", Int32Regs,
652 Float32Regs, int_nvvm_f2ui_rm>;
653 def INT_NVVM_F2UI_RP_FTZ : F_MATH_1<"cvt.rpi.ftz.u32.f32 \t$dst, $src0;",
654 Int32Regs, Float32Regs, int_nvvm_f2ui_rp_ftz>;
655 def INT_NVVM_F2UI_RP : F_MATH_1<"cvt.rpi.u32.f32 \t$dst, $src0;", Int32Regs,
656 Float32Regs, int_nvvm_f2ui_rp>;
658 def INT_NVVM_I2F_RN : F_MATH_1<"cvt.rn.f32.s32 \t$dst, $src0;", Float32Regs,
659 Int32Regs, int_nvvm_i2f_rn>;
660 def INT_NVVM_I2F_RZ : F_MATH_1<"cvt.rz.f32.s32 \t$dst, $src0;", Float32Regs,
661 Int32Regs, int_nvvm_i2f_rz>;
662 def INT_NVVM_I2F_RM : F_MATH_1<"cvt.rm.f32.s32 \t$dst, $src0;", Float32Regs,
663 Int32Regs, int_nvvm_i2f_rm>;
664 def INT_NVVM_I2F_RP : F_MATH_1<"cvt.rp.f32.s32 \t$dst, $src0;", Float32Regs,
665 Int32Regs, int_nvvm_i2f_rp>;
667 def INT_NVVM_UI2F_RN : F_MATH_1<"cvt.rn.f32.u32 \t$dst, $src0;", Float32Regs,
668 Int32Regs, int_nvvm_ui2f_rn>;
669 def INT_NVVM_UI2F_RZ : F_MATH_1<"cvt.rz.f32.u32 \t$dst, $src0;", Float32Regs,
670 Int32Regs, int_nvvm_ui2f_rz>;
671 def INT_NVVM_UI2F_RM : F_MATH_1<"cvt.rm.f32.u32 \t$dst, $src0;", Float32Regs,
672 Int32Regs, int_nvvm_ui2f_rm>;
673 def INT_NVVM_UI2F_RP : F_MATH_1<"cvt.rp.f32.u32 \t$dst, $src0;", Float32Regs,
674 Int32Regs, int_nvvm_ui2f_rp>;
676 def INT_NVVM_LOHI_I2D : F_MATH_2<"mov.b64 \t$dst, {{$src0, $src1}};",
677 Float64Regs, Int32Regs, Int32Regs, int_nvvm_lohi_i2d>;
679 def INT_NVVM_D2I_LO : F_MATH_1<!strconcat("{{\n\t",
680 !strconcat(".reg .b32 %temp; \n\t",
681 !strconcat("mov.b64 \t{$dst, %temp}, $src0;\n\t",
683 Int32Regs, Float64Regs, int_nvvm_d2i_lo>;
684 def INT_NVVM_D2I_HI : F_MATH_1<!strconcat("{{\n\t",
685 !strconcat(".reg .b32 %temp; \n\t",
686 !strconcat("mov.b64 \t{%temp, $dst}, $src0;\n\t",
688 Int32Regs, Float64Regs, int_nvvm_d2i_hi>;
690 def INT_NVVM_F2LL_RN_FTZ : F_MATH_1<"cvt.rni.ftz.s64.f32 \t$dst, $src0;",
691 Int64Regs, Float32Regs, int_nvvm_f2ll_rn_ftz>;
692 def INT_NVVM_F2LL_RN : F_MATH_1<"cvt.rni.s64.f32 \t$dst, $src0;", Int64Regs,
693 Float32Regs, int_nvvm_f2ll_rn>;
694 def INT_NVVM_F2LL_RZ_FTZ : F_MATH_1<"cvt.rzi.ftz.s64.f32 \t$dst, $src0;",
695 Int64Regs, Float32Regs, int_nvvm_f2ll_rz_ftz>;
696 def INT_NVVM_F2LL_RZ : F_MATH_1<"cvt.rzi.s64.f32 \t$dst, $src0;", Int64Regs,
697 Float32Regs, int_nvvm_f2ll_rz>;
698 def INT_NVVM_F2LL_RM_FTZ : F_MATH_1<"cvt.rmi.ftz.s64.f32 \t$dst, $src0;",
699 Int64Regs, Float32Regs, int_nvvm_f2ll_rm_ftz>;
700 def INT_NVVM_F2LL_RM : F_MATH_1<"cvt.rmi.s64.f32 \t$dst, $src0;", Int64Regs,
701 Float32Regs, int_nvvm_f2ll_rm>;
702 def INT_NVVM_F2LL_RP_FTZ : F_MATH_1<"cvt.rpi.ftz.s64.f32 \t$dst, $src0;",
703 Int64Regs, Float32Regs, int_nvvm_f2ll_rp_ftz>;
704 def INT_NVVM_F2LL_RP : F_MATH_1<"cvt.rpi.s64.f32 \t$dst, $src0;", Int64Regs,
705 Float32Regs, int_nvvm_f2ll_rp>;
707 def INT_NVVM_F2ULL_RN_FTZ : F_MATH_1<"cvt.rni.ftz.u64.f32 \t$dst, $src0;",
708 Int64Regs, Float32Regs, int_nvvm_f2ull_rn_ftz>;
709 def INT_NVVM_F2ULL_RN : F_MATH_1<"cvt.rni.u64.f32 \t$dst, $src0;", Int64Regs,
710 Float32Regs, int_nvvm_f2ull_rn>;
711 def INT_NVVM_F2ULL_RZ_FTZ : F_MATH_1<"cvt.rzi.ftz.u64.f32 \t$dst, $src0;",
712 Int64Regs, Float32Regs, int_nvvm_f2ull_rz_ftz>;
713 def INT_NVVM_F2ULL_RZ : F_MATH_1<"cvt.rzi.u64.f32 \t$dst, $src0;", Int64Regs,
714 Float32Regs, int_nvvm_f2ull_rz>;
715 def INT_NVVM_F2ULL_RM_FTZ : F_MATH_1<"cvt.rmi.ftz.u64.f32 \t$dst, $src0;",
716 Int64Regs, Float32Regs, int_nvvm_f2ull_rm_ftz>;
717 def INT_NVVM_F2ULL_RM : F_MATH_1<"cvt.rmi.u64.f32 \t$dst, $src0;", Int64Regs,
718 Float32Regs, int_nvvm_f2ull_rm>;
719 def INT_NVVM_F2ULL_RP_FTZ : F_MATH_1<"cvt.rpi.ftz.u64.f32 \t$dst, $src0;",
720 Int64Regs, Float32Regs, int_nvvm_f2ull_rp_ftz>;
721 def INT_NVVM_F2ULL_RP : F_MATH_1<"cvt.rpi.u64.f32 \t$dst, $src0;", Int64Regs,
722 Float32Regs, int_nvvm_f2ull_rp>;
724 def INT_NVVM_D2LL_RN : F_MATH_1<"cvt.rni.s64.f64 \t$dst, $src0;", Int64Regs,
725 Float64Regs, int_nvvm_d2ll_rn>;
726 def INT_NVVM_D2LL_RZ : F_MATH_1<"cvt.rzi.s64.f64 \t$dst, $src0;", Int64Regs,
727 Float64Regs, int_nvvm_d2ll_rz>;
728 def INT_NVVM_D2LL_RM : F_MATH_1<"cvt.rmi.s64.f64 \t$dst, $src0;", Int64Regs,
729 Float64Regs, int_nvvm_d2ll_rm>;
730 def INT_NVVM_D2LL_RP : F_MATH_1<"cvt.rpi.s64.f64 \t$dst, $src0;", Int64Regs,
731 Float64Regs, int_nvvm_d2ll_rp>;
733 def INT_NVVM_D2ULL_RN : F_MATH_1<"cvt.rni.u64.f64 \t$dst, $src0;", Int64Regs,
734 Float64Regs, int_nvvm_d2ull_rn>;
735 def INT_NVVM_D2ULL_RZ : F_MATH_1<"cvt.rzi.u64.f64 \t$dst, $src0;", Int64Regs,
736 Float64Regs, int_nvvm_d2ull_rz>;
737 def INT_NVVM_D2ULL_RM : F_MATH_1<"cvt.rmi.u64.f64 \t$dst, $src0;", Int64Regs,
738 Float64Regs, int_nvvm_d2ull_rm>;
739 def INT_NVVM_D2ULL_RP : F_MATH_1<"cvt.rpi.u64.f64 \t$dst, $src0;", Int64Regs,
740 Float64Regs, int_nvvm_d2ull_rp>;
742 def INT_NVVM_LL2F_RN : F_MATH_1<"cvt.rn.f32.s64 \t$dst, $src0;", Float32Regs,
743 Int64Regs, int_nvvm_ll2f_rn>;
744 def INT_NVVM_LL2F_RZ : F_MATH_1<"cvt.rz.f32.s64 \t$dst, $src0;", Float32Regs,
745 Int64Regs, int_nvvm_ll2f_rz>;
746 def INT_NVVM_LL2F_RM : F_MATH_1<"cvt.rm.f32.s64 \t$dst, $src0;", Float32Regs,
747 Int64Regs, int_nvvm_ll2f_rm>;
748 def INT_NVVM_LL2F_RP : F_MATH_1<"cvt.rp.f32.s64 \t$dst, $src0;", Float32Regs,
749 Int64Regs, int_nvvm_ll2f_rp>;
750 def INT_NVVM_ULL2F_RN : F_MATH_1<"cvt.rn.f32.u64 \t$dst, $src0;", Float32Regs,
751 Int64Regs, int_nvvm_ull2f_rn>;
752 def INT_NVVM_ULL2F_RZ : F_MATH_1<"cvt.rz.f32.u64 \t$dst, $src0;", Float32Regs,
753 Int64Regs, int_nvvm_ull2f_rz>;
754 def INT_NVVM_ULL2F_RM : F_MATH_1<"cvt.rm.f32.u64 \t$dst, $src0;", Float32Regs,
755 Int64Regs, int_nvvm_ull2f_rm>;
756 def INT_NVVM_ULL2F_RP : F_MATH_1<"cvt.rp.f32.u64 \t$dst, $src0;", Float32Regs,
757 Int64Regs, int_nvvm_ull2f_rp>;
759 def INT_NVVM_LL2D_RN : F_MATH_1<"cvt.rn.f64.s64 \t$dst, $src0;", Float64Regs,
760 Int64Regs, int_nvvm_ll2d_rn>;
761 def INT_NVVM_LL2D_RZ : F_MATH_1<"cvt.rz.f64.s64 \t$dst, $src0;", Float64Regs,
762 Int64Regs, int_nvvm_ll2d_rz>;
763 def INT_NVVM_LL2D_RM : F_MATH_1<"cvt.rm.f64.s64 \t$dst, $src0;", Float64Regs,
764 Int64Regs, int_nvvm_ll2d_rm>;
765 def INT_NVVM_LL2D_RP : F_MATH_1<"cvt.rp.f64.s64 \t$dst, $src0;", Float64Regs,
766 Int64Regs, int_nvvm_ll2d_rp>;
767 def INT_NVVM_ULL2D_RN : F_MATH_1<"cvt.rn.f64.u64 \t$dst, $src0;", Float64Regs,
768 Int64Regs, int_nvvm_ull2d_rn>;
769 def INT_NVVM_ULL2D_RZ : F_MATH_1<"cvt.rz.f64.u64 \t$dst, $src0;", Float64Regs,
770 Int64Regs, int_nvvm_ull2d_rz>;
771 def INT_NVVM_ULL2D_RM : F_MATH_1<"cvt.rm.f64.u64 \t$dst, $src0;", Float64Regs,
772 Int64Regs, int_nvvm_ull2d_rm>;
773 def INT_NVVM_ULL2D_RP : F_MATH_1<"cvt.rp.f64.u64 \t$dst, $src0;", Float64Regs,
774 Int64Regs, int_nvvm_ull2d_rp>;
776 def INT_NVVM_F2H_RN_FTZ : F_MATH_1<!strconcat("{{\n\t",
777 !strconcat(".reg .b16 %temp;\n\t",
778 !strconcat("cvt.rn.ftz.f16.f32 \t%temp, $src0;\n\t",
779 !strconcat("mov.b16 \t$dst, %temp;\n",
781 Int16Regs, Float32Regs, int_nvvm_f2h_rn_ftz>;
782 def INT_NVVM_F2H_RN : F_MATH_1<!strconcat("{{\n\t",
783 !strconcat(".reg .b16 %temp;\n\t",
784 !strconcat("cvt.rn.f16.f32 \t%temp, $src0;\n\t",
785 !strconcat("mov.b16 \t$dst, %temp;\n",
787 Int16Regs, Float32Regs, int_nvvm_f2h_rn>;
789 def INT_NVVM_H2F : F_MATH_1<!strconcat("{{\n\t",
790 !strconcat(".reg .b16 %temp;\n\t",
791 !strconcat("mov.b16 \t%temp, $src0;\n\t",
792 !strconcat("cvt.f32.f16 \t$dst, %temp;\n\t",
794 Float32Regs, Int16Regs, int_nvvm_h2f>;
800 def INT_NVVM_BITCAST_F2I : F_MATH_1<"mov.b32 \t$dst, $src0;", Int32Regs,
801 Float32Regs, int_nvvm_bitcast_f2i>;
802 def INT_NVVM_BITCAST_I2F : F_MATH_1<"mov.b32 \t$dst, $src0;", Float32Regs,
803 Int32Regs, int_nvvm_bitcast_i2f>;
805 def INT_NVVM_BITCAST_LL2D : F_MATH_1<"mov.b64 \t$dst, $src0;", Float64Regs,
806 Int64Regs, int_nvvm_bitcast_ll2d>;
807 def INT_NVVM_BITCAST_D2LL : F_MATH_1<"mov.b64 \t$dst, $src0;", Int64Regs,
808 Float64Regs, int_nvvm_bitcast_d2ll>;
810 //-----------------------------------
812 //-----------------------------------
814 class ATOMIC_GLOBAL_CHK <dag ops, dag frag>
815 : PatFrag<ops, frag, [{
816 return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_GLOBAL);
818 class ATOMIC_SHARED_CHK <dag ops, dag frag>
819 : PatFrag<ops, frag, [{
820 return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_SHARED);
822 class ATOMIC_GENERIC_CHK <dag ops, dag frag>
823 : PatFrag<ops, frag, [{
824 return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_GENERIC);
827 multiclass F_ATOMIC_2_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
828 string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
829 Operand IMMType, SDNode IMM, Predicate Pred> {
830 def reg : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, regclass:$b),
835 !strconcat(" \t$dst, [$addr], $b;", ""))))),
836 [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b))]>,
838 def imm : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, IMMType:$b),
843 !strconcat(" \t$dst, [$addr], $b;", ""))))),
844 [(set regclass:$dst, (IntOp ptrclass:$addr, IMM:$b))]>,
847 multiclass F_ATOMIC_2<NVPTXRegClass regclass, string SpaceStr, string TypeStr,
848 string OpcStr, PatFrag IntOp, Operand IMMType, SDNode IMM, Predicate Pred> {
849 defm p32 : F_ATOMIC_2_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
850 IntOp, IMMType, IMM, Pred>;
851 defm p64 : F_ATOMIC_2_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
852 IntOp, IMMType, IMM, Pred>;
855 // has 2 operands, neg the second one
856 multiclass F_ATOMIC_2_NEG_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
857 string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
858 Operand IMMType, Predicate Pred> {
859 def reg : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, regclass:$b),
860 !strconcat("{{ \n\t",
861 !strconcat(".reg \t.s",
863 !strconcat(" temp; \n\t",
866 !strconcat(" \ttemp, $b; \n\t",
872 !strconcat(" \t$dst, [$addr], temp; \n\t",
873 !strconcat("}}", "")))))))))))))),
874 [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b))]>,
877 multiclass F_ATOMIC_2_NEG<NVPTXRegClass regclass, string SpaceStr,
878 string TypeStr, string OpcStr, PatFrag IntOp, Operand IMMType,
880 defm p32: F_ATOMIC_2_NEG_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
881 IntOp, IMMType, Pred> ;
882 defm p64: F_ATOMIC_2_NEG_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
883 IntOp, IMMType, Pred> ;
887 multiclass F_ATOMIC_3_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
888 string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
889 Operand IMMType, Predicate Pred> {
890 def reg : NVPTXInst<(outs regclass:$dst),
891 (ins ptrclass:$addr, regclass:$b, regclass:$c),
896 !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
898 (IntOp ptrclass:$addr, regclass:$b, regclass:$c))]>,
900 def imm1 : NVPTXInst<(outs regclass:$dst),
901 (ins ptrclass:$addr, IMMType:$b, regclass:$c),
906 !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
907 [(set regclass:$dst, (IntOp ptrclass:$addr, imm:$b, regclass:$c))]>,
909 def imm2 : NVPTXInst<(outs regclass:$dst),
910 (ins ptrclass:$addr, regclass:$b, IMMType:$c),
915 !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
916 [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b, imm:$c))]>,
918 def imm3 : NVPTXInst<(outs regclass:$dst),
919 (ins ptrclass:$addr, IMMType:$b, IMMType:$c),
924 !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
925 [(set regclass:$dst, (IntOp ptrclass:$addr, imm:$b, imm:$c))]>,
928 multiclass F_ATOMIC_3<NVPTXRegClass regclass, string SpaceStr, string TypeStr,
929 string OpcStr, PatFrag IntOp, Operand IMMType, Predicate Pred> {
930 defm p32 : F_ATOMIC_3_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
931 IntOp, IMMType, Pred>;
932 defm p64 : F_ATOMIC_3_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
933 IntOp, IMMType, Pred>;
938 def atomic_load_add_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
939 (atomic_load_add_32 node:$a, node:$b)>;
940 def atomic_load_add_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
941 (atomic_load_add_32 node:$a, node:$b)>;
942 def atomic_load_add_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
943 (atomic_load_add_32 node:$a, node:$b)>;
944 def atomic_load_add_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
945 (atomic_load_add_64 node:$a, node:$b)>;
946 def atomic_load_add_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
947 (atomic_load_add_64 node:$a, node:$b)>;
948 def atomic_load_add_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
949 (atomic_load_add_64 node:$a, node:$b)>;
950 def atomic_load_add_f32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
951 (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
952 def atomic_load_add_f32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
953 (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
954 def atomic_load_add_f32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
955 (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
957 defm INT_PTX_ATOM_ADD_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".add",
958 atomic_load_add_32_g, i32imm, imm, hasAtomRedG32>;
959 defm INT_PTX_ATOM_ADD_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".add",
960 atomic_load_add_32_s, i32imm, imm, hasAtomRedS32>;
961 defm INT_PTX_ATOM_ADD_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".add",
962 atomic_load_add_32_gen, i32imm, imm, hasAtomRedGen32>;
963 defm INT_PTX_ATOM_ADD_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
964 ".add", atomic_load_add_32_gen, i32imm, imm, useAtomRedG32forGen32>;
966 defm INT_PTX_ATOM_ADD_G_64 : F_ATOMIC_2<Int64Regs, ".global", ".u64", ".add",
967 atomic_load_add_64_g, i64imm, imm, hasAtomRedG64>;
968 defm INT_PTX_ATOM_ADD_S_64 : F_ATOMIC_2<Int64Regs, ".shared", ".u64", ".add",
969 atomic_load_add_64_s, i64imm, imm, hasAtomRedS64>;
970 defm INT_PTX_ATOM_ADD_GEN_64 : F_ATOMIC_2<Int64Regs, "", ".u64", ".add",
971 atomic_load_add_64_gen, i64imm, imm, hasAtomRedGen64>;
972 defm INT_PTX_ATOM_ADD_GEN_64_USE_G : F_ATOMIC_2<Int64Regs, ".global", ".u64",
973 ".add", atomic_load_add_64_gen, i64imm, imm, useAtomRedG64forGen64>;
975 defm INT_PTX_ATOM_ADD_G_F32 : F_ATOMIC_2<Float32Regs, ".global", ".f32", ".add",
976 atomic_load_add_f32_g, f32imm, fpimm, hasAtomAddF32>;
977 defm INT_PTX_ATOM_ADD_S_F32 : F_ATOMIC_2<Float32Regs, ".shared", ".f32", ".add",
978 atomic_load_add_f32_s, f32imm, fpimm, hasAtomAddF32>;
979 defm INT_PTX_ATOM_ADD_GEN_F32 : F_ATOMIC_2<Float32Regs, "", ".f32", ".add",
980 atomic_load_add_f32_gen, f32imm, fpimm, hasAtomAddF32>;
984 def atomic_load_sub_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
985 (atomic_load_sub_32 node:$a, node:$b)>;
986 def atomic_load_sub_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
987 (atomic_load_sub_32 node:$a, node:$b)>;
988 def atomic_load_sub_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
989 (atomic_load_sub_32 node:$a, node:$b)>;
990 def atomic_load_sub_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
991 (atomic_load_sub_64 node:$a, node:$b)>;
992 def atomic_load_sub_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
993 (atomic_load_sub_64 node:$a, node:$b)>;
994 def atomic_load_sub_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
995 (atomic_load_sub_64 node:$a, node:$b)>;
997 defm INT_PTX_ATOM_SUB_G_32 : F_ATOMIC_2_NEG<Int32Regs, ".global", "32", ".add",
998 atomic_load_sub_32_g, i32imm, hasAtomRedG32>;
999 defm INT_PTX_ATOM_SUB_G_64 : F_ATOMIC_2_NEG<Int64Regs, ".global", "64", ".add",
1000 atomic_load_sub_64_g, i64imm, hasAtomRedG64>;
1001 defm INT_PTX_ATOM_SUB_GEN_32 : F_ATOMIC_2_NEG<Int32Regs, "", "32", ".add",
1002 atomic_load_sub_32_gen, i32imm, hasAtomRedGen32>;
1003 defm INT_PTX_ATOM_SUB_GEN_32_USE_G : F_ATOMIC_2_NEG<Int32Regs, ".global", "32",
1004 ".add", atomic_load_sub_32_gen, i32imm, useAtomRedG32forGen32>;
1005 defm INT_PTX_ATOM_SUB_S_32 : F_ATOMIC_2_NEG<Int32Regs, ".shared", "32", ".add",
1006 atomic_load_sub_32_s, i32imm, hasAtomRedS32>;
1007 defm INT_PTX_ATOM_SUB_S_64 : F_ATOMIC_2_NEG<Int64Regs, ".shared", "64", ".add",
1008 atomic_load_sub_64_s, i64imm, hasAtomRedS64>;
1009 defm INT_PTX_ATOM_SUB_GEN_64 : F_ATOMIC_2_NEG<Int64Regs, "", "64", ".add",
1010 atomic_load_sub_64_gen, i64imm, hasAtomRedGen64>;
1011 defm INT_PTX_ATOM_SUB_GEN_64_USE_G : F_ATOMIC_2_NEG<Int64Regs, ".global", "64",
1012 ".add", atomic_load_sub_64_gen, i64imm, useAtomRedG64forGen64>;
1016 def atomic_swap_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1017 (atomic_swap_32 node:$a, node:$b)>;
1018 def atomic_swap_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1019 (atomic_swap_32 node:$a, node:$b)>;
1020 def atomic_swap_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1021 (atomic_swap_32 node:$a, node:$b)>;
1022 def atomic_swap_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1023 (atomic_swap_64 node:$a, node:$b)>;
1024 def atomic_swap_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1025 (atomic_swap_64 node:$a, node:$b)>;
1026 def atomic_swap_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1027 (atomic_swap_64 node:$a, node:$b)>;
1029 defm INT_PTX_ATOM_SWAP_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".exch",
1030 atomic_swap_32_g, i32imm, imm, hasAtomRedG32>;
1031 defm INT_PTX_ATOM_SWAP_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".exch",
1032 atomic_swap_32_s, i32imm, imm, hasAtomRedS32>;
1033 defm INT_PTX_ATOM_SWAP_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".exch",
1034 atomic_swap_32_gen, i32imm, imm, hasAtomRedGen32>;
1035 defm INT_PTX_ATOM_SWAP_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1036 ".exch", atomic_swap_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1037 defm INT_PTX_ATOM_SWAP_G_64 : F_ATOMIC_2<Int64Regs, ".global", ".b64", ".exch",
1038 atomic_swap_64_g, i64imm, imm, hasAtomRedG64>;
1039 defm INT_PTX_ATOM_SWAP_S_64 : F_ATOMIC_2<Int64Regs, ".shared", ".b64", ".exch",
1040 atomic_swap_64_s, i64imm, imm, hasAtomRedS64>;
1041 defm INT_PTX_ATOM_SWAP_GEN_64 : F_ATOMIC_2<Int64Regs, "", ".b64", ".exch",
1042 atomic_swap_64_gen, i64imm, imm, hasAtomRedGen64>;
1043 defm INT_PTX_ATOM_SWAP_GEN_64_USE_G : F_ATOMIC_2<Int64Regs, ".global", ".b64",
1044 ".exch", atomic_swap_64_gen, i64imm, imm, useAtomRedG64forGen64>;
1048 def atomic_load_max_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b)
1049 , (atomic_load_max_32 node:$a, node:$b)>;
1050 def atomic_load_max_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1051 (atomic_load_max_32 node:$a, node:$b)>;
1052 def atomic_load_max_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1053 (atomic_load_max_32 node:$a, node:$b)>;
1054 def atomic_load_umax_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1055 (atomic_load_umax_32 node:$a, node:$b)>;
1056 def atomic_load_umax_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1057 (atomic_load_umax_32 node:$a, node:$b)>;
1058 def atomic_load_umax_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1059 (atomic_load_umax_32 node:$a, node:$b)>;
1061 defm INT_PTX_ATOM_LOAD_MAX_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".s32",
1062 ".max", atomic_load_max_32_g, i32imm, imm, hasAtomRedG32>;
1063 defm INT_PTX_ATOM_LOAD_MAX_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".s32",
1064 ".max", atomic_load_max_32_s, i32imm, imm, hasAtomRedS32>;
1065 defm INT_PTX_ATOM_LOAD_MAX_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".s32", ".max",
1066 atomic_load_max_32_gen, i32imm, imm, hasAtomRedGen32>;
1067 defm INT_PTX_ATOM_LOAD_MAX_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1068 ".s32", ".max", atomic_load_max_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1069 defm INT_PTX_ATOM_LOAD_UMAX_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1070 ".max", atomic_load_umax_32_g, i32imm, imm, hasAtomRedG32>;
1071 defm INT_PTX_ATOM_LOAD_UMAX_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32",
1072 ".max", atomic_load_umax_32_s, i32imm, imm, hasAtomRedS32>;
1073 defm INT_PTX_ATOM_LOAD_UMAX_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".max",
1074 atomic_load_umax_32_gen, i32imm, imm, hasAtomRedGen32>;
1075 defm INT_PTX_ATOM_LOAD_UMAX_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1076 ".u32", ".max", atomic_load_umax_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1080 def atomic_load_min_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1081 (atomic_load_min_32 node:$a, node:$b)>;
1082 def atomic_load_min_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1083 (atomic_load_min_32 node:$a, node:$b)>;
1084 def atomic_load_min_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1085 (atomic_load_min_32 node:$a, node:$b)>;
1086 def atomic_load_umin_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1087 (atomic_load_umin_32 node:$a, node:$b)>;
1088 def atomic_load_umin_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1089 (atomic_load_umin_32 node:$a, node:$b)>;
1090 def atomic_load_umin_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1091 (atomic_load_umin_32 node:$a, node:$b)>;
1093 defm INT_PTX_ATOM_LOAD_MIN_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".s32",
1094 ".min", atomic_load_min_32_g, i32imm, imm, hasAtomRedG32>;
1095 defm INT_PTX_ATOM_LOAD_MIN_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".s32",
1096 ".min", atomic_load_min_32_s, i32imm, imm, hasAtomRedS32>;
1097 defm INT_PTX_ATOM_LOAD_MIN_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".s32", ".min",
1098 atomic_load_min_32_gen, i32imm, imm, hasAtomRedGen32>;
1099 defm INT_PTX_ATOM_LOAD_MIN_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1100 ".s32", ".min", atomic_load_min_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1101 defm INT_PTX_ATOM_LOAD_UMIN_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1102 ".min", atomic_load_umin_32_g, i32imm, imm, hasAtomRedG32>;
1103 defm INT_PTX_ATOM_LOAD_UMIN_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32",
1104 ".min", atomic_load_umin_32_s, i32imm, imm, hasAtomRedS32>;
1105 defm INT_PTX_ATOM_LOAD_UMIN_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".min",
1106 atomic_load_umin_32_gen, i32imm, imm, hasAtomRedGen32>;
1107 defm INT_PTX_ATOM_LOAD_UMIN_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1108 ".u32", ".min", atomic_load_umin_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1110 // atom_inc atom_dec
1112 def atomic_load_inc_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1113 (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
1114 def atomic_load_inc_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1115 (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
1116 def atomic_load_inc_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1117 (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
1118 def atomic_load_dec_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1119 (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
1120 def atomic_load_dec_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1121 (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
1122 def atomic_load_dec_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1123 (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
1125 defm INT_PTX_ATOM_INC_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".inc",
1126 atomic_load_inc_32_g, i32imm, imm, hasAtomRedG32>;
1127 defm INT_PTX_ATOM_INC_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".inc",
1128 atomic_load_inc_32_s, i32imm, imm, hasAtomRedS32>;
1129 defm INT_PTX_ATOM_INC_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".inc",
1130 atomic_load_inc_32_gen, i32imm, imm, hasAtomRedGen32>;
1131 defm INT_PTX_ATOM_INC_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1132 ".inc", atomic_load_inc_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1133 defm INT_PTX_ATOM_DEC_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".dec",
1134 atomic_load_dec_32_g, i32imm, imm, hasAtomRedG32>;
1135 defm INT_PTX_ATOM_DEC_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".dec",
1136 atomic_load_dec_32_s, i32imm, imm, hasAtomRedS32>;
1137 defm INT_PTX_ATOM_DEC_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".dec",
1138 atomic_load_dec_32_gen, i32imm, imm, hasAtomRedGen32>;
1139 defm INT_PTX_ATOM_DEC_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1140 ".dec", atomic_load_dec_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1144 def atomic_load_and_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1145 (atomic_load_and_32 node:$a, node:$b)>;
1146 def atomic_load_and_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1147 (atomic_load_and_32 node:$a, node:$b)>;
1148 def atomic_load_and_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1149 (atomic_load_and_32 node:$a, node:$b)>;
1151 defm INT_PTX_ATOM_AND_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".and",
1152 atomic_load_and_32_g, i32imm, imm, hasAtomRedG32>;
1153 defm INT_PTX_ATOM_AND_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".and",
1154 atomic_load_and_32_s, i32imm, imm, hasAtomRedS32>;
1155 defm INT_PTX_ATOM_AND_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".and",
1156 atomic_load_and_32_gen, i32imm, imm, hasAtomRedGen32>;
1157 defm INT_PTX_ATOM_AND_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1158 ".and", atomic_load_and_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1162 def atomic_load_or_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1163 (atomic_load_or_32 node:$a, node:$b)>;
1164 def atomic_load_or_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1165 (atomic_load_or_32 node:$a, node:$b)>;
1166 def atomic_load_or_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1167 (atomic_load_or_32 node:$a, node:$b)>;
1169 defm INT_PTX_ATOM_OR_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".or",
1170 atomic_load_or_32_g, i32imm, imm, hasAtomRedG32>;
1171 defm INT_PTX_ATOM_OR_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".or",
1172 atomic_load_or_32_gen, i32imm, imm, hasAtomRedGen32>;
1173 defm INT_PTX_ATOM_OR_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1174 ".or", atomic_load_or_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1175 defm INT_PTX_ATOM_OR_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".or",
1176 atomic_load_or_32_s, i32imm, imm, hasAtomRedS32>;
1180 def atomic_load_xor_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1181 (atomic_load_xor_32 node:$a, node:$b)>;
1182 def atomic_load_xor_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1183 (atomic_load_xor_32 node:$a, node:$b)>;
1184 def atomic_load_xor_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1185 (atomic_load_xor_32 node:$a, node:$b)>;
1187 defm INT_PTX_ATOM_XOR_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".xor",
1188 atomic_load_xor_32_g, i32imm, imm, hasAtomRedG32>;
1189 defm INT_PTX_ATOM_XOR_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".xor",
1190 atomic_load_xor_32_s, i32imm, imm, hasAtomRedS32>;
1191 defm INT_PTX_ATOM_XOR_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".xor",
1192 atomic_load_xor_32_gen, i32imm, imm, hasAtomRedGen32>;
1193 defm INT_PTX_ATOM_XOR_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1194 ".xor", atomic_load_xor_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1198 def atomic_cmp_swap_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b, node:$c),
1199 (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
1200 def atomic_cmp_swap_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b, node:$c),
1201 (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
1202 def atomic_cmp_swap_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b, node:$c),
1203 (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
1204 def atomic_cmp_swap_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b, node:$c),
1205 (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
1206 def atomic_cmp_swap_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b, node:$c),
1207 (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
1208 def atomic_cmp_swap_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b, node:$c),
1209 (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
1211 defm INT_PTX_ATOM_CAS_G_32 : F_ATOMIC_3<Int32Regs, ".global", ".b32", ".cas",
1212 atomic_cmp_swap_32_g, i32imm, hasAtomRedG32>;
1213 defm INT_PTX_ATOM_CAS_S_32 : F_ATOMIC_3<Int32Regs, ".shared", ".b32", ".cas",
1214 atomic_cmp_swap_32_s, i32imm, hasAtomRedS32>;
1215 defm INT_PTX_ATOM_CAS_GEN_32 : F_ATOMIC_3<Int32Regs, "", ".b32", ".cas",
1216 atomic_cmp_swap_32_gen, i32imm, hasAtomRedGen32>;
1217 defm INT_PTX_ATOM_CAS_GEN_32_USE_G : F_ATOMIC_3<Int32Regs, ".global", ".b32",
1218 ".cas", atomic_cmp_swap_32_gen, i32imm, useAtomRedG32forGen32>;
1219 defm INT_PTX_ATOM_CAS_G_64 : F_ATOMIC_3<Int64Regs, ".global", ".b64", ".cas",
1220 atomic_cmp_swap_64_g, i64imm, hasAtomRedG64>;
1221 defm INT_PTX_ATOM_CAS_S_64 : F_ATOMIC_3<Int64Regs, ".shared", ".b64", ".cas",
1222 atomic_cmp_swap_64_s, i64imm, hasAtomRedS64>;
1223 defm INT_PTX_ATOM_CAS_GEN_64 : F_ATOMIC_3<Int64Regs, "", ".b64", ".cas",
1224 atomic_cmp_swap_64_gen, i64imm, hasAtomRedGen64>;
1225 defm INT_PTX_ATOM_CAS_GEN_64_USE_G : F_ATOMIC_3<Int64Regs, ".global", ".b64",
1226 ".cas", atomic_cmp_swap_64_gen, i64imm, useAtomRedG64forGen64>;
1229 //-----------------------------------
1230 // Read Special Registers
1231 //-----------------------------------
1232 class F_SREG<string OpStr, NVPTXRegClass regclassOut, Intrinsic IntOp> :
1233 NVPTXInst<(outs regclassOut:$dst), (ins),
1235 [(set regclassOut:$dst, (IntOp))]>;
1237 def INT_PTX_SREG_TID_X : F_SREG<"mov.u32 \t$dst, %tid.x;", Int32Regs,
1238 int_nvvm_read_ptx_sreg_tid_x>;
1239 def INT_PTX_SREG_TID_Y : F_SREG<"mov.u32 \t$dst, %tid.y;", Int32Regs,
1240 int_nvvm_read_ptx_sreg_tid_y>;
1241 def INT_PTX_SREG_TID_Z : F_SREG<"mov.u32 \t$dst, %tid.z;", Int32Regs,
1242 int_nvvm_read_ptx_sreg_tid_z>;
1244 def INT_PTX_SREG_NTID_X : F_SREG<"mov.u32 \t$dst, %ntid.x;", Int32Regs,
1245 int_nvvm_read_ptx_sreg_ntid_x>;
1246 def INT_PTX_SREG_NTID_Y : F_SREG<"mov.u32 \t$dst, %ntid.y;", Int32Regs,
1247 int_nvvm_read_ptx_sreg_ntid_y>;
1248 def INT_PTX_SREG_NTID_Z : F_SREG<"mov.u32 \t$dst, %ntid.z;", Int32Regs,
1249 int_nvvm_read_ptx_sreg_ntid_z>;
1251 def INT_PTX_SREG_CTAID_X : F_SREG<"mov.u32 \t$dst, %ctaid.x;", Int32Regs,
1252 int_nvvm_read_ptx_sreg_ctaid_x>;
1253 def INT_PTX_SREG_CTAID_Y : F_SREG<"mov.u32 \t$dst, %ctaid.y;", Int32Regs,
1254 int_nvvm_read_ptx_sreg_ctaid_y>;
1255 def INT_PTX_SREG_CTAID_Z : F_SREG<"mov.u32 \t$dst, %ctaid.z;", Int32Regs,
1256 int_nvvm_read_ptx_sreg_ctaid_z>;
1258 def INT_PTX_SREG_NCTAID_X : F_SREG<"mov.u32 \t$dst, %nctaid.x;", Int32Regs,
1259 int_nvvm_read_ptx_sreg_nctaid_x>;
1260 def INT_PTX_SREG_NCTAID_Y : F_SREG<"mov.u32 \t$dst, %nctaid.y;", Int32Regs,
1261 int_nvvm_read_ptx_sreg_nctaid_y>;
1262 def INT_PTX_SREG_NCTAID_Z : F_SREG<"mov.u32 \t$dst, %nctaid.z;", Int32Regs,
1263 int_nvvm_read_ptx_sreg_nctaid_z>;
1265 def INT_PTX_SREG_WARPSIZE : F_SREG<"mov.u32 \t$dst, WARP_SZ;", Int32Regs,
1266 int_nvvm_read_ptx_sreg_warpsize>;
1269 //-----------------------------------
1270 // Support for ldu on sm_20 or later
1271 //-----------------------------------
1274 // @TODO: Revisit this, Changed imemAny to imem
1275 multiclass LDU_G<string TyStr, NVPTXRegClass regclass, Intrinsic IntOp> {
1276 def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
1277 !strconcat("ldu.global.", TyStr),
1278 [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDU]>;
1279 def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
1280 !strconcat("ldu.global.", TyStr),
1281 [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDU]>;
1282 def avar: NVPTXInst<(outs regclass:$result), (ins imem:$src),
1283 !strconcat("ldu.global.", TyStr),
1284 [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
1286 def ari : NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
1287 !strconcat("ldu.global.", TyStr),
1288 [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDU]>;
1289 def ari64 : NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
1290 !strconcat("ldu.global.", TyStr),
1291 [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDU]>;
1294 defm INT_PTX_LDU_GLOBAL_i8 : LDU_G<"u8 \t$result, [$src];", Int8Regs,
1295 int_nvvm_ldu_global_i>;
1296 defm INT_PTX_LDU_GLOBAL_i16 : LDU_G<"u16 \t$result, [$src];", Int16Regs,
1297 int_nvvm_ldu_global_i>;
1298 defm INT_PTX_LDU_GLOBAL_i32 : LDU_G<"u32 \t$result, [$src];", Int32Regs,
1299 int_nvvm_ldu_global_i>;
1300 defm INT_PTX_LDU_GLOBAL_i64 : LDU_G<"u64 \t$result, [$src];", Int64Regs,
1301 int_nvvm_ldu_global_i>;
1302 defm INT_PTX_LDU_GLOBAL_f32 : LDU_G<"f32 \t$result, [$src];", Float32Regs,
1303 int_nvvm_ldu_global_f>;
1304 defm INT_PTX_LDU_GLOBAL_f64 : LDU_G<"f64 \t$result, [$src];", Float64Regs,
1305 int_nvvm_ldu_global_f>;
1306 defm INT_PTX_LDU_GLOBAL_p32 : LDU_G<"u32 \t$result, [$src];", Int32Regs,
1307 int_nvvm_ldu_global_p>;
1308 defm INT_PTX_LDU_GLOBAL_p64 : LDU_G<"u64 \t$result, [$src];", Int64Regs,
1309 int_nvvm_ldu_global_p>;
1313 // Elementized vector ldu
1314 multiclass VLDU_G_ELE_V2<string TyStr, NVPTXRegClass regclass> {
1315 def _32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1316 (ins Int32Regs:$src),
1317 !strconcat("ldu.global.", TyStr), []>;
1318 def _64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1319 (ins Int64Regs:$src),
1320 !strconcat("ldu.global.", TyStr), []>;
1323 multiclass VLDU_G_ELE_V4<string TyStr, NVPTXRegClass regclass> {
1324 def _32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1325 regclass:$dst4), (ins Int32Regs:$src),
1326 !strconcat("ldu.global.", TyStr), []>;
1327 def _64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1328 regclass:$dst4), (ins Int64Regs:$src),
1329 !strconcat("ldu.global.", TyStr), []>;
1332 defm INT_PTX_LDU_G_v2i8_ELE
1333 : VLDU_G_ELE_V2<"v2.u8 \t{{$dst1, $dst2}}, [$src];", Int8Regs>;
1334 defm INT_PTX_LDU_G_v2i16_ELE
1335 : VLDU_G_ELE_V2<"v2.u16 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
1336 defm INT_PTX_LDU_G_v2i32_ELE
1337 : VLDU_G_ELE_V2<"v2.u32 \t{{$dst1, $dst2}}, [$src];", Int32Regs>;
1338 defm INT_PTX_LDU_G_v2f32_ELE
1339 : VLDU_G_ELE_V2<"v2.f32 \t{{$dst1, $dst2}}, [$src];", Float32Regs>;
1340 defm INT_PTX_LDU_G_v2i64_ELE
1341 : VLDU_G_ELE_V2<"v2.u64 \t{{$dst1, $dst2}}, [$src];", Int64Regs>;
1342 defm INT_PTX_LDU_G_v2f64_ELE
1343 : VLDU_G_ELE_V2<"v2.f64 \t{{$dst1, $dst2}}, [$src];", Float64Regs>;
1344 defm INT_PTX_LDU_G_v4i8_ELE
1345 : VLDU_G_ELE_V4<"v4.u8 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int8Regs>;
1346 defm INT_PTX_LDU_G_v4i16_ELE
1347 : VLDU_G_ELE_V4<"v4.u16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
1349 defm INT_PTX_LDU_G_v4i32_ELE
1350 : VLDU_G_ELE_V4<"v4.u32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
1352 defm INT_PTX_LDU_G_v4f32_ELE
1353 : VLDU_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
1357 //-----------------------------------
1358 // Support for ldg on sm_35 or later
1359 //-----------------------------------
1361 def ldg_i8 : PatFrag<(ops node:$ptr), (int_nvvm_ldg_global_i node:$ptr), [{
1362 MemIntrinsicSDNode *M = cast<MemIntrinsicSDNode>(N);
1363 return M->getMemoryVT() == MVT::i8;
1366 multiclass LDG_G<string TyStr, NVPTXRegClass regclass, Intrinsic IntOp> {
1367 def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
1368 !strconcat("ld.global.nc.", TyStr),
1369 [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDG]>;
1370 def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
1371 !strconcat("ld.global.nc.", TyStr),
1372 [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDG]>;
1373 def avar: NVPTXInst<(outs regclass:$result), (ins imem:$src),
1374 !strconcat("ld.global.nc.", TyStr),
1375 [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
1377 def ari : NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
1378 !strconcat("ld.global.nc.", TyStr),
1379 [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDG]>;
1380 def ari64 : NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
1381 !strconcat("ld.global.nc.", TyStr),
1382 [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDG]>;
1385 multiclass LDG_G_NOINTRIN<string TyStr, NVPTXRegClass regclass, PatFrag IntOp> {
1386 def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
1387 !strconcat("ld.global.nc.", TyStr),
1388 [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDG]>;
1389 def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
1390 !strconcat("ld.global.nc.", TyStr),
1391 [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDG]>;
1392 def avar: NVPTXInst<(outs regclass:$result), (ins imem:$src),
1393 !strconcat("ld.global.nc.", TyStr),
1394 [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
1396 def ari : NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
1397 !strconcat("ld.global.nc.", TyStr),
1398 [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDG]>;
1399 def ari64 : NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
1400 !strconcat("ld.global.nc.", TyStr),
1401 [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDG]>;
1404 defm INT_PTX_LDG_GLOBAL_i8
1405 : LDG_G_NOINTRIN<"u8 \t$result, [$src];", Int16Regs, ldg_i8>;
1406 defm INT_PTX_LDG_GLOBAL_i16
1407 : LDG_G<"u16 \t$result, [$src];", Int16Regs, int_nvvm_ldg_global_i>;
1408 defm INT_PTX_LDG_GLOBAL_i32
1409 : LDG_G<"u32 \t$result, [$src];", Int32Regs, int_nvvm_ldg_global_i>;
1410 defm INT_PTX_LDG_GLOBAL_i64
1411 : LDG_G<"u64 \t$result, [$src];", Int64Regs, int_nvvm_ldg_global_i>;
1412 defm INT_PTX_LDG_GLOBAL_f32
1413 : LDG_G<"f32 \t$result, [$src];", Float32Regs, int_nvvm_ldg_global_f>;
1414 defm INT_PTX_LDG_GLOBAL_f64
1415 : LDG_G<"f64 \t$result, [$src];", Float64Regs, int_nvvm_ldg_global_f>;
1416 defm INT_PTX_LDG_GLOBAL_p32
1417 : LDG_G<"u32 \t$result, [$src];", Int32Regs, int_nvvm_ldg_global_p>;
1418 defm INT_PTX_LDG_GLOBAL_p64
1419 : LDG_G<"u64 \t$result, [$src];", Int64Regs, int_nvvm_ldg_global_p>;
1423 // Elementized vector ldg
1424 multiclass VLDG_G_ELE_V2<string TyStr, NVPTXRegClass regclass> {
1425 def _32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1426 (ins Int32Regs:$src),
1427 !strconcat("ld.global.nc.", TyStr), []>;
1428 def _64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1429 (ins Int64Regs:$src),
1430 !strconcat("ld.global.nc.", TyStr), []>;
1433 multiclass VLDG_G_ELE_V4<string TyStr, NVPTXRegClass regclass> {
1434 def _32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2,
1435 regclass:$dst3, regclass:$dst4), (ins Int32Regs:$src),
1436 !strconcat("ld.global.nc.", TyStr), []>;
1437 def _64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2,
1438 regclass:$dst3, regclass:$dst4), (ins Int64Regs:$src),
1439 !strconcat("ld.global.nc.", TyStr), []>;
1442 // FIXME: 8-bit LDG should be fixed once LDG/LDU nodes are made into proper loads.
1443 defm INT_PTX_LDG_G_v2i8_ELE
1444 : VLDG_G_ELE_V2<"v2.u8 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
1445 defm INT_PTX_LDG_G_v2i16_ELE
1446 : VLDG_G_ELE_V2<"v2.u16 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
1447 defm INT_PTX_LDG_G_v2i32_ELE
1448 : VLDG_G_ELE_V2<"v2.u32 \t{{$dst1, $dst2}}, [$src];", Int32Regs>;
1449 defm INT_PTX_LDG_G_v2f32_ELE
1450 : VLDG_G_ELE_V2<"v2.f32 \t{{$dst1, $dst2}}, [$src];", Float32Regs>;
1451 defm INT_PTX_LDG_G_v2i64_ELE
1452 : VLDG_G_ELE_V2<"v2.u64 \t{{$dst1, $dst2}}, [$src];", Int64Regs>;
1453 defm INT_PTX_LDG_G_v2f64_ELE
1454 : VLDG_G_ELE_V2<"v2.f64 \t{{$dst1, $dst2}}, [$src];", Float64Regs>;
1455 defm INT_PTX_LDG_G_v4i8_ELE
1456 : VLDG_G_ELE_V4<"v4.u8 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
1457 defm INT_PTX_LDG_G_v4i16_ELE
1458 : VLDG_G_ELE_V4<"v4.u16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
1459 defm INT_PTX_LDG_G_v4i32_ELE
1460 : VLDG_G_ELE_V4<"v4.u32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int32Regs>;
1461 defm INT_PTX_LDG_G_v4f32_ELE
1462 : VLDG_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Float32Regs>;
1465 multiclass NG_TO_G<string Str, Intrinsic Intrin> {
1466 def _yes : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1467 !strconcat("cvta.", !strconcat(Str, ".u32 \t$result, $src;")),
1468 [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>,
1469 Requires<[hasGenericLdSt]>;
1470 def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1471 !strconcat("cvta.", !strconcat(Str, ".u64 \t$result, $src;")),
1472 [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>,
1473 Requires<[hasGenericLdSt]>;
1475 // @TODO: Are these actually needed? I believe global addresses will be copied
1476 // to register values anyway.
1477 /*def __addr_yes : NVPTXInst<(outs Int32Regs:$result), (ins imemAny:$src),
1478 !strconcat("cvta.", !strconcat(Str, ".u32 \t$result, $src;")),
1479 [(set Int32Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>,
1480 Requires<[hasGenericLdSt]>;
1481 def __addr_yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins imemAny:$src),
1482 !strconcat("cvta.", !strconcat(Str, ".u64 \t$result, $src;")),
1483 [(set Int64Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>,
1484 Requires<[hasGenericLdSt]>;*/
1486 def _no : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1487 "mov.u32 \t$result, $src;",
1488 [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>;
1489 def _no_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1490 "mov.u64 \t$result, $src;",
1491 [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>;
1493 // @TODO: Are these actually needed? I believe global addresses will be copied
1494 // to register values anyway.
1495 /*def _addr_no : NVPTXInst<(outs Int32Regs:$result), (ins imem:$src),
1496 "mov.u32 \t$result, $src;",
1497 [(set Int32Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>;
1498 def _addr_no_64 : NVPTXInst<(outs Int64Regs:$result), (ins imem:$src),
1499 "mov.u64 \t$result, $src;",
1500 [(set Int64Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>;*/
1503 multiclass G_TO_NG<string Str, Intrinsic Intrin> {
1504 def _yes : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1505 !strconcat("cvta.to.", !strconcat(Str, ".u32 \t$result, $src;")),
1506 [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>,
1507 Requires<[hasGenericLdSt]>;
1508 def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1509 !strconcat("cvta.to.", !strconcat(Str, ".u64 \t$result, $src;")),
1510 [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>,
1511 Requires<[hasGenericLdSt]>;
1512 def _no : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1513 "mov.u32 \t$result, $src;",
1514 [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>;
1515 def _no_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1516 "mov.u64 \t$result, $src;",
1517 [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>;
1520 defm cvta_local : NG_TO_G<"local", int_nvvm_ptr_local_to_gen>;
1521 defm cvta_shared : NG_TO_G<"shared", int_nvvm_ptr_shared_to_gen>;
1522 defm cvta_global : NG_TO_G<"global", int_nvvm_ptr_global_to_gen>;
1523 defm cvta_const : NG_TO_G<"const", int_nvvm_ptr_constant_to_gen>;
1525 defm cvta_to_local : G_TO_NG<"local", int_nvvm_ptr_gen_to_local>;
1526 defm cvta_to_shared : G_TO_NG<"shared", int_nvvm_ptr_gen_to_shared>;
1527 defm cvta_to_global : G_TO_NG<"global", int_nvvm_ptr_gen_to_global>;
1528 defm cvta_to_const : G_TO_NG<"const", int_nvvm_ptr_gen_to_constant>;
1531 // nvvm.ptr.gen.to.param
1532 def nvvm_ptr_gen_to_param : NVPTXInst<(outs Int32Regs:$result),
1533 (ins Int32Regs:$src),
1534 "mov.u32 \t$result, $src;",
1535 [(set Int32Regs:$result,
1536 (int_nvvm_ptr_gen_to_param Int32Regs:$src))]>;
1537 def nvvm_ptr_gen_to_param_64 : NVPTXInst<(outs Int64Regs:$result),
1538 (ins Int64Regs:$src),
1539 "mov.u64 \t$result, $src;",
1540 [(set Int64Regs:$result,
1541 (int_nvvm_ptr_gen_to_param Int64Regs:$src))]>;
1544 // nvvm.move intrinsicc
1545 def nvvm_move_i8 : NVPTXInst<(outs Int8Regs:$r), (ins Int8Regs:$s),
1546 "mov.b16 \t$r, $s;",
1548 (int_nvvm_move_i8 Int8Regs:$s))]>;
1549 def nvvm_move_i16 : NVPTXInst<(outs Int16Regs:$r), (ins Int16Regs:$s),
1550 "mov.b16 \t$r, $s;",
1552 (int_nvvm_move_i16 Int16Regs:$s))]>;
1553 def nvvm_move_i32 : NVPTXInst<(outs Int32Regs:$r), (ins Int32Regs:$s),
1554 "mov.b32 \t$r, $s;",
1556 (int_nvvm_move_i32 Int32Regs:$s))]>;
1557 def nvvm_move_i64 : NVPTXInst<(outs Int64Regs:$r), (ins Int64Regs:$s),
1558 "mov.b64 \t$r, $s;",
1560 (int_nvvm_move_i64 Int64Regs:$s))]>;
1561 def nvvm_move_float : NVPTXInst<(outs Float32Regs:$r), (ins Float32Regs:$s),
1562 "mov.f32 \t$r, $s;",
1563 [(set Float32Regs:$r,
1564 (int_nvvm_move_float Float32Regs:$s))]>;
1565 def nvvm_move_double : NVPTXInst<(outs Float64Regs:$r), (ins Float64Regs:$s),
1566 "mov.f64 \t$r, $s;",
1567 [(set Float64Regs:$r,
1568 (int_nvvm_move_double Float64Regs:$s))]>;
1569 def nvvm_move_ptr32 : NVPTXInst<(outs Int32Regs:$r), (ins Int32Regs:$s),
1570 "mov.u32 \t$r, $s;",
1572 (int_nvvm_move_ptr Int32Regs:$s))]>;
1573 def nvvm_move_ptr64 : NVPTXInst<(outs Int64Regs:$r), (ins Int64Regs:$s),
1574 "mov.u64 \t$r, $s;",
1576 (int_nvvm_move_ptr Int64Regs:$s))]>;
1578 // @TODO: Are these actually needed, or will we always just see symbols
1579 // copied to registers first?
1580 /*def nvvm_move_sym32 : NVPTXInst<(outs Int32Regs:$r), (ins imem:$s),
1581 "mov.u32 \t$r, $s;",
1583 (int_nvvm_move_ptr texternalsym:$s))]>;
1584 def nvvm_move_sym64 : NVPTXInst<(outs Int64Regs:$r), (ins imem:$s),
1585 "mov.u64 \t$r, $s;",
1587 (int_nvvm_move_ptr texternalsym:$s))]>;*/
1590 // MoveParam %r1, param
1591 // ptr_local_to_gen %r2, %r1
1592 // ptr_gen_to_local %r3, %r2
1596 // @TODO: Revisit this. There is a type
1597 // contradiction between iPTRAny and iPTR for the addr defs, so the move_sym
1598 // instructions are not currently defined. However, we can use the ptr
1599 // variants and the asm printer will do the right thing.
1600 def : Pat<(i64 (int_nvvm_ptr_gen_to_local (int_nvvm_ptr_local_to_gen
1601 (MoveParam texternalsym:$src)))),
1602 (nvvm_move_ptr64 texternalsym:$src)>;
1603 def : Pat<(i32 (int_nvvm_ptr_gen_to_local (int_nvvm_ptr_local_to_gen
1604 (MoveParam texternalsym:$src)))),
1605 (nvvm_move_ptr32 texternalsym:$src)>;
1608 //-----------------------------------
1609 // Compiler Error Warn
1610 // - Just ignore them in codegen
1611 //-----------------------------------
1613 def INT_NVVM_COMPILER_WARN_32 : NVPTXInst<(outs), (ins Int32Regs:$a),
1614 "// llvm.nvvm.compiler.warn()",
1615 [(int_nvvm_compiler_warn Int32Regs:$a)]>;
1616 def INT_NVVM_COMPILER_WARN_64 : NVPTXInst<(outs), (ins Int64Regs:$a),
1617 "// llvm.nvvm.compiler.warn()",
1618 [(int_nvvm_compiler_warn Int64Regs:$a)]>;
1619 def INT_NVVM_COMPILER_ERROR_32 : NVPTXInst<(outs), (ins Int32Regs:$a),
1620 "// llvm.nvvm.compiler.error()",
1621 [(int_nvvm_compiler_error Int32Regs:$a)]>;
1622 def INT_NVVM_COMPILER_ERROR_64 : NVPTXInst<(outs), (ins Int64Regs:$a),
1623 "// llvm.nvvm.compiler.error()",
1624 [(int_nvvm_compiler_error Int64Regs:$a)]>;
1628 //===-- Old PTX Back-end Intrinsics ---------------------------------------===//
1630 // These intrinsics are handled to retain compatibility with the old backend.
1632 // PTX Special Purpose Register Accessor Intrinsics
1634 class PTX_READ_SPECIAL_REGISTER_R64<string regname, Intrinsic intop>
1635 : NVPTXInst<(outs Int64Regs:$d), (ins),
1636 !strconcat(!strconcat("mov.u64\t$d, %", regname), ";"),
1637 [(set Int64Regs:$d, (intop))]>;
1639 class PTX_READ_SPECIAL_REGISTER_R32<string regname, Intrinsic intop>
1640 : NVPTXInst<(outs Int32Regs:$d), (ins),
1641 !strconcat(!strconcat("mov.u32\t$d, %", regname), ";"),
1642 [(set Int32Regs:$d, (intop))]>;
1644 // TODO Add read vector-version of special registers
1646 def PTX_READ_TID_X : PTX_READ_SPECIAL_REGISTER_R32<"tid.x",
1647 int_ptx_read_tid_x>;
1648 def PTX_READ_TID_Y : PTX_READ_SPECIAL_REGISTER_R32<"tid.y",
1649 int_ptx_read_tid_y>;
1650 def PTX_READ_TID_Z : PTX_READ_SPECIAL_REGISTER_R32<"tid.z",
1651 int_ptx_read_tid_z>;
1652 def PTX_READ_TID_W : PTX_READ_SPECIAL_REGISTER_R32<"tid.w",
1653 int_ptx_read_tid_w>;
1655 def PTX_READ_NTID_X : PTX_READ_SPECIAL_REGISTER_R32<"ntid.x",
1656 int_ptx_read_ntid_x>;
1657 def PTX_READ_NTID_Y : PTX_READ_SPECIAL_REGISTER_R32<"ntid.y",
1658 int_ptx_read_ntid_y>;
1659 def PTX_READ_NTID_Z : PTX_READ_SPECIAL_REGISTER_R32<"ntid.z",
1660 int_ptx_read_ntid_z>;
1661 def PTX_READ_NTID_W : PTX_READ_SPECIAL_REGISTER_R32<"ntid.w",
1662 int_ptx_read_ntid_w>;
1664 def PTX_READ_LANEID : PTX_READ_SPECIAL_REGISTER_R32<"laneid",
1665 int_ptx_read_laneid>;
1666 def PTX_READ_WARPID : PTX_READ_SPECIAL_REGISTER_R32<"warpid",
1667 int_ptx_read_warpid>;
1668 def PTX_READ_NWARPID : PTX_READ_SPECIAL_REGISTER_R32<"nwarpid",
1669 int_ptx_read_nwarpid>;
1671 def PTX_READ_CTAID_X : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.x",
1672 int_ptx_read_ctaid_x>;
1673 def PTX_READ_CTAID_Y : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.y",
1674 int_ptx_read_ctaid_y>;
1675 def PTX_READ_CTAID_Z : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.z",
1676 int_ptx_read_ctaid_z>;
1677 def PTX_READ_CTAID_W : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.w",
1678 int_ptx_read_ctaid_w>;
1680 def PTX_READ_NCTAID_X : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.x",
1681 int_ptx_read_nctaid_x>;
1682 def PTX_READ_NCTAID_Y : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.y",
1683 int_ptx_read_nctaid_y>;
1684 def PTX_READ_NCTAID_Z : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.z",
1685 int_ptx_read_nctaid_z>;
1686 def PTX_READ_NCTAID_W : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.w",
1687 int_ptx_read_nctaid_w>;
1689 def PTX_READ_SMID : PTX_READ_SPECIAL_REGISTER_R32<"smid",
1691 def PTX_READ_NSMID : PTX_READ_SPECIAL_REGISTER_R32<"nsmid",
1692 int_ptx_read_nsmid>;
1693 def PTX_READ_GRIDID : PTX_READ_SPECIAL_REGISTER_R32<"gridid",
1694 int_ptx_read_gridid>;
1696 def PTX_READ_LANEMASK_EQ
1697 : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_eq", int_ptx_read_lanemask_eq>;
1698 def PTX_READ_LANEMASK_LE
1699 : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_le", int_ptx_read_lanemask_le>;
1700 def PTX_READ_LANEMASK_LT
1701 : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_lt", int_ptx_read_lanemask_lt>;
1702 def PTX_READ_LANEMASK_GE
1703 : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_ge", int_ptx_read_lanemask_ge>;
1704 def PTX_READ_LANEMASK_GT
1705 : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_gt", int_ptx_read_lanemask_gt>;
1708 : PTX_READ_SPECIAL_REGISTER_R32<"clock", int_ptx_read_clock>;
1709 def PTX_READ_CLOCK64
1710 : PTX_READ_SPECIAL_REGISTER_R64<"clock64", int_ptx_read_clock64>;
1712 def PTX_READ_PM0 : PTX_READ_SPECIAL_REGISTER_R32<"pm0", int_ptx_read_pm0>;
1713 def PTX_READ_PM1 : PTX_READ_SPECIAL_REGISTER_R32<"pm1", int_ptx_read_pm1>;
1714 def PTX_READ_PM2 : PTX_READ_SPECIAL_REGISTER_R32<"pm2", int_ptx_read_pm2>;
1715 def PTX_READ_PM3 : PTX_READ_SPECIAL_REGISTER_R32<"pm3", int_ptx_read_pm3>;
1717 // PTX Parallel Synchronization and Communication Intrinsics
1719 def PTX_BAR_SYNC : NVPTXInst<(outs), (ins i32imm:$i), "bar.sync\t$i;",
1720 [(int_ptx_bar_sync imm:$i)]>;