# HG changeset patch
# User Christian Humer <christian.humer@gmail.com>
# Date 1361794479 -3600
# Node ID a4a8909a6096e8e21d781f75653cfbd962d437db
# Parent  6e4fb0ccebb1230587e60fc7c6d1a18e5356d597# Parent  c052cfe3cae3f10698b0b79c11a4cfa57839aea7
Merge.

diff -r 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/Address.java	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CodeCacheProvider.java
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CodeCacheProvider.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CodeCacheProvider.java	Mon Feb 25 13:14:39 2013 +0100
@@ -60,10 +60,8 @@
 
     /**
      * Gets the register configuration to use when compiling a given method.
-     * 
-     * @param method the top level method of a compilation
      */
-    RegisterConfig lookupRegisterConfig(ResolvedJavaMethod method);
+    RegisterConfig lookupRegisterConfig();
 
     /**
      * Custom area on the stack of each compiled method that the VM can use for its own purposes.
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CodeUtil.java
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CodeUtil.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CodeUtil.java	Mon Feb 25 13:14:39 2013 +0100
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2010, 2011, 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,6 +22,7 @@
  */
 package com.oracle.graal.api.code;
 
+import java.lang.reflect.*;
 import java.util.*;
 
 import com.oracle.graal.api.meta.*;
@@ -317,12 +318,12 @@
     public static CallingConvention getCallingConvention(CodeCacheProvider codeCache, CallingConvention.Type type, ResolvedJavaMethod method, boolean stackOnly) {
         Signature sig = method.getSignature();
         JavaType retType = sig.getReturnType(null);
-        JavaType[] argTypes = new JavaType[sig.getParameterCount(false)];
+        JavaType[] argTypes = new JavaType[sig.getParameterCount(!Modifier.isStatic(method.getModifiers()))];
         for (int i = 0; i < argTypes.length; i++) {
             argTypes[i] = sig.getParameterType(i, null);
         }
 
-        RegisterConfig registerConfig = codeCache.lookupRegisterConfig(method);
+        RegisterConfig registerConfig = codeCache.lookupRegisterConfig();
         return registerConfig.getCallingConvention(type, retType, argTypes, codeCache.getTarget(), stackOnly);
     }
 }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CompilationResult.java
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CompilationResult.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CompilationResult.java	Mon Feb 25 13:14:39 2013 +0100
@@ -196,25 +196,6 @@
     }
 
     /**
-     * Labels some inline data in the code.
-     */
-    public static final class InlineData extends CodeAnnotation {
-
-        private static final long serialVersionUID = 305997507263827108L;
-        public final int size;
-
-        public InlineData(int position, int size) {
-            super(position);
-            this.size = size;
-        }
-
-        @Override
-        public String toString() {
-            return getClass().getSimpleName() + "@" + position + ": size=" + size;
-        }
-    }
-
-    /**
      * Describes a table of signed offsets embedded in the code. The offsets are relative to the
      * starting address of the table. This type of table maybe generated when translating a
      * multi-way branch based on a key value from a dense value set (e.g. the {@code tableswitch}
@@ -256,43 +237,6 @@
     }
 
     /**
-     * Describes a table of key and offset pairs. The offset in each table entry is relative to the
-     * address of the table. This type of table maybe generated when translating a multi-way branch
-     * based on a key value from a sparse value set (e.g. the {@code lookupswitch} JVM instruction).
-     */
-    public static final class LookupTable extends CodeAnnotation {
-
-        private static final long serialVersionUID = 8367952567559116160L;
-
-        /**
-         * The number of entries in the table.
-         */
-        public final int npairs;
-
-        /**
-         * The size (in bytes) of entry's key.
-         */
-        public final int keySize;
-
-        /**
-         * The size (in bytes) of entry's offset value.
-         */
-        public final int offsetSize;
-
-        public LookupTable(int position, int npairs, int keySize, int offsetSize) {
-            super(position);
-            this.npairs = npairs;
-            this.keySize = keySize;
-            this.offsetSize = offsetSize;
-        }
-
-        @Override
-        public String toString() {
-            return getClass().getSimpleName() + "@" + position + ": [npairs=" + npairs + ", keySize=" + keySize + ", offsetSize=" + offsetSize + "]";
-        }
-    }
-
-    /**
      * Represents exception handler information for a specific code position. It includes the catch
      * code position as well as the caught exception type.
      */
@@ -349,8 +293,6 @@
     private int customStackAreaOffset = -1;
     private int registerRestoreEpilogueOffset = -1;
 
-    private CalleeSaveLayout calleeSaveLayout;
-
     /**
      * The buffer containing the emitted machine code.
      */
@@ -410,15 +352,6 @@
     }
 
     /**
-     * Sets the info on callee-saved registers used by this method.
-     * 
-     * @param csl the register-saving info.
-     */
-    public void setCalleeSaveLayout(CalleeSaveLayout csl) {
-        calleeSaveLayout = csl;
-    }
-
-    /**
      * Records a reference to the data section in the code section (e.g. to load an integer or
      * floating point constant).
      * 
@@ -537,13 +470,6 @@
     }
 
     /**
-     * @return the layout information for callee-saved registers used by this method.
-     */
-    public CalleeSaveLayout getCalleeSaveLayout() {
-        return calleeSaveLayout;
-    }
-
-    /**
      * @return the machine code generated for this method
      */
     public byte[] getTargetCode() {
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/TargetDescription.java
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/TargetDescription.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/TargetDescription.java	Mon Feb 25 13:14:39 2013 +0100
@@ -87,18 +87,7 @@
      */
     public final int implicitNullCheckLimit;
 
-    /**
-     * Specifies how {@code long} and {@code double} constants are to be stored in
-     * {@linkplain BytecodeFrame frames}. This is useful for VMs such as HotSpot where convention
-     * the interpreter uses is that the second local holds the first raw word of the native long or
-     * double representation. This is actually reasonable, since locals and stack arrays grow
-     * downwards in all implementations. If, on some machine, the interpreter's Java locals or stack
-     * were to grow upwards, the embedded doubles would be word-swapped.)
-     */
-    public final boolean debugInfoDoubleWordsInSecondSlot;
-
-    public TargetDescription(Architecture arch, boolean isMP, int stackAlignment, int stackBias, int implicitNullCheckLimit, int pageSize, int cacheAlignment, boolean inlineObjects,
-                    boolean debugInfoDoubleWordsInSecondSlot) {
+    public TargetDescription(Architecture arch, boolean isMP, int stackAlignment, int stackBias, int implicitNullCheckLimit, int pageSize, int cacheAlignment, boolean inlineObjects) {
         this.arch = arch;
         this.pageSize = pageSize;
         this.isMP = isMP;
@@ -109,7 +98,6 @@
         this.implicitNullCheckLimit = implicitNullCheckLimit;
         this.cacheAlignment = cacheAlignment;
         this.inlineObjects = inlineObjects;
-        this.debugInfoDoubleWordsInSecondSlot = debugInfoDoubleWordsInSecondSlot;
     }
 
     /**
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.api.meta.test/src/com/oracle/graal/api/meta/test/TestResolvedJavaType.java
--- a/graal/com.oracle.graal.api.meta.test/src/com/oracle/graal/api/meta/test/TestResolvedJavaType.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.api.meta.test/src/com/oracle/graal/api/meta/test/TestResolvedJavaType.java	Mon Feb 25 13:14:39 2013 +0100
@@ -29,6 +29,7 @@
 
 import java.lang.annotation.*;
 import java.lang.reflect.*;
+import java.net.*;
 import java.util.*;
 
 import org.junit.*;
@@ -541,4 +542,36 @@
             }
         }
     }
+
+    @Test
+    public void memberClassesTest() {
+        for (Class c : classes) {
+            ResolvedJavaType type = runtime.lookupJavaType(c);
+            assertEquals(c.isLocalClass(), type.isLocal());
+            assertEquals(c.isMemberClass(), type.isMember());
+            Class enclc = c.getEnclosingClass();
+            ResolvedJavaType enclt = type.getEnclosingType();
+            assertFalse(enclc == null ^ enclt == null);
+            if (enclc != null) {
+                assertEquals(enclt, runtime.lookupJavaType(enclc));
+            }
+        }
+    }
+
+    @Test
+    public void classFilePathTest() {
+        for (Class c : classes) {
+            ResolvedJavaType type = runtime.lookupJavaType(c);
+            URL path = type.getClassFilePath();
+            if (type.isPrimitive() || type.isArray()) {
+                assertEquals(null, path);
+            } else {
+                assertNotNull(path);
+                String pathString = path.getPath();
+                if (type.isLocal() || type.isMember()) {
+                    assertTrue(pathString.indexOf('$') > 0);
+                }
+            }
+        }
+    }
 }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/MetaAccessProvider.java
--- a/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/MetaAccessProvider.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/MetaAccessProvider.java	Mon Feb 25 13:14:39 2013 +0100
@@ -87,10 +87,11 @@
 
     /**
      * Reads a value of this kind using a base address and a displacement.
-     *
+     * 
      * @param base the base address from which the value is read
      * @param displacement the displacement within the object in bytes
-     * @return the read value encapsulated in a {@link Constant} object, or {@code null} if the value cannot be read.
+     * @return the read value encapsulated in a {@link Constant} object, or {@code null} if the
+     *         value cannot be read.
      */
     Constant readUnsafeConstant(Kind kind, Object base, long displacement);
 }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/MetaUtil.java
--- a/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/MetaUtil.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/MetaUtil.java	Mon Feb 25 13:14:39 2013 +0100
@@ -135,7 +135,7 @@
                 String prefix = "";
                 Class<?> enclosingClass = clazz;
                 while ((enclosingClass = enclosingClass.getEnclosingClass()) != null) {
-                    prefix = prefix + enclosingClass.getSimpleName() + ".";
+                    prefix = enclosingClass.getSimpleName() + "." + prefix;
                 }
                 return prefix + simpleName;
             }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/ResolvedJavaType.java
--- a/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/ResolvedJavaType.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/ResolvedJavaType.java	Mon Feb 25 13:14:39 2013 +0100
@@ -24,6 +24,7 @@
 
 import java.lang.annotation.*;
 import java.lang.reflect.*;
+import java.net.*;
 
 /**
  * Represents a resolved Java type. Types include primitives, objects, {@code void}, and arrays
@@ -257,5 +258,20 @@
     /**
      * Returns the class file path - if available - of this type, or {@code null}.
      */
