Mercurial > hg > graal-compiler
changeset 11744:0ba840567fba
Merge
| author | Andreas Woess <andreas.woess@jku.at> |
|---|---|
| date | Sat, 21 Sep 2013 04:10:53 +0200 |
| parents | ed54ddfa393d (diff) c287d13cb8b0 (current diff) |
| children | ea4e9cbaa0c9 |
| files | |
| diffstat | 8 files changed, 526 insertions(+), 681 deletions(-) [+] |
line wrap: on
line diff
--- a/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Sat Sep 21 04:01:09 2013 +0200 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Sat Sep 21 04:10:53 2013 +0200 @@ -24,13 +24,13 @@ import static com.oracle.graal.api.code.ValueUtil.*; -import com.oracle.graal.api.code.AbstractAddress; import com.oracle.graal.api.code.Register; import com.oracle.graal.api.code.RegisterConfig; import com.oracle.graal.api.code.TargetDescription; 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.nodes.calc.Condition; import com.oracle.graal.graph.GraalInternalError; import com.oracle.graal.lir.Variable; @@ -40,6 +40,62 @@ super(target); } + public enum ConditionOperator { + // @formatter:off + + // Signed integer operators + S_EQ("eq"), + S_NE("ne"), + S_LT("lt"), + S_LE("le"), + S_GT("gt"), + S_GE("ge"), + + // Unsigned integer operators + U_EQ("eq"), + U_NE("ne"), + U_LO("lo"), + U_LS("ls"), + U_HI("hi"), + U_HS("hs"), + + // Bit-size integer operators + B_EQ("eq"), + B_NE("ne"), + + // Floating-point operators + F_EQ("eq"), + F_NE("ne"), + F_LT("lt"), + F_LE("le"), + F_GT("gt"), + F_GE("ge"), + + // Floating-point operators accepting NaN + F_EQU("equ"), + F_NEU("neu"), + F_LTU("ltu"), + F_LEU("leu"), + F_GTU("gtu"), + F_GEU("geu"), + + // Floating-point operators testing for NaN + F_NUM("num"), + F_NAN("nan"); + + // @formatter:on + + private final String operator; + + private ConditionOperator(String op) { + this.operator = op; + } + + public String getOperator() { + return operator; + } + } + public static class StandardFormat { protected Kind valueKind; @@ -257,10 +313,23 @@ space = ss; } + public String emitAddress(Variable var, Value val) { + return ("[" + emitRegister(var) + " + " + val + "]"); + } + @Override - public String emit() { - return (space.getStateName() + "." + typeForKind(valueKind) + - emitRegister(dest) + emitRegister(source1) + emitValue(source2) + ";"); + public String emitRegister(Variable var) { + return ("%r" + var.index); + } + + public String emit(boolean isLoad) { + if (isLoad) { + return (space.getStateName() + "." + typeForKind(valueKind) + " " + + emitRegister(dest) + ", " + emitAddress(source1, source2) + ";"); + } else { + return (space.getStateName() + "." + typeForKind(valueKind) + " " + + emitAddress(dest, source2) + ", " + emitRegister(source1) + ";"); + } } } @@ -456,10 +525,10 @@ } public void emit(PTXAssembler asm) { - asm.emitString("ld." + super.emit()); + asm.emitString("ld." + super.emit(true)); } } - + public static class St extends LoadStoreFormat { public St(PTXStateSpace space, Variable dst, Variable src1, Value src2) { @@ -467,102 +536,30 @@ } public void emit(PTXAssembler asm) { - asm.emitString("ld." + super.emit()); + asm.emitString("st." + super.emit(false)); } } - + public final void exit() { emitString("exit;" + " " + ""); } -/* - public final void ld_global_b8(Register d, Register a, long immOff) { - emitString("ld.global.b8" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - public final void ld_global_b16(Register d, Register a, long immOff) { - emitString("ld.global.b16" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - public final void ld_global_b32(Register d, Register a, long immOff) { - emitString("ld.global.b32" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - public final void ld_global_b64(Register d, Register a, long immOff) { - emitString("ld.global.b64" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - public final void ld_global_u8(Register d, Register a, long immOff) { - emitString("ld.global.u8" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - public final void ld_global_u16(Register d, Register a, long immOff) { - emitString("ld.global.u16" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - public final void ld_global_u32(Register d, Register a, long immOff) { - emitString("ld.global.u32" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - public final void ld_global_u64(Register d, Register a, long immOff) { - emitString("ld.global.u64" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } + public static class Param extends SingleOperandFormat { - public final void ld_global_s8(Register d, Register a, long immOff) { - emitString("ld.global.s8" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - public final void ld_global_s16(Register d, Register a, long immOff) { - emitString("ld.global.s16" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - public final void ld_global_s32(Register d, Register a, long immOff) { - emitString("ld.global.s32" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } + private boolean lastParameter; - public final void ld_global_s64(Register d, Register a, long immOff) { - emitString("ld.global.s64" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - public final void ld_global_f32(Register d, Register a, long immOff) { - emitString("ld.global.f32" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - public final void ld_global_f64(Register d, Register a, long immOff) { - emitString("ld.global.f64" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - // Load from state space to destination register - public final void ld_from_state_space(String s, Register d, Register a, long immOff) { - emitString("ld" + s + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - // Load return address from return parameter which is in .param state space - public final void ld_return_address(String s, Register d, Register a, long immOff) { - emitString("ld.param." + s + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - */ + public Param(Variable d, boolean lastParam) { + super(d, null); + setLastParameter(lastParam); + } - 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 param_8_decl(Register d, boolean lastParam) { - emitString(".param" + " " + ".s8" + " " + d + (lastParam ? "" : ",")); - } + public void setLastParameter(boolean value) { + lastParameter = value; + } - public final void param_16_decl(Register d, boolean lastParam) { - emitString(".param" + " " + ".s16" + " " + d + (lastParam ? "" : ",")); - } - - public final void param_u16_decl(Register d, boolean lastParam) { - emitString(".param" + " " + ".s16" + " " + d + (lastParam ? "" : ",")); - } - - public final void param_32_decl(Register d, boolean lastParam) { - emitString(".param" + " " + ".s32" + " " + d + (lastParam ? "" : ",")); - } - - public final void param_64_decl(Register d, boolean lastParam) { - emitString(".param" + " " + ".s64" + " " + d + (lastParam ? "" : ",")); + public void emit(PTXAssembler asm) { + asm.emitString(".param " + typeForKind(dest.getKind()) + emitVariable(dest) + (lastParameter ? "" : ",")); + } } public final void popc_b32(Register d, Register a) { @@ -581,317 +578,159 @@ emitString("ret.uni;" + " " + ""); } - public final void setp_eq_f32(Register a, Register b, int p) { - emitString("setp.eq.f32" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_ne_f32(Register a, Register b, int p) { - emitString("setp.ne.f32" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_lt_f32(Register a, Register b, int p) { - emitString("setp.lt.f32" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_le_f32(Register a, Register b, int p) { - emitString("setp.le.f32" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } + public static class Setp { - public final void setp_gt_f32(Register a, Register b, int p) { - emitString("setp.gt.f32" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_ge_f32(Register a, Register b, int p) { - emitString("setp.ge.f32" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_eq_f32(float f32, Register b, int p) { - emitString("setp.eq.f32" + " " + "%p" + p + ", " + f32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_ne_f32(float f32, Register b, int p) { - emitString("setp.ne.f32" + " " + "%p" + p + ", " + f32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_lt_f32(float f32, Register b, int p) { - emitString("setp.lt.f32" + " " + "%p" + p + ", " + f32 + ", %r" + b.encoding() + ";" + ""); - } + private ConditionOperator operator; + private Value first, second; + private Kind kind; + private int predicate; - public final void setp_le_f32(float f32, Register b, int p) { - emitString("setp.le.f32" + " " + "%p" + p + ", " + f32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_gt_f32(float f32, Register b, int p) { - emitString("setp.gt.f32" + " " + "%p" + p + ", " + f32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_ge_f32(float f32, Register b, int p) { - emitString("setp.ge.f32" + " " + "%p" + p + ", " + f32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_eq_f64(double f64, Register b, int p) { - emitString("setp.eq.f64" + " " + "%p" + p + ", " + f64 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_ne_f64(double f64, Register b, int p) { - emitString("setp.ne.f64" + " " + "%p" + p + ", " + f64 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_lt_f64(double f64, Register b, int p) { - emitString("setp.lt.f64" + " " + "%p" + p + ", " + f64 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_le_f64(double f64, Register b, int p) { - emitString("setp.le.f64" + " " + "%p" + p + ", " + f64 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_gt_f64(double f64, Register b, int p) { - emitString("setp.gt.f64" + " " + "%p" + p + ", " + f64 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_ge_f64(double f64, Register b, int p) { - emitString("setp.ge.f64" + " " + "%p" + p + ", " + f64 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_eq_s64(Register a, Register b, int p) { - emitString("setp.eq.s64" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } + public Setp(Condition condition, Value first, Value second, int predicateRegisterNumber) { + setFirst(first); + setSecond(second); + setPredicate(predicateRegisterNumber); + setKind(); + setConditionOperator(operatorForConditon(condition)); + } - public final void setp_eq_s64(long s64, Register b, int p) { - emitString("setp.eq.s64" + " " + "%p" + p + ", " + s64 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_eq_s32(Register a, Register b, int p) { - emitString("setp.eq.s32" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_ne_s32(Register a, Register b, int p) { - emitString("setp.ne.s32" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_lt_s32(Register a, Register b, int p) { - emitString("setp.lt.s32" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_le_s32(Register a, Register b, int p) { - emitString("setp.le.s32" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } + public void setFirst(Value v) { + first = v; + } - public final void setp_gt_s32(Register a, Register b, int p) { - emitString("setp.gt.s32" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_ge_s32(Register a, Register b, int p) { - emitString("setp.ge.s32" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_eq_s32(Register a, int s32, int p) { - emitString("setp.eq.s32" + " " + "%p" + p + ", %r" + a.encoding() + ", " + s32 + ";" + ""); - } - - public final void setp_ne_s32(Register a, int s32, int p) { - emitString("setp.ne.s32" + " " + "%p" + p + ", %r" + a.encoding() + ", " + s32 + ";" + ""); - } - - public final void setp_lt_s32(Register a, int s32, int p) { - emitString("setp.lt.s32" + " " + "%p" + p + ", %r" + a.encoding() + ", " + s32 + ";" + ""); - } + public void setSecond(Value v) { + second = v; + } - public final void setp_le_s32(Register a, int s32, int p) { - emitString("setp.le.s32" + " " + "%p" + p + ", %r" + a.encoding() + ", " + s32 + ";" + ""); - } - - public final void setp_gt_s32(Register a, int s32, int p) { - emitString("setp.gt.s32" + " " + "%p" + p + ", %r" + a.encoding() + ", " + s32 + ";" + ""); - } - - public final void setp_ge_s32(Register a, int s32, int p) { - emitString("setp.ge.s32" + " " + "%p" + p + ", %r" + a.encoding() + ", " + s32 + ";" + ""); - } - - public final void setp_eq_s32(int s32, Register b, int p) { - emitString("setp.eq.s32" + " " + "%p" + p + ", " + s32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_ne_s32(int s32, Register b, int p) { - emitString("setp.ne.s32" + " " + "%p" + p + ", " + s32 + ", %r" + b.encoding() + ";" + ""); - } + public void setPredicate(int p) { + predicate = p; + } - public final void setp_lt_s32(int s32, Register b, int p) { - emitString("setp.lt.s32" + " " + "%p" + p + ", " + s32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_le_s32(int s32, Register b, int p) { - emitString("setp.le.s32" + " " + "%p" + p + ", " + s32 + ", %r" + b.encoding() + ";" + ""); - } + public void setConditionOperator(ConditionOperator co) { + operator = co; + } - public final void setp_gt_s32(int s32, Register b, int p) { - emitString("setp.gt.s32" + " " + "%p" + p + ", " + s32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_ge_s32(int s32, Register b, int p) { - emitString("setp.ge.s32" + " " + "%p" + p + ", " + s32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_eq_u32(Register a, Register b, int p) { - emitString("setp.eq.u32" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } + private ConditionOperator operatorForConditon(Condition condition) { + char typeChar = kind.getTypeChar(); - public final void setp_ne_u32(Register a, Register b, int p) { - emitString("setp.ne.u32" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_lt_u32(Register a, Register b, int p) { - emitString("setp.lt.u32" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_le_u32(Register a, Register b, int p) { - emitString("setp.le.u32" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_gt_u32(Register a, Register b, int p) { - emitString("setp.gt.u32" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_ge_u32(Register a, Register b, int p) { - emitString("setp.ge.u32" + " " + "%p" + p + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_eq_u32(Register a, int u32, int p) { - emitString("setp.eq.u32" + " " + "%p" + p + ", %r" + a.encoding() + ", " + u32 + ";" + ""); - } - - public final void setp_ne_u32(Register a, int u32, int p) { - emitString("setp.ne.u32" + " " + "%p" + p + ", %r" + a.encoding() + ", " + u32 + ";" + ""); - } - - public final void setp_lt_u32(Register a, int u32, int p) { - emitString("setp.lt.u32" + " " + "%p" + p + ", %r" + a.encoding() + ", " + u32 + ";" + ""); - } - - public final void setp_le_u32(Register a, int u32, int p) { - emitString("setp.le.u32" + " " + "%p" + p + ", %r" + a.encoding() + ", " + u32 + ";" + ""); - } + switch (typeChar) { + case 'z': + case 'c': + case 'a': + // unsigned + switch (condition) { + case EQ: return ConditionOperator.U_EQ; + case NE: return ConditionOperator.U_NE; + case LT: return ConditionOperator.U_LO; + case LE: return ConditionOperator.U_LS; + case GT: return ConditionOperator.U_HI; + case GE: return ConditionOperator.U_HS; + default: + throw GraalInternalError.shouldNotReachHere(); + } + case 'b': + case 's': + case 'i': + case 'j': + // signed + switch (condition) { + case EQ: return ConditionOperator.S_EQ; + case NE: return ConditionOperator.S_NE; + case LT: return ConditionOperator.S_LT; + case LE: return ConditionOperator.S_LE; + case GT: return ConditionOperator.S_GT; + case GE: return ConditionOperator.S_GE; + default: + throw GraalInternalError.shouldNotReachHere(); + } + case 'f': + case 'd': + // floating point - do these need to accept NaN?? + switch (condition) { + case EQ: return ConditionOperator.F_EQ; + case NE: return ConditionOperator.F_NE; + case LT: return ConditionOperator.F_LT; + case LE: return ConditionOperator.F_LE; + case GT: return ConditionOperator.F_GT; + case GE: return ConditionOperator.F_GE; + default: + throw GraalInternalError.shouldNotReachHere(); + } + default: + throw GraalInternalError.shouldNotReachHere(); + } + } - public final void setp_gt_u32(Register a, int u32, int p) { - emitString("setp.gt.u32" + " " + "%p" + p + ", %r" + a.encoding() + ", " + u32 + ";" + ""); - } - - public final void setp_ge_u32(Register a, int u32, int p) { - emitString("setp.ge.u32" + " " + "%p" + p + ", %r" + a.encoding() + ", " + u32 + ";" + ""); - } - - public final void setp_eq_u32(int u32, Register b, int p) { - emitString("setp.eq.u32" + " " + "%p" + p + ", " + u32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_ne_u32(int u32, Register b, int p) { - emitString("setp.ne.u32" + " " + "%p" + p + ", " + u32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_lt_u32(int u32, Register b, int p) { - emitString("setp.lt.u32" + " " + "%p" + p + ", " + u32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_le_u32(int u32, Register b, int p) { - emitString("setp.le.u32" + " " + "%p" + p + ", " + u32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_gt_u32(int u32, Register b, int p) { - emitString("setp.gt.u32" + " " + "%p" + p + ", " + u32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_ge_u32(int u32, Register b, int p) { - emitString("setp.ge.u32" + " " + "%p" + p + ", " + u32 + ", %r" + b.encoding() + ";" + ""); - } - - // Store in global state space -/* - public final void st_global_b8(Register a, long immOff, Register b) { - emitString("st.global.b8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_b16(Register a, long immOff, Register b) { - emitString("st.global.b16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } + public void setKind() { + assert isConstant(first) && isConstant(second) == false; - public final void st_global_b32(Register a, long immOff, Register b) { - emitString("st.global.b32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_b64(Register a, long immOff, Register b) { - emitString("st.global.b64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } + if (isConstant(first)) { + kind = second.getKind(); + } else { + kind = first.getKind(); + } + } + + public String emitValue(Value v) { + assert v != null; - public final void st_global_u8(Register a, long immOff, Register b) { - emitString("st.global.u8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_u16(Register a, long immOff, Register b) { - emitString("st.global.u16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } + if (isConstant(v)) { + return (", " + emitConstant(v)); + } else { + return (", " + emitVariable((Variable) v)); + } + } - public final void st_global_u32(Register a, long immOff, Register b) { - emitString("st.global.u32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_u64(Register a, long immOff, Register b) { - emitString("st.global.u64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } + public String typeForKind(Kind k) { + switch (k.getTypeChar()) { + case 'z': + return "u8"; + case 'b': + return "s8"; + case 's': + return "s16"; + case 'c': + return "u16"; + case 'i': + return "s32"; + case 'f': + return "f32"; + case 'j': + return "s64"; + case 'd': + return "f64"; + case 'a': + return "u64"; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } - public final void st_global_s8(Register a, long immOff, Register b) { - emitString("st.global.s8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } + public String emitConstant(Value v) { + Constant constant = (Constant) v; - public final void st_global_s16(Register a, long immOff, Register b) { - emitString("st.global.s16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } + switch (v.getKind().getTypeChar()) { + case 'i': + return (String.valueOf((int) constant.asLong())); + case 'f': + return (String.valueOf(constant.asFloat())); + case 'j': + return (String.valueOf(constant.asLong())); + case 'd': + return (String.valueOf(constant.asDouble())); + default: + throw GraalInternalError.shouldNotReachHere(); + } + } + + public String emitVariable(Variable v) { + return ("%r" + v.index); + } - public final void st_global_s32(Register a, long immOff, Register b) { - emitString("st.global.s32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_s64(Register a, long immOff, Register b) { - emitString("st.global.s64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); + public void emit(PTXAssembler asm) { + asm.emitString("setp." + operator.getOperator() + "." + typeForKind(kind) + + " %p" + predicate + emitValue(first) + emitValue(second) + ";"); + } } - - public final void st_global_f32(Register a, long immOff, Register b) { - emitString("st.global.f32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_f64(Register a, long immOff, Register b) { - emitString("st.global.f64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - // Store return value - public final void st_global_return_value_s8(Register a, long immOff, Register b) { - emitString("st.global.s8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_return_value_s32(Register a, long immOff, Register b) { - emitString("st.global.s32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_return_value_s64(Register a, long immOff, Register b) { - emitString("st.global.s64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_return_value_f32(Register a, long immOff, Register b) { - emitString("st.global.f32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_return_value_f64(Register a, long immOff, Register b) { - emitString("st.global.f64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_return_value_u32(Register a, long immOff, Register b) { - emitString("st.global.u32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_return_value_u64(Register a, long immOff, Register b) { - emitString("st.global.u64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } -*/ @Override public PTXAddress makeAddress(Register base, int displacement) { throw GraalInternalError.shouldNotReachHere();
--- a/graal/com.oracle.graal.compiler.hsail/src/com/oracle/graal/compiler/hsail/HSAILBackend.java Sat Sep 21 04:01:09 2013 +0200 +++ b/graal/com.oracle.graal.compiler.hsail/src/com/oracle/graal/compiler/hsail/HSAILBackend.java Sat Sep 21 04:10:53 2013 +0200 @@ -59,9 +59,12 @@ paramTypeMap.put("HotSpotResolvedPrimitiveType<long>", "s64"); } + /** + * Use the HSAIL register set when the compilation target is HSAIL. + */ @Override public FrameMap newFrameMap() { - return new HSAILFrameMap(runtime(), target, runtime().lookupRegisterConfig()); + return new HSAILFrameMap(runtime(), target, new HSAILRegisterConfig()); } @Override
--- a/graal/com.oracle.graal.compiler.hsail/src/com/oracle/graal/compiler/hsail/HSAILCompilationResult.java Sat Sep 21 04:01:09 2013 +0200 +++ b/graal/com.oracle.graal.compiler.hsail/src/com/oracle/graal/compiler/hsail/HSAILCompilationResult.java Sat Sep 21 04:10:53 2013 +0200 @@ -23,33 +23,25 @@ package com.oracle.graal.compiler.hsail; -import static com.oracle.graal.api.code.CodeUtil.*; - +import java.lang.reflect.*; import java.util.logging.*; import com.oracle.graal.api.code.*; -import com.oracle.graal.api.code.CallingConvention.*; -import com.oracle.graal.api.runtime.Graal; -import com.oracle.graal.compiler.GraalCompiler; -import com.oracle.graal.java.GraphBuilderConfiguration; -import com.oracle.graal.java.GraphBuilderPhase; -import com.oracle.graal.nodes.StructuredGraph; -import com.oracle.graal.nodes.spi.GraalCodeCacheProvider; -import com.oracle.graal.phases.OptimisticOptimizations; -import com.oracle.graal.phases.PhasePlan; -import com.oracle.graal.phases.PhasePlan.PhasePosition; +import com.oracle.graal.api.code.CallingConvention.Type; import com.oracle.graal.api.meta.*; -import com.oracle.graal.nodes.type.*; -import com.oracle.graal.graph.GraalInternalError; +import com.oracle.graal.api.runtime.*; +import com.oracle.graal.compiler.*; +import com.oracle.graal.compiler.target.*; import com.oracle.graal.debug.*; +import com.oracle.graal.graph.*; +import com.oracle.graal.hsail.*; +import com.oracle.graal.java.*; import com.oracle.graal.nodes.*; +import com.oracle.graal.nodes.spi.*; +import com.oracle.graal.nodes.type.*; import com.oracle.graal.phases.*; +import com.oracle.graal.phases.PhasePlan.PhasePosition; import com.oracle.graal.phases.tiers.*; -import com.oracle.graal.nodes.spi.Replacements; -import com.oracle.graal.compiler.target.Backend; -import com.oracle.graal.hsail.*; - -import java.lang.reflect.Method; /** * Class that represents a HSAIL compilation result. Includes the compiled HSAIL code. @@ -96,6 +88,33 @@ return getHSAILCompilationResult(graph); } + /** + * HSAIL doesn't have a calling convention as such. Function arguments are actually passed in + * memory but then loaded into registers in the function body. This routine makes sure that + * arguments to a kernel or function are loaded (by the kernel or function body) into registers + * of the appropriate sizes. For example, int and float parameters should appear in S registers, + * whereas double and long parameters should appear in d registers. + */ + public static CallingConvention getHSAILCallingConvention(CallingConvention.Type type, TargetDescription target, ResolvedJavaMethod method, boolean stackOnly) { + Signature sig = method.getSignature(); + JavaType retType = sig.getReturnType(null); + int sigCount = sig.getParameterCount(false); + JavaType[] argTypes; + int argIndex = 0; + if (!Modifier.isStatic(method.getModifiers())) { + argTypes = new JavaType[sigCount + 1]; + argTypes[argIndex++] = method.getDeclaringClass(); + } else { + argTypes = new JavaType[sigCount]; + } + for (int i = 0; i < sigCount; i++) { + argTypes[argIndex++] = sig.getParameterType(i, null); + } + + RegisterConfig registerConfig = new HSAILRegisterConfig(); + return registerConfig.getCallingConvention(type, retType, argTypes, target, stackOnly); + } + public static HSAILCompilationResult getHSAILCompilationResult(StructuredGraph graph) { Debug.dump(graph, "Graph"); TargetDescription target = new TargetDescription(new HSAIL(), true, 8, 0, true); @@ -105,7 +124,7 @@ phasePlan.addPhase(PhasePosition.AFTER_PARSING, graphBuilderPhase); phasePlan.addPhase(PhasePosition.AFTER_PARSING, new HSAILPhase()); new HSAILPhase().apply(graph); - CallingConvention cc = getCallingConvention(runtime, Type.JavaCallee, graph.method(), false); + CallingConvention cc = getHSAILCallingConvention(Type.JavaCallee, target, graph.method(), false); try { CompilationResult compResult = GraalCompiler.compileGraph(graph, cc, graph.method(), runtime, replacements, hsailBackend, target, null, phasePlan, OptimisticOptimizations.NONE, new SpeculationLog(), suitesProvider.getDefaultSuites(), new CompilationResult());
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/GraalCompilerTest.java Sat Sep 21 04:01:09 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/GraalCompilerTest.java Sat Sep 21 04:10:53 2013 +0200 @@ -25,9 +25,11 @@ import static com.oracle.graal.api.code.CodeUtil.*; import static com.oracle.graal.phases.GraalOptions.*; +import java.io.*; import java.lang.reflect.*; import java.util.*; import java.util.concurrent.*; +import java.util.concurrent.atomic.*; import org.junit.*; import org.junit.internal.*; @@ -169,7 +171,7 @@ return parse(getMethod(methodName)); } - private static int compilationId = 0; + private static AtomicInteger compilationId = new AtomicInteger(); /** * Compares two given objects for {@linkplain Assert#assertEquals(Object, Object) equality}. @@ -206,11 +208,45 @@ } } + @SuppressWarnings("serial") + public static class MultiCauseAssertionError extends AssertionError { + + private Throwable[] causes; + + public MultiCauseAssertionError(String message, Throwable... causes) { + super(message); + this.causes = causes; + } + + @Override + public void printStackTrace(PrintStream out) { + super.printStackTrace(out); + int num = 0; + for (Throwable cause : causes) { + if (cause != null) { + out.print("cause " + (num++)); + cause.printStackTrace(out); + } + } + } + + @Override + public void printStackTrace(PrintWriter out) { + super.printStackTrace(out); + int num = 0; + for (Throwable cause : causes) { + if (cause != null) { + out.print("cause " + (num++) + ": "); + cause.printStackTrace(out); + } + } + } + } + protected void testN(int n, final String name, final Object... args) { - final Throwable[] errors = new Throwable[n]; + final List<Throwable> errors = new ArrayList<>(n); Thread[] threads = new Thread[n]; for (int i = 0; i < n; i++) { - final int idx = i; Thread t = new Thread(i + ":" + name) { @Override @@ -218,26 +254,23 @@ try { test(name, args); } catch (Throwable e) { - errors[idx] = e; + errors.add(e); } } }; threads[i] = t; t.start(); } - int failed = 0; for (int i = 0; i < n; i++) { try { threads[i].join(); } catch (InterruptedException e) { - errors[i] = e; - } - if (errors[i] != null) { - errors[i].printStackTrace(); - failed++; + errors.add(e); } } - Assert.assertTrue(failed + " of " + n + " failed", failed == 0); + if (!errors.isEmpty()) { + throw new MultiCauseAssertionError(errors.size() + " failures", errors.toArray(new Throwable[errors.size()])); + } } protected Object referenceInvoke(Method method, Object receiver, Object... args) throws IllegalAccessException, IllegalArgumentException, InvocationTargetException { @@ -421,7 +454,7 @@ } } - final int id = compilationId++; + final int id = compilationId.incrementAndGet(); InstalledCode installedCode = Debug.scope("Compiling", new Object[]{runtime, new DebugDumpScope(String.valueOf(id), true)}, new Callable<InstalledCode>() {
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.hsail/src/com/oracle/graal/hsail/HSAILRegisterConfig.java Sat Sep 21 04:10:53 2013 +0200 @@ -0,0 +1,188 @@ +/* + * Copyright (c) 2009, 2012, 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.hsail; + +import com.oracle.graal.api.code.*; + +import com.oracle.graal.api.code.CallingConvention.Type; +import com.oracle.graal.api.meta.*; +import com.oracle.graal.graph.*; + +import static com.oracle.graal.hsail.HSAIL.*; + +/** + * This class defines a higher level interface for the register allocator to be able to access info + * about the {@link HSAIL} register set. + * + * Note: In HSAIL, the number of registers of each type is actually a variable number that must + * satisfy the equation: Total num of registers = No. of S registers + 2 * (No. of D registers) + 4 + * * (No. of Q registers) = 128 In other words we can have up to 128S or 64 D or 32Q or a blend. + * + * For now we haven't implemented support for a variable sized register file. Instead we've fixed + * the number of registers of each type so that they satisfy the above equation. See {@link HSAIL} + * for more details. + */ +public class HSAILRegisterConfig implements RegisterConfig { + + private final Register[] allocatable = {s0, s1, s2, s3, s4, s5, s6, /* s7, */s8, s9, s10, s11, s12, s13, s14, s15, s16, s17, s18, s19, s20, s21, s22, s23, s24, s25, s26, s27, s28, s29, s30, s31, + d0, d1, d2, d3, d4, d5, d6, d7, d8, d9, d10, d11, d12, d13, d14, d15, q0, q1, q2, q3, q4, q5, q6, q7, q8, q9, q10, q11, q12, q13, q14, q15}; + + private final Register[] regBitness32 = {s0, s1, s2, s3, s4, s5, s6, s7, s8, s9, s10, s11, s12, s13, s14, s15, s16, s17, s18, s19, s20, s21, s22, s23, s24, s25, s26, s27, s28, s29, s30, s31}; + private final Register[] regBitness64 = {d0, d1, d2, d3, d4, d5, d6, d7, d8, d9, d10, d11, d12, d13, d14, d15}; + private final RegisterAttributes[] attributesMap = RegisterAttributes.createMap(this, HSAIL.allRegisters); + + @Override + public Register getReturnRegister(Kind kind) { + switch (kind) { + case Boolean: + case Byte: + case Char: + case Short: + case Int: + case Long: + case Object: + return s0; + case Float: + case Double: + return d0; + case Void: + case Illegal: + return null; + default: + throw new UnsupportedOperationException("no return register for type " + kind); + } + } + + @Override + public Register getFrameRegister() { + // TODO Auto-generated method stub + return null; + } + + @Override + public CallingConvention getCallingConvention(Type type, JavaType returnType, JavaType[] parameterTypes, TargetDescription target, boolean stackOnly) { + return callingConvention(regBitness32, regBitness64, returnType, parameterTypes, type, target, stackOnly); + } + + private CallingConvention callingConvention(Register[] generalParameterRegisters, Register[] longParameterRegisters, JavaType returnType, JavaType[] parameterTypes, Type type, + TargetDescription target, boolean stackOnly) { + AllocatableValue[] locations = new AllocatableValue[parameterTypes.length]; + + int currentRegs32 = 0; + int currentRegs64 = 0; + int currentStackOffset = 0; + + for (int i = 0; i < parameterTypes.length; i++) { + final Kind kind = parameterTypes[i].getKind(); + + switch (kind) { + case Byte: + case Boolean: + case Short: + case Char: + case Int: + case Float: + + if (!stackOnly && currentRegs32 < generalParameterRegisters.length) { + Register register = generalParameterRegisters[currentRegs32++]; + locations[i] = register.asValue(kind); + } + break; + case Long: + case Object: + case Double: + if (!stackOnly && currentRegs64 < longParameterRegisters.length) { + Register register = longParameterRegisters[currentRegs64++]; + locations[i] = register.asValue(kind); + } + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + + if (locations[i] == null) { + locations[i] = StackSlot.get(kind.getStackKind(), currentStackOffset, !type.out); + currentStackOffset += Math.max(target.arch.getSizeInBytes(kind), target.wordSize); + } + } + + Kind returnKind = returnType == null ? Kind.Void : returnType.getKind(); + AllocatableValue returnLocation = returnKind == Kind.Void ? Value.ILLEGAL : getReturnRegister(returnKind).asValue(returnKind); + return new CallingConvention(currentStackOffset, returnLocation, locations); + } + + @Override + public Register[] getCallingConventionRegisters(Type type, Kind kind) { + throw new GraalInternalError("getCallingConventinoRegisters unimplemented"); + } + + @Override + public Register[] getAllocatableRegisters() { + return allocatable.clone(); + } + + @Override + public Register[] getAllocatableRegisters(PlatformKind kind) { + Kind k = (Kind) kind; + switch (k) { + case Int: + case Short: + case Byte: + case Float: + return regBitness32.clone(); + case Long: + case Double: + case Object: + return regBitness64.clone(); + + default: + throw new GraalInternalError("unknown register allocation"); + } + } + + @Override + public Register[] getCallerSaveRegisters() { + // TODO Auto-generated method stub + return new Register[0]; + } + + @Override + public CalleeSaveLayout getCalleeSaveLayout() { + return null; + } + + @Override + public RegisterAttributes[] getAttributesMap() { + return attributesMap.clone(); + } + + @Override + public Register getRegisterForRole(int id) { + throw new UnsupportedOperationException(); + } + + public HSAILRegisterConfig() { + + } +}
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXCompare.java Sat Sep 21 04:01:09 2013 +0200 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXCompare.java Sat Sep 21 04:10:53 2013 +0200 @@ -22,10 +22,10 @@ */ package com.oracle.graal.lir.ptx; +import static com.oracle.graal.asm.ptx.PTXAssembler.*; import static com.oracle.graal.api.code.ValueUtil.*; import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; -import com.oracle.graal.api.code.*; import com.oracle.graal.api.meta.*; import com.oracle.graal.asm.ptx.*; import com.oracle.graal.graph.*; @@ -55,7 +55,7 @@ @Override public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { - emit(tasm, masm, opcode, condition, x, y, predRegNum); + emit(masm, opcode, condition, x, y, predRegNum); } @Override @@ -67,42 +67,18 @@ } } - public static void emit(TargetMethodAssembler tasm, PTXAssembler masm, - PTXCompare opcode, Condition condition, - Value x, Value y, int p) { + public static void emit(PTXAssembler masm, PTXCompare opcode, + Condition condition, Value x, Value y, int p) { if (isConstant(x)) { + new Setp(condition, x, y, p).emit(masm); + } else if (isConstant(y)) { switch (opcode) { case ICMP: - emitCompareConstReg(masm, condition, tasm.asIntConst(x), asIntReg(y), p); - break; - case FCMP: - emitCompareConstReg(masm, condition, tasm.asFloatConst(x), asFloatReg(y), p); - break; - case DCMP: - emitCompareConstReg(masm, condition, tasm.asDoubleConst(x), asDoubleReg(y), p); - break; - default: - throw GraalInternalError.shouldNotReachHere(); - } - } else if (isConstant(y)) { - Register a = asIntReg(x); - int b = tasm.asIntConst(y); - switch (opcode) { - case ICMP: - emitCompareRegConst(masm, condition, a, b, p); + new Setp(condition, x, y, p).emit(masm); break; case ACMP: if (((Constant) y).isNull()) { - switch (condition) { - case EQ: - masm.setp_eq_s32(a, b, p); - break; - case NE: - masm.setp_ne_s32(a, b, p); - break; - default: - throw GraalInternalError.shouldNotReachHere(); - } + new Setp(condition, x, y, p).emit(masm); } else { throw GraalInternalError.shouldNotReachHere("Only null object constants are allowed in comparisons"); } @@ -111,183 +87,8 @@ throw GraalInternalError.shouldNotReachHere(); } } else { - switch (opcode) { - case ICMP: - emitCompareRegReg(masm, condition, asIntReg(x), asIntReg(y), p); - break; - case LCMP: - emitCompareRegReg(masm, condition, asLongReg(x), asLongReg(y), p); - break; - case FCMP: - emitCompareRegReg(masm, condition, asFloatReg(x), asFloatReg(y), p); - break; - case DCMP: - emitCompareRegReg(masm, condition, asDoubleReg(x), asDoubleReg(y), p); - break; - default: - throw GraalInternalError.shouldNotReachHere("missing: " + opcode); - } - } - } - - private static void emitCompareConstReg(PTXAssembler masm, Condition condition, float a, Register b, int p) { - switch (condition) { - case EQ: - masm.setp_eq_f32(a, b, p); - break; - case NE: - masm.setp_ne_f32(a, b, p); - break; - case LT: - masm.setp_lt_f32(a, b, p); - break; - case LE: - masm.setp_le_f32(a, b, p); - break; - case GT: - masm.setp_gt_f32(a, b, p); - break; - case GE: - masm.setp_ge_f32(a, b, p); - break; - default: - throw GraalInternalError.shouldNotReachHere(); - } - } - - private static void emitCompareConstReg(PTXAssembler masm, Condition condition, double a, Register b, int p) { - switch (condition) { - case EQ: - masm.setp_eq_f64(a, b, p); - break; - case NE: - masm.setp_ne_f64(a, b, p); - break; - case LT: - masm.setp_lt_f64(a, b, p); - break; - case LE: - masm.setp_le_f64(a, b, p); - break; - case GT: - masm.setp_gt_f64(a, b, p); - break; - case GE: - masm.setp_ge_f64(a, b, p); - break; - default: - throw GraalInternalError.shouldNotReachHere(); + new Setp(condition, x, y, p).emit(masm); } } - private static void emitCompareConstReg(PTXAssembler masm, Condition condition, int a, Register b, int p) { - switch (condition) { - case EQ: - masm.setp_eq_s32(a, b, p); - break; - case NE: - masm.setp_ne_s32(a, b, p); - break; - case LT: - masm.setp_lt_s32(a, b, p); - break; - case LE: - masm.setp_le_s32(a, b, p); - break; - case GT: - masm.setp_gt_s32(a, b, p); - break; - case GE: - masm.setp_ge_s32(a, b, p); - break; - case AT: - masm.setp_gt_u32(a, b, p); - break; - case AE: - masm.setp_ge_u32(a, b, p); - break; - case BT: - masm.setp_lt_u32(a, b, p); - break; - case BE: - masm.setp_le_u32(a, b, p); - break; - default: - throw GraalInternalError.shouldNotReachHere(); - } - } - - private static void emitCompareRegConst(PTXAssembler masm, Condition condition, Register a, int b, int p) { - switch (condition) { - case EQ: - masm.setp_eq_s32(a, b, p); - break; - case NE: - masm.setp_ne_s32(a, b, p); - break; - case LT: - masm.setp_lt_s32(a, b, p); - break; - case LE: - masm.setp_le_s32(a, b, p); - break; - case GT: - masm.setp_gt_s32(a, b, p); - break; - case GE: - masm.setp_ge_s32(a, b, p); - break; - case AT: - masm.setp_gt_u32(a, b, p); - break; - case AE: - masm.setp_ge_u32(a, b, p); - break; - case BT: - masm.setp_lt_u32(a, b, p); - break; - case BE: - masm.setp_le_u32(a, b, p); - break; - default: - throw GraalInternalError.shouldNotReachHere(); - } - } - - private static void emitCompareRegReg(PTXAssembler masm, Condition condition, Register a, Register b, int p) { - switch (condition) { - case EQ: - masm.setp_eq_s32(a, b, p); - break; - case NE: - masm.setp_ne_s32(a, b, p); - break; - case LT: - masm.setp_lt_s32(a, b, p); - break; - case LE: - masm.setp_le_s32(a, b, p); - break; - case GT: - masm.setp_gt_s32(a, b, p); - break; - case GE: - masm.setp_ge_s32(a, b, p); - break; - case AT: - masm.setp_gt_u32(a, b, p); - break; - case AE: - masm.setp_ge_u32(a, b, p); - break; - case BT: - masm.setp_lt_u32(a, b, p); - break; - case BE: - masm.setp_le_u32(a, b, p); - break; - default: - throw GraalInternalError.shouldNotReachHere(); - } - } }
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java Sat Sep 21 04:01:09 2013 +0200 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java Sat Sep 21 04:10:53 2013 +0200 @@ -22,11 +22,11 @@ */ package com.oracle.graal.lir.ptx; -import static com.oracle.graal.api.code.ValueUtil.*; +import static com.oracle.graal.asm.ptx.PTXAssembler.*; import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; +import static com.oracle.graal.nodes.calc.Condition.*; import com.oracle.graal.api.code.CompilationResult.JumpTable; -import com.oracle.graal.api.code.*; import com.oracle.graal.api.meta.*; import com.oracle.graal.asm.*; import com.oracle.graal.asm.ptx.*; @@ -167,29 +167,20 @@ @Override public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { - if (key.getKind() == Kind.Int) { - Register intKey = asIntReg(key); + Kind keyKind = key.getKind(); + + if (keyKind == Kind.Int || keyKind == Kind.Long) { 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, predRegNum); + new Setp(EQ, keyConstants[i], key, predRegNum).emit(masm); masm.bra(masm.nameOf(keyTargets[i].label()), predRegNum); } - } else if (key.getKind() == Kind.Long) { - Register longKey = asLongReg(key); + } else if (keyKind == Kind.Object) { for (int i = 0; i < keyConstants.length; i++) { - masm.setp_eq_s64(tasm.asLongConst(keyConstants[i]), longKey, predRegNum); - masm.bra(masm.nameOf(keyTargets[i].label()), predRegNum); - } - } 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, predRegNum); + PTXMove.move(tasm, masm, scratch, keyConstants[i]); + new Setp(EQ, keyConstants[i], scratch, predRegNum).emit(masm); masm.bra(keyTargets[i].label().toString(), predRegNum); } } else { @@ -234,21 +225,21 @@ @Override public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { - tableswitch(tasm, masm, lowKey, defaultTarget, targets, asIntReg(index), asLongReg(scratch), predRegNum); + tableswitch(tasm, masm, lowKey, defaultTarget, targets, index, scratch, predRegNum); } } @SuppressWarnings("unused") - private static void tableswitch(TargetMethodAssembler tasm, PTXAssembler masm, int lowKey, LabelRef defaultTarget, LabelRef[] targets, Register value, Register scratch, int predNum) { + private static void tableswitch(TargetMethodAssembler tasm, PTXAssembler masm, int lowKey, LabelRef defaultTarget, LabelRef[] targets, Value value, Value scratch, int predNum) { 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 // new Sub(value, value, lowKey).emit(masm); - masm.setp_gt_s32(value, highKey - lowKey, predNum); + new Setp(GT, value, Constant.forInt(highKey - lowKey), predNum).emit(masm); } else { - masm.setp_gt_s32(value, highKey, predNum); + new Setp(GT, value, Constant.forInt(highKey), predNum).emit(masm); } // Jump to default target if index is not within the jump table
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXParameterOp.java Sat Sep 21 04:01:09 2013 +0200 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXParameterOp.java Sat Sep 21 04:10:53 2013 +0200 @@ -23,12 +23,11 @@ package com.oracle.graal.lir.ptx; -import static com.oracle.graal.api.code.ValueUtil.*; +import static com.oracle.graal.asm.ptx.PTXAssembler.*; import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; import com.oracle.graal.api.meta.*; import com.oracle.graal.asm.ptx.*; -import com.oracle.graal.graph.*; import com.oracle.graal.lir.*; import com.oracle.graal.lir.asm.*; @@ -46,35 +45,7 @@ // Emit parameter directives for arguments int argCount = params.length; for (int i = 0; i < argCount; i++) { - Kind paramKind = params[i].getKind(); - switch (paramKind) { - case Byte: - masm.param_8_decl(asRegister(params[i]), (i == (argCount - 1))); - break; - case Short: - masm.param_16_decl(asRegister(params[i]), (i == (argCount - 1))); - break; - case Char: - masm.param_u16_decl(asRegister(params[i]), (i == (argCount - 1))); - break; - case Int: - masm.param_32_decl(asIntReg(params[i]), (i == (argCount - 1))); - break; - case Long: - masm.param_64_decl(asLongReg(params[i]), (i == (argCount - 1))); - break; - case Float: - masm.param_32_decl(asFloatReg(params[i]), (i == (argCount - 1))); - break; - case Double: - masm.param_64_decl(asDoubleReg(params[i]), (i == (argCount - 1))); - break; - case Object: - masm.param_64_decl(asObjectReg(params[i]), (i == (argCount - 1))); - break; - default : - throw GraalInternalError.shouldNotReachHere("unhandled parameter type " + paramKind.toString()); - } + new Param((Variable) params[i], (i == (argCount - 1))).emit(masm); } } }