Mercurial > hg > truffle
changeset 8995:585cc62fcdc5
PTX enhancements - arithmetic, control, float, integer math, control and basic switch
line wrap: on
line diff
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/ValueUtil.java Wed Apr 10 17:10:28 2013 +0200 +++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/ValueUtil.java Wed Apr 10 18:51:21 2013 -0400 @@ -89,13 +89,19 @@ } public static Register asIntReg(Value value) { - assert value.getKind() == Kind.Int; - return asRegister(value); + if (value.getKind() != Kind.Int) { + throw new InternalError("needed Int got: " + value.getKind()); + } else { + return asRegister(value); + } } public static Register asLongReg(Value value) { - assert value.getKind() == Kind.Long : value.getKind(); - return asRegister(value); + if (value.getKind() != Kind.Long) { + throw new InternalError("needed Long got: " + value.getKind()); + } else { + return asRegister(value); + } } public static Register asObjectReg(Value value) {
--- a/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Wed Apr 10 17:10:28 2013 +0200 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Wed Apr 10 18:51:21 2013 -0400 @@ -26,8 +26,7 @@ public class PTXAssembler extends AbstractPTXAssembler { - @SuppressWarnings("unused") - public PTXAssembler(TargetDescription target, RegisterConfig registerConfig) { + public PTXAssembler(TargetDescription target, @SuppressWarnings("unused") RegisterConfig registerConfig) { super(target); } @@ -35,6 +34,18 @@ 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() + ";" + ""); } @@ -59,6 +70,14 @@ 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() + ";" + ""); } @@ -123,7 +142,63 @@ emitString("bra.uni" + " " + tgt + ";" + ""); } - public final void div_s16(Register d, Register a, Register b) { + 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() + ";" + ""); } @@ -143,10 +218,30 @@ 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() + ";" + ""); } @@ -255,6 +350,11 @@ emitString("mov.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); } + public final void mov_u64(@SuppressWarnings("unused") Register d, + @SuppressWarnings("unused") 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() + ";" + ""); } @@ -319,6 +419,14 @@ 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() + ";" + ""); } @@ -343,6 +451,14 @@ 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() + ";" + ""); } @@ -367,6 +483,14 @@ 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() + ";" + ""); } @@ -379,6 +503,42 @@ 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() + ";" + ""); } @@ -443,6 +603,86 @@ 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() + ";" + ""); } @@ -587,6 +827,54 @@ 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() + ";" + ""); } @@ -631,8 +919,8 @@ 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) { @@ -691,6 +979,14 @@ 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() + ";" + ""); } @@ -711,10 +1007,22 @@ 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() + ";" + ""); } @@ -727,6 +1035,14 @@ 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() + ";" + ""); } @@ -739,6 +1055,30 @@ 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);
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/BasicPTXTest.java Wed Apr 10 17:10:28 2013 +0200 +++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/BasicPTXTest.java Wed Apr 10 18:51:21 2013 -0400 @@ -1,5 +1,5 @@ /* - * Copyright (c) 2011, Oracle and/or its affiliates. All rights reserved. + * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved. * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. * * This code is free software; you can redistribute it and/or modify it @@ -22,32 +22,18 @@ */ package com.oracle.graal.compiler.ptx.test; -import java.lang.reflect.*; - -import org.junit.*; +import java.lang.reflect.Method; -import com.oracle.graal.api.code.*; -import com.oracle.graal.api.meta.*; -import com.oracle.graal.api.runtime.*; -import com.oracle.graal.compiler.*; -import com.oracle.graal.compiler.ptx.*; -import com.oracle.graal.compiler.test.*; -import com.oracle.graal.debug.*; -import com.oracle.graal.java.*; -import com.oracle.graal.nodes.*; -import com.oracle.graal.nodes.type.*; -import com.oracle.graal.phases.*; -import com.oracle.graal.phases.PhasePlan.PhasePosition; -import com.oracle.graal.ptx.*; +import org.junit.Test; /** * Test class for small Java methods compiled to PTX kernels. */ -public class BasicPTXTest extends GraalCompilerTest { +public class BasicPTXTest extends PTXTestBase { @Test public void testAdd() { - test("testAddSnippet"); + compile("testAddSnippet"); } public static int testAddSnippet(int a) { @@ -56,48 +42,21 @@ @Test public void testArray() { - test("testArraySnippet"); + compile("testArraySnippet"); } public static int testArraySnippet(int[] array) { return array[0]; } - private CompilationResult test(String snippet) { - StructuredGraph graph = parse(snippet); - Debug.dump(graph, "Graph"); - TargetDescription target = new TargetDescription(new PTX(), true, 1, 0, true); - PTXBackend ptxBackend = new PTXBackend(Graal.getRequiredCapability(CodeCacheProvider.class), target); - PhasePlan phasePlan = new PhasePlan(); - GraphBuilderPhase graphBuilderPhase = new GraphBuilderPhase(runtime, GraphBuilderConfiguration.getDefault(), OptimisticOptimizations.NONE); - phasePlan.addPhase(PhasePosition.AFTER_PARSING, graphBuilderPhase); - phasePlan.addPhase(PhasePosition.AFTER_PARSING, new PTXPhase()); - new PTXPhase().apply(graph); - CompilationResult result = GraalCompiler.compileMethod(runtime, replacements, ptxBackend, target, graph.method(), graph, null, phasePlan, OptimisticOptimizations.NONE, new SpeculationLog()); - return result; - } - - private static class PTXPhase extends Phase { - - @Override - protected void run(StructuredGraph graph) { - for (LocalNode local : graph.getNodes(LocalNode.class)) { - if (local.kind() == Kind.Object) { - local.setStamp(StampFactory.declaredNonNull(local.objectStamp().type())); - } - } - } - - } - public static void main(String[] args) { - BasicPTXTest basicPTXTest = new BasicPTXTest(); + BasicPTXTest test = new BasicPTXTest(); Method[] methods = BasicPTXTest.class.getMethods(); for (Method m : methods) { if (m.getAnnotation(Test.class) != null) { String name = m.getName() + "Snippet"; // CheckStyle: stop system..print check - System.out.println(name + ": \n" + new String(basicPTXTest.test(name).getTargetCode())); + System.out.println(name + ": \n" + new String(test.compile(name).getTargetCode())); // CheckStyle: resume system..print check } }
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ControlTest.java Wed Apr 10 18:51:21 2013 -0400 @@ -0,0 +1,95 @@ +/* + * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ +package com.oracle.graal.compiler.ptx.test; + +import org.junit.Test; + +import java.lang.reflect.Method; + +public class ControlTest extends PTXTestBase { + + @Test + public void testControl() { + compile("testSwitch1I"); + compile("testStatic"); + compile("testCall"); + compile("testLookupSwitch1I"); + } + + public static int testSwitch1I(int a) { + switch (a) { + case 1: + return 2; + case 2: + return 3; + default: + return 4; + } + } + + public static int testLookupSwitch1I(int a) { + switch (a) { + case 0: return 1; + case 1: return 2; + case 2: return 3; + case 3: return 1; + case 4: return 2; + case 5: return 3; + case 6: return 1; + case 7: return 2; + case 8: return 3; + case 9: return 1; + case 10: return 2; + case 11: return 3; + default: return -1; + } + } + + @SuppressWarnings("unused") + private static Object foo = null; + + public static boolean testStatic(Object o) { + foo = o; + return true; + } + + private static int method(int a, int b) { + return a + b; + } + + public static int testCall(@SuppressWarnings("unused") Object o, int a, int b) { + return method(a, b); + } + + public static void main(String[] args) { + ControlTest test = new ControlTest(); + for (Method m : ControlTest.class.getMethods()) { + String name = m.getName(); + if (m.getAnnotation(Test.class) == null && name.startsWith("test")) { + // CheckStyle: stop system..print check + System.out.println(name + ": \n" + new String(test.compile(name).getTargetCode())); + // CheckStyle: resume system..print check + } + } + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/FloatPTXTest.java Wed Apr 10 18:51:21 2013 -0400 @@ -0,0 +1,245 @@ +/* + * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ +package com.oracle.graal.compiler.ptx.test; + +import java.lang.reflect.Method; + +import org.junit.*; + + +/* PTX ISA 3.1 - 8.7.3 Floating-Point Instructions */ +public class FloatPTXTest extends PTXTestBase { + + @Test + public void testAdd() { + compile("testAdd2F"); + compile("testAdd2D"); + compile("testAddFConst"); + compile("testAddConstF"); + compile("testAddDConst"); + compile("testAddConstD"); + } + + public static float testAdd2F(float a, float b) { + return a + b; + } + + public static double testAdd2D(double a, double b) { + return a + b; + } + + public static float testAddFConst(float a) { + return a + 32.0F; + } + + public static float testAddConstF(float a) { + return 32.0F + a; + } + + public static double testAddDConst(double a) { + return a + 32.0; + } + + public static double testAddConstD(double a) { + return 32.0 + a; + } + + @Test + public void testSub() { + compile("testSub2F"); + compile("testSub2D"); + compile("testSubFConst"); + compile("testSubConstF"); + compile("testSubDConst"); + compile("testSubConstD"); + } + + public static float testSub2F(float a, float b) { + return a - b; + } + + public static double testSub2D(double a, double b) { + return a - b; + } + + public static float testSubFConst(float a) { + return a - 32.0F; + } + + public static float testSubConstF(float a) { + return 32.0F - a; + } + + public static double testSubDConst(double a) { + return a - 32.0; + } + + public static double testSubConstD(double a) { + return 32.0 - a; + } + + @Test + public void testMul() { + compile("testMul2F"); + compile("testMul2D"); + compile("testMulFConst"); + compile("testMulConstF"); + compile("testMulDConst"); + compile("testMulConstD"); + } + + public static float testMul2F(float a, float b) { + return a * b; + } + + public static double testMul2D(double a, double b) { + return a * b; + } + + public static float testMulFConst(float a) { + return a * 32.0F; + } + + public static float testMulConstF(float a) { + return 32.0F * a; + } + + public static double testMulDConst(double a) { + return a * 32.0; + } + + public static double testMulConstD(double a) { + return 32.0 * a; + } + + @Test + public void testDiv() { + compile("testDiv2F"); + compile("testDiv2D"); + compile("testDivFConst"); + compile("testDivConstF"); + compile("testDivDConst"); + compile("testDivConstD"); + } + + public static float testDiv2F(float a, float b) { + return a / b; + } + + public static double testDiv2D(double a, double b) { + return a / b; + } + + public static float testDivFConst(float a) { + return a / 32.0F; + } + + public static float testDivConstF(float a) { + return 32.0F / a; + } + + public static double testDivDConst(double a) { + return a / 32.0; + } + + public static double testDivConstD(double a) { + return 32.0 / a; + } + + @Test + public void testNeg() { + compile("testNeg2F"); + compile("testNeg2D"); + } + + public static float testNeg2F(float a) { + return -a; + } + + public static double testNeg2D(double a) { + return -a; + } + + @Test + public void testRem() { + // need linkage to PTX remainder() + // compile("testRem2F"); + // compile("testRem2D"); + } + + public static float testRem2F(float a, float b) { + return a % b; + } + + public static double testRem2D(double a, double b) { + return a % b; + } + + @Test + public void testFloatConversion() { + compile("testF2I"); + compile("testF2L"); + compile("testF2D"); + compile("testD2I"); + compile("testD2L"); + compile("testD2F"); + } + + public static int testF2I(float a) { + return (int) a; + } + + public static long testF2L(float a) { + return (long) a; + } + + public static double testF2D(float a) { + return a; + } + + public static int testD2I(double a) { + return (int) a; + } + + public static long testD2L(double a) { + return (long) a; + } + + public static float testD2F(double a) { + return (float) a; + } + + public static void main(String[] args) { + FloatPTXTest test = new FloatPTXTest(); + for (Method m : FloatPTXTest.class.getMethods()) { + String name = m.getName(); + if (m.getAnnotation(Test.class) == null && + name.startsWith("test") && + name.startsWith("testRem") == false) { + // CheckStyle: stop system..print check + System.out.println(name + ": \n" + new String(test.compile(name).getTargetCode())); + // CheckStyle: resume system..print check + } + } + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/IntegerPTXTest.java Wed Apr 10 18:51:21 2013 -0400 @@ -0,0 +1,192 @@ +/* + * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ +package com.oracle.graal.compiler.ptx.test; + +import org.junit.Test; + +import java.lang.reflect.Method; + + +public class IntegerPTXTest extends PTXTestBase { + + @Test + public void testAdd() { + compile("testAdd2I"); + compile("testAdd2L"); + compile("testAdd2B"); + compile("testAddIConst"); + compile("testAddConstI"); + } + + public static int testAdd2I(int a, int b) { + return a + b; + } + + public static long testAdd2L(long a, long b) { + return a + b; + } + + public static int testAdd2B(byte a, byte b) { + return a + b; + } + + public static int testAddIConst(int a) { + return a + 32; + } + + public static int testAddConstI(int a) { + return 32 + a; + } + + @Test + public void testSub() { + compile("testSub2I"); + compile("testSub2L"); + compile("testSubIConst"); + compile("testSubConstI"); + } + + public static int testSub2I(int a, int b) { + return a - b; + } + + public static long testSub2L(long a, long b) { + return a - b; + } + + public static int testSubIConst(int a) { + return a - 32; + } + + public static int testSubConstI(int a) { + return 32 - a; + } + + @Test + public void testMul() { + compile("testMul2I"); + compile("testMul2L"); + compile("testMulIConst"); + compile("testMulConstI"); + } + + public static int testMul2I(int a, int b) { + return a * b; + } + + public static long testMul2L(long a, long b) { + return a * b; + } + + public static int testMulIConst(int a) { + return a * 32; + } + + public static int testMulConstI(int a) { + return 32 * a; + } + + @Test + public void testDiv() { + compile("testDiv2I"); + compile("testDiv2L"); + compile("testDivIConst"); + compile("testDivConstI"); + } + + public static int testDiv2I(int a, int b) { + return a / b; + } + + public static long testDiv2L(long a, long b) { + return a / b; + } + + public static int testDivIConst(int a) { + return a / 32; + } + + public static int testDivConstI(int a) { + return 32 / a; + } + + @Test + public void testRem() { + compile("testRem2I"); + compile("testRem2L"); + } + + public static int testRem2I(int a, int b) { + return a % b; + } + + public static long testRem2L(long a, long b) { + return a % b; + } + + @Test + public void testIntConversion() { + compile("testI2L"); + compile("testL2I"); + compile("testI2C"); + compile("testI2B"); + compile("testI2F"); + compile("testI2D"); + } + + public static long testI2L(int a) { + return a; + } + + public static char testI2C(int a) { + return (char) a; + } + + public static byte testI2B(int a) { + return (byte) a; + } + + public static float testI2F(int a) { + return a; + } + + public static double testI2D(int a) { + return a; + } + + public static int testL2I(long a) { + return (int) a; + } + + public static void main(String[] args) { + IntegerPTXTest test = new IntegerPTXTest(); + for (Method m : IntegerPTXTest.class.getMethods()) { + String name = m.getName(); + if (m.getAnnotation(Test.class) == null && name.startsWith("test")) { + // CheckStyle: stop system..print check + System.out.println(name + ": \n" + new String(test.compile(name).getTargetCode())); + // CheckStyle: resume system..print check + } + } + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/LogicPTXTest.java Wed Apr 10 18:51:21 2013 -0400 @@ -0,0 +1,138 @@ +/* + * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ +package com.oracle.graal.compiler.ptx.test; + +import java.lang.reflect.Method; + +import org.junit.Test; + + +/* PTX ISA 3.1 - 8.7.5 Logic and Shift Instructions */ +public class LogicPTXTest extends PTXTestBase { + + @Test + public void testAnd() { + compile("testAnd2I"); + compile("testAnd2L"); + } + + public static int testAnd2I(int a, int b) { + return a & b; + } + + public static long testAnd2L(long a, long b) { + return a & b; + } + + @Test + public void testOr() { + compile("testOr2I"); + compile("testOr2L"); + } + + public static int testOr2I(int a, int b) { + return a | b; + } + + public static long testOr2L(long a, long b) { + return a | b; + } + + @Test + public void testXor() { + compile("testXor2I"); + compile("testXor2L"); + } + + public static int testXor2I(int a, int b) { + return a ^ b; + } + + public static long testXor2L(long a, long b) { + return a ^ b; + } + + @Test + public void testNot() { + compile("testNot1I"); + compile("testNot1L"); + } + + public static int testNot1I(int a) { + return ~a; + } + + public static long testNot1L(long a) { + return ~a; + } + + @Test + public void testShiftLeft() { + compile("testShiftLeft2I"); + compile("testShiftLeft2L"); + } + + public static int testShiftLeft2I(int a, int b) { + return a << b; + } + + public static long testShiftLeft2L(long a, int b) { + return a << b; + } + + @Test + public void testShiftRight() { + compile("testShiftRight2I"); + compile("testShiftRight2L"); + compile("testUnsignedShiftRight2I"); + compile("testUnsignedShiftRight2L"); + } + + public static int testShiftRight2I(int a, int b) { + return a >> b; + } + + public static long testShiftRight2L(long a, int b) { + return a >> b; + } + + public static int testUnsignedShiftRight2I(int a, int b) { + return a >>> b; + } + + public static long testUnsignedShiftRight2L(long a, long b) { + return a >>> b; + } + + public static void main(String[] args) { + LogicPTXTest test = new LogicPTXTest(); + for (Method m : LogicPTXTest.class.getMethods()) { + String name = m.getName(); + if (m.getAnnotation(Test.class) == null && name.startsWith("test")) { + // CheckStyle: stop system..print check + System.out.println(name + ": \n" + new String(test.compile(name).getTargetCode())); + // CheckStyle: resume system..print check + } + } + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/PTXPhase.java Wed Apr 10 18:51:21 2013 -0400 @@ -0,0 +1,45 @@ +/* + * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ +package com.oracle.graal.compiler.ptx.test; + +import com.oracle.graal.api.meta.Kind; +import com.oracle.graal.nodes.LocalNode; +import com.oracle.graal.nodes.StructuredGraph; +import com.oracle.graal.nodes.type.StampFactory; +import com.oracle.graal.phases.Phase; + + +public class PTXPhase extends Phase { + @Override + protected void run(StructuredGraph graph) { + /* + * Assume that null checks would be done on the CPU caller side prior + * to copying data onto the GPU. + */ + for (LocalNode local : graph.getNodes(LocalNode.class)) { + if (local.kind() == Kind.Object) { + local.setStamp(StampFactory.declaredNonNull(local.objectStamp().type())); + } + } + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/PTXTestBase.java Wed Apr 10 18:51:21 2013 -0400 @@ -0,0 +1,62 @@ +/* + * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ +package com.oracle.graal.compiler.ptx.test; + +import com.oracle.graal.api.code.CodeCacheProvider; +import com.oracle.graal.api.code.CompilationResult; +import com.oracle.graal.api.code.SpeculationLog; +import com.oracle.graal.api.code.TargetDescription; +import com.oracle.graal.api.runtime.Graal; +import com.oracle.graal.compiler.GraalCompiler; +import com.oracle.graal.compiler.ptx.PTXBackend; +import com.oracle.graal.compiler.test.GraalCompilerTest; +import com.oracle.graal.debug.Debug; +import com.oracle.graal.java.GraphBuilderConfiguration; +import com.oracle.graal.java.GraphBuilderPhase; +import com.oracle.graal.hotspot.HotSpotGraalRuntime; +import com.oracle.graal.nodes.StructuredGraph; +import com.oracle.graal.phases.OptimisticOptimizations; +import com.oracle.graal.phases.PhasePlan; +import com.oracle.graal.phases.PhasePlan.PhasePosition; +import com.oracle.graal.ptx.PTX; + +public abstract class PTXTestBase extends GraalCompilerTest { + + protected CompilationResult compile(String test) { + StructuredGraph graph = parse(test); + Debug.dump(graph, "Graph"); + TargetDescription target = new TargetDescription(new PTX(), true, 1, 0, true); + PTXBackend ptxBackend = new PTXBackend(Graal.getRequiredCapability(CodeCacheProvider.class), target); + PhasePlan phasePlan = new PhasePlan(); + GraphBuilderPhase graphBuilderPhase = new GraphBuilderPhase(runtime, GraphBuilderConfiguration.getDefault(), + OptimisticOptimizations.NONE); + phasePlan.addPhase(PhasePosition.AFTER_PARSING, graphBuilderPhase); + phasePlan.addPhase(PhasePosition.AFTER_PARSING, new PTXPhase()); + new PTXPhase().apply(graph); + CompilationResult result = GraalCompiler.compileMethod(runtime, HotSpotGraalRuntime.getInstance().getReplacements(), + ptxBackend, target, graph.method(), graph, null, phasePlan, + OptimisticOptimizations.NONE, new SpeculationLog()); + return result; + } + +}
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java Wed Apr 10 17:10:28 2013 +0200 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java Wed Apr 10 18:51:21 2013 -0400 @@ -94,7 +94,7 @@ codeBuffer.emitString(""); // XXX For now declare one predicate and all registers - codeBuffer.emitString(" .reg .pred %p;"); + codeBuffer.emitString(" .reg .pred %p,%q;"); codeBuffer.emitString(" .reg .u32 %r<16>;"); // Emit code for the LIR
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java Wed Apr 10 17:10:28 2013 +0200 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java Wed Apr 10 18:51:21 2013 -0400 @@ -28,35 +28,65 @@ import static com.oracle.graal.lir.ptx.PTXBitManipulationOp.IntrinsicOpcode.*; import static com.oracle.graal.lir.ptx.PTXCompare.*; -import com.oracle.graal.api.code.*; -import com.oracle.graal.api.meta.*; -import com.oracle.graal.asm.*; -import com.oracle.graal.compiler.gen.*; -import com.oracle.graal.compiler.target.*; -import com.oracle.graal.graph.*; -import com.oracle.graal.lir.*; +import com.oracle.graal.api.code.AllocatableValue; +import com.oracle.graal.api.code.CodeCacheProvider; +import com.oracle.graal.api.code.DeoptimizationAction; +import com.oracle.graal.api.code.RuntimeCallTarget; +import com.oracle.graal.api.code.StackSlot; +import com.oracle.graal.api.code.TargetDescription; +import com.oracle.graal.api.code.RuntimeCallTarget.Descriptor; +import com.oracle.graal.api.meta.Constant; +import com.oracle.graal.api.meta.Kind; +import com.oracle.graal.api.meta.ResolvedJavaMethod; +import com.oracle.graal.api.meta.Value; +import com.oracle.graal.asm.NumUtil; +import com.oracle.graal.compiler.gen.LIRGenerator; +import com.oracle.graal.compiler.target.LIRGenLowerable; +import com.oracle.graal.graph.GraalInternalError; +import com.oracle.graal.lir.FrameMap; +import com.oracle.graal.lir.LIR; +import com.oracle.graal.lir.LIRFrameState; +import com.oracle.graal.lir.LIRInstruction; +import com.oracle.graal.lir.LIRValueUtil; +import com.oracle.graal.lir.LabelRef; import com.oracle.graal.lir.StandardOp.JumpOp; -import com.oracle.graal.lir.ptx.*; +import com.oracle.graal.lir.Variable; +import com.oracle.graal.lir.ptx.PTXAddressValue; import com.oracle.graal.lir.ptx.PTXArithmetic.Op1Stack; import com.oracle.graal.lir.ptx.PTXArithmetic.Op2Reg; import com.oracle.graal.lir.ptx.PTXArithmetic.Op2Stack; import com.oracle.graal.lir.ptx.PTXArithmetic.ShiftOp; +import com.oracle.graal.lir.ptx.PTXBitManipulationOp; import com.oracle.graal.lir.ptx.PTXCompare.CompareOp; import com.oracle.graal.lir.ptx.PTXControlFlow.BranchOp; +import com.oracle.graal.lir.ptx.PTXControlFlow.CondMoveOp; +import com.oracle.graal.lir.ptx.PTXControlFlow.FloatCondMoveOp; import com.oracle.graal.lir.ptx.PTXControlFlow.ReturnOp; +import com.oracle.graal.lir.ptx.PTXControlFlow.SequentialSwitchOp; +import com.oracle.graal.lir.ptx.PTXControlFlow.TableSwitchOp; import com.oracle.graal.lir.ptx.PTXMove.LoadOp; import com.oracle.graal.lir.ptx.PTXMove.MoveFromRegOp; import com.oracle.graal.lir.ptx.PTXMove.MoveToRegOp; import com.oracle.graal.lir.ptx.PTXMove.StoreOp; -import com.oracle.graal.nodes.*; -import com.oracle.graal.nodes.calc.*; -import com.oracle.graal.nodes.java.*; +import com.oracle.graal.nodes.BreakpointNode; +import com.oracle.graal.nodes.DeoptimizingNode; +import com.oracle.graal.nodes.DirectCallTargetNode; +import com.oracle.graal.nodes.IndirectCallTargetNode; +import com.oracle.graal.nodes.SafepointNode; +import com.oracle.graal.nodes.StructuredGraph; +import com.oracle.graal.nodes.ValueNode; +import com.oracle.graal.nodes.calc.Condition; +import com.oracle.graal.nodes.calc.ConvertNode; +import com.oracle.graal.nodes.java.CompareAndSwapNode; /** * This class implements the PTX specific portion of the LIR generator. */ public class PTXLIRGenerator extends LIRGenerator { + public static final Descriptor ARITHMETIC_FREM = new Descriptor("arithmeticFrem", false, float.class, float.class, float.class); + public static final Descriptor ARITHMETIC_DREM = new Descriptor("arithmeticDrem", false, double.class, double.class, double.class); + public static class PTXSpillMoveFactory implements LIR.SpillMoveFactory { @Override @@ -190,6 +220,18 @@ append(new CompareOp(ICMP, cond, left, right)); append(new BranchOp(cond, label)); break; + case Long: + append(new CompareOp(LCMP, cond, left, right)); + append(new BranchOp(cond, label)); + break; + case Float: + append(new CompareOp(FCMP, cond, left, right)); + append(new BranchOp(cond, label)); + break; + case Double: + append(new CompareOp(DCMP, cond, left, right)); + append(new BranchOp(cond, label)); + break; case Object: append(new CompareOp(ACMP, cond, left, right)); append(new BranchOp(cond, label)); @@ -211,7 +253,67 @@ @Override public Variable emitConditionalMove(Value left, Value right, Condition cond, boolean unorderedIsTrue, Value trueValue, Value falseValue) { - throw new InternalError("NYI"); + boolean mirrored = emitCompare(cond, left, right); + Condition finalCondition = mirrored ? cond.mirror() : cond; + + Variable result = newVariable(trueValue.getKind()); + switch (left.getKind().getStackKind()) { + case Int: + case Long: + case Object: + append(new CondMoveOp(result, finalCondition, load(trueValue), loadNonConst(falseValue))); + break; + case Float: + case Double: + append(new FloatCondMoveOp(result, finalCondition, unorderedIsTrue, load(trueValue), load(falseValue))); + break; + default: + throw GraalInternalError.shouldNotReachHere("missing: " + left.getKind()); + } + return result; + } + + /** + * This method emits the compare instruction, and may reorder the operands. It returns true if + * it did so. + * + * @param a the left operand of the comparison + * @param b the right operand of the comparison + * @return true if the left and right operands were switched, false otherwise + */ + private boolean emitCompare(Condition cond, Value a, Value b) { + Variable left; + Value right; + boolean mirrored; + if (LIRValueUtil.isVariable(b)) { + left = load(b); + right = loadNonConst(a); + mirrored = true; + } else { + left = load(a); + right = loadNonConst(b); + mirrored = false; + } + switch (left.getKind().getStackKind()) { + case Int: + append(new CompareOp(ICMP, cond, left, right)); + break; + case Long: + append(new CompareOp(LCMP, cond, left, right)); + break; + case Object: + append(new CompareOp(ACMP, cond, left, right)); + break; + case Float: + append(new CompareOp(FCMP, cond, left, right)); + break; + case Double: + append(new CompareOp(DCMP, cond, left, right)); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + return mirrored; } @Override @@ -226,6 +328,12 @@ case Int: append(new Op1Stack(INEG, result, input)); break; + case Float: + append(new Op1Stack(FNEG, result, input)); + break; + case Double: + append(new Op1Stack(DNEG, result, input)); + break; default: throw GraalInternalError.shouldNotReachHere(); } @@ -239,8 +347,17 @@ case Int: append(new Op2Stack(IADD, result, a, loadNonConst(b))); break; + case Long: + append(new Op2Stack(LADD, result, a, loadNonConst(b))); + break; + case Float: + append(new Op2Stack(FADD, result, a, loadNonConst(b))); + break; + case Double: + append(new Op2Stack(DADD, result, a, loadNonConst(b))); + break; default: - throw GraalInternalError.shouldNotReachHere(); + throw GraalInternalError.shouldNotReachHere("missing: " + a.getKind() + " prim: " + a.getKind().isPrimitive()); } return result; } @@ -252,8 +369,17 @@ case Int: append(new Op2Stack(ISUB, result, a, loadNonConst(b))); break; + case Long: + append(new Op2Stack(LSUB, result, a, loadNonConst(b))); + break; + case Float: + append(new Op2Stack(FSUB, result, a, loadNonConst(b))); + break; + case Double: + append(new Op2Stack(DSUB, result, a, loadNonConst(b))); + break; default: - throw GraalInternalError.shouldNotReachHere(); + throw GraalInternalError.shouldNotReachHere("missing: " + a.getKind()); } return result; } @@ -265,8 +391,17 @@ case Int: append(new Op2Reg(IMUL, result, a, loadNonConst(b))); break; - default: - throw GraalInternalError.shouldNotReachHere(); + case Long: + append(new Op2Reg(LMUL, result, a, loadNonConst(b))); + break; + case Float: + append(new Op2Stack(FMUL, result, a, loadNonConst(b))); + break; + case Double: + append(new Op2Stack(DMUL, result, a, loadNonConst(b))); + break; + default: + throw GraalInternalError.shouldNotReachHere("missing: " + a.getKind()); } return result; } @@ -279,12 +414,40 @@ @Override public Value emitDiv(Value a, Value b, DeoptimizingNode deopting) { - throw new InternalError("NYI"); + Variable result = newVariable(a.getKind()); + switch (a.getKind()) { + case Int: + append(new Op2Reg(IDIV, result, a, loadNonConst(b))); + break; + case Long: + append(new Op2Reg(LDIV, result, a, loadNonConst(b))); + break; + case Float: + append(new Op2Stack(FDIV, result, a, loadNonConst(b))); + break; + case Double: + append(new Op2Stack(DDIV, result, a, loadNonConst(b))); + break; + default: + throw GraalInternalError.shouldNotReachHere("missing: " + a.getKind()); + } + return result; } @Override public Value emitRem(Value a, Value b, DeoptimizingNode deopting) { - throw new InternalError("NYI"); + Variable result = newVariable(a.getKind()); + switch (a.getKind()) { + case Int: + append(new Op2Reg(IREM, result, a, loadNonConst(b))); + break; + case Long: + append(new Op2Reg(LREM, result, a, loadNonConst(b))); + break; + default: + throw GraalInternalError.shouldNotReachHere("missing: " + a.getKind()); + } + return result; } @Override @@ -304,6 +467,58 @@ case Int: append(new Op2Stack(IAND, result, a, loadNonConst(b))); break; + case Long: + append(new Op2Stack(LAND, result, a, loadNonConst(b))); + break; + + default: + throw GraalInternalError.shouldNotReachHere("missing: " + a.getKind()); + } + return result; + } + + @Override + public Variable emitOr(Value a, Value b) { + Variable result = newVariable(a.getKind()); + switch (a.getKind()) { + case Int: + append(new Op2Stack(IOR, result, a, loadNonConst(b))); + break; + case Long: + append(new Op2Stack(LOR, result, a, loadNonConst(b))); + break; + default: + throw GraalInternalError.shouldNotReachHere("missing: " + a.getKind()); + } + return result; + } + + @Override + public Variable emitXor(Value a, Value b) { + Variable result = newVariable(a.getKind()); + switch (a.getKind()) { + case Int: + append(new Op2Stack(IXOR, result, a, loadNonConst(b))); + break; + case Long: + append(new Op2Stack(LXOR, result, a, loadNonConst(b))); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + return result; + } + + @Override + public Variable emitShl(Value a, Value b) { + Variable result = newVariable(a.getKind()); + switch (a.getKind()) { + case Int: + append(new Op2Stack(ISHL, result, a, loadNonConst(b))); + break; + case Long: + append(new Op1Stack(LSHL, result, loadNonConst(b))); + break; default: throw GraalInternalError.shouldNotReachHere(); } @@ -311,23 +526,19 @@ } @Override - public Variable emitOr(Value a, Value b) { - throw new InternalError("NYI"); - } - - @Override - public Variable emitXor(Value a, Value b) { - throw new InternalError("NYI"); - } - - @Override - public Variable emitShl(Value a, Value b) { - throw new InternalError("NYI"); - } - - @Override public Variable emitShr(Value a, Value b) { - throw new InternalError("NYI"); + Variable result = newVariable(a.getKind()); + switch (a.getKind()) { + case Int: + append(new Op2Stack(ISHR, result, a, loadNonConst(b))); + break; + case Long: + append(new Op1Stack(LSHR, result, loadNonConst(b))); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + return result; } @Override @@ -337,15 +548,87 @@ case Int: append(new ShiftOp(IUSHR, result, a, b)); break; + case Long: + append(new ShiftOp(LUSHR, result, a, b)); + break; default: - GraalInternalError.shouldNotReachHere(); + throw GraalInternalError.shouldNotReachHere(); } return result; } @Override public Variable emitConvert(ConvertNode.Op opcode, Value inputVal) { - throw new InternalError("NYI"); + Variable input = load(inputVal); + Variable result = newVariable(opcode.to); + switch (opcode) { + case I2L: + append(new Unary2Op(I2L, result, input)); + break; + case L2I: + append(new Unary1Op(L2I, result, input)); + break; + case I2B: + append(new Unary2Op(I2B, result, input)); + break; + case I2C: + append(new Unary1Op(I2C, result, input)); + break; + case I2S: + append(new Unary2Op(I2S, result, input)); + break; + case F2D: + append(new Unary2Op(F2D, result, input)); + break; + case D2F: + append(new Unary2Op(D2F, result, input)); + break; + case I2F: + append(new Unary2Op(I2F, result, input)); + break; + case I2D: + append(new Unary2Op(I2D, result, input)); + break; + case F2I: + append(new Unary2Op(F2I, result, input)); + break; + case D2I: + append(new Unary2Op(D2I, result, input)); + break; + case L2F: + append(new Unary2Op(L2F, result, input)); + break; + case L2D: + append(new Unary2Op(L2D, result, input)); + break; + case F2L: + append(new Unary2Op(F2L, result, input)); + break; + case D2L: + append(new Unary2Op(D2L, result, input)); + break; + case MOV_I2F: + append(new Unary2Op(MOV_I2F, result, input)); + break; + case MOV_L2D: + append(new Unary2Op(MOV_L2D, result, input)); + break; + case MOV_F2I: + append(new Unary2Op(MOV_F2I, result, input)); + break; + case MOV_D2L: + append(new Unary2Op(MOV_D2L, result, input)); + break; + case UNSIGNED_I2L: + // Instructions that move or generate 32-bit register values also set the upper 32 + // bits of the register to zero. + // Consequently, there is no need for a special zero-extension move. + emitMove(result, input); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + return result; } @Override @@ -434,7 +717,14 @@ @Override protected void emitSequentialSwitch(Constant[] keyConstants, LabelRef[] keyTargets, LabelRef defaultTarget, Value key) { - throw new InternalError("NYI"); + // Making a copy of the switch value is necessary because jump table destroys the input + // value + if (key.getKind() == Kind.Int || key.getKind() == Kind.Long) { + append(new SequentialSwitchOp(keyConstants, keyTargets, defaultTarget, key, Value.ILLEGAL)); + } else { + assert key.getKind() == Kind.Object : key.getKind(); + append(new SequentialSwitchOp(keyConstants, keyTargets, defaultTarget, key, newVariable(Kind.Object))); + } } @Override @@ -444,7 +734,10 @@ @Override protected void emitTableSwitch(int lowKey, LabelRef defaultTarget, LabelRef[] targets, Value key) { - throw new InternalError("NYI"); + // Making a copy of the switch value is necessary because jump table destroys the input + // value + Variable tmp = emitMove(key); + append(new TableSwitchOp(lowKey, defaultTarget, targets, tmp, newVariable(target.wordKind))); } @Override @@ -459,13 +752,14 @@ @Override public void visitSafepointNode(SafepointNode i) { + // LIRFrameState info = state(); + // append(new PTXSafepointOp(info, runtime().config, this)); throw new InternalError("NYI"); } @Override public void emitUnwind(Value operand) { - // TODO Auto-generated method stub - + throw new InternalError("NYI"); } @Override
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXArithmetic.java Wed Apr 10 17:10:28 2013 +0200 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXArithmetic.java Wed Apr 10 18:51:21 2013 -0400 @@ -36,9 +36,9 @@ public enum PTXArithmetic { IADD, ISUB, IMUL, IDIV, IDIVREM, IREM, IUDIV, IUREM, IAND, IOR, IXOR, ISHL, ISHR, IUSHR, LADD, LSUB, LMUL, LDIV, LDIVREM, LREM, LUDIV, LUREM, LAND, LOR, LXOR, LSHL, LSHR, LUSHR, - FADD, FSUB, FMUL, FDIV, FAND, FOR, FXOR, - DADD, DSUB, DMUL, DDIV, DAND, DOR, DXOR, - INEG, LNEG, + FADD, FSUB, FMUL, FDIV, FREM, FAND, FOR, FXOR, + DADD, DSUB, DMUL, DDIV, DREM, DAND, DOR, DXOR, + INEG, LNEG, FNEG, DNEG, I2L, L2I, I2B, I2C, I2S, F2D, D2F, I2F, I2D, F2I, D2I, @@ -46,6 +46,47 @@ MOV_I2F, MOV_L2D, MOV_F2I, MOV_D2L; + /** + * Unary operation with separate source and destination operand. + */ + public static class Unary2Op extends PTXLIRInstruction { + @Opcode private final PTXArithmetic opcode; + @Def({REG}) protected AllocatableValue result; + @Use({REG, STACK}) protected AllocatableValue x; + + public Unary2Op(PTXArithmetic opcode, AllocatableValue result, AllocatableValue x) { + this.opcode = opcode; + this.result = result; + this.x = x; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + PTXMove.move(tasm, masm, result, x); + emit(tasm, masm, opcode, result, x, null); + } + } + + /** + * Unary operation with single operand for source and destination. + */ + public static class Unary1Op extends PTXLIRInstruction { + @Opcode private final PTXArithmetic opcode; + @Def({REG, HINT}) protected AllocatableValue result; + @Use({REG, STACK}) protected AllocatableValue x; + + public Unary1Op(PTXArithmetic opcode, AllocatableValue result, AllocatableValue x) { + this.opcode = opcode; + this.result = result; + this.x = x; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + emit(tasm, masm, opcode, result); + } + } + public static class Op1Reg extends PTXLIRInstruction { @Opcode private final PTXArithmetic opcode; @Def({REG, HINT}) protected Value result; @@ -213,23 +254,70 @@ } } - - @SuppressWarnings("unused") - protected static void emit(TargetMethodAssembler tasm, PTXAssembler masm, PTXArithmetic opcode, Value result) { + protected static void emit(@SuppressWarnings("unused") TargetMethodAssembler tasm, + PTXAssembler masm, PTXArithmetic opcode, Value result) { switch (opcode) { - default: throw GraalInternalError.shouldNotReachHere(); + case L2I: masm.and_b32(asIntReg(result), asIntReg(result), 0xFFFFFFFF); break; + case I2C: masm.and_b16(asIntReg(result), asIntReg(result), (short) 0xFFFF); break; + default: + throw GraalInternalError.shouldNotReachHere("missing: " + opcode); } } public static void emit(TargetMethodAssembler tasm, PTXAssembler masm, PTXArithmetic opcode, Value dst, Value src, LIRFrameState info) { int exceptionOffset = -1; if (isRegister(src)) { - Register a = asIntReg(src); - Register d = asIntReg(dst); switch (opcode) { - case INEG: masm.neg_s32(d, a); break; + case INEG: + masm.neg_s32(asIntReg(dst), asIntReg(src)); + break; + case I2L: + masm.cvt_s64_s32(asLongReg(dst), asIntReg(src)); + break; + case I2C: + masm.cvt_b16_s32(asIntReg(dst), asIntReg(src)); + break; + case I2B: + masm.cvt_s8_s32(asIntReg(dst), asIntReg(src)); + break; + case I2F: + masm.cvt_f32_s32(asFloatReg(dst), asIntReg(src)); + break; + case I2D: + masm.cvt_f64_s32(asDoubleReg(dst), asIntReg(src)); + break; + case FNEG: + masm.neg_f32(asFloatReg(dst), asFloatReg(src)); + break; + case DNEG: + masm.neg_f64(asDoubleReg(dst), asDoubleReg(src)); + break; + case F2I: + masm.cvt_s32_f32(asIntReg(dst), asFloatReg(src)); + break; + case F2L: + masm.cvt_s64_f32(asLongReg(dst), asFloatReg(src)); + break; + case F2D: + masm.cvt_f64_f32(asDoubleReg(dst), asFloatReg(src)); + break; + case D2I: + masm.cvt_s32_f64(asIntReg(dst), asDoubleReg(src)); + break; + case D2L: + masm.cvt_s64_f64(asLongReg(dst), asDoubleReg(src)); + break; + case D2F: + masm.cvt_f32_f64(asFloatReg(dst), asDoubleReg(src)); + break; + case LSHL: + masm.shl_s64(asLongReg(dst), asLongReg(dst), asIntReg(src)); + break; + case LSHR: + masm.shr_s64(asLongReg(dst), asLongReg(dst), asIntReg(src)); + break; default: - throw GraalInternalError.shouldNotReachHere(); + throw GraalInternalError.shouldNotReachHere("missing: " + opcode); } } else if (isConstant(src)) { switch (opcode) { @@ -252,33 +340,74 @@ public static void emit(TargetMethodAssembler tasm, PTXAssembler masm, PTXArithmetic opcode, Value dst, Value src1, Value src2, LIRFrameState info) { int exceptionOffset = -1; if (isConstant(src1)) { - int a = tasm.asIntConst(src1); - Register b = asIntReg(src2); - Register d = asIntReg(dst); switch (opcode) { - case ISUB: masm.sub_s32(d, a, b); break; - case IAND: masm.and_b32(d, b, a); break; - default: throw GraalInternalError.shouldNotReachHere(); + case ISUB: masm.sub_s32(asIntReg(dst), tasm.asIntConst(src1), asIntReg(src2)); break; + case IAND: masm.and_b32(asIntReg(dst), asIntReg(src2), tasm.asIntConst(src1)); break; + case IDIV: masm.div_s32(asIntReg(dst), tasm.asIntConst(src1), asIntReg(src2)); break; + case FSUB: masm.sub_f32(asFloatReg(dst), tasm.asFloatConst(src1), asFloatReg(src2)); break; + case FDIV: masm.div_f32(asFloatReg(dst), tasm.asFloatConst(src1), asFloatReg(src2)); break; + case DSUB: masm.sub_f64(asDoubleReg(dst), tasm.asDoubleConst(src1), asDoubleReg(src2)); break; + case DDIV: masm.div_f64(asDoubleReg(dst), tasm.asDoubleConst(src1), asDoubleReg(src2)); break; + default: + throw GraalInternalError.shouldNotReachHere(); } } else if (isConstant(src2)) { - Register a = asIntReg(src1); - int b = tasm.asIntConst(src2); - Register d = asIntReg(dst); switch (opcode) { - case IADD: masm.add_s32(d, a, b); break; - case IAND: masm.and_b32(d, a, b); break; - case IUSHR: masm.shr_u32(d, a, b); break; - default: throw GraalInternalError.shouldNotReachHere(); + case IADD: masm.add_s32(asIntReg(dst), asIntReg(src1), tasm.asIntConst(src2)); break; + case ISUB: masm.sub_s32(asIntReg(dst), asIntReg(src1), tasm.asIntConst(src2)); break; + case IMUL: masm.mul_s32(asIntReg(dst), asIntReg(src1), tasm.asIntConst(src2)); break; + case IAND: masm.and_b32(asIntReg(dst), asIntReg(src1), tasm.asIntConst(src2)); break; + case ISHL: masm.shl_s32(asIntReg(dst), asIntReg(src1), tasm.asIntConst(src2)); break; + case ISHR: masm.shr_s32(asIntReg(dst), asIntReg(src1), tasm.asIntConst(src2)); break; + case IUSHR: masm.shr_u32(asIntReg(dst), asIntReg(src1), tasm.asIntConst(src2)); break; + case IXOR: masm.xor_b32(asIntReg(dst), asIntReg(src1), tasm.asIntConst(src2)); break; + case LXOR: masm.xor_b64(asLongReg(dst), asLongReg(src1), tasm.asLongConst(src2)); break; + case LUSHR: masm.shr_u64(asLongReg(dst), asLongReg(src1), tasm.asLongConst(src2)); break; + case FADD: masm.add_f32(asFloatReg(dst), asFloatReg(src1), tasm.asFloatConst(src2)); break; + case FMUL: masm.mul_f32(asFloatReg(dst), asFloatReg(src1), tasm.asFloatConst(src2)); break; + case FDIV: masm.div_f32(asFloatReg(dst), asFloatReg(src1), tasm.asFloatConst(src2)); break; + case DADD: masm.add_f64(asDoubleReg(dst), asDoubleReg(src1), tasm.asDoubleConst(src2)); break; + case DMUL: masm.mul_f64(asDoubleReg(dst), asDoubleReg(src1), tasm.asDoubleConst(src2)); break; + case DDIV: masm.div_f64(asDoubleReg(dst), asDoubleReg(src1), tasm.asDoubleConst(src2)); break; + default: + throw GraalInternalError.shouldNotReachHere(); } } else { - Register a = asIntReg(src1); - Register b = asIntReg(src2); - Register d = asIntReg(dst); switch (opcode) { - case IADD: masm.add_s32(d, a, b); break; - case ISUB: masm.sub_s32(d, a, b); break; - case IMUL: masm.mul_s32(d, a, b); break; - default: throw GraalInternalError.shouldNotReachHere(); + case IADD: masm.add_s32(asIntReg(dst), asIntReg(src1), asIntReg(src2)); break; + case ISUB: masm.sub_s32(asIntReg(dst), asIntReg(src1), asIntReg(src2)); break; + case IMUL: masm.mul_s32(asIntReg(dst), asIntReg(src1), asIntReg(src2)); break; + case IDIV: masm.div_s32(asIntReg(dst), asIntReg(src1), asIntReg(src2)); break; + case IAND: masm.and_b32(asIntReg(dst), asIntReg(src1), asIntReg(src2)); break; + case IOR: masm.or_b32(asIntReg(dst), asIntReg(src1), asIntReg(src2)); break; + case IXOR: masm.xor_b32(asIntReg(dst), asIntReg(src1), asIntReg(src2)); break; + case ISHL: masm.shl_s32(asIntReg(dst), asIntReg(src1), asIntReg(src2)); break; + case ISHR: masm.shr_s32(asIntReg(dst), asIntReg(src1), asIntReg(src2)); break; + case IUSHR: masm.shr_u32(asIntReg(dst), asIntReg(src1), asIntReg(src2)); break; + case IREM: masm.rem_s32(asIntReg(dst), asIntReg(src1), asIntReg(src2)); break; + case LADD: masm.add_s64(asLongReg(dst), asLongReg(src1), asLongReg(src2)); break; + case LSUB: masm.sub_s64(asLongReg(dst), asLongReg(src1), asLongReg(src2)); break; + case LMUL: masm.mul_s64(asLongReg(dst), asLongReg(src1), asLongReg(src2)); break; + case LDIV: masm.div_s64(asLongReg(dst), asLongReg(src1), asLongReg(src2)); break; + case LAND: masm.and_b64(asLongReg(dst), asLongReg(src1), asLongReg(src2)); break; + case LOR: masm.or_b64(asLongReg(dst), asLongReg(src1), asLongReg(src2)); break; + case LXOR: masm.xor_b64(asLongReg(dst), asLongReg(src1), asLongReg(src2)); break; + case LSHL: masm.shl_s64(asLongReg(dst), asLongReg(src1), asLongReg(src2)); break; + case LSHR: masm.shr_s64(asLongReg(dst), asLongReg(src1), asLongReg(src2)); break; + case LUSHR: masm.shr_u64(asLongReg(dst), asLongReg(src1), asIntReg(src2)); break; + case LREM: masm.rem_s64(asLongReg(dst), asLongReg(src1), asLongReg(src2)); break; + case FADD: masm.add_f32(asFloatReg(dst), asFloatReg(src1), asFloatReg(src2)); break; + case FSUB: masm.sub_f32(asFloatReg(dst), asFloatReg(src1), asFloatReg(src2)); break; + case FMUL: masm.mul_f32(asFloatReg(dst), asFloatReg(src1), asFloatReg(src2)); break; + case FDIV: masm.div_f32(asFloatReg(dst), asFloatReg(src1), asFloatReg(src2)); break; + case FREM: masm.div_f32(asFloatReg(dst), asFloatReg(src1), asFloatReg(src2)); break; + case DADD: masm.add_f64(asDoubleReg(dst), asDoubleReg(src1), asDoubleReg(src2)); break; + case DSUB: masm.sub_f64(asDoubleReg(dst), asDoubleReg(src1), asDoubleReg(src2)); break; + case DMUL: masm.mul_f64(asDoubleReg(dst), asDoubleReg(src1), asDoubleReg(src2)); break; + case DDIV: masm.div_f64(asDoubleReg(dst), asDoubleReg(src1), asDoubleReg(src2)); break; + case DREM: masm.div_f64(asDoubleReg(dst), asDoubleReg(src1), asDoubleReg(src2)); break; + default: + throw GraalInternalError.shouldNotReachHere("missing: " + opcode); } } @@ -289,9 +418,11 @@ } private static void verifyKind(PTXArithmetic opcode, Value result, Value x, Value y) { - assert (opcode.name().startsWith("I") && result.getKind() == Kind.Int && x.getKind().getStackKind() == Kind.Int && y.getKind().getStackKind() == Kind.Int) + if (((opcode.name().startsWith("I") && result.getKind() == Kind.Int && x.getKind().getStackKind() == Kind.Int && y.getKind().getStackKind() == Kind.Int) || (opcode.name().startsWith("L") && result.getKind() == Kind.Long && x.getKind() == Kind.Long && y.getKind() == Kind.Long) || (opcode.name().startsWith("F") && result.getKind() == Kind.Float && x.getKind() == Kind.Float && y.getKind() == Kind.Float) - || (opcode.name().startsWith("D") && result.getKind() == Kind.Double && x.getKind() == Kind.Double && y.getKind() == Kind.Double); + || (opcode.name().startsWith("D") && result.getKind() == Kind.Double && x.getKind() == Kind.Double && y.getKind() == Kind.Double)) == false) { + throw GraalInternalError.shouldNotReachHere("opcode: " + opcode.name() + " x: " + x.getKind() + " y: " + y.getKind()); + } } }
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXCompare.java Wed Apr 10 17:10:28 2013 +0200 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXCompare.java Wed Apr 10 18:51:21 2013 -0400 @@ -65,11 +65,15 @@ public static void emit(TargetMethodAssembler tasm, PTXAssembler masm, PTXCompare opcode, Condition condition, Value x, Value y) { if (isConstant(x)) { - int a = tasm.asIntConst(x); - Register b = asIntReg(y); switch (opcode) { case ICMP: - emitCompareConstReg(masm, condition, a, b); + emitCompareConstReg(masm, condition, tasm.asIntConst(x), asIntReg(y)); + break; + case FCMP: + emitCompareConstReg(masm, condition, tasm.asFloatConst(x), asFloatReg(y)); + break; + case DCMP: + emitCompareConstReg(masm, condition, tasm.asDoubleConst(x), asDoubleReg(y)); break; default: throw GraalInternalError.shouldNotReachHere(); @@ -101,18 +105,75 @@ throw GraalInternalError.shouldNotReachHere(); } } else { - Register a = asIntReg(x); - Register b = asIntReg(y); switch (opcode) { case ICMP: - emitCompareRegReg(masm, condition, a, b); + emitCompareRegReg(masm, condition, asIntReg(x), asIntReg(y)); + break; + case LCMP: + emitCompareRegReg(masm, condition, asLongReg(x), asLongReg(y)); + break; + case FCMP: + emitCompareRegReg(masm, condition, asFloatReg(x), asFloatReg(y)); + break; + case DCMP: + emitCompareRegReg(masm, condition, asDoubleReg(x), asDoubleReg(y)); break; default: - throw GraalInternalError.shouldNotReachHere(); + throw GraalInternalError.shouldNotReachHere("missing: " + opcode); } } } + private static void emitCompareConstReg(PTXAssembler masm, Condition condition, float a, Register b) { + switch (condition) { + case EQ: + masm.setp_eq_f32(a, b); + break; + case NE: + masm.setp_ne_f32(a, b); + break; + case LT: + masm.setp_lt_f32(a, b); + break; + case LE: + masm.setp_le_f32(a, b); + break; + case GT: + masm.setp_gt_f32(a, b); + break; + case GE: + masm.setp_ge_f32(a, b); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } + + private static void emitCompareConstReg(PTXAssembler masm, Condition condition, double a, Register b) { + switch (condition) { + case EQ: + masm.setp_eq_f64(a, b); + break; + case NE: + masm.setp_ne_f64(a, b); + break; + case LT: + masm.setp_lt_f64(a, b); + break; + case LE: + masm.setp_le_f64(a, b); + break; + case GT: + masm.setp_gt_f64(a, b); + break; + case GE: + masm.setp_ge_f64(a, b); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } + private static void emitCompareConstReg(PTXAssembler masm, Condition condition, int a, Register b) { switch (condition) { case EQ:
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java Wed Apr 10 17:10:28 2013 +0200 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java Wed Apr 10 18:51:21 2013 -0400 @@ -22,14 +22,28 @@ */ package com.oracle.graal.lir.ptx; +import static com.oracle.graal.api.code.ValueUtil.asIntReg; +import static com.oracle.graal.api.code.ValueUtil.asLongReg; +import static com.oracle.graal.api.code.ValueUtil.asObjectReg; import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; -import com.oracle.graal.api.meta.*; -import com.oracle.graal.asm.*; -import com.oracle.graal.asm.ptx.*; -import com.oracle.graal.lir.*; -import com.oracle.graal.lir.asm.*; -import com.oracle.graal.nodes.calc.*; +import com.oracle.graal.api.code.Register; +import com.oracle.graal.api.code.CompilationResult.JumpTable; +import com.oracle.graal.api.meta.Constant; +import com.oracle.graal.api.meta.Kind; +import com.oracle.graal.api.meta.Value; +import com.oracle.graal.asm.Buffer; +import com.oracle.graal.asm.Label; +import com.oracle.graal.asm.NumUtil; +import com.oracle.graal.asm.ptx.AbstractPTXAssembler; +import com.oracle.graal.asm.ptx.PTXAssembler; +import com.oracle.graal.graph.GraalInternalError; +import com.oracle.graal.lir.LabelRef; +import com.oracle.graal.lir.StandardOp; +import com.oracle.graal.lir.StandardOp.FallThroughOp; +import com.oracle.graal.lir.Variable; +import com.oracle.graal.lir.asm.TargetMethodAssembler; +import com.oracle.graal.nodes.calc.Condition; public class PTXControlFlow { @@ -64,7 +78,7 @@ public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { masm.at(); Label l = destination.label(); - l.addPatchAt(tasm.asm.codeBuffer.position()); + // l.addPatchAt(tasm.asm.codeBuffer.position()); String target = l.isBound() ? "L" + l.toString() : AbstractPTXAssembler.UNBOUND_TARGET; masm.bra(target); } @@ -80,4 +94,176 @@ condition = condition.negate(); } } + + @SuppressWarnings("unused") + public static class CondMoveOp extends PTXLIRInstruction { + @Def({REG, HINT}) protected Value result; + @Alive({REG}) protected Value trueValue; + @Use({REG, STACK, CONST}) protected Value falseValue; + private final Condition condition; + + public CondMoveOp(Variable result, Condition condition, Variable trueValue, Value falseValue) { + this.result = result; + this.condition = condition; + this.trueValue = trueValue; + this.falseValue = falseValue; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + // cmove(tasm, masm, result, false, condition, false, trueValue, falseValue); + // see 8.3 Predicated Execution p. 61 of PTX ISA 3.1 + throw new InternalError("NYI"); + } + } + + @SuppressWarnings("unused") + public static class FloatCondMoveOp extends PTXLIRInstruction { + @Def({REG}) protected Value result; + @Alive({REG}) protected Value trueValue; + @Alive({REG}) protected Value falseValue; + private final Condition condition; + private final boolean unorderedIsTrue; + + public FloatCondMoveOp(Variable result, Condition condition, boolean unorderedIsTrue, Variable trueValue, Variable falseValue) { + this.result = result; + this.condition = condition; + this.unorderedIsTrue = unorderedIsTrue; + this.trueValue = trueValue; + this.falseValue = falseValue; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + // cmove(tasm, masm, result, true, condition, unorderedIsTrue, trueValue, falseValue); + // see 8.3 Predicated Execution p. 61 of PTX ISA 3.1 + throw new InternalError("NYI"); + } + } + + public static class SequentialSwitchOp extends PTXLIRInstruction implements FallThroughOp { + @Use({CONST}) protected Constant[] keyConstants; + private final LabelRef[] keyTargets; + private LabelRef defaultTarget; + @Alive({REG}) protected Value key; + @Temp({REG, ILLEGAL}) protected Value scratch; + + public SequentialSwitchOp(Constant[] keyConstants, LabelRef[] keyTargets, LabelRef defaultTarget, + Value key, Value scratch) { + assert keyConstants.length == keyTargets.length; + this.keyConstants = keyConstants; + this.keyTargets = keyTargets; + this.defaultTarget = defaultTarget; + this.key = key; + this.scratch = scratch; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + if (key.getKind() == Kind.Int) { + Register intKey = asIntReg(key); + for (int i = 0; i < keyConstants.length; i++) { + if (tasm.runtime.needsDataPatch(keyConstants[i])) { + tasm.recordDataReferenceInCode(keyConstants[i], 0, true); + } + long lc = keyConstants[i].asLong(); + assert NumUtil.isInt(lc); + masm.setp_eq_s32((int) lc, intKey); + masm.at(); + Label l = keyTargets[i].label(); + l.addPatchAt(tasm.asm.codeBuffer.position()); + String target = l.isBound() ? "L" + l.toString() : AbstractPTXAssembler.UNBOUND_TARGET; + masm.bra(target); + } + } else if (key.getKind() == Kind.Long) { + Register longKey = asLongReg(key); + for (int i = 0; i < keyConstants.length; i++) { + masm.setp_eq_s64(tasm.asLongConst(keyConstants[i]), longKey); + masm.at(); + Label l = keyTargets[i].label(); + l.addPatchAt(tasm.asm.codeBuffer.position()); + String target = l.isBound() ? "L" + l.toString() : AbstractPTXAssembler.UNBOUND_TARGET; + masm.bra(target); + } + } else if (key.getKind() == Kind.Object) { + Register intKey = asObjectReg(key); + Register temp = asObjectReg(scratch); + for (int i = 0; i < keyConstants.length; i++) { + PTXMove.move(tasm, masm, temp.asValue(Kind.Object), keyConstants[i]); + masm.setp_eq_u32(intKey, temp); + masm.at(); + masm.bra(keyTargets[i].label().toString()); + } + } else { + throw new GraalInternalError("sequential switch only supported for int, long and object"); + } + if (defaultTarget != null) { + masm.jmp(defaultTarget.label()); + } else { + // masm.hlt(); + } + } + + @Override + public LabelRef fallThroughTarget() { + return defaultTarget; + } + + @Override + public void setFallThroughTarget(LabelRef target) { + defaultTarget = target; + } + } + + public static class TableSwitchOp extends PTXLIRInstruction { + private final int lowKey; + private final LabelRef defaultTarget; + private final LabelRef[] targets; + @Alive protected Value index; + @Temp protected Value scratch; + + public TableSwitchOp(final int lowKey, final LabelRef defaultTarget, final LabelRef[] targets, + Variable index, Variable scratch) { + this.lowKey = lowKey; + this.defaultTarget = defaultTarget; + this.targets = targets; + this.index = index; + this.scratch = scratch; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + tableswitch(tasm, masm, lowKey, defaultTarget, targets, asIntReg(index), asLongReg(scratch)); + } + } + + @SuppressWarnings("unused") + private static void tableswitch(TargetMethodAssembler tasm, PTXAssembler masm, int lowKey, + LabelRef defaultTarget, LabelRef[] targets, + Register value, Register scratch) { + Buffer buf = masm.codeBuffer; + // Compare index against jump table bounds + int highKey = lowKey + targets.length - 1; + if (lowKey != 0) { + // subtract the low value from the switch value + masm.sub_s32(value, value, lowKey); + masm.setp_gt_s32(value, highKey - lowKey); + } else { + masm.setp_gt_s32(value, highKey); + } + + // Jump to default target if index is not within the jump table + if (defaultTarget != null) { + masm.at(); + masm.bra(defaultTarget.label().toString()); + } + + // address of jump table + int tablePos = buf.position(); + + JumpTable jt = new JumpTable(tablePos, lowKey, highKey, 4); + tasm.compilationResult.addAnnotation(jt); + + // PTX: unimp: tableswitch extract + } }
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java Wed Apr 10 17:10:28 2013 +0200 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java Wed Apr 10 18:51:21 2013 -0400 @@ -165,8 +165,14 @@ case Int: masm.st_global_s32(addr.getBase(), addr.getDisplacement(), asRegister(input)); break; + case Byte: + masm.st_global_s8(addr.getBase(), addr.getDisplacement(), asRegister(input)); + break; + case Object: + masm.st_global_s32(addr.getBase(), addr.getDisplacement(), asRegister(input)); + break; default: - throw GraalInternalError.shouldNotReachHere(); + throw GraalInternalError.shouldNotReachHere("missing: " + address.getKind()); } } } @@ -250,11 +256,20 @@ case Int: masm.mov_s32(asRegister(result), asRegister(input)); break; + case Long: + masm.mov_s64(asRegister(result), asRegister(input)); + break; + case Float: + masm.mov_f32(asRegister(result), asRegister(input)); + break; + case Double: + masm.mov_f64(asRegister(result), asRegister(input)); + break; case Object: masm.mov_u64(asRegister(result), asRegister(input)); break; default: - throw GraalInternalError.shouldNotReachHere("kind=" + result.getKind()); + throw GraalInternalError.shouldNotReachHere("missing: " + input.getKind()); } } @@ -266,8 +281,24 @@ } masm.mov_s32(asRegister(result), input.asInt()); break; + case Long: + if (tasm.runtime.needsDataPatch(input)) { + tasm.recordDataReferenceInCode(input, 0, true); + } + masm.mov_s64(asRegister(result), input.asLong()); + break; + case Object: + if (input.isNull()) { + masm.mov_u64(asRegister(result), 0x0L); + } else if (tasm.target.inlineObjects) { + tasm.recordDataReferenceInCode(input, 0, true); + masm.mov_u64(asRegister(result), 0xDEADDEADDEADDEADL); + } else { + masm.mov_u64(asRegister(result), tasm.recordDataReferenceInCode(input, 0, false)); + } + break; default: - throw GraalInternalError.shouldNotReachHere(); + throw GraalInternalError.shouldNotReachHere("missing: " + input.getKind()); } }
--- a/graal/com.oracle.graal.lir/src/com/oracle/graal/lir/asm/TargetMethodAssembler.java Wed Apr 10 17:10:28 2013 +0200 +++ b/graal/com.oracle.graal.lir/src/com/oracle/graal/lir/asm/TargetMethodAssembler.java Wed Apr 10 18:51:21 2013 -0400 @@ -167,6 +167,36 @@ } /** + * Returns the float value of any constant that can be represented by a 32-bit float value. + */ + public float asFloatConst(Value value) { + assert (value.getKind().getStackKind() == Kind.Float && isConstant(value)); + Constant constant = (Constant) value; + assert !runtime.needsDataPatch(constant) : constant + " should be in a DataPatch"; + return constant.asFloat(); + } + + /** + * Returns the long value of any constant that can be represented by a 64-bit long value. + */ + public long asLongConst(Value value) { + assert (value.getKind().getStackKind() == Kind.Long && isConstant(value)); + Constant constant = (Constant) value; + assert !runtime.needsDataPatch(constant) : constant + " should be in a DataPatch"; + return constant.asLong(); + } + + /** + * Returns the double value of any constant that can be represented by a 64-bit float value. + */ + public double asDoubleConst(Value value) { + assert (value.getKind().getStackKind() == Kind.Double && isConstant(value)); + Constant constant = (Constant) value; + assert !runtime.needsDataPatch(constant) : constant + " should be in a DataPatch"; + return constant.asDouble(); + } + + /** * Returns the address of a float constant that is embedded as a data references into the code. */ public AbstractAddress asFloatConstRef(Value value) {
--- a/graal/com.oracle.graal.replacements/src/com/oracle/graal/replacements/nodes/MathIntrinsicNode.java Wed Apr 10 17:10:28 2013 +0200 +++ b/graal/com.oracle.graal.replacements/src/com/oracle/graal/replacements/nodes/MathIntrinsicNode.java Wed Apr 10 18:51:21 2013 -0400 @@ -107,6 +107,8 @@ return ConstantNode.forDouble(Math.cos(value), graph()); case TAN: return ConstantNode.forDouble(Math.tan(value), graph()); + default: + throw GraalInternalError.shouldNotReachHere(); } } return this;