-    String getClassFilePath();
+    URL getClassFilePath();
+
+    /**
+     * Returns {@code true} if the type is a local type.
+     */
+    boolean isLocal();
+
+    /**
+     * Returns {@code true} if the type is a member type.
+     */
+    boolean isMember();
+
+    /**
+     * Returns the enclosing type of this type, if it exists, or {@code null}.
+     */
+    ResolvedJavaType getEnclosingType();
 }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.asm.amd64.test/src/com/oracle/graal/asm/amd64/test/SimpleAssemblerTest.java	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.asm.amd64/src/com/oracle/graal/asm/amd64/AMD64Assembler.java	Mon Feb 25 13:14:39 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,15 +47,38 @@
     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 {
-        zero(0x4, "|zero|"), notZero(0x5, "|nzero|"), equal(0x4, "="), notEqual(0x5, "!="), less(0xc, "<"), lessEqual(0xe, "<="), greater(0xf, ">"), greaterEqual(0xd, ">="), below(0x2, "|<|"), belowEqual(
-                        0x6, "|<=|"), above(0x7, "|>|"), aboveEqual(0x3, "|>=|"), overflow(0x0, "|of|"), noOverflow(0x1, "|nof|"), carrySet(0x2, "|carry|"), carryClear(0x3, "|ncarry|"), negative(0x8,
-                        "|neg|"), positive(0x9, "|pos|"), parity(0xa, "|par|"), noParity(0xb, "|npar|");
+        Zero(0x4, "|zero|"),
+        NotZero(0x5, "|nzero|"),
+        Equal(0x4, "="),
+        NotEqual(0x5, "!="),
+        Less(0xc, "<"),
+        LessEqual(0xe, "<="),
+        Greater(0xf, ">"),
+        GreaterEqual(0xd, ">="),
+        Below(0x2, "|<|"),
+        BelowEqual(0x6, "|<=|"),
+        Above(0x7, "|>|"),
+        AboveEqual(0x3, "|>=|"),
+        Overflow(0x0, "|of|"),
+        NoOverflow(0x1, "|nof|"),
+        CarrySet(0x2, "|carry|"),
+        CarryClear(0x3, "|ncarry|"),
+        Negative(0x8, "|neg|"),
+        Positive(0x9, "|pos|"),
+        Parity(0xa, "|par|"),
+        NoParity(0xb, "|npar|");
 
-        public final int value;
-        public final String operator;
+        private final int value;
+        private final String operator;
 
         private ConditionFlag(int value, String operator) {
             this.value = value;
@@ -64,49 +87,58 @@
 
         public ConditionFlag negate() {
             switch (this) {
-                case zero:
-                    return notZero;
-                case notZero:
-                    return zero;
-                case equal:
-                    return notEqual;
-                case notEqual:
-                    return equal;
-                case less:
-                    return greaterEqual;
-                case lessEqual:
-                    return greater;
-                case greater:
-                    return lessEqual;
-                case greaterEqual:
-                    return less;
-                case below:
-                    return aboveEqual;
-                case belowEqual:
-                    return above;
-                case above:
-                    return belowEqual;
-                case aboveEqual:
-                    return below;
-                case overflow:
-                    return noOverflow;
-                case noOverflow:
-                    return overflow;
-                case carrySet:
-                    return carryClear;
-                case carryClear:
-                    return carrySet;
-                case negative:
-                    return positive;
-                case positive:
-                    return negative;
-                case parity:
-                    return noParity;
-                case noParity:
-                    return parity;
+                case Zero:
+                    return NotZero;
+                case NotZero:
+                    return Zero;
+                case Equal:
+                    return NotEqual;
+                case NotEqual:
+                    return Equal;
+                case Less:
+                    return GreaterEqual;
+                case LessEqual:
+                    return Greater;
+                case Greater:
+                    return LessEqual;
+                case GreaterEqual:
+                    return Less;
+                case Below:
+                    return AboveEqual;
+                case BelowEqual:
+                    return Above;
+                case Above:
+                    return BelowEqual;
+                case AboveEqual:
+                    return Below;
+                case Overflow:
+                    return NoOverflow;
+                case NoOverflow:
+                    return Overflow;
+                case CarrySet:
+                    return CarryClear;
+                case CarryClear:
+                    return CarrySet;
+                case Negative:
+                    return Positive;
+                case Positive:
+                    return Negative;
+                case Parity:
+                    return NoParity;
+                case NoParity:
+                    return Parity;
             }
             throw new IllegalArgumentException();
         }
+
+        public int getValue() {
+            return value;
+        }
+
+        @Override
+        public String toString() {
+            return operator;
+        }
     }
 
     /**
@@ -155,16 +187,6 @@
         return r.encoding & 0x7;
     }
 
-    private void emitArithB(int op1, int op2, Register dst, int imm8) {
-        assert dst.isByte() : "must have byte register";
-        assert isUByte(op1) && isUByte(op2) : "wrong opcode";
-        assert isUByte(imm8) : "not a byte";
-        assert (op1 & 0x01) == 0 : "should be 8bit operation";
-        emitByte(op1);
-        emitByte(op2 | encode(dst));
-        emitByte(imm8);
-    }
-
     private void emitArith(int op1, int op2, Register dst, int imm32) {
         assert isUByte(op1) && isUByte(op2) : "wrong opcode";
         assert (op1 & 0x01) == 1 : "should be 32bit operation";
@@ -181,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)) {
@@ -201,21 +223,16 @@
         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) {
             assert frameRegister != null : "cannot use register " + Register.Frame + " in assembler with null register configuration";
             base = frameRegister;
-            // } else if (base == Register.CallerFrame) {
-            // assert frameRegister != null : "cannot use register " + Register.Frame +
-            // " in assembler with null register configuration";
-            // base = frameRegister;
-            // disp += targetMethod.frameSize() + 8;
         }
 
         // Encode the registers as needed in the fields they are used in
@@ -223,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()) {
@@ -316,23 +329,17 @@
         }
     }
 
-    public final void addl(Address dst, int imm32) {
+    public final void addl(AMD64Address dst, int imm32) {
         prefix(dst);
         emitArithOperand(0x81, rax, dst, imm32);
     }
 
-    public final void addl(Address dst, Register src) {
-        prefix(dst, src);
-        emitByte(0x01);
-        emitOperandHelper(src, dst);
-    }
-
     public final void addl(Register dst, int imm32) {
         prefix(dst);
         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);
@@ -386,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);
@@ -404,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);
@@ -418,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);
@@ -436,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);
@@ -449,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);
@@ -462,26 +469,18 @@
         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);
     }
 
-    public final void bswapl(Register reg) { // bswap
+    public final void bswapl(Register reg) {
         int encode = prefixAndEncode(reg.encoding);
         emitByte(0x0F);
         emitByte(0xC8 | encode);
     }
 
-    public final void btli(Address src, int imm8) {
-        prefixq(src);
-        emitByte(0x0F);
-        emitByte(0xBA);
-        emitOperandHelper(rsp, src);
-        emitByte(imm8);
-    }
-
     public final void cdql() {
         emitByte(0x99);
     }
@@ -489,31 +488,17 @@
     public final void cmovl(ConditionFlag cc, Register dst, Register src) {
         int encode = prefixAndEncode(dst.encoding, src.encoding);
         emitByte(0x0F);
-        emitByte(0x40 | cc.value);
+        emitByte(0x40 | cc.getValue());
         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.value);
+        emitByte(0x40 | cc.getValue());
         emitOperandHelper(dst, src);
     }
 
-    public final void cmpb(Address dst, int imm8) {
-        prefix(dst);
-        emitByte(0x80);
-        emitOperandHelper(rdi, dst);
-        emitByte(imm8);
-    }
-
-    public final void cmpl(Address dst, int imm32) {
-        prefix(dst);
-        emitByte(0x81);
-        emitOperandHelper(rdi, dst);
-        emitInt(imm32);
-    }
-
     public final void cmpl(Register dst, int imm32) {
         prefix(dst);
         emitArith(0x81, 0xF8, dst, imm32);
@@ -524,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);
@@ -533,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.
@@ -543,7 +528,7 @@
             movl(rax, adr);
             if (reg != rax) {
                 Label l = new Label();
-                jcc(ConditionFlag.notEqual, l);
+                jccb(ConditionFlag.NotEqual, l);
                 movl(adr, reg);
                 bind(l);
             }
@@ -556,43 +541,6 @@
         }
     }
 
-    public final void comisd(Register dst, Address src) {
-        assert dst.isFpu();
-        // NOTE: dbx seems to decode this as comiss even though the
-        // 0x66 is there. Strangly ucomisd comes out correct
-        emitByte(0x66);
-        comiss(dst, src);
-    }
-
-    public final void comiss(Register dst, Address src) {
-        assert dst.isFpu();
-
-        prefix(src, dst);
-        emitByte(0x0F);
-        emitByte(0x2F);
-        emitOperandHelper(dst, src);
-    }
-
-    public final void cvtdq2pd(Register dst, Register src) {
-        assert dst.isFpu();
-        assert src.isFpu();
-
-        emitByte(0xF3);
-        int encode = prefixAndEncode(dst.encoding, src.encoding);
-        emitByte(0x0F);
-        emitByte(0xE6);
-        emitByte(0xC0 | encode);
-    }
-
-    public final void cvtdq2ps(Register dst, Register src) {
-        assert dst.isFpu();
-        assert src.isFpu();
-        int encode = prefixAndEncode(dst.encoding, src.encoding);
-        emitByte(0x0F);
-        emitByte(0x5B);
-        emitByte(0xC0 | encode);
-    }
-
     public final void cvtsd2ss(Register dst, Register src) {
         assert dst.isFpu();
         assert src.isFpu();
@@ -649,14 +597,13 @@
         emitByte(0xC0 | encode);
     }
 
-    public final void decl(Address dst) {
-        // Don't use it directly. Use Macrodecrement() instead.
+    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);
@@ -675,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);
@@ -730,32 +677,31 @@
         }
     }
 
-    public final void incl(Address dst) {
-        // Don't use it directly. Use Macroincrement() instead.
+    protected final void incl(AMD64Address dst) {
         prefix(dst);
         emitByte(0xFF);
         emitOperandHelper(rax, dst);
     }
 
-    public final void jcc(ConditionFlag cc, int jumpTarget, boolean forceDisp32) {
+    private void jcc(ConditionFlag cc, int jumpTarget, boolean forceDisp32) {
         int shortSize = 2;
         int longSize = 6;
         long disp = jumpTarget - codeBuffer.position();
         if (!forceDisp32 && isByte(disp - shortSize)) {
             // 0111 tttn #8-bit disp
-            emitByte(0x70 | cc.value);
+            emitByte(0x70 | cc.getValue());
             emitByte((int) ((disp - shortSize) & 0xFF));
         } else {
             // 0000 1111 1000 tttn #32-bit disp
             assert isInt(disp - longSize) : "must be 32bit offset (call4)";
             emitByte(0x0F);
-            emitByte(0x80 | cc.value);
+            emitByte(0x80 | cc.getValue());
             emitInt((int) (disp - longSize));
         }
     }
 
     public final void jcc(ConditionFlag cc, Label l) {
-        assert (0 <= cc.value) && (cc.value < 16) : "illegal cc";
+        assert (0 <= cc.getValue()) && (cc.getValue() < 16) : "illegal cc";
         if (l.isBound()) {
             jcc(cc, l.position(), false);
         } else {
@@ -765,7 +711,7 @@
             // an 8-bit displacement
             l.addPatchAt(codeBuffer.position());
             emitByte(0x0F);
-            emitByte(0x80 | cc.value);
+            emitByte(0x80 | cc.getValue());
             emitInt(0);
         }
 
@@ -778,22 +724,16 @@
             assert isByte(entry - (codeBuffer.position() + shortSize)) : "Dispacement too large for a short jmp";
             long disp = entry - codeBuffer.position();
             // 0111 tttn #8-bit disp
-            emitByte(0x70 | cc.value);
+            emitByte(0x70 | cc.getValue());
             emitByte((int) ((disp - shortSize) & 0xFF));
         } else {
 
             l.addPatchAt(codeBuffer.position());
-            emitByte(0x70 | cc.value);
+            emitByte(0x70 | cc.getValue());
             emitByte(0);
         }
     }
 
-    public final void jmp(Address adr) {
-        prefix(adr);
-        emitByte(0xFF);
-        emitOperandHelper(rsp, adr);
-    }
-
     public final void jmp(int jumpTarget, boolean forceDisp32) {
         int shortSize = 2;
         int longSize = 5;
@@ -845,18 +785,12 @@
         }
     }
 
-    public final void leaq(Register dst, Address src) {
+    public final void leaq(Register dst, AMD64Address src) {
         prefixq(src, dst);
         emitByte(0x8D);
         emitOperandHelper(dst, src);
     }
 
-    public final void enter(int imm16, int imm8) {
-        emitByte(0xC8);
-        emitShort(imm16);
-        emitByte(imm8);
-    }
-
     public final void leave() {
         emitByte(0xC9);
     }
@@ -870,17 +804,6 @@
         }
     }
 
-    // Emit mfence instruction
-    public final void mfence() {
-        emitByte(0x0F);
-        emitByte(0xAE);
-        emitByte(0xF0);
-    }
-
-    public final void mov(Register dst, Register src) {
-        movq(dst, src);
-    }
-
     public final void movapd(Register dst, Register src) {
         assert dst.isFpu();
         assert src.isFpu();
@@ -930,20 +853,14 @@
         emitByte(0xC0 | dstenc << 3 | srcenc);
     }
 
-    public final void movb(Register dst, Address src) {
-        prefix(src, dst); // , true)
-        emitByte(0x8A);
-        emitOperandHelper(dst, src);
-    }
-
-    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);
@@ -969,63 +886,6 @@
         }
     }
 
-    public final void movdqa(Register dst, Address src) {
-        assert dst.isFpu();
-        emitByte(0x66);
-        prefix(src, dst);
-        emitByte(0x0F);
-        emitByte(0x6F);
-        emitOperandHelper(dst, src);
-    }
-
-    public final void movdqa(Register dst, Register src) {
-        assert dst.isFpu();
-        emitByte(0x66);
-        int encode = prefixqAndEncode(dst.encoding, src.encoding);
-        emitByte(0x0F);
-        emitByte(0x6F);
-        emitByte(0xC0 | encode);
-    }
-
-    public final void movdqa(Address dst, Register src) {
-        assert src.isFpu();
-        emitByte(0x66);
-        prefix(dst, src);
-        emitByte(0x0F);
-        emitByte(0x7F);
-        emitOperandHelper(src, dst);
-    }
-
-    public final void movdqu(Register dst, Address src) {
-        assert dst.isFpu();
-        emitByte(0xF3);
-        prefix(src, dst);
-        emitByte(0x0F);
-        emitByte(0x6F);
-        emitOperandHelper(dst, src);
-    }
-
-    public final void movdqu(Register dst, Register src) {
-        assert dst.isFpu();
-        assert src.isFpu();
-
-        emitByte(0xF3);
-        int encode = prefixqAndEncode(dst.encoding, src.encoding);
-        emitByte(0x0F);
-        emitByte(0x6F);
-        emitByte(0xC0 | encode);
-    }
-
-    public final void movdqu(Address dst, Register src) {
-        assert src.isFpu();
-
-        emitByte(0xF3);
-        prefix(dst, src);
-        emitByte(0x0F);
-        emitByte(0x7F);
-        emitOperandHelper(src, dst);
-    }
-
     public final void movl(Register dst, int imm32) {
         int encode = prefixAndEncode(dst.encoding);
         emitByte(0xB8 | encode);
@@ -1038,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);
@@ -1060,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);
@@ -1072,16 +932,7 @@
         emitOperandHelper(dst, src);
     }
 
-    public final void movlpd(Address dst, Register src) {
-        assert src.isFpu();
-        emitByte(0x66);
-        prefix(dst, src);
-        emitByte(0x0F);
-        emitByte(0x13);
-        emitOperandHelper(src, dst);
-    }
-
-    public final void movq(Register dst, Address src) {
+    public final void movq(Register dst, AMD64Address src) {
         if (dst.isFpu()) {
             emitByte(0xF3);
             prefixq(src, dst);
@@ -1101,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);
@@ -1115,14 +966,14 @@
         }
     }
 
-    public final void movsxb(Register dst, Address src) { // movsxb
+    public final void movsxb(Register dst, AMD64Address src) {
         prefix(src, dst);
         emitByte(0x0F);
         emitByte(0xBE);
         emitOperandHelper(dst, src);
     }
 
-    public final void movsxb(Register dst, Register src) { // movsxb
+    public final void movsxb(Register dst, Register src) {
         int encode = prefixAndEncode(dst.encoding, src.encoding, true);
         emitByte(0x0F);
         emitByte(0xBE);
@@ -1139,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);
@@ -1148,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);
@@ -1167,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);
@@ -1176,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);
@@ -1185,40 +1036,21 @@
         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);
         emitOperandHelper(dst, src);
     }
 
-    public final void movsxw(Register dst, Register src) { // movsxw
+    public final void movsxw(Register dst, Register src) {
         int encode = prefixAndEncode(dst.encoding, src.encoding);
         emitByte(0x0F);
         emitByte(0xBF);
         emitByte(0xC0 | encode);
     }
 
-    public final void movsxw(Register dst, Address src) { // movsxw
-        prefix(src, dst);
-        emitByte(0x0F);
-        emitByte(0xBF);
-        emitOperandHelper(dst, src);
-    }
-
-    public final void movzxd(Register dst, Register src) { // movzxd
-        int encode = prefixAndEncode(dst.encoding, src.encoding);
-        emitByte(0x63);
-        emitByte(0xC0 | encode);
-    }
-
-    public final void movzxd(Register dst, Address src) { // movzxd
-        prefix(src, dst);
-        emitByte(0x63);
-        emitOperandHelper(dst, src);
-    }
-
-    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);
@@ -1226,55 +1058,21 @@
         emitShort(imm16);
     }
 
-    public final void movw(Register dst, Address src) {
-        emitByte(0x66);
-        prefix(src, dst);
-        emitByte(0x8B);
-        emitOperandHelper(dst, src);
-    }
-
-    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 movzxb(Register dst, Address src) { // movzxb
-        prefix(src, dst);
-        emitByte(0x0F);
-        emitByte(0xB6);
-        emitOperandHelper(dst, src);
-    }
-
-    public final void movzxb(Register dst, Register src) { // movzxb
-        int encode = prefixAndEncode(dst.encoding, src.encoding, true);
-        emitByte(0x0F);
-        emitByte(0xB6);
-        emitByte(0xC0 | encode);
-    }
-
-    public final void movzxl(Register dst, Address src) { // movzxw
+    public final void movzxl(Register dst, AMD64Address src) {
         prefix(src, dst);
         emitByte(0x0F);
         emitByte(0xB7);
         emitOperandHelper(dst, src);
     }
 
-    public final void movzxl(Register dst, Register src) { // movzxw
-        int encode = prefixAndEncode(dst.encoding, src.encoding);
-        emitByte(0x0F);
-        emitByte(0xB7);
-        emitByte(0xC0 | encode);
-    }
-
-    public final void mull(Address src) {
-        prefix(src);
-        emitByte(0xF7);
-        emitOperandHelper(rsp, 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);
@@ -1294,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);
@@ -1530,25 +1328,12 @@
         }
     }
 
-    public final void notl(Register dst) {
-        int encode = prefixAndEncode(dst.encoding);
-        emitByte(0xF7);
-        emitByte(0xD0 | encode);
-    }
-
-    public final void orl(Address dst, int imm32) {
-        prefix(dst);
-        emitByte(0x81);
-        emitOperandHelper(rcx, dst);
-        emitInt(imm32);
-    }
-
     public final void orl(Register dst, int imm32) {
         prefix(dst);
         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);
@@ -1559,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);
@@ -1575,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);
@@ -1591,223 +1376,16 @@
         emitByte(0xC0 | encode);
     }
 
-    // generic
     public final void pop(Register dst) {
         int encode = prefixAndEncode(dst.encoding);
         emitByte(0x58 | encode);
     }
 
-    public final void prefetchPrefix(Address src) {
-        prefix(src);
-        emitByte(0x0F);
-    }
-
-    public final void prefetchnta(Address src) {
-        prefetchPrefix(src);
-        emitByte(0x18);
-        emitOperandHelper(rax, src); // 0, src
-    }
-
-    public final void prefetchr(Address src) {
-        prefetchPrefix(src);
-        emitByte(0x0D);
-        emitOperandHelper(rax, src); // 0, src
-    }
-
-    public final void prefetcht0(Address src) {
-        prefetchPrefix(src);
-        emitByte(0x18);
-        emitOperandHelper(rcx, src); // 1, src
-
-    }
-
-    public final void prefetcht1(Address src) {
-        prefetchPrefix(src);
-        emitByte(0x18);
-        emitOperandHelper(rdx, src); // 2, src
-    }
-
-    public final void prefetcht2(Address src) {
-        prefetchPrefix(src);
-        emitByte(0x18);
-        emitOperandHelper(rbx, src); // 3, src
-    }
-
-    public final void prefetchw(Address src) {
-        prefetchPrefix(src);
-        emitByte(0x0D);
-        emitOperandHelper(rcx, src); // 1, src
-    }
-
-    public final void pshufd(Register dst, Register src, int mode) {
-        assert dst.isFpu();
-        assert src.isFpu();
-        assert isUByte(mode) : "invalid value";
-
-        emitByte(0x66);
-        int encode = prefixAndEncode(dst.encoding, src.encoding);
-        emitByte(0x0F);
-        emitByte(0x70);
-        emitByte(0xC0 | encode);
-        emitByte(mode & 0xFF);
-    }
-
-    public final void pshufd(Register dst, Address src, int mode) {
-        assert dst.isFpu();
-        assert isUByte(mode) : "invalid value";
-
-        emitByte(0x66);
-        prefix(src, dst);
-        emitByte(0x0F);
-        emitByte(0x70);
-        emitOperandHelper(dst, src);
-        emitByte(mode & 0xFF);
-
-    }
-
-    public final void pshuflw(Register dst, Register src, int mode) {
-        assert dst.isFpu();
-        assert src.isFpu();
-        assert isUByte(mode) : "invalid value";
-
-        emitByte(0xF2);
-        int encode = prefixAndEncode(dst.encoding, src.encoding);
-        emitByte(0x0F);
-        emitByte(0x70);
-        emitByte(0xC0 | encode);
-        emitByte(mode & 0xFF);
-    }
-
-    public final void pshuflw(Register dst, Address src, int mode) {
-        assert dst.isFpu();
-        assert isUByte(mode) : "invalid value";
-
-        emitByte(0xF2);
-        prefix(src, dst); // QQ new
-        emitByte(0x0F);
-        emitByte(0x70);
-        emitOperandHelper(dst, src);
-        emitByte(mode & 0xFF);
-    }
-
-    public final void psrlq(Register dst, int shift) {
-        assert dst.isFpu();
-        // HMM Table D-1 says sse2 or mmx
-
-        int encode = prefixqAndEncode(xmm2.encoding, dst.encoding);
-        emitByte(0x66);
-        emitByte(0x0F);
-        emitByte(0x73);
-        emitByte(0xC0 | encode);
-        emitByte(shift);
-    }
-
-    public final void punpcklbw(Register dst, Register src) {
-        assert dst.isFpu();
-        assert src.isFpu();
-        emitByte(0x66);
-        int encode = prefixAndEncode(dst.encoding, src.encoding);
-        emitByte(0x0F);
-        emitByte(0x60);
-        emitByte(0xC0 | encode);
-    }
-
-    public final void push(int imm32) {
-        // in 64bits we push 64bits onto the stack but only
-        // take a 32bit immediate
-        emitByte(0x68);
-        emitInt(imm32);
-    }
-
     public final void push(Register src) {
         int encode = prefixAndEncode(src.encoding);
         emitByte(0x50 | encode);
     }
 
-    public final void pushf() {
-        emitByte(0x9C);
-    }
-
-    public final void pxor(Register dst, Address src) {
-        assert dst.isFpu();
-
-        emitByte(0x66);
-        prefix(src, dst);
-        emitByte(0x0F);
-        emitByte(0xEF);
-        emitOperandHelper(dst, src);
-    }
-
-    public final void pxor(Register dst, Register src) {
-        assert dst.isFpu();
-        assert src.isFpu();
-
-        emitByte(0x66);
-        int encode = prefixAndEncode(dst.encoding, src.encoding);
-        emitByte(0x0F);
-        emitByte(0xEF);
-        emitByte(0xC0 | encode);
-
-    }
-
-    public final void rcll(Register dst, int imm8) {
-        assert isShiftCount(imm8) : "illegal shift count";
-        int encode = prefixAndEncode(dst.encoding);
-        if (imm8 == 1) {
-            emitByte(0xD1);
-            emitByte(0xD0 | encode);
-        } else {
-            emitByte(0xC1);
-            emitByte(0xD0 | encode);
-            emitByte(imm8);
-        }
-    }
-
-    public final void pause() {
-        emitByte(0xF3);
-        emitByte(0x90);
-    }
-
-    // Copies data from [X86.rsi] to [X86.rdi] using X86.rcx heap words.
-    public final void repeatMoveWords() {
-        emitByte(0xF3);
-        emitByte(Prefix.REXW);
-        emitByte(0xA5);
-    }
-
-    // Copies data from [X86.rsi] to [X86.rdi] using X86.rcx bytes.
-    public final void repeatMoveBytes() {
-        emitByte(0xF3);
-        emitByte(Prefix.REXW);
-        emitByte(0xA4);
-    }
-
-    // sets X86.rcx pointer sized words with X86.rax, value at [edi]
-    // generic
-    public final void repSet() { // repSet
-        emitByte(0xF3);
-        // STOSQ
-        emitByte(Prefix.REXW);
-        emitByte(0xAB);
-    }
-
-    // scans X86.rcx pointer sized words at [edi] for occurance of X86.rax,
-    // generic
-    public final void repneScan() { // repneScan
-        emitByte(0xF2);
-        // SCASQ
-        emitByte(Prefix.REXW);
-        emitByte(0xAF);
-    }
-
-    // scans X86.rcx 4 byte words at [edi] for occurance of X86.rax,
-    // generic
-    public final void repneScanl() { // repneScan
-        emitByte(0xF2);
-        // SCASL
-        emitByte(0xAF);
-    }
-
     public final void ret(int imm16) {
         if (imm16 == 0) {
             emitByte(0xC3);
@@ -1836,35 +1414,6 @@
         emitByte(0xF8 | encode);
     }
 
-    public final void sbbl(Address dst, int imm32) {
-        prefix(dst);
-        emitArithOperand(0x81, rbx, dst, imm32);
-    }
-
-    public final void sbbl(Register dst, int imm32) {
-        prefix(dst);
-        emitArith(0x81, 0xD8, dst, imm32);
-    }
-
-    public final void sbbl(Register dst, Address src) {
-        prefix(src, dst);
-        emitByte(0x1B);
-        emitOperandHelper(dst, src);
-    }
-
-    public final void sbbl(Register dst, Register src) {
-        prefixAndEncode(dst.encoding, src.encoding);
-        emitArith(0x1B, 0xC0, dst, src);
-    }
-
-    public final void setb(ConditionFlag cc, Register dst) {
-        assert 0 <= cc.value && cc.value < 16 : "illegal cc";
-        int encode = prefixAndEncode(dst.encoding, true);
-        emitByte(0x0F);
-        emitByte(0x90 | cc.value);
-        emitByte(0xC0 | encode);
-    }
-
     public final void shll(Register dst, int imm8) {
         assert isShiftCount(imm8) : "illegal shift count";
         int encode = prefixAndEncode(dst.encoding);
@@ -1898,11 +1447,6 @@
         emitByte(0xE8 | encode);
     }
 
-    // copies a single word from [esi] to [edi]
-    public final void smovl() {
-        emitByte(0xA5);
-    }
-
     public final void sqrtsd(Register dst, Register src) {
         assert dst.isFpu();
         assert src.isFpu();
@@ -1915,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);
@@ -1933,13 +1477,7 @@
         emitArith(0x81, 0xE8, dst, imm32);
     }
 
-    public final void subl(Address dst, Register src) {
-        prefix(dst, src);
-        emitByte(0x29);
-        emitOperandHelper(src, dst);
-    }
-
-    public final void subl(Register dst, Address src) {
+    public final void subl(Register dst, AMD64Address src) {
         prefix(src, dst);
         emitByte(0x2B);
         emitOperandHelper(dst, src);
@@ -1960,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);
@@ -1980,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);
@@ -1990,11 +1528,6 @@
         emitOperandHelper(dst, src);
     }
 
-    public final void testb(Register dst, int imm8) {
-        prefixAndEncode(dst.encoding, true);
-        emitArithB(0xF6, 0xC0, dst, imm8);
-    }
-
     public final void testl(Register dst, int imm32) {
         // not using emitArith because test
         // doesn't support sign-extension of
@@ -2015,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);
@@ -2034,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);
@@ -2052,33 +1585,12 @@
         emitByte(0xC0 | encode);
     }
 
-    public final void xaddl(Address dst, Register src) {
-        assert src.isFpu();
-
-        prefix(dst, src);
-        emitByte(0x0F);
-        emitByte(0xC1);
-        emitOperandHelper(src, dst);
-    }
-
-    public final void xchgl(Register dst, Address src) { // xchg
-        prefix(src, dst);
-        emitByte(0x87);
-        emitOperandHelper(dst, src);
-    }
-
-    public final void xchgl(Register dst, Register src) {
-        int encode = prefixAndEncode(dst.encoding, src.encoding);
-        emitByte(0x87);
-        emitByte(0xc0 | encode);
-    }
-
     public final void xorl(Register dst, int imm32) {
         prefix(dst);
         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);
@@ -2094,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);
     }
@@ -2107,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);
@@ -2120,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);
     }
@@ -2133,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);
@@ -2146,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);
     }
@@ -2159,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);
@@ -2167,29 +1679,25 @@
         emitOperandHelper(dst, src);
     }
 
-    // 32bit only pieces of the assembler
-
-    public final void decl(Register dst) {
-        // Don't use it directly. Use Macrodecrementl() instead.
+    protected final void decl(Register dst) {
         // Use two-byte form (one-byte form is a REX prefix in 64-bit mode)
         int encode = prefixAndEncode(dst.encoding);
         emitByte(0xFF);
         emitByte(0xC8 | encode);
     }
 
-    public final void incl(Register dst) {
-        // Don't use it directly. Use Macroincrementl() instead.
+    protected final void incl(Register dst) {
         // Use two-byte form (one-byte from is a REX prefix in 64-bit mode)
         int encode = prefixAndEncode(dst.encoding);
         emitByte(0xFF);
         emitByte(0xC0 | encode);
     }
 
-    int prefixAndEncode(int regEnc) {
+    private int prefixAndEncode(int regEnc) {
         return prefixAndEncode(regEnc, false);
     }
 
-    int prefixAndEncode(int regEnc, boolean byteinst) {
+    private int prefixAndEncode(int regEnc, boolean byteinst) {
         if (regEnc >= 8) {
             emitByte(Prefix.REXB);
             return regEnc - 8;
@@ -2199,7 +1707,7 @@
         return regEnc;
     }
 
-    int prefixqAndEncode(int regEnc) {
+    private int prefixqAndEncode(int regEnc) {
         if (regEnc < 8) {
             emitByte(Prefix.REXW);
             return regEnc;
@@ -2209,11 +1717,11 @@
         }
     }
 
-    int prefixAndEncode(int dstEnc, int srcEnc) {
+    private int prefixAndEncode(int dstEnc, int srcEnc) {
         return prefixAndEncode(dstEnc, srcEnc, false);
     }
 
-    int prefixAndEncode(int dstEncoding, int srcEncoding, boolean byteinst) {
+    private int prefixAndEncode(int dstEncoding, int srcEncoding, boolean byteinst) {
         int srcEnc = srcEncoding;
         int dstEnc = dstEncoding;
         if (dstEnc < 8) {
@@ -2275,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);
@@ -2289,7 +1797,7 @@
         }
     }
 
-    private void prefixq(Address adr) {
+    private void prefixq(AMD64Address adr) {
         if (needsRex(adr.getBase())) {
             if (needsRex(adr.getIndex())) {
                 emitByte(Prefix.REXWXB);
@@ -2305,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())) {
@@ -2337,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())) {
@@ -2369,23 +1877,12 @@
         }
     }
 
-    public final void addq(Address dst, int imm32) {
-        prefixq(dst);
-        emitArithOperand(0x81, rax, dst, imm32);
-    }
-
-    public final void addq(Address dst, Register src) {
-        prefixq(dst, src);
-        emitByte(0x01);
-        emitOperandHelper(src, dst);
-    }
-
     public final void addq(Register dst, int imm32) {
         prefixqAndEncode(dst.encoding);
         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);
@@ -2401,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);
@@ -2426,47 +1923,34 @@
     public final void cmovq(ConditionFlag cc, Register dst, Register src) {
         int encode = prefixqAndEncode(dst.encoding, src.encoding);
         emitByte(0x0F);
-        emitByte(0x40 | cc.value);
+        emitByte(0x40 | cc.getValue());
         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.value);
+        emitByte(0x40 | cc.getValue());
         emitOperandHelper(dst, src);
     }
 
-    public final void cmpq(Address dst, int imm32) {
-        prefixq(dst);
-        emitByte(0x81);
-        emitOperandHelper(rdi, dst);
-        emitInt(imm32);
-    }
-
     public final void cmpq(Register dst, int imm32) {
         prefixqAndEncode(dst.encoding);
         emitArith(0x81, 0xF8, dst, imm32);
     }
 
-    public final void cmpq(Address dst, Register src) {
-        prefixq(dst, src);
-        emitByte(0x3B);
-        emitOperandHelper(src, dst);
-    }
-
     public final void cmpq(Register dst, Register src) {
         prefixqAndEncode(dst.encoding, src.encoding);
         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);
@@ -2509,16 +1993,14 @@
         emitByte(0xC0 | encode);
     }
 
-    public final void decq(Register dst) {
-        // Don't use it directly. Use Macrodecrementq() instead.
+    protected final void decq(Register dst) {
         // Use two-byte form (one-byte from is a REX prefix in 64-bit mode)
         int encode = prefixqAndEncode(dst.encoding);
         emitByte(0xFF);
         emitByte(0xC8 | encode);
     }
 
-    public final void decq(Address dst) {
-        // Don't use it directly. Use Macrodecrementq() instead.
+    protected final void decq(AMD64Address dst) {
         prefixq(dst);
         emitByte(0xFF);
         emitOperandHelper(rcx, dst);
@@ -2564,13 +2046,6 @@
         emitByte(0xC0 | encode);
     }
 
-    public final void incq(Address dst) {
-        // Don't use it directly. Use Macroincrementq() instead.
-        prefixq(dst);
-        emitByte(0xFF);
-        emitOperandHelper(rax, dst);
-    }
-
     public final void movq(Register dst, long imm64) {
         int encode = prefixqAndEncode(dst.encoding);
         emitByte(0xB8 | encode);
@@ -2600,38 +2075,14 @@
         }
     }
 
-    public final void movsbq(Register dst, Address src) {
-        prefixq(src, dst);
-        emitByte(0x0F);
-        emitByte(0xBE);
-        emitOperandHelper(dst, src);
-    }
-
-    public final void movsbq(Register dst, Register src) {
-        int encode = prefixqAndEncode(dst.encoding, src.encoding);
-        emitByte(0x0F);
-        emitByte(0xBE);
-        emitByte(0xC0 | encode);
-    }
-
-    public final void movslq(Register dst, int imm32) {
-        int encode = prefixqAndEncode(dst.encoding);
-        emitByte(0xC7 | encode);
-        emitInt(imm32);
-        // dbx shows movslq(X86.rcx, 3) as movq $0x0000000049000000,(%X86.rbx)
-        // and movslq(X86.r8, 3); as movl $0x0000000048000000,(%X86.rbx)
-        // as a result we shouldn't use until tested at runtime...
-        throw new InternalError("untested");
-    }
-
-    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);
@@ -2643,73 +2094,18 @@
         emitByte(0xC0 | encode);
     }
 
-    public final void movswq(Register dst, Address src) {
-        prefixq(src, dst);
-        emitByte(0x0F);
-        emitByte(0xBF);
-        emitOperandHelper(dst, src);
-    }
-
-    public final void movswq(Register dst, Register src) {
-        int encode = prefixqAndEncode(dst.encoding, src.encoding);
-        emitByte(0x0F);
-        emitByte(0xBF);
-        emitByte(0xC0 | encode);
-    }
-
-    public final void movzbq(Register dst, Address src) {
-        prefixq(src, dst);
-        emitByte(0x0F);
-        emitByte(0xB6);
-        emitOperandHelper(dst, src);
-    }
-
-    public final void movzbq(Register dst, Register src) {
-        int encode = prefixqAndEncode(dst.encoding, src.encoding);
-        emitByte(0x0F);
-        emitByte(0xB6);
-        emitByte(0xC0 | encode);
-    }
-
-    public final void movzwq(Register dst, Address src) {
-        prefixq(src, dst);
-        emitByte(0x0F);
-        emitByte(0xB7);
-        emitOperandHelper(dst, src);
-    }
-
-    public final void movzwq(Register dst, Register src) {
-        int encode = prefixqAndEncode(dst.encoding, src.encoding);
-        emitByte(0x0F);
-        emitByte(0xB7);
-        emitByte(0xC0 | encode);
-    }
-
     public final void negq(Register dst) {
         int encode = prefixqAndEncode(dst.encoding);
         emitByte(0xF7);
         emitByte(0xD8 | encode);
     }
 
-    public final void notq(Register dst) {
-        int encode = prefixqAndEncode(dst.encoding);
-        emitByte(0xF7);
-        emitByte(0xD0 | encode);
-    }
-
-    public final void orq(Address dst, int imm32) {
-        prefixq(dst);
-        emitByte(0x81);
-        emitOperandHelper(rcx, dst);
-        emitInt(imm32);
-    }
-
     public final void orq(Register dst, int imm32) {
         prefixqAndEncode(dst.encoding);
         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);
@@ -2720,31 +2116,6 @@
         emitArith(0x0B, 0xC0, dst, src);
     }
 
-    public final void popq(Address dst) {
-        prefixq(dst);
-        emitByte(0x8F);
-        emitOperandHelper(rax, dst);
-    }
-
-    public final void pushq(Address src) {
-        prefixq(src);
-        emitByte(0xFF);
-        emitOperandHelper(rsi, src);
-    }
-
-    public final void rclq(Register dst, int imm8) {
-        assert isShiftCount(imm8 >> 1) : "illegal shift count";
-        int encode = prefixqAndEncode(dst.encoding);
-        if (imm8 == 1) {
-            emitByte(0xD1);
-            emitByte(0xD0 | encode);
-        } else {
-            emitByte(0xC1);
-            emitByte(0xD0 | encode);
-            emitByte(imm8);
-        }
-    }
-
     public final void sarq(Register dst, int imm8) {
         assert isShiftCount(imm8 >> 1) : "illegal shift count";
         int encode = prefixqAndEncode(dst.encoding);
@@ -2797,41 +2168,12 @@
         emitByte(0xE8 | encode);
     }
 
-    public final void sqrtsd(Register dst, Address src) {
-        assert dst.isFpu();
-
-        emitByte(0xF2);
-        prefix(src, dst);
-        emitByte(0x0F);
-        emitByte(0x51);
-        emitOperandHelper(dst, src);
-    }
-
-    public final void subq(Address dst, int imm32) {
-        prefixq(dst);
-        if (isByte(imm32)) {
-            emitByte(0x83);
-            emitOperandHelper(rbp, dst);
-            emitByte(imm32 & 0xFF);
-        } else {
-            emitByte(0x81);
-            emitOperandHelper(rbp, dst);
-            emitInt(imm32);
-        }
-    }
-
     public final void subq(Register dst, int imm32) {
         prefixqAndEncode(dst.encoding);
         emitArith(0x81, 0xE8, dst, imm32);
     }
 
-    public final void subq(Address dst, Register src) {
-        prefixq(dst, src);
-        emitByte(0x29);
-        emitOperandHelper(src, dst);
-    }
-
-    public final void subq(Register dst, Address src) {
+    public final void subq(Register dst, AMD64Address src) {
         prefixq(src, dst);
         emitByte(0x2B);
         emitOperandHelper(dst, src);
@@ -2863,31 +2205,12 @@
         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);
     }
 
-    public final void xaddq(Address dst, Register src) {
-        prefixq(dst, src);
-        emitByte(0x0F);
-        emitByte(0xC1);
-        emitOperandHelper(src, dst);
-    }
-
-    public final void xchgq(Register dst, Address src) {
-        prefixq(src, dst);
-        emitByte(0x87);
-        emitOperandHelper(dst, src);
-    }
-
-    public final void xchgq(Register dst, Register src) {
-        int encode = prefixqAndEncode(dst.encoding, src.encoding);
-        emitByte(0x87);
-        emitByte(0xc0 | encode);
-    }
-
     public final void xorq(Register dst, int imm32) {
         prefixqAndEncode(dst.encoding);
         emitArith(0x81, 0xF0, dst, imm32);
@@ -2898,12 +2221,10 @@
         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);
-
     }
 
     public final void membar(int barriers) {
@@ -2919,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
             }
         }
     }
@@ -2960,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
@@ -2970,26 +2291,6 @@
         }
     }
 
-    public void pushfq() {
-        emitByte(0x9c);
-    }
-
-    public void popfq() {
-        emitByte(0x9D);
-    }
-
-    /**
-     * Makes sure that a subsequent {@linkplain #call} does not fail the alignment check.
-     */
-    public final void alignForPatchableDirectCall() {
-        int dispStart = codeBuffer.position() + 1;
-        int mask = target.wordSize - 1;
-        if ((dispStart & ~mask) != ((dispStart + 3) & ~mask)) {
-            nop(target.wordSize - (dispStart & mask));
-            assert ((codeBuffer.position() + 1) & mask) == 0;
-        }
-    }
-
     /**
      * Emits a direct call instruction. Note that the actual call target is not specified, because
      * all calls need patching anyway. Therefore, 0 is emitted as the call target, and the user is
@@ -3006,69 +2307,67 @@
         emitByte(0xD0 | encode);
     }
 
-    public void int3() {
+    public final void int3() {
         emitByte(0xCC);
     }
 
-    public void enter(short imm16, byte imm8) {
-        emitByte(0xC8);
-        // appended:
-        emitByte(imm16 & 0xff);
-        emitByte((imm16 >> 8) & 0xff);
-        emitByte(imm8);
-    }
-
     private void emitx87(int b1, int b2, int i) {
         assert 0 <= i && i < 8 : "illegal stack offset";
         emitByte(b1);
         emitByte(b2 + i);
     }
 
-    public void fld(Address src) {
+    public final void fld(AMD64Address src) {
         emitByte(0xDD);
         emitOperandHelper(rax, src);
     }
 
-    public void fld(int i) {
-        emitx87(0xD9, 0xC0, i);
-    }
-
-    public void fldln2() {
+    public final void fldln2() {
         emitByte(0xD9);
         emitByte(0xED);
     }
 
-    public void fldlg2() {
+    public final void fldlg2() {
         emitByte(0xD9);
         emitByte(0xEC);
     }
 
-    public void fyl2x() {
+    public final void fyl2x() {
         emitByte(0xD9);
         emitByte(0xF1);
     }
 
-    public void fstp(Address src) {
+    public final void fstp(AMD64Address src) {
         emitByte(0xDD);
         emitOperandHelper(rbx, src);
     }
 
-    public void fsin() {
+    public final void fsin() {
         emitByte(0xD9);
         emitByte(0xFE);
     }
 
-    public void fcos() {
+    public final void fcos() {
         emitByte(0xD9);
         emitByte(0xFF);
     }
 
-    public void fptan() {
+    public final void fptan() {
         emitByte(0xD9);
         emitByte(0xF2);
     }
 
-    public void fstp(int i) {
+    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 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.asm.amd64/src/com/oracle/graal/asm/amd64/AMD64MacroAssembler.java	Mon Feb 25 13:14:39 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
@@ -27,7 +27,6 @@
 import com.oracle.graal.amd64.*;
 import com.oracle.graal.api.code.*;
 import com.oracle.graal.api.meta.*;
-import com.oracle.graal.asm.*;
 
 /**
  * This class implements commonly used X86 code patterns.
@@ -38,25 +37,11 @@
         super(target, registerConfig);
     }
 
-    public void pushptr(Address src) {
-        pushq(src);
-    }
-
-    public void popptr(Address src) {
-        popq(src);
-    }
-
-    public void xorptr(Register dst, Register src) {
+    public final void xorptr(Register dst, Register src) {
         xorq(dst, src);
     }
 
-    public void xorptr(Register dst, Address src) {
-        xorq(dst, src);
-    }
-
-    // 64 bit versions
-
-    public void decrementq(Register reg, int value) {
+    public final void decrementq(Register reg, int value) {
         if (value == Integer.MIN_VALUE) {
             subq(reg, value);
             return;
@@ -94,83 +79,19 @@
         }
     }
 
-    // These are mostly for initializing null
-    public void movptr(Address dst, int src) {
+    public final void movptr(AMD64Address dst, int src) {
         movslq(dst, src);
     }
 
-    public final void cmp32(Register src1, int imm) {
-        cmpl(src1, imm);
-    }
-
-    public final void cmp32(Register src1, Address src2) {
-        cmpl(src1, src2);
-    }
-
-    public void cmpsd2int(Register opr1, Register opr2, Register dst, boolean unorderedIsLess) {
-        assert opr1.isFpu() && opr2.isFpu();
-        ucomisd(opr1, opr2);
-
-        Label l = new Label();
-        if (unorderedIsLess) {
-            movl(dst, -1);
-            jcc(AMD64Assembler.ConditionFlag.parity, l);
-            jcc(AMD64Assembler.ConditionFlag.below, l);
-            movl(dst, 0);
-            jcc(AMD64Assembler.ConditionFlag.equal, l);
-            incrementl(dst, 1);
-        } else { // unordered is greater
-            movl(dst, 1);
-            jcc(AMD64Assembler.ConditionFlag.parity, l);
-            jcc(AMD64Assembler.ConditionFlag.above, l);
-            movl(dst, 0);
-            jcc(AMD64Assembler.ConditionFlag.equal, l);
-            decrementl(dst, 1);
-        }
-        bind(l);
-    }
-
-    public void cmpss2int(Register opr1, Register opr2, Register dst, boolean unorderedIsLess) {
-        assert opr1.isFpu();
-        assert opr2.isFpu();
-        ucomiss(opr1, opr2);
-
-        Label l = new Label();
-        if (unorderedIsLess) {
-            movl(dst, -1);
-            jcc(AMD64Assembler.ConditionFlag.parity, l);
-            jcc(AMD64Assembler.ConditionFlag.below, l);
-            movl(dst, 0);
-            jcc(AMD64Assembler.ConditionFlag.equal, l);
-            incrementl(dst, 1);
-        } else { // unordered is greater
-            movl(dst, 1);
-            jcc(AMD64Assembler.ConditionFlag.parity, l);
-            jcc(AMD64Assembler.ConditionFlag.above, l);
-            movl(dst, 0);
-            jcc(AMD64Assembler.ConditionFlag.equal, l);
-            decrementl(dst, 1);
-        }
-        bind(l);
-    }
-
-    public void cmpptr(Register src1, Register src2) {
+    public final void cmpptr(Register src1, Register src2) {
         cmpq(src1, src2);
     }
 
-    public void cmpptr(Register src1, Address src2) {
+    public final void cmpptr(Register src1, AMD64Address src2) {
         cmpq(src1, src2);
     }
 
-    public void cmpptr(Register src1, int src2) {
-        cmpq(src1, src2);
-    }
-
-    public void cmpptr(Address src1, int src2) {
-        cmpq(src1, src2);
-    }
-
-    public void decrementl(Register reg, int value) {
+    public final void decrementl(Register reg, int value) {
         if (value == Integer.MIN_VALUE) {
             subl(reg, value);
             return;
@@ -189,7 +110,7 @@
         }
     }
 
-    public void decrementl(Address dst, int value) {
+    public final void decrementl(AMD64Address dst, int value) {
         if (value == Integer.MIN_VALUE) {
             subl(dst, value);
             return;
@@ -208,7 +129,7 @@
         }
     }
 
-    public void incrementl(Register reg, int value) {
+    public final void incrementl(Register reg, int value) {
         if (value == Integer.MIN_VALUE) {
             addl(reg, value);
             return;
@@ -227,7 +148,7 @@
         }
     }
 
-    public void incrementl(Address dst, int value) {
+    public final void incrementl(AMD64Address dst, int value) {
         if (value == Integer.MIN_VALUE) {
             addl(dst, value);
             return;
@@ -246,21 +167,20 @@
         }
     }
 
-    public void signExtendByte(Register reg) {
+    public final void signExtendByte(Register reg) {
         if (reg.isByte()) {
-            movsxb(reg, reg); // movsxb
+            movsxb(reg, reg);
         } else {
             shll(reg, 24);
             sarl(reg, 24);
         }
     }
 
-    public void signExtendShort(Register reg) {
-        movsxw(reg, reg); // movsxw
+    public final void signExtendShort(Register reg) {
+        movsxw(reg, reg);
     }
 
-    // Support optimal SSE move instructions.
-    public void movflt(Register dst, Register src) {
+    public final void movflt(Register dst, Register src) {
         assert dst.isFpu() && src.isFpu();
         if (UseXmmRegToRegMoveAll) {
             movaps(dst, src);
@@ -269,17 +189,17 @@
         }
     }
 
-    public void movflt(Register dst, Address src) {
+    public final void movflt(Register dst, AMD64Address src) {
         assert dst.isFpu();
         movss(dst, src);
     }
 
-    public void movflt(Address dst, Register src) {
+    public final void movflt(AMD64Address dst, Register src) {
         assert src.isFpu();
         movss(dst, src);
     }
 
-    public void movdbl(Register dst, Register src) {
+    public final void movdbl(Register dst, Register src) {
         assert dst.isFpu() && src.isFpu();
         if (UseXmmRegToRegMoveAll) {
             movapd(dst, src);
@@ -288,7 +208,7 @@
         }
     }
 
-    public void movdbl(Register dst, Address src) {
+    public final void movdbl(Register dst, AMD64Address src) {
         assert dst.isFpu();
         if (UseXmmLoadAndClearUpper) {
             movsd(dst, src);
@@ -297,29 +217,20 @@
         }
     }
 
-    public void movdbl(Address dst, Register src) {
-        assert src.isFpu();
-        movsd(dst, src);
-    }
-
     /**
      * Non-atomic write of a 64-bit constant to memory. Do not use if the address might be a
      * volatile field!
      */
-    public 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));
     }
 
