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,42 **** 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 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) { --- 24,53 ---- import com.oracle.graal.api.code.*; public class PTXAssembler extends AbstractPTXAssembler { 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,66 **** --- 68,85 ---- 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,130 **** --- 140,205 ---- 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,154 **** --- 216,249 ---- 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,262 **** --- 348,361 ---- 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,326 **** --- 416,433 ---- 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,350 **** --- 448,465 ---- 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,374 **** --- 480,497 ---- 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,386 **** --- 500,545 ---- 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,450 **** --- 600,689 ---- 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,594 **** --- 824,881 ---- 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,640 **** 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 st_global_b8(Register a, long immOff, Register b) { emitString("st.global.b8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } --- 916,927 ---- 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, 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,698 **** --- 976,993 ---- 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,722 **** --- 1004,1029 ---- 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,734 **** --- 1032,1049 ---- 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,746 **** --- 1052,1085 ---- 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); }