Mercurial > hg > truffle

changeset 8168:58b1820ff19a

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
Convert PTXAddress to CompositeValue.
author Roland Schatz <roland.schatz@oracle.com>
date Thu, 07 Mar 2013 16:33:27 +0100
parents 7f57c30575c8
children 0fdbde04cca3
files graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAddress.java graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXAddressValue.java graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXBitManipulationOp.java graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java graal/com.oracle.graal.ptx/src/com/oracle/graal/ptx/PTXAddress.java
diffstat 7 files changed, 218 insertions(+), 157 deletions(-) [+]
line wrap: on
line diff
--- /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;
+    }
+}
--- 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
--- 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));
     }
--- /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);
+    }
+}
--- 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;
--- 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");
     }
 }
--- 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;
-    }
-}