-    public void xchgptr(Register src1, Register src2) {
-        xchgq(src1, src2);
-    }
-
-    public void flog(Register dest, Register value, boolean base10) {
+    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 {
@@ -329,40 +240,39 @@
         movsd(tmp, value);
         fld(tmp);
         fyl2x();
-        fstp(tmp);
-        movsd(dest, tmp);
-        addq(AMD64.rsp, 8);
+        trigEpilogue(dest, tmp);
     }
 
-    public void fsin(Register dest, Register value) {
-        ftrig(dest, value, 's');
+    public final void fsin(Register dest, Register value) {
+        AMD64Address tmp = trigPrologue(value);
+        fsin();
+        trigEpilogue(dest, tmp);
     }
 
-    public void fcos(Register dest, Register value) {
-        ftrig(dest, value, 'c');
+    public final void fcos(Register dest, Register value) {
+        AMD64Address tmp = trigPrologue(value);
+        fcos();
+        trigEpilogue(dest, tmp);
     }
 
-    public void ftan(Register dest, Register value) {
-        ftrig(dest, value, 't');
+    public final void ftan(Register dest, Register value) {
+        AMD64Address tmp = trigPrologue(value);
+        fptan();
+        fstp(0); // ftan pushes 1.0 in addition to the actual result, pop
+        trigEpilogue(dest, tmp);
     }
 
-    private void ftrig(Register dest, Register value, char op) {
-        assert dest.isFpu() && value.isFpu();
-
-        Address tmp = new Address(Kind.Double, AMD64.RSP);
+    private AMD64Address trigPrologue(Register value) {
+        assert value.isFpu();
+        AMD64Address tmp = new AMD64Address(Kind.Double, AMD64.RSP);
         subq(AMD64.rsp, 8);
         movsd(tmp, value);
         fld(tmp);
-        if (op == 's') {
-            fsin();
-        } else if (op == 'c') {
-            fcos();
-        } else if (op == 't') {
-            fptan();
-            fstp(0); // ftan pushes 1.0 in addition to the actual result, pop
-        } else {
-            throw new InternalError("should not reach here");
-        }
+        return tmp;
+    }
+
+    private void trigEpilogue(Register dest, AMD64Address tmp) {
+        assert dest.isFpu();
         fstp(tmp);
         movsd(dest, tmp);
         addq(AMD64.rsp, 8);
@@ -375,19 +285,19 @@
      * @param csl the description of the CSA
      * @param frameToCSA offset from the frame pointer to the CSA
      */
-    public void save(CalleeSaveLayout csl, int frameToCSA) {
+    public final void save(CalleeSaveLayout csl, int frameToCSA) {
         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);
         }
     }
 
-    public void restore(CalleeSaveLayout csl, int frameToCSA) {
+    public final void restore(CalleeSaveLayout csl, int frameToCSA) {
         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 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.asm.sparc/src/com/oracle/graal/asm/sparc/SPARCAssembler.java	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.asm.test/src/com/oracle/graal/asm/test/AssemblerTest.java
--- a/graal/com.oracle.graal.asm.test/src/com/oracle/graal/asm/test/AssemblerTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.asm.test/src/com/oracle/graal/asm/test/AssemblerTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -47,7 +47,7 @@
 
     protected InstalledCode assembleMethod(Method m, CodeGenTest test) {
         ResolvedJavaMethod method = codeCache.lookupJavaMethod(m);
-        RegisterConfig registerConfig = codeCache.lookupRegisterConfig(method);
+        RegisterConfig registerConfig = codeCache.lookupRegisterConfig();
         CallingConvention cc = CodeUtil.getCallingConvention(codeCache, CallingConvention.Type.JavaCallee, method, false);
 
         CompilationResult compResult = new CompilationResult();
diff -r 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.asm/src/com/oracle/graal/asm/AbstractAssembler.java	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.amd64/src/com/oracle/graal/compiler/amd64/AMD64LIRGenerator.java	Mon Feb 25 13:14:39 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;
@@ -164,9 +169,7 @@
             IndexedLocationNode indexedLoc = (IndexedLocationNode) location;
 
             index = operand(indexedLoc.index());
-            if (indexedLoc.indexScalingEnabled()) {
-                scale = target().sizeInBytes(location.getValueKind());
-            }
+            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
@@ -184,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
@@ -730,7 +733,7 @@
     public void emitDeoptimizeOnOverflow(DeoptimizationAction action, DeoptimizationReason reason, Object deoptInfo) {
         LIRFrameState info = state();
         LabelRef stubEntry = createDeoptStub(action, reason, info, deoptInfo);
-        append(new BranchOp(ConditionFlag.overflow, stubEntry, info));
+        append(new BranchOp(ConditionFlag.Overflow, stubEntry, info));
     }
 
     @Override
@@ -888,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 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/BasicPTXTest.java
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/BasicPTXTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -0,0 +1,65 @@
+/*
+ * Copyright (c) 2011, 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.compiler.ptx.test;
+
+import org.junit.*;
+
+import com.oracle.graal.api.code.*;
+import com.oracle.graal.api.runtime.*;
+import com.oracle.graal.compiler.*;
+import com.oracle.graal.compiler.ptx.*;
+import com.oracle.graal.compiler.test.*;
+import com.oracle.graal.debug.*;
+import com.oracle.graal.java.*;
+import com.oracle.graal.nodes.*;
+import com.oracle.graal.phases.*;
+import com.oracle.graal.phases.PhasePlan.*;
+import com.oracle.graal.ptx.*;
+
+/**
+ * Test class for small Java methods compiled to PTX kernels.
+ */
+public class BasicPTXTest extends GraalCompilerTest {
+
+    @Test
+    public void test1() {
+        test("test1Snippet");
+    }
+
+    @SuppressWarnings("all")
+    public static int test1Snippet(int a) {
+        return a + 1;
+    }
+
+    private void test(String snippet) {
+        StructuredGraph graph = parse(snippet);
+        Debug.dump(graph, "Graph");
+        TargetDescription target = new TargetDescription(new PTX(), true, 1, 0, 0, 0, 0, true);
+        PTXBackend ptxBackend = new PTXBackend(Graal.getRequiredCapability(CodeCacheProvider.class), target);
+        PhasePlan phasePlan = new PhasePlan();
+        GraphBuilderPhase graphBuilderPhase = new GraphBuilderPhase(runtime, GraphBuilderConfiguration.getDefault(), OptimisticOptimizations.NONE);
+        phasePlan.addPhase(PhasePosition.AFTER_PARSING, graphBuilderPhase);
+        CompilationResult result = GraalCompiler.compileMethod(runtime, ptxBackend, target, graph.method(), graph, null, phasePlan, OptimisticOptimizations.NONE);
+        System.out.println("result=" + result);
+    }
+}
diff -r 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java	Mon Feb 25 13:14:39 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,54 +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());
-            if (indexedLoc.indexScalingEnabled()) {
-                scale = target().sizeInBytes(location.getValueKind());
-            }
+            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 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.sparc/src/com/oracle/graal/compiler/sparc/SPARCLIRGenerator.java	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/BoxingEliminationTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/BoxingEliminationTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/BoxingEliminationTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -126,13 +126,13 @@
                 }
 
                 Assumptions assumptions = new Assumptions(false);
-                new InliningPhase(null, runtime(), hints, assumptions, null, phasePlan, OptimisticOptimizations.ALL).apply(graph);
-                new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
+                new InliningPhase(runtime(), hints, assumptions, null, phasePlan, OptimisticOptimizations.ALL).apply(graph);
+                new CanonicalizerPhase(runtime(), assumptions).apply(graph);
                 Debug.dump(graph, "Graph");
                 new BoxingEliminationPhase(runtime()).apply(graph);
                 Debug.dump(graph, "Graph");
                 new ExpandBoxingNodesPhase(pool).apply(graph);
-                new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
+                new CanonicalizerPhase(runtime(), assumptions).apply(graph);
                 new DeadCodeEliminationPhase().apply(graph);
                 StructuredGraph referenceGraph = parse(referenceSnippet);
                 assertEquals(referenceGraph, graph);
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/CompareCanonicalizerTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/CompareCanonicalizerTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/CompareCanonicalizerTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -35,7 +35,7 @@
 
     private StructuredGraph getCanonicalizedGraph(String name) {
         StructuredGraph graph = parse(name);
-        new CanonicalizerPhase(null, runtime(), null).apply(graph);
+        new CanonicalizerPhase(runtime(), null).apply(graph);
         return graph;
     }
 
@@ -53,7 +53,7 @@
             assertEquals(referenceGraph, graph);
         }
         Assumptions assumptions = new Assumptions(false);
-        new CanonicalizerPhase(null, runtime(), assumptions).apply(referenceGraph);
+        new CanonicalizerPhase(runtime(), assumptions).apply(referenceGraph);
         for (int i = 1; i < 4; i++) {
             StructuredGraph graph = getCanonicalizedGraph("canonicalCompare" + i);
             assertEquals(referenceGraph, graph);
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/CompiledMethodTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/CompiledMethodTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/CompiledMethodTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -56,7 +56,7 @@
     public void test1() {
         Method method = getMethod("testMethod");
         final StructuredGraph graph = parse(method);
-        new CanonicalizerPhase(null, runtime(), new Assumptions(false)).apply(graph);
+        new CanonicalizerPhase(runtime(), new Assumptions(false)).apply(graph);
         new DeadCodeEliminationPhase().apply(graph);
 
         for (Node node : graph.getNodes()) {
@@ -113,7 +113,7 @@
         ResolvedJavaMethod javaMethod = runtime.lookupJavaMethod(method);
         StructuredGraph graph = new StructuredGraph(javaMethod);
         new GraphBuilderPhase(runtime, GraphBuilderConfiguration.getSnippetDefault(), OptimisticOptimizations.NONE).apply(graph);
-        new CanonicalizerPhase(null, runtime, new Assumptions(false)).apply(graph);
+        new CanonicalizerPhase(runtime, new Assumptions(false)).apply(graph);
         new DeadCodeEliminationPhase().apply(graph);
 
         for (Node node : graph.getNodes()) {
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/DegeneratedLoopsTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/DegeneratedLoopsTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/DegeneratedLoopsTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -84,7 +84,7 @@
                 for (Invoke invoke : graph.getInvokes()) {
                     invoke.intrinsify(null);
                 }
-                new CanonicalizerPhase(null, runtime(), new Assumptions(false)).apply(graph);
+                new CanonicalizerPhase(runtime(), new Assumptions(false)).apply(graph);
                 StructuredGraph referenceGraph = parse(REFERENCE_SNIPPET);
                 Debug.dump(referenceGraph, "Graph");
                 assertEquals(referenceGraph, graph);
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/GraalCompilerTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/GraalCompilerTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/GraalCompilerTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -32,6 +32,7 @@
 import com.oracle.graal.api.meta.*;
 import com.oracle.graal.api.runtime.*;
 import com.oracle.graal.compiler.*;
+import com.oracle.graal.compiler.target.*;
 import com.oracle.graal.debug.*;
 import com.oracle.graal.graph.*;
 import com.oracle.graal.graph.Node.Verbosity;
@@ -68,11 +69,13 @@
 
     protected final GraalCodeCacheProvider runtime;
     protected final GraalCompiler graalCompiler;
+    protected final Backend backend;
 
     public GraalCompilerTest() {
         DebugEnvironment.initialize(System.out);
         this.runtime = Graal.getRequiredCapability(GraalCodeCacheProvider.class);
         this.graalCompiler = Graal.getRequiredCapability(GraalCompiler.class);
+        this.backend = Graal.getRequiredCapability(Backend.class);
     }
 
     protected void assertEquals(StructuredGraph expected, StructuredGraph graph) {
@@ -305,8 +308,6 @@
             Assert.assertTrue("expected " + expect.exception, actual.exception != null);
             Assert.assertEquals(expect.exception.getClass(), actual.exception.getClass());
         } else {
-            // System.out.println(name + "(" + Arrays.toString(args) + "): expected=" +
-            // expect.returnValue + ", actual=" + actual.returnValue);
             assertEquals(expect.returnValue, actual.returnValue);
         }
     }
@@ -342,8 +343,6 @@
             if (cached != null) {
                 if (cached.isValid()) {
                     return cached;
-                } else {
-                    // System.out.println(cached.getMethod() + " was invalidated");
                 }
 
             }
@@ -363,7 +362,7 @@
                 GraphBuilderPhase graphBuilderPhase = new GraphBuilderPhase(runtime, GraphBuilderConfiguration.getDefault(), OptimisticOptimizations.ALL);
                 phasePlan.addPhase(PhasePosition.AFTER_PARSING, graphBuilderPhase);
                 editPhasePlan(method, graph, phasePlan);
-                CompilationResult compResult = graalCompiler.compileMethod(method, graph, null, phasePlan, OptimisticOptimizations.ALL);
+                CompilationResult compResult = GraalCompiler.compileMethod(runtime(), backend, runtime().getTarget(), method, graph, null, phasePlan, OptimisticOptimizations.ALL);
                 if (printCompilation) {
                     TTY.println(String.format("@%-6d Graal %-70s %-45s %-50s | %4dms %5dB", id, "", "", "", System.currentTimeMillis() - start, compResult.getTargetCodeSize()));
                 }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/IfBoxingEliminationTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/IfBoxingEliminationTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/IfBoxingEliminationTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -89,19 +89,19 @@
                 }
 
                 Assumptions assumptions = new Assumptions(false);
-                new InliningPhase(null, runtime(), hints, assumptions, null, phasePlan, OptimisticOptimizations.ALL).apply(graph);
-                new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
+                new InliningPhase(runtime(), hints, assumptions, null, phasePlan, OptimisticOptimizations.ALL).apply(graph);
+                new CanonicalizerPhase(runtime(), assumptions).apply(graph);
                 new PhiStampPhase().apply(graph);
-                new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
+                new CanonicalizerPhase(runtime(), assumptions).apply(graph);
                 Debug.dump(graph, "Graph");
                 new BoxingEliminationPhase(runtime()).apply(graph);
                 Debug.dump(graph, "Graph");
                 new ExpandBoxingNodesPhase(pool).apply(graph);
-                new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
-                new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
+                new CanonicalizerPhase(runtime(), assumptions).apply(graph);
+                new CanonicalizerPhase(runtime(), assumptions).apply(graph);
                 new DeadCodeEliminationPhase().apply(graph);
                 StructuredGraph referenceGraph = parse(REFERENCE_SNIPPET);
-                new CanonicalizerPhase(null, runtime(), assumptions).apply(referenceGraph);
+                new CanonicalizerPhase(runtime(), assumptions).apply(referenceGraph);
                 new DeadCodeEliminationPhase().apply(referenceGraph);
 
                 assertEquals(referenceGraph, graph);
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/IfCanonicalizerTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/IfCanonicalizerTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/IfCanonicalizerTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -144,7 +144,7 @@
             n.replaceFirstInput(local, constant);
         }
         Debug.dump(graph, "Graph");
-        new CanonicalizerPhase(null, runtime(), new Assumptions(false)).apply(graph);
+        new CanonicalizerPhase(runtime(), new Assumptions(false)).apply(graph);
         StructuredGraph referenceGraph = parse(REFERENCE_SNIPPET);
         assertEquals(referenceGraph, graph);
     }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/InvokeExceptionTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/InvokeExceptionTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/InvokeExceptionTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -65,8 +65,8 @@
             hints.add(invoke);
         }
         Assumptions assumptions = new Assumptions(false);
-        new InliningPhase(null, runtime(), hints, assumptions, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL).apply(graph);
-        new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
+        new InliningPhase(runtime(), hints, assumptions, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL).apply(graph);
+        new CanonicalizerPhase(runtime(), assumptions).apply(graph);
         new DeadCodeEliminationPhase().apply(graph);
     }
 }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/InvokeHintsTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/InvokeHintsTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/InvokeHintsTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -76,8 +76,8 @@
         }
 
         Assumptions assumptions = new Assumptions(false);
-        new InliningPhase(null, runtime(), hints, assumptions, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL).apply(graph);
-        new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
+        new InliningPhase(runtime(), hints, assumptions, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL).apply(graph);
+        new CanonicalizerPhase(runtime(), assumptions).apply(graph);
         new DeadCodeEliminationPhase().apply(graph);
         StructuredGraph referenceGraph = parse(REFERENCE_SNIPPET);
         assertEquals(referenceGraph, graph);
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/LoopUnswitchTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/LoopUnswitchTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/LoopUnswitchTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -133,8 +133,8 @@
         }
 
         Assumptions assumptions = new Assumptions(false);
