Mercurial > hg > truffle
changeset 7839:838293a77af7
Make Address class platform specific.
line wrap: on
line diff
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.amd64/src/com/oracle/graal/amd64/AMD64Address.java Fri Feb 22 12:15:14 2013 +0100 @@ -0,0 +1,203 @@ +/* + * Copyright (c) 2010, 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.amd64; + +import static com.oracle.graal.api.code.ValueUtil.*; + +import com.oracle.graal.api.code.*; +import com.oracle.graal.api.meta.*; + +/** + * Represents an address in target machine memory, specified via some combination of a base + * register, an index register, a displacement and a scale. Note that the base and index registers + * may be a variable that will get a register assigned later by the register allocator. + */ +public final class AMD64Address extends Address { + + private static final long serialVersionUID = -4101548147426595051L; + + private final Value[] baseIndex; + private final Scale scale; + private final int displacement; + + /** + * Creates an {@link AMD64Address} with given base register, no scaling and no displacement. + * + * @param kind the kind of the value being addressed + * @param base the base register + */ + public AMD64Address(Kind kind, Value base) { + this(kind, base, ILLEGAL, Scale.Times1, 0); + } + + /** + * Creates an {@link AMD64Address} with given base register, no scaling and a given + * displacement. + * + * @param kind the kind of the value being addressed + * @param base the base register + * @param displacement the displacement + */ + public AMD64Address(Kind kind, Value base, int displacement) { + this(kind, base, ILLEGAL, Scale.Times1, displacement); + } + + /** + * Creates an {@link AMD64Address} with given base and index registers, scaling and + * displacement. This is the most general constructor. + * + * @param kind the kind of the value being addressed + * @param base the base register + * @param index the index register + * @param scale the scaling factor + * @param displacement the displacement + */ + public AMD64Address(Kind kind, Value base, Value index, Scale scale, int displacement) { + super(kind); + this.baseIndex = new Value[2]; + this.setBase(base); + this.setIndex(index); + this.scale = scale; + this.displacement = displacement; + + assert !isConstant(base) && !isStackSlot(base); + assert !isConstant(index) && !isStackSlot(index); + } + + /** + * A scaling factor used in the SIB addressing mode. + */ + public enum Scale { + Times1(1, 0), Times2(2, 1), Times4(4, 2), Times8(8, 3); + + private Scale(int value, int log2) { + this.value = value; + this.log2 = log2; + } + + /** + * The value (or multiplier) of this scale. + */ + public final int value; + + /** + * The {@linkplain #value value} of this scale log 2. + */ + public final int log2; + + public static Scale fromInt(int scale) { + switch (scale) { + case 1: + return Times1; + case 2: + return Times2; + case 4: + return Times4; + case 8: + return Times8; + default: + throw new IllegalArgumentException(String.valueOf(scale)); + } + } + } + + @Override + public Value[] components() { + return baseIndex; + } + + @Override + public String toString() { + StringBuilder s = new StringBuilder(); + s.append(getKind().getJavaName()).append("["); + String sep = ""; + if (isLegal(getBase())) { + s.append(getBase()); + sep = " + "; + } + if (isLegal(getIndex())) { + s.append(sep).append(getIndex()).append(" * ").append(getScale().value); + sep = " + "; + } + if (getDisplacement() < 0) { + s.append(" - ").append(-getDisplacement()); + } else if (getDisplacement() > 0) { + s.append(sep).append(getDisplacement()); + } + s.append("]"); + return s.toString(); + } + + @Override + public boolean equals(Object obj) { + if (obj instanceof AMD64Address) { + AMD64Address addr = (AMD64Address) obj; + return getKind() == addr.getKind() && getDisplacement() == addr.getDisplacement() && getBase().equals(addr.getBase()) && getScale() == addr.getScale() && + getIndex().equals(addr.getIndex()); + } + return false; + } + + @Override + public int hashCode() { + return getBase().hashCode() ^ getIndex().hashCode() ^ (getDisplacement() << 4) ^ (getScale().value << 8) ^ (getKind().ordinal() << 12); + } + + /** + * @return Base register that defines the start of the address computation. If not present, is + * denoted by {@link Value#ILLEGAL}. + */ + public Value getBase() { + return baseIndex[0]; + } + + public void setBase(Value base) { + this.baseIndex[0] = base; + } + + /** + * @return Index register, the value of which (possibly scaled by {@link #scale}) is added to + * {@link #getBase}. If not present, is denoted by {@link Value#ILLEGAL}. + */ + public Value getIndex() { + return baseIndex[1]; + } + + public void setIndex(Value index) { + this.baseIndex[1] = index; + } + + /** + * @return Scaling factor for indexing, dependent on target operand size. + */ + public Scale getScale() { + return scale; + } + + /** + * @return Optional additive displacement. + */ + public int getDisplacement() { + return displacement; + } +}
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/Address.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/Address.java Fri Feb 22 12:15:14 2013 +0100 @@ -1,5 +1,5 @@ /* - * Copyright (c) 2010, 2012, Oracle and/or its affiliates. All rights reserved. + * Copyright (c) 2010, 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,16 +22,13 @@ */ package com.oracle.graal.api.code; -import static com.oracle.graal.api.code.ValueUtil.*; - import com.oracle.graal.api.meta.*; /** - * Represents an address in target machine memory, specified via some combination of a base - * register, an index register, a displacement and a scale. Note that the base and index registers - * may be a variable that will get a register assigned later by the register allocator. + * Base class to represent an address in target machine memory. The concrete representation of the + * address is platform dependent. */ -public final class Address extends Value { +public abstract class Address extends Value { private static final long serialVersionUID = -1003772042519945089L; @@ -39,168 +36,20 @@ * A sentinel value used as a place holder in an instruction stream for an address that will be * patched. */ - public static final Address Placeholder = new Address(Kind.Illegal, Value.ILLEGAL); - - private Value base; - private Value index; - private final Scale scale; - private final int displacement; - - /** - * Creates an {@link Address} with given base register, no scaling and no displacement. + /* + * @SuppressWarnings("serial") public static final Address Placeholder = new + * Address(Kind.Illegal) { * - * @param kind the kind of the value being addressed - * @param base the base register - */ - public Address(Kind kind, Value base) { - this(kind, base, ILLEGAL, Scale.Times1, 0); - } - - /** - * Creates an {@link Address} with given base register, no scaling and a given displacement. - * - * @param kind the kind of the value being addressed - * @param base the base register - * @param displacement the displacement + * @Override public String toString() { return "[<placeholder>]"; } }; */ - public Address(Kind kind, Value base, int displacement) { - this(kind, base, ILLEGAL, Scale.Times1, displacement); - } - /** - * Creates an {@link Address} with given base and index registers, scaling and displacement. - * This is the most general constructor. - * - * @param kind the kind of the value being addressed - * @param base the base register - * @param index the index register - * @param scale the scaling factor - * @param displacement the displacement - */ - public Address(Kind kind, Value base, Value index, Scale scale, int displacement) { + public Address(Kind kind) { super(kind); - this.setBase(base); - this.setIndex(index); - this.scale = scale; - this.displacement = displacement; - - assert !isConstant(base) && !isStackSlot(base); - assert !isConstant(index) && !isStackSlot(index); } /** - * A scaling factor used in complex addressing modes such as those supported by x86 platforms. + * The values that this address is composed of. Used by the register allocator to manipulate + * addresses in a platform independent way. */ - public enum Scale { - Times1(1, 0), Times2(2, 1), Times4(4, 2), Times8(8, 3); - - private Scale(int value, int log2) { - this.value = value; - this.log2 = log2; - } - - /** - * The value (or multiplier) of this scale. - */ - public final int value; - - /** - * The {@linkplain #value value} of this scale log 2. - */ - public final int log2; - - public static Scale fromInt(int scale) { - switch (scale) { - case 1: - return Times1; - case 2: - return Times2; - case 4: - return Times4; - case 8: - return Times8; - default: - throw new IllegalArgumentException(String.valueOf(scale)); - } - } - } - - @Override - public String toString() { - if (this == Placeholder) { - return "[<placeholder>]"; - } - - StringBuilder s = new StringBuilder(); - s.append(getKind().getJavaName()).append("["); - String sep = ""; - if (isLegal(getBase())) { - s.append(getBase()); - sep = " + "; - } - if (isLegal(getIndex())) { - s.append(sep).append(getIndex()).append(" * ").append(getScale().value); - sep = " + "; - } - if (getDisplacement() < 0) { - s.append(" - ").append(-getDisplacement()); - } else if (getDisplacement() > 0) { - s.append(sep).append(getDisplacement()); - } - s.append("]"); - return s.toString(); - } - - @Override - public boolean equals(Object obj) { - if (obj instanceof Address) { - Address addr = (Address) obj; - return getKind() == addr.getKind() && getDisplacement() == addr.getDisplacement() && getBase().equals(addr.getBase()) && getScale() == addr.getScale() && - getIndex().equals(addr.getIndex()); - } - return false; - } - - @Override - public int hashCode() { - return getBase().hashCode() ^ getIndex().hashCode() ^ (getDisplacement() << 4) ^ (getScale().value << 8) ^ (getKind().ordinal() << 12); - } - - /** - * @return Base register that defines the start of the address computation. If not present, is - * denoted by {@link Value#ILLEGAL}. - */ - public Value getBase() { - return base; - } - - public void setBase(Value base) { - this.base = base; - } - - /** - * @return Index register, the value of which (possibly scaled by {@link #scale}) is added to - * {@link #base}. If not present, is denoted by {@link Value#ILLEGAL}. - */ - public Value getIndex() { - return index; - } - - public void setIndex(Value index) { - this.index = index; - } - - /** - * @return Scaling factor for indexing, dependent on target operand size. - */ - public Scale getScale() { - return scale; - } - - /** - * @return Optional additive displacement. - */ - public int getDisplacement() { - return displacement; - } + public abstract Value[] components(); }
--- a/graal/com.oracle.graal.asm.amd64.test/src/com/oracle/graal/asm/amd64/test/SimpleAssemblerTest.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.asm.amd64.test/src/com/oracle/graal/asm/amd64/test/SimpleAssemblerTest.java Fri Feb 22 12:15:14 2013 +0100 @@ -57,7 +57,7 @@ AMD64MacroAssembler asm = new AMD64MacroAssembler(target, registerConfig); Register ret = registerConfig.getReturnRegister(Kind.Double); compResult.recordDataReference(asm.codeBuffer.position(), Constant.forDouble(84.72), 8, false); - asm.movdbl(ret, Address.Placeholder); + asm.movdbl(ret, asm.getPlaceholder()); asm.ret(0); return asm.codeBuffer; }
--- a/graal/com.oracle.graal.asm.amd64/src/com/oracle/graal/asm/amd64/AMD64Assembler.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.asm.amd64/src/com/oracle/graal/asm/amd64/AMD64Assembler.java Fri Feb 22 12:15:14 2013 +0100 @@ -1,5 +1,5 @@ /* - * Copyright (c) 2009, 2012, Oracle and/or its affiliates. All rights reserved. + * Copyright (c) 2009, 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 @@ -47,6 +47,12 @@ private static final int MinEncodingNeedsRex = 8; /** + * A sentinel value used as a place holder in an instruction stream for an address that will be + * patched. + */ + private static final AMD64Address Placeholder = new AMD64Address(Kind.Illegal, rip.asValue()); + + /** * The x86 condition codes used for conditional jumps/moves. */ public enum ConditionFlag { @@ -197,7 +203,7 @@ } // immediate-to-memory forms - private void emitArithOperand(int op1, Register rm, Address adr, int imm32) { + private void emitArithOperand(int op1, Register rm, AMD64Address adr, int imm32) { assert (op1 & 0x01) == 1 : "should be 32bit operation"; assert (op1 & 0x02) == 0 : "sign-extension bit should not be set"; if (isByte(imm32)) { @@ -217,11 +223,11 @@ emitByte(op2 | encode(dst) << 3 | encode(src)); } - protected void emitOperandHelper(Register reg, Address addr) { + protected void emitOperandHelper(Register reg, AMD64Address addr) { Register base = isLegal(addr.getBase()) ? asRegister(addr.getBase()) : Register.None; Register index = isLegal(addr.getIndex()) ? asRegister(addr.getIndex()) : Register.None; - Address.Scale scale = addr.getScale(); + AMD64Address.Scale scale = addr.getScale(); int disp = addr.getDisplacement(); if (base == Register.Frame) { @@ -234,15 +240,11 @@ assert reg != Register.None; int regenc = encode(reg) << 3; - if (base == AMD64.rip) { + if (base == AMD64.rip) { // also matches Placeholder // [00 000 101] disp32 + assert index == Register.None : "cannot use RIP relative addressing with index register"; emitByte(0x05 | regenc); emitInt(disp); - } else if (addr == Address.Placeholder) { - // [00 000 101] disp32 - emitByte(0x05 | regenc); - emitInt(0); - } else if (base.isValid()) { int baseenc = base.isValid() ? encode(base) : 0; if (index.isValid()) { @@ -327,7 +329,7 @@ } } - public final void addl(Address dst, int imm32) { + public final void addl(AMD64Address dst, int imm32) { prefix(dst); emitArithOperand(0x81, rax, dst, imm32); } @@ -337,7 +339,7 @@ emitArith(0x81, 0xC0, dst, imm32); } - public final void addl(Register dst, Address src) { + public final void addl(Register dst, AMD64Address src) { prefix(src, dst); emitByte(0x03); emitOperandHelper(dst, src); @@ -391,7 +393,7 @@ emitByte(0xC0 | encode); } - public final void addsd(Register dst, Address src) { + public final void addsd(Register dst, AMD64Address src) { assert dst.isFpu(); emitByte(0xF2); prefix(src, dst); @@ -409,7 +411,7 @@ emitByte(0xC0 | encode); } - public final void addss(Register dst, Address src) { + public final void addss(Register dst, AMD64Address src) { assert dst.isFpu(); emitByte(0xF3); prefix(src, dst); @@ -423,7 +425,7 @@ emitArith(0x81, 0xE0, dst, imm32); } - public final void andl(Register dst, Address src) { + public final void andl(Register dst, AMD64Address src) { prefix(src, dst); emitByte(0x23); emitOperandHelper(dst, src); @@ -441,7 +443,7 @@ emitByte(0xC0 | encode); } - public final void bsfq(Register dst, Address src) { + public final void bsfq(Register dst, AMD64Address src) { prefixq(src, dst); emitByte(0xBC); emitOperandHelper(dst, src); @@ -454,7 +456,7 @@ emitByte(0xC0 | encode); } - public final void bsrq(Register dst, Address src) { + public final void bsrq(Register dst, AMD64Address src) { prefixq(src, dst); emitByte(0xBD); emitOperandHelper(dst, src); @@ -467,7 +469,7 @@ emitByte(0xC0 | encode); } - public final void bsrl(Register dst, Address src) { + public final void bsrl(Register dst, AMD64Address src) { prefix(src, dst); emitByte(0xBD); emitOperandHelper(dst, src); @@ -490,7 +492,7 @@ emitByte(0xC0 | encode); } - public final void cmovl(ConditionFlag cc, Register dst, Address src) { + public final void cmovl(ConditionFlag cc, Register dst, AMD64Address src) { prefix(src, dst); emitByte(0x0F); emitByte(0x40 | cc.getValue()); @@ -507,7 +509,7 @@ emitArith(0x3B, 0xC0, dst, src); } - public final void cmpl(Register dst, Address src) { + public final void cmpl(Register dst, AMD64Address src) { prefix(src, dst); emitByte(0x3B); emitOperandHelper(dst, src); @@ -516,7 +518,7 @@ // The 32-bit cmpxchg compares the value at adr with the contents of X86.rax, // and stores reg into adr if so; otherwise, the value at adr is loaded into X86.rax,. // The ZF is set if the compared values were equal, and cleared otherwise. - public final void cmpxchgl(Register reg, Address adr) { // cmpxchg + public final void cmpxchgl(Register reg, AMD64Address adr) { // cmpxchg if ((Atomics & 2) != 0) { // caveat: no instructionmark, so this isn't relocatable. // Emit a synthetic, non-atomic, CAS equivalent. @@ -595,13 +597,13 @@ emitByte(0xC0 | encode); } - protected final void decl(Address dst) { + protected final void decl(AMD64Address dst) { prefix(dst); emitByte(0xFF); emitOperandHelper(rcx, dst); } - public final void divsd(Register dst, Address src) { + public final void divsd(Register dst, AMD64Address src) { assert dst.isFpu(); emitByte(0xF2); prefix(src, dst); @@ -620,7 +622,7 @@ emitByte(0xC0 | encode); } - public final void divss(Register dst, Address src) { + public final void divss(Register dst, AMD64Address src) { assert dst.isFpu(); emitByte(0xF3); prefix(src, dst); @@ -675,7 +677,7 @@ } } - protected final void incl(Address dst) { + protected final void incl(AMD64Address dst) { prefix(dst); emitByte(0xFF); emitOperandHelper(rax, dst); @@ -783,7 +785,7 @@ } } - public final void leaq(Register dst, Address src) { + public final void leaq(Register dst, AMD64Address src) { prefixq(src, dst); emitByte(0x8D); emitOperandHelper(dst, src); @@ -851,14 +853,14 @@ emitByte(0xC0 | dstenc << 3 | srcenc); } - public final void movb(Address dst, int imm8) { + public final void movb(AMD64Address dst, int imm8) { prefix(dst); emitByte(0xC6); emitOperandHelper(rax, dst); emitByte(imm8); } - public final void movb(Address dst, Register src) { + public final void movb(AMD64Address dst, Register src) { assert src.isByte() : "must have byte register"; prefix(dst, src); // , true) emitByte(0x88); @@ -896,20 +898,20 @@ emitByte(0xC0 | encode); } - public final void movl(Register dst, Address src) { + public final void movl(Register dst, AMD64Address src) { prefix(src, dst); emitByte(0x8B); emitOperandHelper(dst, src); } - public final void movl(Address dst, int imm32) { + public final void movl(AMD64Address dst, int imm32) { prefix(dst); emitByte(0xC7); emitOperandHelper(rax, dst); emitInt(imm32); } - public final void movl(Address dst, Register src) { + public final void movl(AMD64Address dst, Register src) { prefix(dst, src); emitByte(0x89); emitOperandHelper(src, dst); @@ -918,10 +920,10 @@ /** * New CPUs require use of movsd and movss to avoid partial register stall when loading from * memory. But for old Opteron use movlpd instead of movsd. The selection is done in - * {@link AMD64MacroAssembler#movdbl(Register, Address)} and + * {@link AMD64MacroAssembler#movdbl(Register, AMD64Address)} and * {@link AMD64MacroAssembler#movflt(Register, Register)}. */ - public final void movlpd(Register dst, Address src) { + public final void movlpd(Register dst, AMD64Address src) { assert dst.isFpu(); emitByte(0x66); prefix(src, dst); @@ -930,7 +932,7 @@ emitOperandHelper(dst, src); } - public final void movq(Register dst, Address src) { + public final void movq(Register dst, AMD64Address src) { if (dst.isFpu()) { emitByte(0xF3); prefixq(src, dst); @@ -950,7 +952,7 @@ emitByte(0xC0 | encode); } - public final void movq(Address dst, Register src) { + public final void movq(AMD64Address dst, Register src) { if (src.isFpu()) { emitByte(0x66); prefixq(dst, src); @@ -964,7 +966,7 @@ } } - public final void movsxb(Register dst, Address src) { + public final void movsxb(Register dst, AMD64Address src) { prefix(src, dst); emitByte(0x0F); emitByte(0xBE); @@ -988,7 +990,7 @@ emitByte(0xC0 | encode); } - public final void movsd(Register dst, Address src) { + public final void movsd(Register dst, AMD64Address src) { assert dst.isFpu(); emitByte(0xF2); prefix(src, dst); @@ -997,7 +999,7 @@ emitOperandHelper(dst, src); } - public final void movsd(Address dst, Register src) { + public final void movsd(AMD64Address dst, Register src) { assert src.isFpu(); emitByte(0xF2); prefix(dst, src); @@ -1016,7 +1018,7 @@ emitByte(0xC0 | encode); } - public final void movss(Register dst, Address src) { + public final void movss(Register dst, AMD64Address src) { assert dst.isFpu(); emitByte(0xF3); prefix(src, dst); @@ -1025,7 +1027,7 @@ emitOperandHelper(dst, src); } - public final void movss(Address dst, Register src) { + public final void movss(AMD64Address dst, Register src) { assert src.isFpu(); emitByte(0xF3); prefix(dst, src); @@ -1034,7 +1036,7 @@ emitOperandHelper(src, dst); } - public final void movswl(Register dst, Address src) { + public final void movswl(Register dst, AMD64Address src) { prefix(src, dst); emitByte(0x0F); emitByte(0xBF); @@ -1048,7 +1050,7 @@ emitByte(0xC0 | encode); } - public final void movw(Address dst, int imm16) { + public final void movw(AMD64Address dst, int imm16) { emitByte(0x66); // switch to 16-bit mode prefix(dst); emitByte(0xC7); @@ -1056,21 +1058,21 @@ emitShort(imm16); } - public final void movw(Address dst, Register src) { + public final void movw(AMD64Address dst, Register src) { emitByte(0x66); prefix(dst, src); emitByte(0x89); emitOperandHelper(src, dst); } - public final void movzxl(Register dst, Address src) { + public final void movzxl(Register dst, AMD64Address src) { prefix(src, dst); emitByte(0x0F); emitByte(0xB7); emitOperandHelper(dst, src); } - public final void mulsd(Register dst, Address src) { + public final void mulsd(Register dst, AMD64Address src) { assert dst.isFpu(); emitByte(0xF2); prefix(src, dst); @@ -1090,7 +1092,7 @@ emitByte(0xC0 | encode); } - public final void mulss(Register dst, Address src) { + public final void mulss(Register dst, AMD64Address src) { assert dst.isFpu(); emitByte(0xF3); @@ -1331,7 +1333,7 @@ emitArith(0x81, 0xC8, dst, imm32); } - public final void orl(Register dst, Address src) { + public final void orl(Register dst, AMD64Address src) { prefix(src, dst); emitByte(0x0B); emitOperandHelper(dst, src); @@ -1342,7 +1344,7 @@ emitArith(0x0B, 0xC0, dst, src); } - public final void popcntl(Register dst, Address src) { + public final void popcntl(Register dst, AMD64Address src) { emitByte(0xF3); prefix(src, dst); emitByte(0x0F); @@ -1358,7 +1360,7 @@ emitByte(0xC0 | encode); } - public final void popcntq(Register dst, Address src) { + public final void popcntq(Register dst, AMD64Address src) { emitByte(0xF3); prefixq(src, dst); emitByte(0x0F); @@ -1457,7 +1459,7 @@ emitByte(0xC0 | encode); } - public final void subl(Address dst, int imm32) { + public final void subl(AMD64Address dst, int imm32) { prefix(dst); if (isByte(imm32)) { emitByte(0x83); @@ -1475,7 +1477,7 @@ emitArith(0x81, 0xE8, dst, imm32); } - public final void subl(Register dst, Address src) { + public final void subl(Register dst, AMD64Address src) { prefix(src, dst); emitByte(0x2B); emitOperandHelper(dst, src); @@ -1496,7 +1498,7 @@ emitByte(0xC0 | encode); } - public final void subsd(Register dst, Address src) { + public final void subsd(Register dst, AMD64Address src) { assert dst.isFpu(); emitByte(0xF2); @@ -1516,7 +1518,7 @@ emitByte(0xC0 | encode); } - public final void subss(Register dst, Address src) { + public final void subss(Register dst, AMD64Address src) { assert dst.isFpu(); emitByte(0xF3); @@ -1546,13 +1548,13 @@ emitArith(0x85, 0xC0, dst, src); } - public final void testl(Register dst, Address src) { + public final void testl(Register dst, AMD64Address src) { prefix(src, dst); emitByte(0x85); emitOperandHelper(dst, src); } - public final void ucomisd(Register dst, Address src) { + public final void ucomisd(Register dst, AMD64Address src) { assert dst.isFpu(); emitByte(0x66); ucomiss(dst, src); @@ -1565,7 +1567,7 @@ ucomiss(dst, src); } - public final void ucomiss(Register dst, Address src) { + public final void ucomiss(Register dst, AMD64Address src) { assert dst.isFpu(); prefix(src, dst); @@ -1588,7 +1590,7 @@ emitArith(0x81, 0xF0, dst, imm32); } - public final void xorl(Register dst, Address src) { + public final void xorl(Register dst, AMD64Address src) { prefix(src, dst); emitByte(0x33); emitOperandHelper(dst, src); @@ -1604,7 +1606,7 @@ andps(dst, src); } - public final void andpd(Register dst, Address src) { + public final void andpd(Register dst, AMD64Address src) { emitByte(0x66); andps(dst, src); } @@ -1617,7 +1619,7 @@ emitByte(0xC0 | encode); } - public final void andps(Register dst, Address src) { + public final void andps(Register dst, AMD64Address src) { assert dst.isFpu(); prefix(src, dst); emitByte(0x0F); @@ -1630,7 +1632,7 @@ orps(dst, src); } - public final void orpd(Register dst, Address src) { + public final void orpd(Register dst, AMD64Address src) { emitByte(0x66); orps(dst, src); } @@ -1643,7 +1645,7 @@ emitByte(0xC0 | encode); } - public final void orps(Register dst, Address src) { + public final void orps(Register dst, AMD64Address src) { assert dst.isFpu(); prefix(src, dst); emitByte(0x0F); @@ -1656,7 +1658,7 @@ xorps(dst, src); } - public final void xorpd(Register dst, Address src) { + public final void xorpd(Register dst, AMD64Address src) { emitByte(0x66); xorps(dst, src); } @@ -1669,7 +1671,7 @@ emitByte(0xC0 | encode); } - public final void xorps(Register dst, Address src) { + public final void xorps(Register dst, AMD64Address src) { assert dst.isFpu(); prefix(src, dst); emitByte(0x0F); @@ -1781,7 +1783,7 @@ return isRegister(value) && asRegister(value).encoding >= MinEncodingNeedsRex; } - private void prefix(Address adr) { + private void prefix(AMD64Address adr) { if (needsRex(adr.getBase())) { if (needsRex(adr.getIndex())) { emitByte(Prefix.REXXB); @@ -1795,7 +1797,7 @@ } } - private void prefixq(Address adr) { + private void prefixq(AMD64Address adr) { if (needsRex(adr.getBase())) { if (needsRex(adr.getIndex())) { emitByte(Prefix.REXWXB); @@ -1811,7 +1813,7 @@ } } - private void prefix(Address adr, Register reg) { + private void prefix(AMD64Address adr, Register reg) { if (reg.encoding < 8) { if (needsRex(adr.getBase())) { if (needsRex(adr.getIndex())) { @@ -1843,7 +1845,7 @@ } } - private void prefixq(Address adr, Register src) { + private void prefixq(AMD64Address adr, Register src) { if (src.encoding < 8) { if (needsRex(adr.getBase())) { if (needsRex(adr.getIndex())) { @@ -1880,7 +1882,7 @@ emitArith(0x81, 0xC0, dst, imm32); } - public final void addq(Register dst, Address src) { + public final void addq(Register dst, AMD64Address src) { prefixq(src, dst); emitByte(0x03); emitOperandHelper(dst, src); @@ -1896,7 +1898,7 @@ emitArith(0x81, 0xE0, dst, imm32); } - public final void andq(Register dst, Address src) { + public final void andq(Register dst, AMD64Address src) { prefixq(src, dst); emitByte(0x23); emitOperandHelper(dst, src); @@ -1925,7 +1927,7 @@ emitByte(0xC0 | encode); } - public final void cmovq(ConditionFlag cc, Register dst, Address src) { + public final void cmovq(ConditionFlag cc, Register dst, AMD64Address src) { prefixq(src, dst); emitByte(0x0F); emitByte(0x40 | cc.getValue()); @@ -1942,13 +1944,13 @@ emitArith(0x3B, 0xC0, dst, src); } - public final void cmpq(Register dst, Address src) { + public final void cmpq(Register dst, AMD64Address src) { prefixq(src, dst); emitByte(0x3B); emitOperandHelper(dst, src); } - public final void cmpxchgq(Register reg, Address adr) { + public final void cmpxchgq(Register reg, AMD64Address adr) { prefixq(adr, reg); emitByte(0x0F); emitByte(0xB1); @@ -1998,7 +2000,7 @@ emitByte(0xC8 | encode); } - protected final void decq(Address dst) { + protected final void decq(AMD64Address dst) { prefixq(dst); emitByte(0xFF); emitOperandHelper(rcx, dst); @@ -2073,14 +2075,14 @@ } } - public final void movslq(Address dst, int imm32) { + public final void movslq(AMD64Address dst, int imm32) { prefixq(dst); emitByte(0xC7); emitOperandHelper(rax, dst); emitInt(imm32); } - public final void movslq(Register dst, Address src) { + public final void movslq(Register dst, AMD64Address src) { prefixq(src, dst); emitByte(0x63); emitOperandHelper(dst, src); @@ -2103,7 +2105,7 @@ emitArith(0x81, 0xC8, dst, imm32); } - public final void orq(Register dst, Address src) { + public final void orq(Register dst, AMD64Address src) { prefixq(src, dst); emitByte(0x0B); emitOperandHelper(dst, src); @@ -2171,7 +2173,7 @@ emitArith(0x81, 0xE8, dst, imm32); } - public final void subq(Register dst, Address src) { + public final void subq(Register dst, AMD64Address src) { prefixq(src, dst); emitByte(0x2B); emitOperandHelper(dst, src); @@ -2203,7 +2205,7 @@ emitArith(0x85, 0xC0, dst, src); } - public final void testq(Register dst, Address src) { + public final void testq(Register dst, AMD64Address src) { prefixq(src, dst); emitByte(0x85); emitOperandHelper(dst, src); @@ -2219,7 +2221,7 @@ emitArith(0x33, 0xC0, dst, src); } - public final void xorq(Register dst, Address src) { + public final void xorq(Register dst, AMD64Address src) { prefixq(src, dst); emitByte(0x33); emitOperandHelper(dst, src); @@ -2238,7 +2240,7 @@ // the code where this idiom is used, in particular the // orderAccess code. lock(); - addl(new Address(Word, RSP, 0), 0); // Assert the lock# signal here + addl(new AMD64Address(Word, RSP, 0), 0); // Assert the lock# signal here } } } @@ -2279,7 +2281,7 @@ } public void nullCheck(Register r) { - testl(AMD64.rax, new Address(Word, r.asValue(Word), 0)); + testl(AMD64.rax, new AMD64Address(Word, r.asValue(Word), 0)); } @Override @@ -2315,7 +2317,7 @@ emitByte(b2 + i); } - public final void fld(Address src) { + public final void fld(AMD64Address src) { emitByte(0xDD); emitOperandHelper(rax, src); } @@ -2335,7 +2337,7 @@ emitByte(0xF1); } - public final void fstp(Address src) { + public final void fstp(AMD64Address src) { emitByte(0xDD); emitOperandHelper(rbx, src); } @@ -2358,4 +2360,14 @@ public final void fstp(int i) { emitx87(0xDD, 0xD8, i); } + + @Override + public AMD64Address makeAddress(Kind kind, Value base, int displacement) { + return new AMD64Address(kind, base, displacement); + } + + @Override + public AMD64Address getPlaceholder() { + return Placeholder; + } }
--- a/graal/com.oracle.graal.asm.amd64/src/com/oracle/graal/asm/amd64/AMD64MacroAssembler.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.asm.amd64/src/com/oracle/graal/asm/amd64/AMD64MacroAssembler.java Fri Feb 22 12:15:14 2013 +0100 @@ -1,5 +1,5 @@ /* - * Copyright (c) 2009, 2011, Oracle and/or its affiliates. All rights reserved. + * Copyright (c) 2009, 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 @@ -79,7 +79,7 @@ } } - public final void movptr(Address dst, int src) { + public final void movptr(AMD64Address dst, int src) { movslq(dst, src); } @@ -87,7 +87,7 @@ cmpq(src1, src2); } - public final void cmpptr(Register src1, Address src2) { + public final void cmpptr(Register src1, AMD64Address src2) { cmpq(src1, src2); } @@ -110,7 +110,7 @@ } } - public final void decrementl(Address dst, int value) { + public final void decrementl(AMD64Address dst, int value) { if (value == Integer.MIN_VALUE) { subl(dst, value); return; @@ -148,7 +148,7 @@ } } - public final void incrementl(Address dst, int value) { + public final void incrementl(AMD64Address dst, int value) { if (value == Integer.MIN_VALUE) { addl(dst, value); return; @@ -189,12 +189,12 @@ } } - public final void movflt(Register dst, Address src) { + public final void movflt(Register dst, AMD64Address src) { assert dst.isFpu(); movss(dst, src); } - public final void movflt(Address dst, Register src) { + public final void movflt(AMD64Address dst, Register src) { assert src.isFpu(); movss(dst, src); } @@ -208,7 +208,7 @@ } } - public final void movdbl(Register dst, Address src) { + public final void movdbl(Register dst, AMD64Address src) { assert dst.isFpu(); if (UseXmmLoadAndClearUpper) { movsd(dst, src); @@ -221,8 +221,8 @@ * Non-atomic write of a 64-bit constant to memory. Do not use if the address might be a * volatile field! */ - public final void movlong(Address dst, long src) { - Address high = new Address(dst.getKind(), dst.getBase(), dst.getIndex(), dst.getScale(), dst.getDisplacement() + 4); + public final void movlong(AMD64Address dst, long src) { + AMD64Address high = new AMD64Address(dst.getKind(), dst.getBase(), dst.getIndex(), dst.getScale(), dst.getDisplacement() + 4); movl(dst, (int) (src & 0xFFFFFFFF)); movl(high, (int) (src >> 32)); } @@ -230,7 +230,7 @@ public final void flog(Register dest, Register value, boolean base10) { assert dest.isFpu() && value.isFpu(); - Address tmp = new Address(Kind.Double, AMD64.RSP); + AMD64Address tmp = new AMD64Address(Kind.Double, AMD64.RSP); if (base10) { fldlg2(); } else { @@ -244,34 +244,34 @@ } public final void fsin(Register dest, Register value) { - Address tmp = trigPrologue(value); + AMD64Address tmp = trigPrologue(value); fsin(); trigEpilogue(dest, tmp); } public final void fcos(Register dest, Register value) { - Address tmp = trigPrologue(value); + AMD64Address tmp = trigPrologue(value); fcos(); trigEpilogue(dest, tmp); } public final void ftan(Register dest, Register value) { - Address tmp = trigPrologue(value); + AMD64Address tmp = trigPrologue(value); fptan(); fstp(0); // ftan pushes 1.0 in addition to the actual result, pop trigEpilogue(dest, tmp); } - private Address trigPrologue(Register value) { + private AMD64Address trigPrologue(Register value) { assert value.isFpu(); - Address tmp = new Address(Kind.Double, AMD64.RSP); + AMD64Address tmp = new AMD64Address(Kind.Double, AMD64.RSP); subq(AMD64.rsp, 8); movsd(tmp, value); fld(tmp); return tmp; } - private void trigEpilogue(Register dest, Address tmp) { + private void trigEpilogue(Register dest, AMD64Address tmp) { assert dest.isFpu(); fstp(tmp); movsd(dest, tmp); @@ -289,7 +289,7 @@ RegisterValue frame = frameRegister.asValue(); for (Register r : csl.registers) { int offset = csl.offsetOf(r); - movq(new Address(target.wordKind, frame, frameToCSA + offset), r); + movq(new AMD64Address(target.wordKind, frame, frameToCSA + offset), r); } } @@ -297,7 +297,7 @@ RegisterValue frame = frameRegister.asValue(); for (Register r : csl.registers) { int offset = csl.offsetOf(r); - movq(r, new Address(target.wordKind, frame, frameToCSA + offset)); + movq(r, new AMD64Address(target.wordKind, frame, frameToCSA + offset)); } } }
--- a/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Fri Feb 22 12:15:14 2013 +0100 @@ -23,6 +23,8 @@ package com.oracle.graal.asm.ptx; import com.oracle.graal.api.code.*; +import com.oracle.graal.api.meta.*; +import com.oracle.graal.ptx.*; public class PTXAssembler extends AbstractPTXAssembler { @@ -179,55 +181,55 @@ emitString("ld.global.b8" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } - public final void ld_global_b16(Register d, Register a, int immOff) { + public final void ld_global_b16(Register d, Register a, long immOff) { emitString("ld.global.b16" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } - public final void ld_global_b32(Register d, Register a, int immOff) { + public final void ld_global_b32(Register d, Register a, long immOff) { emitString("ld.global.b32" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } - public final void ld_global_b64(Register d, Register a, int immOff) { + public final void ld_global_b64(Register d, Register a, long immOff) { emitString("ld.global.b64" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } - public final void ld_global_u8(Register d, Register a, int immOff) { + public final void ld_global_u8(Register d, Register a, long immOff) { emitString("ld.global.u8" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } - public final void ld_global_u16(Register d, Register a, int immOff) { + public final void ld_global_u16(Register d, Register a, long immOff) { emitString("ld.global.u16" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } - public final void ld_global_u32(Register d, Register a, int immOff) { + public final void ld_global_u32(Register d, Register a, long immOff) { emitString("ld.global.u32" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } - public final void ld_global_u64(Register d, Register a, int immOff) { + public final void ld_global_u64(Register d, Register a, long immOff) { emitString("ld.global.u64" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } - public final void ld_global_s8(Register d, Register a, int immOff) { + public final void ld_global_s8(Register d, Register a, long immOff) { emitString("ld.global.s8" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } - public final void ld_global_s16(Register d, Register a, int immOff) { + public final void ld_global_s16(Register d, Register a, long immOff) { emitString("ld.global.s16" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } - public final void ld_global_s32(Register d, Register a, int immOff) { + public final void ld_global_s32(Register d, Register a, long immOff) { emitString("ld.global.s32" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } - public final void ld_global_s64(Register d, Register a, int immOff) { + public final void ld_global_s64(Register d, Register a, long immOff) { emitString("ld.global.s64" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } - public final void ld_global_f32(Register d, Register a, int immOff) { + public final void ld_global_f32(Register d, Register a, long immOff) { emitString("ld.global.f32" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } - public final void ld_global_f64(Register d, Register a, int immOff) { + public final void ld_global_f64(Register d, Register a, long immOff) { emitString("ld.global.f64" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } @@ -635,59 +637,59 @@ emitString("shr.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); } - public final void st_global_b8(Register a, int immOff, Register b) { + 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, int immOff, Register b) { + public final void st_global_b16(Register a, long immOff, Register b) { emitString("st.global.b16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } - public final void st_global_b32(Register a, int immOff, Register b) { + 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, int immOff, Register b) { + public final void st_global_b64(Register a, long immOff, Register b) { emitString("st.global.b64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } - public final void st_global_u8(Register a, int immOff, Register b) { + 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, int immOff, Register b) { + public final void st_global_u16(Register a, long immOff, Register b) { emitString("st.global.u16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } - public final void st_global_u32(Register a, int immOff, Register b) { + 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, int immOff, Register b) { + public final void st_global_u64(Register a, long immOff, Register b) { emitString("st.global.u64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } - public final void st_global_s8(Register a, int immOff, Register b) { + public final void st_global_s8(Register a, long immOff, Register b) { emitString("st.global.s8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } - public final void st_global_s16(Register a, int immOff, Register b) { + public final void st_global_s16(Register a, long immOff, Register b) { emitString("st.global.s16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } - public final void st_global_s32(Register a, int immOff, Register b) { + 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, int immOff, Register b) { + public final void st_global_s64(Register a, long immOff, Register b) { emitString("st.global.s64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } - public final void st_global_f32(Register a, int immOff, Register b) { + 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, int immOff, Register b) { + public final void st_global_f64(Register a, long immOff, Register b) { emitString("st.global.f64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); } @@ -738,4 +740,15 @@ public final void sub_sat_s32(Register d, int s32, Register b) { emitString("sub.sat.s32" + " " + "%r" + d.encoding() + ", " + s32 + ", %r" + b.encoding() + ";" + ""); } + + @Override + public PTXAddress makeAddress(Kind kind, Value base, int displacement) { + return new PTXAddress(kind, base, displacement); + } + + @Override + public PTXAddress getPlaceholder() { + // TODO Auto-generated method stub + return null; + } }
--- a/graal/com.oracle.graal.asm.sparc/src/com/oracle/graal/asm/sparc/SPARCAssembler.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.asm.sparc/src/com/oracle/graal/asm/sparc/SPARCAssembler.java Fri Feb 22 12:15:14 2013 +0100 @@ -23,6 +23,7 @@ package com.oracle.graal.asm.sparc; import com.oracle.graal.api.code.*; +import com.oracle.graal.api.meta.*; import com.oracle.graal.asm.*; import com.oracle.graal.sparc.*; @@ -52,4 +53,16 @@ protected void patchJumpTarget(int branch, int jumpTarget) { // SPARC: Implement patching of jump target. } + + @Override + public Address makeAddress(Kind kind, Value base, int displacement) { + // SPARC: Implement address calculation. + return null; + } + + @Override + public Address getPlaceholder() { + // SPARC: Implement address patching. + return null; + } }
--- a/graal/com.oracle.graal.asm/src/com/oracle/graal/asm/AbstractAssembler.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.asm/src/com/oracle/graal/asm/AbstractAssembler.java Fri Feb 22 12:15:14 2013 +0100 @@ -25,6 +25,7 @@ import java.nio.*; import com.oracle.graal.api.code.*; +import com.oracle.graal.api.meta.*; /** * The platform-independent base class for the assembler. @@ -83,4 +84,15 @@ protected final void emitString0(String x) { codeBuffer.emitString0(x); } + + /** + * This is used by the TargetMethodAssembler to convert a {@link StackSlot} to an + * {@link Address}. + */ + public abstract Address makeAddress(Kind kind, Value base, int displacement); + + /** + * Returns a target specific placeholder address that can be used for code patching. + */ + public abstract Address getPlaceholder(); }
--- a/graal/com.oracle.graal.compiler.amd64/src/com/oracle/graal/compiler/amd64/AMD64LIRGenerator.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.compiler.amd64/src/com/oracle/graal/compiler/amd64/AMD64LIRGenerator.java Fri Feb 22 12:15:14 2013 +0100 @@ -137,6 +137,11 @@ } @Override + public Address makeAddress(Kind kind, Value base, int displacement) { + return new AMD64Address(kind, base, displacement); + } + + @Override public Address makeAddress(LocationNode location, ValueNode object) { Value base = operand(object); Value index = Value.ILLEGAL; @@ -182,7 +187,7 @@ } } - return new Address(location.getValueKind(), base, index, Address.Scale.fromInt(scale), displacement); + return new AMD64Address(location.getValueKind(), base, index, AMD64Address.Scale.fromInt(scale), displacement); } @Override @@ -886,9 +891,9 @@ if (isConstant(index) && NumUtil.isInt(asConstant(index).asLong() + displacement)) { assert !runtime.needsDataPatch(asConstant(index)); displacement += (int) asConstant(index).asLong(); - address = new Address(kind, load(operand(node.object())), displacement); + address = new AMD64Address(kind, load(operand(node.object())), displacement); } else { - address = new Address(kind, load(operand(node.object())), load(index), Address.Scale.Times1, displacement); + address = new AMD64Address(kind, load(operand(node.object())), load(index), AMD64Address.Scale.Times1, displacement); } RegisterValue rax = AMD64.rax.asValue(kind);
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java Fri Feb 22 12:15:14 2013 +0100 @@ -25,8 +25,8 @@ import static com.oracle.graal.api.code.ValueUtil.*; import static com.oracle.graal.lir.ptx.PTXArithmetic.*; +import static com.oracle.graal.lir.ptx.PTXBitManipulationOp.IntrinsicOpcode.*; import static com.oracle.graal.lir.ptx.PTXCompare.*; -import static com.oracle.graal.lir.ptx.PTXBitManipulationOp.IntrinsicOpcode.*; import com.oracle.graal.api.code.*; import com.oracle.graal.api.meta.*; @@ -36,7 +36,11 @@ import com.oracle.graal.graph.*; import com.oracle.graal.lir.*; import com.oracle.graal.lir.StandardOp.JumpOp; -import com.oracle.graal.lir.ptx.PTXBitManipulationOp; +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.*; import com.oracle.graal.lir.ptx.PTXCompare.CompareOp; import com.oracle.graal.lir.ptx.PTXControlFlow.BranchOp; import com.oracle.graal.lir.ptx.PTXControlFlow.ReturnOp; @@ -45,11 +49,10 @@ 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.Condition; -import com.oracle.graal.nodes.calc.ConvertNode; -import com.oracle.graal.nodes.extended.IndexedLocationNode; -import com.oracle.graal.nodes.extended.LocationNode; +import com.oracle.graal.nodes.calc.*; +import com.oracle.graal.nodes.extended.*; import com.oracle.graal.nodes.java.*; +import com.oracle.graal.ptx.*; /** * This class implements the PTX specific portion of the LIR generator. @@ -97,52 +100,44 @@ } @Override + public Address makeAddress(Kind kind, Value base, int displacement) { + return new PTXAddress(kind, base, displacement); + } + + @Override public Address makeAddress(LocationNode location, ValueNode object) { Value base = operand(object); - Value index = Value.ILLEGAL; - int scale = 1; - int displacement = location.displacement(); + long displacement = location.displacement(); if (isConstant(base)) { if (asConstant(base).isNull()) { base = Value.ILLEGAL; } else if (asConstant(base).getKind() != Kind.Object) { - long newDisplacement = displacement + asConstant(base).asLong(); - if (NumUtil.isInt(newDisplacement)) { - assert !runtime.needsDataPatch(asConstant(base)); - displacement = (int) newDisplacement; - base = Value.ILLEGAL; - } else { - Value newBase = newVariable(Kind.Long); - emitMove(base, newBase); - base = newBase; - } + displacement += asConstant(base).asLong(); + base = Value.ILLEGAL; } } if (location instanceof IndexedLocationNode) { IndexedLocationNode indexedLoc = (IndexedLocationNode) location; - index = operand(indexedLoc.index()); - scale = indexedLoc.indexScaling(); + Value index = operand(indexedLoc.index()); + int scale = indexedLoc.indexScaling(); if (isConstant(index)) { - long newDisplacement = displacement + asConstant(index).asLong() * scale; - // only use the constant index if the resulting displacement fits into a 32 bit - // offset - if (NumUtil.isInt(newDisplacement)) { - displacement = (int) newDisplacement; - index = Value.ILLEGAL; + displacement += asConstant(index).asLong() * scale; + } else { + if (scale != 1) { + index = emitMul(index, Constant.forInt(scale)); + } + if (base == Value.ILLEGAL) { + base = index; } else { - // create a temporary variable for the index, the pointer load cannot handle a - // constant index - Value newIndex = newVariable(Kind.Long); - emitMove(index, newIndex); - index = newIndex; + base = emitAdd(base, index); } } } - return new Address(location.getValueKind(), base, index, Address.Scale.fromInt(scale), displacement); + return new PTXAddress(location.getValueKind(), base, displacement); } @Override
--- a/graal/com.oracle.graal.compiler.sparc/src/com/oracle/graal/compiler/sparc/SPARCLIRGenerator.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.compiler.sparc/src/com/oracle/graal/compiler/sparc/SPARCLIRGenerator.java Fri Feb 22 12:15:14 2013 +0100 @@ -215,6 +215,11 @@ } @Override + public Address makeAddress(Kind kind, Value base, int displacement) { + return null; + } + + @Override public Address makeAddress(LocationNode location, ValueNode object) { // SPARC: Auto-generated method stub return null;
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotBackend.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotBackend.java Fri Feb 22 12:15:14 2013 +0100 @@ -109,8 +109,8 @@ public void visitExceptionObject(ExceptionObjectNode x) { HotSpotVMConfig config = runtime().config; RegisterValue thread = runtime().threadRegister().asValue(); - Address exceptionAddress = new Address(Kind.Object, thread, config.threadExceptionOopOffset); - Address pcAddress = new Address(Kind.Long, thread, config.threadExceptionPcOffset); + Address exceptionAddress = new AMD64Address(Kind.Object, thread, config.threadExceptionOopOffset); + Address pcAddress = new AMD64Address(Kind.Long, thread, config.threadExceptionPcOffset); Value exception = emitLoad(exceptionAddress, false); emitStore(exceptionAddress, Constant.NULL_OBJECT, false); emitStore(pcAddress, Constant.LONG_0, false); @@ -132,9 +132,9 @@ if (ValueUtil.isConstant(index) && NumUtil.isInt(ValueUtil.asConstant(index).asLong() + disp)) { assert !runtime.needsDataPatch(asConstant(index)); disp += (int) ValueUtil.asConstant(index).asLong(); - address = new Address(kind, load(operand(x.object())), disp); + address = new AMD64Address(kind, load(operand(x.object())), disp); } else { - address = new Address(kind, load(operand(x.object())), load(index), Address.Scale.Times1, disp); + address = new AMD64Address(kind, load(operand(x.object())), load(index), AMD64Address.Scale.Times1, disp); } RegisterValue rax = AMD64.rax.asValue(kind); @@ -187,7 +187,7 @@ disp -= frameSize; } tasm.blockComment("[stack overflow check]"); - asm.movq(new Address(asm.target.wordKind, AMD64.RSP, -disp), AMD64.rax); + asm.movq(new AMD64Address(asm.target.wordKind, AMD64.RSP, -disp), AMD64.rax); } } } @@ -208,7 +208,7 @@ if (GraalOptions.ZapStackOnMethodEntry) { final int intSize = 4; for (int i = 0; i < frameSize / intSize; ++i) { - asm.movl(new Address(Kind.Int, rsp.asValue(), i * intSize), 0xC1C1C1C1); + asm.movl(new AMD64Address(Kind.Int, rsp.asValue(), i * intSize), 0xC1C1C1C1); } } CalleeSaveLayout csl = frameMap.registerConfig.getCalleeSaveLayout(); @@ -248,12 +248,12 @@ if (config.isPollingPageFar) { asm.movq(scratch, config.safepointPollingAddress + offset); tasm.recordMark(Marks.MARK_POLL_RETURN_FAR); - asm.movq(scratch, new Address(tasm.target.wordKind, scratch.asValue())); + asm.movq(scratch, new AMD64Address(tasm.target.wordKind, scratch.asValue())); } else { tasm.recordMark(Marks.MARK_POLL_RETURN_NEAR); // The C++ code transforms the polling page offset into an RIP displacement // to the real address at that offset in the polling page. - asm.movq(scratch, new Address(tasm.target.wordKind, rip.asValue(), offset)); + asm.movq(scratch, new AMD64Address(tasm.target.wordKind, rip.asValue(), offset)); } } } @@ -294,7 +294,7 @@ Register inlineCacheKlass = rax; // see definition of IC_Klass in // c1_LIRAssembler_x86.cpp Register receiver = asRegister(cc.getArgument(0)); - Address src = new Address(target.wordKind, receiver.asValue(), config.hubOffset); + AMD64Address src = new AMD64Address(target.wordKind, receiver.asValue(), config.hubOffset); asm.cmpq(inlineCacheKlass, src); asm.jcc(ConditionFlag.NotEqual, unverifiedStub);
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64SafepointOp.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64SafepointOp.java Fri Feb 22 12:15:14 2013 +0100 @@ -25,6 +25,7 @@ import static com.oracle.graal.amd64.AMD64.*; import static com.oracle.graal.phases.GraalOptions.*; +import com.oracle.graal.amd64.*; import com.oracle.graal.api.code.*; import com.oracle.graal.asm.amd64.*; import com.oracle.graal.hotspot.*; @@ -58,13 +59,13 @@ asm.movq(scratch, config.safepointPollingAddress + offset); tasm.recordMark(Marks.MARK_POLL_FAR); tasm.recordSafepoint(pos, state); - asm.movq(scratch, new Address(tasm.target.wordKind, scratch.asValue())); + asm.movq(scratch, new AMD64Address(tasm.target.wordKind, scratch.asValue())); } else { tasm.recordMark(Marks.MARK_POLL_NEAR); tasm.recordSafepoint(pos, state); // The C++ code transforms the polling page offset into an RIP displacement // to the real address at that offset in the polling page. - asm.movq(scratch, new Address(tasm.target.wordKind, rip.asValue(), offset)); + asm.movq(scratch, new AMD64Address(tasm.target.wordKind, rip.asValue(), offset)); } } }
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/CurrentThread.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/CurrentThread.java Fri Feb 22 12:15:14 2013 +0100 @@ -42,7 +42,7 @@ public void generate(LIRGeneratorTool gen) { HotSpotGraalRuntime runtime = HotSpotGraalRuntime.getInstance(); Register thread = runtime.getRuntime().threadRegister(); - gen.setResult(this, gen.emitLoad(new Address(Kind.Object, thread.asValue(gen.target().wordKind), runtime.getConfig().threadObjectOffset), false)); + gen.setResult(this, gen.emitLoad(gen.makeAddress(Kind.Object, thread.asValue(gen.target().wordKind), runtime.getConfig().threadObjectOffset), false)); } @NodeIntrinsic
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/TailcallNode.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/TailcallNode.java Fri Feb 22 12:15:14 2013 +0100 @@ -69,7 +69,7 @@ parameters.add(frameState.localAt(slot)); } Value[] args = gen.visitInvokeArguments(cc, parameters); - Value entry = gen.emitLoad(new Address(Kind.Long, gen.operand(target), config.nmethodEntryOffset), false); + Value entry = gen.emitLoad(gen.makeAddress(Kind.Long, gen.operand(target), config.nmethodEntryOffset), false); HotSpotLIRGenerator hsgen = (HotSpotLIRGenerator) gen; hsgen.emitTailcall(args, entry); }
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/WriteBarrier.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/WriteBarrier.java Fri Feb 22 12:15:14 2013 +0100 @@ -22,7 +22,6 @@ */ package com.oracle.graal.hotspot.nodes; -import com.oracle.graal.api.code.*; import com.oracle.graal.api.meta.*; import com.oracle.graal.hotspot.*; import com.oracle.graal.nodes.*; @@ -46,6 +45,6 @@ } else { base = gen.emitAdd(base, Constant.forLong(config.cardtableStartAddress)); } - gen.emitStore(new Address(Kind.Boolean, base, displacement), Constant.FALSE, false); + gen.emitStore(gen.makeAddress(Kind.Boolean, base, displacement), Constant.FALSE, false); } }
--- a/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Arithmetic.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Arithmetic.java Fri Feb 22 12:15:14 2013 +0100 @@ -400,46 +400,46 @@ case LSHR: masm.sarq(asLongReg(dst), tasm.asIntConst(src) & 63); break; case LUSHR:masm.shrq(asLongReg(dst), tasm.asIntConst(src) & 63); break; - case FADD: masm.addss(asFloatReg(dst), tasm.asFloatConstRef(src)); break; - case FSUB: masm.subss(asFloatReg(dst), tasm.asFloatConstRef(src)); break; - case FMUL: masm.mulss(asFloatReg(dst), tasm.asFloatConstRef(src)); break; - case FAND: masm.andps(asFloatReg(dst), tasm.asFloatConstRef(src, 16)); break; - case FOR: masm.orps(asFloatReg(dst), tasm.asFloatConstRef(src, 16)); break; - case FXOR: masm.xorps(asFloatReg(dst), tasm.asFloatConstRef(src, 16)); break; - case FDIV: masm.divss(asFloatReg(dst), tasm.asFloatConstRef(src)); break; + case FADD: masm.addss(asFloatReg(dst), (AMD64Address) tasm.asFloatConstRef(src)); break; + case FSUB: masm.subss(asFloatReg(dst), (AMD64Address) tasm.asFloatConstRef(src)); break; + case FMUL: masm.mulss(asFloatReg(dst), (AMD64Address) tasm.asFloatConstRef(src)); break; + case FAND: masm.andps(asFloatReg(dst), (AMD64Address) tasm.asFloatConstRef(src, 16)); break; + case FOR: masm.orps(asFloatReg(dst), (AMD64Address) tasm.asFloatConstRef(src, 16)); break; + case FXOR: masm.xorps(asFloatReg(dst), (AMD64Address) tasm.asFloatConstRef(src, 16)); break; + case FDIV: masm.divss(asFloatReg(dst), (AMD64Address) tasm.asFloatConstRef(src)); break; - case DADD: masm.addsd(asDoubleReg(dst), tasm.asDoubleConstRef(src)); break; - case DSUB: masm.subsd(asDoubleReg(dst), tasm.asDoubleConstRef(src)); break; - case DMUL: masm.mulsd(asDoubleReg(dst), tasm.asDoubleConstRef(src)); break; - case DDIV: masm.divsd(asDoubleReg(dst), tasm.asDoubleConstRef(src)); break; - case DAND: masm.andpd(asDoubleReg(dst), tasm.asDoubleConstRef(src, 16)); break; - case DOR: masm.orpd(asDoubleReg(dst), tasm.asDoubleConstRef(src, 16)); break; - case DXOR: masm.xorpd(asDoubleReg(dst), tasm.asDoubleConstRef(src, 16)); break; + case DADD: masm.addsd(asDoubleReg(dst), (AMD64Address) tasm.asDoubleConstRef(src)); break; + case DSUB: masm.subsd(asDoubleReg(dst), (AMD64Address) tasm.asDoubleConstRef(src)); break; + case DMUL: masm.mulsd(asDoubleReg(dst), (AMD64Address) tasm.asDoubleConstRef(src)); break; + case DDIV: masm.divsd(asDoubleReg(dst), (AMD64Address) tasm.asDoubleConstRef(src)); break; + case DAND: masm.andpd(asDoubleReg(dst), (AMD64Address) tasm.asDoubleConstRef(src, 16)); break; + case DOR: masm.orpd(asDoubleReg(dst), (AMD64Address) tasm.asDoubleConstRef(src, 16)); break; + case DXOR: masm.xorpd(asDoubleReg(dst), (AMD64Address) tasm.asDoubleConstRef(src, 16)); break; default: throw GraalInternalError.shouldNotReachHere(); } } else { switch (opcode) { - case IADD: masm.addl(asIntReg(dst), tasm.asIntAddr(src)); break; - case ISUB: masm.subl(asIntReg(dst), tasm.asIntAddr(src)); break; - case IAND: masm.andl(asIntReg(dst), tasm.asIntAddr(src)); break; - case IOR: masm.orl(asIntReg(dst), tasm.asIntAddr(src)); break; - case IXOR: masm.xorl(asIntReg(dst), tasm.asIntAddr(src)); break; + case IADD: masm.addl(asIntReg(dst), (AMD64Address) tasm.asIntAddr(src)); break; + case ISUB: masm.subl(asIntReg(dst), (AMD64Address) tasm.asIntAddr(src)); break; + case IAND: masm.andl(asIntReg(dst), (AMD64Address) tasm.asIntAddr(src)); break; + case IOR: masm.orl(asIntReg(dst), (AMD64Address) tasm.asIntAddr(src)); break; + case IXOR: masm.xorl(asIntReg(dst), (AMD64Address) tasm.asIntAddr(src)); break; - case LADD: masm.addq(asLongReg(dst), tasm.asLongAddr(src)); break; - case LSUB: masm.subq(asLongReg(dst), tasm.asLongAddr(src)); break; - case LAND: masm.andq(asLongReg(dst), tasm.asLongAddr(src)); break; - case LOR: masm.orq(asLongReg(dst), tasm.asLongAddr(src)); break; - case LXOR: masm.xorq(asLongReg(dst), tasm.asLongAddr(src)); break; + case LADD: masm.addq(asLongReg(dst), (AMD64Address) tasm.asLongAddr(src)); break; + case LSUB: masm.subq(asLongReg(dst), (AMD64Address) tasm.asLongAddr(src)); break; + case LAND: masm.andq(asLongReg(dst), (AMD64Address) tasm.asLongAddr(src)); break; + case LOR: masm.orq(asLongReg(dst), (AMD64Address) tasm.asLongAddr(src)); break; + case LXOR: masm.xorq(asLongReg(dst), (AMD64Address) tasm.asLongAddr(src)); break; - case FADD: masm.addss(asFloatReg(dst), tasm.asFloatAddr(src)); break; - case FSUB: masm.subss(asFloatReg(dst), tasm.asFloatAddr(src)); break; - case FMUL: masm.mulss(asFloatReg(dst), tasm.asFloatAddr(src)); break; - case FDIV: masm.divss(asFloatReg(dst), tasm.asFloatAddr(src)); break; + case FADD: masm.addss(asFloatReg(dst), (AMD64Address) tasm.asFloatAddr(src)); break; + case FSUB: masm.subss(asFloatReg(dst), (AMD64Address) tasm.asFloatAddr(src)); break; + case FMUL: masm.mulss(asFloatReg(dst), (AMD64Address) tasm.asFloatAddr(src)); break; + case FDIV: masm.divss(asFloatReg(dst), (AMD64Address) tasm.asFloatAddr(src)); break; - case DADD: masm.addsd(asDoubleReg(dst), tasm.asDoubleAddr(src)); break; - case DSUB: masm.subsd(asDoubleReg(dst), tasm.asDoubleAddr(src)); break; - case DMUL: masm.mulsd(asDoubleReg(dst), tasm.asDoubleAddr(src)); break; - case DDIV: masm.divsd(asDoubleReg(dst), tasm.asDoubleAddr(src)); break; + case DADD: masm.addsd(asDoubleReg(dst), (AMD64Address) tasm.asDoubleAddr(src)); break; + case DSUB: masm.subsd(asDoubleReg(dst), (AMD64Address) tasm.asDoubleAddr(src)); break; + case DMUL: masm.mulsd(asDoubleReg(dst), (AMD64Address) tasm.asDoubleAddr(src)); break; + case DDIV: masm.divsd(asDoubleReg(dst), (AMD64Address) tasm.asDoubleAddr(src)); break; default: throw GraalInternalError.shouldNotReachHere(); } } @@ -455,7 +455,7 @@ tasm.stubs.add(slowPath); switch (result.getKind()) { case Int: masm.cmpl(asIntReg(result), Integer.MIN_VALUE); break; - case Long: masm.cmpq(asLongReg(result), tasm.asLongConstRef(Constant.forLong(java.lang.Long.MIN_VALUE))); break; + case Long: masm.cmpq(asLongReg(result), (AMD64Address) tasm.asLongConstRef(Constant.forLong(java.lang.Long.MIN_VALUE))); break; default: throw GraalInternalError.shouldNotReachHere(); } masm.jcc(ConditionFlag.Equal, slowPath.start); @@ -477,8 +477,8 @@ public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) { masm.bind(start); switch (x.getKind()) { - case Float: masm.ucomiss(asFloatReg(x), tasm.asFloatConstRef(Constant.FLOAT_0)); break; - case Double: masm.ucomisd(asDoubleReg(x), tasm.asDoubleConstRef(Constant.DOUBLE_0)); break; + case Float: masm.ucomiss(asFloatReg(x), (AMD64Address) tasm.asFloatConstRef(Constant.FLOAT_0)); break; + case Double: masm.ucomisd(asDoubleReg(x), (AMD64Address) tasm.asDoubleConstRef(Constant.DOUBLE_0)); break; default: throw GraalInternalError.shouldNotReachHere(); } Label nan = new Label();
--- a/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64BitManipulationOp.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64BitManipulationOp.java Fri Feb 22 12:15:14 2013 +0100 @@ -22,6 +22,7 @@ */ package com.oracle.graal.lir.amd64; +import com.oracle.graal.amd64.*; import com.oracle.graal.api.code.*; import com.oracle.graal.api.meta.*; import com.oracle.graal.asm.amd64.*; @@ -47,7 +48,7 @@ public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) { Register dst = ValueUtil.asIntReg(result); if (ValueUtil.isAddress(input)) { - Address src = ValueUtil.asAddress(input); + AMD64Address src = (AMD64Address) ValueUtil.asAddress(input); switch (opcode) { case IPOPCNT: masm.popcntl(dst, src);
--- a/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Compare.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Compare.java Fri Feb 22 12:15:14 2013 +0100 @@ -25,6 +25,7 @@ import static com.oracle.graal.api.code.ValueUtil.*; import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; +import com.oracle.graal.amd64.*; import com.oracle.graal.api.meta.*; import com.oracle.graal.asm.amd64.*; import com.oracle.graal.graph.*; @@ -81,17 +82,17 @@ } else { throw GraalInternalError.shouldNotReachHere("Only null object constants are allowed in comparisons"); } - case FCMP: masm.ucomiss(asFloatReg(x), tasm.asFloatConstRef(y)); break; - case DCMP: masm.ucomisd(asDoubleReg(x), tasm.asDoubleConstRef(y)); break; + case FCMP: masm.ucomiss(asFloatReg(x), (AMD64Address) tasm.asFloatConstRef(y)); break; + case DCMP: masm.ucomisd(asDoubleReg(x), (AMD64Address) tasm.asDoubleConstRef(y)); break; default: throw GraalInternalError.shouldNotReachHere(); } } else { switch (opcode) { - case ICMP: masm.cmpl(asIntReg(x), tasm.asIntAddr(y)); break; - case LCMP: masm.cmpq(asLongReg(x), tasm.asLongAddr(y)); break; - case ACMP: masm.cmpptr(asObjectReg(x), tasm.asObjectAddr(y)); break; - case FCMP: masm.ucomiss(asFloatReg(x), tasm.asFloatAddr(y)); break; - case DCMP: masm.ucomisd(asDoubleReg(x), tasm.asDoubleAddr(y)); break; + case ICMP: masm.cmpl(asIntReg(x), (AMD64Address) tasm.asIntAddr(y)); break; + case LCMP: masm.cmpq(asLongReg(x), (AMD64Address) tasm.asLongAddr(y)); break; + case ACMP: masm.cmpptr(asObjectReg(x), (AMD64Address) tasm.asObjectAddr(y)); break; + case FCMP: masm.ucomiss(asFloatReg(x), (AMD64Address) tasm.asFloatAddr(y)); break; + case DCMP: masm.ucomisd(asDoubleReg(x), (AMD64Address) tasm.asDoubleAddr(y)); break; default: throw GraalInternalError.shouldNotReachHere(); } }
--- a/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64ControlFlow.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64ControlFlow.java Fri Feb 22 12:15:14 2013 +0100 @@ -26,8 +26,8 @@ import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; import com.oracle.graal.amd64.*; +import com.oracle.graal.amd64.AMD64Address.Scale; import com.oracle.graal.api.code.*; -import com.oracle.graal.api.code.Address.Scale; import com.oracle.graal.api.code.CompilationResult.JumpTable; import com.oracle.graal.api.meta.*; import com.oracle.graal.asm.*; @@ -167,7 +167,7 @@ } else if (key.getKind() == Kind.Long) { Register longKey = asLongReg(key); for (int i = 0; i < keyConstants.length; i++) { - masm.cmpq(longKey, tasm.asLongConstRef(keyConstants[i])); + masm.cmpq(longKey, (AMD64Address) tasm.asLongConstRef(keyConstants[i])); masm.jcc(ConditionFlag.Equal, keyTargets[i].label()); } } else if (key.getKind() == Kind.Object) { @@ -338,11 +338,11 @@ // Set scratch to address of jump table int leaPos = buf.position(); - masm.leaq(scratch, new Address(tasm.target.wordKind, AMD64.rip.asValue(), 0)); + masm.leaq(scratch, new AMD64Address(tasm.target.wordKind, AMD64.rip.asValue(), 0)); int afterLea = buf.position(); // Load jump table entry into scratch and jump to it - masm.movslq(value, new Address(Kind.Int, scratch.asValue(), value.asValue(), Scale.Times4, 0)); + masm.movslq(value, new AMD64Address(Kind.Int, scratch.asValue(), value.asValue(), Scale.Times4, 0)); masm.addq(scratch, value); masm.jmp(scratch); @@ -354,7 +354,7 @@ // Patch LEA instruction above now that we know the position of the jump table int jumpTablePos = buf.position(); buf.setPosition(leaPos); - masm.leaq(scratch, new Address(tasm.target.wordKind, AMD64.rip.asValue(), jumpTablePos - afterLea)); + masm.leaq(scratch, new AMD64Address(tasm.target.wordKind, AMD64.rip.asValue(), jumpTablePos - afterLea)); buf.setPosition(jumpTablePos); // Emit jump table entries @@ -413,9 +413,10 @@ default: throw GraalInternalError.shouldNotReachHere(); } } else { + AMD64Address addr = (AMD64Address) tasm.asAddress(other); switch (other.getKind()) { - case Int: masm.cmovl(cond, asRegister(result), tasm.asAddress(other)); break; - case Long: masm.cmovq(cond, asRegister(result), tasm.asAddress(other)); break; + case Int: masm.cmovl(cond, asRegister(result), addr); break; + case Long: masm.cmovq(cond, asRegister(result), addr); break; default: throw GraalInternalError.shouldNotReachHere(); } }
--- a/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Move.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Move.java Fri Feb 22 12:15:14 2013 +0100 @@ -131,7 +131,7 @@ @Override public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) { - load(tasm, masm, result, (Address) address, state); + load(tasm, masm, result, (AMD64Address) address, state); } } @@ -149,7 +149,7 @@ @Override public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) { - store(tasm, masm, (Address) address, input, state); + store(tasm, masm, (AMD64Address) address, input, state); } } @@ -165,7 +165,7 @@ @Override public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) { - masm.leaq(asLongReg(result), tasm.asAddress(address)); + masm.leaq(asLongReg(result), (AMD64Address) tasm.asAddress(address)); } } @@ -217,7 +217,7 @@ @Override public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) { - compareAndSwap(tasm, masm, result, (Address) address, cmpValue, newValue); + compareAndSwap(tasm, masm, result, (AMD64Address) address, cmpValue, newValue); } } @@ -265,23 +265,25 @@ } private static void reg2stack(TargetMethodAssembler tasm, AMD64MacroAssembler masm, Value result, Value input) { + AMD64Address dest = (AMD64Address) tasm.asAddress(result); switch (input.getKind()) { - case Int: masm.movl(tasm.asAddress(result), asRegister(input)); break; - case Long: masm.movq(tasm.asAddress(result), asRegister(input)); break; - case Float: masm.movflt(tasm.asAddress(result), asFloatReg(input)); break; - case Double: masm.movsd(tasm.asAddress(result), asDoubleReg(input)); break; - case Object: masm.movq(tasm.asAddress(result), asRegister(input)); break; + case Int: masm.movl(dest, asRegister(input)); break; + case Long: masm.movq(dest, asRegister(input)); break; + case Float: masm.movflt(dest, asFloatReg(input)); break; + case Double: masm.movsd(dest, asDoubleReg(input)); break; + case Object: masm.movq(dest, asRegister(input)); break; default: throw GraalInternalError.shouldNotReachHere(); } } private static void stack2reg(TargetMethodAssembler tasm, AMD64MacroAssembler masm, Value result, Value input) { + AMD64Address src = (AMD64Address) tasm.asAddress(input); switch (input.getKind()) { - case Int: masm.movl(asRegister(result), tasm.asAddress(input)); break; - case Long: masm.movq(asRegister(result), tasm.asAddress(input)); break; - case Float: masm.movflt(asFloatReg(result), tasm.asAddress(input)); break; - case Double: masm.movdbl(asDoubleReg(result), tasm.asAddress(input)); break; - case Object: masm.movq(asRegister(result), tasm.asAddress(input)); break; + case Int: masm.movl(asRegister(result), src); break; + case Long: masm.movq(asRegister(result), src); break; + case Float: masm.movflt(asFloatReg(result), src); break; + case Double: masm.movdbl(asDoubleReg(result), src); break; + case Object: masm.movq(asRegister(result), src); break; default: throw GraalInternalError.shouldNotReachHere(); } } @@ -316,7 +318,7 @@ assert !tasm.runtime.needsDataPatch(input); masm.xorps(asFloatReg(result), asFloatReg(result)); } else { - masm.movflt(asFloatReg(result), tasm.asFloatConstRef(input)); + masm.movflt(asFloatReg(result), (AMD64Address) tasm.asFloatConstRef(input)); } break; case Double: @@ -325,7 +327,7 @@ assert !tasm.runtime.needsDataPatch(input); masm.xorpd(asDoubleReg(result), asDoubleReg(result)); } else { - masm.movdbl(asDoubleReg(result), tasm.asDoubleConstRef(input)); + masm.movdbl(asDoubleReg(result), (AMD64Address) tasm.asDoubleConstRef(input)); } break; case Object: @@ -338,7 +340,7 @@ tasm.recordDataReferenceInCode(input, 0, true); masm.movq(asRegister(result), 0xDEADDEADDEADDEADL); } else { - masm.movq(asRegister(result), tasm.recordDataReferenceInCode(input, 0, false)); + masm.movq(asRegister(result), (AMD64Address) tasm.recordDataReferenceInCode(input, 0, false)); } break; default: @@ -348,14 +350,15 @@ private static void const2stack(TargetMethodAssembler tasm, AMD64MacroAssembler masm, Value result, Constant input) { assert !tasm.runtime.needsDataPatch(input); + AMD64Address dest = (AMD64Address) tasm.asAddress(result); switch (input.getKind().getStackKind()) { - case Int: masm.movl(tasm.asAddress(result), input.asInt()); break; - case Long: masm.movlong(tasm.asAddress(result), input.asLong()); break; - case Float: masm.movl(tasm.asAddress(result), floatToRawIntBits(input.asFloat())); break; - case Double: masm.movlong(tasm.asAddress(result), doubleToRawLongBits(input.asDouble())); break; + case Int: masm.movl(dest, input.asInt()); break; + case Long: masm.movlong(dest, input.asLong()); break; + case Float: masm.movl(dest, floatToRawIntBits(input.asFloat())); break; + case Double: masm.movlong(dest, doubleToRawLongBits(input.asDouble())); break; case Object: if (input.isNull()) { - masm.movlong(tasm.asAddress(result), 0L); + masm.movlong(dest, 0L); } else { throw GraalInternalError.shouldNotReachHere("Non-null object constants must be in register"); } @@ -366,7 +369,7 @@ } - public static void load(TargetMethodAssembler tasm, AMD64MacroAssembler masm, Value result, Address loadAddr, LIRFrameState info) { + public static void load(TargetMethodAssembler tasm, AMD64MacroAssembler masm, Value result, AMD64Address loadAddr, LIRFrameState info) { if (info != null) { tasm.recordImplicitException(masm.codeBuffer.position(), info); } @@ -384,7 +387,7 @@ } } - public static void store(TargetMethodAssembler tasm, AMD64MacroAssembler masm, Address storeAddr, Value input, LIRFrameState info) { + public static void store(TargetMethodAssembler tasm, AMD64MacroAssembler masm, AMD64Address storeAddr, Value input, LIRFrameState info) { if (info != null) { tasm.recordImplicitException(masm.codeBuffer.position(), info); } @@ -435,7 +438,7 @@ } } - protected static void compareAndSwap(TargetMethodAssembler tasm, AMD64MacroAssembler masm, Value result, Address address, Value cmpValue, Value newValue) { + protected static void compareAndSwap(TargetMethodAssembler tasm, AMD64MacroAssembler masm, Value result, AMD64Address address, Value cmpValue, Value newValue) { assert asRegister(cmpValue) == AMD64.rax && asRegister(result) == AMD64.rax; if (tasm.target.isMP) {
--- a/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64TestOp.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64TestOp.java Fri Feb 22 12:15:14 2013 +0100 @@ -25,6 +25,7 @@ import static com.oracle.graal.api.code.ValueUtil.*; import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; +import com.oracle.graal.amd64.*; import com.oracle.graal.api.meta.*; import com.oracle.graal.asm.amd64.*; import com.oracle.graal.graph.*; @@ -77,10 +78,10 @@ } else { switch (x.getKind()) { case Int: - masm.testl(asIntReg(x), tasm.asIntAddr(y)); + masm.testl(asIntReg(x), (AMD64Address) tasm.asIntAddr(y)); break; case Long: - masm.testq(asLongReg(x), tasm.asLongAddr(y)); + masm.testq(asLongReg(x), (AMD64Address) tasm.asLongAddr(y)); break; default: throw GraalInternalError.shouldNotReachHere();
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java Fri Feb 22 12:15:14 2013 +0100 @@ -33,6 +33,7 @@ import com.oracle.graal.lir.LIRInstruction.Opcode; import com.oracle.graal.lir.StandardOp.MoveOp; import com.oracle.graal.lir.asm.*; +import com.oracle.graal.ptx.*; public class PTXMove { @@ -131,7 +132,7 @@ @Override public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { - load(tasm, masm, result, (Address) address, state); + load(tasm, masm, result, (PTXAddress) address, state); } } @@ -149,7 +150,7 @@ @Override public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { - store(tasm, masm, (Address) address, input, state); + store(tasm, masm, (PTXAddress) address, input, state); } } @@ -238,9 +239,9 @@ } @SuppressWarnings("unused") - public static void load(TargetMethodAssembler tasm, PTXAssembler masm, Value result, Address loadAddr, LIRFrameState info) { + public static void load(TargetMethodAssembler tasm, PTXAssembler masm, Value result, PTXAddress loadAddr, LIRFrameState info) { Register a = asRegister(loadAddr.getBase()); - int immOff = loadAddr.getDisplacement(); + long immOff = loadAddr.getDisplacement(); switch (loadAddr.getKind()) { case Int: masm.ld_global_s32(asRegister(result), a, immOff); @@ -254,9 +255,9 @@ } @SuppressWarnings("unused") - public static void store(TargetMethodAssembler tasm, PTXAssembler masm, Address storeAddr, Value input, LIRFrameState info) { + public static void store(TargetMethodAssembler tasm, PTXAssembler masm, PTXAddress storeAddr, Value input, LIRFrameState info) { Register a = asRegister(storeAddr.getBase()); - int immOff = storeAddr.getDisplacement(); + long immOff = storeAddr.getDisplacement(); if (isRegister(input)) { switch (storeAddr.getKind()) { case Int:
--- a/graal/com.oracle.graal.lir/src/com/oracle/graal/lir/LIRInstructionClass.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.lir/src/com/oracle/graal/lir/LIRInstructionClass.java Fri Feb 22 12:15:14 2013 +0100 @@ -361,8 +361,10 @@ private static void doAddress(Address address, OperandMode mode, EnumSet<OperandFlag> flags, ValueProcedure proc) { assert flags.contains(OperandFlag.ADDR); - address.setBase(proc.doValue(address.getBase(), mode, LIRInstruction.ADDRESS_FLAGS)); - address.setIndex(proc.doValue(address.getIndex(), mode, LIRInstruction.ADDRESS_FLAGS)); + Value[] components = address.components(); + for (int i = 0; i < components.length; i++) { + components[i] = proc.doValue(components[i], mode, LIRInstruction.ADDRESS_FLAGS); + } } public final Value forEachRegisterHint(LIRInstruction obj, OperandMode mode, ValueProcedure proc) {
--- a/graal/com.oracle.graal.lir/src/com/oracle/graal/lir/asm/TargetMethodAssembler.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.lir/src/com/oracle/graal/lir/asm/TargetMethodAssembler.java Fri Feb 22 12:15:14 2013 +0100 @@ -155,7 +155,7 @@ int pos = asm.codeBuffer.position(); Debug.log("Data reference in code: pos = %d, data = %s", pos, data.toString()); compilationResult.recordDataReference(pos, data, alignment, inlined); - return Address.Placeholder; + return asm.getPlaceholder(); } /** @@ -233,7 +233,7 @@ public Address asAddress(Value value) { if (isStackSlot(value)) { StackSlot slot = (StackSlot) value; - return new Address(slot.getKind(), frameMap.registerConfig.getFrameRegister().asValue(), frameMap.offsetForStackSlot(slot)); + return asm.makeAddress(slot.getKind(), frameMap.registerConfig.getFrameRegister().asValue(), frameMap.offsetForStackSlot(slot)); } return (Address) value; }
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/spi/LIRGeneratorTool.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/spi/LIRGeneratorTool.java Fri Feb 22 12:15:14 2013 +0100 @@ -65,6 +65,16 @@ public abstract Address makeAddress(LocationNode location, ValueNode object); + public abstract Address makeAddress(Kind kind, Value base, int displacement); + + public Address makeAddress(Kind kind, Value base) { + return makeAddress(kind, base, 0); + } + + public Address makeAddress(Kind kind, int address) { + return makeAddress(kind, Value.ILLEGAL, address); + } + public abstract Value emitMove(Value input); public abstract void emitMove(Value src, Value dst);
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.ptx/src/com/oracle/graal/ptx/PTXAddress.java Fri Feb 22 12:15:14 2013 +0100 @@ -0,0 +1,123 @@ +/* + * 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.ptx; + +import static com.oracle.graal.api.code.ValueUtil.*; + +import com.oracle.graal.api.code.*; +import com.oracle.graal.api.meta.*; + +/** + * Represents an address in target machine memory, specified via some combination of a base register + * and a displacement. + */ +public final class PTXAddress extends Address { + + private static final long serialVersionUID = 8343625682010474837L; + + private final Value[] base; + private final long displacement; + + /** + * Creates an {@link PTXAddress} with given base register and no displacement. + * + * @param kind the kind of the value being addressed + * @param base the base register + */ + public PTXAddress(Kind kind, Value base) { + this(kind, base, 0); + } + + /** + * Creates an {@link PTXAddress} with given base register and a displacement. This is the most + * general constructor. + * + * @param kind the kind of the value being addressed + * @param base the base register + * @param displacement the displacement + */ + public PTXAddress(Kind kind, Value base, long displacement) { + super(kind); + this.base = new Value[1]; + this.setBase(base); + this.displacement = displacement; + + assert !isConstant(base) && !isStackSlot(base); + } + + @Override + public Value[] components() { + return base; + } + + @Override + public String toString() { + StringBuilder s = new StringBuilder(); + s.append(getKind().getJavaName()).append("["); + String sep = ""; + if (isLegal(getBase())) { + s.append(getBase()); + sep = " + "; + } + if (getDisplacement() < 0) { + s.append(" - ").append(-getDisplacement()); + } else if (getDisplacement() > 0) { + s.append(sep).append(getDisplacement()); + } + s.append("]"); + return s.toString(); + } + + @Override + public boolean equals(Object obj) { + if (obj instanceof PTXAddress) { + PTXAddress addr = (PTXAddress) obj; + return getKind() == addr.getKind() && getDisplacement() == addr.getDisplacement() && getBase().equals(addr.getBase()); + } + return false; + } + + @Override + public int hashCode() { + return getBase().hashCode() ^ ((int) getDisplacement() << 4) ^ (getKind().ordinal() << 12); + } + + /** + * @return Base register that defines the start of the address computation. If not present, is + * denoted by {@link Value#ILLEGAL}. + */ + public Value getBase() { + return base[0]; + } + + public void setBase(Value base) { + this.base[0] = base; + } + + /** + * @return Optional additive displacement. + */ + public long getDisplacement() { + return displacement; + } +}
--- a/graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/nodes/DirectReadNode.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/nodes/DirectReadNode.java Fri Feb 22 12:15:14 2013 +0100 @@ -22,7 +22,6 @@ */ package com.oracle.graal.snippets.nodes; -import com.oracle.graal.api.code.*; import com.oracle.graal.api.meta.*; import com.oracle.graal.nodes.*; import com.oracle.graal.nodes.extended.*; @@ -46,7 +45,7 @@ @Override public void generate(LIRGeneratorTool gen) { - gen.setResult(this, gen.emitLoad(new Address(readKind, gen.operand(address)), false)); + gen.setResult(this, gen.emitLoad(gen.makeAddress(readKind, gen.operand(address)), false)); } @NodeIntrinsic
--- a/graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/nodes/DirectStoreNode.java Thu Feb 21 14:24:47 2013 -0800 +++ b/graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/nodes/DirectStoreNode.java Fri Feb 22 12:15:14 2013 +0100 @@ -22,7 +22,6 @@ */ package com.oracle.graal.snippets.nodes; -import com.oracle.graal.api.code.*; import com.oracle.graal.api.meta.*; import com.oracle.graal.nodes.*; import com.oracle.graal.nodes.extended.*; @@ -49,7 +48,7 @@ @Override public void generate(LIRGeneratorTool gen) { Value v = gen.operand(value); - gen.emitStore(new Address(kind, gen.operand(address)), v, false); + gen.emitStore(gen.makeAddress(kind, gen.operand(address)), v, false); } /*
--- a/mx/projects Thu Feb 21 14:24:47 2013 -0800 +++ b/mx/projects Fri Feb 22 12:15:14 2013 +0100 @@ -147,7 +147,7 @@ # graal.lir.ptx project@com.oracle.graal.lir.ptx@subDir=graal project@com.oracle.graal.lir.ptx@sourceDirs=src -project@com.oracle.graal.lir.ptx@dependencies=com.oracle.graal.lir,com.oracle.graal.asm.ptx +project@com.oracle.graal.lir.ptx@dependencies=com.oracle.graal.asm.ptx,com.oracle.graal.lir project@com.oracle.graal.lir.ptx@checkstyle=com.oracle.graal.graph project@com.oracle.graal.lir.ptx@javaCompliance=1.7 @@ -245,7 +245,7 @@ # graal.compiler.ptx.test project@com.oracle.graal.compiler.ptx.test@subDir=graal project@com.oracle.graal.compiler.ptx.test@sourceDirs=src -project@com.oracle.graal.compiler.ptx.test@dependencies=com.oracle.graal.compiler.ptx,com.oracle.graal.compiler.test,com.oracle.graal.ptx +project@com.oracle.graal.compiler.ptx.test@dependencies=com.oracle.graal.compiler.ptx,com.oracle.graal.compiler.test project@com.oracle.graal.compiler.ptx.test@checkstyle=com.oracle.graal.graph project@com.oracle.graal.compiler.ptx.test@javaCompliance=1.7 @@ -328,7 +328,7 @@ # graal.asm.ptx project@com.oracle.graal.asm.ptx@subDir=graal project@com.oracle.graal.asm.ptx@sourceDirs=src -project@com.oracle.graal.asm.ptx@dependencies=com.oracle.graal.asm +project@com.oracle.graal.asm.ptx@dependencies=com.oracle.graal.asm,com.oracle.graal.ptx project@com.oracle.graal.asm.ptx@checkstyle=com.oracle.graal.graph project@com.oracle.graal.asm.ptx@javaCompliance=1.7