# HG changeset patch
# User Roland Schatz <roland.schatz@oracle.com>
# Date 1361531714 -3600
# Node ID 838293a77af75ae859b3449efed9d2442bc5f266
# Parent  a063308816d9deb88e3c6d98d3c0943a690db163
Make Address class platform specific.

diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.amd64/src/com/oracle/graal/amd64/AMD64Address.java
--- /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;
+    }
+}
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/Address.java
--- 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();
 }
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.asm.amd64.test/src/com/oracle/graal/asm/amd64/test/SimpleAssemblerTest.java
--- 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;
             }
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.asm.amd64/src/com/oracle/graal/asm/amd64/AMD64Assembler.java
--- 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;
+    }
 }
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.asm.amd64/src/com/oracle/graal/asm/amd64/AMD64MacroAssembler.java
--- 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));
         }
     }
 }
diff -r a063308816d9 -r 838293a77af7 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 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;
+    }
 }
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.asm.sparc/src/com/oracle/graal/asm/sparc/SPARCAssembler.java
--- 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;
+    }
 }
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.asm/src/com/oracle/graal/asm/AbstractAssembler.java
--- 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();
 }
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.compiler.amd64/src/com/oracle/graal/compiler/amd64/AMD64LIRGenerator.java
--- 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);
diff -r a063308816d9 -r 838293a77af7 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 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
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.compiler.sparc/src/com/oracle/graal/compiler/sparc/SPARCLIRGenerator.java
--- 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;
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotBackend.java
--- 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);
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64SafepointOp.java
--- 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));
         }
     }
 }
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/CurrentThread.java
--- 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
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/TailcallNode.java
--- 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);
     }
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/WriteBarrier.java
--- 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);
     }
 }
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Arithmetic.java
--- 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();
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64BitManipulationOp.java
--- 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);
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Compare.java
--- 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();
             }
         }
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64ControlFlow.java
--- 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();
             }
         }
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Move.java
--- 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) {
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64TestOp.java
--- 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();
diff -r a063308816d9 -r 838293a77af7 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 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:
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.lir/src/com/oracle/graal/lir/LIRInstructionClass.java
--- 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) {
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.lir/src/com/oracle/graal/lir/asm/TargetMethodAssembler.java
--- 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;
     }
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/spi/LIRGeneratorTool.java
--- 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);
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.ptx/src/com/oracle/graal/ptx/PTXAddress.java
--- /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;
+    }
+}
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/nodes/DirectReadNode.java
--- 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
diff -r a063308816d9 -r 838293a77af7 graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/nodes/DirectStoreNode.java
--- 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);
     }
 
     /*
diff -r a063308816d9 -r 838293a77af7 mx/projects
--- 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