-        new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
-        new CanonicalizerPhase(null, runtime(), assumptions).apply(referenceGraph);
+        new CanonicalizerPhase(runtime(), assumptions).apply(graph);
+        new CanonicalizerPhase(runtime(), assumptions).apply(referenceGraph);
         Debug.scope("Test", new DebugDumpScope("Test:" + snippet), new Runnable() {
 
             @Override
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/MonitorGraphTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/MonitorGraphTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/MonitorGraphTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -93,8 +93,8 @@
             hints.add(invoke);
         }
         Assumptions assumptions = new Assumptions(false);
-        new InliningPhase(null, runtime(), hints, assumptions, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL).apply(graph);
-        new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
+        new InliningPhase(runtime(), hints, assumptions, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL).apply(graph);
+        new CanonicalizerPhase(runtime(), assumptions).apply(graph);
         new DeadCodeEliminationPhase().apply(graph);
         return graph;
     }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ReassociateAndCanonicalTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ReassociateAndCanonicalTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ReassociateAndCanonicalTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -244,9 +244,9 @@
     private <T extends Node & Node.IterableNodeType> void test(String test, String ref) {
         StructuredGraph testGraph = parse(test);
         Assumptions assumptions = new Assumptions(false);
-        new CanonicalizerPhase(null, runtime(), assumptions).apply(testGraph);
+        new CanonicalizerPhase(runtime(), assumptions).apply(testGraph);
         StructuredGraph refGraph = parse(ref);
-        new CanonicalizerPhase(null, runtime(), assumptions).apply(refGraph);
+        new CanonicalizerPhase(runtime(), assumptions).apply(refGraph);
         assertEquals(testGraph, refGraph);
     }
 }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ScalarTypeSystemTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ScalarTypeSystemTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ScalarTypeSystemTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -165,11 +165,10 @@
         // No debug scope to reduce console noise for @Test(expected = ...) tests
         StructuredGraph graph = parse(snippet);
         Debug.dump(graph, "Graph");
