changeset 11590:797b5ea46a90

Merge
author Mick Jordan <mick.jordan@oracle.com>
date Sat, 07 Sep 2013 12:23:40 -0700
parents 2afda67175e9 (current diff) e12218338164 (diff)
children e460aa80fa46
files graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeUsagesList.java mxtool/.pylintrc
diffstat 61 files changed, 1771 insertions(+), 945 deletions(-) [+]
line wrap: on
line diff
--- a/graal/com.oracle.graal.api.meta.test/src/com/oracle/graal/api/meta/test/TestResolvedJavaType.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.api.meta.test/src/com/oracle/graal/api/meta/test/TestResolvedJavaType.java	Sat Sep 07 12:23:40 2013 -0700
@@ -594,6 +594,7 @@
         "getDeclaredMethods",
         "getDeclaredConstructors",
         "isInitialized",
+        "isLinked",
         "getEncoding",
         "hasFinalizableSubclass",
         "hasFinalizer",
--- a/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/ResolvedJavaMethod.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/ResolvedJavaMethod.java	Sat Sep 07 12:23:40 2013 -0700
@@ -34,9 +34,14 @@
 
     /**
      * Returns the bytecode of this method, if the method has code. The returned byte array does not
-     * contain breakpoints or non-Java bytecodes.
+     * contain breakpoints or non-Java bytecodes. This may return null if the
+     * {@link #getDeclaringClass() holder} is not {@link ResolvedJavaType#isLinked() linked}.
      * 
-     * @return the bytecode of the method, or {@code null} if {@code getCodeSize() == 0}
+     * The contained constant pool indices may not be the ones found in the original class file but
+     * they can be used with the Graal API (e.g. methods in {@link ConstantPool}).
+     * 
+     * @return the bytecode of the method, or {@code null} if {@code getCodeSize() == 0} or if the
+     *         code is not ready.
      */
     byte[] getCode();
 
--- a/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/ResolvedJavaType.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/ResolvedJavaType.java	Sat Sep 07 12:23:40 2013 -0700
@@ -112,7 +112,8 @@
     int getModifiers();
 
     /**
-     * Checks whether this type is initialized.
+     * Checks whether this type is initialized. If a type is initialized it implies that is was
+     * {@link #isLinked() linked} and that the static initializer has run.
      * 
      * @return {@code true} if this type is initialized
      */
@@ -124,6 +125,14 @@
     void initialize();
 
     /**
+     * Checks whether this type is linked and verified. When a type is linked the static initializer
+     * has not necessarily run. An {@link #isInitialized() initialized} type is always linked.
+     * 
+     * @return {@code true} if this type is linked
+     */
+    boolean isLinked();
+
+    /**
      * Determines if this type is either the same as, or is a superclass or superinterface of, the
      * type represented by the specified parameter. This method is identical to
      * {@link Class#isAssignableFrom(Class)} in terms of the value return for this type.
--- a/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java	Sat Sep 07 12:23:40 2013 -0700
@@ -272,69 +272,69 @@
     }
 
     public final void ld_global_b8(Register d, Register a, long immOff) {
-        emitString("ld.global.b8" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.b8" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_b16(Register d, Register a, long immOff) {
-        emitString("ld.global.b16" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.b16" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_b32(Register d, Register a, long immOff) {
-        emitString("ld.global.b32" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.b32" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_b64(Register d, Register a, long immOff) {
-        emitString("ld.global.b64" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.b64" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_u8(Register d, Register a, long immOff) {
-        emitString("ld.global.u8" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.u8" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_u16(Register d, Register a, long immOff) {
-        emitString("ld.global.u16" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.u16" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_u32(Register d, Register a, long immOff) {
-        emitString("ld.global.u32" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.u32" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_u64(Register d, Register a, long immOff) {
-        emitString("ld.global.u64" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.u64" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_s8(Register d, Register a, long immOff) {
-        emitString("ld.global.s8" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.s8" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_s16(Register d, Register a, long immOff) {
-        emitString("ld.global.s16" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.s16" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_s32(Register d, Register a, long immOff) {
-        emitString("ld.global.s32" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.s32" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_s64(Register d, Register a, long immOff) {
-        emitString("ld.global.s64" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.s64" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_f32(Register d, Register a, long immOff) {
-        emitString("ld.global.f32" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.f32" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_f64(Register d, Register a, long immOff) {
-        emitString("ld.global.f64" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.f64" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + "");
     }
 
     // Load from state space to destination register
     public final void ld_from_state_space(String s, Register d, Register a, long immOff) {
-        emitString("ld" + s + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
+        emitString("ld" + s + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + "");
     }
 
     // Load return address from return parameter which is in .param state space
     public final void ld_return_address(String s, Register d, Register a, long immOff) {
-        emitString("ld.param." + s + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.param." + s + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + "");
     }
 
     public final void mov_b16(Register d, Register a) {
@@ -429,68 +429,68 @@
         emitString("mov.f64" + " " + "%r" + d.encoding() + ", " + f64 + ";" + "");
     }
 
-    public final void mul_f32(Register d, Register a, Register b) {
-        emitString("mul.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+    public final void mul_lo_f32(Register d, Register a, Register b) {
+        emitString("mul.lo.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
     }
 
-    public final void mul_f64(Register d, Register a, Register b) {
-        emitString("smul.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+    public final void mul_lo_f64(Register d, Register a, Register b) {
+        emitString("mul.lo.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
     }
 
-    public final void mul_s16(Register d, Register a, Register b) {
-        emitString("mul.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+    public final void mul_lo_s16(Register d, Register a, Register b) {
+        emitString("mul.lo.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
     }
 
-    public final void mul_s32(Register d, Register a, Register b) {
-        emitString("mul.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+    public final void mul_lo_s32(Register d, Register a, Register b) {
+        emitString("mul.lo.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
     }
 
-    public final void mul_s64(Register d, Register a, Register b) {
-        emitString("mul.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+    public final void mul_lo_s64(Register d, Register a, Register b) {
+        emitString("mul.lo.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
     }
 
-    public final void mul_s16(Register d, Register a, short s16) {
-        emitString("mul.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + "");
+    public final void mul_lo_s16(Register d, Register a, short s16) {
+        emitString("mul.lo.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + "");
     }
 
-    public final void mul_s32(Register d, Register a, int s32) {
-        emitString("mul.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + "");
+    public final void mul_lo_s32(Register d, Register a, int s32) {
+        emitString("mul.lo.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + "");
     }
 
-    public final void mul_s64(Register d, Register a, long s64) {
-        emitString("mul.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + "");
+    public final void mul_lo_s64(Register d, Register a, long s64) {
+        emitString("mul.lo.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + "");
     }
 
-    public final void mul_f32(Register d, Register a, float f32) {
-        emitString("mul.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f32 + ";" + "");
+    public final void mul_lo_f32(Register d, Register a, float f32) {
+        emitString("mul.lo.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f32 + ";" + "");
     }
 
-    public final void mul_f64(Register d, Register a, double f64) {
-        emitString("mul.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f64 + ";" + "");
+    public final void mul_lo_f64(Register d, Register a, double f64) {
+        emitString("mul.lo.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f64 + ";" + "");
     }
 
-    public final void mul_u16(Register d, Register a, Register b) {
-        emitString("mul.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+    public final void mul_lo_u16(Register d, Register a, Register b) {
+        emitString("mul.lo.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
     }
 
-    public final void mul_u32(Register d, Register a, Register b) {
-        emitString("mul.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+    public final void mul_lo_u32(Register d, Register a, Register b) {
+        emitString("mul.lo.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
     }
 
-    public final void mul_u64(Register d, Register a, Register b) {
-        emitString("mul.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+    public final void mul_lo_u64(Register d, Register a, Register b) {
+        emitString("mul.lo.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
     }
 
-    public final void mul_u16(Register d, Register a, short u16) {
-        emitString("mul.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + "");
+    public final void mul_lo_u16(Register d, Register a, short u16) {
+        emitString("mul.lo.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + "");
     }
 
-    public final void mul_u32(Register d, Register a, int u32) {
-        emitString("mul.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
+    public final void mul_lo_u32(Register d, Register a, int u32) {
+        emitString("mul.lo.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
     }
 
-    public final void mul_u64(Register d, Register a, long u64) {
-        emitString("mul.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + "");
+    public final void mul_lo_u64(Register d, Register a, long u64) {
+        emitString("mul.lo.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + "");
     }
 
     public final void neg_f32(Register d, Register a) {
@@ -550,15 +550,15 @@
     }
 
     public final void param_8_decl(Register d, boolean lastParam) {
-        emitString(".param" + " " + ".s8" + " " + d.toString() + (lastParam ? "" : ","));
+        emitString(".param" + " " + ".s8" + " " + d + (lastParam ? "" : ","));
     }
 
     public final void param_32_decl(Register d, boolean lastParam) {
-        emitString(".param" + " " + ".s32" + " " + d.toString() + (lastParam ? "" : ","));
+        emitString(".param" + " " + ".s32" + " " + d + (lastParam ? "" : ","));
     }
 
     public final void param_64_decl(Register d, boolean lastParam) {
-        emitString(".param" + " " + ".s64" + " " + d.toString() + (lastParam ? "" : ","));
+        emitString(".param" + " " + ".s64" + " " + d + (lastParam ? "" : ","));
     }
 
     public final void popc_b32(Register d, Register a) {
@@ -849,54 +849,32 @@
         emitString("setp.ge.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + "");
     }
 
-    public final void shl_s16(Register d, Register a, Register b) {
-        emitString("shl.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
-    }
-
-    public final void shl_s32(Register d, Register a, Register b) {
-        emitString("shl.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+    // Shift left - only types supported are .b16, .b32 and .b64
+    public final void shl_b16(Register d, Register a, Register b) {
+        emitString("shl.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
     }
 
-    public final void shl_s64(Register d, Register a, Register b) {
-        emitString("shl.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+    public final void shl_b32(Register d, Register a, Register b) {
+        emitString("shl.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
     }
 
-    public final void shl_s16(Register d, Register a, int u32) {
-        emitString("shl.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
-    }
-
-    public final void shl_s32(Register d, Register a, int u32) {
-        emitString("shl.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
-    }
-
-    public final void shl_s64(Register d, Register a, int u32) {
-        emitString("shl.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
+    public final void shl_b64(Register d, Register a, Register b) {
+        emitString("shl.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
     }
 
-    public final void shl_u16(Register d, Register a, Register b) {
-        emitString("shl.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
-    }
-
-    public final void shl_u32(Register d, Register a, Register b) {
-        emitString("shl.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
-    }
-
-    public final void shl_u64(Register d, Register a, Register b) {
-        emitString("shl.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
+    public final void shl_b16_const(Register d, Register a, int b) {
+        emitString("shl.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b + ";" + "");
     }
 
-    public final void shl_u16(Register d, Register a, int u32) {
-        emitString("shl.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
+    public final void shl_b32_const(Register d, Register a, int b) {
+        emitString("shl.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b + ";" + "");
     }
 
-    public final void shl_u32(Register d, Register a, int u32) {
-        emitString("shl.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
+    public final void shl_b64_const(Register d, Register a, int b) {
+        emitString("shl.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b + ";" + "");
     }
 
-    public final void shl_u64(Register d, Register a, int u32) {
-        emitString("shl.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + "");
-    }
-
+    // Shift Right instruction
     public final void shr_s16(Register d, Register a, Register b) {
         emitString("shr.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
     }
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/IntegerPTXTest.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/IntegerPTXTest.java	Sat Sep 07 12:23:40 2013 -0700
@@ -22,7 +22,7 @@
  */
 package com.oracle.graal.compiler.ptx.test;
 
-import org.junit.Test;
+import org.junit.*;
 
 import java.lang.reflect.Method;
 
