graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java
Print this page
rev 8661 : Graal PTX enhancements
rev 8662 : [mq]: ptx_switch
rev 8663 : [mq]: ptx_lookup_switch
@@ -24,19 +24,30 @@
import com.oracle.graal.api.code.*;
public class PTXAssembler extends AbstractPTXAssembler {
- @SuppressWarnings("unused")
public PTXAssembler(TargetDescription target, RegisterConfig registerConfig) {
super(target);
}
public final void at() {
emitString("@%p" + " " + "");
}
+ public final void atq() {
+ emitString("@%q" + " " + "");
+ }
+
+ public final void add_f32(Register d, Register a, Register b) {
+ emitString("add.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void add_f64(Register d, Register a, Register b) {
+ emitString("add.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
public final void add_s16(Register d, Register a, Register b) {
emitString("add.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
}
public final void add_s32(Register d, Register a, Register b) {
@@ -57,10 +68,18 @@
public final void add_s64(Register d, Register a, long s64) {
emitString("add.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + "");
}
+ public final void add_f32(Register d, Register a, float f32) {
+ emitString("add.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f32 + ";" + "");
+ }
+
+ public final void add_f64(Register d, Register a, double f64) {
+ emitString("add.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f64 + ";" + "");
+ }
+
public final void add_u16(Register d, Register a, Register b) {
emitString("add.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
}
public final void add_u32(Register d, Register a, Register b) {
@@ -121,10 +140,66 @@
public final void bra_uni(String tgt) {
emitString("bra.uni" + " " + tgt + ";" + "");
}
+ public final void cvt_s32_f32(Register d, Register a) {
+ emitString("cvt.s32.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
+ }
+
+ public final void cvt_s64_f32(Register d, Register a) {
+ emitString("cvt.s64.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
+ }
+
+ public final void cvt_f64_f32(Register d, Register a) {
+ emitString("cvt.f64.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
+ }
+
+ public final void cvt_f32_f64(Register d, Register a) {
+ emitString("cvt.f32.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
+ }
+
+ public final void cvt_s32_f64(Register d, Register a) {
+ emitString("cvt.s32.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
+ }
+
+ public final void cvt_s64_f64(Register d, Register a) {
+ emitString("cvt.s64.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
+ }
+
+ public final void cvt_f32_s32(Register d, Register a) {
+ emitString("cvt.f32.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
+ }
+
+ public final void cvt_f64_s32(Register d, Register a) {
+ emitString("cvt.f64.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
+ }
+
+ public final void cvt_s8_s32(Register d, Register a) {
+ emitString("cvt.s8.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
+ }
+
+ public final void cvt_b16_s32(Register d, Register a) {
+ emitString("cvt.b16.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
+ }
+
+ public final void cvt_s64_s32(Register d, Register a) {
+ emitString("cvt.s64.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
+ }
+
+ public final void cvt_s32_s64(Register d, Register a) {
+ emitString("cvt.s32.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
+ }
+
+ public final void div_f32(Register d, Register a, Register b) {
+ emitString("div.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void div_f64(Register d, Register a, Register b) {
+ emitString("div.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
public final void div_s16(Register d, Register a, Register b) {
emitString("div.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
}
public final void div_s32(Register d, Register a, Register b) {
@@ -141,14 +216,34 @@
public final void div_s32(Register d, Register a, int s32) {
emitString("div.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + "");
}
+ public final void div_s32(Register d, int s32, Register b) {
+ emitString("div.s32" + " " + "%r" + d.encoding() + ", " + s32 + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void div_f32(Register d, float f32, Register b) {
+ emitString("div.f32" + " " + "%r" + d.encoding() + ", " + f32 + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void div_f64(Register d, double f64, Register b) {
+ emitString("div.f64" + " " + "%r" + d.encoding() + ", " + f64 + ", %r" + b.encoding() + ";" + "");
+ }
+
public final void div_s64(Register d, Register a, long s64) {
emitString("div.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + "");
}
+ public final void div_f32(Register d, Register a, float f32) {
+ emitString("div.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f32 + ";" + "");
+ }
+
+ public final void div_f64(Register d, Register a, double f64) {
+ emitString("div.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f64 + ";" + "");
+ }
+
public final void div_u16(Register d, Register a, Register b) {
emitString("div.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
}
public final void div_u32(Register d, Register a, Register b) {
@@ -253,10 +348,14 @@
public final void mov_u64(Register d, Register a) {
emitString("mov.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
}
+ public final void mov_u64(Register d, AbstractAddress a) {
+ // emitString("mov.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
+ }
+
public final void mov_s16(Register d, Register a) {
emitString("mov.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
}
public final void mov_s32(Register d, Register a) {
@@ -317,10 +416,18 @@
public final void mov_f64(Register d, double f64) {
emitString("mov.f64" + " " + "%r" + d.encoding() + ", " + f64 + ";" + "");
}
+ public final void mul_f32(Register d, Register a, Register b) {
+ emitString("mul.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void mul_f64(Register d, Register a, Register b) {
+ emitString("smul.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
public final void mul_s16(Register d, Register a, Register b) {
emitString("mul.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
}
public final void mul_s32(Register d, Register a, Register b) {
@@ -341,10 +448,18 @@
public final void mul_s64(Register d, Register a, long s64) {
emitString("mul.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + "");
}
+ public final void mul_f32(Register d, Register a, float f32) {
+ emitString("mul.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f32 + ";" + "");
+ }
+
+ public final void mul_f64(Register d, Register a, double f64) {
+ emitString("mul.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f64 + ";" + "");
+ }
+
public final void mul_u16(Register d, Register a, Register b) {
emitString("mul.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
}
public final void mul_u32(Register d, Register a, Register b) {
@@ -365,10 +480,18 @@
public final void mul_u64(Register d, Register a, long u64) {
emitString("mul.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + "");
}
+ public final void neg_f32(Register d, Register a) {
+ emitString("neg.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
+ }
+
+ public final void neg_f64(Register d, Register a) {
+ emitString("neg.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
+ }
+
public final void neg_s16(Register d, Register a) {
emitString("neg.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
}
public final void neg_s32(Register d, Register a) {
@@ -377,10 +500,46 @@
public final void neg_s64(Register d, Register a) {
emitString("neg.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
}
+ public final void not_s16(Register d, Register a) {
+ emitString("not.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
+ }
+
+ public final void not_s32(Register d, Register a) {
+ emitString("not.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
+ }
+
+ public final void not_s64(Register d, Register a) {
+ emitString("not.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
+ }
+
+ public final void or_b16(Register d, Register a, Register b) {
+ emitString("or.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void or_b32(Register d, Register a, Register b) {
+ emitString("or.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void or_b64(Register d, Register a, Register b) {
+ emitString("or.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void or_b16(Register d, Register a, short b16) {
+ emitString("or.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b16 + ";" + "");
+ }
+
+ public final void or_b32(Register d, Register a, int b32) {
+ emitString("or.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b32 + ";" + "");
+ }
+
+ public final void or_b64(Register d, Register a, long b64) {
+ emitString("or.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b64 + ";" + "");
+ }
+
public final void popc_b32(Register d, Register a) {
emitString("popc.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
}
public final void popc_b64(Register d, Register a) {
@@ -441,10 +600,90 @@
public final void ret_uni() {
emitString("ret.uni;" + " " + "");
}
+ public final void setp_eq_f32(Register a, Register b) {
+ emitString("setp.eq.f32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_ne_f32(Register a, Register b) {
+ emitString("setp.ne.f32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_lt_f32(Register a, Register b) {
+ emitString("setp.lt.f32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_le_f32(Register a, Register b) {
+ emitString("setp.le.f32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_gt_f32(Register a, Register b) {
+ emitString("setp.gt.f32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_ge_f32(Register a, Register b) {
+ emitString("setp.ge.f32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_eq_f32(float f32, Register b) {
+ emitString("setp.eq.f32" + " " + "%p" + ", " + f32 + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_ne_f32(float f32, Register b) {
+ emitString("setp.ne.f32" + " " + "%p" + ", " + f32 + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_lt_f32(float f32, Register b) {
+ emitString("setp.lt.f32" + " " + "%p" + ", " + f32 + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_le_f32(float f32, Register b) {
+ emitString("setp.le.f32" + " " + "%p" + ", " + f32 + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_gt_f32(float f32, Register b) {
+ emitString("setp.gt.f32" + " " + "%p" + ", " + f32 + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_ge_f32(float f32, Register b) {
+ emitString("setp.ge.f32" + " " + "%p" + ", " + f32 + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_eq_f64(double f64, Register b) {
+ emitString("setp.eq.f64" + " " + "%p" + ", " + f64 + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_ne_f64(double f64, Register b) {
+ emitString("setp.ne.f64" + " " + "%p" + ", " + f64 + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_lt_f64(double f64, Register b) {
+ emitString("setp.lt.f64" + " " + "%p" + ", " + f64 + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_le_f64(double f64, Register b) {
+ emitString("setp.le.f64" + " " + "%p" + ", " + f64 + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_gt_f64(double f64, Register b) {
+ emitString("setp.gt.f64" + " " + "%p" + ", " + f64 + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_ge_f64(double f64, Register b) {
+ emitString("setp.ge.f64" + " " + "%p" + ", " + f64 + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_eq_s64(Register a, Register b) {
+ emitString("setp.eq.s64" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void setp_eq_s64(long s64, Register b) {
+ emitString("setp.eq.s64" + " " + "%p" + ", " + s64 + ", %r" + b.encoding() + ";" + "");
+ }
+
public final void setp_eq_s32(Register a, Register b) {
emitString("setp.eq.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
}
public final void setp_ne_s32(Register a, Register b) {
@@ -585,10 +824,58 @@
public final void setp_ge_u32(int u32, Register b) {
emitString("setp.ge.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + "");
}
+ public final void shl_s16(Register d, Register a, Register b) {
+ emitString("shl.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void shl_s32(Register d, Register a, Register b) {
+ emitString("shl.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void shl_s64(Register d, Register a, Register b) {
+ emitString("shl.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void shl_s16(Register d, Register a, int u32) {
+ emitString("shl.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
+ }
+
+ public final void shl_s32(Register d, Register a, int u32) {
+ emitString("shl.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
+ }
+
+ public final void shl_s64(Register d, Register a, int u32) {
+ emitString("shl.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
+ }
+
+ public final void shl_u16(Register d, Register a, Register b) {
+ emitString("shl.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void shl_u32(Register d, Register a, Register b) {
+ emitString("shl.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void shl_u64(Register d, Register a, Register b) {
+ emitString("shl.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void shl_u16(Register d, Register a, int u32) {
+ emitString("shl.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
+ }
+
+ public final void shl_u32(Register d, Register a, int u32) {
+ emitString("shl.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
+ }
+
+ public final void shl_u64(Register d, Register a, int u32) {
+ emitString("shl.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
+ }
+
public final void shr_s16(Register d, Register a, Register b) {
emitString("shr.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
}
public final void shr_s32(Register d, Register a, Register b) {
@@ -629,12 +916,12 @@
public final void shr_u32(Register d, Register a, int u32) {
emitString("shr.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
}
- public final void shr_u64(Register d, Register a, int u32) {
- emitString("shr.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
+ public final void shr_u64(Register d, Register a, long u64) {
+ emitString("shr.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + "");
}
public final void st_global_b8(Register a, long immOff, Register b) {
emitString("st.global.b8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
}
@@ -689,10 +976,18 @@
public final void st_global_f64(Register a, long immOff, Register b) {
emitString("st.global.f64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
}
+ public final void sub_f32(Register d, Register a, Register b) {
+ emitString("sub.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void sub_f64(Register d, Register a, Register b) {
+ emitString("sub.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
public final void sub_s16(Register d, Register a, Register b) {
emitString("sub.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
}
public final void sub_s32(Register d, Register a, Register b) {
@@ -709,14 +1004,26 @@
public final void sub_s32(Register d, Register a, int s32) {
emitString("sub.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + "");
}
+ public final void sub_s64(Register d, Register a, int s32) {
+ emitString("sub.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + "");
+ }
+
public final void sub_s64(Register d, Register a, long s64) {
emitString("sub.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + "");
}
+ public final void sub_f32(Register d, Register a, float f32) {
+ emitString("sub.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f32 + ";" + "");
+ }
+
+ public final void sub_f64(Register d, Register a, double f64) {
+ emitString("sub.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f64 + ";" + "");
+ }
+
public final void sub_s16(Register d, short s16, Register b) {
emitString("sub.s16" + " " + "%r" + d.encoding() + ", " + s16 + ", %r" + b.encoding() + ";" + "");
}
public final void sub_s32(Register d, int s32, Register b) {
@@ -725,10 +1032,18 @@
public final void sub_s64(Register d, long s64, Register b) {
emitString("sub.s64" + " " + "%r" + d.encoding() + ", " + s64 + ", %r" + b.encoding() + ";" + "");
}
+ public final void sub_f32(Register d, float f32, Register b) {
+ emitString("sub.f32" + " " + "%r" + d.encoding() + ", %r" + b.encoding() + ", " + f32 + ";" + "");
+ }
+
+ public final void sub_f64(Register d, double f64, Register b) {
+ emitString("sub.f64" + " " + "%r" + d.encoding() + ", %r" + b.encoding() + ", " + f64 + ";" + "");
+ }
+
public final void sub_sat_s32(Register d, Register a, Register b) {
emitString("sub.sat.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
}
public final void sub_sat_s32(Register d, Register a, int s32) {
@@ -737,10 +1052,34 @@
public final void sub_sat_s32(Register d, int s32, Register b) {
emitString("sub.sat.s32" + " " + "%r" + d.encoding() + ", " + s32 + ", %r" + b.encoding() + ";" + "");
}
+ public final void xor_b16(Register d, Register a, Register b) {
+ emitString("xor.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void xor_b32(Register d, Register a, Register b) {
+ emitString("xor.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void xor_b64(Register d, Register a, Register b) {
+ emitString("xor.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+ }
+
+ public final void xor_b16(Register d, Register a, short b16) {
+ emitString("xor.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b16 + ";" + "");
+ }
+
+ public final void xor_b32(Register d, Register a, int b32) {
+ emitString("xor.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b32 + ";" + "");
+ }
+
+ public final void xor_b64(Register d, Register a, long b64) {
+ emitString("xor.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b64 + ";" + "");
+ }
+
@Override
public PTXAddress makeAddress(Register base, int displacement) {
return new PTXAddress(base, displacement);
}