-        // TypeSystemTest.outputGraph(graph);
         Assumptions assumptions = new Assumptions(false);
-        new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
+        new CanonicalizerPhase(runtime(), assumptions).apply(graph);
         new ConditionalEliminationPhase(runtime()).apply(graph);
-        new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
+        new CanonicalizerPhase(runtime(), assumptions).apply(graph);
         StructuredGraph referenceGraph = parse(referenceSnippet);
         assertEquals(referenceGraph, graph);
     }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/StampCanonicalizerTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/StampCanonicalizerTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/StampCanonicalizerTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -98,7 +98,7 @@
 
     private void testZeroReturn(String methodName) {
         StructuredGraph graph = parse(methodName);
-        new CanonicalizerPhase(null, runtime(), new Assumptions(false)).apply(graph);
+        new CanonicalizerPhase(runtime(), new Assumptions(false)).apply(graph);
         new DeadCodeEliminationPhase().apply(graph);
         assertConstantReturn(graph, 0);
     }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/StraighteningTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/StraighteningTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/StraighteningTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -89,7 +89,7 @@
         // No debug scope to reduce console noise for @Test(expected = ...) tests
         StructuredGraph graph = parse(snippet);
         Debug.dump(graph, "Graph");
-        new CanonicalizerPhase(null, runtime(), new Assumptions(false)).apply(graph);
+        new CanonicalizerPhase(runtime(), new Assumptions(false)).apply(graph);
         StructuredGraph referenceGraph = parse(REFERENCE_SNIPPET);
         assertEquals(referenceGraph, graph);
     }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/TypeSystemTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/TypeSystemTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/TypeSystemTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -186,13 +186,13 @@
         StructuredGraph graph = parse(snippet);
         Debug.dump(graph, "Graph");
         Assumptions assumptions = new Assumptions(false);
-        new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
+        new CanonicalizerPhase(runtime(), assumptions).apply(graph);
         new ConditionalEliminationPhase(runtime()).apply(graph);
-        new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
+        new CanonicalizerPhase(runtime(), assumptions).apply(graph);
         // a second canonicalizer is needed to process nested MaterializeNodes
-        new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
+        new CanonicalizerPhase(runtime(), assumptions).apply(graph);
         StructuredGraph referenceGraph = parse(referenceSnippet);
-        new CanonicalizerPhase(null, runtime(), assumptions).apply(referenceGraph);
+        new CanonicalizerPhase(runtime(), assumptions).apply(referenceGraph);
         assertEquals(referenceGraph, graph);
     }
 
@@ -242,9 +242,9 @@
         StructuredGraph graph = parse(snippet);
         Debug.dump(graph, "Graph");
         Assumptions assumptions = new Assumptions(false);
-        new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
+        new CanonicalizerPhase(runtime(), assumptions).apply(graph);
         new ConditionalEliminationPhase(runtime()).apply(graph);
-        new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
+        new CanonicalizerPhase(runtime(), assumptions).apply(graph);
         Debug.dump(graph, "Graph");
         Assert.assertFalse("shouldn't have nodes of type " + clazz, graph.getNodes(clazz).iterator().hasNext());
     }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ea/EscapeAnalysisTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ea/EscapeAnalysisTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ea/EscapeAnalysisTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -205,9 +205,9 @@
             }
 
             Assumptions assumptions = new Assumptions(false);
-            new InliningPhase(null, runtime(), null, assumptions, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL).apply(graph);
+            new InliningPhase(runtime(), null, assumptions, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL).apply(graph);
             new DeadCodeEliminationPhase().apply(graph);
-            new PartialEscapeAnalysisPhase(null, runtime(), assumptions, iterativeEscapeAnalysis).apply(graph);
+            new PartialEscapeAnalysisPhase(runtime(), assumptions, iterativeEscapeAnalysis).apply(graph);
             Assert.assertEquals(1, graph.getNodes(ReturnNode.class).count());
             ReturnNode returnNode = graph.getNodes(ReturnNode.class).first();
             if (expectedConstantResult != null) {
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ea/PartialEscapeAnalysisTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ea/PartialEscapeAnalysisTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ea/PartialEscapeAnalysisTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -150,14 +150,14 @@
                 n.node().setProbability(100000);
             }
             Assumptions assumptions = new Assumptions(false);
-            new InliningPhase(null, runtime(), null, assumptions, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL).apply(graph);
+            new InliningPhase(runtime(), null, assumptions, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL).apply(graph);
             new DeadCodeEliminationPhase().apply(graph);
-            new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
-            new PartialEscapeAnalysisPhase(null, runtime(), assumptions, false).apply(graph);
+            new CanonicalizerPhase(runtime(), assumptions).apply(graph);
+            new PartialEscapeAnalysisPhase(runtime(), assumptions, false).apply(graph);
 
             new CullFrameStatesPhase().apply(graph);
             new DeadCodeEliminationPhase().apply(graph);
-            new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
+            new CanonicalizerPhase(runtime(), assumptions).apply(graph);
             return graph;
         } catch (AssertionFailedError t) {
             throw new RuntimeException(t.getMessage() + "\n" + getCanonicalGraphString(graph), t);
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/inlining/InliningTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/inlining/InliningTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/inlining/InliningTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -136,9 +136,9 @@
                 Assumptions assumptions = new Assumptions(true);
                 new ComputeProbabilityPhase().apply(graph);
                 Debug.dump(graph, "Graph");
-                new InliningPhase(null, runtime(), null, assumptions, null, phasePlan, OptimisticOptimizations.ALL).apply(graph);
+                new InliningPhase(runtime(), null, assumptions, null, phasePlan, OptimisticOptimizations.ALL).apply(graph);
                 Debug.dump(graph, "Graph");
-                new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
+                new CanonicalizerPhase(runtime(), assumptions).apply(graph);
                 new DeadCodeEliminationPhase().apply(graph);
                 return graph;
             }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/GraalCompiler.java
