Gitiles
Code Review Sign In
gerrit-public.fairphone.software / toolchain / llvm-project / f901b719214d0bd4d42d85e9b5bbbc95d1ee3479 / . / llvm / lib / Target / NVPTX / NVPTXIntrinsics.td
blob: 49e2568dfa2c50406c4a4cc579bc39f7698b3bf0 [file] [log] [blame]
Justin Holewinskiae556d32012-05-04 20:18:50 +0000[diff] [blame]1//===- NVPTXIntrinsics.td - PTX Intrinsics Instructions -------*- tblgen -*-==//
2//
3// The LLVM Compiler Infrastructure
4//
5// This file is distributed under the University of Illinois Open Source
6// License. See LICENSE.TXT for details.
7//
8//===----------------------------------------------------------------------===//
9
10def immFloat0 : PatLeaf<(fpimm), [{
11 float f = (float)N->getValueAPF().convertToFloat();
12 return (f==0.0f);
13}]>;
14
15def immFloat1 : PatLeaf<(fpimm), [{
16 float f = (float)N->getValueAPF().convertToFloat();
17 return (f==1.0f);
18}]>;
19
20def immDouble0 : PatLeaf<(fpimm), [{
21 double d = (double)N->getValueAPF().convertToDouble();
22 return (d==0.0);
23}]>;
24
25def immDouble1 : PatLeaf<(fpimm), [{
26 double d = (double)N->getValueAPF().convertToDouble();
27 return (d==1.0);
28}]>;
29
30
31
32//-----------------------------------
33// Synchronization Functions
34//-----------------------------------
35def INT_CUDA_SYNCTHREADS : NVPTXInst<(outs), (ins),
36 "bar.sync \t0;",
37 [(int_cuda_syncthreads)]>;
38def INT_BARRIER0 : NVPTXInst<(outs), (ins),
39 "bar.sync \t0;",
40 [(int_nvvm_barrier0)]>;
41def INT_BARRIER0_POPC : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred),
42 !strconcat("{{ \n\t",
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))]>;
48def INT_BARRIER0_AND : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred),
49 !strconcat("{{ \n\t",
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))]>;
57def INT_BARRIER0_OR : NVPTXInst<(outs Int32Regs:$dst), (ins Int32Regs:$pred),
58 !strconcat("{{ \n\t",
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))]>;
66
67
68//-----------------------------------
69// Explicit Memory Fence Functions
70//-----------------------------------
71class MEMBAR<string StrOp, Intrinsic IntOP> :
72 NVPTXInst<(outs), (ins),
73 StrOp, [(IntOP)]>;
74
75def INT_MEMBAR_CTA : MEMBAR<"membar.cta;", int_nvvm_membar_cta>;
76def INT_MEMBAR_GL : MEMBAR<"membar.gl;", int_nvvm_membar_gl>;
77def INT_MEMBAR_SYS : MEMBAR<"membar.sys;", int_nvvm_membar_sys>;
78
79
80//-----------------------------------
81// Math Functions
82//-----------------------------------
83
84// Map min(1.0, max(0.0, x)) to sat(x)
85multiclass SAT<NVPTXRegClass regclass, Operand fimm, Intrinsic IntMinOp,
86 Intrinsic IntMaxOp, PatLeaf f0, PatLeaf f1, string OpStr> {
87
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),
91 OpStr,
92 [(set regclass:$dst, (IntMinOp f1:$srcf0 ,
93 (IntMaxOp f0:$srcf1, regclass:$src)))]>;
94
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),
98 OpStr,
99 [(set regclass:$dst, (IntMinOp f1:$srcf0 ,
100 (IntMaxOp regclass:$src, f0:$srcf1)))]>;
101
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),
105 OpStr,
106 [(set regclass:$dst, (IntMinOp
107 (IntMaxOp f0:$srcf0, regclass:$src), f1:$srcf1))]>;
108
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),
112 OpStr,
113 [(set regclass:$dst, (IntMinOp
114 (IntMaxOp regclass:$src, f0:$srcf0), f1:$srcf1))]>;
115
116}
117// Note that max(0.0, min(x, 1.0)) cannot be mapped to sat(x) because when x
118// is NaN
119// max(0.0, min(x, 1.0)) is 1.0 while sat(x) is 0.
120// Same story for fmax, fmin.
121
122defm 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">;
125defm 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">;
128
129
130// We need a full string for OpcStr here because we need to deal with case like
131// INT_PTX_RECIP.
132class F_MATH_1<string OpcStr, NVPTXRegClass target_regclass,
133 NVPTXRegClass src_regclass, Intrinsic IntOP>
134 : NVPTXInst<(outs target_regclass:$dst), (ins src_regclass:$src0),
135 OpcStr,
136 [(set target_regclass:$dst, (IntOP src_regclass:$src0))]>;
137
138// We need a full string for OpcStr here because we need to deal with the case
139// like INT_PTX_NATIVE_POWR_F.
140class 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),
144 OpcStr,
145 [(set t_regclass:$dst, (IntOP s0_regclass:$src0, s1_regclass:$src1))]>;
146
147class 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),
152 OpcStr,
153 [(set t_regclass:$dst,
154 (IntOP s0_regclass:$src0, s1_regclass:$src1, s2_regclass:$src2))]>;
155
156//
157// MISC
158//
159
160def INT_NVVM_CLZ_I : F_MATH_1<"clz.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
161 int_nvvm_clz_i>;
162def INT_NVVM_CLZ_LL : F_MATH_1<"clz.b64 \t$dst, $src0;", Int32Regs, Int64Regs,
163 int_nvvm_clz_ll>;
164
165def INT_NVVM_POPC_I : F_MATH_1<"popc.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
166 int_nvvm_popc_i>;
167def INT_NVVM_POPC_LL : F_MATH_1<"popc.b64 \t$dst, $src0;", Int32Regs, Int64Regs,
168 int_nvvm_popc_ll>;
169
170def INT_NVVM_PRMT : F_MATH_3<"prmt.b32 \t$dst, $src0, $src1, $src2;", Int32Regs,
171 Int32Regs, Int32Regs, Int32Regs, int_nvvm_prmt>;
172
173//
174// Min Max
175//
176
177def INT_NVVM_MIN_I : F_MATH_2<"min.s32 \t$dst, $src0, $src1;", Int32Regs,
178 Int32Regs, Int32Regs, int_nvvm_min_i>;
179def INT_NVVM_MIN_UI : F_MATH_2<"min.u32 \t$dst, $src0, $src1;", Int32Regs,
180 Int32Regs, Int32Regs, int_nvvm_min_ui>;
181
182def INT_NVVM_MIN_LL : F_MATH_2<"min.s64 \t$dst, $src0, $src1;", Int64Regs,
183 Int64Regs, Int64Regs, int_nvvm_min_ll>;
184def INT_NVVM_MIN_ULL : F_MATH_2<"min.u64 \t$dst, $src0, $src1;", Int64Regs,
185 Int64Regs, Int64Regs, int_nvvm_min_ull>;
186
187def INT_NVVM_MAX_I : F_MATH_2<"max.s32 \t$dst, $src0, $src1;", Int32Regs,
188 Int32Regs, Int32Regs, int_nvvm_max_i>;
189def INT_NVVM_MAX_UI : F_MATH_2<"max.u32 \t$dst, $src0, $src1;", Int32Regs,
190 Int32Regs, Int32Regs, int_nvvm_max_ui>;
191
192def INT_NVVM_MAX_LL : F_MATH_2<"max.s64 \t$dst, $src0, $src1;", Int64Regs,
193 Int64Regs, Int64Regs, int_nvvm_max_ll>;
194def INT_NVVM_MAX_ULL : F_MATH_2<"max.u64 \t$dst, $src0, $src1;", Int64Regs,
195 Int64Regs, Int64Regs, int_nvvm_max_ull>;
196
197def INT_NVVM_FMIN_F : F_MATH_2<"min.f32 \t$dst, $src0, $src1;", Float32Regs,
198 Float32Regs, Float32Regs, int_nvvm_fmin_f>;
199def INT_NVVM_FMIN_FTZ_F : F_MATH_2<"min.ftz.f32 \t$dst, $src0, $src1;",
200 Float32Regs, Float32Regs, Float32Regs, int_nvvm_fmin_ftz_f>;
201
202def INT_NVVM_FMAX_F : F_MATH_2<"max.f32 \t$dst, $src0, $src1;", Float32Regs,
203 Float32Regs, Float32Regs, int_nvvm_fmax_f>;
204def INT_NVVM_FMAX_FTZ_F : F_MATH_2<"max.ftz.f32 \t$dst, $src0, $src1;",
205 Float32Regs, Float32Regs, Float32Regs, int_nvvm_fmax_ftz_f>;
206
207def INT_NVVM_FMIN_D : F_MATH_2<"min.f64 \t$dst, $src0, $src1;", Float64Regs,
208 Float64Regs, Float64Regs, int_nvvm_fmin_d>;
209def INT_NVVM_FMAX_D : F_MATH_2<"max.f64 \t$dst, $src0, $src1;", Float64Regs,
210 Float64Regs, Float64Regs, int_nvvm_fmax_d>;
211
212//
213// Multiplication
214//
215
216def INT_NVVM_MULHI_I : F_MATH_2<"mul.hi.s32 \t$dst, $src0, $src1;", Int32Regs,
217 Int32Regs, Int32Regs, int_nvvm_mulhi_i>;
218def INT_NVVM_MULHI_UI : F_MATH_2<"mul.hi.u32 \t$dst, $src0, $src1;", Int32Regs,
219 Int32Regs, Int32Regs, int_nvvm_mulhi_ui>;
220
221def INT_NVVM_MULHI_LL : F_MATH_2<"mul.hi.s64 \t$dst, $src0, $src1;", Int64Regs,
222 Int64Regs, Int64Regs, int_nvvm_mulhi_ll>;
223def INT_NVVM_MULHI_ULL : F_MATH_2<"mul.hi.u64 \t$dst, $src0, $src1;", Int64Regs,
224 Int64Regs, Int64Regs, int_nvvm_mulhi_ull>;
225
226def 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>;
228def INT_NVVM_MUL_RN_F : F_MATH_2<"mul.rn.f32 \t$dst, $src0, $src1;",
229 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rn_f>;
230def 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>;
232def INT_NVVM_MUL_RZ_F : F_MATH_2<"mul.rz.f32 \t$dst, $src0, $src1;",
233 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rz_f>;
234def 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>;
236def INT_NVVM_MUL_RM_F : F_MATH_2<"mul.rm.f32 \t$dst, $src0, $src1;",
237 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rm_f>;
238def 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>;
240def INT_NVVM_MUL_RP_F : F_MATH_2<"mul.rp.f32 \t$dst, $src0, $src1;",
241 Float32Regs, Float32Regs, Float32Regs, int_nvvm_mul_rp_f>;
242
243def INT_NVVM_MUL_RN_D : F_MATH_2<"mul.rn.f64 \t$dst, $src0, $src1;",
244 Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rn_d>;
245def INT_NVVM_MUL_RZ_D : F_MATH_2<"mul.rz.f64 \t$dst, $src0, $src1;",
246 Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rz_d>;
247def INT_NVVM_MUL_RM_D : F_MATH_2<"mul.rm.f64 \t$dst, $src0, $src1;",
248 Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rm_d>;
249def INT_NVVM_MUL_RP_D : F_MATH_2<"mul.rp.f64 \t$dst, $src0, $src1;",
250 Float64Regs, Float64Regs, Float64Regs, int_nvvm_mul_rp_d>;
251
252def INT_NVVM_MUL24_I : F_MATH_2<"mul24.lo.s32 \t$dst, $src0, $src1;",
253 Int32Regs, Int32Regs, Int32Regs, int_nvvm_mul24_i>;
254def INT_NVVM_MUL24_UI : F_MATH_2<"mul24.lo.u32 \t$dst, $src0, $src1;",
255 Int32Regs, Int32Regs, Int32Regs, int_nvvm_mul24_ui>;
256
257//
258// Div
259//
260
261def 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>;
264def INT_NVVM_DIV_APPROX_F : F_MATH_2<"div.approx.f32 \t$dst, $src0, $src1;",
265 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_approx_f>;
266
267def 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>;
269def INT_NVVM_DIV_RN_F : F_MATH_2<"div.rn.f32 \t$dst, $src0, $src1;",
270 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rn_f>;
271def 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>;
273def INT_NVVM_DIV_RZ_F : F_MATH_2<"div.rz.f32 \t$dst, $src0, $src1;",
274 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rz_f>;
275def 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>;
277def INT_NVVM_DIV_RM_F : F_MATH_2<"div.rm.f32 \t$dst, $src0, $src1;",
278 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rm_f>;
279def 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>;
281def INT_NVVM_DIV_RP_F : F_MATH_2<"div.rp.f32 \t$dst, $src0, $src1;",
282 Float32Regs, Float32Regs, Float32Regs, int_nvvm_div_rp_f>;
283
284def INT_NVVM_DIV_RN_D : F_MATH_2<"div.rn.f64 \t$dst, $src0, $src1;",
285 Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rn_d>;
286def INT_NVVM_DIV_RZ_D : F_MATH_2<"div.rz.f64 \t$dst, $src0, $src1;",
287 Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rz_d>;
288def INT_NVVM_DIV_RM_D : F_MATH_2<"div.rm.f64 \t$dst, $src0, $src1;",
289 Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rm_d>;
290def INT_NVVM_DIV_RP_D : F_MATH_2<"div.rp.f64 \t$dst, $src0, $src1;",
291 Float64Regs, Float64Regs, Float64Regs, int_nvvm_div_rp_d>;
292
293//
294// Brev
295//
296
297def INT_NVVM_BREV32 : F_MATH_1<"brev.b32 \t$dst, $src0;", Int32Regs, Int32Regs,
298 int_nvvm_brev32>;
299def INT_NVVM_BREV64 : F_MATH_1<"brev.b64 \t$dst, $src0;", Int64Regs, Int64Regs,
300 int_nvvm_brev64>;
301
302//
303// Sad
304//
305
306def INT_NVVM_SAD_I : F_MATH_3<"sad.s32 \t$dst, $src0, $src1, $src2;",
307 Int32Regs, Int32Regs, Int32Regs, Int32Regs, int_nvvm_sad_i>;
308def INT_NVVM_SAD_UI : F_MATH_3<"sad.u32 \t$dst, $src0, $src1, $src2;",
309 Int32Regs, Int32Regs, Int32Regs, Int32Regs, int_nvvm_sad_ui>;
310
311//
312// Floor Ceil
313//
314
315def INT_NVVM_FLOOR_FTZ_F : F_MATH_1<"cvt.rmi.ftz.f32.f32 \t$dst, $src0;",
316 Float32Regs, Float32Regs, int_nvvm_floor_ftz_f>;
317def INT_NVVM_FLOOR_F : F_MATH_1<"cvt.rmi.f32.f32 \t$dst, $src0;",
318 Float32Regs, Float32Regs, int_nvvm_floor_f>;
319def INT_NVVM_FLOOR_D : F_MATH_1<"cvt.rmi.f64.f64 \t$dst, $src0;",
320 Float64Regs, Float64Regs, int_nvvm_floor_d>;
321
322def INT_NVVM_CEIL_FTZ_F : F_MATH_1<"cvt.rpi.ftz.f32.f32 \t$dst, $src0;",
323 Float32Regs, Float32Regs, int_nvvm_ceil_ftz_f>;
324def INT_NVVM_CEIL_F : F_MATH_1<"cvt.rpi.f32.f32 \t$dst, $src0;",
325 Float32Regs, Float32Regs, int_nvvm_ceil_f>;
326def INT_NVVM_CEIL_D : F_MATH_1<"cvt.rpi.f64.f64 \t$dst, $src0;",
327 Float64Regs, Float64Regs, int_nvvm_ceil_d>;
328
329//
330// Abs
331//
332
333def INT_NVVM_ABS_I : F_MATH_1<"abs.s32 \t$dst, $src0;", Int32Regs, Int32Regs,
334 int_nvvm_abs_i>;
335def INT_NVVM_ABS_LL : F_MATH_1<"abs.s64 \t$dst, $src0;", Int64Regs, Int64Regs,
336 int_nvvm_abs_ll>;
337
338def INT_NVVM_FABS_FTZ_F : F_MATH_1<"abs.ftz.f32 \t$dst, $src0;", Float32Regs,
339 Float32Regs, int_nvvm_fabs_ftz_f>;
340def INT_NVVM_FABS_F : F_MATH_1<"abs.f32 \t$dst, $src0;", Float32Regs,
341 Float32Regs, int_nvvm_fabs_f>;
342
343def INT_NVVM_FABS_D : F_MATH_1<"abs.f64 \t$dst, $src0;", Float64Regs,
344 Float64Regs, int_nvvm_fabs_d>;
345
346//
347// Round
348//
349
350def INT_NVVM_ROUND_FTZ_F : F_MATH_1<"cvt.rni.ftz.f32.f32 \t$dst, $src0;",
351 Float32Regs, Float32Regs, int_nvvm_round_ftz_f>;
352def INT_NVVM_ROUND_F : F_MATH_1<"cvt.rni.f32.f32 \t$dst, $src0;", Float32Regs,
353 Float32Regs, int_nvvm_round_f>;
354
355def INT_NVVM_ROUND_D : F_MATH_1<"cvt.rni.f64.f64 \t$dst, $src0;", Float64Regs,
356 Float64Regs, int_nvvm_round_d>;
357
358//
359// Trunc
360//
361
362def INT_NVVM_TRUNC_FTZ_F : F_MATH_1<"cvt.rzi.ftz.f32.f32 \t$dst, $src0;",
363 Float32Regs, Float32Regs, int_nvvm_trunc_ftz_f>;
364def INT_NVVM_TRUNC_F : F_MATH_1<"cvt.rzi.f32.f32 \t$dst, $src0;", Float32Regs,
365 Float32Regs, int_nvvm_trunc_f>;
366
367def INT_NVVM_TRUNC_D : F_MATH_1<"cvt.rzi.f64.f64 \t$dst, $src0;", Float64Regs,
368 Float64Regs, int_nvvm_trunc_d>;
369
370//
371// Saturate
372//
373
374def INT_NVVM_SATURATE_FTZ_F : F_MATH_1<"cvt.sat.ftz.f32.f32 \t$dst, $src0;",
375 Float32Regs, Float32Regs, int_nvvm_saturate_ftz_f>;
376def INT_NVVM_SATURATE_F : F_MATH_1<"cvt.sat.f32.f32 \t$dst, $src0;",
377 Float32Regs, Float32Regs, int_nvvm_saturate_f>;
378
379def INT_NVVM_SATURATE_D : F_MATH_1<"cvt.sat.f64.f64 \t$dst, $src0;",
380 Float64Regs, Float64Regs, int_nvvm_saturate_d>;
381
382//
383// Exp2 Log2
384//
385
386def 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>;
388def INT_NVVM_EX2_APPROX_F : F_MATH_1<"ex2.approx.f32 \t$dst, $src0;",
389 Float32Regs, Float32Regs, int_nvvm_ex2_approx_f>;
390def INT_NVVM_EX2_APPROX_D : F_MATH_1<"ex2.approx.f64 \t$dst, $src0;",
391 Float64Regs, Float64Regs, int_nvvm_ex2_approx_d>;
392
393def 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>;
395def INT_NVVM_LG2_APPROX_F : F_MATH_1<"lg2.approx.f32 \t$dst, $src0;",
396 Float32Regs, Float32Regs, int_nvvm_lg2_approx_f>;
397def INT_NVVM_LG2_APPROX_D : F_MATH_1<"lg2.approx.f64 \t$dst, $src0;",
398 Float64Regs, Float64Regs, int_nvvm_lg2_approx_d>;
399
400//
401// Sin Cos
402//
403
404def 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>;
406def INT_NVVM_SIN_APPROX_F : F_MATH_1<"sin.approx.f32 \t$dst, $src0;",
407 Float32Regs, Float32Regs, int_nvvm_sin_approx_f>;
408
409def 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>;
411def INT_NVVM_COS_APPROX_F : F_MATH_1<"cos.approx.f32 \t$dst, $src0;",
412 Float32Regs, Float32Regs, int_nvvm_cos_approx_f>;
413
414//
415// Fma
416//
417
418def 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>;
421def 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>;
423def 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>;
426def 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>;
428def 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>;
431def 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>;
433def 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>;
436def 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>;
438
439def 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>;
441def 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>;
443def 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>;
445def 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>;
447
448//
449// Rcp
450//
451
452def 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>;
454def INT_NVVM_RCP_RN_F : F_MATH_1<"rcp.rn.f32 \t$dst, $src0;",
455 Float32Regs, Float32Regs, int_nvvm_rcp_rn_f>;
456def 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>;
458def INT_NVVM_RCP_RZ_F : F_MATH_1<"rcp.rz.f32 \t$dst, $src0;",
459 Float32Regs, Float32Regs, int_nvvm_rcp_rz_f>;
460def 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>;
462def INT_NVVM_RCP_RM_F : F_MATH_1<"rcp.rm.f32 \t$dst, $src0;",
463 Float32Regs, Float32Regs, int_nvvm_rcp_rm_f>;
464def 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>;
466def INT_NVVM_RCP_RP_F : F_MATH_1<"rcp.rp.f32 \t$dst, $src0;",
467 Float32Regs, Float32Regs, int_nvvm_rcp_rp_f>;
468
469def INT_NVVM_RCP_RN_D : F_MATH_1<"rcp.rn.f64 \t$dst, $src0;", Float64Regs,
470 Float64Regs, int_nvvm_rcp_rn_d>;
471def INT_NVVM_RCP_RZ_D : F_MATH_1<"rcp.rz.f64 \t$dst, $src0;", Float64Regs,
472 Float64Regs, int_nvvm_rcp_rz_d>;
473def INT_NVVM_RCP_RM_D : F_MATH_1<"rcp.rm.f64 \t$dst, $src0;", Float64Regs,
474 Float64Regs, int_nvvm_rcp_rm_d>;
475def INT_NVVM_RCP_RP_D : F_MATH_1<"rcp.rp.f64 \t$dst, $src0;", Float64Regs,
476 Float64Regs, int_nvvm_rcp_rp_d>;
477
478def 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>;
480
481//
482// Sqrt
483//
484
485def 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>;
487def INT_NVVM_SQRT_RN_F : F_MATH_1<"sqrt.rn.f32 \t$dst, $src0;", Float32Regs,
488 Float32Regs, int_nvvm_sqrt_rn_f>;
489def 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>;
491def INT_NVVM_SQRT_RZ_F : F_MATH_1<"sqrt.rz.f32 \t$dst, $src0;", Float32Regs,
492 Float32Regs, int_nvvm_sqrt_rz_f>;
493def 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>;
495def INT_NVVM_SQRT_RM_F : F_MATH_1<"sqrt.rm.f32 \t$dst, $src0;", Float32Regs,
496 Float32Regs, int_nvvm_sqrt_rm_f>;
497def 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>;
499def INT_NVVM_SQRT_RP_F : F_MATH_1<"sqrt.rp.f32 \t$dst, $src0;", Float32Regs,
500 Float32Regs, int_nvvm_sqrt_rp_f>;
501def 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>;
503def INT_NVVM_SQRT_APPROX_F : F_MATH_1<"sqrt.approx.f32 \t$dst, $src0;",
504 Float32Regs, Float32Regs, int_nvvm_sqrt_approx_f>;
505
506def INT_NVVM_SQRT_RN_D : F_MATH_1<"sqrt.rn.f64 \t$dst, $src0;", Float64Regs,
507 Float64Regs, int_nvvm_sqrt_rn_d>;
508def INT_NVVM_SQRT_RZ_D : F_MATH_1<"sqrt.rz.f64 \t$dst, $src0;", Float64Regs,
509 Float64Regs, int_nvvm_sqrt_rz_d>;
510def INT_NVVM_SQRT_RM_D : F_MATH_1<"sqrt.rm.f64 \t$dst, $src0;", Float64Regs,
511 Float64Regs, int_nvvm_sqrt_rm_d>;
512def INT_NVVM_SQRT_RP_D : F_MATH_1<"sqrt.rp.f64 \t$dst, $src0;", Float64Regs,
513 Float64Regs, int_nvvm_sqrt_rp_d>;
514
515//
516// Rsqrt
517//
518
519def INT_NVVM_RSQRT_APPROX_FTZ_F
520 : F_MATH_1<"rsqrt.approx.ftz.f32 \t$dst, $src0;", Float32Regs, Float32Regs,
521 int_nvvm_rsqrt_approx_ftz_f>;
522def INT_NVVM_RSQRT_APPROX_F : F_MATH_1<"rsqrt.approx.f32 \t$dst, $src0;",
523 Float32Regs, Float32Regs, int_nvvm_rsqrt_approx_f>;
524def INT_NVVM_RSQRT_APPROX_D : F_MATH_1<"rsqrt.approx.f64 \t$dst, $src0;",
525 Float64Regs, Float64Regs, int_nvvm_rsqrt_approx_d>;
526
527//
528// Add
529//
530
531def INT_NVVM_ADD_RN_FTZ_F : F_MATH_2<"add.rn.ftz.f32 \t$dst, $src0, $src1;",
532 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rn_ftz_f>;
533def INT_NVVM_ADD_RN_F : F_MATH_2<"add.rn.f32 \t$dst, $src0, $src1;",
534 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rn_f>;
535def INT_NVVM_ADD_RZ_FTZ_F : F_MATH_2<"add.rz.ftz.f32 \t$dst, $src0, $src1;",
536 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rz_ftz_f>;
537def INT_NVVM_ADD_RZ_F : F_MATH_2<"add.rz.f32 \t$dst, $src0, $src1;",
538 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rz_f>;
539def INT_NVVM_ADD_RM_FTZ_F : F_MATH_2<"add.rm.ftz.f32 \t$dst, $src0, $src1;",
540 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rm_ftz_f>;
541def INT_NVVM_ADD_RM_F : F_MATH_2<"add.rm.f32 \t$dst, $src0, $src1;",
542 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rm_f>;
543def INT_NVVM_ADD_RP_FTZ_F : F_MATH_2<"add.rp.ftz.f32 \t$dst, $src0, $src1;",
544 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rp_ftz_f>;
545def INT_NVVM_ADD_RP_F : F_MATH_2<"add.rp.f32 \t$dst, $src0, $src1;",
546 Float32Regs, Float32Regs, Float32Regs, int_nvvm_add_rp_f>;
547
548def INT_NVVM_ADD_RN_D : F_MATH_2<"add.rn.f64 \t$dst, $src0, $src1;",
549 Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rn_d>;
550def INT_NVVM_ADD_RZ_D : F_MATH_2<"add.rz.f64 \t$dst, $src0, $src1;",
551 Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rz_d>;
552def INT_NVVM_ADD_RM_D : F_MATH_2<"add.rm.f64 \t$dst, $src0, $src1;",
553 Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rm_d>;
554def INT_NVVM_ADD_RP_D : F_MATH_2<"add.rp.f64 \t$dst, $src0, $src1;",
555 Float64Regs, Float64Regs, Float64Regs, int_nvvm_add_rp_d>;
556
557//
558// Convert
559//
560
561def INT_NVVM_D2F_RN_FTZ : F_MATH_1<"cvt.rn.ftz.f32.f64 \t$dst, $src0;",
562 Float32Regs, Float64Regs, int_nvvm_d2f_rn_ftz>;
563def INT_NVVM_D2F_RN : F_MATH_1<"cvt.rn.f32.f64 \t$dst, $src0;",
564 Float32Regs, Float64Regs, int_nvvm_d2f_rn>;
565def INT_NVVM_D2F_RZ_FTZ : F_MATH_1<"cvt.rz.ftz.f32.f64 \t$dst, $src0;",
566 Float32Regs, Float64Regs, int_nvvm_d2f_rz_ftz>;
567def INT_NVVM_D2F_RZ : F_MATH_1<"cvt.rz.f32.f64 \t$dst, $src0;",
568 Float32Regs, Float64Regs, int_nvvm_d2f_rz>;
569def INT_NVVM_D2F_RM_FTZ : F_MATH_1<"cvt.rm.ftz.f32.f64 \t$dst, $src0;",
570 Float32Regs, Float64Regs, int_nvvm_d2f_rm_ftz>;
571def INT_NVVM_D2F_RM : F_MATH_1<"cvt.rm.f32.f64 \t$dst, $src0;",
572 Float32Regs, Float64Regs, int_nvvm_d2f_rm>;
573def INT_NVVM_D2F_RP_FTZ : F_MATH_1<"cvt.rp.ftz.f32.f64 \t$dst, $src0;",
574 Float32Regs, Float64Regs, int_nvvm_d2f_rp_ftz>;
575def INT_NVVM_D2F_RP : F_MATH_1<"cvt.rp.f32.f64 \t$dst, $src0;",
576 Float32Regs, Float64Regs, int_nvvm_d2f_rp>;
577
578def INT_NVVM_D2I_RN : F_MATH_1<"cvt.rni.s32.f64 \t$dst, $src0;",
579 Int32Regs, Float64Regs, int_nvvm_d2i_rn>;
580def INT_NVVM_D2I_RZ : F_MATH_1<"cvt.rzi.s32.f64 \t$dst, $src0;",
581 Int32Regs, Float64Regs, int_nvvm_d2i_rz>;
582def INT_NVVM_D2I_RM : F_MATH_1<"cvt.rmi.s32.f64 \t$dst, $src0;",
583 Int32Regs, Float64Regs, int_nvvm_d2i_rm>;
584def INT_NVVM_D2I_RP : F_MATH_1<"cvt.rpi.s32.f64 \t$dst, $src0;",
585 Int32Regs, Float64Regs, int_nvvm_d2i_rp>;
586
587def INT_NVVM_D2UI_RN : F_MATH_1<"cvt.rni.u32.f64 \t$dst, $src0;",
588 Int32Regs, Float64Regs, int_nvvm_d2ui_rn>;
589def INT_NVVM_D2UI_RZ : F_MATH_1<"cvt.rzi.u32.f64 \t$dst, $src0;",
590 Int32Regs, Float64Regs, int_nvvm_d2ui_rz>;
591def INT_NVVM_D2UI_RM : F_MATH_1<"cvt.rmi.u32.f64 \t$dst, $src0;",
592 Int32Regs, Float64Regs, int_nvvm_d2ui_rm>;
593def INT_NVVM_D2UI_RP : F_MATH_1<"cvt.rpi.u32.f64 \t$dst, $src0;",
594 Int32Regs, Float64Regs, int_nvvm_d2ui_rp>;
595
596def INT_NVVM_I2D_RN : F_MATH_1<"cvt.rn.f64.s32 \t$dst, $src0;",
597 Float64Regs, Int32Regs, int_nvvm_i2d_rn>;
598def INT_NVVM_I2D_RZ : F_MATH_1<"cvt.rz.f64.s32 \t$dst, $src0;",
599 Float64Regs, Int32Regs, int_nvvm_i2d_rz>;
600def INT_NVVM_I2D_RM : F_MATH_1<"cvt.rm.f64.s32 \t$dst, $src0;",
601 Float64Regs, Int32Regs, int_nvvm_i2d_rm>;
602def INT_NVVM_I2D_RP : F_MATH_1<"cvt.rp.f64.s32 \t$dst, $src0;",
603 Float64Regs, Int32Regs, int_nvvm_i2d_rp>;
604
605def INT_NVVM_UI2D_RN : F_MATH_1<"cvt.rn.f64.u32 \t$dst, $src0;",
606 Float64Regs, Int32Regs, int_nvvm_ui2d_rn>;
607def INT_NVVM_UI2D_RZ : F_MATH_1<"cvt.rz.f64.u32 \t$dst, $src0;",
608 Float64Regs, Int32Regs, int_nvvm_ui2d_rz>;
609def INT_NVVM_UI2D_RM : F_MATH_1<"cvt.rm.f64.u32 \t$dst, $src0;",
610 Float64Regs, Int32Regs, int_nvvm_ui2d_rm>;
611def INT_NVVM_UI2D_RP : F_MATH_1<"cvt.rp.f64.u32 \t$dst, $src0;",
612 Float64Regs, Int32Regs, int_nvvm_ui2d_rp>;
613
614def INT_NVVM_F2I_RN_FTZ : F_MATH_1<"cvt.rni.ftz.s32.f32 \t$dst, $src0;",
615 Int32Regs, Float32Regs, int_nvvm_f2i_rn_ftz>;
616def INT_NVVM_F2I_RN : F_MATH_1<"cvt.rni.s32.f32 \t$dst, $src0;", Int32Regs,
617 Float32Regs, int_nvvm_f2i_rn>;
618def INT_NVVM_F2I_RZ_FTZ : F_MATH_1<"cvt.rzi.ftz.s32.f32 \t$dst, $src0;",
619 Int32Regs, Float32Regs, int_nvvm_f2i_rz_ftz>;
620def INT_NVVM_F2I_RZ : F_MATH_1<"cvt.rzi.s32.f32 \t$dst, $src0;", Int32Regs,
621 Float32Regs, int_nvvm_f2i_rz>;
622def INT_NVVM_F2I_RM_FTZ : F_MATH_1<"cvt.rmi.ftz.s32.f32 \t$dst, $src0;",
623 Int32Regs, Float32Regs, int_nvvm_f2i_rm_ftz>;
624def INT_NVVM_F2I_RM : F_MATH_1<"cvt.rmi.s32.f32 \t$dst, $src0;", Int32Regs,
625 Float32Regs, int_nvvm_f2i_rm>;
626def INT_NVVM_F2I_RP_FTZ : F_MATH_1<"cvt.rpi.ftz.s32.f32 \t$dst, $src0;",
627 Int32Regs, Float32Regs, int_nvvm_f2i_rp_ftz>;
628def INT_NVVM_F2I_RP : F_MATH_1<"cvt.rpi.s32.f32 \t$dst, $src0;", Int32Regs,
629 Float32Regs, int_nvvm_f2i_rp>;
630
631def INT_NVVM_F2UI_RN_FTZ : F_MATH_1<"cvt.rni.ftz.u32.f32 \t$dst, $src0;",
632 Int32Regs, Float32Regs, int_nvvm_f2ui_rn_ftz>;
633def INT_NVVM_F2UI_RN : F_MATH_1<"cvt.rni.u32.f32 \t$dst, $src0;", Int32Regs,
634 Float32Regs, int_nvvm_f2ui_rn>;
635def INT_NVVM_F2UI_RZ_FTZ : F_MATH_1<"cvt.rzi.ftz.u32.f32 \t$dst, $src0;",
636 Int32Regs, Float32Regs, int_nvvm_f2ui_rz_ftz>;
637def INT_NVVM_F2UI_RZ : F_MATH_1<"cvt.rzi.u32.f32 \t$dst, $src0;", Int32Regs,
638 Float32Regs, int_nvvm_f2ui_rz>;
639def INT_NVVM_F2UI_RM_FTZ : F_MATH_1<"cvt.rmi.ftz.u32.f32 \t$dst, $src0;",
640 Int32Regs, Float32Regs, int_nvvm_f2ui_rm_ftz>;
641def INT_NVVM_F2UI_RM : F_MATH_1<"cvt.rmi.u32.f32 \t$dst, $src0;", Int32Regs,
642 Float32Regs, int_nvvm_f2ui_rm>;
643def INT_NVVM_F2UI_RP_FTZ : F_MATH_1<"cvt.rpi.ftz.u32.f32 \t$dst, $src0;",
644 Int32Regs, Float32Regs, int_nvvm_f2ui_rp_ftz>;
645def INT_NVVM_F2UI_RP : F_MATH_1<"cvt.rpi.u32.f32 \t$dst, $src0;", Int32Regs,
646 Float32Regs, int_nvvm_f2ui_rp>;
647
648def INT_NVVM_I2F_RN : F_MATH_1<"cvt.rn.f32.s32 \t$dst, $src0;", Float32Regs,
649 Int32Regs, int_nvvm_i2f_rn>;
650def INT_NVVM_I2F_RZ : F_MATH_1<"cvt.rz.f32.s32 \t$dst, $src0;", Float32Regs,
651 Int32Regs, int_nvvm_i2f_rz>;
652def INT_NVVM_I2F_RM : F_MATH_1<"cvt.rm.f32.s32 \t$dst, $src0;", Float32Regs,
653 Int32Regs, int_nvvm_i2f_rm>;
654def INT_NVVM_I2F_RP : F_MATH_1<"cvt.rp.f32.s32 \t$dst, $src0;", Float32Regs,
655 Int32Regs, int_nvvm_i2f_rp>;
656
657def INT_NVVM_UI2F_RN : F_MATH_1<"cvt.rn.f32.u32 \t$dst, $src0;", Float32Regs,
658 Int32Regs, int_nvvm_ui2f_rn>;
659def INT_NVVM_UI2F_RZ : F_MATH_1<"cvt.rz.f32.u32 \t$dst, $src0;", Float32Regs,
660 Int32Regs, int_nvvm_ui2f_rz>;
661def INT_NVVM_UI2F_RM : F_MATH_1<"cvt.rm.f32.u32 \t$dst, $src0;", Float32Regs,
662 Int32Regs, int_nvvm_ui2f_rm>;
663def INT_NVVM_UI2F_RP : F_MATH_1<"cvt.rp.f32.u32 \t$dst, $src0;", Float32Regs,
664 Int32Regs, int_nvvm_ui2f_rp>;
665
666def INT_NVVM_LOHI_I2D : F_MATH_2<"mov.b64 \t$dst, {{$src0, $src1}};",
667 Float64Regs, Int32Regs, Int32Regs, int_nvvm_lohi_i2d>;
668
669def INT_NVVM_D2I_LO : F_MATH_1<!strconcat("{{\n\t",
670 !strconcat(".reg .b32 %temp; \n\t",
671 !strconcat("mov.b64 \t{$dst, %temp}, $src0;\n\t",
672 "}}"))),
673 Int32Regs, Float64Regs, int_nvvm_d2i_lo>;
674def INT_NVVM_D2I_HI : F_MATH_1<!strconcat("{{\n\t",
675 !strconcat(".reg .b32 %temp; \n\t",
676 !strconcat("mov.b64 \t{%temp, $dst}, $src0;\n\t",
677 "}}"))),
678 Int32Regs, Float64Regs, int_nvvm_d2i_hi>;
679
680def INT_NVVM_F2LL_RN_FTZ : F_MATH_1<"cvt.rni.ftz.s64.f32 \t$dst, $src0;",
681 Int64Regs, Float32Regs, int_nvvm_f2ll_rn_ftz>;
682def INT_NVVM_F2LL_RN : F_MATH_1<"cvt.rni.s64.f32 \t$dst, $src0;", Int64Regs,
683 Float32Regs, int_nvvm_f2ll_rn>;
684def INT_NVVM_F2LL_RZ_FTZ : F_MATH_1<"cvt.rzi.ftz.s64.f32 \t$dst, $src0;",
685 Int64Regs, Float32Regs, int_nvvm_f2ll_rz_ftz>;
686def INT_NVVM_F2LL_RZ : F_MATH_1<"cvt.rzi.s64.f32 \t$dst, $src0;", Int64Regs,
687 Float32Regs, int_nvvm_f2ll_rz>;
688def INT_NVVM_F2LL_RM_FTZ : F_MATH_1<"cvt.rmi.ftz.s64.f32 \t$dst, $src0;",
689 Int64Regs, Float32Regs, int_nvvm_f2ll_rm_ftz>;
690def INT_NVVM_F2LL_RM : F_MATH_1<"cvt.rmi.s64.f32 \t$dst, $src0;", Int64Regs,
691 Float32Regs, int_nvvm_f2ll_rm>;
692def INT_NVVM_F2LL_RP_FTZ : F_MATH_1<"cvt.rpi.ftz.s64.f32 \t$dst, $src0;",
693 Int64Regs, Float32Regs, int_nvvm_f2ll_rp_ftz>;
694def INT_NVVM_F2LL_RP : F_MATH_1<"cvt.rpi.s64.f32 \t$dst, $src0;", Int64Regs,
695 Float32Regs, int_nvvm_f2ll_rp>;
696
697def INT_NVVM_F2ULL_RN_FTZ : F_MATH_1<"cvt.rni.ftz.u64.f32 \t$dst, $src0;",
698 Int64Regs, Float32Regs, int_nvvm_f2ull_rn_ftz>;
699def INT_NVVM_F2ULL_RN : F_MATH_1<"cvt.rni.u64.f32 \t$dst, $src0;", Int64Regs,
700 Float32Regs, int_nvvm_f2ull_rn>;
701def INT_NVVM_F2ULL_RZ_FTZ : F_MATH_1<"cvt.rzi.ftz.u64.f32 \t$dst, $src0;",
702 Int64Regs, Float32Regs, int_nvvm_f2ull_rz_ftz>;
703def INT_NVVM_F2ULL_RZ : F_MATH_1<"cvt.rzi.u64.f32 \t$dst, $src0;", Int64Regs,
704 Float32Regs, int_nvvm_f2ull_rz>;
705def INT_NVVM_F2ULL_RM_FTZ : F_MATH_1<"cvt.rmi.ftz.u64.f32 \t$dst, $src0;",
706 Int64Regs, Float32Regs, int_nvvm_f2ull_rm_ftz>;
707def INT_NVVM_F2ULL_RM : F_MATH_1<"cvt.rmi.u64.f32 \t$dst, $src0;", Int64Regs,
708 Float32Regs, int_nvvm_f2ull_rm>;
709def INT_NVVM_F2ULL_RP_FTZ : F_MATH_1<"cvt.rpi.ftz.u64.f32 \t$dst, $src0;",
710 Int64Regs, Float32Regs, int_nvvm_f2ull_rp_ftz>;
711def INT_NVVM_F2ULL_RP : F_MATH_1<"cvt.rpi.u64.f32 \t$dst, $src0;", Int64Regs,
712 Float32Regs, int_nvvm_f2ull_rp>;
713
714def INT_NVVM_D2LL_RN : F_MATH_1<"cvt.rni.s64.f64 \t$dst, $src0;", Int64Regs,
715 Float64Regs, int_nvvm_d2ll_rn>;
716def INT_NVVM_D2LL_RZ : F_MATH_1<"cvt.rzi.s64.f64 \t$dst, $src0;", Int64Regs,
717 Float64Regs, int_nvvm_d2ll_rz>;
718def INT_NVVM_D2LL_RM : F_MATH_1<"cvt.rmi.s64.f64 \t$dst, $src0;", Int64Regs,
719 Float64Regs, int_nvvm_d2ll_rm>;
720def INT_NVVM_D2LL_RP : F_MATH_1<"cvt.rpi.s64.f64 \t$dst, $src0;", Int64Regs,
721 Float64Regs, int_nvvm_d2ll_rp>;
722
723def INT_NVVM_D2ULL_RN : F_MATH_1<"cvt.rni.u64.f64 \t$dst, $src0;", Int64Regs,
724 Float64Regs, int_nvvm_d2ull_rn>;
725def INT_NVVM_D2ULL_RZ : F_MATH_1<"cvt.rzi.u64.f64 \t$dst, $src0;", Int64Regs,
726 Float64Regs, int_nvvm_d2ull_rz>;
727def INT_NVVM_D2ULL_RM : F_MATH_1<"cvt.rmi.u64.f64 \t$dst, $src0;", Int64Regs,
728 Float64Regs, int_nvvm_d2ull_rm>;
729def INT_NVVM_D2ULL_RP : F_MATH_1<"cvt.rpi.u64.f64 \t$dst, $src0;", Int64Regs,
730 Float64Regs, int_nvvm_d2ull_rp>;
731
732def INT_NVVM_LL2F_RN : F_MATH_1<"cvt.rn.f32.s64 \t$dst, $src0;", Float32Regs,
733 Int64Regs, int_nvvm_ll2f_rn>;
734def INT_NVVM_LL2F_RZ : F_MATH_1<"cvt.rz.f32.s64 \t$dst, $src0;", Float32Regs,
735 Int64Regs, int_nvvm_ll2f_rz>;
736def INT_NVVM_LL2F_RM : F_MATH_1<"cvt.rm.f32.s64 \t$dst, $src0;", Float32Regs,
737 Int64Regs, int_nvvm_ll2f_rm>;
738def INT_NVVM_LL2F_RP : F_MATH_1<"cvt.rp.f32.s64 \t$dst, $src0;", Float32Regs,
739 Int64Regs, int_nvvm_ll2f_rp>;
740def INT_NVVM_ULL2F_RN : F_MATH_1<"cvt.rn.f32.u64 \t$dst, $src0;", Float32Regs,
741 Int64Regs, int_nvvm_ull2f_rn>;
742def INT_NVVM_ULL2F_RZ : F_MATH_1<"cvt.rz.f32.u64 \t$dst, $src0;", Float32Regs,
743 Int64Regs, int_nvvm_ull2f_rz>;
744def INT_NVVM_ULL2F_RM : F_MATH_1<"cvt.rm.f32.u64 \t$dst, $src0;", Float32Regs,
745 Int64Regs, int_nvvm_ull2f_rm>;
746def INT_NVVM_ULL2F_RP : F_MATH_1<"cvt.rp.f32.u64 \t$dst, $src0;", Float32Regs,
747 Int64Regs, int_nvvm_ull2f_rp>;
748
749def INT_NVVM_LL2D_RN : F_MATH_1<"cvt.rn.f64.s64 \t$dst, $src0;", Float64Regs,
750 Int64Regs, int_nvvm_ll2d_rn>;
751def INT_NVVM_LL2D_RZ : F_MATH_1<"cvt.rz.f64.s64 \t$dst, $src0;", Float64Regs,
752 Int64Regs, int_nvvm_ll2d_rz>;
753def INT_NVVM_LL2D_RM : F_MATH_1<"cvt.rm.f64.s64 \t$dst, $src0;", Float64Regs,
754 Int64Regs, int_nvvm_ll2d_rm>;
755def INT_NVVM_LL2D_RP : F_MATH_1<"cvt.rp.f64.s64 \t$dst, $src0;", Float64Regs,
756 Int64Regs, int_nvvm_ll2d_rp>;
757def INT_NVVM_ULL2D_RN : F_MATH_1<"cvt.rn.f64.u64 \t$dst, $src0;", Float64Regs,
758 Int64Regs, int_nvvm_ull2d_rn>;
759def INT_NVVM_ULL2D_RZ : F_MATH_1<"cvt.rz.f64.u64 \t$dst, $src0;", Float64Regs,
760 Int64Regs, int_nvvm_ull2d_rz>;
761def INT_NVVM_ULL2D_RM : F_MATH_1<"cvt.rm.f64.u64 \t$dst, $src0;", Float64Regs,
762 Int64Regs, int_nvvm_ull2d_rm>;
763def INT_NVVM_ULL2D_RP : F_MATH_1<"cvt.rp.f64.u64 \t$dst, $src0;", Float64Regs,
764 Int64Regs, int_nvvm_ull2d_rp>;
765
766def INT_NVVM_F2H_RN_FTZ : F_MATH_1<!strconcat("{{\n\t",
767 !strconcat(".reg .b16 %temp;\n\t",
768 !strconcat("cvt.rn.ftz.f16.f32 \t%temp, $src0;\n\t",
769 !strconcat("mov.b16 \t$dst, %temp;\n",
770 "}}")))),
771 Int16Regs, Float32Regs, int_nvvm_f2h_rn_ftz>;
772def INT_NVVM_F2H_RN : F_MATH_1<!strconcat("{{\n\t",
773 !strconcat(".reg .b16 %temp;\n\t",
774 !strconcat("cvt.rn.f16.f32 \t%temp, $src0;\n\t",
775 !strconcat("mov.b16 \t$dst, %temp;\n",
776 "}}")))),
777 Int16Regs, Float32Regs, int_nvvm_f2h_rn>;
778
779def INT_NVVM_H2F : F_MATH_1<!strconcat("{{\n\t",
780 !strconcat(".reg .b16 %temp;\n\t",
781 !strconcat("mov.b16 \t%temp, $src0;\n\t",
782 !strconcat("cvt.f32.f16 \t$dst, %temp;\n\t",
783 "}}")))),
784 Float32Regs, Int16Regs, int_nvvm_h2f>;
785
786//
787// Bitcast
788//
789
790def INT_NVVM_BITCAST_F2I : F_MATH_1<"mov.b32 \t$dst, $src0;", Int32Regs,
791 Float32Regs, int_nvvm_bitcast_f2i>;
792def INT_NVVM_BITCAST_I2F : F_MATH_1<"mov.b32 \t$dst, $src0;", Float32Regs,
793 Int32Regs, int_nvvm_bitcast_i2f>;
794
795def INT_NVVM_BITCAST_LL2D : F_MATH_1<"mov.b64 \t$dst, $src0;", Float64Regs,
796 Int64Regs, int_nvvm_bitcast_ll2d>;
797def INT_NVVM_BITCAST_D2LL : F_MATH_1<"mov.b64 \t$dst, $src0;", Int64Regs,
798 Float64Regs, int_nvvm_bitcast_d2ll>;
799
800//-----------------------------------
801// Atomic Functions
802//-----------------------------------
803
804class ATOMIC_GLOBAL_CHK <dag ops, dag frag>
805 : PatFrag<ops, frag, [{
806 return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_GLOBAL);
807}]>;
808class ATOMIC_SHARED_CHK <dag ops, dag frag>
809 : PatFrag<ops, frag, [{
810 return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_SHARED);
811}]>;
812class ATOMIC_GENERIC_CHK <dag ops, dag frag>
813 : PatFrag<ops, frag, [{
814 return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_GENERIC);
815}]>;
816
817multiclass F_ATOMIC_2_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
818 string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
819 Operand IMMType, SDNode IMM, Predicate Pred> {
820 def reg : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, regclass:$b),
821 !strconcat("atom",
822 !strconcat(SpaceStr,
823 !strconcat(OpcStr,
824 !strconcat(TypeStr,
825 !strconcat(" \t$dst, [$addr], $b;", ""))))),
826 [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b))]>,
827 Requires<[Pred]>;
828 def imm : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, IMMType:$b),
829 !strconcat("atom",
830 !strconcat(SpaceStr,
831 !strconcat(OpcStr,
832 !strconcat(TypeStr,
833 !strconcat(" \t$dst, [$addr], $b;", ""))))),
834 [(set regclass:$dst, (IntOp ptrclass:$addr, IMM:$b))]>,
835 Requires<[Pred]>;
836}
837multiclass F_ATOMIC_2<NVPTXRegClass regclass, string SpaceStr, string TypeStr,
838 string OpcStr, PatFrag IntOp, Operand IMMType, SDNode IMM, Predicate Pred> {
839 defm p32 : F_ATOMIC_2_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
840 IntOp, IMMType, IMM, Pred>;
841 defm p64 : F_ATOMIC_2_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
842 IntOp, IMMType, IMM, Pred>;
843}
844
845// has 2 operands, neg the second one
846multiclass F_ATOMIC_2_NEG_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
847 string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
848 Operand IMMType, Predicate Pred> {
849 def reg : NVPTXInst<(outs regclass:$dst), (ins ptrclass:$addr, regclass:$b),
850 !strconcat("{{ \n\t",
851 !strconcat(".reg \t.s",
852 !strconcat(TypeStr,
853 !strconcat(" temp; \n\t",
854 !strconcat("neg.s",
855 !strconcat(TypeStr,
856 !strconcat(" \ttemp, $b; \n\t",
857 !strconcat("atom",
858 !strconcat(SpaceStr,
859 !strconcat(OpcStr,
860 !strconcat(".u",
861 !strconcat(TypeStr,
862 !strconcat(" \t$dst, [$addr], temp; \n\t",
863 !strconcat("}}", "")))))))))))))),
864 [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b))]>,
865 Requires<[Pred]>;
866}
867multiclass F_ATOMIC_2_NEG<NVPTXRegClass regclass, string SpaceStr,
868 string TypeStr, string OpcStr, PatFrag IntOp, Operand IMMType,
869 Predicate Pred> {
870 defm p32: F_ATOMIC_2_NEG_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
871 IntOp, IMMType, Pred> ;
872 defm p64: F_ATOMIC_2_NEG_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
873 IntOp, IMMType, Pred> ;
874}
875
876// has 3 operands
877multiclass F_ATOMIC_3_imp<NVPTXRegClass ptrclass, NVPTXRegClass regclass,
878 string SpaceStr, string TypeStr, string OpcStr, PatFrag IntOp,
879 Operand IMMType, Predicate Pred> {
880 def reg : NVPTXInst<(outs regclass:$dst),
881 (ins ptrclass:$addr, regclass:$b, regclass:$c),
882 !strconcat("atom",
883 !strconcat(SpaceStr,
884 !strconcat(OpcStr,
885 !strconcat(TypeStr,
886 !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
887 [(set regclass:$dst,
888 (IntOp ptrclass:$addr, regclass:$b, regclass:$c))]>,
889 Requires<[Pred]>;
890 def imm1 : NVPTXInst<(outs regclass:$dst),
891 (ins ptrclass:$addr, IMMType:$b, regclass:$c),
892 !strconcat("atom",
893 !strconcat(SpaceStr,
894 !strconcat(OpcStr,
895 !strconcat(TypeStr,
896 !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
897 [(set regclass:$dst, (IntOp ptrclass:$addr, imm:$b, regclass:$c))]>,
898 Requires<[Pred]>;
899 def imm2 : NVPTXInst<(outs regclass:$dst),
900 (ins ptrclass:$addr, regclass:$b, IMMType:$c),
901 !strconcat("atom",
902 !strconcat(SpaceStr,
903 !strconcat(OpcStr,
904 !strconcat(TypeStr,
905 !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
906 [(set regclass:$dst, (IntOp ptrclass:$addr, regclass:$b, imm:$c))]>,
907 Requires<[Pred]>;
908 def imm3 : NVPTXInst<(outs regclass:$dst),
909 (ins ptrclass:$addr, IMMType:$b, IMMType:$c),
910 !strconcat("atom",
911 !strconcat(SpaceStr,
912 !strconcat(OpcStr,
913 !strconcat(TypeStr,
914 !strconcat(" \t$dst, [$addr], $b, $c;", ""))))),
915 [(set regclass:$dst, (IntOp ptrclass:$addr, imm:$b, imm:$c))]>,
916 Requires<[Pred]>;
917}
918multiclass F_ATOMIC_3<NVPTXRegClass regclass, string SpaceStr, string TypeStr,
919 string OpcStr, PatFrag IntOp, Operand IMMType, Predicate Pred> {
920 defm p32 : F_ATOMIC_3_imp<Int32Regs, regclass, SpaceStr, TypeStr, OpcStr,
921 IntOp, IMMType, Pred>;
922 defm p64 : F_ATOMIC_3_imp<Int64Regs, regclass, SpaceStr, TypeStr, OpcStr,
923 IntOp, IMMType, Pred>;
924}
925
926// atom_add
927
928def atomic_load_add_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
929 (atomic_load_add_32 node:$a, node:$b)>;
930def atomic_load_add_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
931 (atomic_load_add_32 node:$a, node:$b)>;
932def atomic_load_add_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
933 (atomic_load_add_32 node:$a, node:$b)>;
934def atomic_load_add_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
935 (atomic_load_add_64 node:$a, node:$b)>;
936def atomic_load_add_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
937 (atomic_load_add_64 node:$a, node:$b)>;
938def atomic_load_add_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
939 (atomic_load_add_64 node:$a, node:$b)>;
940def atomic_load_add_f32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
941 (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
942def atomic_load_add_f32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
943 (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
944def atomic_load_add_f32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
945 (int_nvvm_atomic_load_add_f32 node:$a, node:$b)>;
946
947defm INT_PTX_ATOM_ADD_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".add",
948 atomic_load_add_32_g, i32imm, imm, hasAtomRedG32>;
949defm INT_PTX_ATOM_ADD_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".add",
950 atomic_load_add_32_s, i32imm, imm, hasAtomRedS32>;
951defm INT_PTX_ATOM_ADD_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".add",
952 atomic_load_add_32_gen, i32imm, imm, hasAtomRedGen32>;
953defm INT_PTX_ATOM_ADD_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
954 ".add", atomic_load_add_32_gen, i32imm, imm, useAtomRedG32forGen32>;
955
956defm INT_PTX_ATOM_ADD_G_64 : F_ATOMIC_2<Int64Regs, ".global", ".u64", ".add",
957 atomic_load_add_64_g, i64imm, imm, hasAtomRedG64>;
958defm INT_PTX_ATOM_ADD_S_64 : F_ATOMIC_2<Int64Regs, ".shared", ".u64", ".add",
959 atomic_load_add_64_s, i64imm, imm, hasAtomRedS64>;
960defm INT_PTX_ATOM_ADD_GEN_64 : F_ATOMIC_2<Int64Regs, "", ".u64", ".add",
961 atomic_load_add_64_gen, i64imm, imm, hasAtomRedGen64>;
962defm INT_PTX_ATOM_ADD_GEN_64_USE_G : F_ATOMIC_2<Int64Regs, ".global", ".u64",
963 ".add", atomic_load_add_64_gen, i64imm, imm, useAtomRedG64forGen64>;
964
965defm INT_PTX_ATOM_ADD_G_F32 : F_ATOMIC_2<Float32Regs, ".global", ".f32", ".add",
966 atomic_load_add_f32_g, f32imm, fpimm, hasAtomAddF32>;
967defm INT_PTX_ATOM_ADD_S_F32 : F_ATOMIC_2<Float32Regs, ".shared", ".f32", ".add",
968 atomic_load_add_f32_s, f32imm, fpimm, hasAtomAddF32>;
969defm INT_PTX_ATOM_ADD_GEN_F32 : F_ATOMIC_2<Float32Regs, "", ".f32", ".add",
970 atomic_load_add_f32_gen, f32imm, fpimm, hasAtomAddF32>;
971
972// atom_sub
973
974def atomic_load_sub_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
975 (atomic_load_sub_32 node:$a, node:$b)>;
976def atomic_load_sub_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
977 (atomic_load_sub_32 node:$a, node:$b)>;
978def atomic_load_sub_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
979 (atomic_load_sub_32 node:$a, node:$b)>;
980def atomic_load_sub_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
981 (atomic_load_sub_64 node:$a, node:$b)>;
982def atomic_load_sub_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
983 (atomic_load_sub_64 node:$a, node:$b)>;
984def atomic_load_sub_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
985 (atomic_load_sub_64 node:$a, node:$b)>;
986
987defm INT_PTX_ATOM_SUB_G_32 : F_ATOMIC_2_NEG<Int32Regs, ".global", "32", ".add",
988 atomic_load_sub_32_g, i32imm, hasAtomRedG32>;
989defm INT_PTX_ATOM_SUB_G_64 : F_ATOMIC_2_NEG<Int64Regs, ".global", "64", ".add",
990 atomic_load_sub_64_g, i64imm, hasAtomRedG64>;
991defm INT_PTX_ATOM_SUB_GEN_32 : F_ATOMIC_2_NEG<Int32Regs, "", "32", ".add",
992 atomic_load_sub_32_gen, i32imm, hasAtomRedGen32>;
993defm INT_PTX_ATOM_SUB_GEN_32_USE_G : F_ATOMIC_2_NEG<Int32Regs, ".global", "32",
994 ".add", atomic_load_sub_32_gen, i32imm, useAtomRedG32forGen32>;
995defm INT_PTX_ATOM_SUB_S_32 : F_ATOMIC_2_NEG<Int32Regs, ".shared", "32", ".add",
996 atomic_load_sub_32_s, i32imm, hasAtomRedS32>;
997defm INT_PTX_ATOM_SUB_S_64 : F_ATOMIC_2_NEG<Int64Regs, ".shared", "64", ".add",
998 atomic_load_sub_64_s, i64imm, hasAtomRedS64>;
999defm INT_PTX_ATOM_SUB_GEN_64 : F_ATOMIC_2_NEG<Int64Regs, "", "64", ".add",
1000 atomic_load_sub_64_gen, i64imm, hasAtomRedGen64>;
1001defm INT_PTX_ATOM_SUB_GEN_64_USE_G : F_ATOMIC_2_NEG<Int64Regs, ".global", "64",
1002 ".add", atomic_load_sub_64_gen, i64imm, useAtomRedG64forGen64>;
1003
1004// atom_swap
1005
1006def atomic_swap_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1007 (atomic_swap_32 node:$a, node:$b)>;
1008def atomic_swap_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1009 (atomic_swap_32 node:$a, node:$b)>;
1010def atomic_swap_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1011 (atomic_swap_32 node:$a, node:$b)>;
1012def atomic_swap_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1013 (atomic_swap_64 node:$a, node:$b)>;
1014def atomic_swap_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1015 (atomic_swap_64 node:$a, node:$b)>;
1016def atomic_swap_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1017 (atomic_swap_64 node:$a, node:$b)>;
1018
1019defm INT_PTX_ATOM_SWAP_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".exch",
1020 atomic_swap_32_g, i32imm, imm, hasAtomRedG32>;
1021defm INT_PTX_ATOM_SWAP_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".exch",
1022 atomic_swap_32_s, i32imm, imm, hasAtomRedS32>;
1023defm INT_PTX_ATOM_SWAP_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".exch",
1024 atomic_swap_32_gen, i32imm, imm, hasAtomRedGen32>;
1025defm INT_PTX_ATOM_SWAP_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1026 ".exch", atomic_swap_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1027defm INT_PTX_ATOM_SWAP_G_64 : F_ATOMIC_2<Int64Regs, ".global", ".b64", ".exch",
1028 atomic_swap_64_g, i64imm, imm, hasAtomRedG64>;
1029defm INT_PTX_ATOM_SWAP_S_64 : F_ATOMIC_2<Int64Regs, ".shared", ".b64", ".exch",
1030 atomic_swap_64_s, i64imm, imm, hasAtomRedS64>;
1031defm INT_PTX_ATOM_SWAP_GEN_64 : F_ATOMIC_2<Int64Regs, "", ".b64", ".exch",
1032 atomic_swap_64_gen, i64imm, imm, hasAtomRedGen64>;
1033defm INT_PTX_ATOM_SWAP_GEN_64_USE_G : F_ATOMIC_2<Int64Regs, ".global", ".b64",
1034 ".exch", atomic_swap_64_gen, i64imm, imm, useAtomRedG64forGen64>;
1035
1036// atom_max
1037
1038def atomic_load_max_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b)
1039 , (atomic_load_max_32 node:$a, node:$b)>;
1040def atomic_load_max_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1041 (atomic_load_max_32 node:$a, node:$b)>;
1042def atomic_load_max_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1043 (atomic_load_max_32 node:$a, node:$b)>;
1044def atomic_load_umax_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1045 (atomic_load_umax_32 node:$a, node:$b)>;
1046def atomic_load_umax_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1047 (atomic_load_umax_32 node:$a, node:$b)>;
1048def atomic_load_umax_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1049 (atomic_load_umax_32 node:$a, node:$b)>;
1050
1051defm INT_PTX_ATOM_LOAD_MAX_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".s32",
1052 ".max", atomic_load_max_32_g, i32imm, imm, hasAtomRedG32>;
1053defm INT_PTX_ATOM_LOAD_MAX_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".s32",
1054 ".max", atomic_load_max_32_s, i32imm, imm, hasAtomRedS32>;
1055defm INT_PTX_ATOM_LOAD_MAX_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".s32", ".max",
1056 atomic_load_max_32_gen, i32imm, imm, hasAtomRedGen32>;
1057defm INT_PTX_ATOM_LOAD_MAX_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1058 ".s32", ".max", atomic_load_max_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1059defm INT_PTX_ATOM_LOAD_UMAX_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1060 ".max", atomic_load_umax_32_g, i32imm, imm, hasAtomRedG32>;
1061defm INT_PTX_ATOM_LOAD_UMAX_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32",
1062 ".max", atomic_load_umax_32_s, i32imm, imm, hasAtomRedS32>;
1063defm INT_PTX_ATOM_LOAD_UMAX_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".max",
1064 atomic_load_umax_32_gen, i32imm, imm, hasAtomRedGen32>;
1065defm INT_PTX_ATOM_LOAD_UMAX_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1066 ".u32", ".max", atomic_load_umax_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1067
1068// atom_min
1069
1070def atomic_load_min_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1071 (atomic_load_min_32 node:$a, node:$b)>;
1072def atomic_load_min_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1073 (atomic_load_min_32 node:$a, node:$b)>;
1074def atomic_load_min_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1075 (atomic_load_min_32 node:$a, node:$b)>;
1076def atomic_load_umin_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1077 (atomic_load_umin_32 node:$a, node:$b)>;
1078def atomic_load_umin_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1079 (atomic_load_umin_32 node:$a, node:$b)>;
1080def atomic_load_umin_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1081 (atomic_load_umin_32 node:$a, node:$b)>;
1082
1083defm INT_PTX_ATOM_LOAD_MIN_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".s32",
1084 ".min", atomic_load_min_32_g, i32imm, imm, hasAtomRedG32>;
1085defm INT_PTX_ATOM_LOAD_MIN_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".s32",
1086 ".min", atomic_load_min_32_s, i32imm, imm, hasAtomRedS32>;
1087defm INT_PTX_ATOM_LOAD_MIN_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".s32", ".min",
1088 atomic_load_min_32_gen, i32imm, imm, hasAtomRedGen32>;
1089defm INT_PTX_ATOM_LOAD_MIN_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1090 ".s32", ".min", atomic_load_min_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1091defm INT_PTX_ATOM_LOAD_UMIN_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1092 ".min", atomic_load_umin_32_g, i32imm, imm, hasAtomRedG32>;
1093defm INT_PTX_ATOM_LOAD_UMIN_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32",
1094 ".min", atomic_load_umin_32_s, i32imm, imm, hasAtomRedS32>;
1095defm INT_PTX_ATOM_LOAD_UMIN_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".min",
1096 atomic_load_umin_32_gen, i32imm, imm, hasAtomRedGen32>;
1097defm INT_PTX_ATOM_LOAD_UMIN_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global",
1098 ".u32", ".min", atomic_load_umin_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1099
1100// atom_inc atom_dec
1101
1102def atomic_load_inc_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1103 (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
1104def atomic_load_inc_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1105 (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
1106def atomic_load_inc_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1107 (int_nvvm_atomic_load_inc_32 node:$a, node:$b)>;
1108def atomic_load_dec_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1109 (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
1110def atomic_load_dec_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1111 (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
1112def atomic_load_dec_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1113 (int_nvvm_atomic_load_dec_32 node:$a, node:$b)>;
1114
1115defm INT_PTX_ATOM_INC_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".inc",
1116 atomic_load_inc_32_g, i32imm, imm, hasAtomRedG32>;
1117defm INT_PTX_ATOM_INC_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".inc",
1118 atomic_load_inc_32_s, i32imm, imm, hasAtomRedS32>;
1119defm INT_PTX_ATOM_INC_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".inc",
1120 atomic_load_inc_32_gen, i32imm, imm, hasAtomRedGen32>;
1121defm INT_PTX_ATOM_INC_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1122 ".inc", atomic_load_inc_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1123defm INT_PTX_ATOM_DEC_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".u32", ".dec",
1124 atomic_load_dec_32_g, i32imm, imm, hasAtomRedG32>;
1125defm INT_PTX_ATOM_DEC_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".u32", ".dec",
1126 atomic_load_dec_32_s, i32imm, imm, hasAtomRedS32>;
1127defm INT_PTX_ATOM_DEC_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".u32", ".dec",
1128 atomic_load_dec_32_gen, i32imm, imm, hasAtomRedGen32>;
1129defm INT_PTX_ATOM_DEC_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".u32",
1130 ".dec", atomic_load_dec_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1131
1132// atom_and
1133
1134def atomic_load_and_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1135 (atomic_load_and_32 node:$a, node:$b)>;
1136def atomic_load_and_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1137 (atomic_load_and_32 node:$a, node:$b)>;
1138def atomic_load_and_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1139 (atomic_load_and_32 node:$a, node:$b)>;
1140
1141defm INT_PTX_ATOM_AND_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".and",
1142 atomic_load_and_32_g, i32imm, imm, hasAtomRedG32>;
1143defm INT_PTX_ATOM_AND_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".and",
1144 atomic_load_and_32_s, i32imm, imm, hasAtomRedS32>;
1145defm INT_PTX_ATOM_AND_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".and",
1146 atomic_load_and_32_gen, i32imm, imm, hasAtomRedGen32>;
1147defm INT_PTX_ATOM_AND_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1148 ".and", atomic_load_and_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1149
1150// atom_or
1151
1152def atomic_load_or_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1153 (atomic_load_or_32 node:$a, node:$b)>;
1154def atomic_load_or_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1155 (atomic_load_or_32 node:$a, node:$b)>;
1156def atomic_load_or_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1157 (atomic_load_or_32 node:$a, node:$b)>;
1158
1159defm INT_PTX_ATOM_OR_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".or",
1160 atomic_load_or_32_g, i32imm, imm, hasAtomRedG32>;
1161defm INT_PTX_ATOM_OR_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".or",
1162 atomic_load_or_32_gen, i32imm, imm, hasAtomRedGen32>;
1163defm INT_PTX_ATOM_OR_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1164 ".or", atomic_load_or_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1165defm INT_PTX_ATOM_OR_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".or",
1166 atomic_load_or_32_s, i32imm, imm, hasAtomRedS32>;
1167
1168// atom_xor
1169
1170def atomic_load_xor_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b),
1171 (atomic_load_xor_32 node:$a, node:$b)>;
1172def atomic_load_xor_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b),
1173 (atomic_load_xor_32 node:$a, node:$b)>;
1174def atomic_load_xor_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b),
1175 (atomic_load_xor_32 node:$a, node:$b)>;
1176
1177defm INT_PTX_ATOM_XOR_G_32 : F_ATOMIC_2<Int32Regs, ".global", ".b32", ".xor",
1178 atomic_load_xor_32_g, i32imm, imm, hasAtomRedG32>;
1179defm INT_PTX_ATOM_XOR_S_32 : F_ATOMIC_2<Int32Regs, ".shared", ".b32", ".xor",
1180 atomic_load_xor_32_s, i32imm, imm, hasAtomRedS32>;
1181defm INT_PTX_ATOM_XOR_GEN_32 : F_ATOMIC_2<Int32Regs, "", ".b32", ".xor",
1182 atomic_load_xor_32_gen, i32imm, imm, hasAtomRedGen32>;
1183defm INT_PTX_ATOM_XOR_GEN_32_USE_G : F_ATOMIC_2<Int32Regs, ".global", ".b32",
1184 ".xor", atomic_load_xor_32_gen, i32imm, imm, useAtomRedG32forGen32>;
1185
1186// atom_cas
1187
1188def atomic_cmp_swap_32_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b, node:$c),
1189 (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
1190def atomic_cmp_swap_32_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b, node:$c),
1191 (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
1192def atomic_cmp_swap_32_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b, node:$c),
1193 (atomic_cmp_swap_32 node:$a, node:$b, node:$c)>;
1194def atomic_cmp_swap_64_g: ATOMIC_GLOBAL_CHK<(ops node:$a, node:$b, node:$c),
1195 (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
1196def atomic_cmp_swap_64_s: ATOMIC_SHARED_CHK<(ops node:$a, node:$b, node:$c),
1197 (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
1198def atomic_cmp_swap_64_gen: ATOMIC_GENERIC_CHK<(ops node:$a, node:$b, node:$c),
1199 (atomic_cmp_swap_64 node:$a, node:$b, node:$c)>;
1200
1201defm INT_PTX_ATOM_CAS_G_32 : F_ATOMIC_3<Int32Regs, ".global", ".b32", ".cas",
1202 atomic_cmp_swap_32_g, i32imm, hasAtomRedG32>;
1203defm INT_PTX_ATOM_CAS_S_32 : F_ATOMIC_3<Int32Regs, ".shared", ".b32", ".cas",
1204 atomic_cmp_swap_32_s, i32imm, hasAtomRedS32>;
1205defm INT_PTX_ATOM_CAS_GEN_32 : F_ATOMIC_3<Int32Regs, "", ".b32", ".cas",
1206 atomic_cmp_swap_32_gen, i32imm, hasAtomRedGen32>;
1207defm INT_PTX_ATOM_CAS_GEN_32_USE_G : F_ATOMIC_3<Int32Regs, ".global", ".b32",
1208 ".cas", atomic_cmp_swap_32_gen, i32imm, useAtomRedG32forGen32>;
1209defm INT_PTX_ATOM_CAS_G_64 : F_ATOMIC_3<Int64Regs, ".global", ".b64", ".cas",
1210 atomic_cmp_swap_64_g, i64imm, hasAtomRedG64>;
1211defm INT_PTX_ATOM_CAS_S_64 : F_ATOMIC_3<Int64Regs, ".shared", ".b64", ".cas",
1212 atomic_cmp_swap_64_s, i64imm, hasAtomRedS64>;
1213defm INT_PTX_ATOM_CAS_GEN_64 : F_ATOMIC_3<Int64Regs, "", ".b64", ".cas",
1214 atomic_cmp_swap_64_gen, i64imm, hasAtomRedGen64>;
1215defm INT_PTX_ATOM_CAS_GEN_64_USE_G : F_ATOMIC_3<Int64Regs, ".global", ".b64",
1216 ".cas", atomic_cmp_swap_64_gen, i64imm, useAtomRedG64forGen64>;
1217
1218
1219//-----------------------------------
1220// Read Special Registers
1221//-----------------------------------
1222class F_SREG<string OpStr, NVPTXRegClass regclassOut, Intrinsic IntOp> :
1223 NVPTXInst<(outs regclassOut:$dst), (ins),
1224 OpStr,
1225 [(set regclassOut:$dst, (IntOp))]>;
1226
1227def INT_PTX_SREG_TID_X : F_SREG<"mov.u32 \t$dst, %tid.x;", Int32Regs,
1228 int_nvvm_read_ptx_sreg_tid_x>;
1229def INT_PTX_SREG_TID_Y : F_SREG<"mov.u32 \t$dst, %tid.y;", Int32Regs,
1230 int_nvvm_read_ptx_sreg_tid_y>;
1231def INT_PTX_SREG_TID_Z : F_SREG<"mov.u32 \t$dst, %tid.z;", Int32Regs,
1232 int_nvvm_read_ptx_sreg_tid_z>;
1233
1234def INT_PTX_SREG_NTID_X : F_SREG<"mov.u32 \t$dst, %ntid.x;", Int32Regs,
1235 int_nvvm_read_ptx_sreg_ntid_x>;
1236def INT_PTX_SREG_NTID_Y : F_SREG<"mov.u32 \t$dst, %ntid.y;", Int32Regs,
1237 int_nvvm_read_ptx_sreg_ntid_y>;
1238def INT_PTX_SREG_NTID_Z : F_SREG<"mov.u32 \t$dst, %ntid.z;", Int32Regs,
1239 int_nvvm_read_ptx_sreg_ntid_z>;
1240
1241def INT_PTX_SREG_CTAID_X : F_SREG<"mov.u32 \t$dst, %ctaid.x;", Int32Regs,
1242 int_nvvm_read_ptx_sreg_ctaid_x>;
1243def INT_PTX_SREG_CTAID_Y : F_SREG<"mov.u32 \t$dst, %ctaid.y;", Int32Regs,
1244 int_nvvm_read_ptx_sreg_ctaid_y>;
1245def INT_PTX_SREG_CTAID_Z : F_SREG<"mov.u32 \t$dst, %ctaid.z;", Int32Regs,
1246 int_nvvm_read_ptx_sreg_ctaid_z>;
1247
1248def INT_PTX_SREG_NCTAID_X : F_SREG<"mov.u32 \t$dst, %nctaid.x;", Int32Regs,
1249 int_nvvm_read_ptx_sreg_nctaid_x>;
1250def INT_PTX_SREG_NCTAID_Y : F_SREG<"mov.u32 \t$dst, %nctaid.y;", Int32Regs,
1251 int_nvvm_read_ptx_sreg_nctaid_y>;
1252def INT_PTX_SREG_NCTAID_Z : F_SREG<"mov.u32 \t$dst, %nctaid.z;", Int32Regs,
1253 int_nvvm_read_ptx_sreg_nctaid_z>;
1254
1255def INT_PTX_SREG_WARPSIZE : F_SREG<"mov.u32 \t$dst, WARP_SZ;", Int32Regs,
1256 int_nvvm_read_ptx_sreg_warpsize>;
1257
1258
1259//-----------------------------------
1260// Support for ldu on sm_20 or later
1261//-----------------------------------
1262
1263// Scalar
1264// @TODO: Revisit this, Changed imemAny to imem
1265multiclass LDU_G<string TyStr, NVPTXRegClass regclass, Intrinsic IntOp> {
1266 def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
1267 !strconcat("ldu.global.", TyStr),
1268 [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDU]>;
1269 def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
1270 !strconcat("ldu.global.", TyStr),
1271 [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDU]>;
1272 def avar: NVPTXInst<(outs regclass:$result), (ins imem:$src),
1273 !strconcat("ldu.global.", TyStr),
1274 [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
1275 Requires<[hasLDU]>;
1276 def ari : NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
1277 !strconcat("ldu.global.", TyStr),
1278 [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDU]>;
1279 def ari64 : NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
1280 !strconcat("ldu.global.", TyStr),
1281 [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDU]>;
1282}
1283
1284defm INT_PTX_LDU_GLOBAL_i8 : LDU_G<"u8 \t$result, [$src];", Int8Regs,
1285int_nvvm_ldu_global_i>;
1286defm INT_PTX_LDU_GLOBAL_i16 : LDU_G<"u16 \t$result, [$src];", Int16Regs,
1287int_nvvm_ldu_global_i>;
1288defm INT_PTX_LDU_GLOBAL_i32 : LDU_G<"u32 \t$result, [$src];", Int32Regs,
1289int_nvvm_ldu_global_i>;
1290defm INT_PTX_LDU_GLOBAL_i64 : LDU_G<"u64 \t$result, [$src];", Int64Regs,
1291int_nvvm_ldu_global_i>;
1292defm INT_PTX_LDU_GLOBAL_f32 : LDU_G<"f32 \t$result, [$src];", Float32Regs,
1293int_nvvm_ldu_global_f>;
1294defm INT_PTX_LDU_GLOBAL_f64 : LDU_G<"f64 \t$result, [$src];", Float64Regs,
1295int_nvvm_ldu_global_f>;
1296defm INT_PTX_LDU_GLOBAL_p32 : LDU_G<"u32 \t$result, [$src];", Int32Regs,
1297int_nvvm_ldu_global_p>;
1298defm INT_PTX_LDU_GLOBAL_p64 : LDU_G<"u64 \t$result, [$src];", Int64Regs,
1299int_nvvm_ldu_global_p>;
1300
1301// vector
1302
1303// Elementized vector ldu
1304multiclass VLDU_G_ELE_V2<string TyStr, NVPTXRegClass regclass> {
1305 def _32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1306 (ins Int32Regs:$src),
1307 !strconcat("ldu.global.", TyStr), []>;
1308 def _64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1309 (ins Int64Regs:$src),
1310 !strconcat("ldu.global.", TyStr), []>;
1311}
1312
1313multiclass VLDU_G_ELE_V4<string TyStr, NVPTXRegClass regclass> {
1314 def _32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1315 regclass:$dst4), (ins Int32Regs:$src),
1316 !strconcat("ldu.global.", TyStr), []>;
1317 def _64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2, regclass:$dst3,
1318 regclass:$dst4), (ins Int64Regs:$src),
1319 !strconcat("ldu.global.", TyStr), []>;
1320}
1321
1322defm INT_PTX_LDU_G_v2i8_ELE
1323 : VLDU_G_ELE_V2<"v2.u8 \t{{$dst1, $dst2}}, [$src];", Int8Regs>;
1324defm INT_PTX_LDU_G_v2i16_ELE
1325 : VLDU_G_ELE_V2<"v2.u16 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
1326defm INT_PTX_LDU_G_v2i32_ELE
1327 : VLDU_G_ELE_V2<"v2.u32 \t{{$dst1, $dst2}}, [$src];", Int32Regs>;
1328defm INT_PTX_LDU_G_v2f32_ELE
1329 : VLDU_G_ELE_V2<"v2.f32 \t{{$dst1, $dst2}}, [$src];", Float32Regs>;
1330defm INT_PTX_LDU_G_v2i64_ELE
1331 : VLDU_G_ELE_V2<"v2.u64 \t{{$dst1, $dst2}}, [$src];", Int64Regs>;
1332defm INT_PTX_LDU_G_v2f64_ELE
1333 : VLDU_G_ELE_V2<"v2.f64 \t{{$dst1, $dst2}}, [$src];", Float64Regs>;
1334defm INT_PTX_LDU_G_v4i8_ELE
1335 : VLDU_G_ELE_V4<"v4.u8 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int8Regs>;
1336defm INT_PTX_LDU_G_v4i16_ELE
1337 : VLDU_G_ELE_V4<"v4.u16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
1338 Int16Regs>;
1339defm INT_PTX_LDU_G_v4i32_ELE
1340 : VLDU_G_ELE_V4<"v4.u32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
1341 Int32Regs>;
1342defm INT_PTX_LDU_G_v4f32_ELE
1343 : VLDU_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];",
1344 Float32Regs>;
1345
Justin Holewinskibe8dc642013-02-12 14:18:49 +0000[diff] [blame]1346
1347//-----------------------------------
1348// Support for ldg on sm_35 or later
1349//-----------------------------------
1350
1351def ldg_i8 : PatFrag<(ops node:$ptr), (int_nvvm_ldg_global_i node:$ptr), [{
1352 MemIntrinsicSDNode *M = cast<MemIntrinsicSDNode>(N);
1353 return M->getMemoryVT() == MVT::i8;
1354}]>;
1355
1356multiclass LDG_G<string TyStr, NVPTXRegClass regclass, Intrinsic IntOp> {
1357 def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
1358 !strconcat("ld.global.nc.", TyStr),
1359 [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDG]>;
1360 def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
1361 !strconcat("ld.global.nc.", TyStr),
1362 [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDG]>;
1363 def avar: NVPTXInst<(outs regclass:$result), (ins imem:$src),
1364 !strconcat("ld.global.nc.", TyStr),
1365 [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
1366 Requires<[hasLDG]>;
1367 def ari : NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
1368 !strconcat("ld.global.nc.", TyStr),
1369 [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDG]>;
1370 def ari64 : NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
1371 !strconcat("ld.global.nc.", TyStr),
1372 [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDG]>;
Justin Holewinskiae556d32012-05-04 20:18:50 +0000[diff] [blame]1373}
1374
Justin Holewinskibe8dc642013-02-12 14:18:49 +0000[diff] [blame]1375multiclass LDG_G_NOINTRIN<string TyStr, NVPTXRegClass regclass, PatFrag IntOp> {
1376 def areg: NVPTXInst<(outs regclass:$result), (ins Int32Regs:$src),
1377 !strconcat("ld.global.nc.", TyStr),
1378 [(set regclass:$result, (IntOp Int32Regs:$src))]>, Requires<[hasLDG]>;
1379 def areg64: NVPTXInst<(outs regclass:$result), (ins Int64Regs:$src),
1380 !strconcat("ld.global.nc.", TyStr),
1381 [(set regclass:$result, (IntOp Int64Regs:$src))]>, Requires<[hasLDG]>;
1382 def avar: NVPTXInst<(outs regclass:$result), (ins imem:$src),
1383 !strconcat("ld.global.nc.", TyStr),
1384 [(set regclass:$result, (IntOp (Wrapper tglobaladdr:$src)))]>,
1385 Requires<[hasLDG]>;
1386 def ari : NVPTXInst<(outs regclass:$result), (ins MEMri:$src),
1387 !strconcat("ld.global.nc.", TyStr),
1388 [(set regclass:$result, (IntOp ADDRri:$src))]>, Requires<[hasLDG]>;
1389 def ari64 : NVPTXInst<(outs regclass:$result), (ins MEMri64:$src),
1390 !strconcat("ld.global.nc.", TyStr),
1391 [(set regclass:$result, (IntOp ADDRri64:$src))]>, Requires<[hasLDG]>;
Justin Holewinskiae556d32012-05-04 20:18:50 +0000[diff] [blame]1392}
1393
Justin Holewinskibe8dc642013-02-12 14:18:49 +0000[diff] [blame]1394defm INT_PTX_LDG_GLOBAL_i8
1395 : LDG_G_NOINTRIN<"u8 \t$result, [$src];", Int16Regs, ldg_i8>;
1396defm INT_PTX_LDG_GLOBAL_i16
1397 : LDG_G<"u16 \t$result, [$src];", Int16Regs, int_nvvm_ldg_global_i>;
1398defm INT_PTX_LDG_GLOBAL_i32
1399 : LDG_G<"u32 \t$result, [$src];", Int32Regs, int_nvvm_ldg_global_i>;
1400defm INT_PTX_LDG_GLOBAL_i64
1401 : LDG_G<"u64 \t$result, [$src];", Int64Regs, int_nvvm_ldg_global_i>;
1402defm INT_PTX_LDG_GLOBAL_f32
1403 : LDG_G<"f32 \t$result, [$src];", Float32Regs, int_nvvm_ldg_global_f>;
1404defm INT_PTX_LDG_GLOBAL_f64
1405 : LDG_G<"f64 \t$result, [$src];", Float64Regs, int_nvvm_ldg_global_f>;
1406defm INT_PTX_LDG_GLOBAL_p32
1407 : LDG_G<"u32 \t$result, [$src];", Int32Regs, int_nvvm_ldg_global_p>;
1408defm INT_PTX_LDG_GLOBAL_p64
1409 : LDG_G<"u64 \t$result, [$src];", Int64Regs, int_nvvm_ldg_global_p>;
1410
1411// vector
1412
1413// Elementized vector ldg
1414multiclass VLDG_G_ELE_V2<string TyStr, NVPTXRegClass regclass> {
1415 def _32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1416 (ins Int32Regs:$src),
1417 !strconcat("ld.global.nc.", TyStr), []>;
1418 def _64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2),
1419 (ins Int64Regs:$src),
1420 !strconcat("ld.global.nc.", TyStr), []>;
1421}
1422
1423multiclass VLDG_G_ELE_V4<string TyStr, NVPTXRegClass regclass> {
1424 def _32: NVPTXInst<(outs regclass:$dst1, regclass:$dst2,
1425 regclass:$dst3, regclass:$dst4), (ins Int32Regs:$src),
1426 !strconcat("ld.global.nc.", TyStr), []>;
1427 def _64: NVPTXInst<(outs regclass:$dst1, regclass:$dst2,
1428 regclass:$dst3, regclass:$dst4), (ins Int64Regs:$src),
1429 !strconcat("ld.global.nc.", TyStr), []>;
1430}
1431
1432// FIXME: 8-bit LDG should be fixed once LDG/LDU nodes are made into proper loads.
1433defm INT_PTX_LDG_G_v2i8_ELE
1434 : VLDG_G_ELE_V2<"v2.u8 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
1435defm INT_PTX_LDG_G_v2i16_ELE
1436 : VLDG_G_ELE_V2<"v2.u16 \t{{$dst1, $dst2}}, [$src];", Int16Regs>;
1437defm INT_PTX_LDG_G_v2i32_ELE
1438 : VLDG_G_ELE_V2<"v2.u32 \t{{$dst1, $dst2}}, [$src];", Int32Regs>;
1439defm INT_PTX_LDG_G_v2f32_ELE
1440 : VLDG_G_ELE_V2<"v2.f32 \t{{$dst1, $dst2}}, [$src];", Float32Regs>;
1441defm INT_PTX_LDG_G_v2i64_ELE
1442 : VLDG_G_ELE_V2<"v2.u64 \t{{$dst1, $dst2}}, [$src];", Int64Regs>;
1443defm INT_PTX_LDG_G_v2f64_ELE
1444 : VLDG_G_ELE_V2<"v2.f64 \t{{$dst1, $dst2}}, [$src];", Float64Regs>;
1445defm INT_PTX_LDG_G_v4i8_ELE
1446 : VLDG_G_ELE_V4<"v4.u8 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
1447defm INT_PTX_LDG_G_v4i16_ELE
1448 : VLDG_G_ELE_V4<"v4.u16 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int16Regs>;
1449defm INT_PTX_LDG_G_v4i32_ELE
1450 : VLDG_G_ELE_V4<"v4.u32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Int32Regs>;
1451defm INT_PTX_LDG_G_v4f32_ELE
1452 : VLDG_G_ELE_V4<"v4.f32 \t{{$dst1, $dst2, $dst3, $dst4}}, [$src];", Float32Regs>;
Justin Holewinskiae556d32012-05-04 20:18:50 +0000[diff] [blame]1453
1454
1455multiclass NG_TO_G<string Str, Intrinsic Intrin> {
1456 def _yes : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1457 !strconcat("cvta.", !strconcat(Str, ".u32 \t$result, $src;")),
1458 [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>,
1459 Requires<[hasGenericLdSt]>;
1460 def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1461 !strconcat("cvta.", !strconcat(Str, ".u64 \t$result, $src;")),
1462 [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>,
1463 Requires<[hasGenericLdSt]>;
1464
1465// @TODO: Are these actually needed? I believe global addresses will be copied
1466// to register values anyway.
1467 /*def __addr_yes : NVPTXInst<(outs Int32Regs:$result), (ins imemAny:$src),
1468 !strconcat("cvta.", !strconcat(Str, ".u32 \t$result, $src;")),
1469 [(set Int32Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>,
1470 Requires<[hasGenericLdSt]>;
1471 def __addr_yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins imemAny:$src),
1472 !strconcat("cvta.", !strconcat(Str, ".u64 \t$result, $src;")),
1473 [(set Int64Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>,
1474 Requires<[hasGenericLdSt]>;*/
1475
1476 def _no : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1477 "mov.u32 \t$result, $src;",
1478 [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>;
1479 def _no_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1480 "mov.u64 \t$result, $src;",
1481 [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>;
1482
1483// @TODO: Are these actually needed? I believe global addresses will be copied
1484// to register values anyway.
1485 /*def _addr_no : NVPTXInst<(outs Int32Regs:$result), (ins imem:$src),
1486 "mov.u32 \t$result, $src;",
1487 [(set Int32Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>;
1488 def _addr_no_64 : NVPTXInst<(outs Int64Regs:$result), (ins imem:$src),
1489 "mov.u64 \t$result, $src;",
1490 [(set Int64Regs:$result, (Intrin (Wrapper tglobaladdr:$src)))]>;*/
1491}
1492
1493multiclass G_TO_NG<string Str, Intrinsic Intrin> {
1494 def _yes : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1495 !strconcat("cvta.to.", !strconcat(Str, ".u32 \t$result, $src;")),
1496 [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>,
1497 Requires<[hasGenericLdSt]>;
1498 def _yes_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1499 !strconcat("cvta.to.", !strconcat(Str, ".u64 \t$result, $src;")),
1500 [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>,
1501 Requires<[hasGenericLdSt]>;
1502 def _no : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1503 "mov.u32 \t$result, $src;",
1504 [(set Int32Regs:$result, (Intrin Int32Regs:$src))]>;
1505 def _no_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1506 "mov.u64 \t$result, $src;",
1507 [(set Int64Regs:$result, (Intrin Int64Regs:$src))]>;
1508}
1509
1510defm cvta_local : NG_TO_G<"local", int_nvvm_ptr_local_to_gen>;
1511defm cvta_shared : NG_TO_G<"shared", int_nvvm_ptr_shared_to_gen>;
1512defm cvta_global : NG_TO_G<"global", int_nvvm_ptr_global_to_gen>;
1513
1514defm cvta_to_local : G_TO_NG<"local", int_nvvm_ptr_gen_to_local>;
1515defm cvta_to_shared : G_TO_NG<"shared", int_nvvm_ptr_gen_to_shared>;
1516defm cvta_to_global : G_TO_NG<"global", int_nvvm_ptr_gen_to_global>;
1517
1518def cvta_const : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1519 "mov.u32 \t$result, $src;",
1520 [(set Int32Regs:$result, (int_nvvm_ptr_constant_to_gen Int32Regs:$src))]>;
1521def cvta_const_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1522 "mov.u64 \t$result, $src;",
1523 [(set Int64Regs:$result, (int_nvvm_ptr_constant_to_gen Int64Regs:$src))]>;
1524
1525
1526
1527// @TODO: Revisit this. There is a type
1528// contradiction between iPTRAny and iPTR for the def.
1529/*def cvta_const_addr : NVPTXInst<(outs Int32Regs:$result), (ins imemAny:$src),
1530 "mov.u32 \t$result, $src;",
1531 [(set Int32Regs:$result, (int_nvvm_ptr_constant_to_gen
1532 (Wrapper tglobaladdr:$src)))]>;
1533def cvta_const_addr_64 : NVPTXInst<(outs Int64Regs:$result), (ins imemAny:$src),
1534 "mov.u64 \t$result, $src;",
1535 [(set Int64Regs:$result, (int_nvvm_ptr_constant_to_gen
1536 (Wrapper tglobaladdr:$src)))]>;*/
1537
1538
1539def cvta_to_const : NVPTXInst<(outs Int32Regs:$result), (ins Int32Regs:$src),
1540 "mov.u32 \t$result, $src;",
1541 [(set Int32Regs:$result, (int_nvvm_ptr_gen_to_constant Int32Regs:$src))]>;
1542def cvta_to_const_64 : NVPTXInst<(outs Int64Regs:$result), (ins Int64Regs:$src),
1543 "mov.u64 \t$result, $src;",
1544 [(set Int64Regs:$result, (int_nvvm_ptr_gen_to_constant Int64Regs:$src))]>;
1545
1546
1547// nvvm.ptr.gen.to.param
1548def nvvm_ptr_gen_to_param : NVPTXInst<(outs Int32Regs:$result),
1549 (ins Int32Regs:$src),
1550 "mov.u32 \t$result, $src;",
1551 [(set Int32Regs:$result,
1552 (int_nvvm_ptr_gen_to_param Int32Regs:$src))]>;
1553def nvvm_ptr_gen_to_param_64 : NVPTXInst<(outs Int64Regs:$result),
1554 (ins Int64Regs:$src),
1555 "mov.u64 \t$result, $src;",
1556 [(set Int64Regs:$result,
1557 (int_nvvm_ptr_gen_to_param Int64Regs:$src))]>;
1558
1559
1560// nvvm.move intrinsicc
1561def nvvm_move_i8 : NVPTXInst<(outs Int8Regs:$r), (ins Int8Regs:$s),
1562 "mov.b16 \t$r, $s;",
1563 [(set Int8Regs:$r,
1564 (int_nvvm_move_i8 Int8Regs:$s))]>;
1565def nvvm_move_i16 : NVPTXInst<(outs Int16Regs:$r), (ins Int16Regs:$s),
1566 "mov.b16 \t$r, $s;",
1567 [(set Int16Regs:$r,
1568 (int_nvvm_move_i16 Int16Regs:$s))]>;
1569def nvvm_move_i32 : NVPTXInst<(outs Int32Regs:$r), (ins Int32Regs:$s),
1570 "mov.b32 \t$r, $s;",
1571 [(set Int32Regs:$r,
1572 (int_nvvm_move_i32 Int32Regs:$s))]>;
1573def nvvm_move_i64 : NVPTXInst<(outs Int64Regs:$r), (ins Int64Regs:$s),
1574 "mov.b64 \t$r, $s;",
1575 [(set Int64Regs:$r,
1576 (int_nvvm_move_i64 Int64Regs:$s))]>;
1577def nvvm_move_float : NVPTXInst<(outs Float32Regs:$r), (ins Float32Regs:$s),
1578 "mov.f32 \t$r, $s;",
1579 [(set Float32Regs:$r,
1580 (int_nvvm_move_float Float32Regs:$s))]>;
1581def nvvm_move_double : NVPTXInst<(outs Float64Regs:$r), (ins Float64Regs:$s),
1582 "mov.f64 \t$r, $s;",
1583 [(set Float64Regs:$r,
1584 (int_nvvm_move_double Float64Regs:$s))]>;
1585def nvvm_move_ptr32 : NVPTXInst<(outs Int32Regs:$r), (ins Int32Regs:$s),
1586 "mov.u32 \t$r, $s;",
1587 [(set Int32Regs:$r,
1588 (int_nvvm_move_ptr Int32Regs:$s))]>;
1589def nvvm_move_ptr64 : NVPTXInst<(outs Int64Regs:$r), (ins Int64Regs:$s),
1590 "mov.u64 \t$r, $s;",
1591 [(set Int64Regs:$r,
1592 (int_nvvm_move_ptr Int64Regs:$s))]>;
1593
1594// @TODO: Are these actually needed, or will we always just see symbols
1595// copied to registers first?
1596/*def nvvm_move_sym32 : NVPTXInst<(outs Int32Regs:$r), (ins imem:$s),
1597 "mov.u32 \t$r, $s;",
1598 [(set Int32Regs:$r,
1599 (int_nvvm_move_ptr texternalsym:$s))]>;
1600def nvvm_move_sym64 : NVPTXInst<(outs Int64Regs:$r), (ins imem:$s),
1601 "mov.u64 \t$r, $s;",
1602 [(set Int64Regs:$r,
1603 (int_nvvm_move_ptr texternalsym:$s))]>;*/
1604
1605
1606// MoveParam %r1, param
1607// ptr_local_to_gen %r2, %r1
1608// ptr_gen_to_local %r3, %r2
1609// ->
1610// mov %r1, param
1611
1612// @TODO: Revisit this. There is a type
1613// contradiction between iPTRAny and iPTR for the addr defs, so the move_sym
1614// instructions are not currently defined. However, we can use the ptr
1615// variants and the asm printer will do the right thing.
1616def : Pat<(i64 (int_nvvm_ptr_gen_to_local (int_nvvm_ptr_local_to_gen
1617 (MoveParam texternalsym:$src)))),
1618 (nvvm_move_ptr64 texternalsym:$src)>;
1619def : Pat<(i32 (int_nvvm_ptr_gen_to_local (int_nvvm_ptr_local_to_gen
1620 (MoveParam texternalsym:$src)))),
1621 (nvvm_move_ptr32 texternalsym:$src)>;
1622
1623
1624//-----------------------------------
1625// Compiler Error Warn
1626// - Just ignore them in codegen
1627//-----------------------------------
1628
1629def INT_NVVM_COMPILER_WARN_32 : NVPTXInst<(outs), (ins Int32Regs:$a),
1630 "// llvm.nvvm.compiler.warn()",
1631 [(int_nvvm_compiler_warn Int32Regs:$a)]>;
1632def INT_NVVM_COMPILER_WARN_64 : NVPTXInst<(outs), (ins Int64Regs:$a),
1633 "// llvm.nvvm.compiler.warn()",
1634 [(int_nvvm_compiler_warn Int64Regs:$a)]>;
1635def INT_NVVM_COMPILER_ERROR_32 : NVPTXInst<(outs), (ins Int32Regs:$a),
1636 "// llvm.nvvm.compiler.error()",
1637 [(int_nvvm_compiler_error Int32Regs:$a)]>;
1638def INT_NVVM_COMPILER_ERROR_64 : NVPTXInst<(outs), (ins Int64Regs:$a),
1639 "// llvm.nvvm.compiler.error()",
1640 [(int_nvvm_compiler_error Int64Regs:$a)]>;
1641
1642
1643
1644//===-- Old PTX Back-end Intrinsics ---------------------------------------===//
1645
1646// These intrinsics are handled to retain compatibility with the old backend.
1647
1648// PTX Special Purpose Register Accessor Intrinsics
1649
1650class PTX_READ_SPECIAL_REGISTER_R64<string regname, Intrinsic intop>
1651 : NVPTXInst<(outs Int64Regs:$d), (ins),
1652 !strconcat(!strconcat("mov.u64\t$d, %", regname), ";"),
1653 [(set Int64Regs:$d, (intop))]>;
1654
1655class PTX_READ_SPECIAL_REGISTER_R32<string regname, Intrinsic intop>
1656 : NVPTXInst<(outs Int32Regs:$d), (ins),
1657 !strconcat(!strconcat("mov.u32\t$d, %", regname), ";"),
1658 [(set Int32Regs:$d, (intop))]>;
1659
1660// TODO Add read vector-version of special registers
1661
1662def PTX_READ_TID_X : PTX_READ_SPECIAL_REGISTER_R32<"tid.x",
1663 int_ptx_read_tid_x>;
1664def PTX_READ_TID_Y : PTX_READ_SPECIAL_REGISTER_R32<"tid.y",
1665 int_ptx_read_tid_y>;
1666def PTX_READ_TID_Z : PTX_READ_SPECIAL_REGISTER_R32<"tid.z",
1667 int_ptx_read_tid_z>;
1668def PTX_READ_TID_W : PTX_READ_SPECIAL_REGISTER_R32<"tid.w",
1669 int_ptx_read_tid_w>;
1670
1671def PTX_READ_NTID_X : PTX_READ_SPECIAL_REGISTER_R32<"ntid.x",
1672 int_ptx_read_ntid_x>;
1673def PTX_READ_NTID_Y : PTX_READ_SPECIAL_REGISTER_R32<"ntid.y",
1674 int_ptx_read_ntid_y>;
1675def PTX_READ_NTID_Z : PTX_READ_SPECIAL_REGISTER_R32<"ntid.z",
1676 int_ptx_read_ntid_z>;
1677def PTX_READ_NTID_W : PTX_READ_SPECIAL_REGISTER_R32<"ntid.w",
1678 int_ptx_read_ntid_w>;
1679
1680def PTX_READ_LANEID : PTX_READ_SPECIAL_REGISTER_R32<"laneid",
1681 int_ptx_read_laneid>;
1682def PTX_READ_WARPID : PTX_READ_SPECIAL_REGISTER_R32<"warpid",
1683 int_ptx_read_warpid>;
1684def PTX_READ_NWARPID : PTX_READ_SPECIAL_REGISTER_R32<"nwarpid",
1685 int_ptx_read_nwarpid>;
1686
1687def PTX_READ_CTAID_X : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.x",
1688 int_ptx_read_ctaid_x>;
1689def PTX_READ_CTAID_Y : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.y",
1690 int_ptx_read_ctaid_y>;
1691def PTX_READ_CTAID_Z : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.z",
1692 int_ptx_read_ctaid_z>;
1693def PTX_READ_CTAID_W : PTX_READ_SPECIAL_REGISTER_R32<"ctaid.w",
1694 int_ptx_read_ctaid_w>;
1695
1696def PTX_READ_NCTAID_X : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.x",
1697 int_ptx_read_nctaid_x>;
1698def PTX_READ_NCTAID_Y : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.y",
1699 int_ptx_read_nctaid_y>;
1700def PTX_READ_NCTAID_Z : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.z",
1701 int_ptx_read_nctaid_z>;
1702def PTX_READ_NCTAID_W : PTX_READ_SPECIAL_REGISTER_R32<"nctaid.w",
1703 int_ptx_read_nctaid_w>;
1704
1705def PTX_READ_SMID : PTX_READ_SPECIAL_REGISTER_R32<"smid",
1706 int_ptx_read_smid>;
1707def PTX_READ_NSMID : PTX_READ_SPECIAL_REGISTER_R32<"nsmid",
1708 int_ptx_read_nsmid>;
1709def PTX_READ_GRIDID : PTX_READ_SPECIAL_REGISTER_R32<"gridid",
1710 int_ptx_read_gridid>;
1711
1712def PTX_READ_LANEMASK_EQ
1713 : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_eq", int_ptx_read_lanemask_eq>;
1714def PTX_READ_LANEMASK_LE
1715 : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_le", int_ptx_read_lanemask_le>;
1716def PTX_READ_LANEMASK_LT
1717 : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_lt", int_ptx_read_lanemask_lt>;
1718def PTX_READ_LANEMASK_GE
1719 : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_ge", int_ptx_read_lanemask_ge>;
1720def PTX_READ_LANEMASK_GT
1721 : PTX_READ_SPECIAL_REGISTER_R32<"lanemask_gt", int_ptx_read_lanemask_gt>;
1722
1723def PTX_READ_CLOCK
1724 : PTX_READ_SPECIAL_REGISTER_R32<"clock", int_ptx_read_clock>;
1725def PTX_READ_CLOCK64
1726 : PTX_READ_SPECIAL_REGISTER_R64<"clock64", int_ptx_read_clock64>;
1727
1728def PTX_READ_PM0 : PTX_READ_SPECIAL_REGISTER_R32<"pm0", int_ptx_read_pm0>;
1729def PTX_READ_PM1 : PTX_READ_SPECIAL_REGISTER_R32<"pm1", int_ptx_read_pm1>;
1730def PTX_READ_PM2 : PTX_READ_SPECIAL_REGISTER_R32<"pm2", int_ptx_read_pm2>;
1731def PTX_READ_PM3 : PTX_READ_SPECIAL_REGISTER_R32<"pm3", int_ptx_read_pm3>;
1732
1733// PTX Parallel Synchronization and Communication Intrinsics
1734
1735def PTX_BAR_SYNC : NVPTXInst<(outs), (ins i32imm:$i), "bar.sync\t$i;",
1736 [(int_ptx_bar_sync imm:$i)]>;