@@ -32,10 +32,19 @@
     @Test
     public void testAdd() {
 
+        Integer r4 = (Integer) invoke(compile("testAdd2I"), 18, 24);
+        if (r4 == null) {
+            printReport("testAdd2I FAILED");
+        } else if (r4.intValue() == testAdd2I(18, 24)) {
+            printReport("testAdd2I PASSED");
+        } else {
+            printReport("testAdd2I FAILED");
+        }
+
         Long r2 = (Long) invoke(compile("testAdd2L"), (long) 12, (long) 6);
         if (r2 == null) {
             printReport("testAdd2L FAILED");
-        } else if (r2.longValue() == 18) {
+        } else if (r2.longValue() == testAdd2L(12, 6)) {
             printReport("testAdd2L PASSED");
         } else {
             printReport("testAdd2L FAILED");
@@ -43,10 +52,10 @@
 
         //invoke(compile("testAdd2B"), (byte) 6, (byte) 4);
 
-        Integer r4 = (Integer) invoke(compile("testAddIConst"), 5);
+        r4 = (Integer) invoke(compile("testAddIConst"), 5);
         if (r4 == null) {
             printReport("testAddIConst FAILED");
-        } else if (r4.intValue() == 37) {
+        } else if (r4.intValue() == testAddIConst(5)) {
             printReport("testAddIConst PASSED");
         } else {
             printReport("testAddIConst FAILED");
@@ -55,20 +64,12 @@
         r4 = (Integer) invoke(compile("testAddConstI"), 7);
         if (r4 == null) {
             printReport("testAddConstI FAILED");
-        } else if (r4.intValue() == 39) {
+        } else if (r4.intValue() == testAddConstI(7)) {
             printReport("testAddConstI PASSED");
         } else {
             printReport("testAddConstI FAILED");
         }
 
-        r4 = (Integer) invoke(compile("testAdd2I"), 18, 24);
-        if (r4 == null) {
-            printReport("testAdd2I FAILED");
-        } else if (r4.intValue() == 42) {
-            printReport("testAdd2I PASSED");
-        } else {
-            printReport("testAdd2I FAILED");
-        }
     }
 
     public static int testAdd2I(int a, int b) {
@@ -93,20 +94,21 @@
 
     @Test
     public void testSub() {
-        Long r2 = (Long) invoke(compile("testSub2L"), (long) 12, (long) 6);
-        if (r2 == null) {
-            printReport("testSub2I FAILED (null return value)");
-        } else if (r2.longValue() == 6) {
+
+        Integer r1 = (Integer) invoke(compile("testSub2I"), 18, 4);
+
+        if (r1 == null) {
+            printReport("testSub2I FAILED");
+        } else if (r1.intValue() == testSub2I(18, 4)) {
             printReport("testSub2I PASSED");
         } else {
             printReport("testSub2I FAILED");
         }
 
-        Integer r1 = (Integer) invoke(compile("testSub2I"), 18, 4);
-
-        if (r1 == null) {
-            printReport("testSub2I FAILED");
-        } else if (r1.intValue() == 14) {
+        Long r2 = (Long) invoke(compile("testSub2L"), (long) 12, (long) 6);
+        if (r2 == null) {
+            printReport("testSub2I FAILED (null return value)");
+        } else if (r2.longValue() == testSub2L(12, 6)) {
             printReport("testSub2I PASSED");
         } else {
             printReport("testSub2I FAILED");
@@ -115,7 +117,7 @@
         r1 = (Integer) invoke(compile("testSubIConst"), 35);
         if (r1 == null) {
             printReport("testSubIConst FAILED");
-        } else if (r1.intValue() == 3) {
+        } else if (r1.intValue() == testSubIConst(35)) {
             printReport("testSubIConst PASSED");
         } else {
             printReport("testSubIConst FAILED");
@@ -124,7 +126,7 @@
         r1 = (Integer) invoke(compile("testSubConstI"), 12);
         if (r1 == null) {
             printReport("testSubConstI FAILED");
-        } else if (r1.intValue() == 20) {
+        } else if (r1.intValue() == testSubConstI(12)) {
             printReport("testSubConstI PASSED");
         } else {
             printReport("testSubConstI FAILED");
@@ -149,10 +151,42 @@
 
     @Test
     public void testMul() {
-        invoke(compile("testMul2I"), 8, 4);
-        invoke(compile("testMul2L"), (long) 12, (long) 6);
-        invoke(compile("testMulIConst"), 4);
-        invoke(compile("testMulConstI"), 5);
+
+        Integer r1 = (Integer) invoke(compile("testMul2I"), 8, 4);
+        if (r1 == null) {
+            printReport("testMul2I FAILED");
+        } else if (r1.intValue() == testMul2I(8, 4)) {
+            printReport("testMul2I PASSED");
+        } else {
+            printReport("testMul2I FAILED");
+        }
+
+        Long r2 = (Long) invoke(compile("testMul2L"), (long) 12, (long) 6);
+        if (r2 == null) {
+            printReport("testMul2L FAILED");
+        } else if (r2.longValue() == testMul2L(12, 6)) {
+            printReport("testMul2L PASSED");
+        } else {
+            printReport("testMul2L FAILED");
+        }
+
+        r1 = (Integer) invoke(compile("testMulIConst"), 4);
+        if (r1 == null) {
+            printReport("testMulIConst FAILED");
+        } else if (r1.intValue() == testMulIConst(4)) {
+            printReport("testMulIConst PASSED");
+        } else {
+            printReport("testMulIConst FAILED");
+        }
+
+        r1 = (Integer) invoke(compile("testMulConstI"), 5);
+        if (r1 == null) {
+            printReport("testMulConstI FAILED");
+        } else if (r1.intValue() == testMulConstI(5)) {
+            printReport("testMulConstI PASSED");
+        } else {
+            printReport("testMulConstI FAILED");
+        }
     }
 
     public static int testMul2I(int a, int b) {
@@ -170,12 +204,44 @@
     public static int testMulConstI(int a) {
         return 32 * a;
     }
+
     @Test
     public void testDiv() {
-        invoke(compile("testDiv2I"), 8, 4);
-        invoke(compile("testDiv2L"), (long) 12, (long) 6);
-        invoke(compile("testDivIConst"), 64);
-        invoke(compile("testDivConstI"), 8);
+        Integer r1 = (Integer) invoke(compile("testDiv2I"), 8, 4);
+        if (r1 == null) {
+            printReport("testDiv2I FAILED (null value returned)");
+        } else if (r1.intValue() == testDiv2I(8, 4)) {
+            printReport("testDiv2I PASSED");
+        } else {
+            printReport("testDiv2I FAILED");
+        }
+
+        Long r2 = (Long) invoke(compile("testDiv2L"), (long) 12, (long) 6);
+        if (r2 == null) {
+            printReport("testDiv2L FAILED (null value returned)");
+        } else if (r2.longValue() == testDiv2L(12, 6)) {
+            printReport("testDiv2L PASSED");
+        } else {
+            printReport("testDiv2L FAILED");
+        }
+
+        r1 = (Integer) invoke(compile("testDivIConst"), 64);
+        if (r1 == null) {
+            printReport("testDivIConst FAILED (null value returned)");
+        } else if (r1.intValue() == testDivIConst(64)) {
+            printReport("testDivIConst PASSED");
+        } else {
+            printReport("testDivIConst FAILED");
+        }
+
+        r1 = (Integer) invoke(compile("testDivConstI"), 8);
+        if (r1 == null) {
+            printReport("testDivConstI FAILED (null value returned)");
+        } else if (r1.intValue() == testDivConstI(8)) {
+            printReport("testDivConstI PASSED");
+        } else {
+            printReport("testDivConstI FAILED");
+        }
     }
 
     public static int testDiv2I(int a, int b) {
@@ -196,8 +262,23 @@
 
     @Test
     public void testRem() {
-        invoke(compile("testRem2I"), 8, 4);
-        invoke(compile("testRem2L"), (long) 12, (long) 6);
+        Integer r1 = (Integer) invoke(compile("testRem2I"), 8, 4);
+        if (r1 == null) {
+            printReport("testRem2I FAILED (null value returned)");
+        } else if (r1.intValue() == testRem2I(8, 4)) {
+            printReport("testRem2I PASSED");
+        } else {
+            printReport("testRem2I FAILED");
+        }
+
+        Long r2 = (Long) invoke(compile("testRem2L"), (long) 12, (long) 6);
+        if (r2 == null) {
+            printReport("testRem2L FAILED (null value returned)");
+        } else if (r1.longValue() == testRem2L(12, 6)) {
+            printReport("testRem2L PASSED");
+        } else {
+            printReport("testRem2L FAILED");
+        }
     }
 
     public static int testRem2I(int a, int b) {
@@ -207,11 +288,27 @@
     public static long testRem2L(long a, long b) {
         return a % b;
     }
-
+    @Ignore
     @Test
     public void testIntConversion() {
-        invoke(compile("testI2L"), 8);
-        invoke(compile("testL2I"), (long) 12);
+        Long r1 = (Long) invoke(compile("testI2L"), 8);
+        if (r1 == null) {
+            printReport("testI2L FAILED (null value returned)");
+        } else if (r1.longValue() == testI2L(8)) {
+            printReport("testI2L PASSED");
+        } else {
+            printReport("testI2L FAILED");
+        }
+
+        Integer r2 = (Integer) invoke(compile("testL2I"), (long) 12);
+        if (r2 == null) {
+            printReport("testL2I FAILED (null value returned)");
+        } else if (r1.longValue() == testL2I(12)) {
+            printReport("testL2I PASSED");
+        } else {
+            printReport("testL2I FAILED");
+        }
+
         // invoke(compile("testI2C"), 65);
         // invoke(compile("testI2B"), 9);
         // invoke(compile("testI2F"), 17);
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java	Sat Sep 07 12:23:40 2013 -0700
@@ -174,7 +174,7 @@
         }
 
         for (Integer i : signed32) {
-            codeBuffer.emitString("  .reg .s32 %r" + i.intValue() + ";");
+            codeBuffer.emitString(".reg .s32 %r" + i.intValue() + ";");
         }
         for (Integer i : signed64) {
             codeBuffer.emitString(".reg .s64 %r" + i.intValue() + ";");
--- a/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/alloc/LinearScan.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/alloc/LinearScan.java	Sat Sep 07 12:23:40 2013 -0700
@@ -686,9 +686,8 @@
             List<LIRInstruction> instructions = ir.lir(block);
             int numInst = instructions.size();
 
-            // iterate all instructions of the block. skip the first because it is always a label
-            assert !instructions.get(0).hasOperands() : "first operation must always be a label";
-            for (int j = 1; j < numInst; j++) {
+            // iterate all instructions of the block
+            for (int j = 0; j < numInst; j++) {
                 final LIRInstruction op = instructions.get(j);
 
                 ValueProcedure useProc = new ValueProcedure() {
@@ -1174,10 +1173,8 @@
             }
 
             // iterate all instructions of the block in reverse order.
-            // skip the first instruction because it is always a label
             // definitions of intervals are processed before uses
-            assert !instructions.get(0).hasOperands() : "first operation must always be a label";
-            for (int j = instructions.size() - 1; j >= 1; j--) {
+            for (int j = instructions.size() - 1; j >= 0; j--) {
                 final LIRInstruction op = instructions.get(j);
                 final int opId = op.id();
 
--- a/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/gen/LIRGenerator.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/gen/LIRGenerator.java	Sat Sep 07 12:23:40 2013 -0700
@@ -40,7 +40,6 @@
 import com.oracle.graal.lir.*;
 import com.oracle.graal.lir.StandardOp.JumpOp;
 import com.oracle.graal.lir.StandardOp.LabelOp;
-import com.oracle.graal.lir.StandardOp.ParametersOp;
 import com.oracle.graal.nodes.*;
 import com.oracle.graal.nodes.PhiNode.PhiType;
 import com.oracle.graal.nodes.calc.*;
@@ -405,7 +404,7 @@
     }
 
     public void emitIncomingValues(Value[] params) {
-        append(new ParametersOp(params));
+        ((LabelOp) lir.lir(currentBlock).get(0)).setIncomingValues(params);
     }
 
     @Override
--- a/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/Graph.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/Graph.java	Sat Sep 07 12:23:40 2013 -0700
@@ -40,6 +40,16 @@
 
     private final ArrayList<Node> nodes;
 
+    /**
+     * Records the modification count for nodes. This is only used in assertions.
+     */
+    private int[] nodeModCounts;
+
+    /**
+     * Records the modification count for nodes' usage lists. This is only used in assertions.
+     */
+    private int[] nodeUsageModCounts;
+
     // these two arrays contain one entry for each NodeClass, indexed by NodeClass.iterableId.
     // they contain the first and last pointer to a linked list of all nodes with this type.
     private final ArrayList<Node> nodeCacheFirst;
@@ -87,6 +97,18 @@
         this(null);
     }
 
+    static final boolean MODIFICATION_COUNTS_ENABLED = assertionsEnabled();
+
+    /**
+     * Determines if assertions are enabled for the {@link Graph} class.
+     */
+    @SuppressWarnings("all")
+    private static boolean assertionsEnabled() {
+        boolean enabled = false;
+        assert enabled = true;
+        return enabled;
+    }
+
     /**
      * Creates an empty Graph with a given name.
      * 
@@ -97,6 +119,58 @@
         nodeCacheFirst = new ArrayList<>(NodeClass.cacheSize());
         nodeCacheLast = new ArrayList<>(NodeClass.cacheSize());
         this.name = name;
+        if (MODIFICATION_COUNTS_ENABLED) {
+            nodeModCounts = new int[nodes.size()];
+            nodeUsageModCounts = new int[nodes.size()];
+        }
+    }
+
+    int extractOriginalNodeId(Node node) {
+        int id = node.id;
+        if (id <= Node.DELETED_ID_START) {
+            id = Node.DELETED_ID_START - id;
+        }
+        return id;
+    }
+
+    int modCount(Node node) {
+        int id = extractOriginalNodeId(node);
+        if (id >= 0 && id < nodeModCounts.length) {
+            return nodeModCounts[id];
+        }
+        return 0;
+    }
+
+    void incModCount(Node node) {
+        int id = extractOriginalNodeId(node);
+        if (id >= 0) {
+            if (id >= nodeModCounts.length) {
+                nodeModCounts = Arrays.copyOf(nodeModCounts, id + 30);
+            }
+            nodeModCounts[id]++;
+        } else {
+            assert false;
+        }
+    }
+
+    int usageModCount(Node node) {
+        int id = extractOriginalNodeId(node);
+        if (id >= 0 && id < nodeUsageModCounts.length) {
+            return nodeUsageModCounts[node.id];
+        }
+        return 0;
+    }
+
+    void incUsageModCount(Node node) {
+        int id = extractOriginalNodeId(node);
+        if (id >= 0) {
+            if (id >= nodeUsageModCounts.length) {
+                nodeUsageModCounts = Arrays.copyOf(nodeUsageModCounts, id + 30);
+            }
+            nodeUsageModCounts[id]++;
+        } else {
+            assert false;
+        }
     }
 
     /**
--- a/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/Node.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/Node.java	Sat Sep 07 12:23:40 2013 -0700
@@ -22,6 +22,8 @@
  */
 package com.oracle.graal.graph;
 
+import static com.oracle.graal.graph.Graph.*;
+
 import java.lang.annotation.*;
 import java.util.*;
 
@@ -119,9 +121,18 @@
     // therefore points to the next Node of the same type.
     Node typeCacheNext;
 
-    private NodeUsagesList usages;
+    private static final int INLINE_USAGE_COUNT = 2;
+    private static final Node[] NO_NODES = {};
+
+    /**
+     * Head of usage list. The elements of the usage list in order are {@link #usage0},
+     * {@link #usage1} and {@link #extraUsages}. The first null entry terminates the list.
+     */
+    private Node usage0;
+    private Node usage1;
+    private Node[] extraUsages = NO_NODES;
+
     private Node predecessor;
-    private int modCount;
 
     public Node() {
         this.graph = null;
@@ -156,8 +167,211 @@
         return getNodeClass().getSuccessorIterable(this);
     }
 
+    class NodeUsageIterator implements Iterator<Node> {
+
+        private final int expectedModCount = usageModCount();
+        int index = -1;
+        Node current;
+
+        private void advance() {
+            assert index == -1 || current != null;
+            current = null;
+            index++;
+            if (index == 0) {
+                current = usage0;
+            } else if (index == 1) {
+                current = usage1;
+            } else {
+                if (index - INLINE_USAGE_COUNT < extraUsages.length) {
+                    current = extraUsages[index - INLINE_USAGE_COUNT];
+                }
+            }
+        }
+
+        public NodeUsageIterator() {
+            advance();
+        }
+
+        public boolean hasNext() {
+            assert expectedModCount == usageModCount();
+            return current != null;
+        }
+
+        public Node next() {
+            assert expectedModCount == usageModCount();
+            Node result = current;
+            if (result == null) {
+                throw new NoSuchElementException();
+            }
+            advance();
+            return result;
+        }
+
+        public void remove() {
+            throw new UnsupportedOperationException();
+        }
+    }
+
+    class NodeUsageIterable extends AbstractNodeIterable<Node> {
+
+        public NodeUsageIterator iterator() {
+            return new NodeUsageIterator();
+        }
+
+        @Override
+        public boolean isEmpty() {
+            return usage0 == null;
+        }
+
+        @Override
+        public boolean isNotEmpty() {
+            return usage0 != null;
+        }
+
+        @Override
+        public int count() {
+            if (usage0 == null) {
+                return 0;
+            }
+            if (usage1 == null) {
+                return 1;
+            }
+            return 2 + indexOfLastNonNull(extraUsages) + 1;
+        }
+    }
+
+    /**
+     * Gets the list of nodes that use this node (e.g., as an input).
+     */
     public final NodeIterable<Node> usages() {
-        return usages;
+        return new NodeUsageIterable();
+    }
+
+    /**
+     * Finds the index of the last non-null entry in a node array. The search assumes that all
+     * non-null entries precede the first null entry in the array.
+     * 
+     * @param nodes the array to search
+     * @return the index of the last non-null entry in {@code nodes} if it exists, else -1
+     */
+    private static int indexOfLastNonNull(Node[] nodes) {
+        if (nodes.length == 0 || nodes[0] == null) {
+            return -1;
+        }
+        if (nodes[nodes.length - 1] != null) {
+            return nodes.length - 1;
+        }
+
+        // binary search
+        int low = 0;
+        int high = nodes.length - 1;
+        while (true) {
+            int mid = (low + high) >>> 1;
+            if (nodes[mid] == null) {
+                if (nodes[mid - 1] != null) {
+                    return mid - 1;
+                }
+                high = mid - 1;
+            } else {
+                if (mid == nodes.length - 1 || nodes[mid + 1] == null) {
+                    return mid;
+                }
+                low = mid + 1;
+            }
+        }
+    }
+
+    /**
+     * Adds a given node to this node's {@linkplain #usages() usages}.
+     * 
+     * @param node the node to add
+     */
+    private void addUsage(Node node) {
+        incUsageModCount();
+        if (usage0 == null) {
+            usage0 = node;
+        } else if (usage1 == null) {
+            usage1 = node;
+        } else {
+            int length = extraUsages.length;
+            if (length == 0) {
+                extraUsages = new Node[4];
+                extraUsages[0] = node;
+            } else {
+                int lastNonNull = indexOfLastNonNull(extraUsages);
+                if (lastNonNull == length - 1) {
+                    extraUsages = Arrays.copyOf(extraUsages, length * 2 + 1);
+                    extraUsages[length] = node;
+                } else if (lastNonNull == -1) {
+                    extraUsages[0] = node;
+                } else {
+                    extraUsages[lastNonNull + 1] = node;
+                }
+            }
+        }
+    }
+
+    /**
+     * Removes a given node from this node's {@linkplain #usages() usages}.
+     * 
+     * @param node the node to remove
+     * @return whether or not {@code usage} was in the usage list
+     */
+    private boolean removeUsage(Node node) {
+        // It is critical that this method maintains the invariant that
+        // the usage list has no null element preceding a non-null element
+        incUsageModCount();
+        if (usage0 == node) {
+            if (usage1 != null) {
+                int lastNonNull = indexOfLastNonNull(extraUsages);
+                if (lastNonNull >= 0) {
+                    usage0 = extraUsages[lastNonNull];
+                    extraUsages[lastNonNull] = null;
+                } else {
+                    // usage1 is the last element
+                    usage0 = usage1;
+                    usage1 = null;
+                }
+            } else {
+                // usage0 is the last element
+                usage0 = null;
+            }
+            return true;
+        }
+        if (usage1 == node) {
+            int lastNonNull = indexOfLastNonNull(extraUsages);
+            if (lastNonNull >= 0) {
+                usage1 = extraUsages[lastNonNull];
+                extraUsages[lastNonNull] = null;
+            } else {
+                // usage1 is the last element
+                usage1 = null;
+            }
+            return true;
+        }
+        int lastNonNull = indexOfLastNonNull(extraUsages);
+        if (lastNonNull >= 0) {
+            for (int i = 0; i <= lastNonNull; ++i) {
+                Node n = extraUsages[i];
+                if (n == node) {
+                    if (i < lastNonNull) {
+                        extraUsages[i] = extraUsages[lastNonNull];
+                        extraUsages[lastNonNull] = null;
+                    } else {
+                        extraUsages[i] = null;
+                    }
+                    return true;
+                }
+            }
+        }
+        return false;
+    }
+
+    private void clearUsages() {
+        incUsageModCount();
+        usage0 = null;
+        usage1 = null;
+        extraUsages = NO_NODES;
     }
 
     public final Node predecessor() {
@@ -165,11 +379,29 @@
     }
 
     final int modCount() {
-        return modCount;
+        if (MODIFICATION_COUNTS_ENABLED && graph != null) {
+            return graph.modCount(this);
+        }
+        return 0;
     }
 
     final void incModCount() {
-        modCount++;
+        if (MODIFICATION_COUNTS_ENABLED && graph != null) {
+            graph.incModCount(this);
+        }
+    }
+
+    final int usageModCount() {
+        if (MODIFICATION_COUNTS_ENABLED && graph != null) {
+            return graph.usageModCount(this);
+        }
+        return 0;
+    }
+
+    final void incUsageModCount() {
+        if (MODIFICATION_COUNTS_ENABLED && graph != null) {
+            graph.incUsageModCount(this);
+        }
     }
 
     public boolean isDeleted() {
@@ -185,7 +417,6 @@
      * newInput: removes this node from oldInput's usages and adds this node to newInput's usages.
      */
     protected void updateUsages(Node oldInput, Node newInput) {
-        assert assertTrue(usages != null, "usages == null while adding %s to %s", newInput, this);
         if (oldInput != newInput) {
             if (oldInput != null) {
                 boolean result = removeThisFromUsages(oldInput);
@@ -196,8 +427,7 @@
                 if (inputChanged != null) {
                     inputChanged.nodeChanged(this);
                 }
-                assert newInput.usages != null : "not yet added? " + newInput;
-                newInput.usages.add(this);
+                newInput.addUsage(this);
             } else if (oldInput != null && oldInput.usages().isEmpty()) {
                 NodeChangedListener nodeChangedListener = graph.usagesDroppedZero;
                 if (nodeChangedListener != null) {
@@ -213,7 +443,6 @@
      * this node to newSuccessor's predecessors.
      */
     protected void updatePredecessor(Node oldSuccessor, Node newSuccessor) {
-        assert assertTrue(usages != null, "usages == null while adding %s to %s", newSuccessor, this);
         if (oldSuccessor != newSuccessor) {
             if (oldSuccessor != null) {
                 assert assertTrue(oldSuccessor.predecessor == this, "wrong predecessor in old successor (%s): %s", oldSuccessor, oldSuccessor.predecessor);
@@ -230,7 +459,6 @@
         assert assertTrue(id == INITIAL_ID, "unexpected id: %d", id);
         this.graph = newGraph;
         newGraph.register(this);
-        usages = new NodeUsagesList();
         for (Node input : inputs()) {
             updateUsages(null, input);
         }
@@ -253,7 +481,7 @@
 
     public void replaceAtUsages(Node other) {
         assert checkReplaceWith(other);
-        for (Node usage : usages) {
+        for (Node usage : usages()) {
             boolean result = usage.getNodeClass().replaceFirstInput(usage, this, other);
             assert assertTrue(result, "not found in inputs, usage: %s", usage);
             if (other != null) {
@@ -261,10 +489,10 @@
                 if (inputChanged != null) {
                     inputChanged.nodeChanged(usage);
                 }
-                other.usages.add(usage);
+                other.addUsage(usage);
             }
         }
-        usages.clear();
+        clearUsages();
     }
 
     public void replaceAtPredecessor(Node other) {
@@ -314,11 +542,7 @@
     }
 
     private boolean removeThisFromUsages(Node n) {
-        if (n.usages.remove(this)) {
-            return true;
-        } else {
-            return false;
-        }
+        return n.removeUsage(this);
     }
 
     public void clearSuccessors() {
@@ -332,7 +556,7 @@
     }
 
     private boolean checkDeletion() {
-        assertTrue(usages.isEmpty(), "cannot delete node %s because of usages: %s", this, usages);
+        assertTrue(usages().isEmpty(), "cannot delete node %s because of usages: %s", this, usages());
         assertTrue(predecessor == null, "cannot delete node %s because of predecessor: %s", this, predecessor);
         return true;
     }
@@ -355,7 +579,7 @@
         NodeClass clazz = getNodeClass();
         clazz.copyInputs(this, newNode);
         for (Node input : inputs()) {
-            input.usages.add(newNode);
+            input.addUsage(newNode);
         }
         return newNode;
     }
@@ -373,9 +597,10 @@
         newNode.typeCacheNext = null;
         newNode.id = INITIAL_ID;
         into.register(newNode);
-        newNode.usages = new NodeUsagesList();
+        newNode.usage0 = null;
+        newNode.usage1 = null;
+        newNode.extraUsages = NO_NODES;
         newNode.predecessor = null;
-        newNode.modCount = 0;
         return newNode;
     }
 
@@ -576,10 +801,10 @@
         }
 
         if (precision > 0) {
-            if (this.usages.count() > 0) {
+            if (!usages().isEmpty()) {
                 formatter.format(" usages={");
                 int z = 0;
-                for (Node usage : this.usages) {
+                for (Node usage : usages()) {
                     if (z != 0) {
                         formatter.format(", ");
                     }
--- a/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeClass.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeClass.java	Sat Sep 07 12:23:40 2013 -0700
@@ -22,6 +22,8 @@
  */
 package com.oracle.graal.graph;
 
+import static com.oracle.graal.graph.Graph.*;
+
 import java.lang.reflect.*;
 import java.util.*;
 import java.util.Map.Entry;
@@ -432,7 +434,7 @@
          */
         private NodeClassIterator(Node node, long[] offsets, int directCount) {
             this.node = node;
-            this.modCount = node.modCount();
+            this.modCount = MODIFICATION_COUNTS_ENABLED ? node.modCount() : 0;
             this.offsets = offsets;
             this.directCount = directCount;
             index = NOT_ITERABLE;
--- a/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeInputList.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeInputList.java	Sat Sep 07 12:23:40 2013 -0700
@@ -39,13 +39,13 @@
 
     public NodeInputList(Node self, T[] elements) {
         super(elements);
-        assert self.usages() == null;
+        assert self.usages().isEmpty();
         this.self = self;
     }
 
     public NodeInputList(Node self, List<? extends T> elements) {
         super(elements);
-        assert self.usages() == null;
+        assert self.usages().isEmpty();
         this.self = self;
     }
 
--- a/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeSuccessorList.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeSuccessorList.java	Sat Sep 07 12:23:40 2013 -0700
@@ -37,7 +37,7 @@
 
     public NodeSuccessorList(Node self, T[] elements) {
         super(elements);
-        assert self.usages() == null;
+        assert self.usages().isEmpty();
         this.self = self;
     }
 
--- a/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeUsagesList.java	Tue Sep 03 16:48:17 2013 -0700
+++ /dev/null	Thu Jan 01 00:00:00 1970 +0000
@@ -1,163 +0,0 @@
-/*
- * 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.graph;
-
-import java.util.*;
-
-import com.oracle.graal.graph.iterators.*;
-
-public final class NodeUsagesList extends AbstractNodeIterable<Node> {
-
-    protected static final Node[] EMPTY_NODE_ARRAY = new Node[0];
-
-    protected Node[] nodes = EMPTY_NODE_ARRAY;
-    private int size;
-    private int modCount;
-
-    NodeUsagesList() {
-        this.size = 0;
-        this.nodes = EMPTY_NODE_ARRAY;
-    }
-
-    NodeUsagesList(Node[] nodes) {
-        this.size = nodes.length;
-        this.nodes = nodes;
-    }
-
-    @Override
-    public boolean isEmpty() {
-        return size == 0;
-    }
-
-    @Override
-    public boolean isNotEmpty() {
-        return size > 0;
-    }
-
-    @Override
-    public int count() {
-        return size;
-    }
-
-    @Override
-    public Iterator<Node> iterator() {
-        return new Iterator<Node>() {
-
-            private final int expectedModCount = NodeUsagesList.this.modCount;
-            private int index = 0;
-
-            @Override
-            public boolean hasNext() {
-                assert expectedModCount == NodeUsagesList.this.modCount;
-                return index < NodeUsagesList.this.size;
-            }
-
-            @Override
-            public Node next() {
-                assert expectedModCount == NodeUsagesList.this.modCount;
-                return NodeUsagesList.this.nodes[index++];
-            }
-
-            @Override
-            public void remove() {
-                throw new UnsupportedOperationException();
-            }
-        };
-    }
-
-    @Override
-    public boolean contains(Node other) {
-        for (int i = 0; i < size; i++) {
-            if (nodes[i] == other) {
-                return true;
-            }
-        }
-        return false;
-    }
-
-    @Override
-    public List<Node> snapshot() {
-        return Arrays.asList(Arrays.copyOf(NodeUsagesList.this.nodes, NodeUsagesList.this.size));
-    }
-
-    private void incModCount() {
-        modCount++;
-    }
-
-    boolean add(Node node) {
-        incModCount();
-        if (size == nodes.length) {
-            nodes = Arrays.copyOf(nodes, nodes.length * 2 + 1);
-        }
-        nodes[size++] = node;
-        return true;
-    }
-
-    void copyAndClear(NodeUsagesList other) {
-        incModCount();
-        other.incModCount();
-        nodes = other.nodes;
-        size = other.size;
-        nodes = EMPTY_NODE_ARRAY;
-        size = 0;
-    }
-
-    void clear() {
-        incModCount();
-        nodes = EMPTY_NODE_ARRAY;
-        size = 0;
-    }
-
-    boolean remove(Node node) {
-        int i = 0;
-        incModCount();
-        while (i < size && nodes[i] != node) {
-            i++;
-        }
-        if (i < size) {
-            i++;
-            while (i < size) {
-                nodes[i - 1] = nodes[i];
-                i++;
-            }
-            nodes[--size] = null;
-            return true;
-        } else {
-            return false;
-        }
-    }
-
-    @Override
-    public String toString() {
-        StringBuilder str = new StringBuilder();
-        str.append('[');
-        for (int i = 0; i < size; i++) {
-            if (i > 0) {
-                str.append(", ");
-            }
-            str.append(nodes[i]);
-        }
-        str.append(']');
-        return str.toString();
-    }
-}
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotBackend.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotBackend.java	Sat Sep 07 12:23:40 2013 -0700
@@ -45,7 +45,7 @@
 import com.oracle.graal.hotspot.stubs.*;
 import com.oracle.graal.lir.*;
 import com.oracle.graal.lir.LIRInstruction.ValueProcedure;
-import com.oracle.graal.lir.StandardOp.ParametersOp;
+import com.oracle.graal.lir.StandardOp.LabelOp;
 import com.oracle.graal.lir.amd64.*;
 import com.oracle.graal.lir.asm.*;
 import com.oracle.graal.nodes.*;
@@ -214,7 +214,7 @@
         };
         for (Block block : lir.codeEmittingOrder()) {
             for (LIRInstruction op : lir.lir(block)) {
-                if (op instanceof ParametersOp) {
+                if (op instanceof LabelOp) {
                     // Don't consider this as a definition
                 } else {
                     op.forEachTemp(defProc);
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotLIRGenerator.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotLIRGenerator.java	Sat Sep 07 12:23:40 2013 -0700
@@ -46,7 +46,6 @@
 import com.oracle.graal.hotspot.nodes.*;
 import com.oracle.graal.hotspot.stubs.*;
 import com.oracle.graal.lir.*;
-import com.oracle.graal.lir.StandardOp.ParametersOp;
 import com.oracle.graal.lir.StandardOp.PlaceholderOp;
 import com.oracle.graal.lir.amd64.*;
 import com.oracle.graal.lir.amd64.AMD64ControlFlow.CondMoveOp;
@@ -162,9 +161,8 @@
             }
         }
         params[params.length - 1] = rbpParam;
-        ParametersOp paramsOp = new ParametersOp(params);
 
-        append(paramsOp);
+        emitIncomingValues(params);
 
         saveRbp = new SaveRbp(new PlaceholderOp(currentBlock, lir.lir(currentBlock).size()));
         append(saveRbp.placeholder);
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/CompilerToVM.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/CompilerToVM.java	Sat Sep 07 12:23:40 2013 -0700
@@ -226,4 +226,6 @@
     void reprofile(long metaspaceMethod);
 
     void invalidateInstalledCode(HotSpotInstalledCode hotspotInstalledCode);
+
+    boolean isTypeLinked(HotSpotResolvedObjectType hotSpotResolvedObjectType);
 }
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/CompilerToVMImpl.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/CompilerToVMImpl.java	Sat Sep 07 12:23:40 2013 -0700
@@ -102,6 +102,8 @@
     @Override
     public native boolean isTypeInitialized(HotSpotResolvedObjectType klass);
 
+    public native boolean isTypeLinked(HotSpotResolvedObjectType hotSpotResolvedObjectType);
+
     @Override
     public native boolean hasFinalizableSubclass(HotSpotResolvedObjectType klass);
 
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedJavaMethod.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedJavaMethod.java	Sat Sep 07 12:23:40 2013 -0700
@@ -138,7 +138,7 @@
         if (codeSize == 0) {
             return null;
         }
-        if (code == null) {
+        if (code == null && graalRuntime().getCompilerToVM().isTypeLinked(holder)) {
             code = graalRuntime().getCompilerToVM().initializeBytecode(metaspaceMethod, new byte[codeSize]);
             assert code.length == codeSize : "expected: " + codeSize + ", actual: " + code.length;
         }
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedObjectType.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedObjectType.java	Sat Sep 07 12:23:40 2013 -0700
@@ -83,6 +83,7 @@
     private ResolvedJavaType[] interfaces;
     private ConstantPool constantPool;
     private boolean isInitialized;
+    private boolean isLinked;
     private ResolvedJavaType arrayOfType;
 
     /**
@@ -292,9 +293,18 @@
     }
 
     @Override
+    public boolean isLinked() {
+        if (!isLinked) {
+            isLinked = graalRuntime().getCompilerToVM().isTypeLinked(this);
+        }
+        return isLinked;
+    }
+
+    @Override
     public void initialize() {
         if (!isInitialized) {
             graalRuntime().getCompilerToVM().initializeType(this);
+            assert graalRuntime().getCompilerToVM().isTypeInitialized(this);
         }
         isInitialized = true;
     }
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedPrimitiveType.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedPrimitiveType.java	Sat Sep 07 12:23:40 2013 -0700
@@ -112,6 +112,10 @@
         return true;
     }
 
+    public boolean isLinked() {
+        return true;
+    }
+
     @Override
     public boolean isInstance(Constant obj) {
         return false;
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/replacements/HotSpotReplacementsUtil.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/replacements/HotSpotReplacementsUtil.java	Sat Sep 07 12:23:40 2013 -0700
@@ -167,7 +167,7 @@
      */
     public static boolean clearPendingException(Word thread) {
         boolean result = thread.readObject(threadPendingExceptionOffset(), PENDING_EXCEPTION_LOCATION) != null;
-        thread.writeObject(threadPendingExceptionOffset(), null);
+        thread.writeObject(threadPendingExceptionOffset(), null, PENDING_EXCEPTION_LOCATION);
         return result;
     }
 
@@ -178,7 +178,7 @@
      */
     public static Object getAndClearObjectResult(Word thread) {
         Object result = thread.readObject(objectResultOffset(), OBJECT_RESULT_LOCATION);
-        thread.writeObject(objectResultOffset(), null);
+        thread.writeObject(objectResultOffset(), null, OBJECT_RESULT_LOCATION);
         return result;
     }
 
--- a/graal/com.oracle.graal.java/src/com/oracle/graal/java/GraphBuilderPhase.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.java/src/com/oracle/graal/java/GraphBuilderPhase.java	Sat Sep 07 12:23:40 2013 -0700
@@ -242,7 +242,7 @@
 
         // remove dead FrameStates
         for (Node n : currentGraph.getNodes(FrameState.class)) {
-            if (n.usages().count() == 0 && n.predecessor() == null) {
+            if (n.usages().isEmpty() && n.predecessor() == null) {
                 n.safeDelete();
             }
         }
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXArithmetic.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXArithmetic.java	Sat Sep 07 12:23:40 2013 -0700
@@ -316,7 +316,7 @@
                     masm.cvt_f32_f64(asFloatReg(dst), asDoubleReg(src));
                     break;
                 case LSHL:
-                    masm.shl_s64(asLongReg(dst), asLongReg(dst), asIntReg(src));
+                    masm.shl_b64(asLongReg(dst), asLongReg(dst), asIntReg(src));
                     break;
                 case LSHR:
                     masm.shr_s64(asLongReg(dst), asLongReg(dst), asIntReg(src));
@@ -360,19 +360,19 @@
             switch (opcode) {
             case IADD:  masm.add_s32(asIntReg(dst),    asIntReg(src1),    tasm.asIntConst(src2));    break;
             case ISUB:  masm.sub_s32(asIntReg(dst),    asIntReg(src1),    tasm.asIntConst(src2));    break;
-            case IMUL:  masm.mul_s32(asIntReg(dst),    asIntReg(src1),    tasm.asIntConst(src2));    break;
+            case IMUL:  masm.mul_lo_s32(asIntReg(dst), asIntReg(src1),    tasm.asIntConst(src2));    break;
             case IAND:  masm.and_b32(asIntReg(dst),    asIntReg(src1),    tasm.asIntConst(src2));    break;
-            case ISHL:  masm.shl_s32(asIntReg(dst),    asIntReg(src1),    tasm.asIntConst(src2));    break;
+            case ISHL:  masm.shl_b32_const(asIntReg(dst), asIntReg(src1), tasm.asIntConst(src2));    break;
             case ISHR:  masm.shr_s32(asIntReg(dst),    asIntReg(src1),    tasm.asIntConst(src2));    break;
             case IUSHR: masm.shr_u32(asIntReg(dst),    asIntReg(src1),    tasm.asIntConst(src2));    break;
             case IXOR:  masm.xor_b32(asIntReg(dst),    asIntReg(src1),    tasm.asIntConst(src2));    break;
             case LXOR:  masm.xor_b64(asLongReg(dst),   asLongReg(src1),   tasm.asLongConst(src2));   break;
             case LUSHR: masm.shr_u64(asLongReg(dst),   asLongReg(src1),   tasm.asLongConst(src2));   break;
             case FADD:  masm.add_f32(asFloatReg(dst),  asFloatReg(src1),  tasm.asFloatConst(src2));  break;
-            case FMUL:  masm.mul_f32(asFloatReg(dst),  asFloatReg(src1),  tasm.asFloatConst(src2));  break;
+            case FMUL:  masm.mul_lo_f32(asFloatReg(dst), asFloatReg(src1), tasm.asFloatConst(src2)); break;
             case FDIV:  masm.div_f32(asFloatReg(dst),  asFloatReg(src1),  tasm.asFloatConst(src2));  break;
             case DADD:  masm.add_f64(asDoubleReg(dst), asDoubleReg(src1), tasm.asDoubleConst(src2)); break;
-            case DMUL:  masm.mul_f64(asDoubleReg(dst), asDoubleReg(src1), tasm.asDoubleConst(src2)); break;
+            case DMUL:  masm.mul_lo_f64(asDoubleReg(dst), asDoubleReg(src1), tasm.asDoubleConst(src2)); break;
             case DDIV:  masm.div_f64(asDoubleReg(dst), asDoubleReg(src1), tasm.asDoubleConst(src2)); break;
             default:
                 throw GraalInternalError.shouldNotReachHere();
@@ -387,34 +387,34 @@
             // case D:  new Mul(Double, dst, src1, src2);
             case IADD:  masm.add_s32(asIntReg(dst),    asIntReg(src1),    asIntReg(src2));    break;
             case ISUB:  masm.sub_s32(asIntReg(dst),    asIntReg(src1),    asIntReg(src2));    break;
-            case IMUL:  masm.mul_s32(asIntReg(dst),    asIntReg(src1),    asIntReg(src2));    break;
+            case IMUL:  masm.mul_lo_s32(asIntReg(dst), asIntReg(src1),    asIntReg(src2));    break;
             case IDIV:  masm.div_s32(asIntReg(dst),    asIntReg(src1),    asIntReg(src2));    break;
             case IAND:  masm.and_b32(asIntReg(dst),    asIntReg(src1),    asIntReg(src2));    break;
             case IOR:    masm.or_b32(asIntReg(dst),    asIntReg(src1),    asIntReg(src2));    break;
             case IXOR:  masm.xor_b32(asIntReg(dst),    asIntReg(src1),    asIntReg(src2));    break;
-            case ISHL:  masm.shl_s32(asIntReg(dst),    asIntReg(src1),    asIntReg(src2));    break;
+            case ISHL:  masm.shl_b32(asIntReg(dst),    asIntReg(src1),    asIntReg(src2));    break;
             case ISHR:  masm.shr_s32(asIntReg(dst),    asIntReg(src1),    asIntReg(src2));    break;
             case IUSHR: masm.shr_u32(asIntReg(dst),    asIntReg(src1),    asIntReg(src2));    break;
             case IREM:  masm.rem_s32(asIntReg(dst),    asIntReg(src1),    asIntReg(src2));    break;
             case LADD:  masm.add_s64(asLongReg(dst),   asLongReg(src1),   asLongReg(src2));   break;
             case LSUB:  masm.sub_s64(asLongReg(dst),   asLongReg(src1),   asLongReg(src2));   break;
-            case LMUL:  masm.mul_s64(asLongReg(dst),   asLongReg(src1),   asLongReg(src2));   break;
+            case LMUL:  masm.mul_lo_s64(asLongReg(dst), asLongReg(src1),  asLongReg(src2));   break;
             case LDIV:  masm.div_s64(asLongReg(dst),   asLongReg(src1),   asLongReg(src2));   break;
             case LAND:  masm.and_b64(asLongReg(dst),   asLongReg(src1),   asLongReg(src2));   break;
-            case LOR:    masm.or_b64(asLongReg(dst),   asLongReg(src1),   asLongReg(src2));   break;
+            case LOR:   masm.or_b64(asLongReg(dst),    asLongReg(src1),   asLongReg(src2));   break;
             case LXOR:  masm.xor_b64(asLongReg(dst),   asLongReg(src1),   asLongReg(src2));   break;
-            case LSHL:  masm.shl_s64(asLongReg(dst),   asLongReg(src1),   asLongReg(src2));   break;
+            case LSHL:  masm.shl_b64(asLongReg(dst),   asLongReg(src1),   asLongReg(src2));   break;
             case LSHR:  masm.shr_s64(asLongReg(dst),   asLongReg(src1),   asLongReg(src2));   break;
             case LUSHR: masm.shr_u64(asLongReg(dst),   asLongReg(src1),   asIntReg(src2));    break;
             case LREM:  masm.rem_s64(asLongReg(dst),   asLongReg(src1),   asLongReg(src2));   break;
             case FADD:  masm.add_f32(asFloatReg(dst),  asFloatReg(src1),  asFloatReg(src2));  break;
             case FSUB:  masm.sub_f32(asFloatReg(dst),  asFloatReg(src1),  asFloatReg(src2));  break;
-            case FMUL:  masm.mul_f32(asFloatReg(dst),  asFloatReg(src1),  asFloatReg(src2));  break;
+            case FMUL:  masm.mul_lo_f32(asFloatReg(dst), asFloatReg(src1), asFloatReg(src2)); break;
             case FDIV:  masm.div_f32(asFloatReg(dst),  asFloatReg(src1),  asFloatReg(src2));  break;
             case FREM:  masm.div_f32(asFloatReg(dst),  asFloatReg(src1),  asFloatReg(src2));  break;
             case DADD:  masm.add_f64(asDoubleReg(dst), asDoubleReg(src1), asDoubleReg(src2)); break;
             case DSUB:  masm.sub_f64(asDoubleReg(dst), asDoubleReg(src1), asDoubleReg(src2)); break;
-            case DMUL:  masm.mul_f64(asDoubleReg(dst), asDoubleReg(src1), asDoubleReg(src2)); break;
+            case DMUL:  masm.mul_lo_f64(asDoubleReg(dst), asDoubleReg(src1), asDoubleReg(src2)); break;
             case DDIV:  masm.div_f64(asDoubleReg(dst), asDoubleReg(src1), asDoubleReg(src2)); break;
             case DREM:  masm.div_f64(asDoubleReg(dst), asDoubleReg(src1), asDoubleReg(src2)); break;
             default:
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMemOp.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMemOp.java	Sat Sep 07 12:23:40 2013 -0700
@@ -158,7 +158,7 @@
                     masm.ld_from_state_space(".param.s16", asRegister(result), addr.getBase(), addr.getDisplacement());
                     break;
                 case Char:
-                    masm.ld_from_state_space(".param.s16", asRegister(result), addr.getBase(), addr.getDisplacement());
+                    masm.ld_from_state_space(".param.u16", asRegister(result), addr.getBase(), addr.getDisplacement());
                     break;
                 case Int:
                     masm.ld_from_state_space(".param.s32", asRegister(result), addr.getBase(), addr.getDisplacement());
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXParameterOp.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXParameterOp.java	Sat Sep 07 12:23:40 2013 -0700
@@ -42,23 +42,23 @@
 
     @Override
     public void emitCode(TargetMethodAssembler tasm) {
-        PTXAssembler ptxasm = (PTXAssembler) tasm.asm;
+        PTXAssembler masm = (PTXAssembler) tasm.asm;
         // Emit parameter directives for arguments
         int argCount = params.length;
         for (int i = 0; i < argCount; i++) {
             Kind paramKind = params[i].getKind();
             switch (paramKind) {
             case Int :
-                ptxasm.param_32_decl(asIntReg(params[i]), (i == (argCount - 1)));
+                masm.param_32_decl(asIntReg(params[i]), (i == (argCount - 1)));
                 break;
             case Long :
-                ptxasm.param_64_decl(asLongReg(params[i]), (i == (argCount - 1)));
+                masm.param_64_decl(asLongReg(params[i]), (i == (argCount - 1)));
                 break;
             case Float :
-                ptxasm.param_32_decl(asFloatReg(params[i]), (i == (argCount - 1)));
+                masm.param_32_decl(asFloatReg(params[i]), (i == (argCount - 1)));
                 break;
             case Double :
-                ptxasm.param_64_decl(asDoubleReg(params[i]), (i == (argCount - 1)));
+                masm.param_64_decl(asDoubleReg(params[i]), (i == (argCount - 1)));
                 break;
             default :
                 throw GraalInternalError.shouldNotReachHere("unhandled parameter type "  + paramKind.toString());
--- a/graal/com.oracle.graal.lir/src/com/oracle/graal/lir/StandardOp.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.lir/src/com/oracle/graal/lir/StandardOp.java	Sat Sep 07 12:23:40 2013 -0700
@@ -54,12 +54,30 @@
      */
     public static class LabelOp extends LIRInstruction {
 
+        private static final Value[] NO_VALUES = new Value[0];
+
+        /**
+         * In the LIR, every register and variable must be defined before it is used. For method
+         * parameters that are passed in fixed registers, exception objects passed to the exception
+         * handler in a fixed register, or any other use of a fixed register not defined in this
+         * method, an artificial definition is necessary. To avoid spill moves to be inserted
+         * between the label at the beginning of a block an an actual definition in the second
+         * instruction of a block, the registers are defined here in the label.
+         */
+        @Def({REG, STACK}) private Value[] incomingValues;
+
         private final Label label;
         private final boolean align;
 
         public LabelOp(Label label, boolean align) {
             this.label = label;
             this.align = align;
+            this.incomingValues = NO_VALUES;
+        }
+
+        public void setIncomingValues(Value[] values) {
+            assert incomingValues.length == 0;
+            incomingValues = values;
         }
 
         @Override
@@ -121,25 +139,6 @@
     }
 
     /**
-     * Meta-operation that defines the incoming method parameters. In the LIR, every register and
-     * variable must be defined before it is used. This operation is the definition point of method
-     * parameters, but is otherwise a no-op. In particular, it is not the actual method prologue.
-     */
-    public static final class ParametersOp extends LIRInstruction {
-
-        @Def({REG, STACK}) protected Value[] params;
-
-        public ParametersOp(Value[] params) {
-            this.params = params;
-        }
-
-        @Override
-        public void emitCode(TargetMethodAssembler tasm) {
-            // No code to emit.
-        }
-    }
-
-    /**
      * Placeholder for a LIR instruction that will be subsequently replaced.
      */
     public static class PlaceholderOp extends LIRInstruction {
--- a/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/LoopFragment.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/LoopFragment.java	Sat Sep 07 12:23:40 2013 -0700
@@ -145,11 +145,11 @@
         }
     }
 
-    protected static NodeBitMap computeNodes(Graph graph, Collection<AbstractBeginNode> blocks) {
+    protected static NodeBitMap computeNodes(Graph graph, Iterable<AbstractBeginNode> blocks) {
         return computeNodes(graph, blocks, Collections.<AbstractBeginNode> emptyList());
     }
 
-    protected static NodeBitMap computeNodes(Graph graph, Collection<AbstractBeginNode> blocks, Collection<AbstractBeginNode> earlyExits) {
+    protected static NodeBitMap computeNodes(Graph graph, Iterable<AbstractBeginNode> blocks, Iterable<AbstractBeginNode> earlyExits) {
         final NodeBitMap nodes = graph.createNodeBitMap(true);
         for (AbstractBeginNode b : blocks) {
             for (Node n : b.getBlockNodes()) {
@@ -229,12 +229,28 @@
         return false;
     }
 
-    public static Collection<AbstractBeginNode> toHirBlocks(Collection<Block> blocks) {
-        List<AbstractBeginNode> hir = new ArrayList<>(blocks.size());
-        for (Block b : blocks) {
-            hir.add(b.getBeginNode());
-        }
-        return hir;
+    public static NodeIterable<AbstractBeginNode> toHirBlocks(final Iterable<Block> blocks) {
+        return new AbstractNodeIterable<AbstractBeginNode>() {
+
+            public Iterator<AbstractBeginNode> iterator() {
+                final Iterator<Block> it = blocks.iterator();
+                return new Iterator<AbstractBeginNode>() {
+
+                    public void remove() {
+                        throw new UnsupportedOperationException();
+                    }
+
+                    public AbstractBeginNode next() {
+                        return it.next().getBeginNode();
+                    }
+
+                    public boolean hasNext() {
+                        return it.hasNext();
+                    }
+                };
+            }
+
+        };
     }
 
     /**
--- a/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/LoopFragmentInside.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/LoopFragmentInside.java	Sat Sep 07 12:23:40 2013 -0700
@@ -231,7 +231,7 @@
         StructuredGraph graph = graph();
         if (endsToMerge.size() == 1) {
             AbstractEndNode end = endsToMerge.get(0);
-            assert end.usages().count() == 0;
+            assert end.usages().isEmpty();
             newExit = graph.add(new BeginNode());
             end.replaceAtPredecessor(newExit);
             end.safeDelete();
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/LoopEndNode.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/LoopEndNode.java	Sat Sep 07 12:23:40 2013 -0700
@@ -72,7 +72,7 @@
     @Override
     public boolean verify() {
         assertTrue(loopBegin != null, "must have a loop begin");
-        assertTrue(usages().count() == 0, "LoopEnds can not be used");
+        assertTrue(usages().isEmpty(), "LoopEnds can not be used");
         return super.verify();
     }
 
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/extended/ReadNode.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/extended/ReadNode.java	Sat Sep 07 12:23:40 2013 -0700
@@ -72,7 +72,7 @@
 
     public static ValueNode canonicalizeRead(ValueNode read, LocationNode location, ValueNode object, CanonicalizerTool tool, boolean compressible) {
         MetaAccessProvider runtime = tool.runtime();
-        if (read.usages().count() == 0) {
+        if (read.usages().isEmpty()) {
             // Read without usages can be savely removed.
             return null;
         }
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/java/InstanceOfDynamicNode.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/java/InstanceOfDynamicNode.java	Sat Sep 07 12:23:40 2013 -0700
@@ -23,9 +23,7 @@
 package com.oracle.graal.nodes.java;
 
 import com.oracle.graal.api.meta.*;
-import com.oracle.graal.graph.*;
 import com.oracle.graal.nodes.*;
-import com.oracle.graal.nodes.calc.*;
 import com.oracle.graal.nodes.spi.*;
 import com.oracle.graal.nodes.type.*;
 
@@ -74,12 +72,4 @@
     public ValueNode mirror() {
         return mirror;
     }
-
-    @Override
-    public boolean verify() {
-        for (Node usage : usages()) {
-            assertTrue(usage instanceof IfNode || usage instanceof FixedGuardNode || usage instanceof ConditionalNode, "unsupported usage: %s", usage);
-        }
-        return super.verify();
-    }
 }
--- a/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/InliningPhase.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/InliningPhase.java	Sat Sep 07 12:23:40 2013 -0700
@@ -512,7 +512,7 @@
             while ((current = nextQueuedNode()) != null) {
                 assert current.isAlive();
 
-                if (current instanceof Invoke) {
+                if (current instanceof Invoke && ((Invoke) current).callTarget() instanceof MethodCallTargetNode) {
                     if (current != start) {
                         invokes.addLast((Invoke) current);
                     }
--- a/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/RemoveValueProxyPhase.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/RemoveValueProxyPhase.java	Sat Sep 07 12:23:40 2013 -0700
@@ -39,7 +39,7 @@
             FrameState stateAfter = exit.stateAfter();
             if (stateAfter != null) {
                 exit.setStateAfter(null);
-                if (stateAfter.usages().count() == 0) {
+                if (stateAfter.usages().isEmpty()) {
                     stateAfter.safeDelete();
                 }
             }
--- a/graal/com.oracle.graal.replacements/src/com/oracle/graal/replacements/SnippetTemplate.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.graal.replacements/src/com/oracle/graal/replacements/SnippetTemplate.java	Sat Sep 07 12:23:40 2013 -0700
@@ -528,8 +528,6 @@
         ReturnNode retNode = null;
         StartNode entryPointNode = snippet.start();
 
-        new DeadCodeEliminationPhase().apply(snippetCopy);
-
         nodes = new ArrayList<>(snippet.getNodeCount());
         for (Node node : snippet.getNodes()) {
             if (node == entryPointNode || node == entryPointNode.stateAfter()) {
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.truffle.api.dsl.test/src/com/oracle/truffle/api/dsl/test/ImplicitCastTest.java	Sat Sep 07 12:23:40 2013 -0700
@@ -0,0 +1,84 @@
+/*
+ * Copyright (c) 2012, 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.truffle.api.dsl.test;
+
+import org.junit.*;
+
+import com.oracle.truffle.api.dsl.*;
+import com.oracle.truffle.api.dsl.test.ImplicitCastTestFactory.ImplicitCast0NodeFactory;
+import com.oracle.truffle.api.dsl.test.NodeContainerTest.Str;
+import com.oracle.truffle.api.dsl.test.TypeSystemTest.TestRootNode;
+import com.oracle.truffle.api.dsl.test.TypeSystemTest.ValueNode;
+import com.oracle.truffle.api.frame.*;
+
+public class ImplicitCastTest {
+
+    @TypeSystem({int.class, boolean.class, String.class, Str.class})
+    static class ImplicitCast0Types {
+
+        @ImplicitCast
+        boolean castInt(int intvalue) {
+            return intvalue == 1 ? true : false;
+        }
+
+        @ImplicitCast
+        boolean castString(String strvalue) {
+            return strvalue.equals("1");
+        }
+
+    }
+
+    @TypeSystemReference(ImplicitCast0Types.class)
+    @NodeChild(value = "operand", type = ImplicitCast0Node.class)
+    abstract static class ImplicitCast0Node extends ValueNode {
+
+        public abstract Object executeEvaluated(VirtualFrame frame, Object value2);
+
+        @Specialization(order = 1)
+        public String op1(String value) throws RuntimeException {
+            return value;
+        }
+
+        @Specialization(order = 2)
+        public boolean op1(boolean value) throws RuntimeException {
+            return value;
+        }
+
+    }
+
+    @Test
+    public void testImplicitCast0() {
+        ImplicitCast0Node node = ImplicitCast0NodeFactory.create(null);
+        TestRootNode<ImplicitCast0Node> root = new TestRootNode<>(node);
+        Assert.assertEquals("2", root.getNode().executeEvaluated(null, "2"));
+        Assert.assertEquals(true, root.getNode().executeEvaluated(null, 1));
+        Assert.assertEquals("1", root.getNode().executeEvaluated(null, "1"));
+        Assert.assertEquals(true, root.getNode().executeEvaluated(null, 1));
+        Assert.assertEquals(true, root.getNode().executeEvaluated(null, true));
+    }
+
+    // TODO assert implicit casts only in one direction
+
+    // test example that covers the most cases
+
+}
--- a/graal/com.oracle.truffle.api.dsl.test/src/com/oracle/truffle/api/dsl/test/PolymorphicTest.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.truffle.api.dsl.test/src/com/oracle/truffle/api/dsl/test/PolymorphicTest.java	Sat Sep 07 12:23:40 2013 -0700
@@ -37,6 +37,17 @@
 
 public class PolymorphicTest {
 
+    private static void assertParent(Node expectedParent, Node child) {
+        Node parent = child.getParent();
+        while (parent != null && parent != expectedParent) {
+            parent = parent.getParent();
+        }
+
+        if (parent != expectedParent) {
+            assertEquals(expectedParent, parent);
+        }
+    }
+
     @Test
     public void testJustSpecialize() {
         TestRootNode<Node1> node = TestHelper.createRoot(Node1Factory.getInstance());
@@ -45,8 +56,8 @@
         assertEquals("(int,boolean)", executeWith(node, 42, false));
         assertEquals("(boolean,int)", executeWith(node, false, 42));
         assertEquals(Kind.SPECIALIZED, node.getNode().getClass().getAnnotation(NodeInfo.class).kind());
-        assertEquals(node.getNode(), node.getNode().getLeft().getParent());
-        assertEquals(node.getNode(), node.getNode().getRight().getParent());
+        assertParent(node.getNode(), node.getNode().getLeft());
+        assertParent(node.getNode(), node.getNode().getRight());
     }
 
     @Test
@@ -55,8 +66,8 @@
         assertEquals("(int,boolean)", executeWith(node, 42, false));
         assertEquals("(int,int)", executeWith(node, 42, 42));
         assertEquals(Kind.POLYMORPHIC, node.getNode().getClass().getAnnotation(NodeInfo.class).kind());
-        assertEquals(node.getNode(), node.getNode().getLeft().getParent());
-        assertEquals(node.getNode(), node.getNode().getRight().getParent());
+        assertParent(node.getNode(), node.getNode().getLeft());
+        assertParent(node.getNode(), node.getNode().getRight());
     }
 
     @Test
@@ -66,8 +77,8 @@
         assertEquals("(boolean,boolean)", executeWith(node, true, false));
         assertEquals("(int,int)", executeWith(node, 42, 42));
         assertEquals(Kind.POLYMORPHIC, node.getNode().getClass().getAnnotation(NodeInfo.class).kind());
-        assertEquals(node.getNode(), node.getNode().getLeft().getParent());
-        assertEquals(node.getNode(), node.getNode().getRight().getParent());
+        assertParent(node.getNode(), node.getNode().getLeft());
+        assertParent(node.getNode(), node.getNode().getRight());
     }
 
     @Test
@@ -78,17 +89,17 @@
         assertEquals("(boolean,boolean)", executeWith(node, true, false));
         assertEquals("(int,int)", executeWith(node, 42, 42));
         assertEquals(Kind.GENERIC, node.getNode().getClass().getAnnotation(NodeInfo.class).kind());
-        assertEquals(node.getNode(), node.getNode().getLeft().getParent());
-        assertEquals(node.getNode(), node.getNode().getRight().getParent());
+        assertParent(node.getNode(), node.getNode().getLeft());
+        assertParent(node.getNode(), node.getNode().getRight());
     }
 
     @Test
     public void testGenericInitial() {
         TestRootNode<Node1> node = TestHelper.createRoot(Node1Factory.getInstance());
-        assertEquals("(generic,generic)", executeWith(node, "", ""));
+        assertEquals("(generic,generic)", executeWith(node, "1", "1"));
         assertEquals(Kind.GENERIC, node.getNode().getClass().getAnnotation(NodeInfo.class).kind());
-        assertEquals(node.getNode(), node.getNode().getLeft().getParent());
-        assertEquals(node.getNode(), node.getNode().getRight().getParent());
+        assertParent(node.getNode(), node.getNode().getLeft());
+        assertParent(node.getNode(), node.getNode().getRight());
     }
 
     @Test
@@ -99,8 +110,8 @@
         assertEquals("(generic,generic)", executeWith(node, "", ""));
         assertEquals(Kind.GENERIC, node.getNode().getClass().getAnnotation(NodeInfo.class).kind());
         /* Assertions for bug GRAAL-425 */
-        assertEquals(node.getNode(), node.getNode().getLeft().getParent());
-        assertEquals(node.getNode(), node.getNode().getRight().getParent());
+        assertParent(node.getNode(), node.getNode().getLeft());
+        assertParent(node.getNode(), node.getNode().getRight());
     }
 
     @SuppressWarnings("unused")
--- a/graal/com.oracle.truffle.api.dsl.test/src/com/oracle/truffle/api/dsl/test/TypeSystemTest.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.truffle.api.dsl.test/src/com/oracle/truffle/api/dsl/test/TypeSystemTest.java	Sat Sep 07 12:23:40 2013 -0700
@@ -48,6 +48,11 @@
             return (int) value;
         }
 
+        @ImplicitCast
+        Str castStr(String strvalue) {
+            return new Str(strvalue);
+        }
+
     }
 
     @TypeSystemReference(SimpleTypes.class)
--- a/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/nodes/Node.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/nodes/Node.java	Sat Sep 07 12:23:40 2013 -0700
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2012, Oracle and/or its affiliates. All rights reserved.
+ * Copyright (c) 2012, 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
@@ -59,6 +59,11 @@
         CompilerAsserts.neverPartOfCompilation();
     }
 
+    protected Node(SourceSection sourceSection) {
+        CompilerAsserts.neverPartOfCompilation();
+        this.sourceSection = sourceSection;
+    }
+
     /**
      * Assigns a link to a guest language source section to this node.
      * 
--- a/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/nodes/RootNode.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/nodes/RootNode.java	Sat Sep 07 12:23:40 2013 -0700
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2012, Oracle and/or its affiliates. All rights reserved.
+ * Copyright (c) 2012, 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
@@ -34,6 +34,13 @@
  */
 public abstract class RootNode extends Node {
 
+    protected RootNode() {
+    }
+
+    protected RootNode(SourceSection sourceSection) {
+        super(sourceSection);
+    }
+
     /**
      * Executes this function using the specified frame and returns the result value.
      * 
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/Utils.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/Utils.java	Sat Sep 07 12:23:40 2013 -0700
@@ -192,6 +192,10 @@
             return true;
         }
 
+        if (isObject(to)) {
+            return true;
+        }
+
         // JLS 5.1.2 widening primitives
         if (Utils.isPrimitive(from) && Utils.isPrimitive(to)) {
             TypeKind fromKind = from.getKind();
@@ -257,6 +261,10 @@
             return isAssignable(context, ((ArrayType) from).getComponentType(), ((ArrayType) to).getComponentType());
         }
 
+        if (from instanceof ArrayType || to instanceof ArrayType) {
+            return false;
+        }
+
         TypeElement fromType = Utils.fromTypeMirror(from);
         TypeElement toType = Utils.fromTypeMirror(to);
         if (fromType == null || toType == null) {
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/ast/CodeTreeBuilder.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/ast/CodeTreeBuilder.java	Sat Sep 07 12:23:40 2013 -0700
@@ -165,14 +165,18 @@
     }
 
     public CodeTreeBuilder startCall(String callSite) {
-        return startCall(null, callSite);
+        return startCall((CodeTree) null, callSite);
     }
 
     public CodeTreeBuilder startCall(String receiver, String callSite) {
+        return startCall(singleString(receiver), callSite);
+    }
+
+    public CodeTreeBuilder startCall(CodeTree receiver, String callSite) {
         if (receiver == null) {
             return startGroup().string(callSite).startParanthesesCommaGroup().endAfter();
         } else {
-            return startGroup().string(receiver).string(".").string(callSite).startParanthesesCommaGroup().endAfter();
+            return startGroup().tree(receiver).string(".").string(callSite).startParanthesesCommaGroup().endAfter();
         }
     }
 
@@ -336,6 +340,10 @@
         return startGroup().string("if ").startParanthesesCommaGroup().endAndWhitespaceAfter().startGroup().endAfter();
     }
 
+    public CodeTreeBuilder startFor() {
+        return startGroup().string("for ").startParantheses().endAndWhitespaceAfter().startGroup().endAfter();
+    }
+
     public boolean startIf(boolean elseIf) {
         if (elseIf) {
             startElseIf();
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/NodeChildData.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/NodeChildData.java	Sat Sep 07 12:23:40 2013 -0700
@@ -49,6 +49,7 @@
         DEFAULT, SHORT_CIRCUIT
     }
 
+    private final NodeData parent;
     private final Element sourceElement;
     private final AnnotationMirror sourceAnnotationMirror;
 
@@ -64,8 +65,9 @@
 
     private NodeData nodeData;
 
-    public NodeChildData(Element sourceElement, AnnotationMirror sourceMirror, String name, TypeMirror nodeType, TypeMirror originalNodeType, Element accessElement, Cardinality cardinality,
-                    ExecutionKind executionKind) {
+    public NodeChildData(NodeData parent, Element sourceElement, AnnotationMirror sourceMirror, String name, TypeMirror nodeType, TypeMirror originalNodeType, Element accessElement,
+                    Cardinality cardinality, ExecutionKind executionKind) {
+        this.parent = parent;
         this.sourceElement = sourceElement;
         this.sourceAnnotationMirror = sourceMirror;
         this.name = name;
@@ -84,6 +86,31 @@
         this.executeWith = executeWith;
     }
 
+    public boolean needsImplicitCast(ProcessorContext context) {
+        if (!parent.needsRewrites(context)) {
+            return false;
+        }
+
+        boolean used = false;
+        SpecializationData generic = parent.getGenericSpecialization();
+        for (ActualParameter param : generic.getParameters()) {
+            if (!param.getSpecification().isSignature()) {
+                continue;
+            }
+            NodeChildData child = parent.findChild(param.getSpecification().getName());
+            if (child == this) {
+                used = true;
+                break;
+            }
+        }
+
+        if (!used) {
+            return false;
+        }
+
+        return getNodeData().getTypeSystem().getImplicitCasts() != null && !getNodeData().getTypeSystem().getImplicitCasts().isEmpty();
+    }
+
     public ExecutableTypeData findExecutableType(ProcessorContext context, TypeData targetType) {
         ExecutableTypeData executableType = nodeData.findExecutableType(targetType, getExecuteWith().size());
         if (executableType == null) {
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/NodeCodeGenerator.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/NodeCodeGenerator.java	Sat Sep 07 12:23:40 2013 -0700
@@ -62,12 +62,17 @@
         return node.getNodeId() + "Factory";
     }
 
+    private static String nodeCastClassName(NodeData node, TypeData type) {
+        String nodeid = resolveNodeId(node);
+        if (type == null) {
+            return nodeid + "ImplicitCast";
+        } else {
+            return Utils.firstLetterUpperCase(Utils.getSimpleName(type.getPrimitiveType())) + "Cast";
+        }
+    }
+
     private static String nodeSpecializationClassName(SpecializationData specialization) {
-        String nodeid = specialization.getNode().getNodeId();
-        if (nodeid.endsWith("Node") && !nodeid.equals("Node")) {
-            nodeid = nodeid.substring(0, nodeid.length() - 4);
-        }
-
+        String nodeid = resolveNodeId(specialization.getNode());
         String name = Utils.firstLetterUpperCase(nodeid);
         name += Utils.firstLetterUpperCase(specialization.getId());
         name += "Node";
@@ -75,10 +80,7 @@
     }
 
     private static String nodePolymorphicClassName(NodeData node, SpecializationData specialization) {
-        String nodeid = node.getNodeId();
-        if (nodeid.endsWith("Node") && !nodeid.equals("Node")) {
-            nodeid = nodeid.substring(0, nodeid.length() - 4);
-        }
+        String nodeid = resolveNodeId(node);
 
         String name = Utils.firstLetterUpperCase(nodeid);
         if (specialization == node.getGenericPolymorphicSpecialization()) {
@@ -89,10 +91,22 @@
         return name;
     }
 
+    private static String resolveNodeId(NodeData node) {
+        String nodeid = node.getNodeId();
+        if (nodeid.endsWith("Node") && !nodeid.equals("Node")) {
+            nodeid = nodeid.substring(0, nodeid.length() - 4);
+        }
+        return nodeid;
+    }
+
     private static String valueNameEvaluated(ActualParameter targetParameter) {
         return valueName(targetParameter) + "Evaluated";
     }
 
+    private static String typeName(ActualParameter param) {
+        return param.getLocalName() + "Type";
+    }
+
     private static String valueName(ActualParameter param) {
         return param.getLocalName();
     }
@@ -147,7 +161,8 @@
         }
     }
 
-    private void addInternalValueParameterNames(CodeTreeBuilder builder, TemplateMethod source, TemplateMethod specialization, String unexpectedValueName, boolean forceFrame, boolean includeImplicit) {
+    private void addInternalValueParameterNames(CodeTreeBuilder builder, TemplateMethod source, TemplateMethod specialization, String unexpectedValueName, boolean forceFrame, boolean includeImplicit,
+                    Map<String, String> customNames) {
         if (forceFrame && specialization.getSpecification().findParameterSpec("frame") != null) {
             builder.string("frameValue");
         }
@@ -166,7 +181,9 @@
 
             ActualParameter sourceParameter = source.findParameter(parameter.getLocalName());
 
-            if (unexpectedValueName != null && parameter.getLocalName().equals(unexpectedValueName)) {
+            if (customNames != null && customNames.containsKey(parameter.getLocalName())) {
+                builder.string(customNames.get(parameter.getLocalName()));
+            } else if (unexpectedValueName != null && parameter.getLocalName().equals(unexpectedValueName)) {
                 builder.cast(parameter.getType(), CodeTreeBuilder.singleString("ex.getResult()"));
             } else if (sourceParameter != null) {
                 builder.string(valueName(sourceParameter, parameter));
@@ -298,10 +315,7 @@
     }
 
     private static String baseClassName(NodeData node) {
-        String nodeid = node.getNodeId();
-        if (nodeid.endsWith("Node") && !nodeid.equals("Node")) {
-            nodeid = nodeid.substring(0, nodeid.length() - 4);
-        }
+        String nodeid = resolveNodeId(node);
         String name = Utils.firstLetterUpperCase(nodeid);
         name += "BaseNode";
         return name;
@@ -362,7 +376,7 @@
         builder.startThrow().startNew(getContext().getType(UnsupportedOperationException.class));
         builder.startCall("createInfo0");
         builder.doubleQuote("Unsupported values");
-        addInternalValueParameterNames(builder, current, current, null, false, true);
+        addInternalValueParameterNames(builder, current, current, null, false, true, null);
         builder.end().end().end();
     }
 
@@ -422,7 +436,11 @@
     }
 
     @Override
+    @SuppressWarnings("unchecked")
     protected void createChildren(NodeData node) {
+        List<CodeTypeElement> casts = new ArrayList<>(getElement().getEnclosedElements());
+        getElement().getEnclosedElements().clear();
+
         Map<NodeData, List<TypeElement>> childTypes = new LinkedHashMap<>();
         if (node.getDeclaredNodes() != null && !node.getDeclaredNodes().isEmpty()) {
             for (NodeData nodeChild : node.getDeclaredNodes()) {
@@ -432,14 +450,35 @@
         }
 
         if (node.needsFactory() || node.getNodeDeclaringChildren().size() > 0) {
-            add(new NodeFactoryFactory(context, childTypes), node);
+            NodeFactoryFactory factory = new NodeFactoryFactory(context, childTypes);
+            add(factory, node);
+            factory.getElement().getEnclosedElements().addAll(casts);
+        }
+    }
+
+    protected CodeTree createCastType(NodeData node, TypeData sourceType, TypeData targetType, boolean expect, CodeTree value) {
+        if (targetType == null) {
+            return value;
+        } else if (sourceType != null && !sourceType.needsCastTo(getContext(), targetType)) {
+            return value;
         }
+
+        CodeTreeBuilder builder = CodeTreeBuilder.createBuilder();
+        String targetMethodName;
+        if (expect) {
+            targetMethodName = TypeSystemCodeGenerator.expectTypeMethodName(targetType);
+        } else {
+            targetMethodName = TypeSystemCodeGenerator.asTypeMethodName(targetType);
+        }
+        startCallTypeSystemMethod(getContext(), builder, node, targetMethodName);
+        builder.tree(value);
+        builder.end().end();
+        return builder.getRoot();
     }
 
     private class NodeFactoryFactory extends ClassElementFactory<NodeData> {
 
         private final Map<NodeData, List<TypeElement>> childTypes;
-
         private CodeTypeElement generatedNode;
 
         public NodeFactoryFactory(ProcessorContext context, Map<NodeData, List<TypeElement>> childElements) {
@@ -853,7 +892,8 @@
                     ExecutableElement getter = (ExecutableElement) child.getAccessElement();
                     CodeExecutableElement method = CodeExecutableElement.clone(getContext().getEnvironment(), getter);
                     method.getModifiers().remove(Modifier.ABSTRACT);
-                    method.createBuilder().startReturn().string("this.").string(child.getName()).end();
+                    CodeTreeBuilder builder = method.createBuilder();
+                    builder.startReturn().string("this.").string(child.getName()).end();
                     clazz.add(method);
                 }
             }
@@ -1107,29 +1147,58 @@
                 builder.startStatement();
                 String fieldName = var.getSimpleName().toString();
 
-                CodeTree fieldInit = CodeTreeBuilder.singleString(var.getSimpleName().toString());
-                builder.string("this.").string(var.getSimpleName().toString());
-
                 NodeChildData child = node.findChild(fieldName);
-                if (child != null) {
-                    CreateCastData createCast = node.findCast(child.getName());
-                    if (createCast != null) {
-                        fieldInit = createTemplateMethodCall(builder, null, node.getGenericSpecialization(), createCast, null, child.getName());
-                    }
-                }
-
-                if (Utils.isAssignable(getContext(), var.asType(), getContext().getTruffleTypes().getNode())) {
-                    builder.string(" = adoptChild(").tree(fieldInit).string(")");
-                } else if (Utils.isAssignable(getContext(), var.asType(), getContext().getTruffleTypes().getNodeArray())) {
-                    builder.string(" = adoptChildren(").tree(fieldInit).string(")");
-                } else {
-                    builder.string(" = ").tree(fieldInit);
-                }
+
+                CodeTree init = createStaticCast(builder, child, fieldName);
+                init = createAdoptChild(builder, var.asType(), init);
+
+                builder.string("this.").string(fieldName).string(" = ").tree(init);
                 builder.end();
             }
             return method;
         }
 
+        private CodeTree createStaticCast(CodeTreeBuilder parent, NodeChildData child, String fieldName) {
+            NodeData parentNode = getModel().getNode();
+            if (child != null) {
+                CreateCastData createCast = parentNode.findCast(child.getName());
+                if (createCast != null) {
+                    return createTemplateMethodCall(parent, null, parentNode.getGenericSpecialization(), createCast, null, fieldName);
+                }
+            }
+            return CodeTreeBuilder.singleString(fieldName);
+        }
+
+        private CodeTree createAdoptChild(CodeTreeBuilder parent, TypeMirror type, CodeTree value) {
+            CodeTreeBuilder builder = new CodeTreeBuilder(parent);
+            if (Utils.isAssignable(getContext(), type, getContext().getTruffleTypes().getNode())) {
+                builder.string("adoptChild(").tree(value).string(")");
+            } else if (Utils.isAssignable(getContext(), type, getContext().getTruffleTypes().getNodeArray())) {
+                builder.string("adoptChildren(").tree(value).string(")");
+            } else {
+                builder.tree(value);
+            }
+            return builder.getRoot();
+        }
+
+        private CodeTree createCopyArray(CodeTreeBuilder parent, NodeChildData child, TypeMirror arrayType, CodeBlock<String> accessElement) {
+            CodeTreeBuilder builder = parent.create();
+            NodeData node = getModel().getNode();
+            builder.string("new ").type(arrayType).string(" {");
+            builder.startCommaGroup();
+            for (ActualParameter parameter : getModel().getParameters()) {
+                NodeChildData foundChild = node.findChild(parameter.getSpecification().getName());
+                if (foundChild == child) {
+                    builder.startGroup();
+                    builder.tree(accessElement.create(builder, String.valueOf(parameter.getIndex())));
+                    builder.end();
+                }
+            }
+            builder.end();
+            builder.end().string("}");
+            return builder.getRoot();
+        }
+
         private CodeExecutableElement createCopyConstructor(CodeTypeElement type, ExecutableElement superConstructor) {
             CodeExecutableElement method = new CodeExecutableElement(null, type.getSimpleName().toString());
             CodeTreeBuilder builder = method.createBuilder();
@@ -1141,34 +1210,22 @@
 
             for (VariableElement var : type.getFields()) {
                 builder.startStatement();
-                String varName = var.getSimpleName().toString();
-                builder.string("this.").string(varName);
-                if (Utils.isAssignable(getContext(), var.asType(), getContext().getTruffleTypes().getNode())) {
-                    builder.string(" = adoptChild(copy.").string(varName).string(")");
-                } else if (Utils.isAssignable(getContext(), var.asType(), getContext().getTruffleTypes().getNodeArray())) {
-                    NodeData node = getModel().getNode();
-                    NodeChildData child = node.findChild(varName);
-                    if (child != null) {
-                        builder.string(" = adoptChildren(");
-                        builder.string("new ").type((child.getNodeType())).string(" {");
-                        builder.startCommaGroup();
-                        for (ActualParameter parameter : getModel().getParameters()) {
-                            NodeChildData foundChild = node.findChild(parameter.getSpecification().getName());
-                            if (foundChild == child) {
-                                builder.startGroup();
-                                builder.string("copy.").string(varName).string("[").string(String.valueOf(parameter.getIndex())).string("]");
-                                builder.end();
-                            }
+                final String varName = var.getSimpleName().toString();
+                final TypeMirror varType = var.asType();
+
+                final String copyAccess = "copy." + varName;
+                CodeTree init = CodeTreeBuilder.singleString(copyAccess);
+                if (Utils.isAssignable(getContext(), var.asType(), getContext().getTruffleTypes().getNodeArray())) {
+                    NodeChildData child = getModel().getNode().findChild(varName);
+                    init = createCopyArray(builder, child, varType, new CodeBlock<String>() {
+
+                        public CodeTree create(CodeTreeBuilder parent, String index) {
+                            return CodeTreeBuilder.singleString(copyAccess + "[" + index + "]");
                         }
-
-                        builder.end().string("})");
-                    } else {
-                        builder.string(" = adoptChildren(copy.").string(varName).string(")");
-                    }
-                } else {
-                    builder.string(" = copy.").string(varName);
+                    });
                 }
-                builder.end();
+                init = createAdoptChild(builder, varType, init);
+                builder.startStatement().string("this.").string(varName).string(" = ").tree(init).end();
             }
             if (getModel().getNode().isPolymorphic()) {
                 builder.statement("this.next0 = adoptChild(copy.next0)");
@@ -1184,7 +1241,8 @@
         }
 
         private CodeVariableElement createChildField(NodeChildData child) {
-            CodeVariableElement var = new CodeVariableElement(child.getNodeType(), child.getName());
+            TypeMirror type = child.getNodeType();
+            CodeVariableElement var = new CodeVariableElement(type, child.getName());
             var.getModifiers().add(Modifier.PROTECTED);
 
             DeclaredType annotationType;
@@ -1223,7 +1281,7 @@
             }
 
             builder.startStatement().string("String message = ").startCall("createInfo0").string("reason");
-            addInternalValueParameterNames(builder, node.getGenericSpecialization(), node.getGenericSpecialization(), null, false, true);
+            addInternalValueParameterNames(builder, node.getGenericSpecialization(), node.getGenericSpecialization(), null, false, true, null);
             builder.end().end();
 
             final String currentNodeVar = currentNode;
@@ -1434,7 +1492,7 @@
                     guardsAnd = " && ";
                 }
 
-                CodeTree cast = createCast(castBuilder, child, valueParam, typeGuard.getType());
+                CodeTree cast = createCast(castBuilder, child, valueParam, typeGuard.getType(), minimumState);
                 if (cast != null) {
                     castBuilder.tree(cast);
                 }
@@ -1526,7 +1584,15 @@
                 builder.string(" || ");
             }
 
-            startCallTypeSystemMethod(getContext(), builder, node, TypeSystemCodeGenerator.isTypeMethodName(targetType));
+            String castMethodName;
+            List<TypeData> types = getModel().getNode().getTypeSystem().lookupSourceTypes(targetType);
+            if (types.size() > 1) {
+                castMethodName = TypeSystemCodeGenerator.isImplicitTypeMethodName(targetType);
+            } else {
+                castMethodName = TypeSystemCodeGenerator.isTypeMethodName(targetType);
+            }
+
+            startCallTypeSystemMethod(getContext(), builder, node, castMethodName);
             builder.string(valueName(source));
             builder.end().end(); // call
 
@@ -1539,7 +1605,7 @@
             return builder.getRoot();
         }
 
-        private CodeTree createCast(CodeTreeBuilder parent, NodeChildData field, ActualParameter source, TypeData targetType) {
+        private CodeTree createCast(CodeTreeBuilder parent, NodeChildData field, ActualParameter source, TypeData targetType, boolean minimumState) {
             NodeData node = field.getNodeData();
             TypeData sourceType = source.getTypeSystemType();
 
@@ -1554,9 +1620,24 @@
                 condition = CodeTreeBuilder.singleString(valueName(shortCircuit));
             }
 
-            CodeTree value = createCallTypeSystemMethod(context, parent, node, TypeSystemCodeGenerator.asTypeMethodName(targetType), CodeTreeBuilder.singleString(valueName(source)));
-
-            return createLazyAssignment(parent, castValueName(source), targetType.getPrimitiveType(), condition, value);
+            String castMethodName;
+            List<TypeData> types = getModel().getNode().getTypeSystem().lookupSourceTypes(targetType);
+            if (types.size() > 1) {
+                castMethodName = TypeSystemCodeGenerator.asImplicitTypeMethodName(targetType);
+            } else {
+                castMethodName = TypeSystemCodeGenerator.asTypeMethodName(targetType);
+            }
+
+            CodeTree value = createCallTypeSystemMethod(context, parent, node, castMethodName, CodeTreeBuilder.singleString(valueName(source)));
+
+            CodeTreeBuilder builder = parent.create();
+            builder.tree(createLazyAssignment(parent, castValueName(source), targetType.getPrimitiveType(), condition, value));
+            if (minimumState && types.size() > 1) {
+                CodeTree castType = createCallTypeSystemMethod(context, parent, node, TypeSystemCodeGenerator.getImplicitClass(targetType), CodeTreeBuilder.singleString(valueName(source)));
+                builder.tree(createLazyAssignment(builder, typeName(source), getContext().getType(Class.class), condition, castType));
+            }
+
+            return builder.getRoot();
         }
 
         private CodeTree createMethodGuard(CodeTreeBuilder parent, String prefix, SpecializationData source, GuardData guard) {
@@ -1595,7 +1676,6 @@
                 builder.tree(createRewriteGeneric(builder, source, current, currentNodeVar));
                 builder.end();
             } else {
-                // simple rewrite
                 if (current.getExceptions().isEmpty()) {
                     builder.tree(createGenericInvoke(builder, source, current, createReplaceCall(builder, current, currentNodeVar, currentNodeVar, null), null));
                 } else {
@@ -1669,7 +1749,7 @@
             }
             if (current.isGeneric()) {
                 builder.startReturn().tree(replace).string(".").startCall(EXECUTE_GENERIC_NAME);
-                addInternalValueParameterNames(builder, source, current, null, current.getNode().needsFrame(), true);
+                addInternalValueParameterNames(builder, source, current, null, current.getNode().needsFrame(), true, null);
                 builder.end().end();
             } else if (current.getMethod() == null) {
                 if (replaceCall != null) {
@@ -1696,7 +1776,19 @@
             } else {
                 replaceCall.startCall("replace");
             }
-            replaceCall.startGroup().startNew(className).string(source).end().end();
+            replaceCall.startGroup().startNew(className).string(source);
+            for (ActualParameter param : current.getParameters()) {
+                if (!param.getSpecification().isSignature()) {
+                    continue;
+                }
+                NodeChildData child = getModel().getNode().findChild(param.getSpecification().getName());
+                List<TypeData> types = child.getNodeData().getTypeSystem().lookupSourceTypes(param.getTypeSystemType());
+                if (types.size() > 1) {
+                    replaceCall.string(typeName(param));
+                }
+            }
+            replaceCall.end().end();
+
             if (message == null) {
                 replaceCall.string("message");
             } else {
@@ -1732,7 +1824,7 @@
 
             builder.startReturn();
             builder.startCall(currentNode + ".next0", executeCachedName(node.getGenericPolymorphicSpecialization()));
-            addInternalValueParameterNames(builder, node.getGenericSpecialization(), node.getGenericSpecialization(), null, true, true);
+            addInternalValueParameterNames(builder, node.getGenericSpecialization(), node.getGenericSpecialization(), null, true, true, null);
             builder.end();
             builder.end();
 
@@ -1776,7 +1868,7 @@
                 executeParameterNames[i] = valueName(executeParameters.get(i));
             }
 
-            builder.tree(createExecuteChildren(builder, executable, specialization, executeParameters, null, true));
+            builder.tree(createExecuteChildren(builder, executable, specialization, executeParameters, null));
 
             CodeTree primaryExecuteCall = createTemplateMethodCall(builder, null, executable, castExecutable, null, executeParameterNames);
             if (needsTry) {
@@ -1827,134 +1919,180 @@
             return createCastType(node, sourceType, castedType.getType(), hasUnexpected, value);
         }
 
-        protected CodeTree createCastType(NodeData node, TypeData sourceType, TypeData targetType, boolean expect, CodeTree value) {
-            if (targetType == null) {
-                return value;
-            } else if (!sourceType.needsCastTo(getContext(), targetType)) {
-                return value;
-            }
-
-            CodeTreeBuilder builder = CodeTreeBuilder.createBuilder();
-            String targetMethodName;
-            if (expect) {
-                targetMethodName = TypeSystemCodeGenerator.expectTypeMethodName(targetType);
-            } else {
-                targetMethodName = TypeSystemCodeGenerator.asTypeMethodName(targetType);
-            }
-            startCallTypeSystemMethod(getContext(), builder, node, targetMethodName);
-
-            builder.tree(value);
-            builder.end().end();
-            return builder.getRoot();
-        }
-
         protected CodeTree createExecuteChildren(CodeTreeBuilder parent, ExecutableTypeData sourceExecutable, SpecializationData specialization, List<ActualParameter> targetParameters,
-                        ActualParameter unexpectedParameter, boolean cast) {
-            NodeData sourceNode = specialization.getNode();
-
-            CodeTreeBuilder builder = new CodeTreeBuilder(parent);
-
+                        ActualParameter unexpectedParameter) {
+            CodeTreeBuilder builder = parent.create();
+            NodeData node = specialization.getNode();
             for (ActualParameter targetParameter : targetParameters) {
-                NodeChildData field = sourceNode.findChild(targetParameter.getSpecification().getName());
+                NodeChildData child = node.findChild(targetParameter.getSpecification().getName());
                 if (!targetParameter.getSpecification().isSignature()) {
                     continue;
                 }
-
                 TypeData targetType = targetParameter.getTypeSystemType();
                 ExecutableTypeData targetExecutable = null;
-                if (field != null) {
-                    targetExecutable = field.findExecutableType(getContext(), targetType);
+                if (child != null) {
+                    targetExecutable = child.findExecutableType(getContext(), targetType);
+                }
+
+                if (targetExecutable == null) {
+                    // TODO what to do? assertion?
+                    continue;
                 }
 
-                ActualParameter sourceParameter = sourceExecutable.findParameter(targetParameter.getLocalName());
-
-                String targetVariableName = valueName(targetParameter);
-
-                CodeTree executionExpression = null;
-                if ((sourceParameter != null && cast) || sourceParameter != null) {
-                    TypeData sourceType = sourceParameter.getTypeSystemType();
-                    if (targetExecutable == null || !sourceType.needsCastTo(getContext(), targetType)) {
-                        if (field != null && field.isShortCircuit() && sourceParameter != null) {
-                            builder.tree(createShortCircuitValue(builder, specialization, field, targetParameter.getPreviousParameter(), unexpectedParameter));
-                        }
-                        builder.startStatement();
-                        builder.type(targetParameter.getType()).string(" ");
-                        builder.string(targetVariableName).string(" = ");
-                        builder.tree(CodeTreeBuilder.singleString(valueNameEvaluated(targetParameter)));
-                        builder.end();
-                        continue;
+                CodeTree executionExpressions = createExecutionExpresssions(builder, child, sourceExecutable, targetExecutable, targetParameter, unexpectedParameter);
+
+                String targetVarName = valueName(targetParameter);
+                CodeTree unexpectedTree = createCatchUnexpectedTree(builder, executionExpressions, targetVarName, specialization, sourceExecutable, targetExecutable, targetParameter,
+                                isShortCircuit(child));
+                CodeTree shortCircuitTree = createShortCircuitTree(builder, unexpectedTree, targetVarName, specialization, targetParameter, unexpectedParameter);
+
+                if (shortCircuitTree == executionExpressions) {
+                    if (containsNewLine(executionExpressions)) {
+                        builder.declaration(sourceExecutable.getType().getPrimitiveType(), targetVarName);
+                        builder.tree(shortCircuitTree);
                     } else {
-                        CodeTree valueTree = CodeTreeBuilder.singleString(valueNameEvaluated(targetParameter));
-                        executionExpression = createExpectExecutableType(sourceNode, sourceType, targetExecutable, valueTree);
+                        builder.startStatement().type(targetParameter.getType()).string(" ").tree(shortCircuitTree).end();
                     }
-                } else if (sourceParameter == null) {
-                    executionExpression = createExecuteChildExpression(builder, field, targetParameter, unexpectedParameter);
-                }
-
-                if (executionExpression != null) {
-                    CodeTreeVariable executionVar = new CodeTreeVariable();
-                    CodeTree shortCircuitTree = createShortCircuitTree(builder, executionVar, targetVariableName, specialization, targetParameter, unexpectedParameter);
-                    CodeTree unexpectedTree = createCatchUnexpectedTree(builder, executionExpression, targetVariableName, specialization, sourceExecutable, targetExecutable, targetParameter,
-                                    shortCircuitTree != executionVar);
-
-                    executionVar.set(unexpectedTree);
+                } else {
                     builder.tree(shortCircuitTree);
                 }
+
+            }
+            return builder.getRoot();
+        }
+
+        private CodeTree createExecutionExpresssions(CodeTreeBuilder parent, NodeChildData child, ExecutableTypeData sourceExecutable, ExecutableTypeData targetExecutable, ActualParameter param,
+                        ActualParameter unexpectedParameter) {
+            CodeTreeBuilder builder = parent.create();
+
+            TypeData type = param.getTypeSystemType();
+            List<TypeData> targetTypes = child.getNodeData().getTypeSystem().lookupSourceTypes(type);
+
+            if (targetTypes.size() > 1) {
+                boolean elseIf = false;
+                int index = 0;
+                for (TypeData typeData : targetTypes) {
+                    if (index < targetTypes.size() - 1) {
+                        elseIf = builder.startIf(elseIf);
+                        builder.string(typeName(param)).string(" == ").typeLiteral(typeData.getPrimitiveType());
+                        builder.end();
+                        builder.startBlock();
+                    } else {
+                        builder.startElseBlock();
+                    }
+
+                    ExecutableTypeData implictExecutableTypeData = child.getNodeData().findExecutableType(typeData, targetExecutable.getEvaluatedCount());
+                    ImplicitCastData cast = child.getNodeData().getTypeSystem().lookupCast(typeData, param.getTypeSystemType());
+                    CodeTree execute = createExecuteExpression(parent, child, sourceExecutable, implictExecutableTypeData, param, unexpectedParameter, cast);
+                    builder.statement(execute);
+                    builder.end();
+                    index++;
+                }
+            } else {
+                builder.tree(createExecuteExpression(parent, child, sourceExecutable, targetExecutable, param, unexpectedParameter, null));
             }
             return builder.getRoot();
         }
 
+        private CodeTree createExecuteExpression(CodeTreeBuilder parent, NodeChildData child, ExecutableTypeData sourceExecutable, ExecutableTypeData targetExecutable,
+                        ActualParameter targetParameter, ActualParameter unexpectedParameter, ImplicitCastData cast) {
+            CodeTreeBuilder builder = parent.create();
+            builder.string(valueName(targetParameter));
+            builder.string(" = ");
+            if (cast != null) {
+                startCallTypeSystemMethod(getContext(), builder, child.getNodeData(), cast.getMethodName());
+            }
+
+            NodeData node = getModel().getNode();
+            ActualParameter sourceParameter = sourceExecutable.findParameter(targetParameter.getLocalName());
+            if (sourceParameter == null) {
+                builder.tree(createExecuteChildExpression(builder, child, targetParameter, targetExecutable, unexpectedParameter));
+            } else {
+                CodeTree var = CodeTreeBuilder.singleString(valueNameEvaluated(targetParameter));
+                builder.tree(createExpectExecutableType(node, sourceParameter.getTypeSystemType(), targetExecutable, var));
+            }
+            if (cast != null) {
+                builder.end().end();
+            }
+
+            return builder.getRoot();
+        }
+
+        private boolean containsNewLine(CodeTree tree) {
+            if (tree.getCodeKind() == CodeTreeKind.NEW_LINE) {
+                return true;
+            }
+
+            for (CodeTree codeTree : tree.getEnclosedElements()) {
+                if (containsNewLine(codeTree)) {
+                    return true;
+                }
+            }
+            return false;
+        }
+
+        private boolean hasUnexpected(ExecutableTypeData target, ActualParameter sourceParameter, ActualParameter targetParameter) {
+            List<TypeData> types = getModel().getNode().getTypeSystem().lookupSourceTypes(targetParameter.getTypeSystemType());
+            NodeChildData child = getModel().getNode().findChild(targetParameter.getSpecification().getName());
+            boolean hasUnexpected = false;
+            for (TypeData type : types) {
+                if (hasUnexpected) {
+                    continue;
+                }
+                ExecutableTypeData execTarget = target;
+                if (type != execTarget.getType()) {
+                    execTarget = child.findExecutableType(getContext(), type);
+                }
+                hasUnexpected = hasUnexpected || hasUnexpectedType(execTarget, sourceParameter, type);
+            }
+            return hasUnexpected;
+        }
+
+        private boolean hasUnexpectedType(ExecutableTypeData target, ActualParameter sourceParameter, TypeData type) {
+            if (sourceParameter == null) {
+                return target.hasUnexpectedValue(getContext());
+            } else {
+                if (sourceParameter.getTypeSystemType().needsCastTo(getContext(), type)) {
+                    return target.hasUnexpectedValue(getContext());
+                }
+                return false;
+            }
+        }
+
         private CodeTree createCatchUnexpectedTree(CodeTreeBuilder parent, CodeTree body, String targetVariableName, SpecializationData specialization, ExecutableTypeData currentExecutable,
                         ExecutableTypeData targetExecutable, ActualParameter param, boolean shortCircuit) {
             CodeTreeBuilder builder = new CodeTreeBuilder(parent);
-            boolean unexpected = targetExecutable.hasUnexpectedValue(getContext());
-            boolean cast = false;
-            if (targetExecutable.getType().needsCastTo(getContext(), param.getTypeSystemType())) {
-                unexpected = true;
-                cast = true;
+            ActualParameter sourceParameter = currentExecutable.findParameter(param.getLocalName());
+            boolean unexpected = hasUnexpected(targetExecutable, sourceParameter, param);
+            if (!unexpected) {
+                return body;
             }
 
-            builder.startStatement();
-
             if (!shortCircuit) {
-                builder.type(param.getType()).string(" ").string(targetVariableName);
+                builder.declaration(param.getType(), targetVariableName);
             }
-
-            if (unexpected) {
-                if (!shortCircuit) {
-                    builder.end();
-                }
-                builder.startTryBlock();
-                builder.startStatement();
-                builder.string(targetVariableName);
-            } else if (shortCircuit) {
-                builder.startStatement();
-                builder.string(targetVariableName);
-            }
-            builder.string(" = ");
-            if (cast) {
-                builder.tree(createCastType(specialization.getNode(), targetExecutable.getType(), param.getTypeSystemType(), true, body));
-            } else {
+            builder.startTryBlock();
+
+            if (containsNewLine(body)) {
                 builder.tree(body);
+            } else {
+                builder.statement(body);
             }
-            builder.end();
-
-            if (unexpected) {
-                builder.end().startCatchBlock(getUnexpectedValueException(), "ex");
-                SpecializationData generic = specialization.getNode().getGenericSpecialization();
-                ActualParameter genericParameter = generic.findParameter(param.getLocalName());
-
-                List<ActualParameter> genericParameters = generic.getParametersAfter(genericParameter);
-                builder.tree(createDeoptimize(builder));
-                builder.tree(createExecuteChildren(parent, currentExecutable, generic, genericParameters, genericParameter, false));
-                if (specialization.isPolymorphic()) {
-                    builder.tree(createReturnOptimizeTypes(builder, currentExecutable, specialization, param));
-                } else {
-                    builder.tree(createReturnExecuteAndSpecialize(builder, currentExecutable, specialization, param,
-                                    "Expected " + param.getLocalName() + " instanceof " + Utils.getSimpleName(param.getType())));
-                }
-                builder.end(); // catch block
+
+            builder.end().startCatchBlock(getUnexpectedValueException(), "ex");
+            SpecializationData generic = specialization.getNode().getGenericSpecialization();
+            ActualParameter genericParameter = generic.findParameter(param.getLocalName());
+
+            List<ActualParameter> genericParameters = generic.getParametersAfter(genericParameter);
+            builder.tree(createDeoptimize(builder));
+            builder.tree(createExecuteChildren(parent, currentExecutable, generic, genericParameters, genericParameter));
+            if (specialization.isPolymorphic()) {
+                builder.tree(createReturnOptimizeTypes(builder, currentExecutable, specialization, param));
+            } else {
+                builder.tree(createReturnExecuteAndSpecialize(builder, currentExecutable, specialization, param,
+                                "Expected " + param.getLocalName() + " instanceof " + Utils.getSimpleName(param.getType())));
             }
+            builder.end(); // catch block
 
             return builder.getRoot();
         }
@@ -1969,7 +2107,7 @@
 
             CodeTreeBuilder execute = new CodeTreeBuilder(builder);
             execute.startCall("next0", executeCachedName(generic));
-            addInternalValueParameterNames(execute, specialization, generic, param.getLocalName(), true, true);
+            addInternalValueParameterNames(execute, specialization, generic, param.getLocalName(), true, true, null);
             execute.end();
 
             TypeData sourceType = generic.getReturnType().getTypeSystemType();
@@ -1980,37 +2118,26 @@
             return builder.getRoot();
         }
 
-        private CodeTree createExecuteChildExpression(CodeTreeBuilder parent, NodeChildData targetField, ActualParameter sourceParameter, ActualParameter unexpectedParameter) {
-            TypeData type = sourceParameter.getTypeSystemType();
-            ExecutableTypeData execType = targetField.findExecutableType(getContext(), type);
-
+        private CodeTree createExecuteChildExpression(CodeTreeBuilder parent, NodeChildData targetChild, ActualParameter targetParameter, ExecutableTypeData targetExecutable,
+                        ActualParameter unexpectedParameter) {
             CodeTreeBuilder builder = new CodeTreeBuilder(parent);
-            if (targetField != null) {
-                Element accessElement = targetField.getAccessElement();
-                if (accessElement == null || accessElement.getKind() == ElementKind.METHOD) {
-                    builder.string("this.").string(targetField.getName());
-                } else if (accessElement.getKind() == ElementKind.FIELD) {
-                    builder.string("this.").string(accessElement.getSimpleName().toString());
-                } else {
-                    throw new AssertionError();
-                }
-                if (sourceParameter.getSpecification().isIndexed()) {
-                    builder.string("[" + sourceParameter.getIndex() + "]");
-                }
+            if (targetChild != null) {
+                builder.tree(createAccessChild(builder, targetChild, targetParameter));
                 builder.string(".");
             }
 
-            builder.startCall(execType.getMethodName());
-
+            builder.startCall(targetExecutable.getMethodName());
+
+            // TODO this should be merged with #createTemplateMethodCall
             int index = 0;
-            for (ActualParameter parameter : execType.getParameters()) {
+            for (ActualParameter parameter : targetExecutable.getParameters()) {
 
                 if (!parameter.getSpecification().isSignature()) {
                     builder.string(parameter.getLocalName());
                 } else {
 
-                    if (index < targetField.getExecuteWith().size()) {
-                        NodeChildData child = targetField.getExecuteWith().get(index);
+                    if (index < targetChild.getExecuteWith().size()) {
+                        NodeChildData child = targetChild.getExecuteWith().get(index);
 
                         ParameterSpec spec = getModel().getSpecification().findParameterSpec(child.getName());
                         List<ActualParameter> specializationParams = getModel().findParameters(spec);
@@ -2049,32 +2176,50 @@
             return builder.getRoot();
         }
 
+        private CodeTree createAccessChild(CodeTreeBuilder parent, NodeChildData targetChild, ActualParameter targetParameter) throws AssertionError {
+            CodeTreeBuilder builder = parent.create();
+            Element accessElement = targetChild.getAccessElement();
+            if (accessElement == null || accessElement.getKind() == ElementKind.METHOD) {
+                builder.string("this.").string(targetChild.getName());
+            } else if (accessElement.getKind() == ElementKind.FIELD) {
+                builder.string("this.").string(accessElement.getSimpleName().toString());
+            } else {
+                throw new AssertionError();
+            }
+            if (targetParameter.getSpecification().isIndexed()) {
+                builder.string("[" + targetParameter.getIndex() + "]");
+            }
+            return builder.getRoot();
+        }
+
         private CodeTree createShortCircuitTree(CodeTreeBuilder parent, CodeTree body, String targetVariableName, SpecializationData specialization, ActualParameter parameter,
                         ActualParameter exceptionParam) {
-            CodeTreeBuilder builder = new CodeTreeBuilder(parent);
-
             NodeChildData forField = specialization.getNode().findChild(parameter.getSpecification().getName());
-            if (forField == null) {
+            if (!isShortCircuit(forField)) {
                 return body;
             }
 
-            if (forField.getExecutionKind() != ExecutionKind.SHORT_CIRCUIT) {
-                return body;
-            }
-
+            CodeTreeBuilder builder = new CodeTreeBuilder(parent);
             ActualParameter shortCircuitParam = specialization.getPreviousParam(parameter);
-
             builder.tree(createShortCircuitValue(builder, specialization, forField, shortCircuitParam, exceptionParam));
-
             builder.declaration(parameter.getType(), targetVariableName, CodeTreeBuilder.createBuilder().defaultValue(parameter.getType()));
             builder.startIf().string(shortCircuitParam.getLocalName()).end();
             builder.startBlock();
-            builder.tree(body);
+
+            if (containsNewLine(body)) {
+                builder.tree(body);
+            } else {
+                builder.statement(body);
+            }
             builder.end();
 
             return builder.getRoot();
         }
 
+        private boolean isShortCircuit(NodeChildData forField) {
+            return forField != null && forField.getExecutionKind() == ExecutionKind.SHORT_CIRCUIT;
+        }
+
         private CodeTree createShortCircuitValue(CodeTreeBuilder parent, SpecializationData specialization, NodeChildData forField, ActualParameter shortCircuitParam, ActualParameter exceptionParam) {
             CodeTreeBuilder builder = new CodeTreeBuilder(parent);
             int shortCircuitIndex = 0;
@@ -2109,7 +2254,7 @@
             CodeTreeBuilder specializeCall = new CodeTreeBuilder(parent);
             specializeCall.startCall(EXECUTE_SPECIALIZE_NAME);
             specializeCall.string(String.valueOf(node.getSpecializations().indexOf(current)));
-            addInternalValueParameterNames(specializeCall, generic, node.getGenericSpecialization(), exceptionParam != null ? exceptionParam.getLocalName() : null, true, true);
+            addInternalValueParameterNames(specializeCall, generic, node.getGenericSpecialization(), exceptionParam != null ? exceptionParam.getLocalName() : null, true, true, null);
             specializeCall.doubleQuote(reason);
             specializeCall.end().end();
 
@@ -2199,6 +2344,233 @@
         }
     }
 
+    private class BaseCastNodeFactory extends ClassElementFactory<NodeData> {
+
+        protected final Set<TypeData> usedTargetTypes;
+
+        public BaseCastNodeFactory(ProcessorContext context, Set<TypeData> usedTargetTypes) {
+            super(context);
+            this.usedTargetTypes = usedTargetTypes;
+        }
+
+        @Override
+        protected CodeTypeElement create(NodeData m) {
+            CodeTypeElement type = createClass(m, modifiers(STATIC), nodeCastClassName(m, null), context.getTruffleTypes().getNode(), false);
+
+            CodeVariableElement delegate = new CodeVariableElement(m.getNodeType(), "delegate");
+            delegate.getModifiers().add(PROTECTED);
+            delegate.getAnnotationMirrors().add(new CodeAnnotationMirror(getContext().getTruffleTypes().getChildAnnotation()));
+
+            type.add(delegate);
+            type.add(createConstructorUsingFields(modifiers(), type));
+            return type;
+        }
+
+        @Override
+        protected void createChildren(NodeData m) {
+            CodeTypeElement type = getElement();
+            type.add(createExecute(EXECUTE_SPECIALIZE_NAME, true));
+            type.add(createExecute(EXECUTE_GENERIC_NAME, false));
+
+            for (ExecutableTypeData targetExecutable : m.getExecutableTypes()) {
+                if (!usedTargetTypes.contains(targetExecutable.getType()) && targetExecutable.hasUnexpectedValue(getContext())) {
+                    continue;
+                }
+                CodeExecutableElement execute = createCastExecute(targetExecutable, targetExecutable, false);
+                CodeExecutableElement expect = createCastExecute(targetExecutable, targetExecutable, true);
+                if (execute != null) {
+                    getElement().add(execute);
+                }
+                if (expect != null) {
+                    getElement().add(expect);
+                }
+            }
+            Set<TypeData> sourceTypes = new TreeSet<>();
+            List<ImplicitCastData> casts = getModel().getTypeSystem().getImplicitCasts();
+            for (ImplicitCastData cast : casts) {
+                sourceTypes.add(cast.getSourceType());
+            }
+
+            CodeTypeElement baseType = getElement();
+            for (TypeData sourceType : sourceTypes) {
+                add(new SpecializedCastNodeFactory(context, baseType, sourceType, usedTargetTypes), getModel());
+            }
+        }
+
+        private CodeExecutableElement createExecute(String name, boolean specialize) {
+            NodeData node = getModel();
+            TypeMirror objectType = node.getTypeSystem().getGenericType();
+            CodeExecutableElement method = new CodeExecutableElement(modifiers(PUBLIC), objectType, name, new CodeVariableElement(objectType, "value"));
+            if (specialize) {
+                method.getModifiers().add(FINAL);
+            }
+            CodeTreeBuilder builder = method.createBuilder();
+
+            List<ImplicitCastData> casts = node.getTypeSystem().getImplicitCasts();
+            boolean elseIf = false;
+            for (ImplicitCastData cast : casts) {
+                elseIf = builder.startIf(elseIf);
+                startCallTypeSystemMethod(context, builder, getModel(), TypeSystemCodeGenerator.isTypeMethodName(cast.getSourceType()));
+                builder.string("value");
+                builder.end().end();
+                builder.end();
+                builder.startBlock();
+
+                if (specialize) {
+                    builder.startStatement().startCall("replace").startNew(nodeCastClassName(getModel(), cast.getSourceType())).string("delegate").end().doubleQuote("Added cast").end().end();
+                }
+                builder.startReturn();
+
+                startCallTypeSystemMethod(context, builder, getModel(), cast.getMethodName());
+                startCallTypeSystemMethod(context, builder, getModel(), TypeSystemCodeGenerator.asTypeMethodName(cast.getSourceType()));
+                builder.string("value");
+                builder.end().end();
+                builder.end().end();
+
+                builder.end();
+                builder.end();
+            }
+
+            builder.startReturn().string("value").end();
+
+            return method;
+        }
+
+        protected CodeExecutableElement createCastExecute(ExecutableTypeData sourceExecutable, ExecutableTypeData targetExecutable, boolean expect) {
+            ImplicitCastData cast = null;
+            if (!sourceExecutable.getType().equals(targetExecutable.getType())) {
+                cast = getModel().getTypeSystem().lookupCast(sourceExecutable.getType(), targetExecutable.getType());
+                if (cast == null) {
+                    return null;
+                }
+            }
+
+            if (expect) {
+                if (targetExecutable.getEvaluatedCount() > 0) {
+                    return null;
+                } else if (Utils.isObject(targetExecutable.getType().getPrimitiveType())) {
+                    return null;
+                }
+            }
+
+            boolean hasTargetUnexpected = targetExecutable.hasUnexpectedValue(getContext());
+            boolean hasSourceUnexpected = sourceExecutable.hasUnexpectedValue(getContext());
+
+            CodeExecutableElement method = copyTemplateMethod(targetExecutable);
+            method.getModifiers().add(PUBLIC);
+
+            CodeTreeBuilder builder = method.createBuilder();
+
+            if (hasSourceUnexpected && cast != null) {
+                builder.startTryBlock();
+            }
+
+            if (expect) {
+                method.getParameters().clear();
+                String expectMethodName;
+                if (hasTargetUnexpected) {
+                    expectMethodName = TypeSystemCodeGenerator.expectTypeMethodName(targetExecutable.getType());
+                } else {
+                    expectMethodName = TypeSystemCodeGenerator.asTypeMethodName(targetExecutable.getType());
+                }
+                method.setSimpleName(CodeNames.of(expectMethodName));
+                method.addParameter(new CodeVariableElement(getModel().getTypeSystem().getGenericType(), "value"));
+            }
+
+            builder.startReturn();
+            CodeTree executeCall;
+            if (expect) {
+                executeCall = createCastType(getModel(), getModel().getTypeSystem().getGenericTypeData(), sourceExecutable.getType(), hasSourceUnexpected, CodeTreeBuilder.singleString("value"));
+            } else {
+                executeCall = createTemplateMethodCall(builder, CodeTreeBuilder.singleString("delegate."), targetExecutable, sourceExecutable, null);
+            }
+            if (cast != null) {
+                startCallTypeSystemMethod(context, builder, getModel(), cast.getMethodName());
+                builder.tree(executeCall);
+                builder.end().end();
+            } else {
+                builder.tree(executeCall);
+            }
+            builder.end();
+
+            if (hasSourceUnexpected && cast != null) {
+                builder.end();
+                builder.startCatchBlock(getContext().getTruffleTypes().getUnexpectedValueException(), "ex");
+                builder.startStatement().startCall("replace").startNew(nodeCastClassName(getModel(), null)).string("delegate").end().doubleQuote("Removed cast").end().end();
+
+                if (hasTargetUnexpected) {
+                    builder.startThrow().string("ex").end();
+                } else {
+                    builder.startThrow().startNew(getContext().getType(AssertionError.class)).end().end();
+                }
+                builder.end();
+            }
+
+            return method;
+        }
+
+        private CodeExecutableElement copyTemplateMethod(TemplateMethod targetExecutable) {
+            CodeExecutableElement method = CodeExecutableElement.clone(getContext().getEnvironment(), targetExecutable.getMethod());
+            method.getModifiers().remove(ABSTRACT);
+            method.getAnnotationMirrors().clear();
+            Modifier visibility = Utils.getVisibility(method.getModifiers());
+            if (visibility != null) {
+                method.getModifiers().remove(visibility);
+            }
+            int index = 0;
+            for (ActualParameter parameter : targetExecutable.getParameters()) {
+                ((CodeVariableElement) method.getParameters().get(index)).setName(parameter.getLocalName());
+                index++;
+            }
+            return method;
+        }
+
+    }
+
+    private class SpecializedCastNodeFactory extends BaseCastNodeFactory {
+
+        private final CodeTypeElement baseType;
+        private final TypeData sourceType;
+
+        public SpecializedCastNodeFactory(ProcessorContext context, CodeTypeElement baseType, TypeData type, Set<TypeData> usedTargetTypes) {
+            super(context, usedTargetTypes);
+            this.baseType = baseType;
+            this.sourceType = type;
+        }
+
+        @Override
+        protected CodeTypeElement create(NodeData m) {
+            CodeTypeElement type = createClass(m, modifiers(PRIVATE, STATIC, FINAL), nodeCastClassName(m, sourceType), baseType.asType(), false);
+            type.add(createConstructorUsingFields(modifiers(), type));
+            return type;
+        }
+
+        @Override
+        protected void createChildren(NodeData node) {
+            for (TypeData targetType : usedTargetTypes) {
+                for (ExecutableTypeData targetExecutable : node.getExecutableTypes()) {
+                    if (targetExecutable.getType().equals(targetType)) {
+                        ExecutableTypeData sourceExecutable = node.findExecutableType(sourceType, targetExecutable.getEvaluatedCount());
+                        if (sourceExecutable == null) {
+                            // TODO what if there is no evaluated version?
+                            continue;
+                        }
+                        CodeExecutableElement execute = createCastExecute(sourceExecutable, targetExecutable, false);
+                        CodeExecutableElement expect = createCastExecute(sourceExecutable, targetExecutable, true);
+                        if (execute != null) {
+                            getElement().add(execute);
+                        }
+                        if (expect != null) {
+                            getElement().add(expect);
+                        }
+                    }
+                }
+
+            }
+        }
+
+    }
+
     private class SpecializedNodeFactory extends NodeBaseFactory {
 
         protected final CodeTypeElement nodeGen;
@@ -2277,15 +2649,31 @@
                 }
 
                 CodeExecutableElement superConstructor = createSuperConstructor(clazz, constructor);
+                CodeTree body = superConstructor.getBodyTree();
+                CodeTreeBuilder builder = superConstructor.createBuilder();
+                builder.tree(body);
 
                 if (superConstructor != null) {
                     if (getModel().isGeneric() && node.isPolymorphic()) {
-                        CodeTree body = superConstructor.getBodyTree();
-                        CodeTreeBuilder builder = superConstructor.createBuilder();
-                        builder.tree(body);
                         builder.statement("this.next0 = null");
                     }
 
+                    for (ActualParameter param : getModel().getParameters()) {
+                        if (!param.getSpecification().isSignature()) {
+                            continue;
+                        }
+                        NodeChildData child = getModel().getNode().findChild(param.getSpecification().getName());
+                        List<TypeData> types = child.getNodeData().getTypeSystem().lookupSourceTypes(param.getTypeSystemType());
+                        if (types.size() > 1) {
+                            clazz.add(new CodeVariableElement(modifiers(PRIVATE, FINAL), getContext().getType(Class.class), typeName(param)));
+                            superConstructor.getParameters().add(new CodeVariableElement(getContext().getType(Class.class), typeName(param)));
+
+                            builder.startStatement();
+                            builder.string("this.").string(typeName(param)).string(" = ").string(typeName(param));
+                            builder.end();
+                        }
+                    }
+
                     clazz.add(superConstructor);
                 }
             }
@@ -2333,7 +2721,7 @@
                 } else {
                     CodeTreeBuilder elseBuilder = new CodeTreeBuilder(builder);
                     elseBuilder.startReturn().startCall("this.next0", executeCachedName(polymorphic));
-                    addInternalValueParameterNames(elseBuilder, polymorphic, polymorphic, null, true, true);
+                    addInternalValueParameterNames(elseBuilder, polymorphic, polymorphic, null, true, true, null);
                     elseBuilder.end().end();
 
                     boolean forceElse = specialization.getExceptions().size() > 0;
@@ -2375,7 +2763,7 @@
             CodeTreeBuilder specializeCall = new CodeTreeBuilder(builder);
             specializeCall.startCall(EXECUTE_SPECIALIZE_NAME);
             specializeCall.string("0");
-            addInternalValueParameterNames(specializeCall, specialization, node.getGenericSpecialization(), null, true, true);
+            addInternalValueParameterNames(specializeCall, specialization, node.getGenericSpecialization(), null, true, true, null);
             specializeCall.startGroup().doubleQuote("Uninitialized polymorphic (").string(" + depth + ").doubleQuote("/" + node.getPolymorphicDepth() + ")").end();
             specializeCall.end().end();
 
@@ -2482,7 +2870,7 @@
                 builder.tree(createDeoptimize(builder));
             }
 
-            builder.tree(createExecuteChildren(builder, executable, specialization, specialization.getParameters(), null, false));
+            builder.tree(createExecuteChildren(builder, executable, specialization, specialization.getParameters(), null));
 
             CodeTree returnSpecialized = null;
 
@@ -2519,26 +2907,25 @@
             CodeTreeBuilder returnBuilder = new CodeTreeBuilder(parent);
             if (specialization.isPolymorphic()) {
                 returnBuilder.startCall("next0", executeCachedName(specialization));
-                addInternalValueParameterNames(returnBuilder, specialization, specialization, null, true, true);
+                addInternalValueParameterNames(returnBuilder, specialization, specialization, null, true, true, null);
                 returnBuilder.end();
             } else if (specialization.isUninitialized()) {
                 returnBuilder.startCall("super", EXECUTE_SPECIALIZE_NAME);
                 returnBuilder.string("0");
-                addInternalValueParameterNames(returnBuilder, specialization, specialization, null, true, true);
+                addInternalValueParameterNames(returnBuilder, specialization, specialization, null, true, true, null);
                 returnBuilder.doubleQuote("Uninitialized monomorphic");
                 returnBuilder.end();
             } else if (specialization.getMethod() == null && !node.needsRewrites(context)) {
                 emitEncounteredSynthetic(builder, specialization);
             } else if (specialization.isGeneric()) {
                 returnBuilder.startCall("super", EXECUTE_GENERIC_NAME);
-                addInternalValueParameterNames(returnBuilder, specialization, specialization, null, node.needsFrame(), true);
+                addInternalValueParameterNames(returnBuilder, specialization, specialization, null, node.needsFrame(), true, null);
                 returnBuilder.end();
             } else {
                 returnBuilder.tree(createTemplateMethodCall(returnBuilder, null, specialization, specialization, null));
             }
 
             if (!returnBuilder.isEmpty()) {
-
                 ExecutableTypeData sourceExecutableType = node.findExecutableType(specialization.getReturnType().getTypeSystemType(), 0);
                 boolean sourceThrowsUnexpected = sourceExecutableType != null && sourceExecutableType.hasUnexpectedValue(getContext());
                 boolean targetSupportsUnexpected = executable.hasUnexpectedValue(getContext());
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/NodeData.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/NodeData.java	Sat Sep 07 12:23:40 2013 -0700
@@ -32,7 +32,7 @@
 import com.oracle.truffle.dsl.processor.template.*;
 import com.oracle.truffle.dsl.processor.typesystem.*;
 
-public class NodeData extends Template {
+public class NodeData extends Template implements Comparable<NodeData> {
 
     private final String nodeId;
     private NodeData declaringNode;
@@ -90,6 +90,15 @@
         return false;
     }
 
+    public boolean needsImplicitCast(ProcessorContext context) {
+        for (NodeChildData child : getChildren()) {
+            if (child.needsImplicitCast(context)) {
+                return true;
+            }
+        }
+        return false;
+    }
+
     public int getPolymorphicDepth() {
         return polymorphicDepth;
     }
@@ -574,4 +583,7 @@
         return null;
     }
 
+    public int compareTo(NodeData o) {
+        return getNodeId().compareTo(o.getNodeId());
+    }
 }
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/NodeParser.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/NodeParser.java	Sat Sep 07 12:23:40 2013 -0700
@@ -332,7 +332,7 @@
         nodeData.setNodeContainer(nodeContainer != null);
         nodeData.setTypeSystem(typeSystem);
         nodeData.setFields(parseFields(typeHierarchy, elements));
-        nodeData.setChildren(parseChildren(elements, typeHierarchy));
+        nodeData.setChildren(parseChildren(nodeData, elements, typeHierarchy));
         nodeData.setExecutableTypes(groupExecutableTypes(new ExecutableTypeMethodParser(context, nodeData).parse(elements)));
 
         // resolveChildren invokes cyclic parsing.
@@ -407,7 +407,7 @@
         }
     }
 
-    private List<NodeChildData> parseChildren(List<? extends Element> elements, final List<TypeElement> typeHierarchy) {
+    private List<NodeChildData> parseChildren(NodeData parent, List<? extends Element> elements, final List<TypeElement> typeHierarchy) {
         Set<String> shortCircuits = new HashSet<>();
         for (ExecutableElement method : ElementFilter.methodsIn(elements)) {
             AnnotationMirror mirror = Utils.findAnnotationMirror(processingEnv, method, ShortCircuit.class);
@@ -472,7 +472,7 @@
                     kind = ExecutionKind.SHORT_CIRCUIT;
                 }
 
-                NodeChildData nodeChild = new NodeChildData(type, childMirror, name, childType, originalChildType, getter, cardinality, kind);
+                NodeChildData nodeChild = new NodeChildData(parent, type, childMirror, name, childType, originalChildType, getter, cardinality, kind);
 
                 parsedChildren.add(nodeChild);
 
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/template/MessageContainer.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/template/MessageContainer.java	Sat Sep 07 12:23:40 2013 -0700
@@ -129,7 +129,7 @@
 
         TypeElement rootEnclosing = Utils.findRootEnclosingType(getMessageElement());
         TypeElement baseEnclosing = Utils.findRootEnclosingType(baseType);
-        if (rootEnclosing == null || !Utils.typeEquals(baseEnclosing.asType(), rootEnclosing.asType()) || this != message.getOriginalContainer()) {
+        if (rootEnclosing == null || !Utils.typeEquals(baseEnclosing.asType(), rootEnclosing.asType())) {
             // redirect message
             MessageContainer original = message.getOriginalContainer();
             messageElement = baseType;
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/ImplicitCastData.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/ImplicitCastData.java	Sat Sep 07 12:23:40 2013 -0700
@@ -26,8 +26,32 @@
 
 public class ImplicitCastData extends TemplateMethod {
 
-    public ImplicitCastData(TemplateMethod method) {
+    private final TypeData sourceType;
+    private final TypeData targetType;
+
+    public ImplicitCastData(TemplateMethod method, TypeData sourceType, TypeData targetType) {
         super(method);
+        this.sourceType = sourceType;
+        this.targetType = targetType;
+    }
+
+    public TypeData getSourceType() {
+        return sourceType;
+    }
+
+    public TypeData getTargetType() {
+        return targetType;
+    }
+
+    @Override
+    public int compareTo(TemplateMethod o) {
+        if (o instanceof ImplicitCastData) {
+            // implicit casts are ordered by source type since
+            // its also the order in which they are checked.
+            TypeData otherSourceType = ((ImplicitCastData) o).getSourceType();
+            return this.sourceType.compareTo(otherSourceType);
+        }
+        return super.compareTo(o);
     }
 
 }
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/ImplicitCastParser.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/ImplicitCastParser.java	Sat Sep 07 12:23:40 2013 -0700
@@ -56,7 +56,20 @@
 
     @Override
     public ImplicitCastData create(TemplateMethod method, boolean invalid) {
-        return new ImplicitCastData(method);
+        if (invalid) {
+            return new ImplicitCastData(method, null, null);
+        }
+
+        ActualParameter target = method.findParameter("targetValue");
+        ActualParameter source = method.findParameter("sourceValue");
+
+        TypeData targetType = target.getTypeSystemType();
+        TypeData sourceType = source.getTypeSystemType();
+
+        if (targetType.equals(sourceType)) {
+            method.addError("Target type and source type of an @%s must not be the same type.", ImplicitCast.class.getSimpleName());
+        }
+
+        return new ImplicitCastData(method, sourceType, targetType);
     }
-
 }
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/TypeSystemCodeGenerator.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/TypeSystemCodeGenerator.java	Sat Sep 07 12:23:40 2013 -0700
@@ -44,10 +44,22 @@
         return "is" + Utils.getTypeId(type.getBoxedType());
     }
 
+    public static String isImplicitTypeMethodName(TypeData type) {
+        return "isImplicit" + Utils.getTypeId(type.getBoxedType());
+    }
+
     public static String asTypeMethodName(TypeData type) {
         return "as" + Utils.getTypeId(type.getBoxedType());
     }
 
+    public static String asImplicitTypeMethodName(TypeData type) {
+        return "asImplicit" + Utils.getTypeId(type.getBoxedType());
+    }
+
+    public static String getImplicitClass(TypeData type) {
+        return "getImplicit" + Utils.getTypeId(type.getBoxedType()) + "Class";
+    }
+
     public static String expectTypeMethodName(TypeData type) {
         return "expect" + Utils.getTypeId(type.getBoxedType());
     }
@@ -100,6 +112,20 @@
                             clazz.add(expect);
                         }
                     }
+
+                    CodeExecutableElement asImplicit = createAsImplicitTypeMethod(type);
+                    if (asImplicit != null) {
+                        clazz.add(asImplicit);
+                    }
+                    CodeExecutableElement isImplicit = createIsImplicitTypeMethod(type);
+                    if (isImplicit != null) {
+                        clazz.add(isImplicit);
+                    }
+
+                    CodeExecutableElement typeIndex = createGetTypeIndex(type);
+                    if (typeIndex != null) {
+                        clazz.add(typeIndex);
+                    }
                 }
             }
 
@@ -133,6 +159,95 @@
             return field;
         }
 
+        private CodeExecutableElement createIsImplicitTypeMethod(TypeData type) {
+            TypeSystemData typeSystem = getModel();
+            List<ImplicitCastData> casts = typeSystem.lookupByTargetType(type);
+            if (casts.isEmpty()) {
+                return null;
+            }
+            CodeExecutableElement method = new CodeExecutableElement(modifiers(PUBLIC), getContext().getType(boolean.class), TypeSystemCodeGenerator.isImplicitTypeMethodName(type));
+            method.addParameter(new CodeVariableElement(getContext().getType(Object.class), LOCAL_VALUE));
+            CodeTreeBuilder builder = method.createBuilder();
+
+            List<TypeData> sourceTypes = typeSystem.lookupSourceTypes(type);
+
+            builder.startReturn();
+            String sep = "";
+            for (TypeData sourceType : sourceTypes) {
+                builder.string(sep);
+                builder.startCall(isTypeMethodName(sourceType)).string(LOCAL_VALUE).end();
+                sep = " || ";
+            }
+            builder.end();
+            return method;
+        }
+
+        private CodeExecutableElement createAsImplicitTypeMethod(TypeData type) {
+            TypeSystemData typeSystem = getModel();
+            List<ImplicitCastData> casts = typeSystem.lookupByTargetType(type);
+            if (casts.isEmpty()) {
+                return null;
+            }
+            CodeExecutableElement method = new CodeExecutableElement(modifiers(PUBLIC), type.getPrimitiveType(), TypeSystemCodeGenerator.asImplicitTypeMethodName(type));
+            method.addParameter(new CodeVariableElement(getContext().getType(Object.class), LOCAL_VALUE));
+
+            List<TypeData> sourceTypes = typeSystem.lookupSourceTypes(type);
+
+            CodeTreeBuilder builder = method.createBuilder();
+            boolean elseIf = false;
+            for (TypeData sourceType : sourceTypes) {
+                elseIf = builder.startIf(elseIf);
+                builder.startCall(isTypeMethodName(sourceType)).string(LOCAL_VALUE).end();
+                builder.end().startBlock();
+
+                builder.startReturn();
+                ImplicitCastData cast = typeSystem.lookupCast(sourceType, type);
+                if (cast != null) {
+                    builder.startCall(cast.getMethodName());
+                }
+                builder.startCall(asTypeMethodName(sourceType)).string(LOCAL_VALUE).end();
+                if (cast != null) {
+                    builder.end();
+                }
+                builder.end();
+                builder.end();
+            }
+
+            builder.startElseBlock();
+            builder.startStatement().startStaticCall(getContext().getTruffleTypes().getCompilerDirectives(), "transferToInterpreter").end().end();
+            builder.startThrow().startNew(getContext().getType(IllegalArgumentException.class)).doubleQuote("Illegal type ").end().end();
+            builder.end();
+            return method;
+        }
+
+        private CodeExecutableElement createGetTypeIndex(TypeData type) {
+            TypeSystemData typeSystem = getModel();
+            List<ImplicitCastData> casts = typeSystem.lookupByTargetType(type);
+            if (casts.isEmpty()) {
+                return null;
+            }
+            CodeExecutableElement method = new CodeExecutableElement(modifiers(PUBLIC), getContext().getType(Class.class), TypeSystemCodeGenerator.getImplicitClass(type));
+            method.addParameter(new CodeVariableElement(getContext().getType(Object.class), LOCAL_VALUE));
+
+            List<TypeData> sourceTypes = typeSystem.lookupSourceTypes(type);
+            CodeTreeBuilder builder = method.createBuilder();
+            boolean elseIf = false;
+            for (TypeData sourceType : sourceTypes) {
+                elseIf = builder.startIf(elseIf);
+                builder.startCall(isTypeMethodName(sourceType)).string(LOCAL_VALUE).end();
+                builder.end().startBlock();
+                builder.startReturn().typeLiteral(sourceType.getPrimitiveType()).end();
+                builder.end();
+            }
+
+            builder.startElseBlock();
+            builder.startStatement().startStaticCall(getContext().getTruffleTypes().getCompilerDirectives(), "transferToInterpreter").end().end();
+            builder.startThrow().startNew(getContext().getType(IllegalArgumentException.class)).doubleQuote("Illegal type ").end().end();
+            builder.end();
+
+            return method;
+        }
+
         private CodeExecutableElement createIsTypeMethod(TypeData type) {
             if (!type.getTypeChecks().isEmpty()) {
                 return null;
@@ -174,8 +289,8 @@
             method.addThrownType(getContext().getTruffleTypes().getUnexpectedValueException());
 
             CodeTreeBuilder body = method.createBuilder();
-            body.startIf().startCall(null, TypeSystemCodeGenerator.isTypeMethodName(expectedType)).string(LOCAL_VALUE).end().end().startBlock();
-            body.startReturn().startCall(null, TypeSystemCodeGenerator.asTypeMethodName(expectedType)).string(LOCAL_VALUE).end().end();
+            body.startIf().startCall(TypeSystemCodeGenerator.isTypeMethodName(expectedType)).string(LOCAL_VALUE).end().end().startBlock();
+            body.startReturn().startCall(TypeSystemCodeGenerator.asTypeMethodName(expectedType)).string(LOCAL_VALUE).end().end();
             body.end(); // if-block
             body.startThrow().startNew(getContext().getTruffleTypes().getUnexpectedValueException()).string(LOCAL_VALUE).end().end();
 
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/TypeSystemData.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/TypeSystemData.java	Sat Sep 07 12:23:40 2013 -0700
@@ -36,6 +36,7 @@
     private List<TypeMirror> primitiveTypeMirrors = new ArrayList<>();
     private List<TypeMirror> boxedTypeMirrors = new ArrayList<>();
 
+    private List<ImplicitCastData> implicitCasts;
     private List<TypeCastData> casts;
     private List<TypeCheckData> checks;
 
@@ -61,6 +62,14 @@
         }
     }
 
+    public void setImplicitCasts(List<ImplicitCastData> implicitCasts) {
+        this.implicitCasts = implicitCasts;
+    }
+
+    public List<ImplicitCastData> getImplicitCasts() {
+        return implicitCasts;
+    }
+
     public void setCasts(List<TypeCastData> casts) {
         this.casts = casts;
     }
@@ -89,6 +98,9 @@
         if (casts != null) {
             sinks.addAll(casts);
         }
+        if (implicitCasts != null) {
+            sinks.addAll(implicitCasts);
+        }
         return sinks;
     }
 
@@ -161,4 +173,55 @@
         return getClass().getSimpleName() + "[template = " + Utils.getSimpleName(getTemplateType()) + ", types = " + types + "]";
     }
 
+    public Set<TypeData> lookupCastSourceTypes() {
+        if (getImplicitCasts() == null) {
+            return null;
+        }
+
+        Set<TypeData> sourceTypes = new TreeSet<>();
+        for (ImplicitCastData cast : getImplicitCasts()) {
+            sourceTypes.add(cast.getSourceType());
+        }
+        return sourceTypes;
+    }
+
+    public List<ImplicitCastData> lookupByTargetType(TypeData targetType) {
+        if (getImplicitCasts() == null) {
+            return Collections.emptyList();
+        }
+        List<ImplicitCastData> foundCasts = new ArrayList<>();
+        for (ImplicitCastData cast : getImplicitCasts()) {
+            if (cast.getTargetType().equals(targetType)) {
+                foundCasts.add(cast);
+            }
+        }
+        return foundCasts;
+    }
+
+    public ImplicitCastData lookupCast(TypeData sourceType, TypeData targetType) {
+        if (getImplicitCasts() == null) {
+            return null;
+        }
+        for (ImplicitCastData cast : getImplicitCasts()) {
+            if (cast.getSourceType().equals(sourceType) && cast.getTargetType().equals(targetType)) {
+                return cast;
+            }
+        }
+        return null;
+    }
+
+    public List<TypeData> lookupSourceTypes(TypeData type) {
+        List<TypeData> sourceTypes = new ArrayList<>();
+        sourceTypes.add(type);
+        if (getImplicitCasts() != null) {
+            for (ImplicitCastData cast : getImplicitCasts()) {
+                if (cast.getTargetType() == type) {
+                    sourceTypes.add(cast.getSourceType());
+                }
+            }
+        }
+        Collections.sort(sourceTypes);
+        return sourceTypes;
+    }
+
 }
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/TypeSystemParser.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/TypeSystemParser.java	Sat Sep 07 12:23:40 2013 -0700
@@ -71,7 +71,7 @@
         }
 
         typeSystem.setTypes(parseTypes(typeSystem));
-        if (typeSystem.getTypes() == null) {
+        if (typeSystem.hasErrors()) {
             return typeSystem;
         }
 
@@ -91,6 +91,8 @@
         if (casts == null || checks == null || implicitCasts == null) {
             return typeSystem;
         }
+
+        typeSystem.setImplicitCasts(implicitCasts);
         typeSystem.setCasts(casts);
         typeSystem.setChecks(checks);
 
@@ -106,6 +108,7 @@
             cast.getTargetType().addTypeCast(cast);
         }
 
+        verifyImplicitCasts(typeSystem);
         verifyGenericTypeChecksAndCasts(typeSystem);
         verifyMethodSignatures(typeSystem);
         verifyNamesUnique(typeSystem);
@@ -113,6 +116,24 @@
         return typeSystem;
     }
 
+    private static void verifyImplicitCasts(TypeSystemData typeSystem) {
+        Set<TypeData> types = new HashSet<>();
+        Set<TypeData> duplicateSourceTypes = new HashSet<>();
+        for (ImplicitCastData cast : typeSystem.getImplicitCasts()) {
+            if (types.contains(cast.getSourceType())) {
+                duplicateSourceTypes.add(cast.getSourceType());
+            }
+            types.add(cast.getSourceType());
+        }
+        for (TypeData duplicateType : duplicateSourceTypes) {
+            for (ImplicitCastData cast : typeSystem.getImplicitCasts()) {
+                if (cast.getSourceType().equals(duplicateType)) {
+                    cast.addError("Duplicate cast source type %s.", Utils.getSimpleName(duplicateType.getPrimitiveType()), ImplicitCast.class.getSimpleName());
+                }
+            }
+        }
+    }
+
     private static void verifyGenericTypeChecksAndCasts(TypeSystemData typeSystem) {
         for (TypeData type : typeSystem.getTypes()) {
             if (!type.getTypeChecks().isEmpty()) {
--- a/graal/com.oracle.truffle.sl/src/com/oracle/truffle/sl/SLTypes.java	Tue Sep 03 16:48:17 2013 -0700
+++ b/graal/com.oracle.truffle.sl/src/com/oracle/truffle/sl/SLTypes.java	Sat Sep 07 12:23:40 2013 -0700
@@ -69,4 +69,9 @@
     public boolean isBigInteger(@SuppressWarnings("unused") int value) {
         return true;
     }
+
+    @ImplicitCast
+    public BigInteger castBigInteger(int value) {
+        return BigInteger.valueOf(value);
+    }
 }
--- a/mx/.pylintrc	Tue Sep 03 16:48:17 2013 -0700
+++ b/mx/.pylintrc	Sat Sep 07 12:23:40 2013 -0700
@@ -15,7 +15,7 @@
 ignore=CVS
 
 # Pickle collected data for later comparisons.
-persistent=yes
+persistent=no
 
 # List of plugins (as comma separated values of python modules names) to load,
 # usually to register additional checkers.
@@ -41,7 +41,10 @@
 disable=attribute-defined-outside-init,arguments-differ,
         bare-except,global-statement,protected-access,redefined-outer-name,
         unused-argument,star-args,pointless-string-statement,old-style-class,
-        too-many-lines,missing-docstring,no-init
+        too-many-lines,missing-docstring,no-init,no-self-use,too-many-statements,
+        too-many-locals,too-few-public-methods,too-many-instance-attributes,
+        too-many-arguments,too-many-branches,too-many-public-methods,
+        abstract-method
 
 [REPORTS]
 
@@ -92,10 +95,10 @@
 class-rgx=[A-Z_][a-zA-Z0-9]+$
 
 # Regular expression which should only match correct function names
-function-rgx=[a-z_][a-zA-Z0-9_]{1,30}$
+function-rgx=[a-z_][a-zA-Z0-9_]{1,40}$
 
 # Regular expression which should only match correct method names
-method-rgx=[a-z_][a-zA-Z0-9_]{2,30}$
+method-rgx=[a-z_][a-zA-Z0-9_]{2,40}$
 
 # Regular expression which should only match correct instance attribute names
 attr-rgx=[a-z_][a-zA-Z0-9_]{1,30}$
--- a/mx/commands.py	Tue Sep 03 16:48:17 2013 -0700
+++ b/mx/commands.py	Sat Sep 07 12:23:40 2013 -0700
@@ -442,6 +442,38 @@
         log.close()
     return ret
 
+def pylint(args):
+    """run pylint (if available) over Python source files"""
+    rcfile = join(_graal_home, 'mx', '.pylintrc')
+    if not exists(rcfile):
+        mx.log('pylint configuration file does not exist: ' + rcfile)
+        return
+
+    try:
+        output = subprocess.check_output(['pylint', '--version'], stderr=subprocess.STDOUT)
+        m = re.match(r'.*pylint (\d+)\.(\d+)\.(\d+).*', output, re.DOTALL)
+        if not m:
+            mx.log('could not determine pylint version from ' + output)
+            return
+        major, minor, micro = (int(m.group(1)), int(m.group(2)), int(m.group(3)))
+        if major < 1:
+            mx.log('require pylint version >= 1 (got {0}.{1}.{2})'.format(major, minor, micro))
+            return
+    except BaseException:
+        mx.log('pylint is not available')
+        return
+
+
+    env = os.environ.copy()
+    env['PYTHONPATH'] = dirname(mx.__file__)
+
+    versioned = subprocess.check_output(['hg', 'locate', '-f'], stderr=subprocess.STDOUT).split(os.linesep)
+    for f in versioned:
+        if f.endswith('.py'):
+            pyfile = f
+            mx.log('Running pylint on ' + pyfile + '...')
+            mx.run(['pylint', '--reports=n', '--rcfile=' + rcfile, pyfile], env=env)
+
 def jdkhome(args, vm=None):
     """print the JDK directory selected for the 'vm' command"""
 
@@ -523,8 +555,8 @@
         assert vm == 'graal', vm
         buildSuffix = 'graal'
 
-    if _installed_jdks:
-        if not mx.ask_yes_no("You are going to build because --installed-jdks is set (" + _installed_jdks + ") - are you sure you want to continue", 'n'):
+    if _installed_jdks and _installed_jdks != _graal_home:
+        if not mx.ask_yes_no("Warning: building while --installed-jdks is set (" + _installed_jdks + ") is not recommanded - are you sure you want to continue", 'n'):
             mx.abort(1)
 
     for build in builds:
@@ -949,6 +981,10 @@
     total = Task('Gate')
     try:
 
+        t = Task('Pylint')
+        pylint([])
+        tasks.append(t.stop())
+
         t = Task('Clean')
         cleanArgs = []
         if not args.cleanNative:
@@ -1207,7 +1243,6 @@
         if len(parts) > 1:
             assert len(parts) == 2
             group = parts[0]
-            print group
             availableBenchmarks.add(group)
 
     _run_benchmark(args, sorted(availableBenchmarks), launcher)
@@ -1333,6 +1368,7 @@
         'hcfdis': [hcfdis, ''],
         'igv' : [igv, ''],
         'jdkhome': [jdkhome, ''],
+        'pylint': [pylint, ''],
         'dacapo': [dacapo, '[VM options] benchmarks...|"all" [DaCapo options]'],
         'scaladacapo': [scaladacapo, '[VM options] benchmarks...|"all" [Scala DaCapo options]'],
         'specjvm2008': [specjvm2008, '[VM options] benchmarks...|"all" [SPECjvm2008 options]'],
--- a/mx/sanitycheck.py	Tue Sep 03 16:48:17 2013 -0700
+++ b/mx/sanitycheck.py	Sat Sep 07 12:23:40 2013 -0700
@@ -77,7 +77,8 @@
     'sunflow':    [           'fastdebug', 'debug'],
     'tomcat':     ['product', 'fastdebug', 'debug'],
     'tradebeans': ['product', 'fastdebug', 'debug'],
-    'tradesoap':  ['product'],
+    # tradesoap is too unreliable for the gate, often crashing with "java.net.BindException: Address already in use"
+    'tradesoap':  [                               ],
     'xalan':      ['product', 'fastdebug', 'debug'],
 }
 
--- a/mxtool/.pylintrc	Tue Sep 03 16:48:17 2013 -0700
+++ /dev/null	Thu Jan 01 00:00:00 1970 +0000
@@ -1,275 +0,0 @@
-[MASTER]
-
-# Specify a configuration file.
-#rcfile=
-
-# Python code to execute, usually for sys.path manipulation such as
-# pygtk.require().
-#init-hook=
-
-# Profiled execution.
-profile=no
-
-# Add files or directories to the blacklist. They should be base names, not
-# paths.
-ignore=CVS
-
-# Pickle collected data for later comparisons.
-persistent=yes
-
-# List of plugins (as comma separated values of python modules names) to load,
-# usually to register additional checkers.
-load-plugins=
-
-
-[MESSAGES CONTROL]
-
-# Enable the message, report, category or checker with the given id(s). You can
-# either give multiple identifier separated by comma (,) or put this option
-# multiple time. See also the "--disable" option for examples.
-#enable=
-
-# Disable the message, report, category or checker with the given id(s). You
-# can either give multiple identifiers separated by comma (,) or put this
-# option multiple times (only on the command line, not in the configuration
-# file where it should appear only once).You can also use "--disable=all" to
-# disable everything first and then reenable specific checks. For example, if
-# you want to run only the similarities checker, you can use "--disable=all
-# --enable=similarities". If you want to run only the classes checker, but have
-# no Warning level messages displayed, use"--disable=all --enable=classes
-# --disable=W"
-disable=attribute-defined-outside-init,
-        bare-except,global-statement,protected-access,
-        redefined-outer-name,unused-argument,star-args,
-        old-style-class,too-many-lines,pointless-string-statement,
-
-[REPORTS]
-
-# Set the output format. Available formats are text, parseable, colorized, msvs
-# (visual studio) and html. You can also give a reporter class, eg
-# mypackage.mymodule.MyReporterClass.
-output-format=text
-
-# Put messages in a separate file for each module / package specified on the
-# command line instead of printing them on stdout. Reports (if any) will be
-# written in a file name "pylint_global.[txt|html]".
-files-output=no
-
-# Tells whether to display a full report or only the messages
-reports=no
-
-# Python expression which should return a note less than 10 (10 is the highest
-# note). You have access to the variables errors warning, statement which
-# respectively contain the number of errors / warnings messages and the total
-# number of statements analyzed. This is used by the global evaluation report
-# (RP0004).
-evaluation=10.0 - ((float(5 * error + warning + refactor + convention) / statement) * 10)
-
-# Add a comment according to your evaluation note. This is used by the global
-# evaluation report (RP0004).
-comment=no
-
-# Template used to display messages. This is a python new-style format string
-# used to format the massage information. See doc for all details
-#msg-template=
-
-
-[BASIC]
-
-# Required attributes for module, separated by a comma
-required-attributes=
-
-# List of builtins function names that should not be used, separated by a comma
-bad-functions=filter,apply,input
-
-# Regular expression which should only match correct module names
-module-rgx=(([a-z_][a-z0-9_]*)|([A-Z][a-zA-Z0-9]+))$
-
-# Regular expression which should only match correct module level names
-const-rgx=_[a-zA-Z0-9_]{2,30}$
-
-# Regular expression which should only match correct class names
-class-rgx=[A-Z_][a-zA-Z0-9]+$
-
-# Regular expression which should only match correct function names
-function-rgx=[a-z_][a-zA-Z0-9_]{1,50}$
-
-# Regular expression which should only match correct method names
-method-rgx=[a-z_][a-zA-Z0-9_]{1,50}$
-
-# Regular expression which should only match correct instance attribute names
-attr-rgx=[a-z_][a-zA-Z0-9_]{2,30}$
-
-# Regular expression which should only match correct argument names
-argument-rgx=[a-z_][a-zA-Z0-9_]{0,30}$
-
-# Regular expression which should only match correct variable names
-variable-rgx=[a-z_][a-zA-Z0-9_]{0,30}$
-
-# Regular expression which should only match correct attribute names in class
-# bodies
-class-attribute-rgx=([A-Za-z_][A-Za-z0-9_]{2,30}|(__.*__))$
-
-# Regular expression which should only match correct list comprehension /
-# generator expression variable names
-inlinevar-rgx=[A-Za-z_][A-Za-z0-9_]*$
-
-# Good variable names which should always be accepted, separated by a comma
-good-names=i,j,k,ex,Run,_
-
-# Bad variable names which should always be refused, separated by a comma
-bad-names=foo,bar,baz,toto,tutu,tata
-
-# Regular expression which should only match function or class names that do
-# not require a docstring.
-no-docstring-rgx=.*
-
-# Minimum line length for functions/classes that require docstrings, shorter
-# ones are exempt.
-docstring-min-length=-1
-
-
-[FORMAT]
-
-# Maximum number of characters on a single line.
-max-line-length=300
-
-# Regexp for a line that is allowed to be longer than the limit.
-ignore-long-lines=^\s*(# )?<?https?://\S+>?$
-
-# Maximum number of lines in a module
-max-module-lines=1000
-
-# String used as indentation unit. This is usually " " (4 spaces) or "\t" (1
-# tab).
-indent-string='    '
-
-
-[MISCELLANEOUS]
-
-# List of note tags to take in consideration, separated by a comma.
-notes=FIXME,XXX,TODO
-
-
-[SIMILARITIES]
-
-# Minimum lines number of a similarity.
-min-similarity-lines=4
-
-# Ignore comments when computing similarities.
-ignore-comments=yes
-
-# Ignore docstrings when computing similarities.
-ignore-docstrings=yes
-
-# Ignore imports when computing similarities.
-ignore-imports=no
-
-
-[TYPECHECK]
-
-# Tells whether missing members accessed in mixin class should be ignored. A
-# mixin class is detected if its name ends with "mixin" (case insensitive).
-ignore-mixin-members=yes
-
-# List of classes names for which member attributes should not be checked
-# (useful for classes with attributes dynamically set).
-ignored-classes=SQLObject
-
-# When zope mode is activated, add a predefined set of Zope acquired attributes
-# to generated-members.
-zope=no
-
-# List of members which are set dynamically and missed by pylint inference
-# system, and so shouldn't trigger E0201 when accessed. Python regular
-# expressions are accepted.
-generated-members=REQUEST,acl_users,aq_parent
-
-
-[VARIABLES]
-
-# Tells whether we should check for unused import in __init__ files.
-init-import=no
-
-# A regular expression matching the beginning of the name of dummy variables
-# (i.e. not used).
-dummy-variables-rgx=_$|dummy
-
-# List of additional names supposed to be defined in builtins. Remember that
-# you should avoid to define new builtins when possible.
-additional-builtins=
-
-
-[CLASSES]
-
-# List of interface methods to ignore, separated by a comma. This is used for
-# instance to not check methods defines in Zope's Interface base class.
-ignore-iface-methods=isImplementedBy,deferred,extends,names,namesAndDescriptions,queryDescriptionFor,getBases,getDescriptionFor,getDoc,getName,getTaggedValue,getTaggedValueTags,isEqualOrExtendedBy,setTaggedValue,isImplementedByInstancesOf,adaptWith,is_implemented_by
-
-# List of method names used to declare (i.e. assign) instance attributes.
-defining-attr-methods=__init__,__new__,setUp
-
-# List of valid names for the first argument in a class method.
-valid-classmethod-first-arg=cls
-
-# List of valid names for the first argument in a metaclass class method.
-valid-metaclass-classmethod-first-arg=mcs
-
-
-[DESIGN]
-
-# Maximum number of arguments for function / method
-max-args=5
-
-# Argument names that match this expression will be ignored. Default to name
-# with leading underscore
-ignored-argument-names=_.*
-
-# Maximum number of locals for function / method body
-max-locals=15
-
-# Maximum number of return / yield for function / method body
-max-returns=6
-
-# Maximum number of branch for function / method body
-max-branches=12
-
-# Maximum number of statements in function / method body
-max-statements=50
-
-# Maximum number of parents for a class (see R0901).
-max-parents=7
-
-# Maximum number of attributes for a class (see R0902).
-max-attributes=7
-
-# Minimum number of public methods for a class (see R0903).
-min-public-methods=2
-
-# Maximum number of public methods for a class (see R0904).
-max-public-methods=20
-
-
-[IMPORTS]
-
-# Deprecated modules which should not be used, separated by a comma
-deprecated-modules=regsub,TERMIOS,Bastion,rexec
-
-# Create a graph of every (i.e. internal and external) dependencies in the
-# given file (report RP0402 must not be disabled)
-import-graph=
-
-# Create a graph of external dependencies in the given file (report RP0402 must
-# not be disabled)
-ext-import-graph=
-
-# Create a graph of internal dependencies in the given file (report RP0402 must
-# not be disabled)
-int-import-graph=
-
-
-[EXCEPTIONS]
-
-# Exceptions that will emit a warning when being caught. Defaults to
-# "Exception"
-overgeneral-exceptions=Exception
--- a/src/gpu/ptx/vm/gpu_ptx.cpp	Tue Sep 03 16:48:17 2013 -0700
+++ b/src/gpu/ptx/vm/gpu_ptx.cpp	Sat Sep 07 12:23:40 2013 -0700
@@ -38,6 +38,7 @@
 gpu::Ptx::cuda_cu_ctx_create_func_t gpu::Ptx::_cuda_cu_ctx_create;
 gpu::Ptx::cuda_cu_ctx_destroy_func_t gpu::Ptx::_cuda_cu_ctx_destroy;
 gpu::Ptx::cuda_cu_ctx_synchronize_func_t gpu::Ptx::_cuda_cu_ctx_synchronize;
+gpu::Ptx::cuda_cu_ctx_set_current_func_t gpu::Ptx::_cuda_cu_ctx_set_current;
 gpu::Ptx::cuda_cu_device_get_count_func_t gpu::Ptx::_cuda_cu_device_get_count;
 gpu::Ptx::cuda_cu_device_get_name_func_t gpu::Ptx::_cuda_cu_device_get_name;
 gpu::Ptx::cuda_cu_device_get_func_t gpu::Ptx::_cuda_cu_device_get;
@@ -87,7 +88,7 @@
     tty->print_cr("Failed to initialize CUDA device");
     return false;
   }
- 
+
   if (TraceGPUInteraction) {
     tty->print_cr("CUDA driver initialization: Success");
   }
@@ -108,7 +109,7 @@
   if (TraceGPUInteraction) {
     tty->print_cr("[CUDA] Number of compute-capable devices found: %d", device_count);
   }
-  
+
   /* Get the handle to the first compute device */
   int device_id = 0;
   /* Compute-capable device handle */
@@ -195,12 +196,6 @@
   jit_options[2] = GRAAL_CU_JIT_MAX_REGISTERS;
   jit_option_values[2] = (void *)(size_t)jit_register_count;
 
-  if (TraceGPUInteraction) {
-    tty->print_cr("[CUDA] PTX Kernel\n%s", code);
-    tty->print_cr("[CUDA] Function name : %s", name);
-
-  }
-
   /* Create CUDA context to compile and execute the kernel */
   int status = _cuda_cu_ctx_create(&_device_context, 0, _cu_device);
 
@@ -213,6 +208,23 @@
     tty->print_cr("[CUDA] Success: Created context for device: %d", _cu_device);
   }
 
+  status = _cuda_cu_ctx_set_current(_device_context);
+
+  if (status != GRAAL_CUDA_SUCCESS) {
+    tty->print_cr("[CUDA] Failed to set current context for device: %d", _cu_device);
+    return NULL;
+  }
+
+  if (TraceGPUInteraction) {
+    tty->print_cr("[CUDA] Success: Set current context for device: %d", _cu_device);
+  }
+
+  if (TraceGPUInteraction) {
+    tty->print_cr("[CUDA] PTX Kernel\n%s", code);
+    tty->print_cr("[CUDA] Function name : %s", name);
+
+  }
+
   /* Load module's data with compiler options */
   status = _cuda_cu_module_load_data_ex(&cu_module, (void*) code, jit_num_options,
                                             jit_options, (void **)jit_option_values);
@@ -220,7 +232,7 @@
     if (status == GRAAL_CUDA_ERROR_NO_BINARY_FOR_GPU) {
       tty->print_cr("[CUDA] Check for malformed PTX kernel or incorrect PTX compilation options");
     }
-    tty->print_cr("[CUDA] *** Error (%d) Failed to load module data with online compiler options for method %s", 
+    tty->print_cr("[CUDA] *** Error (%d) Failed to load module data with online compiler options for method %s",
                   status, name);
     return NULL;
   }
@@ -255,7 +267,7 @@
   unsigned int blockX = 1;
   unsigned int blockY = 1;
   unsigned int blockZ = 1;
-  
+
   struct CUfunc_st* cu_function = (struct CUfunc_st*) kernel;
 
   void * config[5] = {
@@ -366,7 +378,7 @@
   if (cuda_library_name != NULL) {
     char *buffer = (char*)malloc(STD_BUFFER_SIZE);
     void *handle = os::dll_load(cuda_library_name, buffer, STD_BUFFER_SIZE);
-	free(buffer);
+        free(buffer);
     if (handle != NULL) {
       _cuda_cu_init =
         CAST_TO_FN_PTR(cuda_cu_init_func_t, os::dll_lookup(handle, "cuInit"));
@@ -376,6 +388,8 @@
         CAST_TO_FN_PTR(cuda_cu_ctx_destroy_func_t, os::dll_lookup(handle, "cuCtxDestroy"));
       _cuda_cu_ctx_synchronize =
         CAST_TO_FN_PTR(cuda_cu_ctx_synchronize_func_t, os::dll_lookup(handle, "cuCtxSynchronize"));
+      _cuda_cu_ctx_set_current =
+        CAST_TO_FN_PTR(cuda_cu_ctx_set_current_func_t, os::dll_lookup(handle, "cuCtxSetCurrent"));
       _cuda_cu_device_get_count =
         CAST_TO_FN_PTR(cuda_cu_device_get_count_func_t, os::dll_lookup(handle, "cuDeviceGetCount"));
       _cuda_cu_device_get_name =
@@ -416,4 +430,3 @@
   tty->print_cr("Failed to find CUDA linkage");
   return false;
 }
-
--- a/src/gpu/ptx/vm/gpu_ptx.hpp	Tue Sep 03 16:48:17 2013 -0700
+++ b/src/gpu/ptx/vm/gpu_ptx.hpp	Sat Sep 07 12:23:40 2013 -0700
@@ -87,6 +87,7 @@
   typedef int (*cuda_cu_ctx_create_func_t)(void*, int, int);
   typedef int (*cuda_cu_ctx_destroy_func_t)(void*);
   typedef int (*cuda_cu_ctx_synchronize_func_t)(void);
+  typedef int (*cuda_cu_ctx_set_current_func_t)(void*);
   typedef int (*cuda_cu_device_get_count_func_t)(int*);
   typedef int (*cuda_cu_device_get_name_func_t)(char*, int, int);
   typedef int (*cuda_cu_device_get_func_t)(int*, int);
@@ -98,7 +99,7 @@
                                               unsigned int, void*, void**, void**);
   typedef int (*cuda_cu_module_get_function_func_t)(void*, void*, const char*);
   typedef int (*cuda_cu_module_load_data_ex_func_t)(void*, void*, unsigned int, void*, void**);
-  typedef int (*cuda_cu_memalloc_func_t)(void*, unsigned int);
+  typedef int (*cuda_cu_memalloc_func_t)(void*, size_t);
   typedef int (*cuda_cu_memfree_func_t)(gpu::Ptx::CUdeviceptr);
   typedef int (*cuda_cu_memcpy_htod_func_t)(gpu::Ptx::CUdeviceptr, const void*, unsigned int);
   typedef int (*cuda_cu_memcpy_dtoh_func_t)(const void*, gpu::Ptx::CUdeviceptr,  unsigned int);
@@ -120,6 +121,7 @@
   static cuda_cu_memfree_func_t                   _cuda_cu_memfree;
   static cuda_cu_memcpy_htod_func_t               _cuda_cu_memcpy_htod;
   static cuda_cu_memcpy_dtoh_func_t               _cuda_cu_memcpy_dtoh;
+  static cuda_cu_ctx_set_current_func_t           _cuda_cu_ctx_set_current;
 
 protected:
   static void* _device_context;
--- a/src/gpu/ptx/vm/kernelArguments.hpp	Tue Sep 03 16:48:17 2013 -0700
+++ b/src/gpu/ptx/vm/kernelArguments.hpp	Sat Sep 07 12:23:40 2013 -0700
@@ -56,6 +56,7 @@
     _args = args;
     _success = true;
     _bufferOffset = 0;
+    _return_value_ptr = 0;
     if (!is_static) {
       // TODO : Create a device argument for receiver object and add it to _kernelBuffer
       tty->print_cr("{CUDA] ****** TODO: Support for execution of non-static java methods not implemented yet.");
--- a/src/share/vm/graal/graalCompilerToVM.cpp	Tue Sep 03 16:48:17 2013 -0700
+++ b/src/share/vm/graal/graalCompilerToVM.cpp	Sat Sep 07 12:23:40 2013 -0700
@@ -66,7 +66,7 @@
   int code_size = method->code_size();
   jbyte* reconstituted_code = NEW_RESOURCE_ARRAY(jbyte, code_size);
 
-  bool is_rewritten = method->method_holder()->is_rewritten();
+  guarantee(method->method_holder()->is_rewritten(), "Method's holder should be rewritten");
   // iterate over all bytecodes and replace non-Java bytecodes
 
   for (BytecodeStream s(method); s.next() != Bytecodes::_illegal; ) {
@@ -81,7 +81,7 @@
       memcpy(&reconstituted_code[bci+1], s.bcp()+1, len-1);
     }
 
-    if (is_rewritten && len > 1) {
+    if (len > 1) {
       // Restore the big-endian constant pool indexes.
       // Cf. Rewriter::scan_method
       switch (code) {
@@ -554,6 +554,12 @@
   return InstanceKlass::cast(klass)->is_initialized();
 C2V_END
 
+C2V_VMENTRY(jboolean, isTypeLinked,(JNIEnv *, jobject, jobject hotspot_klass))
+  Klass* klass = java_lang_Class::as_Klass(HotSpotResolvedObjectType::javaMirror(hotspot_klass));
+  assert(klass != NULL, "method must not be called for primitive types");
+  return InstanceKlass::cast(klass)->is_linked();
+C2V_END
+
 C2V_VMENTRY(jboolean, hasFinalizableSubclass,(JNIEnv *, jobject, jobject hotspot_klass))
   Klass* klass = java_lang_Class::as_Klass(HotSpotResolvedObjectType::javaMirror(hotspot_klass));
   assert(klass != NULL, "method must not be called for primitive types");
@@ -1198,6 +1204,7 @@
   {CC"getInstanceFields",             CC"("HS_RESOLVED_TYPE")["HS_RESOLVED_FIELD,                       FN_PTR(getInstanceFields)},
   {CC"getMethods",                    CC"("HS_RESOLVED_TYPE")["HS_RESOLVED_METHOD,                      FN_PTR(getMethods)},
   {CC"isTypeInitialized",             CC"("HS_RESOLVED_TYPE")Z",                                        FN_PTR(isTypeInitialized)},
+  {CC"isTypeLinked",                  CC"("HS_RESOLVED_TYPE")Z",                                        FN_PTR(isTypeLinked)},
   {CC"hasFinalizableSubclass",        CC"("HS_RESOLVED_TYPE")Z",                                        FN_PTR(hasFinalizableSubclass)},
   {CC"initializeType",                CC"("HS_RESOLVED_TYPE")V",                                        FN_PTR(initializeType)},
   {CC"getMaxCallTargetOffset",        CC"(J)J",                                                         FN_PTR(getMaxCallTargetOffset)},