--- a/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/GraalCompiler.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/GraalCompiler.java	Mon Feb 25 13:14:39 2013 +0100
@@ -47,50 +47,30 @@
 
 public class GraalCompiler {
 
-    /**
-     * The target that this compiler has been configured for.
-     */
-    public final TargetDescription target;
-
-    /**
-     * The runtime that this compiler has been configured for.
-     */
-    public final GraalCodeCacheProvider runtime;
-
-    /**
-     * The backend that this compiler has been configured for.
-     */
-    public final Backend backend;
-
-    public GraalCompiler(GraalCodeCacheProvider runtime, TargetDescription target, Backend backend) {
-        this.runtime = runtime;
-        this.target = target;
-        this.backend = backend;
-    }
-
-    public CompilationResult compileMethod(final ResolvedJavaMethod method, final StructuredGraph graph, final GraphCache cache, final PhasePlan plan, final OptimisticOptimizations optimisticOpts) {
+    public static CompilationResult compileMethod(final GraalCodeCacheProvider runtime, final Backend backend, final TargetDescription target, final ResolvedJavaMethod method,
+                    final StructuredGraph graph, final GraphCache cache, final PhasePlan plan, final OptimisticOptimizations optimisticOpts) {
         assert (method.getModifiers() & Modifier.NATIVE) == 0 : "compiling native methods is not supported";
 
-        return Debug.scope("GraalCompiler", new Object[]{graph, method, this}, new Callable<CompilationResult>() {
+        return Debug.scope("GraalCompiler", new Object[]{graph, method}, new Callable<CompilationResult>() {
 
             public CompilationResult call() {
                 final Assumptions assumptions = new Assumptions(GraalOptions.OptAssumptions);
                 final LIR lir = Debug.scope("FrontEnd", new Callable<LIR>() {
 
                     public LIR call() {
-                        return emitHIR(graph, assumptions, cache, plan, optimisticOpts);
+                        return emitHIR(runtime, target, graph, assumptions, cache, plan, optimisticOpts);
                     }
                 });
                 final FrameMap frameMap = Debug.scope("BackEnd", lir, new Callable<FrameMap>() {
 
                     public FrameMap call() {
-                        return emitLIR(lir, graph, method);
+                        return emitLIR(backend, target, lir, graph, method);
                     }
                 });
                 return Debug.scope("CodeGen", frameMap, new Callable<CompilationResult>() {
 
                     public CompilationResult call() {
-                        return emitCode(getLeafGraphIdArray(graph), assumptions, method, lir, frameMap);
+                        return emitCode(backend, getLeafGraphIdArray(graph), assumptions, method, lir, frameMap);
                     }
 
                 });
@@ -110,8 +90,13 @@
 
     /**
      * Builds the graph, optimizes it.
+     * 
+     * @param runtime
+     * 
+     * @param target
      */
-    public LIR emitHIR(StructuredGraph graph, Assumptions assumptions, GraphCache cache, PhasePlan plan, OptimisticOptimizations optimisticOpts) {
+    public static LIR emitHIR(GraalCodeCacheProvider runtime, TargetDescription target, StructuredGraph graph, Assumptions assumptions, GraphCache cache, PhasePlan plan,
+                    OptimisticOptimizations optimisticOpts) {
 
         if (graph.start().next() == null) {
             plan.runPhases(PhasePosition.AFTER_PARSING, graph);
@@ -125,16 +110,16 @@
         }
 
         if (GraalOptions.OptCanonicalizer) {
-            new CanonicalizerPhase(target, runtime, assumptions).apply(graph);
+            new CanonicalizerPhase(runtime, assumptions).apply(graph);
         }
 
         if (GraalOptions.Inline && !plan.isPhaseDisabled(InliningPhase.class)) {
-            new InliningPhase(target, runtime, null, assumptions, cache, plan, optimisticOpts).apply(graph);
+            new InliningPhase(runtime, null, assumptions, cache, plan, optimisticOpts).apply(graph);
             new DeadCodeEliminationPhase().apply(graph);
 
             if (GraalOptions.CheckCastElimination && GraalOptions.OptCanonicalizer) {
-                new CanonicalizerPhase(target, runtime, assumptions).apply(graph);
-                new IterativeConditionalEliminationPhase(target, runtime, assumptions).apply(graph);
+                new CanonicalizerPhase(runtime, assumptions).apply(graph);
+                new IterativeConditionalEliminationPhase(runtime, assumptions).apply(graph);
             }
         }
 
@@ -145,19 +130,19 @@
         if (GraalOptions.FullUnroll) {
             new LoopFullUnrollPhase(runtime, assumptions).apply(graph);
             if (GraalOptions.OptCanonicalizer) {
-                new CanonicalizerPhase(target, runtime, assumptions).apply(graph);
+                new CanonicalizerPhase(runtime, assumptions).apply(graph);
             }
         }
 
         if (GraalOptions.OptTailDuplication) {
             new TailDuplicationPhase().apply(graph);
             if (GraalOptions.OptCanonicalizer) {
-                new CanonicalizerPhase(target, runtime, assumptions).apply(graph);
+                new CanonicalizerPhase(runtime, assumptions).apply(graph);
             }
         }
 
         if (GraalOptions.PartialEscapeAnalysis && !plan.isPhaseDisabled(PartialEscapeAnalysisPhase.class)) {
-            new PartialEscapeAnalysisPhase(target, runtime, assumptions, true).apply(graph);
+            new PartialEscapeAnalysisPhase(runtime, assumptions, true).apply(graph);
         }
 
         new LockEliminationPhase().apply(graph);
@@ -169,7 +154,7 @@
         new RemoveValueProxyPhase().apply(graph);
 
         if (GraalOptions.OptCanonicalizer) {
-            new CanonicalizerPhase(target, runtime, assumptions).apply(graph);
+            new CanonicalizerPhase(runtime, assumptions).apply(graph);
         }
 
         new LoweringPhase(target, runtime, assumptions).apply(graph);
@@ -181,7 +166,7 @@
         if (GraalOptions.OptFloatingReads) {
             int mark = graph.getMark();
             new FloatingReadPhase().apply(graph);
-            new CanonicalizerPhase(target, runtime, assumptions, mark, null).apply(graph);
+            new CanonicalizerPhase(runtime, assumptions, mark, null).apply(graph);
             if (GraalOptions.OptReadElimination) {
                 new ReadEliminationPhase().apply(graph);
             }
@@ -189,7 +174,7 @@
         new RemoveValueProxyPhase().apply(graph);
 
         if (GraalOptions.OptCanonicalizer) {
-            new CanonicalizerPhase(target, runtime, assumptions).apply(graph);
+            new CanonicalizerPhase(runtime, assumptions).apply(graph);
         }
 
         if (GraalOptions.OptEliminatePartiallyRedundantGuards) {
@@ -197,7 +182,7 @@
         }
 
         if (GraalOptions.CheckCastElimination && GraalOptions.OptCanonicalizer) {
-            new IterativeConditionalEliminationPhase(target, runtime, assumptions).apply(graph);
+            new IterativeConditionalEliminationPhase(runtime, assumptions).apply(graph);
         }
 
         if (GraalOptions.OptEliminatePartiallyRedundantGuards) {
@@ -240,8 +225,8 @@
 
     }
 
-    public FrameMap emitLIR(final LIR lir, StructuredGraph graph, final ResolvedJavaMethod method) {
-        final FrameMap frameMap = backend.newFrameMap(runtime.lookupRegisterConfig(method));
+    public static FrameMap emitLIR(Backend backend, final TargetDescription target, final LIR lir, StructuredGraph graph, final ResolvedJavaMethod method) {
+        final FrameMap frameMap = backend.newFrameMap();
         final LIRGenerator lirGenerator = backend.newLIRGenerator(graph, frameMap, method, lir);
 
         Debug.scope("LIRGen", lirGenerator, new Runnable() {
@@ -275,7 +260,7 @@
         return frameMap;
     }
 
-    public CompilationResult emitCode(long[] leafGraphIds, Assumptions assumptions, ResolvedJavaMethod method, LIR lir, FrameMap frameMap) {
+    public static CompilationResult emitCode(Backend backend, long[] leafGraphIds, Assumptions assumptions, ResolvedJavaMethod method, LIR lir, FrameMap frameMap) {
         TargetMethodAssembler tasm = backend.newAssembler(frameMap, lir);
         backend.emitCode(tasm, method, lir);
         CompilationResult result = tasm.finishTargetMethod(method, false);
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/target/Backend.java
--- a/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/target/Backend.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/target/Backend.java	Mon Feb 25 13:14:39 2013 +0100
@@ -46,8 +46,8 @@
         return runtime;
     }
 
-    public FrameMap newFrameMap(RegisterConfig registerConfig) {
-        return new FrameMap(runtime, target, registerConfig);
+    public FrameMap newFrameMap() {
+        return new FrameMap(runtime, target, runtime.lookupRegisterConfig());
     }
 
     public abstract LIRGenerator newLIRGenerator(StructuredGraph graph, FrameMap frameMap, ResolvedJavaMethod method, LIR lir);
diff -r 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotBackend.java	Mon Feb 25 13:14:39 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,10 +294,10 @@
             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);
+            asm.jcc(ConditionFlag.NotEqual, unverifiedStub);
         }
 
         asm.align(config.codeEntryAlignment);
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotGraalRuntime.java
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotGraalRuntime.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotGraalRuntime.java	Mon Feb 25 13:14:39 2013 +0100
@@ -51,7 +51,7 @@
         final int stackFrameAlignment = 16;
         final int stackBias = 0;
         final int implicitNullCheckLimit = 4096;
-        return new TargetDescription(new AMD64(), true, stackFrameAlignment, stackBias, implicitNullCheckLimit, config.vmPageSize, wordSize, true, true);
+        return new TargetDescription(new AMD64(), true, stackFrameAlignment, stackBias, implicitNullCheckLimit, config.vmPageSize, wordSize, true);
     }
 
     @Override
diff -r 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64SafepointOp.java	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/CompilationTask.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/CompilationTask.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/CompilationTask.java	Mon Feb 25 13:14:39 2013 +0100
@@ -29,6 +29,7 @@
 
 import com.oracle.graal.api.code.*;
 import com.oracle.graal.api.meta.*;
+import com.oracle.graal.compiler.*;
 import com.oracle.graal.debug.*;
 import com.oracle.graal.debug.internal.*;
 import com.oracle.graal.graph.*;
@@ -143,10 +144,9 @@
                         } else {
                             // Compiling an intrinsic graph - must clone the graph
                             graph = graph.copy();
-                            // System.out.println("compiling intrinsic " + method);
                         }
                         InlinedBytecodes.add(method.getCodeSize());
-                        return graalRuntime.getCompiler().compileMethod(method, graph, graalRuntime.getCache(), plan, optimisticOpts);
+                        return GraalCompiler.compileMethod(graalRuntime.getRuntime(), graalRuntime.getBackend(), graalRuntime.getTarget(), method, graph, graalRuntime.getCache(), plan, optimisticOpts);
                     }
                 });
             } finally {
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/HotSpotGraalRuntime.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/HotSpotGraalRuntime.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/HotSpotGraalRuntime.java	Mon Feb 25 13:14:39 2013 +0100
@@ -31,6 +31,7 @@
 import com.oracle.graal.api.meta.*;
 import com.oracle.graal.api.runtime.*;
 import com.oracle.graal.compiler.*;
+import com.oracle.graal.compiler.target.*;
 import com.oracle.graal.hotspot.bridge.*;
 import com.oracle.graal.hotspot.logging.*;
 import com.oracle.graal.hotspot.meta.*;
@@ -116,6 +117,7 @@
     private volatile HotSpotGraphCache cache;
 
     protected final HotSpotVMConfig config;
+    private final HotSpotBackend backend;
 
     protected HotSpotGraalRuntime() {
         CompilerToVM toVM = new CompilerToVMImpl();
@@ -139,9 +141,9 @@
 
         runtime = createRuntime();
 
-        HotSpotBackend backend = createBackend();
+        backend = createBackend();
         GraalOptions.StackShadowPages = config.stackShadowPages;
-        compiler = new GraalCompiler(getRuntime(), getTarget(), backend);
+        compiler = new GraalCompiler();
         if (GraalOptions.CacheGraphs) {
             cache = new HotSpotGraphCache();
         }
@@ -262,9 +264,19 @@
         if (clazz == DisassemblerProvider.class || clazz == BytecodeDisassemblerProvider.class) {
             return (T) getRuntime();
         }
+        if (clazz == HotSpotRuntime.class) {
+            return (T) runtime;
+        }
         if (clazz == GraalCompiler.class) {
             return (T) getCompiler();
         }
+        if (clazz == Backend.class) {
+            return (T) getBackend();
+        }
         return null;
     }
+
+    public HotSpotBackend getBackend() {
+        return backend;
+    }
 }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/VMToCompilerImpl.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/VMToCompilerImpl.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/VMToCompilerImpl.java	Mon Feb 25 13:14:39 2013 +0100
@@ -36,7 +36,6 @@
 
 import com.oracle.graal.api.code.*;
 import com.oracle.graal.api.meta.*;
-import com.oracle.graal.compiler.*;
 import com.oracle.graal.debug.*;
 import com.oracle.graal.debug.internal.*;
 import com.oracle.graal.hotspot.*;
@@ -143,10 +142,9 @@
         }
 
         // Install intrinsics.
-        GraalCompiler compiler = graalRuntime.getCompiler();
-        final HotSpotRuntime runtime = (HotSpotRuntime) compiler.runtime;
+        final HotSpotRuntime runtime = graalRuntime.getCapability(HotSpotRuntime.class);
         if (GraalOptions.Intrinsify) {
-            Debug.scope("InstallSnippets", new Object[]{new DebugDumpScope("InstallSnippets"), compiler}, new Runnable() {
+            Debug.scope("InstallSnippets", new Object[]{new DebugDumpScope("InstallSnippets")}, new Runnable() {
 
                 @Override
                 public void run() {
@@ -155,7 +153,7 @@
                     Assumptions assumptions = new Assumptions(false);
                     SnippetInstaller installer = new HotSpotSnippetInstaller(runtime, assumptions, runtime.getGraalRuntime().getTarget());
                     GraalIntrinsics.installIntrinsics(installer);
-                    runtime.installSnippets(installer, assumptions);
+                    runtime.installSnippets(graalRuntime.getBackend(), installer, assumptions);
                 }
             });
 
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedObjectType.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedObjectType.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedObjectType.java	Mon Feb 25 13:14:39 2013 +0100
@@ -500,14 +500,24 @@
     }
 
     @Override
-    public String getClassFilePath() {
+    public URL getClassFilePath() {
         Class<?> cls = mirror();
-        String name = cls.getName();
-        int dot = name.lastIndexOf('.');
-        if (dot != -1) {
-            name = name.substring(dot + 1);
-        }
-        URL classFilePath = cls.getResource(name + ".class");
-        return classFilePath == null ? null : classFilePath.getPath();
+        return cls.getResource(MetaUtil.getSimpleName(cls, true).replace('.', '$') + ".class");
+    }
+
+    @Override
+    public boolean isLocal() {
+        return mirror().isLocalClass();
+    }
+
+    @Override
+    public boolean isMember() {
+        return mirror().isMemberClass();
+    }
+
+    @Override
+    public ResolvedJavaType getEnclosingType() {
+        final Class<?> encl = mirror().getEnclosingClass();
+        return encl == null ? null : fromClass(encl);
     }
 }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedPrimitiveType.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedPrimitiveType.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedPrimitiveType.java	Mon Feb 25 13:14:39 2013 +0100
@@ -24,6 +24,7 @@
 
 import java.lang.annotation.*;
 import java.lang.reflect.*;
+import java.net.*;
 
 import com.oracle.graal.api.meta.*;
 import com.oracle.graal.graph.*;
@@ -191,7 +192,22 @@
     }
 
     @Override
-    public String getClassFilePath() {
+    public URL getClassFilePath() {
+        return null;
+    }
+
+    @Override
+    public boolean isLocal() {
+        return false;
+    }
+
+    @Override
+    public boolean isMember() {
+        return false;
+    }
+
+    @Override
+    public ResolvedJavaType getEnclosingType() {
         return null;
     }
 }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotRuntime.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotRuntime.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotRuntime.java	Mon Feb 25 13:14:39 2013 +0100
@@ -51,6 +51,7 @@
 import com.oracle.graal.api.code.Register.RegisterFlag;
 import com.oracle.graal.api.code.RuntimeCallTarget.Descriptor;
 import com.oracle.graal.api.meta.*;
+import com.oracle.graal.compiler.target.*;
 import com.oracle.graal.graph.*;
 import com.oracle.graal.hotspot.*;
 import com.oracle.graal.hotspot.bridge.*;
@@ -328,7 +329,7 @@
 
     protected abstract RegisterConfig createRegisterConfig(boolean globalStubConfig);
 
-    public void installSnippets(SnippetInstaller installer, Assumptions assumptions) {
+    public void installSnippets(Backend backend, SnippetInstaller installer, Assumptions assumptions) {
         if (GraalOptions.IntrinsifyObjectMethods) {
             installer.installSubstitutions(ObjectSubstitutions.class);
         }
@@ -370,8 +371,8 @@
 
         newInstanceStub = new NewInstanceStub(this, assumptions, graalRuntime.getTarget());
         newArrayStub = new NewArrayStub(this, assumptions, graalRuntime.getTarget());
-        newInstanceStub.install(graalRuntime.getCompiler());
-        newArrayStub.install(graalRuntime.getCompiler());
+        newInstanceStub.install(backend, graalRuntime.getCompiler());
+        newArrayStub.install(backend, graalRuntime.getCompiler());
     }
 
     public HotSpotGraalRuntime getGraalRuntime() {
@@ -504,7 +505,7 @@
     }
 
     @Override
-    public RegisterConfig lookupRegisterConfig(ResolvedJavaMethod method) {
+    public RegisterConfig lookupRegisterConfig() {
         return regConfig;
     }
 
@@ -627,7 +628,7 @@
                     graph.addAfterFixed(cas, writeBarrier);
                 } else {
                     // This may be an array store so use an array write barrier
-                    LocationNode location = IndexedLocationNode.create(LocationNode.ANY_LOCATION, cas.expected().kind(), cas.displacement(), cas.offset(), graph, false);
+                    LocationNode location = IndexedLocationNode.create(LocationNode.ANY_LOCATION, cas.expected().kind(), cas.displacement(), cas.offset(), graph, 1);
                     graph.addAfterFixed(cas, graph.add(new ArrayWriteBarrier(cas.object(), location)));
                 }
             }
@@ -679,7 +680,7 @@
         } else if (n instanceof UnsafeLoadNode) {
             UnsafeLoadNode load = (UnsafeLoadNode) n;
             assert load.kind() != Kind.Illegal;
-            IndexedLocationNode location = IndexedLocationNode.create(LocationNode.ANY_LOCATION, load.accessKind(), load.displacement(), load.offset(), graph, false);
+            IndexedLocationNode location = IndexedLocationNode.create(LocationNode.ANY_LOCATION, load.accessKind(), load.displacement(), load.offset(), graph, 1);
             ReadNode memoryRead = graph.add(new ReadNode(load.object(), location, load.stamp()));
             // An unsafe read must not floating outside its block as may float above an explicit
             // null check on its object.
@@ -687,7 +688,7 @@
             graph.replaceFixedWithFixed(load, memoryRead);
         } else if (n instanceof UnsafeStoreNode) {
             UnsafeStoreNode store = (UnsafeStoreNode) n;
-            IndexedLocationNode location = IndexedLocationNode.create(LocationNode.ANY_LOCATION, store.accessKind(), store.displacement(), store.offset(), graph, false);
+            IndexedLocationNode location = IndexedLocationNode.create(LocationNode.ANY_LOCATION, store.accessKind(), store.displacement(), store.offset(), graph, 1);
             ValueNode object = store.object();
             WriteNode write = graph.add(new WriteNode(object, store.value(), location));
             write.setStateAfter(store.stateAfter());
@@ -751,8 +752,9 @@
         }
     }
 
-    private static IndexedLocationNode createArrayLocation(Graph graph, Kind elementKind, ValueNode index) {
-        return IndexedLocationNode.create(LocationNode.getArrayLocation(elementKind), elementKind, getArrayBaseOffset(elementKind), index, graph, true);
+    private IndexedLocationNode createArrayLocation(Graph graph, Kind elementKind, ValueNode index) {
+        int scale = this.graalRuntime.getTarget().sizeInBytes(elementKind);
+        return IndexedLocationNode.create(LocationNode.getArrayLocation(elementKind), elementKind, getArrayBaseOffset(elementKind), index, graph, scale);
     }
 
     private SafeReadNode safeReadArrayLength(ValueNode array) {
diff -r 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/CurrentThread.java	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/TailcallNode.java	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/WriteBarrier.java	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/snippets/ArrayCopyNode.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/snippets/ArrayCopyNode.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/snippets/ArrayCopyNode.java	Mon Feb 25 13:14:39 2013 +0100
@@ -83,9 +83,9 @@
         }
         // the canonicalization before loop unrolling is needed to propagate the length into
         // additions, etc.
-        new CanonicalizerPhase(tool.getTarget(), tool.getRuntime(), tool.assumptions()).apply(snippetGraph);
+        new CanonicalizerPhase(tool.getRuntime(), tool.assumptions()).apply(snippetGraph);
         new LoopFullUnrollPhase(tool.getRuntime(), tool.assumptions()).apply(snippetGraph);
-        new CanonicalizerPhase(tool.getTarget(), tool.getRuntime(), tool.assumptions()).apply(snippetGraph);
+        new CanonicalizerPhase(tool.getRuntime(), tool.assumptions()).apply(snippetGraph);
     }
 
     @Override
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/stubs/Stub.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/stubs/Stub.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/stubs/Stub.java	Mon Feb 25 13:14:39 2013 +0100
@@ -31,6 +31,7 @@
 import com.oracle.graal.api.code.RuntimeCallTarget.Descriptor;
 import com.oracle.graal.api.meta.*;
 import com.oracle.graal.compiler.*;
+import com.oracle.graal.compiler.target.*;
 import com.oracle.graal.debug.*;
 import com.oracle.graal.graph.*;
 import com.oracle.graal.hotspot.*;
@@ -70,7 +71,7 @@
 
     /**
      * Creates a new stub container. The new stub still needs to be
-     * {@linkplain #install(GraalCompiler) installed}.
+     * {@linkplain #install(Backend, GraalCompiler) installed}.
      * 
      * @param descriptor linkage details for a call to the stub
      */
@@ -106,7 +107,7 @@
      * Compiles the code for this stub, installs it and initializes the address used for calls to
      * it.
      */
-    public void install(GraalCompiler compiler) {
+    public void install(Backend backend, GraalCompiler compiler) {
         StructuredGraph graph = (StructuredGraph) stubMethod.getCompilerStorage().get(Graph.class);
 
         Key key = new Key(stubMethod);
@@ -117,7 +118,7 @@
         PhasePlan phasePlan = new PhasePlan();
         GraphBuilderPhase graphBuilderPhase = new GraphBuilderPhase(runtime, GraphBuilderConfiguration.getDefault(), OptimisticOptimizations.ALL);
         phasePlan.addPhase(PhasePosition.AFTER_PARSING, graphBuilderPhase);
-        final CompilationResult compResult = compiler.compileMethod(stubMethod, graph, null, phasePlan, OptimisticOptimizations.ALL);
+        final CompilationResult compResult = GraalCompiler.compileMethod(runtime(), backend, runtime().getTarget(), stubMethod, graph, null, phasePlan, OptimisticOptimizations.ALL);
 
         final CodeInfo[] info = new CodeInfo[1];
         stubCode = Debug.scope("CodeInstall", new Object[]{compiler, stubMethod}, new Callable<InstalledCode>() {
diff -r 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Arithmetic.java	Mon Feb 25 13:14:39 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,10 +455,10 @@
         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);
+        masm.jcc(ConditionFlag.Equal, slowPath.start);
         masm.bind(slowPath.continuation);
     }
 
@@ -477,13 +477,13 @@
         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();
-            masm.jcc(ConditionFlag.parity, nan);
-            masm.jcc(ConditionFlag.below, continuation);
+            masm.jcc(ConditionFlag.Parity, nan);
+            masm.jcc(ConditionFlag.Below, continuation);
 
             // input is > 0 -> return maxInt
             // result register already contains 0x80000000, so subtracting 1 gives 0x7fffffff
diff -r 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64BitManipulationOp.java	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Compare.java	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64ControlFlow.java	Mon Feb 25 13:14:39 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.*;
@@ -162,13 +162,13 @@
                     long lc = keyConstants[i].asLong();
                     assert NumUtil.isInt(lc);
                     masm.cmpl(intKey, (int) lc);
-                    masm.jcc(ConditionFlag.equal, keyTargets[i].label());
+                    masm.jcc(ConditionFlag.Equal, keyTargets[i].label());
                 }
             } 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.jcc(ConditionFlag.equal, keyTargets[i].label());
+                    masm.cmpq(longKey, (AMD64Address) tasm.asLongConstRef(keyConstants[i]));
+                    masm.jcc(ConditionFlag.Equal, keyTargets[i].label());
                 }
             } else if (key.getKind() == Kind.Object) {
                 Register intKey = asObjectReg(key);
@@ -176,7 +176,7 @@
                 for (int i = 0; i < keyConstants.length; i++) {
                     AMD64Move.move(tasm, masm, temp.asValue(Kind.Object), keyConstants[i]);
                     masm.cmpptr(intKey, temp);
-                    masm.jcc(ConditionFlag.equal, keyTargets[i].label());
+                    masm.jcc(ConditionFlag.Equal, keyTargets[i].label());
                 }
             } else {
                 throw new GraalInternalError("sequential switch only supported for int, long and object");
@@ -226,15 +226,15 @@
                 int highKey = highKeys[i];
                 if (lowKey == highKey) {
                     masm.cmpl(asIntReg(key), lowKey);
-                    masm.jcc(ConditionFlag.equal, keyTargets[i].label());
+                    masm.jcc(ConditionFlag.Equal, keyTargets[i].label());
                     skipLowCheck = false;
                 } else {
                     if (!skipLowCheck || (prevHighKey + 1) != lowKey) {
                         masm.cmpl(asIntReg(key), lowKey);
-                        masm.jcc(ConditionFlag.less, actualDefaultTarget);
+                        masm.jcc(ConditionFlag.Less, actualDefaultTarget);
                     }
                     masm.cmpl(asIntReg(key), highKey);
-                    masm.jcc(ConditionFlag.lessEqual, keyTargets[i].label());
+                    masm.jcc(ConditionFlag.LessEqual, keyTargets[i].label());
                     skipLowCheck = true;
                 }
                 prevHighKey = highKey;
@@ -333,16 +333,16 @@
 
         // Jump to default target if index is not within the jump table
         if (defaultTarget != null) {
-            masm.jcc(ConditionFlag.above, defaultTarget.label());
+            masm.jcc(ConditionFlag.Above, defaultTarget.label());
         }
 
         // 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
@@ -380,9 +380,9 @@
     private static void floatJcc(AMD64MacroAssembler masm, ConditionFlag condition, boolean unorderedIsTrue, Label label) {
         Label endLabel = new Label();
         if (unorderedIsTrue && !trueOnUnordered(condition)) {
-            masm.jcc(ConditionFlag.parity, label);
+            masm.jcc(ConditionFlag.Parity, label);
         } else if (!unorderedIsTrue && trueOnUnordered(condition)) {
-            masm.jcc(ConditionFlag.parity, endLabel);
+            masm.jccb(ConditionFlag.Parity, endLabel);
         }
         masm.jcc(condition, label);
         masm.bind(endLabel);
@@ -397,9 +397,9 @@
 
         if (isFloat) {
             if (unorderedIsTrue && !trueOnUnordered(condition)) {
-                cmove(tasm, masm, result, ConditionFlag.parity, trueValue);
+                cmove(tasm, masm, result, ConditionFlag.Parity, trueValue);
             } else if (!unorderedIsTrue && trueOnUnordered(condition)) {
-                cmove(tasm, masm, result, ConditionFlag.parity, falseValue);
+                cmove(tasm, masm, result, ConditionFlag.Parity, falseValue);
             }
         }
     }
@@ -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();
             }
         }
@@ -423,45 +424,45 @@
 
     private static ConditionFlag intCond(Condition cond) {
         switch (cond) {
-            case EQ: return ConditionFlag.equal;
-            case NE: return ConditionFlag.notEqual;
-            case LT: return ConditionFlag.less;
-            case LE: return ConditionFlag.lessEqual;
-            case GE: return ConditionFlag.greaterEqual;
-            case GT: return ConditionFlag.greater;
-            case BE: return ConditionFlag.belowEqual;
-            case AE: return ConditionFlag.aboveEqual;
-            case AT: return ConditionFlag.above;
-            case BT: return ConditionFlag.below;
+            case EQ: return ConditionFlag.Equal;
+            case NE: return ConditionFlag.NotEqual;
+            case LT: return ConditionFlag.Less;
+            case LE: return ConditionFlag.LessEqual;
+            case GE: return ConditionFlag.GreaterEqual;
+            case GT: return ConditionFlag.Greater;
+            case BE: return ConditionFlag.BelowEqual;
+            case AE: return ConditionFlag.AboveEqual;
+            case AT: return ConditionFlag.Above;
+            case BT: return ConditionFlag.Below;
             default: throw GraalInternalError.shouldNotReachHere();
         }
     }
 
     private static ConditionFlag floatCond(Condition cond) {
         switch (cond) {
-            case EQ: return ConditionFlag.equal;
-            case NE: return ConditionFlag.notEqual;
-            case LT: return ConditionFlag.below;
-            case LE: return ConditionFlag.belowEqual;
-            case GE: return ConditionFlag.aboveEqual;
-            case GT: return ConditionFlag.above;
+            case EQ: return ConditionFlag.Equal;
+            case NE: return ConditionFlag.NotEqual;
+            case LT: return ConditionFlag.Below;
+            case LE: return ConditionFlag.BelowEqual;
+            case GE: return ConditionFlag.AboveEqual;
+            case GT: return ConditionFlag.Above;
             default: throw GraalInternalError.shouldNotReachHere();
         }
     }
 
     private static boolean trueOnUnordered(ConditionFlag condition) {
         switch(condition) {
-            case aboveEqual:
-            case notEqual:
-            case above:
-            case less:
-            case overflow:
+            case AboveEqual:
+            case NotEqual:
+            case Above:
+            case Less:
+            case Overflow:
                 return false;
-            case equal:
-            case belowEqual:
-            case below:
-            case greaterEqual:
-            case noOverflow:
+            case Equal:
+            case BelowEqual:
+            case Below:
+            case GreaterEqual:
+            case NoOverflow:
                 return true;
             default:
                 throw GraalInternalError.shouldNotReachHere();
diff -r 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Move.java	Mon Feb 25 13:14:39 2013 +0100
@@ -38,7 +38,6 @@
 import com.oracle.graal.lir.StandardOp.MoveOp;
 import com.oracle.graal.lir.asm.*;
 
-// @formatter:off
 public class AMD64Move {
 
     @Opcode("MOVE")
@@ -132,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);
         }
     }
 
@@ -150,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);
         }
     }
 
