18 *
19 * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
20 * or visit www.oracle.com if you need additional information or have any
21 * questions.
22 */
23 package com.oracle.graal.asm.ptx;
24
25 import com.oracle.graal.api.code.*;
26
27 public class PTXAssembler extends AbstractPTXAssembler {
28
29 @SuppressWarnings("unused")
30 public PTXAssembler(TargetDescription target, RegisterConfig registerConfig) {
31 super(target);
32 }
33
34 public final void at() {
35 emitString("@%p" + " " + "");
36 }
37
38 public final void add_s16(Register d, Register a, Register b) {
39 emitString("add.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
40 }
41
42 public final void add_s32(Register d, Register a, Register b) {
43 emitString("add.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
44 }
45
46 public final void add_s64(Register d, Register a, Register b) {
47 emitString("add.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
48 }
49
50 public final void add_s16(Register d, Register a, short s16) {
51 emitString("add.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + "");
52 }
53
54 public final void add_s32(Register d, Register a, int s32) {
55 emitString("add.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + "");
56 }
57
58 public final void add_s64(Register d, Register a, long s64) {
59 emitString("add.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + "");
60 }
61
62 public final void add_u16(Register d, Register a, Register b) {
63 emitString("add.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
64 }
65
66 public final void add_u32(Register d, Register a, Register b) {
67 emitString("add.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
68 }
69
70 public final void add_u64(Register d, Register a, Register b) {
71 emitString("add.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
72 }
73
74 public final void add_u16(Register d, Register a, short u16) {
75 emitString("add.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + "");
76 }
77
78 public final void add_u32(Register d, Register a, int u32) {
79 emitString("add.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
80 }
81
106 public final void and_b16(Register d, Register a, short b16) {
107 emitString("and.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b16 + ";" + "");
108 }
109
110 public final void and_b32(Register d, Register a, int b32) {
111 emitString("and.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b32 + ";" + "");
112 }
113
114 public final void and_b64(Register d, Register a, long b64) {
115 emitString("and.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b64 + ";" + "");
116 }
117
118 public final void bra(String tgt) {
119 emitString("bra" + " " + tgt + ";" + "");
120 }
121
122 public final void bra_uni(String tgt) {
123 emitString("bra.uni" + " " + tgt + ";" + "");
124 }
125
126 public final void div_s16(Register d, Register a, Register b) {
127 emitString("div.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
128 }
129
130 public final void div_s32(Register d, Register a, Register b) {
131 emitString("div.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
132 }
133
134 public final void div_s64(Register d, Register a, Register b) {
135 emitString("div.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
136 }
137
138 public final void div_s16(Register d, Register a, short s16) {
139 emitString("div.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + "");
140 }
141
142 public final void div_s32(Register d, Register a, int s32) {
143 emitString("div.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + "");
144 }
145
146 public final void div_s64(Register d, Register a, long s64) {
147 emitString("div.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + "");
148 }
149
150 public final void div_u16(Register d, Register a, Register b) {
151 emitString("div.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
152 }
153
154 public final void div_u32(Register d, Register a, Register b) {
155 emitString("div.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
156 }
157
158 public final void div_u64(Register d, Register a, Register b) {
159 emitString("div.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
160 }
161
162 public final void div_u16(Register d, Register a, short u16) {
163 emitString("div.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + "");
164 }
165
166 public final void div_u32(Register d, Register a, int u32) {
167 emitString("div.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
168 }
169
238 public final void mov_b32(Register d, Register a) {
239 emitString("mov.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
240 }
241
242 public final void mov_b64(Register d, Register a) {
243 emitString("mov.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
244 }
245
246 public final void mov_u16(Register d, Register a) {
247 emitString("mov.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
248 }
249
250 public final void mov_u32(Register d, Register a) {
251 emitString("mov.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
252 }
253
254 public final void mov_u64(Register d, Register a) {
255 emitString("mov.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
256 }
257
258 public final void mov_s16(Register d, Register a) {
259 emitString("mov.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
260 }
261
262 public final void mov_s32(Register d, Register a) {
263 emitString("mov.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
264 }
265
266 public final void mov_s64(Register d, Register a) {
267 emitString("mov.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
268 }
269
270 public final void mov_f32(Register d, Register a) {
271 emitString("mov.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
272 }
273
274 public final void mov_f64(Register d, Register a) {
275 emitString("mov.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
276 }
277
302 public final void mov_s16(Register d, short s16) {
303 emitString("mov.s16" + " " + "%r" + d.encoding() + ", " + s16 + ";" + "");
304 }
305
306 public final void mov_s32(Register d, int s32) {
307 emitString("mov.s32" + " " + "%r" + d.encoding() + ", " + s32 + ";" + "");
308 }
309
310 public final void mov_s64(Register d, long s64) {
311 emitString("mov.s64" + " " + "%r" + d.encoding() + ", " + s64 + ";" + "");
312 }
313
314 public final void mov_f32(Register d, float f32) {
315 emitString("mov.f32" + " " + "%r" + d.encoding() + ", " + f32 + ";" + "");
316 }
317
318 public final void mov_f64(Register d, double f64) {
319 emitString("mov.f64" + " " + "%r" + d.encoding() + ", " + f64 + ";" + "");
320 }
321
322 public final void mul_s16(Register d, Register a, Register b) {
323 emitString("mul.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
324 }
325
326 public final void mul_s32(Register d, Register a, Register b) {
327 emitString("mul.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
328 }
329
330 public final void mul_s64(Register d, Register a, Register b) {
331 emitString("mul.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
332 }
333
334 public final void mul_s16(Register d, Register a, short s16) {
335 emitString("mul.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + "");
336 }
337
338 public final void mul_s32(Register d, Register a, int s32) {
339 emitString("mul.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + "");
340 }
341
342 public final void mul_s64(Register d, Register a, long s64) {
343 emitString("mul.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + "");
344 }
345
346 public final void mul_u16(Register d, Register a, Register b) {
347 emitString("mul.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
348 }
349
350 public final void mul_u32(Register d, Register a, Register b) {
351 emitString("mul.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
352 }
353
354 public final void mul_u64(Register d, Register a, Register b) {
355 emitString("mul.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
356 }
357
358 public final void mul_u16(Register d, Register a, short u16) {
359 emitString("mul.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + "");
360 }
361
362 public final void mul_u32(Register d, Register a, int u32) {
363 emitString("mul.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
364 }
365
366 public final void mul_u64(Register d, Register a, long u64) {
367 emitString("mul.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + "");
368 }
369
370 public final void neg_s16(Register d, Register a) {
371 emitString("neg.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
372 }
373
374 public final void neg_s32(Register d, Register a) {
375 emitString("neg.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
376 }
377
378 public final void neg_s64(Register d, Register a) {
379 emitString("neg.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
380 }
381
382 public final void popc_b32(Register d, Register a) {
383 emitString("popc.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
384 }
385
386 public final void popc_b64(Register d, Register a) {
387 emitString("popc.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
388 }
389
390 public final void rem_s16(Register d, Register a, Register b) {
391 emitString("rem.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
392 }
393
394 public final void rem_s32(Register d, Register a, Register b) {
395 emitString("rem.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
396 }
397
398 public final void rem_s64(Register d, Register a, Register b) {
399 emitString("rem.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
400 }
401
426 public final void rem_u16(Register d, Register a, short u16) {
427 emitString("rem.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + "");
428 }
429
430 public final void rem_u32(Register d, Register a, int u32) {
431 emitString("rem.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
432 }
433
434 public final void rem_u64(Register d, Register a, long u64) {
435 emitString("rem.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + "");
436 }
437
438 public final void ret() {
439 emitString("ret;" + " " + "");
440 }
441
442 public final void ret_uni() {
443 emitString("ret.uni;" + " " + "");
444 }
445
446 public final void setp_eq_s32(Register a, Register b) {
447 emitString("setp.eq.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
448 }
449
450 public final void setp_ne_s32(Register a, Register b) {
451 emitString("setp.ne.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
452 }
453
454 public final void setp_lt_s32(Register a, Register b) {
455 emitString("setp.lt.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
456 }
457
458 public final void setp_le_s32(Register a, Register b) {
459 emitString("setp.le.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
460 }
461
462 public final void setp_gt_s32(Register a, Register b) {
463 emitString("setp.gt.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
464 }
465
570 public final void setp_ne_u32(int u32, Register b) {
571 emitString("setp.ne.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + "");
572 }
573
574 public final void setp_lt_u32(int u32, Register b) {
575 emitString("setp.lt.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + "");
576 }
577
578 public final void setp_le_u32(int u32, Register b) {
579 emitString("setp.le.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + "");
580 }
581
582 public final void setp_gt_u32(int u32, Register b) {
583 emitString("setp.gt.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + "");
584 }
585
586 public final void setp_ge_u32(int u32, Register b) {
587 emitString("setp.ge.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + "");
588 }
589
590 public final void shr_s16(Register d, Register a, Register b) {
591 emitString("shr.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
592 }
593
594 public final void shr_s32(Register d, Register a, Register b) {
595 emitString("shr.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
596 }
597
598 public final void shr_s64(Register d, Register a, Register b) {
599 emitString("shr.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
600 }
601
602 public final void shr_s16(Register d, Register a, int u32) {
603 emitString("shr.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
604 }
605
606 public final void shr_s32(Register d, Register a, int u32) {
607 emitString("shr.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
608 }
609
614 public final void shr_u16(Register d, Register a, Register b) {
615 emitString("shr.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
616 }
617
618 public final void shr_u32(Register d, Register a, Register b) {
619 emitString("shr.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
620 }
621
622 public final void shr_u64(Register d, Register a, Register b) {
623 emitString("shr.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
624 }
625
626 public final void shr_u16(Register d, Register a, int u32) {
627 emitString("shr.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
628 }
629
630 public final void shr_u32(Register d, Register a, int u32) {
631 emitString("shr.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
632 }
633
634 public final void shr_u64(Register d, Register a, int u32) {
635 emitString("shr.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
636 }
637
638 public final void st_global_b8(Register a, long immOff, Register b) {
639 emitString("st.global.b8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
640 }
641
642 public final void st_global_b16(Register a, long immOff, Register b) {
643 emitString("st.global.b16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
644 }
645
646 public final void st_global_b32(Register a, long immOff, Register b) {
647 emitString("st.global.b32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
648 }
649
650 public final void st_global_b64(Register a, long immOff, Register b) {
651 emitString("st.global.b64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
652 }
653
654 public final void st_global_u8(Register a, long immOff, Register b) {
655 emitString("st.global.u8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
674 public final void st_global_s16(Register a, long immOff, Register b) {
675 emitString("st.global.s16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
676 }
677
678 public final void st_global_s32(Register a, long immOff, Register b) {
679 emitString("st.global.s32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
680 }
681
682 public final void st_global_s64(Register a, long immOff, Register b) {
683 emitString("st.global.s64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
684 }
685
686 public final void st_global_f32(Register a, long immOff, Register b) {
687 emitString("st.global.f32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
688 }
689
690 public final void st_global_f64(Register a, long immOff, Register b) {
691 emitString("st.global.f64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
692 }
693
694 public final void sub_s16(Register d, Register a, Register b) {
695 emitString("sub.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
696 }
697
698 public final void sub_s32(Register d, Register a, Register b) {
699 emitString("sub.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
700 }
701
702 public final void sub_s64(Register d, Register a, Register b) {
703 emitString("sub.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
704 }
705
706 public final void sub_s16(Register d, Register a, short s16) {
707 emitString("sub.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + "");
708 }
709
710 public final void sub_s32(Register d, Register a, int s32) {
711 emitString("sub.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + "");
712 }
713
714 public final void sub_s64(Register d, Register a, long s64) {
715 emitString("sub.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + "");
716 }
717
718 public final void sub_s16(Register d, short s16, Register b) {
719 emitString("sub.s16" + " " + "%r" + d.encoding() + ", " + s16 + ", %r" + b.encoding() + ";" + "");
720 }
721
722 public final void sub_s32(Register d, int s32, Register b) {
723 emitString("sub.s32" + " " + "%r" + d.encoding() + ", " + s32 + ", %r" + b.encoding() + ";" + "");
724 }
725
726 public final void sub_s64(Register d, long s64, Register b) {
727 emitString("sub.s64" + " " + "%r" + d.encoding() + ", " + s64 + ", %r" + b.encoding() + ";" + "");
728 }
729
730 public final void sub_sat_s32(Register d, Register a, Register b) {
731 emitString("sub.sat.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
732 }
733
734 public final void sub_sat_s32(Register d, Register a, int s32) {
735 emitString("sub.sat.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + "");
736 }
737
738 public final void sub_sat_s32(Register d, int s32, Register b) {
739 emitString("sub.sat.s32" + " " + "%r" + d.encoding() + ", " + s32 + ", %r" + b.encoding() + ";" + "");
740 }
741
742 @Override
743 public PTXAddress makeAddress(Register base, int displacement) {
744 return new PTXAddress(base, displacement);
745 }
746
747 @Override
748 public PTXAddress getPlaceholder() {
749 // TODO Auto-generated method stub
750 return null;
751 }
752 }
|
18 *
19 * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
20 * or visit www.oracle.com if you need additional information or have any
21 * questions.
22 */
23 package com.oracle.graal.asm.ptx;
24
25 import com.oracle.graal.api.code.*;
26
27 public class PTXAssembler extends AbstractPTXAssembler {
28
29 @SuppressWarnings("unused")
30 public PTXAssembler(TargetDescription target, RegisterConfig registerConfig) {
31 super(target);
32 }
33
34 public final void at() {
35 emitString("@%p" + " " + "");
36 }
37
38 public final void add_f32(Register d, Register a, Register b) {
39 emitString("add.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
40 }
41
42 public final void add_f64(Register d, Register a, Register b) {
43 emitString("add.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
44 }
45
46 public final void add_s16(Register d, Register a, Register b) {
47 emitString("add.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
48 }
49
50 public final void add_s32(Register d, Register a, Register b) {
51 emitString("add.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
52 }
53
54 public final void add_s64(Register d, Register a, Register b) {
55 emitString("add.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
56 }
57
58 public final void add_s16(Register d, Register a, short s16) {
59 emitString("add.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + "");
60 }
61
62 public final void add_s32(Register d, Register a, int s32) {
63 emitString("add.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + "");
64 }
65
66 public final void add_s64(Register d, Register a, long s64) {
67 emitString("add.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + "");
68 }
69
70 public final void add_f32(Register d, Register a, float f32) {
71 emitString("add.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f32 + ";" + "");
72 }
73
74 public final void add_f64(Register d, Register a, double f64) {
75 emitString("add.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f64 + ";" + "");
76 }
77
78 public final void add_u16(Register d, Register a, Register b) {
79 emitString("add.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
80 }
81
82 public final void add_u32(Register d, Register a, Register b) {
83 emitString("add.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
84 }
85
86 public final void add_u64(Register d, Register a, Register b) {
87 emitString("add.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
88 }
89
90 public final void add_u16(Register d, Register a, short u16) {
91 emitString("add.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + "");
92 }
93
94 public final void add_u32(Register d, Register a, int u32) {
95 emitString("add.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
96 }
97
122 public final void and_b16(Register d, Register a, short b16) {
123 emitString("and.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b16 + ";" + "");
124 }
125
126 public final void and_b32(Register d, Register a, int b32) {
127 emitString("and.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b32 + ";" + "");
128 }
129
130 public final void and_b64(Register d, Register a, long b64) {
131 emitString("and.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b64 + ";" + "");
132 }
133
134 public final void bra(String tgt) {
135 emitString("bra" + " " + tgt + ";" + "");
136 }
137
138 public final void bra_uni(String tgt) {
139 emitString("bra.uni" + " " + tgt + ";" + "");
140 }
141
142 public final void cvt_s32_f32(Register d, Register a) {
143 emitString("cvt.s32.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
144 }
145
146 public final void cvt_s64_f32(Register d, Register a) {
147 emitString("cvt.s64.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
148 }
149
150 public final void cvt_f64_f32(Register d, Register a) {
151 emitString("cvt.f64.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
152 }
153
154 public final void cvt_f32_f64(Register d, Register a) {
155 emitString("cvt.f32.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
156 }
157
158 public final void cvt_s32_f64(Register d, Register a) {
159 emitString("cvt.s32.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
160 }
161
162 public final void cvt_s64_f64(Register d, Register a) {
163 emitString("cvt.s64.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
164 }
165
166 public final void cvt_f32_s32(Register d, Register a) {
167 emitString("cvt.f32.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
168 }
169
170 public final void cvt_f64_s32(Register d, Register a) {
171 emitString("cvt.f64.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
172 }
173
174 public final void cvt_s8_s32(Register d, Register a) {
175 emitString("cvt.s8.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
176 }
177
178 public final void cvt_b16_s32(Register d, Register a) {
179 emitString("cvt.b16.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
180 }
181
182 public final void cvt_s64_s32(Register d, Register a) {
183 emitString("cvt.s64.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
184 }
185
186 public final void cvt_s32_s64(Register d, Register a) {
187 emitString("cvt.s32.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
188 }
189
190 public final void div_f32(Register d, Register a, Register b) {
191 emitString("div.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
192 }
193
194 public final void div_f64(Register d, Register a, Register b) {
195 emitString("div.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
196 }
197
198 public final void div_s16(Register d, Register a, Register b) {
199 emitString("div.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
200 }
201
202 public final void div_s32(Register d, Register a, Register b) {
203 emitString("div.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
204 }
205
206 public final void div_s64(Register d, Register a, Register b) {
207 emitString("div.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
208 }
209
210 public final void div_s16(Register d, Register a, short s16) {
211 emitString("div.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + "");
212 }
213
214 public final void div_s32(Register d, Register a, int s32) {
215 emitString("div.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + "");
216 }
217
218 public final void div_s32(Register d, int s32, Register b) {
219 emitString("div.s32" + " " + "%r" + d.encoding() + ", " + s32 + ", %r" + b.encoding() + ";" + "");
220 }
221
222 public final void div_f32(Register d, float f32, Register b) {
223 emitString("div.f32" + " " + "%r" + d.encoding() + ", " + f32 + ", %r" + b.encoding() + ";" + "");
224 }
225
226 public final void div_f64(Register d, double f64, Register b) {
227 emitString("div.f64" + " " + "%r" + d.encoding() + ", " + f64 + ", %r" + b.encoding() + ";" + "");
228 }
229
230 public final void div_s64(Register d, Register a, long s64) {
231 emitString("div.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + "");
232 }
233
234 public final void div_f32(Register d, Register a, float f32) {
235 emitString("div.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f32 + ";" + "");
236 }
237
238 public final void div_f64(Register d, Register a, double f64) {
239 emitString("div.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f64 + ";" + "");
240 }
241
242 public final void div_u16(Register d, Register a, Register b) {
243 emitString("div.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
244 }
245
246 public final void div_u32(Register d, Register a, Register b) {
247 emitString("div.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
248 }
249
250 public final void div_u64(Register d, Register a, Register b) {
251 emitString("div.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
252 }
253
254 public final void div_u16(Register d, Register a, short u16) {
255 emitString("div.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + "");
256 }
257
258 public final void div_u32(Register d, Register a, int u32) {
259 emitString("div.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
260 }
261
330 public final void mov_b32(Register d, Register a) {
331 emitString("mov.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
332 }
333
334 public final void mov_b64(Register d, Register a) {
335 emitString("mov.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
336 }
337
338 public final void mov_u16(Register d, Register a) {
339 emitString("mov.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
340 }
341
342 public final void mov_u32(Register d, Register a) {
343 emitString("mov.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
344 }
345
346 public final void mov_u64(Register d, Register a) {
347 emitString("mov.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
348 }
349
350 public final void mov_u64(Register d, AbstractAddress a) {
351 // emitString("mov.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
352 }
353
354 public final void mov_s16(Register d, Register a) {
355 emitString("mov.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
356 }
357
358 public final void mov_s32(Register d, Register a) {
359 emitString("mov.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
360 }
361
362 public final void mov_s64(Register d, Register a) {
363 emitString("mov.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
364 }
365
366 public final void mov_f32(Register d, Register a) {
367 emitString("mov.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
368 }
369
370 public final void mov_f64(Register d, Register a) {
371 emitString("mov.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
372 }
373
398 public final void mov_s16(Register d, short s16) {
399 emitString("mov.s16" + " " + "%r" + d.encoding() + ", " + s16 + ";" + "");
400 }
401
402 public final void mov_s32(Register d, int s32) {
403 emitString("mov.s32" + " " + "%r" + d.encoding() + ", " + s32 + ";" + "");
404 }
405
406 public final void mov_s64(Register d, long s64) {
407 emitString("mov.s64" + " " + "%r" + d.encoding() + ", " + s64 + ";" + "");
408 }
409
410 public final void mov_f32(Register d, float f32) {
411 emitString("mov.f32" + " " + "%r" + d.encoding() + ", " + f32 + ";" + "");
412 }
413
414 public final void mov_f64(Register d, double f64) {
415 emitString("mov.f64" + " " + "%r" + d.encoding() + ", " + f64 + ";" + "");
416 }
417
418 public final void mul_f32(Register d, Register a, Register b) {
419 emitString("mul.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
420 }
421
422 public final void mul_f64(Register d, Register a, Register b) {
423 emitString("smul.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
424 }
425
426 public final void mul_s16(Register d, Register a, Register b) {
427 emitString("mul.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
428 }
429
430 public final void mul_s32(Register d, Register a, Register b) {
431 emitString("mul.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
432 }
433
434 public final void mul_s64(Register d, Register a, Register b) {
435 emitString("mul.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
436 }
437
438 public final void mul_s16(Register d, Register a, short s16) {
439 emitString("mul.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + "");
440 }
441
442 public final void mul_s32(Register d, Register a, int s32) {
443 emitString("mul.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + "");
444 }
445
446 public final void mul_s64(Register d, Register a, long s64) {
447 emitString("mul.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + "");
448 }
449
450 public final void mul_f32(Register d, Register a, float f32) {
451 emitString("mul.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f32 + ";" + "");
452 }
453
454 public final void mul_f64(Register d, Register a, double f64) {
455 emitString("mul.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f64 + ";" + "");
456 }
457
458 public final void mul_u16(Register d, Register a, Register b) {
459 emitString("mul.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
460 }
461
462 public final void mul_u32(Register d, Register a, Register b) {
463 emitString("mul.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
464 }
465
466 public final void mul_u64(Register d, Register a, Register b) {
467 emitString("mul.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
468 }
469
470 public final void mul_u16(Register d, Register a, short u16) {
471 emitString("mul.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + "");
472 }
473
474 public final void mul_u32(Register d, Register a, int u32) {
475 emitString("mul.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
476 }
477
478 public final void mul_u64(Register d, Register a, long u64) {
479 emitString("mul.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + "");
480 }
481
482 public final void neg_f32(Register d, Register a) {
483 emitString("neg.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
484 }
485
486 public final void neg_f64(Register d, Register a) {
487 emitString("neg.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
488 }
489
490 public final void neg_s16(Register d, Register a) {
491 emitString("neg.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
492 }
493
494 public final void neg_s32(Register d, Register a) {
495 emitString("neg.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
496 }
497
498 public final void neg_s64(Register d, Register a) {
499 emitString("neg.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
500 }
501
502 public final void not_s16(Register d, Register a) {
503 emitString("not.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
504 }
505
506 public final void not_s32(Register d, Register a) {
507 emitString("not.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
508 }
509
510 public final void not_s64(Register d, Register a) {
511 emitString("not.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
512 }
513
514 public final void or_b16(Register d, Register a, Register b) {
515 emitString("or.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
516 }
517
518 public final void or_b32(Register d, Register a, Register b) {
519 emitString("or.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
520 }
521
522 public final void or_b64(Register d, Register a, Register b) {
523 emitString("or.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
524 }
525
526 public final void or_b16(Register d, Register a, short b16) {
527 emitString("or.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b16 + ";" + "");
528 }
529
530 public final void or_b32(Register d, Register a, int b32) {
531 emitString("or.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b32 + ";" + "");
532 }
533
534 public final void or_b64(Register d, Register a, long b64) {
535 emitString("or.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b64 + ";" + "");
536 }
537
538 public final void popc_b32(Register d, Register a) {
539 emitString("popc.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
540 }
541
542 public final void popc_b64(Register d, Register a) {
543 emitString("popc.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
544 }
545
546 public final void rem_s16(Register d, Register a, Register b) {
547 emitString("rem.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
548 }
549
550 public final void rem_s32(Register d, Register a, Register b) {
551 emitString("rem.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
552 }
553
554 public final void rem_s64(Register d, Register a, Register b) {
555 emitString("rem.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
556 }
557
582 public final void rem_u16(Register d, Register a, short u16) {
583 emitString("rem.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + "");
584 }
585
586 public final void rem_u32(Register d, Register a, int u32) {
587 emitString("rem.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
588 }
589
590 public final void rem_u64(Register d, Register a, long u64) {
591 emitString("rem.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + "");
592 }
593
594 public final void ret() {
595 emitString("ret;" + " " + "");
596 }
597
598 public final void ret_uni() {
599 emitString("ret.uni;" + " " + "");
600 }
601
602 public final void setp_eq_f32(Register a, Register b) {
603 emitString("setp.eq.f32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
604 }
605
606 public final void setp_ne_f32(Register a, Register b) {
607 emitString("setp.ne.f32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
608 }
609
610 public final void setp_lt_f32(Register a, Register b) {
611 emitString("setp.lt.f32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
612 }
613
614 public final void setp_le_f32(Register a, Register b) {
615 emitString("setp.le.f32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
616 }
617
618 public final void setp_gt_f32(Register a, Register b) {
619 emitString("setp.gt.f32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
620 }
621
622 public final void setp_ge_f32(Register a, Register b) {
623 emitString("setp.ge.f32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
624 }
625
626 public final void setp_eq_f32(float f32, Register b) {
627 emitString("setp.eq.f32" + " " + "%p" + ", " + f32 + ", %r" + b.encoding() + ";" + "");
628 }
629
630 public final void setp_ne_f32(float f32, Register b) {
631 emitString("setp.ne.f32" + " " + "%p" + ", " + f32 + ", %r" + b.encoding() + ";" + "");
632 }
633
634 public final void setp_lt_f32(float f32, Register b) {
635 emitString("setp.lt.f32" + " " + "%p" + ", " + f32 + ", %r" + b.encoding() + ";" + "");
636 }
637
638 public final void setp_le_f32(float f32, Register b) {
639 emitString("setp.le.f32" + " " + "%p" + ", " + f32 + ", %r" + b.encoding() + ";" + "");
640 }
641
642 public final void setp_gt_f32(float f32, Register b) {
643 emitString("setp.gt.f32" + " " + "%p" + ", " + f32 + ", %r" + b.encoding() + ";" + "");
644 }
645
646 public final void setp_ge_f32(float f32, Register b) {
647 emitString("setp.ge.f32" + " " + "%p" + ", " + f32 + ", %r" + b.encoding() + ";" + "");
648 }
649
650 public final void setp_eq_f64(double f64, Register b) {
651 emitString("setp.eq.f64" + " " + "%p" + ", " + f64 + ", %r" + b.encoding() + ";" + "");
652 }
653
654 public final void setp_ne_f64(double f64, Register b) {
655 emitString("setp.ne.f64" + " " + "%p" + ", " + f64 + ", %r" + b.encoding() + ";" + "");
656 }
657
658 public final void setp_lt_f64(double f64, Register b) {
659 emitString("setp.lt.f64" + " " + "%p" + ", " + f64 + ", %r" + b.encoding() + ";" + "");
660 }
661
662 public final void setp_le_f64(double f64, Register b) {
663 emitString("setp.le.f64" + " " + "%p" + ", " + f64 + ", %r" + b.encoding() + ";" + "");
664 }
665
666 public final void setp_gt_f64(double f64, Register b) {
667 emitString("setp.gt.f64" + " " + "%p" + ", " + f64 + ", %r" + b.encoding() + ";" + "");
668 }
669
670 public final void setp_ge_f64(double f64, Register b) {
671 emitString("setp.ge.f64" + " " + "%p" + ", " + f64 + ", %r" + b.encoding() + ";" + "");
672 }
673
674 public final void setp_eq_s32(Register a, Register b) {
675 emitString("setp.eq.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
676 }
677
678 public final void setp_ne_s32(Register a, Register b) {
679 emitString("setp.ne.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
680 }
681
682 public final void setp_lt_s32(Register a, Register b) {
683 emitString("setp.lt.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
684 }
685
686 public final void setp_le_s32(Register a, Register b) {
687 emitString("setp.le.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
688 }
689
690 public final void setp_gt_s32(Register a, Register b) {
691 emitString("setp.gt.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
692 }
693
798 public final void setp_ne_u32(int u32, Register b) {
799 emitString("setp.ne.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + "");
800 }
801
802 public final void setp_lt_u32(int u32, Register b) {
803 emitString("setp.lt.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + "");
804 }
805
806 public final void setp_le_u32(int u32, Register b) {
807 emitString("setp.le.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + "");
808 }
809
810 public final void setp_gt_u32(int u32, Register b) {
811 emitString("setp.gt.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + "");
812 }
813
814 public final void setp_ge_u32(int u32, Register b) {
815 emitString("setp.ge.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + "");
816 }
817
818 public final void shl_s16(Register d, Register a, Register b) {
819 emitString("shl.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
820 }
821
822 public final void shl_s32(Register d, Register a, Register b) {
823 emitString("shl.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
824 }
825
826 public final void shl_s64(Register d, Register a, Register b) {
827 emitString("shl.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
828 }
829
830 public final void shl_s16(Register d, Register a, int u32) {
831 emitString("shl.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
832 }
833
834 public final void shl_s32(Register d, Register a, int u32) {
835 emitString("shl.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
836 }
837
838 public final void shl_s64(Register d, Register a, int u32) {
839 emitString("shl.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
840 }
841
842 public final void shl_u16(Register d, Register a, Register b) {
843 emitString("shl.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
844 }
845
846 public final void shl_u32(Register d, Register a, Register b) {
847 emitString("shl.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
848 }
849
850 public final void shl_u64(Register d, Register a, Register b) {
851 emitString("shl.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
852 }
853
854 public final void shl_u16(Register d, Register a, int u32) {
855 emitString("shl.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
856 }
857
858 public final void shl_u32(Register d, Register a, int u32) {
859 emitString("shl.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
860 }
861
862 public final void shl_u64(Register d, Register a, int u32) {
863 emitString("shl.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
864 }
865
866 public final void shr_s16(Register d, Register a, Register b) {
867 emitString("shr.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
868 }
869
870 public final void shr_s32(Register d, Register a, Register b) {
871 emitString("shr.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
872 }
873
874 public final void shr_s64(Register d, Register a, Register b) {
875 emitString("shr.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
876 }
877
878 public final void shr_s16(Register d, Register a, int u32) {
879 emitString("shr.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
880 }
881
882 public final void shr_s32(Register d, Register a, int u32) {
883 emitString("shr.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
884 }
885
890 public final void shr_u16(Register d, Register a, Register b) {
891 emitString("shr.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
892 }
893
894 public final void shr_u32(Register d, Register a, Register b) {
895 emitString("shr.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
896 }
897
898 public final void shr_u64(Register d, Register a, Register b) {
899 emitString("shr.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
900 }
901
902 public final void shr_u16(Register d, Register a, int u32) {
903 emitString("shr.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
904 }
905
906 public final void shr_u32(Register d, Register a, int u32) {
907 emitString("shr.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
908 }
909
910 public final void shr_u64(Register d, Register a, long u64) {
911 emitString("shr.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + "");
912 }
913
914 public final void st_global_b8(Register a, long immOff, Register b) {
915 emitString("st.global.b8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
916 }
917
918 public final void st_global_b16(Register a, long immOff, Register b) {
919 emitString("st.global.b16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
920 }
921
922 public final void st_global_b32(Register a, long immOff, Register b) {
923 emitString("st.global.b32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
924 }
925
926 public final void st_global_b64(Register a, long immOff, Register b) {
927 emitString("st.global.b64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
928 }
929
930 public final void st_global_u8(Register a, long immOff, Register b) {
931 emitString("st.global.u8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
950 public final void st_global_s16(Register a, long immOff, Register b) {
951 emitString("st.global.s16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
952 }
953
954 public final void st_global_s32(Register a, long immOff, Register b) {
955 emitString("st.global.s32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
956 }
957
958 public final void st_global_s64(Register a, long immOff, Register b) {
959 emitString("st.global.s64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
960 }
961
962 public final void st_global_f32(Register a, long immOff, Register b) {
963 emitString("st.global.f32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
964 }
965
966 public final void st_global_f64(Register a, long immOff, Register b) {
967 emitString("st.global.f64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
968 }
969
970 public final void sub_f32(Register d, Register a, Register b) {
971 emitString("sub.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
972 }
973
974 public final void sub_f64(Register d, Register a, Register b) {
975 emitString("sub.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
976 }
977
978 public final void sub_s16(Register d, Register a, Register b) {
979 emitString("sub.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
980 }
981
982 public final void sub_s32(Register d, Register a, Register b) {
983 emitString("sub.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
984 }
985
986 public final void sub_s64(Register d, Register a, Register b) {
987 emitString("sub.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
988 }
989
990 public final void sub_s16(Register d, Register a, short s16) {
991 emitString("sub.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + "");
992 }
993
994 public final void sub_s32(Register d, Register a, int s32) {
995 emitString("sub.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + "");
996 }
997
998 public final void sub_s64(Register d, Register a, long s64) {
999 emitString("sub.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + "");
1000 }
1001
1002 public final void sub_f32(Register d, Register a, float f32) {
1003 emitString("sub.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f32 + ";" + "");
1004 }
1005
1006 public final void sub_f64(Register d, Register a, double f64) {
1007 emitString("sub.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f64 + ";" + "");
1008 }
1009
1010 public final void sub_s16(Register d, short s16, Register b) {
1011 emitString("sub.s16" + " " + "%r" + d.encoding() + ", " + s16 + ", %r" + b.encoding() + ";" + "");
1012 }
1013
1014 public final void sub_s32(Register d, int s32, Register b) {
1015 emitString("sub.s32" + " " + "%r" + d.encoding() + ", " + s32 + ", %r" + b.encoding() + ";" + "");
1016 }
1017
1018 public final void sub_s64(Register d, long s64, Register b) {
1019 emitString("sub.s64" + " " + "%r" + d.encoding() + ", " + s64 + ", %r" + b.encoding() + ";" + "");
1020 }
1021
1022 public final void sub_f32(Register d, float f32, Register b) {
1023 emitString("sub.f32" + " " + "%r" + d.encoding() + ", %r" + b.encoding() + ", " + f32 + ";" + "");
1024 }
1025
1026 public final void sub_f64(Register d, double f64, Register b) {
1027 emitString("sub.f64" + " " + "%r" + d.encoding() + ", %r" + b.encoding() + ", " + f64 + ";" + "");
1028 }
1029
1030 public final void sub_sat_s32(Register d, Register a, Register b) {
1031 emitString("sub.sat.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
1032 }
1033
1034 public final void sub_sat_s32(Register d, Register a, int s32) {
1035 emitString("sub.sat.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + "");
1036 }
1037
1038 public final void sub_sat_s32(Register d, int s32, Register b) {
1039 emitString("sub.sat.s32" + " " + "%r" + d.encoding() + ", " + s32 + ", %r" + b.encoding() + ";" + "");
1040 }
1041
1042 public final void xor_b16(Register d, Register a, Register b) {
1043 emitString("xor.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
1044 }
1045
1046 public final void xor_b32(Register d, Register a, Register b) {
1047 emitString("xor.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
1048 }
1049
1050 public final void xor_b64(Register d, Register a, Register b) {
1051 emitString("xor.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
1052 }
1053
1054 public final void xor_b16(Register d, Register a, short b16) {
1055 emitString("xor.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b16 + ";" + "");
1056 }
1057
1058 public final void xor_b32(Register d, Register a, int b32) {
1059 emitString("xor.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b32 + ";" + "");
1060 }
1061
1062 public final void xor_b64(Register d, Register a, long b64) {
1063 emitString("xor.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b64 + ";" + "");
1064 }
1065
1066 @Override
1067 public PTXAddress makeAddress(Register base, int displacement) {
1068 return new PTXAddress(base, displacement);
1069 }
1070
1071 @Override
1072 public PTXAddress getPlaceholder() {
1073 // TODO Auto-generated method stub
1074 return null;
1075 }
1076 }
|