# HG changeset patch # User Roland Schatz # Date 1362670407 -3600 # Node ID 58b1820ff19aac77d5c3cd988c69d849cad9a69c # Parent 7f57c30575c897deeb066d8878e45e4ee087ce82 Convert PTXAddress to CompositeValue. diff -r 7f57c30575c8 -r 58b1820ff19a graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAddress.java --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAddress.java Thu Mar 07 16:33:27 2013 +0100 @@ -0,0 +1,72 @@ +/* + * 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.asm.ptx; + +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 implements AbstractAddress { + + private final Register base; + private final long displacement; + + /** + * Creates an {@link PTXAddress} with given base register and no displacement. + * + * @param base the base register + */ + public PTXAddress(Register base) { + this(base, 0); + } + + /** + * Creates an {@link PTXAddress} with given base register and a displacement. This is the most + * general constructor. + * + * @param base the base register + * @param displacement the displacement + */ + public PTXAddress(Register base, long displacement) { + this.base = base; + this.displacement = displacement; + } + + /** + * @return Base register that defines the start of the address computation. If not present, is + * denoted by {@link Value#ILLEGAL}. + */ + public Register getBase() { + return base; + } + + /** + * @return Optional additive displacement. + */ + public long getDisplacement() { + return displacement; + } +} diff -r 7f57c30575c8 -r 58b1820ff19a graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java --- a/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Thu Mar 07 16:20:18 2013 +0100 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Thu Mar 07 16:33:27 2013 +0100 @@ -22,9 +22,10 @@ */ package com.oracle.graal.asm.ptx; +import static com.oracle.graal.api.code.ValueUtil.*; + import com.oracle.graal.api.code.*; import com.oracle.graal.api.meta.*; -import com.oracle.graal.ptx.*; public class PTXAssembler extends AbstractPTXAssembler { @@ -177,7 +178,7 @@ emitString("exit;" + " " + ""); } - public final void ld_global_b8(Register d, Register a, int immOff) { + public final void ld_global_b8(Register d, Register a, long immOff) { emitString("ld.global.b8" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); } @@ -743,7 +744,8 @@ @Override public PTXAddress makeAddress(Kind kind, Value base, int displacement) { - return new PTXAddress(kind, base, displacement); + assert isRegister(base); + return new PTXAddress(asRegister(base), displacement); } @Override diff -r 7f57c30575c8 -r 58b1820ff19a graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java --- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java Thu Mar 07 16:20:18 2013 +0100 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java Thu Mar 07 16:33:27 2013 +0100 @@ -36,11 +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.*; 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; @@ -51,7 +51,6 @@ import com.oracle.graal.nodes.*; import com.oracle.graal.nodes.calc.*; import com.oracle.graal.nodes.java.*; -import com.oracle.graal.ptx.*; /** * This class implements the PTX specific portion of the LIR generator. @@ -114,40 +113,49 @@ } } - private PTXAddress prepareAddress(Kind kind, Value base, int displacement, Value index, int scale) { - Value baseRegister = base; + private PTXAddressValue prepareAddress(Kind kind, Value base, int displacement, Value index, int scale) { + AllocatableValue baseRegister; long finalDisp = displacement; if (isConstant(base)) { if (asConstant(base).isNull()) { - baseRegister = Value.ILLEGAL; + baseRegister = AllocatableValue.UNUSED; } else if (asConstant(base).getKind() != Kind.Object) { finalDisp += asConstant(base).asLong(); - baseRegister = Value.ILLEGAL; + baseRegister = AllocatableValue.UNUSED; + } else { + baseRegister = load(base); } + } else if (base == Value.ILLEGAL) { + baseRegister = AllocatableValue.UNUSED; + } else { + baseRegister = asAllocatable(base); } if (index != Value.ILLEGAL) { if (isConstant(index)) { finalDisp += asConstant(index).asLong() * scale; } else { - Value indexRegister = index; + Value indexRegister; if (scale != 1) { indexRegister = emitMul(index, Constant.forInt(scale)); + } else { + indexRegister = index; } - if (baseRegister == Value.ILLEGAL) { - baseRegister = indexRegister; + + if (baseRegister == AllocatableValue.UNUSED) { + baseRegister = asAllocatable(indexRegister); } else { baseRegister = emitAdd(baseRegister, indexRegister); } } } - return new PTXAddress(kind, baseRegister, finalDisp); + return new PTXAddressValue(kind, baseRegister, finalDisp); } @Override public Variable emitLoad(Kind kind, Value base, int displacement, Value index, int scale, boolean canTrap) { - PTXAddress loadAddress = prepareAddress(kind, base, displacement, index, scale); + PTXAddressValue loadAddress = prepareAddress(kind, base, displacement, index, scale); Variable result = newVariable(loadAddress.getKind()); append(new LoadOp(result, loadAddress, canTrap ? state() : null)); return result; @@ -155,7 +163,7 @@ @Override public void emitStore(Kind kind, Value base, int displacement, Value index, int scale, Value inputVal, boolean canTrap) { - PTXAddress storeAddress = prepareAddress(kind, base, displacement, index, scale); + PTXAddressValue storeAddress = prepareAddress(kind, base, displacement, index, scale); Variable input = load(inputVal); append(new StoreOp(storeAddress, input, canTrap ? state() : null)); } diff -r 7f57c30575c8 -r 58b1820ff19a graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXAddressValue.java --- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXAddressValue.java Thu Mar 07 16:33:27 2013 +0100 @@ -0,0 +1,106 @@ +/* + * 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.lir.ptx; + +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.lir.*; + +/** + * Represents an address in target machine memory, specified via some combination of a base register + * and a displacement. + */ +public final class PTXAddressValue extends CompositeValue { + + private static final long serialVersionUID = 1802222435353022623L; + + @Component({REG, UNUSED}) private AllocatableValue base; + private final long displacement; + + /** + * Creates an {@link PTXAddressValue} with given base register and no displacement. + * + * @param kind the kind of the value being addressed + * @param base the base register + */ + public PTXAddressValue(Kind kind, AllocatableValue base) { + this(kind, base, 0); + } + + /** + * Creates an {@link PTXAddressValue} 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 PTXAddressValue(Kind kind, AllocatableValue base, long displacement) { + super(kind); + this.base = base; + this.displacement = displacement; + + assert !isStackSlot(base); + } + + public PTXAddress toAddress() { + Register baseReg = base == AllocatableValue.UNUSED ? Register.None : asRegister(base); + return new PTXAddress(baseReg, displacement); + } + + @Override + public String toString() { + StringBuilder s = new StringBuilder(); + s.append(getKind().getJavaName()).append("["); + String sep = ""; + if (isLegal(base)) { + s.append(base); + sep = " + "; + } + if (displacement < 0) { + s.append(" - ").append(-displacement); + } else if (displacement > 0) { + s.append(sep).append(displacement); + } + s.append("]"); + return s.toString(); + } + + @Override + public boolean equals(Object obj) { + if (obj instanceof PTXAddressValue) { + PTXAddressValue addr = (PTXAddressValue) obj; + return getKind() == addr.getKind() && displacement == addr.displacement && base.equals(addr.base); + } + return false; + } + + @Override + public int hashCode() { + return base.hashCode() ^ ((int) displacement << 4) ^ (getKind().ordinal() << 12); + } +} diff -r 7f57c30575c8 -r 58b1820ff19a graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXBitManipulationOp.java --- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXBitManipulationOp.java Thu Mar 07 16:20:18 2013 +0100 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXBitManipulationOp.java Thu Mar 07 16:33:27 2013 +0100 @@ -36,7 +36,7 @@ @Opcode private final IntrinsicOpcode opcode; @Def protected Value result; - @Use({OperandFlag.REG, OperandFlag.ADDR}) protected Value input; + @Use({OperandFlag.REG}) protected Value input; public PTXBitManipulationOp(IntrinsicOpcode opcode, Value result, Value input) { this.opcode = opcode; diff -r 7f57c30575c8 -r 58b1820ff19a graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java --- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java Thu Mar 07 16:20:18 2013 +0100 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java Thu Mar 07 16:33:27 2013 +0100 @@ -33,7 +33,6 @@ 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 { @@ -121,10 +120,10 @@ public static class LoadOp extends PTXLIRInstruction { @Def({REG}) protected AllocatableValue result; - @Use({ADDR}) protected PTXAddress address; + @Use({COMPOSITE}) protected PTXAddressValue address; @State protected LIRFrameState state; - public LoadOp(AllocatableValue result, PTXAddress address, LIRFrameState state) { + public LoadOp(AllocatableValue result, PTXAddressValue address, LIRFrameState state) { this.result = result; this.address = address; this.state = state; @@ -132,14 +131,13 @@ @Override public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { - Register a = asRegister(address.getBase()); - long immOff = address.getDisplacement(); + PTXAddress addr = address.toAddress(); switch (address.getKind()) { case Int: - masm.ld_global_s32(asRegister(result), a, immOff); + masm.ld_global_s32(asRegister(result), addr.getBase(), addr.getDisplacement()); break; case Object: - masm.ld_global_u32(asRegister(result), a, immOff); + masm.ld_global_u32(asRegister(result), addr.getBase(), addr.getDisplacement()); break; default: throw GraalInternalError.shouldNotReachHere(); @@ -149,11 +147,11 @@ public static class StoreOp extends PTXLIRInstruction { - @Use({ADDR}) protected PTXAddress address; + @Use({COMPOSITE}) protected PTXAddressValue address; @Use({REG}) protected AllocatableValue input; @State protected LIRFrameState state; - public StoreOp(PTXAddress address, AllocatableValue input, LIRFrameState state) { + public StoreOp(PTXAddressValue address, AllocatableValue input, LIRFrameState state) { this.address = address; this.input = input; this.state = state; @@ -161,13 +159,11 @@ @Override public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { - Register a = asRegister(address.getBase()); - long immOff = address.getDisplacement(); - assert isRegister(input); + PTXAddress addr = address.toAddress(); switch (address.getKind()) { case Int: - masm.st_global_s32(a, immOff, asRegister(input)); + masm.st_global_s32(addr.getBase(), addr.getDisplacement(), asRegister(input)); break; default: throw GraalInternalError.shouldNotReachHere(); @@ -178,9 +174,9 @@ public static class LeaOp extends PTXLIRInstruction { @Def({REG}) protected AllocatableValue result; - @Use({ADDR, UNINITIALIZED}) protected PTXAddress address; + @Use({COMPOSITE, UNINITIALIZED}) protected PTXAddressValue address; - public LeaOp(AllocatableValue result, PTXAddress address) { + public LeaOp(AllocatableValue result, PTXAddressValue address) { this.result = result; this.address = address; } @@ -211,11 +207,11 @@ public static class CompareAndSwapOp extends PTXLIRInstruction { @Def protected AllocatableValue result; - @Use({ADDR}) protected PTXAddress address; + @Use({COMPOSITE}) protected PTXAddressValue address; @Use protected AllocatableValue cmpValue; @Use protected AllocatableValue newValue; - public CompareAndSwapOp(AllocatableValue result, PTXAddress address, AllocatableValue cmpValue, AllocatableValue newValue) { + public CompareAndSwapOp(AllocatableValue result, PTXAddressValue address, AllocatableValue cmpValue, AllocatableValue newValue) { this.result = result; this.address = address; this.cmpValue = cmpValue; @@ -276,7 +272,7 @@ } @SuppressWarnings("unused") - protected static void compareAndSwap(TargetMethodAssembler tasm, PTXAssembler masm, AllocatableValue result, PTXAddress address, AllocatableValue cmpValue, AllocatableValue newValue) { + protected static void compareAndSwap(TargetMethodAssembler tasm, PTXAssembler masm, AllocatableValue result, PTXAddressValue address, AllocatableValue cmpValue, AllocatableValue newValue) { throw new InternalError("NYI"); } } diff -r 7f57c30575c8 -r 58b1820ff19a graal/com.oracle.graal.ptx/src/com/oracle/graal/ptx/PTXAddress.java --- a/graal/com.oracle.graal.ptx/src/com/oracle/graal/ptx/PTXAddress.java Thu Mar 07 16:20:18 2013 +0100 +++ /dev/null Thu Jan 01 00:00:00 1970 +0000 @@ -1,123 +0,0 @@ -/* - * 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; - } -}