@@ -166,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));
         }
     }
 
@@ -218,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);
         }
     }
 
@@ -266,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();
         }
     }
@@ -317,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:
@@ -326,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:
@@ -339,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:
@@ -349,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");
                 }
@@ -367,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);
         }
@@ -385,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);
         }
@@ -436,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 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64TestOp.java	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.lir/src/com/oracle/graal/lir/LIRInstructionClass.java	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.lir/src/com/oracle/graal/lir/asm/TargetMethodAssembler.java	Mon Feb 25 13:14:39 2013 +0100
@@ -65,7 +65,6 @@
     public final FrameContext frameContext;
 
     private List<ExceptionInfo> exceptionInfoList;
-    private int lastSafepointPos;
 
     public TargetMethodAssembler(TargetDescription target, CodeCacheProvider runtime, FrameMap frameMap, AbstractAssembler asm, FrameContext frameContext, List<Code> stubs) {
         this.target = target;
@@ -75,8 +74,6 @@
         this.asm = asm;
         this.compilationResult = new CompilationResult();
         this.frameContext = frameContext;
-        // 0 is a valid pc for safepoints in template methods
-        this.lastSafepointPos = -1;
     }
 
     public void setFrameSize(int frameSize) {
@@ -109,9 +106,6 @@
             }
         }
 
-        // Set the info on callee-saved registers
-        compilationResult.setCalleeSaveLayout(frameMap.registerConfig.getCalleeSaveLayout());
-
         Debug.metric("TargetMethods").increment();
         Debug.metric("CodeBytesEmitted").add(compilationResult.getTargetCodeSize());
         Debug.metric("SafepointsEmitted").add(compilationResult.getSafepoints().size());
@@ -135,8 +129,6 @@
     public void recordImplicitException(int pcOffset, LIRFrameState info) {
         // record an implicit exception point
         if (info != null) {
-            assert lastSafepointPos < pcOffset : lastSafepointPos + "<" + pcOffset;
-            lastSafepointPos = pcOffset;
             compilationResult.recordSafepoint(pcOffset, info.debugInfo());
             assert info.exceptionEdge == null;
         }
@@ -144,23 +136,17 @@
 
     public void recordDirectCall(int posBefore, int posAfter, InvokeTarget callTarget, LIRFrameState info) {
         DebugInfo debugInfo = info != null ? info.debugInfo() : null;
-        assert lastSafepointPos < posAfter;
-        lastSafepointPos = posAfter;
         compilationResult.recordCall(posBefore, posAfter - posBefore, callTarget, debugInfo, true);
     }
 
     public void recordIndirectCall(int posBefore, int posAfter, InvokeTarget callTarget, LIRFrameState info) {
         DebugInfo debugInfo = info != null ? info.debugInfo() : null;
-        assert lastSafepointPos < posAfter;
-        lastSafepointPos = posAfter;
         compilationResult.recordCall(posBefore, posAfter - posBefore, callTarget, debugInfo, false);
     }
 
     public void recordSafepoint(int pos, LIRFrameState info) {
         // safepoints always need debug info
         DebugInfo debugInfo = info.debugInfo();
-        assert lastSafepointPos < pos;
-        lastSafepointPos = pos;
         compilationResult.recordSafepoint(pos, debugInfo);
     }
 
@@ -169,11 +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;
-    }
-
-    public int lastSafepointPos() {
-        return lastSafepointPos;
+        return asm.getPlaceholder();
     }
 
     /**
@@ -251,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 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.loop/src/com/oracle/graal/loop/LoopTransformations.java
--- a/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/LoopTransformations.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/LoopTransformations.java	Mon Feb 25 13:14:39 2013 +0100
@@ -60,7 +60,7 @@
         while (!loopBegin.isDeleted()) {
             int mark = graph.getMark();
             peel(loop);
-            new CanonicalizerPhase(null, runtime, assumptions, mark, null).apply(graph);
+            new CanonicalizerPhase(runtime, assumptions, mark, null).apply(graph);
             if (iterations++ > UNROLL_LIMIT || graph.getNodeCount() > GraalOptions.MaximumDesiredSize * 3) {
                 throw new BailoutException("FullUnroll : Graph seems to grow out of proportion");
             }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/PhiNode.java
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/PhiNode.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/PhiNode.java	Mon Feb 25 13:14:39 2013 +0100
@@ -37,7 +37,8 @@
 
     public static enum PhiType {
         Value(null), // normal value phis
-        Guard(StampFactory.dependency()), Memory(StampFactory.dependency());
+        Guard(StampFactory.dependency()),
+        Memory(StampFactory.dependency());
 
         public final Stamp stamp;
 
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/ConvertNode.java
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/ConvertNode.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/ConvertNode.java	Mon Feb 25 13:14:39 2013 +0100
@@ -35,9 +35,26 @@
 public final class ConvertNode extends FloatingNode implements Canonicalizable, LIRLowerable {
 
     public enum Op {
-        I2L(Int, Long), L2I(Long, Int), I2B(Int, Byte), I2C(Int, Char), I2S(Int, Short), F2D(Float, Double), D2F(Double, Float), I2F(Int, Float), I2D(Int, Double), F2I(Float, Int), D2I(Double, Int), L2F(
-                        Long, Float), L2D(Long, Double), F2L(Float, Long), D2L(Double, Long), UNSIGNED_I2L(Int, Long), MOV_I2F(Int, Float), MOV_L2D(Long, Double), MOV_F2I(Float, Int), MOV_D2L(Double,
-                        Long);
+        I2L(Int, Long),
+        L2I(Long, Int),
+        I2B(Int, Byte),
+        I2C(Int, Char),
+        I2S(Int, Short),
+        F2D(Float, Double),
+        D2F(Double, Float),
+        I2F(Int, Float),
+        I2D(Int, Double),
+        F2I(Float, Int),
+        D2I(Double, Int),
+        L2F(Long, Float),
+        L2D(Long, Double),
+        F2L(Float, Long),
+        D2L(Double, Long),
+        UNSIGNED_I2L(Int, Long),
+        MOV_I2F(Int, Float),
+        MOV_L2D(Long, Double),
+        MOV_F2I(Float, Int),
+        MOV_D2L(Double, Long);
 
         public final Kind from;
         public final Kind to;
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/extended/IndexedLocationNode.java
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/extended/IndexedLocationNode.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/extended/IndexedLocationNode.java	Mon Feb 25 13:14:39 2013 +0100
@@ -28,17 +28,14 @@
 import com.oracle.graal.nodes.spi.*;
 
 /**
- * Extension of a {@linkplain LocationNode location} to include a scaled index or an additional
- * offset.
+ * Extension of a {@linkplain LocationNode location} to include a scaled index. Can represent
+ * locations in the form of [base + index * scale + disp] where base and index are nodes and scale
+ * and disp are integer constants.
  */
 public final class IndexedLocationNode extends LocationNode implements Canonicalizable {
 
-    /**
-     * An offset or index depending on whether {@link #indexScalingEnabled} is true or false
-     * respectively.
-     */
     @Input private ValueNode index;
-    private final boolean indexScalingEnabled;
+    private final int indexScaling;
 
     /**
      * Gets the index or offset of this location.
@@ -52,21 +49,20 @@
     }
 
     /**
-     * @return whether scaling of the index by the value kind's size is enabled (the default) or
-     *         disabled.
+     * @return Constant that is used to scale the index.
      */
-    public boolean indexScalingEnabled() {
-        return indexScalingEnabled;
+    public int indexScaling() {
+        return indexScaling;
     }
 
-    public static IndexedLocationNode create(Object identity, Kind kind, int displacement, ValueNode index, Graph graph, boolean indexScalingEnabled) {
-        return graph.unique(new IndexedLocationNode(identity, kind, index, displacement, indexScalingEnabled));
+    public static IndexedLocationNode create(Object identity, Kind kind, int displacement, ValueNode index, Graph graph, int indexScaling) {
+        return graph.unique(new IndexedLocationNode(identity, kind, index, displacement, indexScaling));
     }
 
-    private IndexedLocationNode(Object identity, Kind kind, ValueNode index, int displacement, boolean indexScalingEnabled) {
+    private IndexedLocationNode(Object identity, Kind kind, ValueNode index, int displacement, int indexScaling) {
         super(identity, kind, displacement);
         this.index = index;
-        this.indexScalingEnabled = indexScalingEnabled;
+        this.indexScaling = indexScaling;
     }
 
     @Override
@@ -74,12 +70,7 @@
         Constant constantIndex = index.asConstant();
         if (constantIndex != null) {
             long constantIndexLong = constantIndex.asLong();
-            if (indexScalingEnabled) {
-                if (tool.target() == null) {
-                    return this;
-                }
-                constantIndexLong *= tool.target().sizeInBytes(getValueKind());
-            }
+            constantIndexLong *= indexScaling;
             constantIndexLong += displacement();
             int constantIndexInt = (int) constantIndexLong;
             if (constantIndexLong == constantIndexInt) {
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/extended/LocationNode.java
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/extended/LocationNode.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/extended/LocationNode.java	Mon Feb 25 13:14:39 2013 +0100
@@ -30,8 +30,9 @@
 import com.oracle.graal.nodes.type.*;
 
 /**
- * A location for a memory access in terms of the kind of value accessed and the displacement (in
- * bytes) from a base object or address.
+ * A location for a memory access in terms of the kind of value accessed and how to access it. The
+ * base version can represent addresses of the form [base + disp] where base is a node and disp is a
+ * constant.
  */
 @NodeInfo(nameTemplate = "Loc {p#locationIdentity/s}")
 public class LocationNode extends FloatingNode implements LIRLowerable, ValueNumberable {
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/spi/CanonicalizerTool.java
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/spi/CanonicalizerTool.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/spi/CanonicalizerTool.java	Mon Feb 25 13:14:39 2013 +0100
@@ -28,8 +28,6 @@
 
 public interface CanonicalizerTool {
 
-    TargetDescription target();
-
     Assumptions assumptions();
 
     MetaAccessProvider runtime();
diff -r 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/spi/LIRGeneratorTool.java	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/CanonicalizerPhase.java
--- a/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/CanonicalizerPhase.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/CanonicalizerPhase.java	Mon Feb 25 13:14:39 2013 +0100
@@ -48,7 +48,6 @@
     public static final DebugMetric METRIC_GLOBAL_VALUE_NUMBERING_HITS = Debug.metric("GlobalValueNumberingHits");
 
     private final int newNodesMark;
-    private final TargetDescription target;
     private final Assumptions assumptions;
     private final MetaAccessProvider runtime;
     private final CustomCanonicalizer customCanonicalizer;
@@ -63,33 +62,31 @@
         ValueNode canonicalize(Node node);
     }
 
-    public CanonicalizerPhase(TargetDescription target, MetaAccessProvider runtime, Assumptions assumptions) {
-        this(target, runtime, assumptions, null, 0, null);
+    public CanonicalizerPhase(MetaAccessProvider runtime, Assumptions assumptions) {
+        this(runtime, assumptions, null, 0, null);
     }
 
     /**
-     * @param target
      * @param runtime
      * @param assumptions
      * @param workingSet the initial working set of nodes on which the canonicalizer works, should
      *            be an auto-grow node bitmap
      * @param customCanonicalizer
      */
-    public CanonicalizerPhase(TargetDescription target, MetaAccessProvider runtime, Assumptions assumptions, Iterable<Node> workingSet, CustomCanonicalizer customCanonicalizer) {
-        this(target, runtime, assumptions, workingSet, 0, customCanonicalizer);
+    public CanonicalizerPhase(MetaAccessProvider runtime, Assumptions assumptions, Iterable<Node> workingSet, CustomCanonicalizer customCanonicalizer) {
+        this(runtime, assumptions, workingSet, 0, customCanonicalizer);
     }
 
     /**
      * @param newNodesMark only the {@linkplain Graph#getNewNodes(int) new nodes} specified by this
      *            mark are processed otherwise all nodes in the graph are processed
      */
-    public CanonicalizerPhase(TargetDescription target, MetaAccessProvider runtime, Assumptions assumptions, int newNodesMark, CustomCanonicalizer customCanonicalizer) {
-        this(target, runtime, assumptions, null, newNodesMark, customCanonicalizer);
+    public CanonicalizerPhase(MetaAccessProvider runtime, Assumptions assumptions, int newNodesMark, CustomCanonicalizer customCanonicalizer) {
+        this(runtime, assumptions, null, newNodesMark, customCanonicalizer);
     }
 
-    public CanonicalizerPhase(TargetDescription target, MetaAccessProvider runtime, Assumptions assumptions, Iterable<Node> workingSet, int newNodesMark, CustomCanonicalizer customCanonicalizer) {
+    public CanonicalizerPhase(MetaAccessProvider runtime, Assumptions assumptions, Iterable<Node> workingSet, int newNodesMark, CustomCanonicalizer customCanonicalizer) {
         this.newNodesMark = newNodesMark;
-        this.target = target;
         this.assumptions = assumptions;
         this.runtime = runtime;
         this.customCanonicalizer = customCanonicalizer;
@@ -108,7 +105,7 @@
         if (newNodesMark > 0) {
             workList.addAll(graph.getNewNodes(newNodesMark));
         }
-        tool = new Tool(workList, runtime, target, assumptions);
+        tool = new Tool(workList, runtime, assumptions);
         processWorkSet(graph);
     }
 
@@ -309,13 +306,11 @@
 
         private final NodeWorkList nodeWorkSet;
         private final MetaAccessProvider runtime;
-        private final TargetDescription target;
         private final Assumptions assumptions;
 
-        public Tool(NodeWorkList nodeWorkSet, MetaAccessProvider runtime, TargetDescription target, Assumptions assumptions) {
+        public Tool(NodeWorkList nodeWorkSet, MetaAccessProvider runtime, Assumptions assumptions) {
             this.nodeWorkSet = nodeWorkSet;
             this.runtime = runtime;
-            this.target = target;
             this.assumptions = assumptions;
         }
 
@@ -326,15 +321,6 @@
         }
 
         /**
-         * @return the current target or {@code null} if no target is available in the current
-         *         context.
-         */
-        @Override
-        public TargetDescription target() {
-            return target;
-        }
-
-        /**
          * @return an object that can be used for recording assumptions or {@code null} if
          *         assumptions are not allowed in the current context.
          */
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/InliningPhase.java
--- a/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/InliningPhase.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/InliningPhase.java	Mon Feb 25 13:14:39 2013 +0100
@@ -47,7 +47,6 @@
      * invoke.bci, method, true);
      */
 
-    private final TargetDescription target;
     private final PhasePlan plan;
 
     private final GraalCodeCacheProvider runtime;
@@ -63,18 +62,15 @@
     private static final DebugMetric metricInliningStoppedByMaxDesiredSize = Debug.metric("InliningStoppedByMaxDesiredSize");
     private static final DebugMetric metricInliningRuns = Debug.metric("Runs");
 
-    public InliningPhase(TargetDescription target, GraalCodeCacheProvider runtime, Collection<Invoke> hints, Assumptions assumptions, GraphCache cache, PhasePlan plan,
-                    OptimisticOptimizations optimisticOpts) {
-        this(target, runtime, assumptions, cache, plan, createInliningPolicy(runtime, assumptions, optimisticOpts, hints), optimisticOpts);
+    public InliningPhase(GraalCodeCacheProvider runtime, Collection<Invoke> hints, Assumptions assumptions, GraphCache cache, PhasePlan plan, OptimisticOptimizations optimisticOpts) {
+        this(runtime, assumptions, cache, plan, createInliningPolicy(runtime, assumptions, optimisticOpts, hints), optimisticOpts);
     }
 
     public void setCustomCanonicalizer(CustomCanonicalizer customCanonicalizer) {
         this.customCanonicalizer = customCanonicalizer;
     }
 
-    public InliningPhase(TargetDescription target, GraalCodeCacheProvider runtime, Assumptions assumptions, GraphCache cache, PhasePlan plan, InliningPolicy inliningPolicy,
-                    OptimisticOptimizations optimisticOpts) {
-        this.target = target;
+    public InliningPhase(GraalCodeCacheProvider runtime, Assumptions assumptions, GraphCache cache, PhasePlan plan, InliningPolicy inliningPolicy, OptimisticOptimizations optimisticOpts) {
         this.runtime = runtime;
         this.assumptions = assumptions;
         this.cache = cache;
@@ -103,7 +99,7 @@
                         Iterable<Node> newNodes = graph.getNewNodes(mark);
                         inliningPolicy.scanInvokes(newNodes);
                         if (GraalOptions.OptCanonicalizer) {
-                            new CanonicalizerPhase(target, runtime, assumptions, invokeUsages, mark, customCanonicalizer).apply(graph);
+                            new CanonicalizerPhase(runtime, assumptions, invokeUsages, mark, customCanonicalizer).apply(graph);
                         }
                         metricInliningPerformed.increment();
                     } catch (BailoutException bailout) {
@@ -142,7 +138,7 @@
             new ComputeProbabilityPhase().apply(newGraph);
         }
         if (GraalOptions.OptCanonicalizer) {
-            new CanonicalizerPhase(target, runtime, assumptions).apply(newGraph);
+            new CanonicalizerPhase(runtime, assumptions).apply(newGraph);
         }
         if (GraalOptions.CullFrameStates) {
             new CullFrameStatesPhase().apply(newGraph);
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/InliningUtil.java
--- a/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/InliningUtil.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/InliningUtil.java	Mon Feb 25 13:14:39 2013 +0100
@@ -696,7 +696,8 @@
                 ResolvedJavaMethod targetMethod = methodCallTarget.targetMethod();
                 ResolvedJavaType leastCommonType = getLeastCommonType();
                 // check if we have a common base type that implements the interface -> in that case
-// we have a vtable entry for the interface method and can use a less expensive virtual call
+                // we have a vtable entry for the interface method and can use a less expensive
+                // virtual call
                 if (!leastCommonType.isInterface() && targetMethod.getDeclaringClass().isAssignableFrom(leastCommonType)) {
                     ResolvedJavaMethod baseClassTargetMethod = leastCommonType.resolveMethod(targetMethod);
                     if (baseClassTargetMethod != null) {
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/IterativeConditionalEliminationPhase.java
--- a/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/IterativeConditionalEliminationPhase.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/IterativeConditionalEliminationPhase.java	Mon Feb 25 13:14:39 2013 +0100
@@ -33,12 +33,10 @@
 
 public class IterativeConditionalEliminationPhase extends Phase {
 
-    private final TargetDescription target;
     private final MetaAccessProvider runtime;
     private final Assumptions assumptions;
 
-    public IterativeConditionalEliminationPhase(TargetDescription target, MetaAccessProvider runtime, Assumptions assumptions) {
-        this.target = target;
+    public IterativeConditionalEliminationPhase(MetaAccessProvider runtime, Assumptions assumptions) {
         this.runtime = runtime;
         this.assumptions = assumptions;
     }
@@ -55,7 +53,7 @@
             if (canonicalizationRoots.isEmpty()) {
                 break;
             }
-            new CanonicalizerPhase(target, runtime, assumptions, canonicalizationRoots, null).apply(graph);
+            new CanonicalizerPhase(runtime, assumptions, canonicalizationRoots, null).apply(graph);
             canonicalizationRoots.clear();
         }
     }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/LoweringPhase.java
--- a/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/LoweringPhase.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/LoweringPhase.java	Mon Feb 25 13:14:39 2013 +0100
@@ -139,7 +139,7 @@
             deferred = false;
             processBlock(schedule.getCFG().getStartBlock(), graph.createNodeBitMap(), null, schedule, processed);
             Debug.dump(graph, "Lowering iteration %d", i++);
-            new CanonicalizerPhase(null, runtime, assumptions, mark, null).apply(graph);
+            new CanonicalizerPhase(runtime, assumptions, mark, null).apply(graph);
 
             if (!deferred && !containsLowerable(graph.getNewNodes(mark))) {
                 // No new lowerable nodes - done!
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.printer/src/com/oracle/graal/printer/CFGPrinterObserver.java
--- a/graal/com.oracle.graal.printer/src/com/oracle/graal/printer/CFGPrinterObserver.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.printer/src/com/oracle/graal/printer/CFGPrinterObserver.java	Mon Feb 25 13:14:39 2013 +0100
@@ -28,7 +28,6 @@
 
 import com.oracle.graal.api.code.*;
 import com.oracle.graal.api.meta.*;
-import com.oracle.graal.compiler.*;
 import com.oracle.graal.compiler.alloc.*;
 import com.oracle.graal.compiler.gen.*;
 import com.oracle.graal.debug.*;
@@ -100,10 +99,6 @@
     private static final AtomicInteger uniqueId = new AtomicInteger();
 
     public void dumpSandboxed(Object object, String message) {
-        GraalCompiler compiler = Debug.contextLookup(GraalCompiler.class);
-        if (compiler == null) {
-            return;
-        }
 
         if (cfgPrinter == null) {
             cfgFile = new File("compilations-" + timestamp + "_" + uniqueId.incrementAndGet() + ".cfg");
@@ -120,7 +115,7 @@
             return;
         }
 
-        cfgPrinter.target = compiler.target;
+        cfgPrinter.target = Debug.contextLookup(TargetDescription.class);
         if (object instanceof LIR) {
             cfgPrinter.lir = (LIR) object;
         } else {
@@ -131,7 +126,7 @@
             cfgPrinter.cfg = cfgPrinter.lir.cfg;
         }
 
-        CodeCacheProvider runtime = compiler.runtime;
+        CodeCacheProvider runtime = Debug.contextLookup(CodeCacheProvider.class);
 
         if (object instanceof BciBlockMapping) {
             BciBlockMapping blockMap = (BciBlockMapping) object;
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.printer/src/com/oracle/graal/printer/HexCodeFile.java
--- a/graal/com.oracle.graal.printer/src/com/oracle/graal/printer/HexCodeFile.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.printer/src/com/oracle/graal/printer/HexCodeFile.java	Mon Feb 25 13:14:39 2013 +0100
@@ -30,7 +30,6 @@
 import com.oracle.graal.api.code.CompilationResult.CodeAnnotation;
 import com.oracle.graal.api.code.CompilationResult.CodeComment;
 import com.oracle.graal.api.code.CompilationResult.JumpTable;
-import com.oracle.graal.api.code.CompilationResult.LookupTable;
 
 /**
  * A HexCodeFile is a textual format for representing a chunk of machine code along with extra
@@ -122,8 +121,6 @@
 
     public final ArrayList<JumpTable> jumpTables = new ArrayList<>();
 
-    public final ArrayList<LookupTable> lookupTables = new ArrayList<>();
-
     public final String isa;
 
     public final int wordWidth;
@@ -169,10 +166,6 @@
             ps.printf("JumpTable %d %d %d %d %s%n", table.position, table.entrySize, table.low, table.high, SECTION_DELIM);
         }
 
-        for (LookupTable table : lookupTables) {
-            ps.printf("LookupTable %d %d %d %d %s%n", table.position, table.npairs, table.keySize, table.keySize, SECTION_DELIM);
-        }
-
         for (Map.Entry<Integer, List<String>> e : comments.entrySet()) {
             int pos = e.getKey();
             for (String comment : e.getValue()) {
@@ -233,9 +226,6 @@
             if (a instanceof JumpTable) {
                 JumpTable table = (JumpTable) a;
                 hcf.jumpTables.add(table);
-            } else if (a instanceof LookupTable) {
-                LookupTable table = (LookupTable) a;
-                hcf.lookupTables.add(table);
             } else if (a instanceof CodeComment) {
                 CodeComment comment = (CodeComment) a;
                 hcf.addComment(comment.position, comment.value);
@@ -422,15 +412,6 @@
                 int low = parseInt(bodyOffset + m.start(3), m.group(3));
                 int high = parseInt(bodyOffset + m.start(4), m.group(4));
                 hcf.jumpTables.add(new JumpTable(pos, low, high, entrySize));
-            } else if (header.equals("LookupTable")) {
-                checkHCF("LookupTable", headerOffset);
-                m = HexCodeFile.LOOKUP_TABLE.matcher(body);
-                check(m.matches(), bodyOffset, "LookupTable does not match pattern " + HexCodeFile.LOOKUP_TABLE);
-                int pos = parseInt(bodyOffset + m.start(1), m.group(1));
-                int npairs = parseInt(bodyOffset + m.start(2), m.group(2));
-                int keySize = parseInt(bodyOffset + m.start(3), m.group(3));
-                int offsetSize = parseInt(bodyOffset + m.start(4), m.group(4));
-                hcf.lookupTables.add(new LookupTable(pos, npairs, keySize, offsetSize));
             } else {
                 error(offset, "Unknown section header: " + header);
             }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.snippets.test/src/com/oracle/graal/snippets/IntrinsificationTest.java
--- a/graal/com.oracle.graal.snippets.test/src/com/oracle/graal/snippets/IntrinsificationTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.snippets.test/src/com/oracle/graal/snippets/IntrinsificationTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -426,9 +426,9 @@
                 Assumptions assumptions = new Assumptions(true);
                 new ComputeProbabilityPhase().apply(graph);
                 Debug.dump(graph, "Graph");
-                new InliningPhase(null, runtime(), null, assumptions, null, phasePlan, OptimisticOptimizations.ALL).apply(graph);
+                new InliningPhase(runtime(), null, assumptions, null, phasePlan, OptimisticOptimizations.ALL).apply(graph);
                 Debug.dump(graph, "Graph");
-                new CanonicalizerPhase(null, runtime(), assumptions).apply(graph);
+                new CanonicalizerPhase(runtime(), assumptions).apply(graph);
                 new DeadCodeEliminationPhase().apply(graph);
 
                 assertNotInGraph(graph, Invoke.class);
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.snippets.test/src/com/oracle/graal/snippets/WordTest.java
--- a/graal/com.oracle.graal.snippets.test/src/com/oracle/graal/snippets/WordTest.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.snippets.test/src/com/oracle/graal/snippets/WordTest.java	Mon Feb 25 13:14:39 2013 +0100
@@ -29,7 +29,6 @@
 import com.oracle.graal.api.code.*;
 import com.oracle.graal.api.meta.*;
 import com.oracle.graal.api.runtime.*;
-import com.oracle.graal.compiler.*;
 import com.oracle.graal.compiler.test.*;
 import com.oracle.graal.nodes.*;
 import com.oracle.graal.snippets.Snippet.SnippetInliningPolicy;
@@ -43,7 +42,7 @@
     private final SnippetInstaller installer;
 
     public WordTest() {
-        TargetDescription target = Graal.getRequiredCapability(GraalCompiler.class).target;
+        TargetDescription target = Graal.getRequiredCapability(CodeCacheProvider.class).getTarget();
         installer = new SnippetInstaller(runtime, new Assumptions(false), target);
     }
 
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/SnippetInstaller.java
--- a/graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/SnippetInstaller.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/SnippetInstaller.java	Mon Feb 25 13:14:39 2013 +0100
@@ -254,7 +254,7 @@
 
                 Debug.dump(graph, "after inlining %s", callee);
                 if (GraalOptions.OptCanonicalizer) {
-                    new CanonicalizerPhase(target, runtime, assumptions).apply(graph);
+                    new CanonicalizerPhase(runtime, assumptions).apply(graph);
                 }
                 substituteCallsOriginal = true;
             } else {
@@ -264,7 +264,7 @@
                     Debug.dump(graph, "after inlining %s", callee);
                     if (GraalOptions.OptCanonicalizer) {
                         new WordTypeRewriterPhase(runtime, target.wordKind).apply(graph);
-                        new CanonicalizerPhase(target, runtime, assumptions).apply(graph);
+                        new CanonicalizerPhase(runtime, assumptions).apply(graph);
                     }
                 }
             }
@@ -276,7 +276,7 @@
 
         new DeadCodeEliminationPhase().apply(graph);
         if (GraalOptions.OptCanonicalizer) {
-            new CanonicalizerPhase(target, runtime, assumptions).apply(graph);
+            new CanonicalizerPhase(runtime, assumptions).apply(graph);
         }
 
         for (LoopEndNode end : graph.getNodes(LoopEndNode.class)) {
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/SnippetProvider.java
--- a/graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/SnippetProvider.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/SnippetProvider.java	Mon Feb 25 13:14:39 2013 +0100
@@ -23,8 +23,9 @@
 package com.oracle.graal.snippets;
 
 import com.oracle.graal.api.code.*;
+import com.oracle.graal.compiler.target.*;
 
 public interface SnippetProvider {
 
-    void installSnippets(SnippetInstaller installer, Assumptions assumptions);
+    void installSnippets(Backend backend, SnippetInstaller installer, Assumptions assumptions);
 }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/SnippetTemplate.java
--- a/graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/SnippetTemplate.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/SnippetTemplate.java	Mon Feb 25 13:14:39 2013 +0100
@@ -285,7 +285,7 @@
             new SnippetIntrinsificationPhase(runtime, new BoxingMethodPool(runtime), false).apply(snippetCopy);
             new WordTypeRewriterPhase(runtime, target.wordKind).apply(snippetCopy);
 
-            new CanonicalizerPhase(null, runtime, assumptions, 0, null).apply(snippetCopy);
+            new CanonicalizerPhase(runtime, assumptions, 0, null).apply(snippetCopy);
         }
 
         // Gather the template parameters
@@ -344,7 +344,7 @@
                     LoopEx loop = new LoopsData(snippetCopy).loop(loopBegin);
                     int mark = snippetCopy.getMark();
                     LoopTransformations.fullUnroll(loop, runtime, null);
-                    new CanonicalizerPhase(null, runtime, assumptions, mark, null).apply(snippetCopy);
+                    new CanonicalizerPhase(runtime, assumptions, mark, null).apply(snippetCopy);
                 }
                 FixedNode explodeLoopNext = explodeLoop.next();
                 explodeLoop.clearSuccessors();
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/nodes/DirectObjectStoreNode.java
--- a/graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/nodes/DirectObjectStoreNode.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/nodes/DirectObjectStoreNode.java	Mon Feb 25 13:14:39 2013 +0100
@@ -62,7 +62,7 @@
     @Override
     public void lower(LoweringTool tool) {
         StructuredGraph graph = (StructuredGraph) this.graph();
-        IndexedLocationNode location = IndexedLocationNode.create(LocationNode.ANY_LOCATION, value.kind(), displacement, offset, graph, false);
+        IndexedLocationNode location = IndexedLocationNode.create(LocationNode.ANY_LOCATION, value.kind(), displacement, offset, graph, 1);
         WriteNode write = graph.add(new WriteNode(object, value, location));
         graph.replaceFixedWithFixed(this, write);
     }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/nodes/DirectReadNode.java	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 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	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.snippets/src/com/oracle/graal/snippets/nodes/DirectStoreNode.java	Mon Feb 25 13:14:39 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 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.virtual/src/com/oracle/graal/virtual/phases/ea/PartialEscapeAnalysisPhase.java
--- a/graal/com.oracle.graal.virtual/src/com/oracle/graal/virtual/phases/ea/PartialEscapeAnalysisPhase.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.virtual/src/com/oracle/graal/virtual/phases/ea/PartialEscapeAnalysisPhase.java	Mon Feb 25 13:14:39 2013 +0100
@@ -40,14 +40,12 @@
 
 public class PartialEscapeAnalysisPhase extends Phase {
 
-    private final TargetDescription target;
     private final MetaAccessProvider runtime;
     private final Assumptions assumptions;
     private CustomCanonicalizer customCanonicalizer;
     private final boolean iterative;
 
-    public PartialEscapeAnalysisPhase(TargetDescription target, MetaAccessProvider runtime, Assumptions assumptions, boolean iterative) {
-        this.target = target;
+    public PartialEscapeAnalysisPhase(MetaAccessProvider runtime, Assumptions assumptions, boolean iterative) {
         this.runtime = runtime;
         this.assumptions = assumptions;
         this.iterative = iterative;
@@ -114,7 +112,7 @@
                         return false;
                     }
                     if (GraalOptions.OptCanonicalizer) {
-                        new CanonicalizerPhase(target, runtime, assumptions, null, customCanonicalizer).apply(graph);
+                        new CanonicalizerPhase(runtime, assumptions, null, customCanonicalizer).apply(graph);
                     }
                     return true;
                 }
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.graal.word/src/com/oracle/graal/word/phases/WordTypeRewriterPhase.java
--- a/graal/com.oracle.graal.word/src/com/oracle/graal/word/phases/WordTypeRewriterPhase.java	Mon Feb 25 13:13:02 2013 +0100
+++ b/graal/com.oracle.graal.word/src/com/oracle/graal/word/phases/WordTypeRewriterPhase.java	Mon Feb 25 13:14:39 2013 +0100
@@ -260,7 +260,7 @@
     }
 
     private static ValueNode readOp(StructuredGraph graph, ValueNode base, ValueNode offset, Invoke invoke, Object locationIdentity) {
-        IndexedLocationNode location = IndexedLocationNode.create(locationIdentity, invoke.node().kind(), 0, offset, graph, false);
+        IndexedLocationNode location = IndexedLocationNode.create(locationIdentity, invoke.node().kind(), 0, offset, graph, 1);
         ReadNode read = graph.add(new ReadNode(base, location, invoke.node().stamp()));
         graph.addBeforeFixed(invoke.node(), read);
         // The read must not float outside its block otherwise it may float above an explicit zero
@@ -270,7 +270,7 @@
     }
 
     private static ValueNode writeOp(StructuredGraph graph, ValueNode base, ValueNode offset, ValueNode value, Invoke invoke, Object locationIdentity) {
-        IndexedLocationNode location = IndexedLocationNode.create(locationIdentity, value.kind(), 0, offset, graph, false);
+        IndexedLocationNode location = IndexedLocationNode.create(locationIdentity, value.kind(), 0, offset, graph, 1);
         WriteNode write = graph.add(new WriteNode(base, value, location));
         graph.addBeforeFixed(invoke.node(), write);
         return write;
diff -r 6e4fb0ccebb1 -r a4a8909a6096 graal/com.oracle.truffle.codegen.processor/src/com/oracle/truffle/codegen/processor/node/NodeParser.java
diff -r 6e4fb0ccebb1 -r a4a8909a6096 mx/eclipse-settings/org.eclipse.jdt.core.prefs
--- a/mx/eclipse-settings/org.eclipse.jdt.core.prefs	Mon Feb 25 13:13:02 2013 +0100
+++ b/mx/eclipse-settings/org.eclipse.jdt.core.prefs	Mon Feb 25 13:14:39 2013 +0100
@@ -132,7 +132,7 @@
 org.eclipse.jdt.core.formatter.align_type_members_on_columns=false
 org.eclipse.jdt.core.formatter.alignment_for_arguments_in_allocation_expression=16
 org.eclipse.jdt.core.formatter.alignment_for_arguments_in_annotation=0
-org.eclipse.jdt.core.formatter.alignment_for_arguments_in_enum_constant=16
+org.eclipse.jdt.core.formatter.alignment_for_arguments_in_enum_constant=48
 org.eclipse.jdt.core.formatter.alignment_for_arguments_in_explicit_constructor_call=16
 org.eclipse.jdt.core.formatter.alignment_for_arguments_in_method_invocation=16
 org.eclipse.jdt.core.formatter.alignment_for_arguments_in_qualified_allocation_expression=16
@@ -140,7 +140,7 @@
 org.eclipse.jdt.core.formatter.alignment_for_binary_expression=16
 org.eclipse.jdt.core.formatter.alignment_for_compact_if=16
 org.eclipse.jdt.core.formatter.alignment_for_conditional_expression=80
-org.eclipse.jdt.core.formatter.alignment_for_enum_constants=0
+org.eclipse.jdt.core.formatter.alignment_for_enum_constants=48
 org.eclipse.jdt.core.formatter.alignment_for_expressions_in_array_initializer=16
 org.eclipse.jdt.core.formatter.alignment_for_method_declaration=0
 org.eclipse.jdt.core.formatter.alignment_for_multiple_fields=16
@@ -149,7 +149,7 @@
 org.eclipse.jdt.core.formatter.alignment_for_resources_in_try=80
 org.eclipse.jdt.core.formatter.alignment_for_selector_in_method_invocation=0
 org.eclipse.jdt.core.formatter.alignment_for_superclass_in_type_declaration=16
-org.eclipse.jdt.core.formatter.alignment_for_superinterfaces_in_enum_declaration=16
+org.eclipse.jdt.core.formatter.alignment_for_superinterfaces_in_enum_declaration=48
 org.eclipse.jdt.core.formatter.alignment_for_superinterfaces_in_type_declaration=16
 org.eclipse.jdt.core.formatter.alignment_for_throws_clause_in_constructor_declaration=16
 org.eclipse.jdt.core.formatter.alignment_for_throws_clause_in_method_declaration=16
diff -r 6e4fb0ccebb1 -r a4a8909a6096 mx/projects
--- a/mx/projects	Mon Feb 25 13:13:02 2013 +0100
+++ b/mx/projects	Mon Feb 25 13:14:39 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
 
@@ -242,6 +242,13 @@
 project@com.oracle.graal.compiler.ptx@checkstyle=com.oracle.graal.graph
 project@com.oracle.graal.compiler.ptx@javaCompliance=1.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
+project@com.oracle.graal.compiler.ptx.test@checkstyle=com.oracle.graal.graph
+project@com.oracle.graal.compiler.ptx.test@javaCompliance=1.7
+
 # graal.compiler.sparc
 project@com.oracle.graal.compiler.sparc@subDir=graal
 project@com.oracle.graal.compiler.sparc@sourceDirs=src
@@ -321,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