Mercurial > hg > truffle

changeset 11494:96e4e5333a25

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
Merge.
author Doug Simon <doug.simon@oracle.com>
date Sat, 31 Aug 2013 09:18:58 +0200
parents 94779c895aad (current diff) 49bb1bc983c6 (diff)
children 3662471dcfaa
files graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotRuntime.java
diffstat 61 files changed, 2446 insertions(+), 604 deletions(-) [+]
line wrap: on
line diff
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CodeCacheProvider.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CodeCacheProvider.java	Sat Aug 31 09:18:58 2013 +0200
@@ -89,12 +89,4 @@
      * Gets a description of the target architecture.
      */
     TargetDescription getTarget();
-
-    /**
-     * Returns the register the runtime uses for maintaining the heap base address. This is mainly
-     * utilized by runtimes which support compressed pointers.
-     * 
-     * @return the register that keeps the heap base address
-     */
-    Register heapBaseRegister();
 }
--- a/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/Constant.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/Constant.java	Sat Aug 31 09:18:58 2013 +0200
@@ -454,4 +454,102 @@
                 throw new IllegalArgumentException(kind.toString());
         }
     }
+
+    /**
+     * Returns the zero value for a given numeric kind.
+     */
+    public static Constant zero(Kind kind) {
+        switch (kind) {
+            case Byte:
+                return forByte((byte) 0);
+            case Char:
+                return forChar((char) 0);
+            case Double:
+                return DOUBLE_0;
+            case Float:
+                return FLOAT_0;
+            case Int:
+                return INT_0;
+            case Long:
+                return LONG_0;
+            case Short:
+                return forShort((short) 0);
+            default:
+                throw new IllegalArgumentException(kind.toString());
+        }
+    }
+
+    /**
+     * Returns the one value for a given numeric kind.
+     */
+    public static Constant one(Kind kind) {
+        switch (kind) {
+            case Byte:
+                return forByte((byte) 1);
+            case Char:
+                return forChar((char) 1);
+            case Double:
+                return DOUBLE_1;
+            case Float:
+                return FLOAT_1;
+            case Int:
+                return INT_1;
+            case Long:
+                return LONG_1;
+            case Short:
+                return forShort((short) 1);
+            default:
+                throw new IllegalArgumentException(kind.toString());
+        }
+    }
+
+    /**
+     * Adds two numeric constants.
+     */
+    public static Constant add(Constant x, Constant y) {
+        assert x.getKind() == y.getKind();
+        switch (x.getKind()) {
+            case Byte:
+                return forByte((byte) (x.asInt() + y.asInt()));
+            case Char:
+                return forChar((char) (x.asInt() + y.asInt()));
+            case Double:
+                return forDouble(x.asDouble() + y.asDouble());
+            case Float:
+                return forFloat(x.asFloat() + y.asFloat());
+            case Int:
+                return forInt(x.asInt() + y.asInt());
+            case Long:
+                return forLong(x.asLong() + y.asLong());
+            case Short:
+                return forShort((short) (x.asInt() + y.asInt()));
+            default:
+                throw new IllegalArgumentException(x.getKind().toString());
+        }
+    }
+
+    /**
+     * Multiplies two numeric constants.
+     */
+    public static Constant mul(Constant x, Constant y) {
+        assert x.getKind() == y.getKind();
+        switch (x.getKind()) {
+            case Byte:
+                return forByte((byte) (x.asInt() * y.asInt()));
+            case Char:
+                return forChar((char) (x.asInt() * y.asInt()));
+            case Double:
+                return forDouble(x.asDouble() * y.asDouble());
+            case Float:
+                return forFloat(x.asFloat() * y.asFloat());
+            case Int:
+                return forInt(x.asInt() * y.asInt());
+            case Long:
+                return forLong(x.asLong() * y.asLong());
+            case Short:
+                return forShort((short) (x.asInt() * y.asInt()));
+            default:
+                throw new IllegalArgumentException(x.getKind().toString());
+        }
+    }
 }
--- a/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java	Sat Aug 31 09:18:58 2013 +0200
@@ -272,59 +272,69 @@
     }
 
     public final void ld_global_b8(Register d, Register a, long immOff) {
-        emitString("ld.global.b8" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.b8" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_b16(Register d, Register a, long immOff) {
-        emitString("ld.global.b16" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.b16" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_b32(Register d, Register a, long immOff) {
-        emitString("ld.global.b32" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.b32" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_b64(Register d, Register a, long immOff) {
-        emitString("ld.global.b64" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.b64" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_u8(Register d, Register a, long immOff) {
-        emitString("ld.global.u8" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.u8" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_u16(Register d, Register a, long immOff) {
-        emitString("ld.global.u16" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.u16" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_u32(Register d, Register a, long immOff) {
-        emitString("ld.global.u32" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.u32" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_u64(Register d, Register a, long immOff) {
-        emitString("ld.global.u64" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.u64" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_s8(Register d, Register a, long immOff) {
-        emitString("ld.global.s8" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.s8" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_s16(Register d, Register a, long immOff) {
-        emitString("ld.global.s16" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.s16" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_s32(Register d, Register a, long immOff) {
-        emitString("ld.global.s32" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.s32" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_s64(Register d, Register a, long immOff) {
-        emitString("ld.global.s64" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.s64" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_f32(Register d, Register a, long immOff) {
-        emitString("ld.global.f32" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.f32" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + immOff + "]" + ";" + "");
     }
 
     public final void ld_global_f64(Register d, Register a, long immOff) {
-        emitString("ld.global.f64" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + "");
+        emitString("ld.global.f64" + " " + "%r" + d.encoding() + ", [" + a.toString() + " + " + 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 + "]" + ";" + "");
+    }
+
+    // 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 + "]" + ";" + "");
     }
 
     public final void mov_b16(Register d, Register a) {
@@ -539,6 +549,18 @@
         emitString("or.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b64 + ";" + "");
     }
 
+    public final void param_8_decl(Register d, boolean lastParam) {
+        emitString(".param" + " " + ".s8" + " " + d.toString() + (lastParam ? "" : ","));
+    }
+
+    public final void param_32_decl(Register d, boolean lastParam) {
+        emitString(".param" + " " + ".s32" + " " + d.toString() + (lastParam ? "" : ","));
+    }
+
+    public final void param_64_decl(Register d, boolean lastParam) {
+        emitString(".param" + " " + ".s64" + " " + d.toString() + (lastParam ? "" : ","));
+    }
+
     public final void popc_b32(Register d, Register a) {
         emitString("popc.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + "");
     }
@@ -923,6 +945,8 @@
         emitString("shr.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + "");
     }
 
+    // Store in global state space
+
     public final void st_global_b8(Register a, long immOff, Register b) {
         emitString("st.global.b8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
     }
@@ -979,6 +1003,37 @@
         emitString("st.global.f64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
     }
 
+    // Store return value
+    public final void st_global_return_value_s8(Register a, long immOff, Register b) {
+        emitString("st.global.s8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
+    }
+
+    public final void st_global_return_value_s32(Register a, long immOff, Register b) {
+        emitString("st.global.s32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
+    }
+
+    public final void st_global_return_value_s64(Register a, long immOff, Register b) {
+        emitString("st.global.s64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
+    }
+
+    public final void st_global_return_value_f32(Register a, long immOff, Register b) {
+        emitString("st.global.f32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
+    }
+
+    public final void st_global_return_value_f64(Register a, long immOff, Register b) {
+        emitString("st.global.f64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
+    }
+
+    public final void st_global_return_value_u32(Register a, long immOff, Register b) {
+        emitString("st.global.u32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
+    }
+
+    public final void st_global_return_value_u64(Register a, long immOff, Register b) {
+        emitString("st.global.u64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + "");
+    }
+
+    // Subtract instruction
+
     public final void sub_f32(Register d, Register a, Register b) {
         emitString("sub.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + "");
     }
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ArrayTest.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ArrayTest.java	Sat Aug 31 09:18:58 2013 +0200
@@ -24,10 +24,11 @@
 
 import java.lang.reflect.Method;
 
-import org.junit.Test;
+import org.junit.*;
 
 public class ArrayTest extends PTXTestBase {
 
+    @Ignore
     @Test
     public void testArray() {
         compile("testArray1I");
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/BasicPTXTest.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/BasicPTXTest.java	Sat Aug 31 09:18:58 2013 +0200
@@ -41,7 +41,7 @@
         invoke(compile("testConstI"));
     }
 
-    public int testConstI() {
+    public static int testConstI() {
         return 42;
     }
 
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ControlTest.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ControlTest.java	Sat Aug 31 09:18:58 2013 +0200
@@ -22,12 +22,13 @@
  */
 package com.oracle.graal.compiler.ptx.test;
 
-import org.junit.Test;
+import org.junit.*;
 
 import java.lang.reflect.Method;
 
 public class ControlTest extends PTXTestBase {
 
+    @Ignore
     @Test
     public void testControl() {
         compile("testSwitch1I");
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/FloatPTXTest.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/FloatPTXTest.java	Sat Aug 31 09:18:58 2013 +0200
@@ -26,18 +26,38 @@
 
 import org.junit.*;
 
+import com.oracle.graal.api.code.CompilationResult;
 
 /* PTX ISA 3.1 - 8.7.3 Floating-Point Instructions */
 public class FloatPTXTest extends PTXTestBase {
 
     @Test
     public void testAdd() {
-        compile("testAdd2F");
-        compile("testAdd2D");
-        compile("testAddFConst");
-        compile("testAddConstF");
-        compile("testAddDConst");
-        compile("testAddConstD");
+        CompilationResult r = compile("testAdd2F");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testAdd2F FAILED");
+        }
+        r = compile("testAdd2D");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testAdd2D FAILED");
+        }
+
+        r = compile("testAddFConst");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testAddFConst FAILED");
+        }
+        r = compile("testAddConstF");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testConstF FAILED");
+        }
+        r = compile("testAddDConst");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testAddDConst FAILED");
+        }
+        r = compile("testAddConstD");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testConstD FAILED");
+        }
     }
 
     public static float testAdd2F(float a, float b) {
@@ -66,12 +86,35 @@
 
     @Test
     public void testSub() {
-        compile("testSub2F");
-        compile("testSub2D");
-        compile("testSubFConst");
-        compile("testSubConstF");
-        compile("testSubDConst");
-        compile("testSubConstD");
+        CompilationResult r = compile("testSub2F");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testSub2F FAILED");
+        }
+
+        r = compile("testSub2D");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testSub2D FAILED");
+        }
+
+        r = compile("testSubFConst");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testSubFConst FAILED");
+        }
+
+        r = compile("testSubConstF");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testSubConstF FAILED");
+        }
+
+        r = compile("testSubDConst");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testSubDconst FAILED");
+        }
+
+        r = compile("testSubConstD");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testConstD FAILED");
+        }
     }
 
     public static float testSub2F(float a, float b) {
@@ -100,12 +143,35 @@
 
     @Test
     public void testMul() {
-        compile("testMul2F");
-        compile("testMul2D");
-        compile("testMulFConst");
-        compile("testMulConstF");
-        compile("testMulDConst");
-        compile("testMulConstD");
+        CompilationResult r = compile("testMul2F");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testAdd2F FAILED");
+        }
+
+        r = compile("testMul2D");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testAdd2F FAILED");
+        }
+
+        r = compile("testMulFConst");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testAdd2F FAILED");
+        }
+
+        r = compile("testMulConstF");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testAdd2F FAILED");
+        }
+
+        r = compile("testMulDConst");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testAdd2F FAILED");
+        }
+
+        r = compile("testMulConstD");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testAdd2F FAILED");
+        }
     }
 
     public static float testMul2F(float a, float b) {
@@ -134,12 +200,35 @@
 
     @Test
     public void testDiv() {
-        compile("testDiv2F");
-        compile("testDiv2D");
-        compile("testDivFConst");
-        compile("testDivConstF");
-        compile("testDivDConst");
-        compile("testDivConstD");
+        CompilationResult r = compile("testDiv2F");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testDiv2F FAILED");
+        }
+
+        r = compile("testDiv2D");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testDiv2D FAILED");
+        }
+
+        r = compile("testDivFConst");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testDivFConst FAILED");
+        }
+
+        r = compile("testDivConstF");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testDivConstF FAILED");
+        }
+
+        r = compile("testDivDConst");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testDivDConst FAILED");
+        }
+
+        r = compile("testDivConstD");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testDivConstD FAILED");
+        }
     }
 
     public static float testDiv2F(float a, float b) {
@@ -168,8 +257,15 @@
 
     @Test
     public void testNeg() {
-        compile("testNeg2F");
-        compile("testNeg2D");
+        CompilationResult r = compile("testNeg2F");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testNeg2F FAILED");
+        }
+
+        r = compile("testNeg2D");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testNeg2D FAILED");
+        }
     }
 
     public static float testNeg2F(float a) {
@@ -195,14 +291,38 @@
         return a % b;
     }
 
+    @Ignore
     @Test
     public void testFloatConversion() {
-        compile("testF2I");
-        compile("testF2L");
-        compile("testF2D");
-        compile("testD2I");
-        compile("testD2L");
-        compile("testD2F");
+        CompilationResult r = compile("testF2I");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of tesF2I FAILED");
+        }
+
+        r = compile("testF2L");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testF2L FAILED");
+        }
+
+        r = compile("testF2D");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testF2D FAILED");
+        }
+
+        r = compile("testD2I");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testD2I FAILED");
+        }
+
+        r = compile("testD2L");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testD2L FAILED");
+        }
+
+        r = compile("testD2F");
+        if (r.getTargetCode() == null) {
+            printReport("Compilation of testD2F FAILED");
+        }
     }
 
     public static int testF2I(float a) {
@@ -229,6 +349,13 @@
         return (float) a;
     }
 
+    public static void printReport(String message) {
+        // CheckStyle: stop system..print check
+        System.out.println(message);
+        // CheckStyle: resume system..print check
+
+    }
+
     public static void main(String[] args) {
         FloatPTXTest test = new FloatPTXTest();
         for (Method m : FloatPTXTest.class.getMethods()) {
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/IntegerPTXTest.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/IntegerPTXTest.java	Sat Aug 31 09:18:58 2013 +0200
@@ -31,11 +31,44 @@
 
     @Test
     public void testAdd() {
-        invoke(compile("testAdd2I"), 8, 4);
-        invoke(compile("testAdd2L"), 12, 6);
-        invoke(compile("testAdd2B"), 6, 4);
-        invoke(compile("testAddIConst"), 5);
-        invoke(compile("testAddConstI"), 7);
+
+        Long r2 = (Long) invoke(compile("testAdd2L"), (long) 12, (long) 6);
+        if (r2 == null) {
+            printReport("testAdd2L FAILED");
+        } else if (r2.longValue() == 18) {
+            printReport("testAdd2L PASSED");
+        } else {
+            printReport("testAdd2L FAILED");
+        }
+
+        //invoke(compile("testAdd2B"), (byte) 6, (byte) 4);
+
+        Integer r4 = (Integer) invoke(compile("testAddIConst"), 5);
+        if (r4 == null) {
+            printReport("testAddIConst FAILED");
+        } else if (r4.intValue() == 37) {
+            printReport("testAddIConst PASSED");
+        } else {
+            printReport("testAddIConst FAILED");
+        }
+
+        r4 = (Integer) invoke(compile("testAddConstI"), 7);
+        if (r4 == null) {
+            printReport("testAddConstI FAILED");
+        } else if (r4.intValue() == 39) {
+            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) {
@@ -60,10 +93,42 @@
 
     @Test
     public void testSub() {
-        invoke(compile("testSub2I"), 8, 4);
-        invoke(compile("testSub2L"), 12, 6);
-        invoke(compile("testSubIConst"), 35);
-        invoke(compile("testSubConstI"), 12);
+        Long r2 = (Long) invoke(compile("testSub2L"), (long) 12, (long) 6);
+        if (r2 == null) {
+            printReport("testSub2I FAILED (null return value)");
+        } else if (r2.longValue() == 6) {
+            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) {
+            printReport("testSub2I PASSED");
+        } else {
+            printReport("testSub2I FAILED");
+        }
+
+        r1 = (Integer) invoke(compile("testSubIConst"), 35);
+        if (r1 == null) {
+            printReport("testSubIConst FAILED");
+        } else if (r1.intValue() == 3) {
+            printReport("testSubIConst PASSED");
+        } else {
+            printReport("testSubIConst FAILED");
+        }
+
+        r1 = (Integer) invoke(compile("testSubConstI"), 12);
+        if (r1 == null) {
+            printReport("testSubConstI FAILED");
+        } else if (r1.intValue() == 20) {
+            printReport("testSubConstI PASSED");
+        } else {
+            printReport("testSubConstI FAILED");
+        }
     }
 
     public static int testSub2I(int a, int b) {
@@ -85,7 +150,7 @@
     @Test
     public void testMul() {
         invoke(compile("testMul2I"), 8, 4);
-        invoke(compile("testMul2L"), 12, 6);
+        invoke(compile("testMul2L"), (long) 12, (long) 6);
         invoke(compile("testMulIConst"), 4);
         invoke(compile("testMulConstI"), 5);
     }
@@ -105,11 +170,10 @@
     public static int testMulConstI(int a) {
         return 32 * a;
     }
-
     @Test
     public void testDiv() {
         invoke(compile("testDiv2I"), 8, 4);
-        invoke(compile("testDiv2L"), 12, 6);
+        invoke(compile("testDiv2L"), (long) 12, (long) 6);
         invoke(compile("testDivIConst"), 64);
         invoke(compile("testDivConstI"), 8);
     }
@@ -133,7 +197,7 @@
     @Test
     public void testRem() {
         invoke(compile("testRem2I"), 8, 4);
-        invoke(compile("testRem2L"), 12, 6);
+        invoke(compile("testRem2L"), (long) 12, (long) 6);
     }
 
     public static int testRem2I(int a, int b) {
@@ -147,11 +211,11 @@
     @Test
     public void testIntConversion() {
         invoke(compile("testI2L"), 8);
-        invoke(compile("testL2I"), 12L);
-        invoke(compile("testI2C"), 65);
-        invoke(compile("testI2B"), 9);
-        invoke(compile("testI2F"), 17);
-        invoke(compile("testI2D"), 22);
+        invoke(compile("testL2I"), (long) 12);
+        // invoke(compile("testI2C"), 65);
+        // invoke(compile("testI2B"), 9);
+        // invoke(compile("testI2F"), 17);
+        // invoke(compile("testI2D"), 22);
     }
 
     public static long testI2L(int a) {
@@ -178,6 +242,13 @@
         return (int) a;
     }
 
+    public static void printReport(String message) {
+        // CheckStyle: stop system..print check
+        System.out.println(message);
+        // CheckStyle: resume system..print check
+
+    }
+
     public static void main(String[] args) {
         IntegerPTXTest test = new IntegerPTXTest();
         for (Method m : IntegerPTXTest.class.getMethods()) {
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/PTXTestBase.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/PTXTestBase.java	Sat Aug 31 09:18:58 2013 +0200
@@ -82,11 +82,11 @@
         return sg;
     }
 
-    protected void invoke(CompilationResult result, Object... args) {
+    protected Object invoke(CompilationResult result, Object... args) {
         try {
             if (((ExternalCompilationResult) result).getEntryPoint() == 0) {
                 Debug.dump(result, "[CUDA] *** Null entry point - Not launching kernel");
-                return;
+                return null;
             }
 
             /* Check if the method compiled is static */
@@ -95,9 +95,11 @@
             Object[] executeArgs = argsWithReceiver((isStatic ? null : this), args);
             HotSpotRuntime hsr = (HotSpotRuntime) runtime;
             InstalledCode installedCode = hsr.addExternalMethod(sg.method(), result, sg);
-            installedCode.executeVarargs(executeArgs);
+            Object r = installedCode.executeVarargs(executeArgs);
+            return r;
         } catch (Throwable th) {
             th.printStackTrace();
+            return null;
         }
     }
 }
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java	Sat Aug 31 09:18:58 2013 +0200
@@ -22,6 +22,10 @@
  */
 package com.oracle.graal.compiler.ptx;
 
+import static com.oracle.graal.api.code.ValueUtil.*;
+
+import java.util.*;
+
 import com.oracle.graal.api.code.*;
 import com.oracle.graal.api.meta.*;
 import com.oracle.graal.asm.*;
@@ -32,6 +36,11 @@
 import com.oracle.graal.lir.asm.*;
 import com.oracle.graal.lir.ptx.*;
 import com.oracle.graal.nodes.*;
+import com.oracle.graal.nodes.cfg.Block;
+import com.oracle.graal.lir.LIRInstruction.OperandFlag;
+import com.oracle.graal.lir.LIRInstruction.OperandMode;
+import com.oracle.graal.lir.LIRInstruction.ValueProcedure;
+import com.oracle.graal.graph.GraalInternalError;
 
 /**
  * PTX specific backend.
@@ -84,67 +93,121 @@
         return tasm;
     }
 
-    @Override
-    public void emitCode(TargetMethodAssembler tasm, LIRGenerator lirGen, ResolvedJavaMethod codeCacheOwner) {
-        // Emit the prologue
+    private static void emitKernelEntry(TargetMethodAssembler tasm, LIRGenerator lirGen,
+                                        ResolvedJavaMethod codeCacheOwner) {
+        // Emit PTX kernel entry text based on PTXParameterOp
+        // instructions in the start block.  Remove the instructions
+        // once kernel entry text and directives are emitted to
+        // facilitate seemless PTX code generation subsequently.
         assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method";
         final String name = codeCacheOwner.getName();
         Buffer codeBuffer = tasm.asm.codeBuffer;
+
+        // Emit initial boiler-plate directives.
         codeBuffer.emitString(".version 1.4");
         codeBuffer.emitString(".target sm_10");
         codeBuffer.emitString0(".entry " + name + " (");
         codeBuffer.emitString("");
 
-        Signature signature = codeCacheOwner.getSignature();
-        int paramCount = signature.getParameterCount(false);
-        // TODO - Revisit this.
-        // Bit-size of registers to be declared and used by the kernel.
-        int regSize = 32;
-        for (int i = 0; i < paramCount; i++) {
-            String param;
-            // No unsigned types in Java. So using .s specifier
-            switch (signature.getParameterKind(i)) {
-                case Boolean:
-                case Byte:
-                    param = ".param .s8 param" + i;
-                    regSize = 8;
-                    break;
-                case Char:
-                case Short:
-                    param = ".param .s16 param" + i;
-                    regSize = 16;
-                    break;
-                case Int:
-                    param = ".param .s32 param" + i;
-                    regSize = 32;
-                    break;
-                case Long:
-                case Float:
-                case Double:
-                case Void:
-                    param = ".param .s64 param" + i;
-                    regSize = 32;
-                    break;
-                default:
-                    // Not sure but specify 64-bit specifier??
-                    param = ".param .s64 param" + i;
-                    break;
+        // Get the start block
+        Block startBlock = lirGen.lir.cfg.getStartBlock();
+        // Keep a list of ParameterOp instructions to delete from the
+        // list of instructions in the block.
+        ArrayList<LIRInstruction> deleteOps = new ArrayList<>();
+
+        // Emit .param arguments to kernel entry based on ParameterOp
+        // instruction.
+        for (LIRInstruction op : lirGen.lir.lir(startBlock)) {
+            if (op instanceof PTXParameterOp) {
+                op.emitCode(tasm);
+                deleteOps.add(op);
             }
-            if (i != (paramCount - 1)) {
-                param += ",";
-            }
-            codeBuffer.emitString(param);
         }
 
+        // Delete ParameterOp instructions.
+        for (LIRInstruction op : deleteOps) {
+            lirGen.lir.lir(startBlock).remove(op);
+        }
+
+        // Start emiting body of the PTX kernel.
         codeBuffer.emitString0(") {");
         codeBuffer.emitString("");
 
-        // XXX For now declare one predicate and all registers
-        codeBuffer.emitString("  .reg .pred %p,%q;");
-        codeBuffer.emitString("  .reg .s" + regSize + " %r<16>;");
+        codeBuffer.emitString(".reg .u64" + " %rax;");
+    }
+
+    // Emit .reg space declarations
+    private static void emitRegisterDecl(TargetMethodAssembler tasm, LIRGenerator lirGen,
+                                         ResolvedJavaMethod codeCacheOwner) {
+        assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method";
+        Buffer codeBuffer = tasm.asm.codeBuffer;
+
+        final SortedSet<Integer> signed32 = new TreeSet<>();
+        final SortedSet<Integer> signed64 = new TreeSet<>();
+
+        ValueProcedure trackRegisterKind = new ValueProcedure() {
+
+            @Override
+            public Value doValue(Value value, OperandMode mode, EnumSet<OperandFlag> flags) {
+                if (isRegister(value)) {
+                    RegisterValue regVal = (RegisterValue) value;
+                    Kind regKind = regVal.getKind();
+                    switch (regKind) {
+                       case Int:
+                           signed32.add(regVal.getRegister().encoding());
+                           break;
+                       case Long:
+                           signed64.add(regVal.getRegister().encoding());
+                           break;
+                       default :
+                           throw GraalInternalError.shouldNotReachHere("unhandled register type "  + value.toString());
+                    }
+                }
+                return value;
+            }
+        };
 
+        for (Block b : lirGen.lir.codeEmittingOrder()) {
+            for (LIRInstruction op : lirGen.lir.lir(b)) {
+                op.forEachOutput(trackRegisterKind);
+            }
+        }
+
+        for (Integer i : signed32) {
+            codeBuffer.emitString("  .reg .s32 %r" + i.intValue() + ";");
+        }
+        for (Integer i : signed64) {
+            codeBuffer.emitString(".reg .s64 %r" + i.intValue() + ";");
+        }
+    }
+
+    @Override
+    public void emitCode(TargetMethodAssembler tasm, LIRGenerator lirGen, ResolvedJavaMethod codeCacheOwner) {
+        assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method";
+        Buffer codeBuffer = tasm.asm.codeBuffer;
+        // Emit the prologue
+        emitKernelEntry(tasm, lirGen, codeCacheOwner);
+
+        // Emit register declarations
+        try {
+            emitRegisterDecl(tasm, lirGen, codeCacheOwner);
+        } catch (GraalInternalError e) {
+            // TODO : Better error handling needs to be done once
+            //        all types of parameters are handled.
+            codeBuffer.setPosition(0);
+            codeBuffer.close(false);
+            return;
+        }
         // Emit code for the LIR
-        lirGen.lir.emitCode(tasm);
+        try {
+            lirGen.lir.emitCode(tasm);
+        } catch (GraalInternalError e) {
+            // TODO : Better error handling needs to be done once
+            //        all types of parameters are handled.
+            codeBuffer.setPosition(0);
+            codeBuffer.close(false);
+            return;
+        }
 
         // Emit the epilogue
         codeBuffer.emitString0("}");
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java	Sat Aug 31 09:18:58 2013 +0200
@@ -47,12 +47,16 @@
 import com.oracle.graal.lir.ptx.PTXControlFlow.CondMoveOp;
 import com.oracle.graal.lir.ptx.PTXControlFlow.FloatCondMoveOp;
 import com.oracle.graal.lir.ptx.PTXControlFlow.ReturnOp;
+import com.oracle.graal.lir.ptx.PTXControlFlow.ReturnNoValOp;
 import com.oracle.graal.lir.ptx.PTXControlFlow.SequentialSwitchOp;
 import com.oracle.graal.lir.ptx.PTXControlFlow.TableSwitchOp;
-import com.oracle.graal.lir.ptx.PTXMove.LoadOp;
 import com.oracle.graal.lir.ptx.PTXMove.MoveFromRegOp;
 import com.oracle.graal.lir.ptx.PTXMove.MoveToRegOp;
-import com.oracle.graal.lir.ptx.PTXMove.StoreOp;
+import com.oracle.graal.lir.ptx.PTXMemOp.LoadOp;
+import com.oracle.graal.lir.ptx.PTXMemOp.StoreOp;
+import com.oracle.graal.lir.ptx.PTXMemOp.LoadParamOp;
+import com.oracle.graal.lir.ptx.PTXMemOp.LoadReturnAddrOp;
+import com.oracle.graal.lir.ptx.PTXMemOp.StoreReturnValOp;
 import com.oracle.graal.nodes.*;
 import com.oracle.graal.nodes.calc.*;
 import com.oracle.graal.nodes.java.*;
@@ -96,6 +100,42 @@
         }
     }
 
+    protected static AllocatableValue toParamKind(AllocatableValue value) {
+        if (value.getKind().getStackKind() != value.getKind()) {
+            // We only have stack-kinds in the LIR, so convert the operand kind for values from the
+            // calling convention.
+            if (isRegister(value)) {
+                return asRegister(value).asValue(value.getKind().getStackKind());
+            } else if (isStackSlot(value)) {
+                return StackSlot.get(value.getKind().getStackKind(), asStackSlot(value).getRawOffset(), asStackSlot(value).getRawAddFrameSize());
+            } else {
+                throw GraalInternalError.shouldNotReachHere();
+            }
+        }
+        return value;
+    }
+
+    @Override
+    public void emitPrologue() {
+        // Need to emit .param directives based on incoming arguments and return value
+        CallingConvention incomingArguments = cc;
+        int argCount = incomingArguments.getArgumentCount();
+        // Additional argument for return value.
+        Value[] params = new Value[argCount + 1];
+        for (int i = 0; i < argCount; i++) {
+            params[i] = toParamKind(incomingArguments.getArgument(i));
+        }
+        // Add the return value as the last parameter.
+        params[argCount] =  incomingArguments.getReturn();
+
+        append(new PTXParameterOp(params));
+        for (LocalNode local : graph.getNodes(LocalNode.class)) {
+            Value param = params[local.index()];
+            assert param.getKind() == local.kind().getStackKind();
+            setResult(local, emitLoadParam(param.getKind(), param, null));
+        }
+    }
+
     @Override
     public Variable emitMove(Value input) {
         Variable result = newVariable(input.getKind());
@@ -250,7 +290,8 @@
     /**
      * This method emits the compare instruction, and may reorder the operands. It returns true if
      * it did so.
-     * 
+     *
+     *
      * @param a the left operand of the comparison
      * @param b the right operand of the comparison
      * @return true if the left and right operands were switched, false otherwise
@@ -705,6 +746,10 @@
         append(new ReturnOp(input));
     }
 
+    private void emitReturnNoVal() {
+        append(new ReturnNoValOp());
+    }
+
     @Override
     protected void emitSequentialSwitch(Constant[] keyConstants, LabelRef[] keyTargets, LabelRef defaultTarget, Value key) {
         // Making a copy of the switch value is necessary because jump table destroys the input
@@ -761,4 +806,38 @@
     public void visitInfopointNode(InfopointNode i) {
         throw new InternalError("NYI");
     }
+
+    public Variable emitLoadParam(Kind kind, Value address, DeoptimizingNode deopting) {
+        PTXAddressValue loadAddress = asAddress(address);
+        Variable result = newVariable(kind);
+        append(new LoadParamOp(kind, result, loadAddress, deopting != null ? state(deopting) : null));
+        return result;
+    }
+
+    public Variable emitLoadReturnAddress(Kind kind, Value address, DeoptimizingNode deopting) {
+        PTXAddressValue loadAddress = asAddress(address);
+        Variable result = newVariable(kind);
+        append(new LoadReturnAddrOp(kind, result, loadAddress, deopting != null ? state(deopting) : null));
+        return result;
+    }
+
+    public void emitStoreReturnValue(Kind kind, Value address, Value inputVal, DeoptimizingNode deopting) {
+        PTXAddressValue storeAddress = asAddress(address);
+        Variable input = load(inputVal);
+        append(new StoreReturnValOp(kind, storeAddress, input, deopting != null ? state(deopting) : null));
+    }
+
+
+    @Override
+    public void visitReturn(ReturnNode x) {
+        AllocatableValue operand = Value.ILLEGAL;
+        if (x.result() != null) {
+            operand = resultOperandFor(x.result().kind());
+            // Load the global memory address from return parameter
+            Variable loadVar = emitLoadReturnAddress(operand.getKind(), operand, null);
+            // Store result in global memory whose location is loadVar
+            emitStoreReturnValue(operand.getKind(), loadVar, operand(x.result()), null);
+        }
+        emitReturnNoVal();
+    }
 }
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXTargetMethodAssembler.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXTargetMethodAssembler.java	Sat Aug 31 09:18:58 2013 +0200
@@ -50,7 +50,7 @@
         ExternalCompilationResult graalCompile = (ExternalCompilationResult) super.finishTargetMethod(graph);
 
         try {
-            if (validDevice) {
+            if ((validDevice) && (graalCompile.getTargetCode() != null)) {
                 long kernel = toGPU.generateKernel(graalCompile.getTargetCode(), method.getName());
                 graalCompile.setEntryPoint(kernel);
             }
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotBackend.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotBackend.java	Sat Aug 31 09:18:58 2013 +0200
@@ -244,10 +244,10 @@
             AMD64Address src = new AMD64Address(receiver, config.hubOffset);
 
             AMD64HotSpotLIRGenerator gen = (AMD64HotSpotLIRGenerator) lirGen;
-            HotSpotRuntime hr = ((HotSpotRuntime) gen.getRuntime());
-            if (hr.config.useCompressedKlassPointers) {
+            AMD64HotSpotRuntime hr = ((AMD64HotSpotRuntime) gen.getRuntime());
+            if (hr.useCompressedKlassPointers()) {
                 Register register = r10;
-                AMD64Move.decodeKlassPointer(asm, register, hr.heapBaseRegister(), src, hr.config.narrowKlassBase, hr.config.narrowKlassShift, hr.config.logKlassAlignment);
+                AMD64HotSpotMove.decodeKlassPointer(asm, register, hr.heapBaseRegister(), src, config.narrowKlassBase, config.narrowKlassShift, config.logKlassAlignment);
                 asm.cmpq(inlineCacheKlass, register);
             } else {
                 asm.cmpq(inlineCacheKlass, src);
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotLIRGenerator.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotLIRGenerator.java	Sat Aug 31 09:18:58 2013 +0200
@@ -38,6 +38,10 @@
 import com.oracle.graal.compiler.gen.*;
 import com.oracle.graal.graph.*;
 import com.oracle.graal.hotspot.*;
+import com.oracle.graal.hotspot.amd64.AMD64HotSpotMove.CompareAndSwapCompressedOp;
+import com.oracle.graal.hotspot.amd64.AMD64HotSpotMove.LoadCompressedPointer;
+import com.oracle.graal.hotspot.amd64.AMD64HotSpotMove.StoreCompressedConstantOp;
+import com.oracle.graal.hotspot.amd64.AMD64HotSpotMove.StoreCompressedPointer;
 import com.oracle.graal.hotspot.meta.*;
 import com.oracle.graal.hotspot.nodes.*;
 import com.oracle.graal.hotspot.stubs.*;
@@ -46,12 +50,9 @@
 import com.oracle.graal.lir.StandardOp.PlaceholderOp;
 import com.oracle.graal.lir.amd64.*;
 import com.oracle.graal.lir.amd64.AMD64ControlFlow.CondMoveOp;
-import com.oracle.graal.lir.amd64.AMD64Move.CompareAndSwapCompressedOp;
 import com.oracle.graal.lir.amd64.AMD64Move.CompareAndSwapOp;
-import com.oracle.graal.lir.amd64.AMD64Move.LoadCompressedPointer;
 import com.oracle.graal.lir.amd64.AMD64Move.LoadOp;
 import com.oracle.graal.lir.amd64.AMD64Move.MoveFromRegOp;
-import com.oracle.graal.lir.amd64.AMD64Move.StoreCompressedPointer;
 import com.oracle.graal.lir.amd64.AMD64Move.StoreConstantOp;
 import com.oracle.graal.lir.amd64.AMD64Move.StoreOp;
 import com.oracle.graal.nodes.*;
@@ -427,12 +428,12 @@
          * algorithms may differ.
          */
         if (isCompressCandidate(access)) {
-            if (runtime().config.useCompressedOops && kind == Kind.Object) {
-                append(new LoadCompressedPointer(kind, result, runtime().heapBaseRegister().asValue(), loadAddress, access != null ? state(access) : null, runtime().config.narrowOopBase,
-                                runtime().config.narrowOopShift, runtime().config.logMinObjAlignment));
-            } else if (runtime().config.useCompressedKlassPointers && kind == Kind.Long) {
-                append(new LoadCompressedPointer(kind, result, runtime().heapBaseRegister().asValue(), loadAddress, access != null ? state(access) : null, runtime().config.narrowKlassBase,
-                                runtime().config.narrowKlassShift, runtime().config.logKlassAlignment));
+            if (runtime().useCompressedOops() && kind == Kind.Object) {
+                append(new LoadCompressedPointer(kind, result, runtime().heapBaseRegister().asValue(), loadAddress, access != null ? state(access) : null, getNarrowOopBase(), getNarrowOopShift(),
+                                getLogMinObjectAlignment()));
+            } else if (runtime().useCompressedKlassPointers() && kind == Kind.Long) {
+                append(new LoadCompressedPointer(kind, result, runtime().heapBaseRegister().asValue(), loadAddress, access != null ? state(access) : null, getNarrowKlassBase(), getNarrowKlassShift(),
+                                getLogKlassAlignment()));
             } else {
                 append(new LoadOp(kind, result, loadAddress, access != null ? state(access) : null));
             }
@@ -449,29 +450,29 @@
         if (isConstant(inputVal)) {
             Constant c = asConstant(inputVal);
             if (canStoreConstant(c)) {
-                if (inputVal.getKind() == Kind.Object) {
-                    append(new StoreConstantOp(kind, storeAddress, c, state, runtime().config.useCompressedOops && isCompressCandidate(access)));
-                } else if (inputVal.getKind() == Kind.Long) {
-                    append(new StoreConstantOp(kind, storeAddress, c, state, runtime().config.useCompressedKlassPointers && isCompressCandidate(access)));
+                if (inputVal.getKind() == Kind.Object && runtime().useCompressedOops() && isCompressCandidate(access)) {
+                    append(new StoreCompressedConstantOp(kind, storeAddress, c, state));
+                } else if (inputVal.getKind() == Kind.Long && runtime().useCompressedKlassPointers() && isCompressCandidate(access)) {
+                    append(new StoreCompressedConstantOp(kind, storeAddress, c, state));
                 } else {
-                    append(new StoreConstantOp(kind, storeAddress, c, state, false));
+                    append(new StoreConstantOp(kind, storeAddress, c, state));
                 }
                 return;
             }
         }
         Variable input = load(inputVal);
         if (isCompressCandidate(access)) {
-            if (runtime().config.useCompressedOops && kind == Kind.Object) {
+            if (runtime().useCompressedOops() && kind == Kind.Object) {
                 if (input.getKind() == Kind.Object) {
                     Variable scratch = newVariable(Kind.Long);
-                    append(new StoreCompressedPointer(kind, storeAddress, input, scratch, state, runtime().config.narrowOopBase, runtime().config.narrowOopShift, runtime().config.logMinObjAlignment));
+                    append(new StoreCompressedPointer(kind, storeAddress, input, scratch, state, getNarrowOopBase(), getNarrowOopShift(), getLogMinObjectAlignment()));
                 } else {
                     // the input oop is already compressed
                     append(new StoreOp(input.getKind(), storeAddress, input, state));
                 }
-            } else if (runtime().config.useCompressedKlassPointers && kind == Kind.Long) {
+            } else if (runtime().useCompressedKlassPointers() && kind == Kind.Long) {
                 Variable scratch = newVariable(Kind.Long);
-                append(new StoreCompressedPointer(kind, storeAddress, input, scratch, state, runtime().config.narrowKlassBase, runtime().config.narrowKlassShift, runtime().config.logKlassAlignment));
+                append(new StoreCompressedPointer(kind, storeAddress, input, scratch, state, getNarrowKlassBase(), getNarrowKlassShift(), getLogKlassAlignment()));
             } else {
                 append(new StoreOp(kind, storeAddress, input, state));
             }
@@ -480,6 +481,30 @@
         }
     }
 
+    private int getLogMinObjectAlignment() {
+        return runtime().config.logMinObjAlignment;
+    }
+
+    private int getNarrowOopShift() {
+        return runtime().config.narrowOopShift;
+    }
+
+    private long getNarrowOopBase() {
+        return runtime().config.narrowOopBase;
+    }
+
+    private int getLogKlassAlignment() {
+        return runtime().config.logKlassAlignment;
+    }
+
+    private int getNarrowKlassShift() {
+        return runtime().config.narrowKlassShift;
+    }
+
+    private long getNarrowKlassBase() {
+        return runtime().config.narrowKlassBase;
+    }
+
     @Override
     public void visitCompareAndSwap(LoweredCompareAndSwapNode node, Value address) {
         Kind kind = node.getNewValue().kind();
@@ -489,9 +514,9 @@
         AMD64AddressValue addressValue = asAddressValue(address);
         RegisterValue raxRes = AMD64.rax.asValue(kind);
         emitMove(raxRes, expected);
-        if (runtime().config.useCompressedOops && node.isCompressible()) {
+        if (runtime().useCompressedOops() && node.isCompressible()) {
             Variable scratch = newVariable(Kind.Long);
-            append(new CompareAndSwapCompressedOp(raxRes, addressValue, raxRes, newValue, scratch, runtime().config.narrowOopBase, runtime().config.narrowOopShift, runtime().config.logMinObjAlignment));
+            append(new CompareAndSwapCompressedOp(raxRes, addressValue, raxRes, newValue, scratch, getNarrowOopBase(), getNarrowOopShift(), getLogMinObjectAlignment()));
         } else {
             append(new CompareAndSwapOp(raxRes, addressValue, raxRes, newValue));
         }
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotMove.java	Sat Aug 31 09:18:58 2013 +0200
@@ -0,0 +1,250 @@
+/*
+ * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved.
+ * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
+ *
+ * This code is free software; you can redistribute it and/or modify it
+ * under the terms of the GNU General Public License version 2 only, as
+ * published by the Free Software Foundation.
+ *
+ * This code is distributed in the hope that it will be useful, but WITHOUT
+ * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+ * FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+ * version 2 for more details (a copy is included in the LICENSE file that
+ * accompanied this code).
+ *
+ * You should have received a copy of the GNU General Public License version
+ * 2 along with this work; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
+ *
+ * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
+ * or visit www.oracle.com if you need additional information or have any
+ * questions.
+ */
+package com.oracle.graal.hotspot.amd64;
+
+import static com.oracle.graal.api.code.ValueUtil.*;
+import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*;
+
+import com.oracle.graal.amd64.*;
+import com.oracle.graal.api.code.*;
+import com.oracle.graal.api.meta.*;
+import com.oracle.graal.asm.*;
+import com.oracle.graal.asm.amd64.*;
+import com.oracle.graal.asm.amd64.AMD64Assembler.ConditionFlag;
+import com.oracle.graal.graph.*;
+import com.oracle.graal.hotspot.meta.*;
+import com.oracle.graal.lir.*;
+import com.oracle.graal.lir.amd64.*;
+import com.oracle.graal.lir.amd64.AMD64Move.LoadOp;
+import com.oracle.graal.lir.amd64.AMD64Move.StoreConstantOp;
+import com.oracle.graal.lir.asm.*;
+
+public class AMD64HotSpotMove {
+
+    public static class StoreCompressedConstantOp extends StoreConstantOp {
+
+        public StoreCompressedConstantOp(Kind kind, AMD64AddressValue address, Constant input, LIRFrameState state) {
+            super(kind, address, input, state);
+        }
+
+        @Override
+        public void emitMemAccess(AMD64MacroAssembler masm) {
+            if (kind == Kind.Long) {
+                if (NumUtil.isInt(input.asLong())) {
+                    masm.movl(address.toAddress(), (int) input.asLong());
+                } else {
+                    throw GraalInternalError.shouldNotReachHere("Cannot store 64-bit constants to memory");
+                }
+            } else if (kind == Kind.Object) {
+                if (input.isNull()) {
+                    masm.movl(address.toAddress(), 0);
+                } else {
+                    throw GraalInternalError.shouldNotReachHere("Cannot store 64-bit constants to memory");
+                }
+            } else {
+                throw GraalInternalError.shouldNotReachHere("Attempt to store compressed constant of wrong type.");
+            }
+        }
+    }
+
+    public static class LoadCompressedPointer extends LoadOp {
+
+        private long base;
+        private int shift;
+        private int alignment;
+        @Alive({REG}) protected AllocatableValue heapBaseRegister;
+
+        public LoadCompressedPointer(Kind kind, AllocatableValue result, AllocatableValue heapBaseRegister, AMD64AddressValue address, LIRFrameState state, long base, int shift, int alignment) {
+            super(kind, result, address, state);
+            this.base = base;
+            this.shift = shift;
+            this.alignment = alignment;
+            this.heapBaseRegister = heapBaseRegister;
+            assert kind == Kind.Object || kind == Kind.Long;
+        }
+
+        @Override
+        public void emitMemAccess(AMD64MacroAssembler masm) {
+            Register resRegister = asRegister(result);
+            masm.movl(resRegister, address.toAddress());
+            if (kind == Kind.Object) {
+                decodePointer(masm, resRegister, asRegister(heapBaseRegister), base, shift, alignment);
+            } else {
+                decodeKlassPointer(masm, resRegister, asRegister(heapBaseRegister), base, shift, alignment);
+            }
+        }
+    }
+
+    public static class StoreCompressedPointer extends AMD64LIRInstruction {
+
+        protected final Kind kind;
+        private long base;
+        private int shift;
+        private int alignment;
+        @Temp({REG}) private AllocatableValue scratch;
+        @Alive({REG}) protected AllocatableValue input;
+        @Alive({COMPOSITE}) protected AMD64AddressValue address;
+        @State protected LIRFrameState state;
+
+        public StoreCompressedPointer(Kind kind, AMD64AddressValue address, AllocatableValue input, AllocatableValue scratch, LIRFrameState state, long base, int shift, int alignment) {
+            this.base = base;
+            this.shift = shift;
+            this.alignment = alignment;
+            this.scratch = scratch;
+            this.kind = kind;
+            this.address = address;
+            this.state = state;
+            this.input = input;
+            assert kind == Kind.Object || kind == Kind.Long;
+        }
+
+        @Override
+        public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) {
+            Register heapBase = ((HotSpotRuntime) tasm.runtime).heapBaseRegister();
+            masm.movq(asRegister(scratch), asRegister(input));
+            if (kind == Kind.Object) {
+                encodePointer(masm, asRegister(scratch), heapBase, base, shift, alignment);
+            } else {
+                encodeKlassPointer(masm, asRegister(scratch), heapBase, base, shift, alignment);
+            }
+            if (state != null) {
+                tasm.recordImplicitException(masm.codeBuffer.position(), state);
+            }
+            masm.movl(address.toAddress(), asRegister(scratch));
+        }
+    }
+
+    @Opcode("CAS")
+    public static class CompareAndSwapCompressedOp extends AMD64LIRInstruction {
+
+        @Def protected AllocatableValue result;
+        @Alive({COMPOSITE}) protected AMD64AddressValue address;
+        @Alive protected AllocatableValue cmpValue;
+        @Alive protected AllocatableValue newValue;
+        @Temp({REG}) protected AllocatableValue scratch;
+
+        private long base;
+        private int shift;
+        private int alignment;
+
+        public CompareAndSwapCompressedOp(AllocatableValue result, AMD64AddressValue address, AllocatableValue cmpValue, AllocatableValue newValue, AllocatableValue scratch, long base, int shift,
+                        int alignment) {
+            this.base = base;
+            this.shift = shift;
+            this.alignment = alignment;
+            this.scratch = scratch;
+            this.result = result;
+            this.address = address;
+            this.cmpValue = cmpValue;
+            this.newValue = newValue;
+            assert cmpValue.getKind() == Kind.Object;
+        }
+
+        @Override
+        public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) {
+            compareAndSwapCompressed(tasm, masm, result, address, cmpValue, newValue, scratch, base, shift, alignment);
+        }
+    }
+
+    protected static void compareAndSwapCompressed(TargetMethodAssembler tasm, AMD64MacroAssembler masm, AllocatableValue result, AMD64AddressValue address, AllocatableValue cmpValue,
+                    AllocatableValue newValue, AllocatableValue scratch, long base, int shift, int alignment) {
+        assert AMD64.rax.equals(asRegister(cmpValue)) && AMD64.rax.equals(asRegister(result));
+        final Register scratchRegister = asRegister(scratch);
+        final Register cmpRegister = asRegister(cmpValue);
+        final Register newRegister = asRegister(newValue);
+        Register heapBase = ((HotSpotRuntime) tasm.runtime).heapBaseRegister();
+        encodePointer(masm, cmpRegister, heapBase, base, shift, alignment);
+        masm.movq(scratchRegister, newRegister);
+        encodePointer(masm, scratchRegister, heapBase, base, shift, alignment);
+        if (tasm.target.isMP) {
+            masm.lock();
+        }
+        masm.cmpxchgl(scratchRegister, address.toAddress());
+    }
+
+    private static void encodePointer(AMD64MacroAssembler masm, Register scratchRegister, Register heapBaseRegister, long base, int shift, int alignment) {
+        // If the base is zero, the uncompressed address has to be shifted right
+        // in order to be compressed.
+        if (base == 0) {
+            if (shift != 0) {
+                assert alignment == shift : "Encode algorithm is wrong";
+                masm.shrq(scratchRegister, alignment);
+            }
+        } else {
+            // Otherwise the heap base, which resides always in register 12, is subtracted
+            // followed by right shift.
+            masm.testq(scratchRegister, scratchRegister);
+            // If the stored reference is null, move the heap to scratch
+            // register and then calculate the compressed oop value.
+            masm.cmovq(ConditionFlag.Equal, scratchRegister, heapBaseRegister);
+            masm.subq(scratchRegister, heapBaseRegister);
+            masm.shrq(scratchRegister, alignment);
+        }
+    }
+
+    private static void decodePointer(AMD64MacroAssembler masm, Register resRegister, Register heapBaseRegister, long base, int shift, int alignment) {
+        // If the base is zero, the compressed address has to be shifted left
+        // in order to be uncompressed.
+        if (base == 0) {
+            if (shift != 0) {
+                assert alignment == shift : "Decode algorithm is wrong";
+                masm.shlq(resRegister, alignment);
+            }
+        } else {
+            Label done = new Label();
+            masm.shlq(resRegister, alignment);
+            masm.jccb(ConditionFlag.Equal, done);
+            // Otherwise the heap base is added to the shifted address.
+            masm.addq(resRegister, heapBaseRegister);
+            masm.bind(done);
+        }
+    }
+
+    private static void encodeKlassPointer(AMD64MacroAssembler masm, Register scratchRegister, Register heapBaseRegister, long base, int shift, int alignment) {
+        if (base != 0) {
+            masm.subq(scratchRegister, heapBaseRegister);
+        }
+        if (shift != 0) {
+            assert alignment == shift : "Encode algorithm is wrong";
+            masm.shrq(scratchRegister, alignment);
+        }
+    }
+
+    private static void decodeKlassPointer(AMD64MacroAssembler masm, Register resRegister, Register heapBaseRegister, long base, int shift, int alignment) {
+        if (shift != 0) {
+            assert alignment == shift : "Decode algorithm is wrong";
+            masm.shlq(resRegister, alignment);
+            if (base != 0) {
+                masm.addq(resRegister, heapBaseRegister);
+            }
+        } else {
+            assert base == 0 : "Sanity";
+        }
+    }
+
+    public static void decodeKlassPointer(AMD64MacroAssembler masm, Register register, Register heapBaseRegister, AMD64Address address, long narrowKlassBase, int narrowKlassShift,
+                    int logKlassAlignment) {
+        masm.movl(register, address);
+        decodeKlassPointer(masm, register, heapBaseRegister, narrowKlassBase, narrowKlassShift, logKlassAlignment);
+    }
+}
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotRuntime.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotRuntime.java	Sat Aug 31 09:18:58 2013 +0200
@@ -345,6 +345,12 @@
     public abstract Register threadRegister();
 
     /**
+     * Returns the register used by the runtime for maintaining the heap base address for compressed
+     * pointers.
+     */
+    public abstract Register heapBaseRegister();
+
+    /**
      * Gets the stack pointer register.
      */
     public abstract Register stackPointerRegister();
@@ -478,6 +484,14 @@
         return Array.getLength(array.asObject());
     }
 
+    public boolean useCompressedOops() {
+        return config.useCompressedOops;
+    }
+
+    public boolean useCompressedKlassPointers() {
+        return config.useCompressedKlassPointers;
+    }
+
     @Override
     public void lower(Node n, LoweringTool tool) {
         StructuredGraph graph = (StructuredGraph) n.graph();
@@ -863,13 +877,13 @@
     private FloatingReadNode createReadHub(StructuredGraph graph, Kind wordKind, ValueNode object, GuardingNode guard) {
         LocationNode location = ConstantLocationNode.create(FINAL_LOCATION, wordKind, config.hubOffset, graph);
         assert !object.isConstant() || object.asConstant().isNull();
-        return graph.add(new FloatingReadNode(object, location, null, StampFactory.forKind(wordKind()), guard, BarrierType.NONE, config.useCompressedKlassPointers));
+        return graph.add(new FloatingReadNode(object, location, null, StampFactory.forKind(wordKind()), guard, BarrierType.NONE, useCompressedKlassPointers()));
     }
 
     private WriteNode createWriteHub(StructuredGraph graph, Kind wordKind, ValueNode object, ValueNode value) {
         LocationNode location = ConstantLocationNode.create(ANY_LOCATION, wordKind, config.hubOffset, graph);
         assert !object.isConstant() || object.asConstant().isNull();
-        return graph.add(new WriteNode(object, value, location, BarrierType.NONE, config.useCompressedKlassPointers));
+        return graph.add(new WriteNode(object, value, location, BarrierType.NONE, useCompressedKlassPointers()));
     }
 
     private static BarrierType getFieldLoadBarrierType(HotSpotResolvedJavaField loadField) {
@@ -927,7 +941,7 @@
     }
 
     public int getScalingFactor(Kind kind) {
-        if (config.useCompressedOops && kind == Kind.Object) {
+        if (useCompressedOops() && kind == Kind.Object) {
             return this.graalRuntime.getTarget().arch.getSizeInBytes(Kind.Int);
         } else {
             return this.graalRuntime.getTarget().arch.getSizeInBytes(kind);
@@ -1134,7 +1148,7 @@
             case Int:
                 return Constant.forInt(base == null ? unsafe.getInt(displacement) : unsafe.getInt(base, displacement));
             case Long:
-                if (displacement == config().hubOffset && this.getGraalRuntime().getRuntime().config.useCompressedKlassPointers) {
+                if (displacement == config().hubOffset && useCompressedKlassPointers()) {
                     if (base == null) {
                         throw new GraalInternalError("Base of object must not be null");
                     } else {
--- a/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Move.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Move.java	Sat Aug 31 09:18:58 2013 +0200
@@ -32,7 +32,6 @@
 import com.oracle.graal.api.meta.*;
 import com.oracle.graal.asm.*;
 import com.oracle.graal.asm.amd64.*;
-import com.oracle.graal.asm.amd64.AMD64Assembler.ConditionFlag;
 import com.oracle.graal.graph.*;
 import com.oracle.graal.lir.*;
 import com.oracle.graal.lir.StandardOp.MoveOp;
@@ -117,34 +116,6 @@
         }
     }
 
-    public static class LoadCompressedPointer extends LoadOp {
-
-        private long base;
-        private int shift;
-        private int alignment;
-        @Alive({REG}) protected AllocatableValue heapBaseRegister;
-
-        public LoadCompressedPointer(Kind kind, AllocatableValue result, AllocatableValue heapBaseRegister, AMD64AddressValue address, LIRFrameState state, long base, int shift, int alignment) {
-            super(kind, result, address, state);
-            this.base = base;
-            this.shift = shift;
-            this.alignment = alignment;
-            this.heapBaseRegister = heapBaseRegister;
-            assert kind == Kind.Object || kind == Kind.Long;
-        }
-
-        @Override
-        public void emitMemAccess(AMD64MacroAssembler masm) {
-            Register resRegister = asRegister(result);
-            masm.movl(resRegister, address.toAddress());
-            if (kind == Kind.Object) {
-                decodePointer(masm, resRegister, asRegister(heapBaseRegister), base, shift, alignment);
-            } else {
-                decodeKlassPointer(masm, resRegister, asRegister(heapBaseRegister), base, shift, alignment);
-            }
-        }
-    }
-
     public static class LoadOp extends MemOp {
 
         @Def({REG}) protected AllocatableValue result;
@@ -188,44 +159,6 @@
         }
     }
 
-    public static class StoreCompressedPointer extends AMD64LIRInstruction {
-
-        protected final Kind kind;
-        private long base;
-        private int shift;
-        private int alignment;
-        @Temp({REG}) private AllocatableValue scratch;
-        @Alive({REG}) protected AllocatableValue input;
-        @Alive({COMPOSITE}) protected AMD64AddressValue address;
-        @State protected LIRFrameState state;
-
-        public StoreCompressedPointer(Kind kind, AMD64AddressValue address, AllocatableValue input, AllocatableValue scratch, LIRFrameState state, long base, int shift, int alignment) {
-            this.base = base;
-            this.shift = shift;
-            this.alignment = alignment;
-            this.scratch = scratch;
-            this.kind = kind;
-            this.address = address;
-            this.state = state;
-            this.input = input;
-            assert kind == Kind.Object || kind == Kind.Long;
-        }
-
-        @Override
-        public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) {
-            masm.movq(asRegister(scratch), asRegister(input));
-            if (kind == Kind.Object) {
-                encodePointer(masm, asRegister(scratch), tasm.runtime.heapBaseRegister(), base, shift, alignment);
-            } else {
-                encodeKlassPointer(masm, asRegister(scratch), tasm.runtime.heapBaseRegister(), base, shift, alignment);
-            }
-            if (state != null) {
-                tasm.recordImplicitException(masm.codeBuffer.position(), state);
-            }
-            masm.movl(address.toAddress(), asRegister(scratch));
-        }
-    }
-
     public static class StoreOp extends MemOp {
 
         @Use({REG}) protected AllocatableValue input;
@@ -271,12 +204,10 @@
     public static class StoreConstantOp extends MemOp {
 
         protected final Constant input;
-        private final boolean compressible;
 
-        public StoreConstantOp(Kind kind, AMD64AddressValue address, Constant input, LIRFrameState state, boolean compressible) {
+        public StoreConstantOp(Kind kind, AMD64AddressValue address, Constant input, LIRFrameState state) {
             super(kind, address, state);
             this.input = input;
-            this.compressible = compressible;
         }
 
         @Override
@@ -295,11 +226,7 @@
                     break;
                 case Long:
                     if (NumUtil.isInt(input.asLong())) {
-                        if (compressible) {
-                            masm.movl(address.toAddress(), (int) input.asLong());
-                        } else {
-                            masm.movslq(address.toAddress(), (int) input.asLong());
-                        }
+                        masm.movslq(address.toAddress(), (int) input.asLong());
                     } else {
                         throw GraalInternalError.shouldNotReachHere("Cannot store 64-bit constants to memory");
                     }
@@ -311,11 +238,7 @@
                     throw GraalInternalError.shouldNotReachHere("Cannot store 64-bit constants to memory");
                 case Object:
                     if (input.isNull()) {
-                        if (compressible) {
-                            masm.movl(address.toAddress(), 0);
-                        } else {
-                            masm.movptr(address.toAddress(), 0);
-                        }
+                        masm.movptr(address.toAddress(), 0);
                     } else {
                         throw GraalInternalError.shouldNotReachHere("Cannot store 64-bit constants to memory");
                     }
@@ -410,38 +333,6 @@
         }
     }
 
-    @Opcode("CAS")
-    public static class CompareAndSwapCompressedOp extends AMD64LIRInstruction {
-
-        @Def protected AllocatableValue result;
-        @Alive({COMPOSITE}) protected AMD64AddressValue address;
-        @Alive protected AllocatableValue cmpValue;
-        @Alive protected AllocatableValue newValue;
-        @Temp({REG}) protected AllocatableValue scratch;
-
-        private long base;
-        private int shift;
-        private int alignment;
-
-        public CompareAndSwapCompressedOp(AllocatableValue result, AMD64AddressValue address, AllocatableValue cmpValue, AllocatableValue newValue, AllocatableValue scratch, long base, int shift,
-                        int alignment) {
-            this.base = base;
-            this.shift = shift;
-            this.alignment = alignment;
-            this.scratch = scratch;
-            this.result = result;
-            this.address = address;
-            this.cmpValue = cmpValue;
-            this.newValue = newValue;
-            assert cmpValue.getKind() == Kind.Object;
-        }
-
-        @Override
-        public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) {
-            compareAndSwapCompressed(tasm, masm, result, address, cmpValue, newValue, scratch, base, shift, alignment);
-        }
-    }
-
     public static void move(TargetMethodAssembler tasm, AMD64MacroAssembler masm, Value result, Value input) {
         if (isRegister(input)) {
             if (isRegister(result)) {
@@ -650,85 +541,4 @@
                 throw GraalInternalError.shouldNotReachHere();
         }
     }
-
-    protected static void compareAndSwapCompressed(TargetMethodAssembler tasm, AMD64MacroAssembler masm, AllocatableValue result, AMD64AddressValue address, AllocatableValue cmpValue,
-                    AllocatableValue newValue, AllocatableValue scratch, long base, int shift, int alignment) {
-        assert AMD64.rax.equals(asRegister(cmpValue)) && AMD64.rax.equals(asRegister(result));
-        final Register scratchRegister = asRegister(scratch);
-        final Register cmpRegister = asRegister(cmpValue);
-        final Register newRegister = asRegister(newValue);
-        encodePointer(masm, cmpRegister, tasm.runtime.heapBaseRegister(), base, shift, alignment);
-        masm.movq(scratchRegister, newRegister);
-        encodePointer(masm, scratchRegister, tasm.runtime.heapBaseRegister(), base, shift, alignment);
-        if (tasm.target.isMP) {
-            masm.lock();
-        }
-        masm.cmpxchgl(scratchRegister, address.toAddress());
-    }
-
-    private static void encodePointer(AMD64MacroAssembler masm, Register scratchRegister, Register heapBaseRegister, long base, int shift, int alignment) {
-        // If the base is zero, the uncompressed address has to be shifted right
-        // in order to be compressed.
-        if (base == 0) {
-            if (shift != 0) {
-                assert alignment == shift : "Encode algorithm is wrong";
-                masm.shrq(scratchRegister, alignment);
-            }
-        } else {
-            // Otherwise the heap base, which resides always in register 12, is subtracted
-            // followed by right shift.
-            masm.testq(scratchRegister, scratchRegister);
-            // If the stored reference is null, move the heap to scratch
-            // register and then calculate the compressed oop value.
-            masm.cmovq(ConditionFlag.Equal, scratchRegister, heapBaseRegister);
-            masm.subq(scratchRegister, heapBaseRegister);
-            masm.shrq(scratchRegister, alignment);
-        }
-    }
-
-    private static void decodePointer(AMD64MacroAssembler masm, Register resRegister, Register heapBaseRegister, long base, int shift, int alignment) {
-        // If the base is zero, the compressed address has to be shifted left
-        // in order to be uncompressed.
-        if (base == 0) {
-            if (shift != 0) {
-                assert alignment == shift : "Decode algorithm is wrong";
-                masm.shlq(resRegister, alignment);
-            }
-        } else {
-            Label done = new Label();
-            masm.shlq(resRegister, alignment);
-            masm.jccb(ConditionFlag.Equal, done);
-            // Otherwise the heap base is added to the shifted address.
-            masm.addq(resRegister, heapBaseRegister);
-            masm.bind(done);
-        }
-    }
-
-    private static void encodeKlassPointer(AMD64MacroAssembler masm, Register scratchRegister, Register heapBaseRegister, long base, int shift, int alignment) {
-        if (base != 0) {
-            masm.subq(scratchRegister, heapBaseRegister);
-        }
-        if (shift != 0) {
-            assert alignment == shift : "Encode algorithm is wrong";
-            masm.shrq(scratchRegister, alignment);
-        }
-    }
-
-    private static void decodeKlassPointer(AMD64MacroAssembler masm, Register resRegister, Register heapBaseRegister, long base, int shift, int alignment) {
-        if (shift != 0) {
-            assert alignment == shift : "Decode algorithm is wrong";
-            masm.shlq(resRegister, alignment);
-            if (base != 0) {
-                masm.addq(resRegister, heapBaseRegister);
-            }
-        } else {
-            assert base == 0 : "Sanity";
-        }
-    }
-
-    public static void decodeKlassPointer(AMD64MacroAssembler masm, Register register, Register heapBaseRegister, AMD64Address address, long narrowKlassBase, int narrowKlassShift,
-                    int logKlassAlignment) {
-        masm.movl(register, address);
-        decodeKlassPointer(masm, register, heapBaseRegister, narrowKlassBase, narrowKlassShift, logKlassAlignment);
-    }
 }
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java	Sat Aug 31 09:18:58 2013 +0200
@@ -55,6 +55,19 @@
         }
     }
 
+    public static class ReturnNoValOp extends PTXLIRInstruction {
+
+        public ReturnNoValOp() { }
+
+        @Override
+        public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) {
+            if (tasm.frameContext != null) {
+                tasm.frameContext.leave(tasm);
+            }
+            masm.ret();
+        }
+    }
+
     public static class BranchOp extends PTXLIRInstruction implements StandardOp.BranchOp {
 
         protected Condition condition;
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMemOp.java	Sat Aug 31 09:18:58 2013 +0200
@@ -0,0 +1,264 @@
+/*
+ * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved.
+ * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
+ *
+ * This code is free software; you can redistribute it and/or modify it
+ * under the terms of the GNU General Public License version 2 only, as
+ * published by the Free Software Foundation.
+ *
+ * This code is distributed in the hope that it will be useful, but WITHOUT
+ * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+ * FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+ * version 2 for more details (a copy is included in the LICENSE file that
+ * accompanied this code).
+ *
+ * You should have received a copy of the GNU General Public License version
+ * 2 along with this work; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
+ *
+ * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
+ * or visit www.oracle.com if you need additional information or have any
+ * questions.
+ */
+package com.oracle.graal.lir.ptx;
+
+import static com.oracle.graal.api.code.ValueUtil.*;
+import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*;
+
+import com.oracle.graal.api.meta.*;
+import com.oracle.graal.asm.ptx.*;
+import com.oracle.graal.graph.*;
+import com.oracle.graal.lir.*;
+import com.oracle.graal.lir.asm.*;
+
+public class PTXMemOp {
+
+    // Load operation from .global state space
+    @Opcode("LOAD")
+    public static class LoadOp extends PTXLIRInstruction {
+
+        private final Kind kind;
+        @Def({REG}) protected AllocatableValue result;
+        @Use({COMPOSITE}) protected PTXAddressValue address;
+        @State protected LIRFrameState state;
+
+        public LoadOp(Kind kind, AllocatableValue result, PTXAddressValue address, LIRFrameState state) {
+            this.kind = kind;
+            this.result = result;
+            this.address = address;
+            this.state = state;
+        }
+
+        @Override
+        public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) {
+            PTXAddress addr = address.toAddress();
+            switch (kind) {
+                case Byte:
+                    masm.ld_global_s8(asRegister(result), addr.getBase(), addr.getDisplacement());
+                    break;
+                case Short:
+                    masm.ld_global_s16(asRegister(result), addr.getBase(), addr.getDisplacement());
+                    break;
+                case Char:
+                    masm.ld_global_u16(asRegister(result), addr.getBase(), addr.getDisplacement());
+                    break;
+                case Int:
+                    masm.ld_global_s32(asRegister(result), addr.getBase(), addr.getDisplacement());
+                    break;
+                case Long:
+                    masm.ld_global_s64(asRegister(result), addr.getBase(), addr.getDisplacement());
+                    break;
+                case Float:
+                    masm.ld_global_f32(asRegister(result), addr.getBase(), addr.getDisplacement());
+                    break;
+                case Double:
+                    masm.ld_global_f64(asRegister(result), addr.getBase(), addr.getDisplacement());
+                    break;
+                case Object:
+                    masm.ld_global_u32(asRegister(result), addr.getBase(), addr.getDisplacement());
+                    break;
+                default:
+                    throw GraalInternalError.shouldNotReachHere();
+            }
+        }
+    }
+
+    // Store operation from .global state space
+    @Opcode("STORE")
+    public static class StoreOp extends PTXLIRInstruction {
+
+        private final Kind kind;
+        @Use({COMPOSITE}) protected PTXAddressValue address;
+        @Use({REG}) protected AllocatableValue input;
+        @State protected LIRFrameState state;
+
+        public StoreOp(Kind kind, PTXAddressValue address, AllocatableValue input, LIRFrameState state) {
+            this.kind = kind;
+            this.address = address;
+            this.input = input;
+            this.state = state;
+        }
+
+        @Override
+        public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) {
+            assert isRegister(input);
+            PTXAddress addr = address.toAddress();
+            switch (kind) {
+                case Byte:
+                    masm.st_global_s8(addr.getBase(), addr.getDisplacement(), asRegister(input));
+                    break;
+                case Short:
+                    masm.st_global_s8(addr.getBase(), addr.getDisplacement(), asRegister(input));
+                    break;
+                case Int:
+                    masm.st_global_s32(addr.getBase(), addr.getDisplacement(), asRegister(input));
+                    break;
+                case Long:
+                    masm.st_global_s64(addr.getBase(), addr.getDisplacement(), asRegister(input));
+                    break;
+                case Float:
+                    masm.st_global_f32(addr.getBase(), addr.getDisplacement(), asRegister(input));
+                    break;
+                case Double:
+                    masm.st_global_f64(addr.getBase(), addr.getDisplacement(), asRegister(input));
+                    break;
+                case Object:
+                    masm.st_global_u64(addr.getBase(), addr.getDisplacement(), asRegister(input));
+                    break;
+                default:
+                    throw GraalInternalError.shouldNotReachHere("missing: " + address.getKind());
+            }
+        }
+    }
+
+    // Load operation from .param state space
+    @Opcode("LOAD")
+    public static class LoadParamOp extends PTXLIRInstruction {
+
+        private final Kind kind;
+        @Def({REG}) protected AllocatableValue result;
+        @Use({COMPOSITE}) protected PTXAddressValue address;
+        @State protected LIRFrameState state;
+
+        public LoadParamOp(Kind kind, AllocatableValue result, PTXAddressValue address, LIRFrameState state) {
+            this.kind = kind;
+            this.result = result;
+            this.address = address;
+            this.state = state;
+        }
+
+        @Override
+        public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) {
+            PTXAddress addr = address.toAddress();
+            switch (kind) {
+                case Byte:
+                    masm.ld_from_state_space(".param.s8", asRegister(result), addr.getBase(), addr.getDisplacement());
+                    break;
+                case Short:
+                    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());
+                    break;
+                case Int:
+                    masm.ld_from_state_space(".param.s32", asRegister(result), addr.getBase(), addr.getDisplacement());
+                    break;
+                case Long:
+                    masm.ld_from_state_space(".param.s64", asRegister(result), addr.getBase(), addr.getDisplacement());
+                    break;
+                case Float:
+                    masm.ld_from_state_space(".param.f32", asRegister(result), addr.getBase(), addr.getDisplacement());
+                    break;
+                case Double:
+                    masm.ld_from_state_space(".param.f64", asRegister(result), addr.getBase(), addr.getDisplacement());
+                    break;
+                case Object:
+                    masm.ld_from_state_space(".param.u64", asRegister(result), addr.getBase(), addr.getDisplacement());
+                    break;
+                default:
+                    throw GraalInternalError.shouldNotReachHere();
+            }
+        }
+    }
+
+    // Load contents of return value pointer from return argument in
+    // .param state space
+    @Opcode("LOAD_RET_ADDR")
+    public static class LoadReturnAddrOp extends PTXLIRInstruction {
+
+        private final Kind kind;
+        @Def({REG}) protected AllocatableValue result;
+        @Use({COMPOSITE}) protected PTXAddressValue address;
+        @State protected LIRFrameState state;
+
+        public LoadReturnAddrOp(Kind kind, AllocatableValue result, PTXAddressValue address, LIRFrameState state) {
+            this.kind = kind;
+            this.result = result;
+            this.address = address;
+            this.state = state;
+        }
+
+        @Override
+        public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) {
+            PTXAddress addr = address.toAddress();
+            switch (kind) {
+                case Int:
+                    masm.ld_return_address("u32", asRegister(result), addr.getBase(), addr.getDisplacement());
+                    break;
+                case Long:
+                    masm.ld_return_address("u64", asRegister(result), addr.getBase(), addr.getDisplacement());
+                    break;
+                default:
+                    throw GraalInternalError.shouldNotReachHere();
+            }
+        }
+    }
+
+    // Store operation from .global state space
+    @Opcode("STORE_RETURN_VALUE")
+    public static class StoreReturnValOp extends PTXLIRInstruction {
+
+        private final Kind kind;
+        @Use({COMPOSITE}) protected PTXAddressValue address;
+        @Use({REG}) protected AllocatableValue input;
+        @State protected LIRFrameState state;
+
+        public StoreReturnValOp(Kind kind, PTXAddressValue address, AllocatableValue input, LIRFrameState state) {
+            this.kind = kind;
+            this.address = address;
+            this.input = input;
+            this.state = state;
+        }
+
+        @Override
+        public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) {
+            assert isRegister(input);
+            PTXAddress addr = address.toAddress();
+            // masm.st_global_return_value_s64(addr.getBase(), addr.getDisplacement(), asRegister(input));
+
+            switch (kind) {
+                case Byte:
+                case Short:
+                    masm.st_global_return_value_s8(addr.getBase(), addr.getDisplacement(), asRegister(input));
+                    break;
+                case Int:
+                    masm.st_global_return_value_s32(addr.getBase(), addr.getDisplacement(), asRegister(input));
+                    break;
+                case Long:
+                    masm.st_global_return_value_s64(addr.getBase(), addr.getDisplacement(), asRegister(input));
+                    break;
+                case Float:
+                    masm.st_global_return_value_f32(addr.getBase(), addr.getDisplacement(), asRegister(input));
+                    break;
+                case Double:
+                    masm.st_global_return_value_f64(addr.getBase(), addr.getDisplacement(), asRegister(input));
+                    break;
+                case Object:
+                    masm.st_global_return_value_u64(addr.getBase(), addr.getDisplacement(), asRegister(input));
+                    break;
+                default:
+                    throw GraalInternalError.shouldNotReachHere("missing: " + address.getKind());
+            }
+        }
+    }
+}
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java	Sat Aug 31 09:18:58 2013 +0200
@@ -116,100 +116,6 @@
         }
     }
 
-    public static class LoadOp extends PTXLIRInstruction {
-
-        private final Kind kind;
-        @Def({REG}) protected AllocatableValue result;
-        @Use({COMPOSITE}) protected PTXAddressValue address;
-        @State protected LIRFrameState state;
-
-        public LoadOp(Kind kind, AllocatableValue result, PTXAddressValue address, LIRFrameState state) {
-            this.kind = kind;
-            this.result = result;
-            this.address = address;
-            this.state = state;
-        }
-
-        @Override
-        public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) {
-            PTXAddress addr = address.toAddress();
-            switch (kind) {
-                case Byte:
-                    masm.ld_global_s8(asRegister(result), addr.getBase(), addr.getDisplacement());
-                    break;
-                case Short:
-                    masm.ld_global_s16(asRegister(result), addr.getBase(), addr.getDisplacement());
-                    break;
-                case Char:
-                    masm.ld_global_u16(asRegister(result), addr.getBase(), addr.getDisplacement());
-                    break;
-                case Int:
-                    masm.ld_global_s32(asRegister(result), addr.getBase(), addr.getDisplacement());
-                    break;
-                case Long:
-                    masm.ld_global_s64(asRegister(result), addr.getBase(), addr.getDisplacement());
-                    break;
-                case Float:
-                    masm.ld_global_f32(asRegister(result), addr.getBase(), addr.getDisplacement());
-                    break;
-                case Double:
-                    masm.ld_global_f64(asRegister(result), addr.getBase(), addr.getDisplacement());
-                    break;
-                case Object:
-                    masm.ld_global_u32(asRegister(result), addr.getBase(), addr.getDisplacement());
-                    break;
-                default:
-                    throw GraalInternalError.shouldNotReachHere();
-            }
-        }
-    }
-
-    public static class StoreOp extends PTXLIRInstruction {
-
-        private final Kind kind;
-        @Use({COMPOSITE}) protected PTXAddressValue address;
-        @Use({REG}) protected AllocatableValue input;
-        @State protected LIRFrameState state;
-
-        public StoreOp(Kind kind, PTXAddressValue address, AllocatableValue input, LIRFrameState state) {
-            this.kind = kind;
-            this.address = address;
-            this.input = input;
-            this.state = state;
-        }
-
-        @Override
-        public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) {
-            assert isRegister(input);
-            PTXAddress addr = address.toAddress();
-            switch (kind) {
-                case Byte:
-                    masm.st_global_s8(addr.getBase(), addr.getDisplacement(), asRegister(input));
-                    break;
-                case Short:
-                    masm.st_global_s8(addr.getBase(), addr.getDisplacement(), asRegister(input));
-                    break;
-                case Int:
-                    masm.st_global_s32(addr.getBase(), addr.getDisplacement(), asRegister(input));
-                    break;
-                case Long:
-                    masm.st_global_s64(addr.getBase(), addr.getDisplacement(), asRegister(input));
-                    break;
-                case Float:
-                    masm.st_global_f32(addr.getBase(), addr.getDisplacement(), asRegister(input));
-                    break;
-                case Double:
-                    masm.st_global_f64(addr.getBase(), addr.getDisplacement(), asRegister(input));
-                    break;
-                case Object:
-                    masm.st_global_s32(addr.getBase(), addr.getDisplacement(), asRegister(input));
-                    break;
-                default:
-                    throw GraalInternalError.shouldNotReachHere("missing: " + address.getKind());
-            }
-        }
-    }
-
     public static class LeaOp extends PTXLIRInstruction {
 
         @Def({REG}) protected AllocatableValue result;
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXParameterOp.java	Sat Aug 31 09:18:58 2013 +0200
@@ -0,0 +1,68 @@
+/*
+ * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved.
+ * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
+ *
+ * This code is free software; you can redistribute it and/or modify it
+ * under the terms of the GNU General Public License version 2 only, as
+ * published by the Free Software Foundation.
+ *
+ * This code is distributed in the hope that it will be useful, but WITHOUT
+ * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+ * FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+ * version 2 for more details (a copy is included in the LICENSE file that
+ * accompanied this code).
+ *
+ * You should have received a copy of the GNU General Public License version
+ * 2 along with this work; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
+ *
+ * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
+ * or visit www.oracle.com if you need additional information or have any
+ * questions.
+ */
+
+package com.oracle.graal.lir.ptx;
+
+import static com.oracle.graal.api.code.ValueUtil.*;
+import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*;
+
+import com.oracle.graal.api.meta.*;
+import com.oracle.graal.asm.ptx.*;
+import com.oracle.graal.graph.*;
+import com.oracle.graal.lir.*;
+import com.oracle.graal.lir.asm.*;
+
+public class PTXParameterOp extends LIRInstruction {
+
+    @Def({REG}) protected Value[] params;
+
+    public PTXParameterOp(Value[] params) {
+        this.params = params;
+    }
+
+    @Override
+    public void emitCode(TargetMethodAssembler tasm) {
+        PTXAssembler ptxasm = (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)));
+                break;
+            case Long :
+                ptxasm.param_64_decl(asLongReg(params[i]), (i == (argCount - 1)));
+                break;
+            case Float :
+                ptxasm.param_32_decl(asFloatReg(params[i]), (i == (argCount - 1)));
+                break;
+            case Double :
+                ptxasm.param_64_decl(asDoubleReg(params[i]), (i == (argCount - 1)));
+                break;
+            default :
+                throw GraalInternalError.shouldNotReachHere("unhandled parameter type "  + paramKind.toString());
+            }
+        }
+    }
+}
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/ConstantNode.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/ConstantNode.java	Sat Aug 31 09:18:58 2013 +0200
@@ -77,6 +77,14 @@
     }
 
     /**
+     * Returns a node for a primitive constant.
+     */
+    public static ConstantNode forPrimitive(Constant constant, Graph graph) {
+        assert constant.getKind() != Kind.Object;
+        return forConstant(constant, null, graph);
+    }
+
+    /**
      * Returns a node for a double constant.
      * 
      * @param d the double value for which to create the instruction
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/PhiNode.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/PhiNode.java	Sat Aug 31 09:18:58 2013 +0200
@@ -34,7 +34,7 @@
  * variable.
  */
 @NodeInfo(nameTemplate = "{p#type/s}Phi({i#values})")
-public final class PhiNode extends FloatingNode implements Canonicalizable, Node.IterableNodeType, GuardingNode {
+public class PhiNode extends FloatingNode implements Canonicalizable, Node.IterableNodeType, GuardingNode {
 
     public static enum PhiType {
         Value(null), // normal value phis
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/BinaryNode.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/BinaryNode.java	Sat Aug 31 09:18:58 2013 +0200
@@ -83,6 +83,57 @@
         }
     }
 
+    public static BinaryNode add(ValueNode x, ValueNode y) {
+        assert x.kind() == y.kind();
+        switch (x.kind()) {
+            case Byte:
+            case Char:
+            case Short:
+            case Int:
+            case Long:
+                return IntegerArithmeticNode.add(x, y);
+            case Float:
+            case Double:
+                return x.graph().unique(new FloatAddNode(x.kind(), x, y, false));
+            default:
+                throw GraalInternalError.shouldNotReachHere();
+        }
+    }
+
+    public static BinaryNode sub(ValueNode x, ValueNode y) {
+        assert x.kind() == y.kind();
+        switch (x.kind()) {
+            case Byte:
+            case Char:
+            case Short:
+            case Int:
+            case Long:
+                return IntegerArithmeticNode.sub(x, y);
+            case Float:
+            case Double:
+                return x.graph().unique(new FloatSubNode(x.kind(), x, y, false));
+            default:
+                throw GraalInternalError.shouldNotReachHere();
+        }
+    }
+
+    public static BinaryNode mul(ValueNode x, ValueNode y) {
+        assert x.kind() == y.kind();
+        switch (x.kind()) {
+            case Byte:
+            case Char:
+            case Short:
+            case Int:
+            case Long:
+                return IntegerArithmeticNode.mul(x, y);
+            case Float:
+            case Double:
+                return x.graph().unique(new FloatMulNode(x.kind(), x, y, false));
+            default:
+                throw GraalInternalError.shouldNotReachHere();
+        }
+    }
+
     public static boolean canTryReassociate(BinaryNode node) {
         return node instanceof IntegerAddNode || node instanceof IntegerSubNode || node instanceof IntegerMulNode || node instanceof AndNode || node instanceof OrNode || node instanceof XorNode;
     }
--- /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/CompilerErrorTest.java	Sat Aug 31 09:18:58 2013 +0200
@@ -0,0 +1,57 @@
+/*
+ * Copyright (c) 2012, 2012, 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 com.oracle.truffle.api.dsl.*;
+import com.oracle.truffle.api.dsl.test.TypeSystemTest.ValueNode;
+
+public class CompilerErrorTest {
+
+    abstract static class Visiblity01 extends ValueNode {
+
+        @Specialization
+        @SuppressWarnings("static-method")
+        @ExpectError("Method annotated with @Specialization must not be private.")
+        private Object s() {
+            return null;
+        }
+
+    }
+
+    @ExpectError("Classes containing a @Specialization annotation must not be private.")
+    private abstract static class Visiblity02 extends ValueNode {
+
+        @Specialization
+        public Object s() {
+            return null;
+        }
+
+    }
+
+    // assert no error
+    @ExpectError({})
+    private abstract static class Visiblity03 extends ValueNode {
+
+    }
+
+}
--- a/graal/com.oracle.truffle.api.dsl.test/src/com/oracle/truffle/api/dsl/test/SpecializationGroupingTest.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.api.dsl.test/src/com/oracle/truffle/api/dsl/test/SpecializationGroupingTest.java	Sat Aug 31 09:18:58 2013 +0200
@@ -26,10 +26,13 @@
 
 import com.oracle.truffle.api.*;
 import com.oracle.truffle.api.dsl.*;
+import com.oracle.truffle.api.dsl.test.SpecializationGroupingTestFactory.TestElseConnectionBug1Factory;
+import com.oracle.truffle.api.dsl.test.SpecializationGroupingTestFactory.TestElseConnectionBug2Factory;
 import com.oracle.truffle.api.dsl.test.SpecializationGroupingTestFactory.TestGroupingFactory;
 import com.oracle.truffle.api.dsl.test.TypeSystemTest.SimpleTypes;
 import com.oracle.truffle.api.dsl.test.TypeSystemTest.TestRootNode;
 import com.oracle.truffle.api.dsl.test.TypeSystemTest.ValueNode;
+import com.oracle.truffle.api.frame.*;
 import com.oracle.truffle.api.nodes.*;
 
 /**
@@ -154,6 +157,83 @@
 
     }
 
+    @Test
+    public void testElseConnectionBug1() {
+        CallTarget target = TestHelper.createCallTarget(TestElseConnectionBug1Factory.create(new GenericInt()));
+        Assert.assertEquals(42, target.call());
+    }
+
+    @SuppressWarnings("unused")
+    @NodeChild(value = "genericChild", type = GenericInt.class)
+    public abstract static class TestElseConnectionBug1 extends ValueNode {
+
+        @Specialization(order = 1, rewriteOn = {SlowPathException.class}, guards = "isInitialized")
+        public int doInteger(int value) throws SlowPathException {
+            throw new SlowPathException();
+        }
+
+        @Specialization(order = 3, guards = "isInitialized")
+        public int doObject(int value) {
+            return value == 42 ? value : 0;
+        }
+
+        @Specialization(order = 4, guards = "!isInitialized")
+        public Object doUninitialized(int value) {
+            throw new AssertionError();
+        }
+
+        boolean isInitialized(int value) {
+            return true;
+        }
+    }
+
+    public static final class GenericInt extends ValueNode {
+
+        @Override
+        public Object execute(VirtualFrame frame) {
+            return executeInt(frame);
+        }
+
+        @Override
+        public int executeInt(VirtualFrame frame) {
+            return 42;
+        }
+
+    }
+
+    @Test
+    public void testElseConnectionBug2() {
+        TestHelper.assertRuns(TestElseConnectionBug2Factory.getInstance(), 42, 42);
+    }
+
+    @SuppressWarnings("unused")
+    @NodeChild
+    public abstract static class TestElseConnectionBug2 extends ValueNode {
+
+        @Specialization(order = 2, guards = "guard0")
+        public int doGuard0(int value) {
+            throw new AssertionError();
+        }
+
+        @Specialization(order = 3, guards = "guard1")
+        public int doGuard1(int value) {
+            throw new AssertionError();
+        }
+
+        @Specialization(order = 4, guards = "!guard0")
+        public int doUninitialized(int value) {
+            return value;
+        }
+
+        boolean guard0(int value) {
+            return false;
+        }
+
+        boolean guard1(int value) {
+            return false;
+        }
+    }
+
     private static class MockAssumption implements Assumption {
 
         int checked;
--- a/graal/com.oracle.truffle.api.dsl.test/src/com/oracle/truffle/api/dsl/test/TestHelper.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.api.dsl.test/src/com/oracle/truffle/api/dsl/test/TestHelper.java	Sat Aug 31 09:18:58 2013 +0200
@@ -22,6 +22,8 @@
  */
 package com.oracle.truffle.api.dsl.test;
 
+import static org.junit.Assert.*;
+
 import java.util.*;
 
 import com.oracle.truffle.api.*;
@@ -75,4 +77,69 @@
         return createCallTarget(node).call(new TestArguments(values));
     }
 
+    static Object array(Object... val) {
+        return val;
+    }
+
+    static <E> List<List<E>> permutations(List<E> list) {
+        return permutations(new ArrayList<E>(), list, new ArrayList<List<E>>());
+    }
+
+    static Object[][] permutations(Object... list) {
+        List<List<Object>> permutations = permutations(Arrays.asList(list));
+
+        Object[][] a = new Object[permutations.size()][];
+        int index = 0;
+        for (List<Object> p : permutations) {
+            a[index] = p.toArray(new Object[p.size()]);
+            index++;
+        }
+
+        return a;
+    }
+
+    static <E> List<List<E>> permutations(List<E> prefix, List<E> suffix, List<List<E>> output) {
+        if (suffix.size() == 1) {
+            ArrayList<E> newElement = new ArrayList<>(prefix);
+            newElement.addAll(suffix);
+            output.add(newElement);
+            return output;
+        }
+
+        for (int i = 0; i < suffix.size(); i++) {
+            List<E> newPrefix = new ArrayList<>(prefix);
+            newPrefix.add(suffix.get(i));
+            List<E> newSuffix = new ArrayList<>(suffix);
+            newSuffix.remove(i);
+            permutations(newPrefix, newSuffix, output);
+        }
+
+        return output;
+    }
+
+    /* Methods tests all test values in combinational order. */
+    static void assertRuns(NodeFactory<? extends ValueNode> factory, Object result, Object... testValues) {
+        // test each run by its own.
+        for (int i = 0; i < testValues.length; i++) {
+            assertValue(createRoot(factory), result, testValues);
+        }
+
+        // test all combinations of the test values
+        List<List<Object>> permuts = permutations(Arrays.asList(testValues));
+        for (List<Object> list : permuts) {
+            TestRootNode<?> root = createRoot(factory);
+            for (Object object : list) {
+                assertValue(root, result, object);
+            }
+        }
+    }
+
+    static void assertValue(TestRootNode<? extends ValueNode> root, Object result, Object testValues) {
+        if (testValues instanceof Object[]) {
+            assertEquals(result, executeWith(root, (Object[]) testValues));
+        } else {
+            assertEquals(result, executeWith(root, testValues));
+        }
+    }
+
 }
--- /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/TypeSystemErrorsTest.java	Sat Aug 31 09:18:58 2013 +0200
@@ -0,0 +1,70 @@
+/*
+ * 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 com.oracle.truffle.api.dsl.*;
+import com.oracle.truffle.api.dsl.test.TypeSystemTest.*;
+
+public class TypeSystemErrorsTest {
+
+    @TypeSystem({int.class, boolean.class})
+    public static class Types0 {
+
+    }
+
+    @ExpectError("Invalid type order. The type(s) [java.lang.String] are inherited from a earlier defined type java.lang.CharSequence.")
+    @TypeSystem({CharSequence.class, String.class})
+    public static class Types1 {
+
+    }
+
+    @TypeSystem({int.class, boolean.class})
+    public static class Types2 {
+
+        @TypeCast
+        @ExpectError("The provided return type \"String\" does not match expected return type \"int\".%")
+        String asInteger(Object value) {
+            return (String) value;
+        }
+
+    }
+
+    @TypeSystem({int.class, boolean.class})
+    public static class Types3 {
+
+        @TypeCast
+        @ExpectError("The provided return type \"boolean\" does not match expected return type \"int\".%")
+        boolean asInteger(Object value) {
+            return (boolean) value;
+        }
+
+    }
+
+    @TypeSystemReference(Types0.class)
+    @NodeChild
+    @ExpectError("The @TypeSystem of the node and the @TypeSystem of the @NodeChild does not match. Types0 != SimpleTypes. ")
+    abstract static class ErrorNode1 extends ValueNode {
+
+    }
+
+}
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.truffle.api.dsl/src/com/oracle/truffle/api/dsl/ExpectError.java	Sat Aug 31 09:18:58 2013 +0200
@@ -0,0 +1,38 @@
+/*
+ * Copyright (c) 2012, 2012, 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.  Oracle designates this
+ * particular file as subject to the "Classpath" exception as provided
+ * by Oracle in the LICENSE file that accompanied this code.
+ *
+ * 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;
+
+import java.lang.annotation.*;
+
+/**
+ * This annotation is internally known by the dsl processor and used to expect errors for testing
+ * purposes. This is not part of public API.
+ */
+@Retention(RetentionPolicy.RUNTIME)
+public @interface ExpectError {
+
+    String[] value();
+
+}
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.truffle.api.dsl/src/com/oracle/truffle/api/dsl/ImplicitCast.java	Sat Aug 31 09:18:58 2013 +0200
@@ -0,0 +1,37 @@
+/*
+ * Copyright (c) 2012, 2012, 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.  Oracle designates this
+ * particular file as subject to the "Classpath" exception as provided
+ * by Oracle in the LICENSE file that accompanied this code.
+ *
+ * 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;
+
+import java.lang.annotation.*;
+
+/**
+ * EXPERIMENTAL Filter feature. May change or disappear without notice. This feature is not
+ * functional yet.
+ */
+@Retention(RetentionPolicy.CLASS)
+@Target({ElementType.METHOD})
+public @interface ImplicitCast {
+
+}
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/AbstractParser.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/AbstractParser.java	Sat Aug 31 09:18:58 2013 +0200
@@ -65,7 +65,7 @@
                 return null;
             }
 
-            model.emitMessages((TypeElement) element, log);
+            model.emitMessages(context, (TypeElement) element, log);
             return filterErrorElements(model);
         } catch (CompileErrorException e) {
             log.message(Kind.WARNING, element, null, null, "The truffle processor could not parse class due to error: %s", e.getMessage());
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/TruffleProcessor.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/TruffleProcessor.java	Sat Aug 31 09:18:58 2013 +0200
@@ -96,7 +96,7 @@
 
     private static void handleThrowable(AnnotationProcessor generator, Throwable t, Element e) {
         String message = "Uncaught error in " + generator.getClass().getSimpleName() + " while processing " + e;
-        generator.getContext().getLog().message(Kind.ERROR, e, null, null, message + ": " + Utils.printException(t));
+        generator.getContext().getEnvironment().getMessager().printMessage(Kind.ERROR, message + ": " + Utils.printException(t), e);
     }
 
     @SuppressWarnings("unchecked")
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/TruffleTypes.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/TruffleTypes.java	Sat Aug 31 09:18:58 2013 +0200
@@ -30,6 +30,7 @@
 
 import com.oracle.truffle.api.*;
 import com.oracle.truffle.api.CompilerDirectives.*;
+import com.oracle.truffle.api.dsl.*;
 import com.oracle.truffle.api.frame.*;
 import com.oracle.truffle.api.nodes.*;
 import com.oracle.truffle.api.nodes.Node.Child;
@@ -56,6 +57,7 @@
     private final TypeMirror compilerAsserts;
     private final DeclaredType slowPath;
     private final DeclaredType truffleOptions;
+    private final TypeElement expectError;
 
     private final List<String> errors = new ArrayList<>();
 
@@ -74,6 +76,11 @@
         nodeInfoKind = getRequired(context, NodeInfo.Kind.class);
         slowPath = getRequired(context, SlowPath.class);
         truffleOptions = getRequired(context, TruffleOptions.class);
+        expectError = (TypeElement) getRequired(context, ExpectError.class).asElement();
+    }
+
+    public TypeElement getExpectError() {
+        return expectError;
     }
 
     public DeclaredType getNodeInfoAnnotation() {
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/Utils.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/Utils.java	Sat Aug 31 09:18:58 2013 +0200
@@ -423,6 +423,8 @@
                 return getSimpleName(mirror);
             case ERROR:
                 throw new CompileErrorException("Type error " + mirror);
+            case EXECUTABLE:
+                return ((ExecutableType) mirror).toString();
             case NONE:
                 return "$none";
             default:
@@ -755,6 +757,10 @@
 
     public static AnnotationMirror findAnnotationMirror(ProcessingEnvironment processingEnv, List<? extends AnnotationMirror> mirrors, Class<?> annotationClass) {
         TypeElement expectedAnnotationType = processingEnv.getElementUtils().getTypeElement(annotationClass.getCanonicalName());
+        return findAnnotationMirror(mirrors, expectedAnnotationType);
+    }
+
+    public static AnnotationMirror findAnnotationMirror(List<? extends AnnotationMirror> mirrors, TypeElement expectedAnnotationType) {
         for (AnnotationMirror mirror : mirrors) {
             DeclaredType annotationType = mirror.getAnnotationType();
             TypeElement actualAnnotationType = (TypeElement) annotationType.asElement();
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/CreateCastParser.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/CreateCastParser.java	Sat Aug 31 09:18:58 2013 +0200
@@ -61,7 +61,7 @@
     }
 
     @Override
-    public CreateCastData create(TemplateMethod method) {
+    public CreateCastData create(TemplateMethod method, boolean invalid) {
         AnnotationMirror mirror = method.getMarkerAnnotation();
         List<String> childNames = Utils.getAnnotationValueList(String.class, mirror, "value");
         CreateCastData cast = new CreateCastData(method, childNames);
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/ExecutableTypeMethodParser.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/ExecutableTypeMethodParser.java	Sat Aug 31 09:18:58 2013 +0200
@@ -73,17 +73,13 @@
 
     @Override
     protected List<TypeMirror> nodeTypeMirrors(NodeData nodeData) {
-        // executable types not yet available
-        if (nodeData.getTypeSystem() == null) {
-            return Collections.emptyList();
-        }
-        List<TypeMirror> types = new ArrayList<>(nodeData.getTypeSystem().getPrimitiveTypeMirrors());
-        types.add(nodeData.getTypeSystem().getVoidType().getPrimitiveType());
+        List<TypeMirror> types = new ArrayList<>(getNode().getTypeSystem().getPrimitiveTypeMirrors());
+        types.add(getNode().getTypeSystem().getVoidType().getPrimitiveType());
         return types;
     }
 
     @Override
-    public ExecutableTypeData create(TemplateMethod method) {
+    public ExecutableTypeData create(TemplateMethod method, boolean invalid) {
         TypeData resolvedType = method.getReturnType().getTypeSystemType();
         return new ExecutableTypeData(method, method.getMethod(), getNode().getTypeSystem(), resolvedType);
     }
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/GenericParser.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/GenericParser.java	Sat Aug 31 09:18:58 2013 +0200
@@ -61,7 +61,7 @@
     }
 
     @Override
-    public SpecializationData create(TemplateMethod method) {
+    public SpecializationData create(TemplateMethod method, boolean invalid) {
         SpecializationData data = new SpecializationData(method, true, false, false);
         return data;
     }
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/NodeMethodParser.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/NodeMethodParser.java	Sat Aug 31 09:18:58 2013 +0200
@@ -85,10 +85,14 @@
         return methodSpec;
     }
 
-    private void addDefaultChildren(boolean shortCircuitsEnabled, String shortCircuitName, MethodSpec methodSpec) {
+    public void addDefaultChildren(boolean shortCircuitsEnabled, String breakName, MethodSpec methodSpec) {
         // children are null when parsing executable types
         if (getNode().getChildren() != null) {
             for (NodeChildData child : getNode().getChildren()) {
+                String valueName = child.getName();
+                if (breakName != null && valueName.equals(breakName)) {
+                    break;
+                }
                 if (child.getExecutionKind() == ExecutionKind.DEFAULT) {
                     ParameterSpec spec = createValueParameterSpec(child.getName(), child.getNodeData(), child.getExecuteWith().size());
                     if (child.getCardinality().isMany()) {
@@ -97,10 +101,6 @@
                     }
                     methodSpec.addRequired(spec);
                 } else if (child.getExecutionKind() == ExecutionKind.SHORT_CIRCUIT) {
-                    String valueName = child.getName();
-                    if (shortCircuitName != null && valueName.equals(shortCircuitName)) {
-                        break;
-                    }
 
                     if (shortCircuitsEnabled) {
                         methodSpec.addRequired(new ParameterSpec(shortCircuitValueName(valueName), getContext().getType(boolean.class)));
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/NodeParser.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/NodeParser.java	Sat Aug 31 09:18:58 2013 +0200
@@ -199,6 +199,10 @@
         }
 
         for (NodeData splittedNode : nodes) {
+            if (templateType.getModifiers().contains(Modifier.PRIVATE) && splittedNode.getSpecializations().size() > 0) {
+                splittedNode.addError("Classes containing a @%s annotation must not be private.", Specialization.class.getSimpleName());
+            }
+
             finalizeSpecializations(elements, splittedNode);
             verifyNode(splittedNode, elements);
             createPolymorphicSpecializations(splittedNode);
@@ -328,11 +332,13 @@
         nodeData.setNodeContainer(nodeContainer != null);
         nodeData.setTypeSystem(typeSystem);
         nodeData.setFields(parseFields(typeHierarchy, elements));
-        parsedNodes.put(Utils.getQualifiedName(templateType), nodeData);
-        // parseChildren invokes cyclic parsing.
         nodeData.setChildren(parseChildren(elements, typeHierarchy));
         nodeData.setExecutableTypes(groupExecutableTypes(new ExecutableTypeMethodParser(context, nodeData).parse(elements)));
 
+        // resolveChildren invokes cyclic parsing.
+        parsedNodes.put(Utils.getQualifiedName(templateType), nodeData);
+        resolveChildren(nodeData);
+
         return nodeData;
     }
 
@@ -380,6 +386,25 @@
         return fields;
     }
 
+    private void resolveChildren(NodeData node) {
+        for (NodeChildData nodeChild : node.getChildren()) {
+            NodeData fieldNodeData = resolveNode(Utils.fromTypeMirror(nodeChild.getNodeType()));
+            nodeChild.setNode(fieldNodeData);
+            if (fieldNodeData == null) {
+                nodeChild.addError("Node type '%s' is invalid or not a valid Node.", Utils.getQualifiedName(nodeChild.getNodeType()));
+            } else if (!Utils.typeEquals(fieldNodeData.getTypeSystem().getTemplateType().asType(), (node.getTypeSystem().getTemplateType().asType()))) {
+                nodeChild.addError("The @%s of the node and the @%s of the @%s does not match. %s != %s. ", TypeSystem.class.getSimpleName(), TypeSystem.class.getSimpleName(),
+                                NodeChild.class.getSimpleName(), Utils.getSimpleName(node.getTypeSystem().getTemplateType()), Utils.getSimpleName(fieldNodeData.getTypeSystem().getTemplateType()));
+            }
+            List<ExecutableTypeData> types = nodeChild.findGenericExecutableTypes(context);
+            if (types.isEmpty()) {
+                AnnotationValue executeWithValue = Utils.getAnnotationValue(nodeChild.getMessageAnnotation(), "executeWith");
+                nodeChild.addError(executeWithValue, "No generic execute method found with %s evaluated arguments for node type %s.", nodeChild.getExecuteWith().size(),
+                                Utils.getSimpleName(nodeChild.getNodeType()));
+            }
+        }
+    }
+
     private List<NodeChildData> parseChildren(List<? extends Element> elements, final List<TypeElement> typeHierarchy) {
         Set<String> shortCircuits = new HashSet<>();
         for (ExecutableElement method : ElementFilter.methodsIn(elements)) {
@@ -454,12 +479,6 @@
                     continue;
                 }
 
-                NodeData fieldNodeData = resolveNode(Utils.fromTypeMirror(childType));
-                nodeChild.setNode(fieldNodeData);
-                if (fieldNodeData == null) {
-                    nodeChild.addError("Node type '%s' is invalid or not a valid Node.", Utils.getQualifiedName(childType));
-                }
-
                 index++;
             }
         }
@@ -512,12 +531,6 @@
             if (child.getNodeData() == null) {
                 continue;
             }
-
-            List<ExecutableTypeData> types = child.findGenericExecutableTypes(context);
-            if (types.isEmpty()) {
-                child.addError(executeWithValue, "No generic execute method found with %s evaluated arguments for node type %s.", executeWith.size(), Utils.getSimpleName(child.getNodeType()));
-                continue;
-            }
         }
 
         return filteredChildren;
@@ -1006,9 +1019,6 @@
         for (TemplateMethod method : nodeData.getAllTemplateMethods()) {
             unusedElements.remove(method.getMethod());
         }
-        if (nodeData.getExtensionElements() != null) {
-            unusedElements.removeAll(nodeData.getExtensionElements());
-        }
 
         for (NodeFieldData field : nodeData.getFields()) {
             if (field.getGetter() != null) {
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/ShortCircuitParser.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/ShortCircuitParser.java	Sat Aug 31 09:18:58 2013 +0200
@@ -58,7 +58,7 @@
     }
 
     @Override
-    public ShortCircuitData create(TemplateMethod method) {
+    public ShortCircuitData create(TemplateMethod method, boolean invalid) {
         String shortCircuitValue = Utils.getAnnotationValue(String.class, method.getMarkerAnnotation(), "value");
 
         if (!shortCircuitValues.contains(shortCircuitValue)) {
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/SpecializationGroup.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/SpecializationGroup.java	Sat Aug 31 09:18:58 2013 +0200
@@ -41,12 +41,14 @@
     private final List<TypeGuard> typeGuards;
     private final List<GuardData> guards;
 
+    private final NodeData node;
     private final SpecializationData specialization;
     private final List<SpecializationGroup> children = new ArrayList<>();
 
     private SpecializationGroup parent;
 
     private SpecializationGroup(SpecializationData data) {
+        this.node = data.getNode();
         this.assumptions = new ArrayList<>();
         this.typeGuards = new ArrayList<>();
         this.guards = new ArrayList<>();
@@ -61,9 +63,11 @@
     }
 
     public SpecializationGroup(List<SpecializationGroup> children, List<String> assumptionMatches, List<TypeGuard> typeGuardsMatches, List<GuardData> guardMatches) {
+        assert !children.isEmpty() : "children must not be empty";
         this.assumptions = assumptionMatches;
         this.typeGuards = typeGuardsMatches;
         this.guards = guardMatches;
+        this.node = children.get(0).node;
         this.specialization = null;
         updateChildren(children);
     }
@@ -106,18 +110,18 @@
     }
 
     private GuardData findNegatedGuardInPrevious(GuardData guard) {
-        SpecializationGroup previous = this;
-        while ((previous = previous.getPreviousGroup()) != null) {
-            List<GuardData> elseConnectedGuards = previous.getElseConnectableGuards();
+        SpecializationGroup previous = this.getPreviousGroup();
+        if (previous == null) {
+            return null;
+        }
+        List<GuardData> elseConnectedGuards = previous.getElseConnectableGuards();
 
-            if (previous == null || previous.getGuards().size() != elseConnectedGuards.size() + 1) {
-                return null;
-            }
-            GuardData previousGuard = previous.getGuards().get(elseConnectedGuards.size());
-            if (guard.getMethod().equals(previousGuard.getMethod())) {
-                assert guard.isNegated() != previousGuard.isNegated();
-                return guard;
-            }
+        if (previous == null || previous.getGuards().size() != elseConnectedGuards.size() + 1) {
+            return null;
+        }
+        GuardData previousGuard = previous.getGuards().get(elseConnectedGuards.size());
+        if (guard.getMethod().equals(previousGuard.getMethod()) && guard.isNegated() != previousGuard.isNegated()) {
+            return guard;
         }
         return null;
     }
@@ -205,16 +209,29 @@
         for (Iterator<GuardData> iterator = guardMatches.iterator(); iterator.hasNext();) {
             GuardData guardMatch = iterator.next();
 
-            List<TypeMirror> guardTypes = TemplateMethod.getSignatureTypes(guardMatch.getParameters());
-            for (int i = 0; i < guardTypes.size(); i++) {
-                TypeMirror guardType = guardTypes.get(i);
-                int signatureIndex = i + 1;
+            int signatureIndex = 0;
+            for (ActualParameter parameter : guardMatch.getParameters()) {
+                signatureIndex++;
+                if (parameter.getSpecification().isSignature()) {
+                    continue;
+                }
+
+                TypeMirror guardType = parameter.getType();
 
                 // object guards can be safely moved up
                 if (Utils.isObject(guardType)) {
                     continue;
                 }
 
+                // generic guards can be safely moved up
+                SpecializationData generic = first.node.getGenericSpecialization();
+                if (generic != null) {
+                    ActualParameter genericParameter = generic.findParameter(parameter.getLocalName());
+                    if (genericParameter != null && Utils.typeEquals(genericParameter.getType(), guardType)) {
+                        continue;
+                    }
+                }
+
                 // signature index required for moving up guards
                 if (containsIndex(typeGuardsMatches, signatureIndex) || (first.getParent() != null && first.getParent().containsTypeGuardIndex(signatureIndex))) {
                     continue;
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/SpecializationListenerParser.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/SpecializationListenerParser.java	Sat Aug 31 09:18:58 2013 +0200
@@ -47,7 +47,7 @@
     }
 
     @Override
-    public SpecializationListenerData create(TemplateMethod method) {
+    public SpecializationListenerData create(TemplateMethod method, boolean invalid) {
         return new SpecializationListenerData(method);
     }
 
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/SpecializationMethodParser.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/SpecializationMethodParser.java	Sat Aug 31 09:18:58 2013 +0200
@@ -44,7 +44,7 @@
     }
 
     @Override
-    public SpecializationData create(TemplateMethod method) {
+    public SpecializationData create(TemplateMethod method, boolean invalid) {
         return parseSpecialization(method);
     }
 
@@ -57,7 +57,6 @@
         int order = Utils.getAnnotationValue(Integer.class, method.getMarkerAnnotation(), "order");
         if (order < 0 && order != Specialization.DEFAULT_ORDER) {
             method.addError("Invalid order attribute %d. The value must be >= 0 or the default value.");
-            return null;
         }
 
         AnnotationValue rewriteValue = Utils.getAnnotationValue(method.getMarkerAnnotation(), "rewriteOn");
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/template/MessageContainer.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/template/MessageContainer.java	Sat Aug 31 09:18:58 2013 +0200
@@ -51,13 +51,21 @@
 
     public abstract Element getMessageElement();
 
-    public final void emitMessages(TypeElement baseElement, Log log) {
-        emitMessagesImpl(baseElement, log, new HashSet<MessageContainer>());
+    public final void emitMessages(ProcessorContext context, TypeElement baseElement, Log log) {
+        emitMessagesImpl(context, baseElement, log, new HashSet<MessageContainer>(), null);
     }
 
-    private void emitMessagesImpl(TypeElement baseElement, Log log, Set<MessageContainer> visitedSinks) {
+    private void emitMessagesImpl(ProcessorContext context, TypeElement baseElement, Log log, Set<MessageContainer> visitedSinks, List<Message> verifiedMessages) {
+        List<Message> childMessages;
+        if (verifiedMessages == null) {
+            childMessages = collectMessagesWithElementChildren(new HashSet<MessageContainer>(), getMessageElement());
+        } else {
+            childMessages = verifiedMessages;
+        }
+        verifyExpectedMessages(context, log, childMessages);
+
         for (Message message : getMessages()) {
-            emitDefault(baseElement, log, message);
+            emitDefault(context, baseElement, log, message);
         }
 
         for (MessageContainer sink : findChildContainers()) {
@@ -66,18 +74,95 @@
             }
 
             visitedSinks.add(sink);
-            sink.emitMessagesImpl(baseElement, log, visitedSinks);
+            if (sink.getMessageElement() == this.getMessageElement()) {
+                sink.emitMessagesImpl(context, baseElement, log, visitedSinks, childMessages);
+            } else {
+                sink.emitMessagesImpl(context, baseElement, log, visitedSinks, null);
+            }
+        }
+    }
+
+    private List<Message> collectMessagesWithElementChildren(Set<MessageContainer> visitedSinks, Element e) {
+        if (visitedSinks.contains(this)) {
+            return Collections.emptyList();
+        }
+        visitedSinks.add(this);
+
+        List<Message> foundMessages = new ArrayList<>();
+        if (Utils.typeEquals(getMessageElement().asType(), e.asType())) {
+            foundMessages.addAll(getMessages());
+        }
+        for (MessageContainer sink : findChildContainers()) {
+            foundMessages.addAll(sink.collectMessagesWithElementChildren(visitedSinks, e));
+        }
+        return foundMessages;
+    }
+
+    private void verifyExpectedMessages(ProcessorContext context, Log log, List<Message> msgs) {
+        TypeElement expectError = context.getTruffleTypes().getExpectError();
+        if (expectError != null) {
+            Element element = getMessageElement();
+            AnnotationMirror mirror = Utils.findAnnotationMirror(element.getAnnotationMirrors(), expectError);
+            if (mirror != null) {
+                List<String> values = Utils.getAnnotationValueList(String.class, mirror, "value");
+                if (values == null) {
+                    values = Collections.emptyList();
+                }
+                if (values.size() != msgs.size()) {
+                    log.message(Kind.ERROR, element, mirror, Utils.getAnnotationValue(mirror, "value"), String.format("Error count expected %s but was %s.", values.size(), msgs.size()));
+                }
+            }
         }
     }
 
-    private void emitDefault(TypeElement baseType, Log log, Message message) {
+    private void emitDefault(ProcessorContext context, TypeElement baseType, Log log, Message message) {
+        Kind kind = message.getKind();
+
+        Element messageElement = getMessageElement();
+        AnnotationMirror messageAnnotation = getMessageAnnotation();
+        AnnotationValue messageValue = getMessageAnnotationValue();
+        if (message.getAnnotationValue() != null) {
+            messageValue = message.getAnnotationValue();
+        }
+
+        String text = message.getText();
+
         TypeElement rootEnclosing = Utils.findRootEnclosingType(getMessageElement());
-        if (rootEnclosing != null && Utils.typeEquals(baseType.asType(), rootEnclosing.asType()) && this == message.getOriginalContainer()) {
-            log.message(message.getKind(), getMessageElement(), getMessageAnnotation(), getMessageAnnotationValue(), message.getText());
-        } else {
+        TypeElement baseEnclosing = Utils.findRootEnclosingType(baseType);
+        if (rootEnclosing == null || !Utils.typeEquals(baseEnclosing.asType(), rootEnclosing.asType()) || this != message.getOriginalContainer()) {
+            // redirect message
             MessageContainer original = message.getOriginalContainer();
-            log.message(message.getKind(), baseType, null, null, wrapText(original.getMessageElement(), original.getMessageAnnotation(), message.getText()));
+            messageElement = baseType;
+            messageAnnotation = null;
+            messageValue = null;
+            text = wrapText(original.getMessageElement(), original.getMessageAnnotation(), message.getText());
         }
+
+        TypeElement expectError = context.getTruffleTypes().getExpectError();
+        if (expectError != null) {
+            AnnotationMirror mirror = Utils.findAnnotationMirror(messageElement.getAnnotationMirrors(), expectError);
+            if (mirror != null) {
+                List<String> expectedTexts = Utils.getAnnotationValueList(String.class, mirror, "value");
+                boolean found = false;
+                for (String expectedText : expectedTexts) {
+                    if (expectedText.endsWith("%") && text.startsWith(expectedText.substring(0, expectedText.length() - 1))) {
+                        found = true;
+                        break;
+                    } else if (text.equals(expectedText)) {
+                        found = true;
+                        break;
+                    }
+                }
+                if (!found) {
+                    log.message(kind, messageElement, mirror, Utils.getAnnotationValue(mirror, "value"), "Message expected one of '%s' but was '%s'.", expectedTexts, text);
+                } else {
+                    return;
+                }
+
+            }
+        }
+
+        log.message(kind, messageElement, messageAnnotation, messageValue, text);
     }
 
     private static String wrapText(Element element, AnnotationMirror mirror, String text) {
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/template/Template.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/template/Template.java	Sat Aug 31 09:18:58 2013 +0200
@@ -27,7 +27,6 @@
 import javax.lang.model.element.*;
 
 import com.oracle.truffle.dsl.processor.*;
-import com.oracle.truffle.dsl.processor.api.element.*;
 import com.oracle.truffle.dsl.processor.typesystem.*;
 
 public abstract class Template extends MessageContainer {
@@ -36,8 +35,6 @@
     private final String templateMethodName;
     private final AnnotationMirror annotation;
 
-    private List<? extends WritableElement> extensionElements;
-
     public Template(TypeElement templateType, String templateMethodName, AnnotationMirror annotation) {
         this.templateType = templateType;
         this.templateMethodName = templateMethodName;
@@ -68,14 +65,6 @@
         return annotation;
     }
 
-    public List<? extends WritableElement> getExtensionElements() {
-        return extensionElements;
-    }
-
-    public void setExtensionElements(List<? extends WritableElement> extensionMethods) {
-        this.extensionElements = extensionMethods;
-    }
-
     @Override
     public String toString() {
         return getClass().getSimpleName() + "[" + Utils.getSimpleName(getTemplateType()) + "]";
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/template/TemplateMethodParser.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/template/TemplateMethodParser.java	Sat Aug 31 09:18:58 2013 +0200
@@ -43,7 +43,7 @@
     protected final T template;
 
     private boolean emitErrors = true;
-    private boolean parseNullOnError = true;
+    private boolean parseNullOnError = false;
 
     public TemplateMethodParser(ProcessorContext context, T template) {
         this.template = template;
@@ -76,7 +76,7 @@
 
     public abstract MethodSpec createSpecification(ExecutableElement method, AnnotationMirror mirror);
 
-    public abstract E create(TemplateMethod method);
+    public abstract E create(TemplateMethod method, boolean invalid);
 
     public abstract boolean isParsable(ExecutableElement method);
 
@@ -104,7 +104,8 @@
             E parsedMethod = parse(method, mirror);
 
             if (method.getModifiers().contains(Modifier.PRIVATE) && emitErrors) {
-                parsedMethod.addError("Method must not be private.");
+                parsedMethod.addError("Method annotated with @%s must not be private.", getAnnotationType().getSimpleName());
+                parsedMethods.add(parsedMethod);
                 valid = false;
                 continue;
             }
@@ -142,7 +143,7 @@
         ActualParameter returnTypeMirror = matchParameter(returnTypeSpec, method.getReturnType(), template, 0, false);
         if (returnTypeMirror == null) {
             if (emitErrors) {
-                E invalidMethod = create(new TemplateMethod(id, template, methodSpecification, method, annotation, returnTypeMirror, Collections.<ActualParameter> emptyList()));
+                E invalidMethod = create(new TemplateMethod(id, template, methodSpecification, method, annotation, returnTypeMirror, Collections.<ActualParameter> emptyList()), true);
                 String expectedReturnType = returnTypeSpec.toSignatureString(true);
                 String actualReturnType = Utils.getSimpleName(method.getReturnType());
 
@@ -163,7 +164,7 @@
         List<ActualParameter> parameters = parseParameters(methodSpecification, parameterTypes);
         if (parameters == null) {
             if (isEmitErrors()) {
-                E invalidMethod = create(new TemplateMethod(id, template, methodSpecification, method, annotation, returnTypeMirror, Collections.<ActualParameter> emptyList()));
+                E invalidMethod = create(new TemplateMethod(id, template, methodSpecification, method, annotation, returnTypeMirror, Collections.<ActualParameter> emptyList()), true);
                 String message = String.format("Method signature %s does not match to the expected signature: \n%s", createActualSignature(methodSpecification, method),
                                 methodSpecification.toSignatureString(method.getSimpleName().toString()));
                 invalidMethod.addError(message);
@@ -173,7 +174,7 @@
             }
         }
 
-        return create(new TemplateMethod(id, template, methodSpecification, method, annotation, returnTypeMirror, parameters));
+        return create(new TemplateMethod(id, template, methodSpecification, method, annotation, returnTypeMirror, parameters), false);
     }
 
     private static String createActualSignature(MethodSpec spec, ExecutableElement method) {
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/GuardData.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/GuardData.java	Sat Aug 31 09:18:58 2013 +0200
@@ -60,7 +60,7 @@
 
     @Override
     public String toString() {
-        return getMethodName() + getParameters().toString();
+        return (negated ? "!" : "") + getMethodName() + getParameters().toString();
     }
 
 }
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/GuardParser.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/GuardParser.java	Sat Aug 31 09:18:58 2013 +0200
@@ -79,7 +79,7 @@
     }
 
     @Override
-    public GuardData create(TemplateMethod method) {
+    public GuardData create(TemplateMethod method, boolean invalid) {
         GuardData guard = new GuardData(method, specialization, negated);
         /*
          * Update parameters in way that parameter specifications match again the node field names
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/ImplicitCastData.java	Sat Aug 31 09:18:58 2013 +0200
@@ -0,0 +1,33 @@
+/*
+ * Copyright (c) 2012, 2012, 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.dsl.processor.typesystem;
+
+import com.oracle.truffle.dsl.processor.template.*;
+
+public class ImplicitCastData extends TemplateMethod {
+
+    public ImplicitCastData(TemplateMethod method) {
+        super(method);
+    }
+
+}
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/ImplicitCastParser.java	Sat Aug 31 09:18:58 2013 +0200
@@ -0,0 +1,62 @@
+/*
+ * Copyright (c) 2012, 2012, 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.dsl.processor.typesystem;
+
+import java.lang.annotation.*;
+import java.util.*;
+
+import javax.lang.model.element.*;
+import javax.lang.model.type.*;
+
+import com.oracle.truffle.api.dsl.*;
+import com.oracle.truffle.dsl.processor.*;
+import com.oracle.truffle.dsl.processor.template.*;
+
+public class ImplicitCastParser extends TypeSystemMethodParser<ImplicitCastData> {
+
+    public ImplicitCastParser(ProcessorContext context, TypeSystemData typeSystem) {
+        super(context, typeSystem);
+    }
+
+    @Override
+    public Class<? extends Annotation> getAnnotationType() {
+        return ImplicitCast.class;
+    }
+
+    @Override
+    public MethodSpec createSpecification(ExecutableElement method, AnnotationMirror mirror) {
+        List<TypeMirror> types = new ArrayList<>();
+        for (TypeData typeData : getTypeSystem().getTypes()) {
+            types.add(typeData.getPrimitiveType());
+        }
+        MethodSpec spec = new MethodSpec(new ParameterSpec("target", types));
+        spec.addRequired(new ParameterSpec("source", types)).setSignature(true);
+        return spec;
+    }
+
+    @Override
+    public ImplicitCastData create(TemplateMethod method, boolean invalid) {
+        return new ImplicitCastData(method);
+    }
+
+}
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/TypeCastParser.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/TypeCastParser.java	Sat Aug 31 09:18:58 2013 +0200
@@ -48,10 +48,25 @@
     }
 
     @Override
-    public TypeCastData create(TemplateMethod method) {
+    public TypeCastData create(TemplateMethod method, boolean invalid) {
+        if (invalid) {
+            return new TypeCastData(method, null, null);
+        }
+
         TypeData targetType = findTypeByMethodName(method, "as");
         ActualParameter parameter = method.findParameter("valueValue");
-        return new TypeCastData(method, parameter.getTypeSystemType(), targetType);
+
+        TypeData sourceType = null;
+        if (parameter != null) {
+            sourceType = getTypeSystem().findTypeData(parameter.getType());
+        }
+        TypeCastData cast = new TypeCastData(method, sourceType, targetType);
+
+        if (targetType != method.getReturnType().getTypeSystemType()) {
+            cast.addError("Cast type %s does not match to the returned type %s.", Utils.getSimpleName(targetType.getPrimitiveType()),
+                            method.getReturnType() != null ? Utils.getSimpleName(method.getReturnType().getTypeSystemType().getPrimitiveType()) : null);
+        }
+        return cast;
     }
 
     @Override
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/TypeCheckParser.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/TypeCheckParser.java	Sat Aug 31 09:18:58 2013 +0200
@@ -48,7 +48,7 @@
     }
 
     @Override
-    public TypeCheckData create(TemplateMethod method) {
+    public TypeCheckData create(TemplateMethod method, boolean invalid) {
         TypeData checkedType = findTypeByMethodName(method, "is");
         assert checkedType != null;
         ActualParameter parameter = method.findParameter("valueValue");
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/TypeSystemData.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/TypeSystemData.java	Sat Aug 31 09:18:58 2013 +0200
@@ -36,6 +36,9 @@
     private List<TypeMirror> primitiveTypeMirrors = new ArrayList<>();
     private List<TypeMirror> boxedTypeMirrors = new ArrayList<>();
 
+    private List<TypeCastData> casts;
+    private List<TypeCheckData> checks;
+
     private TypeMirror genericType;
     private TypeData voidType;
 
@@ -58,6 +61,14 @@
         }
     }
 
+    public void setCasts(List<TypeCastData> casts) {
+        this.casts = casts;
+    }
+
+    public void setChecks(List<TypeCheckData> checks) {
+        this.checks = checks;
+    }
+
     void setGenericType(TypeMirror genericType) {
         this.genericType = genericType;
     }
@@ -72,6 +83,12 @@
         if (types != null) {
             sinks.addAll(types);
         }
+        if (checks != null) {
+            sinks.addAll(checks);
+        }
+        if (casts != null) {
+            sinks.addAll(casts);
+        }
         return sinks;
     }
 
--- a/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/TypeSystemParser.java	Sat Aug 31 09:18:42 2013 +0200
+++ b/graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/TypeSystemParser.java	Sat Aug 31 09:18:58 2013 +0200
@@ -37,7 +37,7 @@
 
 public class TypeSystemParser extends TemplateParser<TypeSystemData> {
 
-    public static final List<Class<TypeSystem>> ANNOTATIONS = Arrays.asList(TypeSystem.class);
+    public static final List<Class<? extends Annotation>> ANNOTATIONS = Arrays.asList(TypeSystem.class, ExpectError.class);
 
     public TypeSystemParser(ProcessorContext c) {
         super(c);
@@ -84,11 +84,17 @@
         verifyExclusiveMethodAnnotation(typeSystem, TypeCast.class, TypeCheck.class);
 
         List<Element> elements = new ArrayList<>(context.getEnvironment().getElementUtils().getAllMembers(templateType));
-
+        List<ImplicitCastData> implicitCasts = new ImplicitCastParser(context, typeSystem).parse(elements);
         List<TypeCastData> casts = new TypeCastParser(context, typeSystem).parse(elements);
         List<TypeCheckData> checks = new TypeCheckParser(context, typeSystem).parse(elements);
 
-        if (casts == null || checks == null) {
+        if (casts == null || checks == null || implicitCasts == null) {
+            return typeSystem;
+        }
+        typeSystem.setCasts(casts);
+        typeSystem.setChecks(checks);
+
+        if (typeSystem.hasErrors()) {
             return typeSystem;
         }
 
--- a/src/gpu/ptx/vm/gpu_ptx.cpp	Sat Aug 31 09:18:42 2013 +0200
+++ b/src/gpu/ptx/vm/gpu_ptx.cpp	Sat Aug 31 09:18:58 2013 +0200
@@ -29,12 +29,14 @@
 #include "utilities/ostream.hpp"
 #include "memory/allocation.hpp"
 #include "memory/allocation.inline.hpp"
+#include "kernelArguments.hpp"
 
 void * gpu::Ptx::_device_context;
+int    gpu::Ptx::_cu_device = 0;
 
 gpu::Ptx::cuda_cu_init_func_t gpu::Ptx::_cuda_cu_init;
 gpu::Ptx::cuda_cu_ctx_create_func_t gpu::Ptx::_cuda_cu_ctx_create;
-gpu::Ptx::cuda_cu_ctx_detach_func_t gpu::Ptx::_cuda_cu_ctx_detach;
+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_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;
@@ -44,6 +46,8 @@
 gpu::Ptx::cuda_cu_launch_kernel_func_t gpu::Ptx::_cuda_cu_launch_kernel;
 gpu::Ptx::cuda_cu_module_get_function_func_t gpu::Ptx::_cuda_cu_module_get_function;
 gpu::Ptx::cuda_cu_module_load_data_ex_func_t gpu::Ptx::_cuda_cu_module_load_data_ex;
+gpu::Ptx::cuda_cu_memcpy_dtoh_func_t gpu::Ptx::_cuda_cu_memcpy_dtoh;
+gpu::Ptx::cuda_cu_memfree_func_t gpu::Ptx::_cuda_cu_memfree;
 
 void gpu::probe_linkage() {
 #if defined(__APPLE__) || defined(LINUX)
@@ -67,9 +71,9 @@
   }
 }
 
-bool gpu::execute_kernel(address kernel, JavaCallArguments * jca) {
+bool gpu::execute_kernel(address kernel, PTXKernelArguments & ptxka, JavaValue& ret) {
   if (gpu::has_gpu_linkage()) {
-    return (gpu::Ptx::execute_kernel(kernel, jca));
+    return (gpu::Ptx::execute_kernel(kernel, ptxka, ret));
   } else {
     return false;
   }
@@ -108,8 +112,7 @@
   /* Get the handle to the first compute device */
   int device_id = 0;
   /* Compute-capable device handle */
-  int cu_device = 0;
-  status = _cuda_cu_device_get(&cu_device, device_id);
+  status = _cuda_cu_device_get(&_cu_device, device_id);
 
   if (status != GRAAL_CUDA_SUCCESS) {
     tty->print_cr("[CUDA] Failed to get handle of first compute-capable device i.e., the one at ordinal: %d", device_id);
@@ -122,42 +125,42 @@
 
   /* Get device attributes */
   int minor, major, unified_addressing;
-  status = _cuda_cu_device_get_attribute(&minor, GRAAL_CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, cu_device);
+  status = _cuda_cu_device_get_attribute(&minor, GRAAL_CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, _cu_device);
 
   if (status != GRAAL_CUDA_SUCCESS) {
-    tty->print_cr("[CUDA] Failed to get minor attribute of device: %d", cu_device);
+    tty->print_cr("[CUDA] Failed to get minor attribute of device: %d", _cu_device);
     return false;
   }
 
-  status = _cuda_cu_device_get_attribute(&major, GRAAL_CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, cu_device);
+  status = _cuda_cu_device_get_attribute(&major, GRAAL_CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, _cu_device);
 
   if (status != GRAAL_CUDA_SUCCESS) {
-    tty->print_cr("[CUDA] Failed to get major attribute of device: %d", cu_device);
+    tty->print_cr("[CUDA] Failed to get major attribute of device: %d", _cu_device);
     return false;
   }
 
   if (TraceGPUInteraction) {
-    tty->print_cr("[CUDA] Compatibility version of device %d: %d.%d", cu_device, major, minor);
+    tty->print_cr("[CUDA] Compatibility version of device %d: %d.%d", _cu_device, major, minor);
   }
 
-  status = _cuda_cu_device_get_attribute(&unified_addressing, GRAAL_CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING, cu_device);
+  status = _cuda_cu_device_get_attribute(&unified_addressing, GRAAL_CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING, _cu_device);
 
   if (status != GRAAL_CUDA_SUCCESS) {
-    tty->print_cr("[CUDA] Failed to query unified addressing mode of device: %d", cu_device);
+    tty->print_cr("[CUDA] Failed to query unified addressing mode of device: %d", _cu_device);
     return false;
   }
 
   if (TraceGPUInteraction) {
-    tty->print_cr("[CUDA] Unified addressing support on device %d: %d", cu_device, unified_addressing);
+    tty->print_cr("[CUDA] Unified addressing support on device %d: %d", _cu_device, unified_addressing);
   }
 
 
   /* Get device name */
   char device_name[256];
-  status = _cuda_cu_device_get_name(device_name, 256, cu_device);
+  status = _cuda_cu_device_get_name(device_name, 256, _cu_device);
 
   if (status != GRAAL_CUDA_SUCCESS) {
-    tty->print_cr("[CUDA] Failed to get name of device: %d", cu_device);
+    tty->print_cr("[CUDA] Failed to get name of device: %d", _cu_device);
     return false;
   }
 
@@ -165,18 +168,6 @@
     tty->print_cr("[CUDA] Using %s", device_name);
   }
 
-  /* Create CUDA context */
-  status = _cuda_cu_ctx_create(&_device_context, 0, cu_device);
-
-  if (status != GRAAL_CUDA_SUCCESS) {
-    tty->print_cr("[CUDA] Failed to create CUDA context for device: %d", cu_device);
-    return false;
-  }
-
-  if (TraceGPUInteraction) {
-    tty->print_cr("[CUDA] Success: Created context for device: %d", cu_device);
-  }
-
   return true;
 }
 
@@ -210,8 +201,20 @@
 
   }
 
+  /* Create CUDA context to compile and execute the kernel */
+  int status = _cuda_cu_ctx_create(&_device_context, 0, _cu_device);
+
+  if (status != GRAAL_CUDA_SUCCESS) {
+    tty->print_cr("[CUDA] Failed to create CUDA context for device: %d", _cu_device);
+    return NULL;
+  }
+
+  if (TraceGPUInteraction) {
+    tty->print_cr("[CUDA] Success: Created context for device: %d", _cu_device);
+  }
+
   /* Load module's data with compiler options */
-  int status = _cuda_cu_module_load_data_ex(&cu_module, code, jit_num_options,
+  status = _cuda_cu_module_load_data_ex(&cu_module, (void*) code, jit_num_options,
                                             jit_options, (void **)jit_option_values);
   if (status != GRAAL_CUDA_SUCCESS) {
     if (status == GRAAL_CUDA_ERROR_NO_BINARY_FOR_GPU) {
@@ -238,10 +241,11 @@
   if (TraceGPUInteraction) {
     tty->print_cr("[CUDA] Got function handle for %s", name);
   }
+
   return cu_function;
 }
 
-bool gpu::Ptx::execute_kernel(address kernel, JavaCallArguments * jca) {
+bool gpu::Ptx::execute_kernel(address kernel, PTXKernelArguments &ptxka, JavaValue &ret) {
   // grid dimensionality
   unsigned int gridX = 1;
   unsigned int gridY = 1;
@@ -252,14 +256,11 @@
   unsigned int blockY = 1;
   unsigned int blockZ = 1;
   
-  int *cu_function = (int *)kernel;
+  struct CUfunc_st* cu_function = (struct CUfunc_st*) kernel;
 
-  char * paramBuffer = (char *) jca->parameters();
-  size_t paramBufferSz = (size_t) jca->size_of_parameters();
-
-  void * config[] = {
-    GRAAL_CU_LAUNCH_PARAM_BUFFER_POINTER, paramBuffer,
-    GRAAL_CU_LAUNCH_PARAM_BUFFER_SIZE, &paramBufferSz,
+  void * config[5] = {
+    GRAAL_CU_LAUNCH_PARAM_BUFFER_POINTER, ptxka._kernelArgBuffer,
+    GRAAL_CU_LAUNCH_PARAM_BUFFER_SIZE, &(ptxka._bufferOffset),
     GRAAL_CU_LAUNCH_PARAM_END
   };
 
@@ -270,10 +271,11 @@
   if (TraceGPUInteraction) {
     tty->print_cr("[CUDA] launching kernel");
   }
+
   int status = _cuda_cu_launch_kernel(cu_function,
                                       gridX, gridY, gridZ,
                                       blockX, blockY, blockZ,
-                                      0, NULL, NULL, config);
+                                      0, NULL, NULL, (void **) &config);
   if (status != GRAAL_CUDA_SUCCESS) {
     tty->print_cr("[CUDA] Failed to launch kernel");
     return false;
@@ -282,7 +284,72 @@
   if (TraceGPUInteraction) {
     tty->print_cr("[CUDA] Success: Kernel Launch");
   }
-  return status == 0;  // GRAAL_CUDA_SUCCESS
+
+  status = _cuda_cu_ctx_synchronize();
+
+  if (status != GRAAL_CUDA_SUCCESS) {
+    tty->print_cr("[CUDA] Failed to synchronize launched kernel (%d)", status);
+    return false;
+  }
+
+  if (TraceGPUInteraction) {
+    tty->print_cr("[CUDA] Success: Synchronized launch kernel");
+  }
+
+
+  // Get the result. TODO: Move this code to get_return_oop()
+  BasicType return_type = ptxka.get_ret_type();
+  switch (return_type) {
+     case T_INT :
+       {
+         int return_val;
+         status = gpu::Ptx::_cuda_cu_memcpy_dtoh(&return_val, ptxka._return_value_ptr, T_INT_BYTE_SIZE);
+         if (status != GRAAL_CUDA_SUCCESS) {
+           tty->print_cr("[CUDA] *** Error (%d) Failed to copy value to device argument", status);
+           return false;
+         }
+         ret.set_jint(return_val);
+       }
+       break;
+     case T_LONG :
+       {
+         long return_val;
+         status = gpu::Ptx::_cuda_cu_memcpy_dtoh(&return_val, ptxka._return_value_ptr, T_LONG_BYTE_SIZE);
+         if (status != GRAAL_CUDA_SUCCESS) {
+           tty->print_cr("[CUDA] *** Error (%d) Failed to copy value to device argument", status);
+           return false;
+         }
+         ret.set_jlong(return_val);
+       }
+       break;
+     default:
+       tty->print_cr("[CUDA] TODO *** Unhandled return type");
+  }
+
+
+  // Free device memory allocated for result
+  status = gpu::Ptx::_cuda_cu_memfree(ptxka._return_value_ptr);
+  if (status != GRAAL_CUDA_SUCCESS) {
+    tty->print_cr("[CUDA] *** Error (%d) Failed to free device memory of return value", status);
+    return false;
+  }
+
+  if (TraceGPUInteraction) {
+    tty->print_cr("[CUDA] Success: Freed device memory of return value");
+  }
+
+  // Destroy context
+  status = gpu::Ptx::_cuda_cu_ctx_destroy(_device_context);
+  if (status != GRAAL_CUDA_SUCCESS) {
+    tty->print_cr("[CUDA] *** Error (%d) Failed to destroy context", status);
+    return false;
+  }
+
+  if (TraceGPUInteraction) {
+    tty->print_cr("[CUDA] Success: Destroy context");
+  }
+
+  return (status == GRAAL_CUDA_SUCCESS);
 }
 
 #if defined(LINUX)
@@ -305,8 +372,8 @@
         CAST_TO_FN_PTR(cuda_cu_init_func_t, os::dll_lookup(handle, "cuInit"));
       _cuda_cu_ctx_create =
         CAST_TO_FN_PTR(cuda_cu_ctx_create_func_t, os::dll_lookup(handle, "cuCtxCreate"));
-      _cuda_cu_ctx_detach =
-        CAST_TO_FN_PTR(cuda_cu_ctx_detach_func_t, os::dll_lookup(handle, "cuCtxDetach"));
+      _cuda_cu_ctx_destroy =
+        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_device_get_count =
@@ -325,6 +392,15 @@
         CAST_TO_FN_PTR(cuda_cu_module_load_data_ex_func_t, os::dll_lookup(handle, "cuModuleLoadDataEx"));
       _cuda_cu_launch_kernel =
         CAST_TO_FN_PTR(cuda_cu_launch_kernel_func_t, os::dll_lookup(handle, "cuLaunchKernel"));
+      _cuda_cu_memalloc =
+        CAST_TO_FN_PTR(cuda_cu_memalloc_func_t, os::dll_lookup(handle, "cuMemAlloc"));
+      _cuda_cu_memfree =
+        CAST_TO_FN_PTR(cuda_cu_memfree_func_t, os::dll_lookup(handle, "cuMemFree"));
+      _cuda_cu_memcpy_htod =
+        CAST_TO_FN_PTR(cuda_cu_memcpy_htod_func_t, os::dll_lookup(handle, "cuMemcpyHtoD"));
+      _cuda_cu_memcpy_dtoh =
+        CAST_TO_FN_PTR(cuda_cu_memcpy_dtoh_func_t, os::dll_lookup(handle, "cuMemcpyDtoH"));
+
       if (TraceGPUInteraction) {
         tty->print_cr("[CUDA] Success: library linkage");
       }
--- a/src/gpu/ptx/vm/gpu_ptx.hpp	Sat Aug 31 09:18:42 2013 +0200
+++ b/src/gpu/ptx/vm/gpu_ptx.hpp	Sat Aug 31 09:18:58 2013 +0200
@@ -25,7 +25,7 @@
 #ifndef GPU_PTX_HPP
 #define GPU_PTX_HPP
 
-/* 
+/*
  * Some useful macro definitions from publicly available cuda.h.
  * These definitions are for convenience.
  */
@@ -44,7 +44,7 @@
  * End of array terminator for the extra parameter to
  * ::cuLaunchKernel
  */
-#define GRAAL_CU_LAUNCH_PARAM_END            ((void *) 0x00)
+#define GRAAL_CU_LAUNCH_PARAM_END            ((void*) 0x00)
 
 /**
  * Indicator that the next value in the  extra parameter to
@@ -55,7 +55,7 @@
  *  extra array, then ::GRAAL_CU_LAUNCH_PARAM_BUFFER_POINTER will have no
  * effect.
  */
-#define GRAAL_CU_LAUNCH_PARAM_BUFFER_POINTER ((void *) 0x01)
+#define GRAAL_CU_LAUNCH_PARAM_BUFFER_POINTER ((void*) 0x01)
 
 /**
  * Indicator that the next value in the  extra parameter to
@@ -65,7 +65,7 @@
  * in the extra array if the value associated with
  * ::GRAAL_CU_LAUNCH_PARAM_BUFFER_SIZE is not zero.
  */
-#define GRAAL_CU_LAUNCH_PARAM_BUFFER_SIZE    ((void *) 0x02)
+#define GRAAL_CU_LAUNCH_PARAM_BUFFER_SIZE    ((void*) 0x02)
 
 class Ptx {
   friend class gpu;
@@ -74,28 +74,39 @@
   static bool probe_linkage();
   static bool initialize_gpu();
   static void * generate_kernel(unsigned char *code, int code_len, const char *name);
-  static bool execute_kernel(address kernel, JavaCallArguments *);
-  
+  static bool execute_kernel(address kernel, PTXKernelArguments & ka, JavaValue &ret);
+public:
+#if defined(__x86_64) || defined(AMD64) || defined(_M_AMD64)
+  typedef unsigned long long CUdeviceptr;
+#else
+  typedef unsigned int CUdeviceptr;
+#endif
+
 private:
   typedef int (*cuda_cu_init_func_t)(unsigned int);
-  typedef int (*cuda_cu_ctx_create_func_t)(void *, int, int);
-  typedef int (*cuda_cu_ctx_detach_func_t)(int *);
-  typedef int (*cuda_cu_ctx_synchronize_func_t)(int *);
-  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);
-  typedef int (*cuda_cu_device_compute_capability_func_t)(int *, int *, int);
-  typedef int (*cuda_cu_device_get_attribute_func_t)(int *, int, int);
-  typedef int (*cuda_cu_launch_kernel_func_t)(void *,
+  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_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);
+  typedef int (*cuda_cu_device_compute_capability_func_t)(int*, int*, int);
+  typedef int (*cuda_cu_device_get_attribute_func_t)(int*, int, int);
+  typedef int (*cuda_cu_launch_kernel_func_t)(struct CUfunc_st*,
                                               unsigned int, unsigned int, unsigned int,
                                               unsigned int, unsigned int, unsigned int,
-                                              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 **);
+                                              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_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);
 
+public:
   static cuda_cu_init_func_t                      _cuda_cu_init;
   static cuda_cu_ctx_create_func_t                _cuda_cu_ctx_create;
-  static cuda_cu_ctx_detach_func_t                _cuda_cu_ctx_detach;
+  static cuda_cu_ctx_destroy_func_t               _cuda_cu_ctx_destroy;
   static cuda_cu_ctx_synchronize_func_t           _cuda_cu_ctx_synchronize;
   static cuda_cu_device_get_count_func_t          _cuda_cu_device_get_count;
   static cuda_cu_device_get_name_func_t           _cuda_cu_device_get_name;
@@ -105,8 +116,13 @@
   static cuda_cu_launch_kernel_func_t             _cuda_cu_launch_kernel;
   static cuda_cu_module_get_function_func_t       _cuda_cu_module_get_function;
   static cuda_cu_module_load_data_ex_func_t       _cuda_cu_module_load_data_ex;
+  static cuda_cu_memalloc_func_t                  _cuda_cu_memalloc;
+  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;
 
 protected:
-  static void * _device_context;
+  static void* _device_context;
+  static int _cu_device;
 };
 #endif // GPU_PTX_HPP
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/gpu/ptx/vm/kernelArguments.cpp	Sat Aug 31 09:18:58 2013 +0200
@@ -0,0 +1,130 @@
+/*
+ * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved.
+ * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
+ *
+ * This code is free software; you can redistribute it and/or modify it
+ * under the terms of the GNU General Public License version 2 only, as
+ * published by the Free Software Foundation.
+ *
+ * This code is distributed in the hope that it will be useful, but WITHOUT
+ * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+ * FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+ * version 2 for more details (a copy is included in the LICENSE file that
+ * accompanied this code).
+ *
+ * You should have received a copy of the GNU General Public License version
+ * 2 along with this work; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
+ *
+ * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
+ * or visit www.oracle.com if you need additional information or have any
+ * questions.
+ *
+ */
+
+#include "precompiled.hpp"
+#include "kernelArguments.hpp"
+#include "runtime/javaCalls.hpp"
+
+gpu::Ptx::cuda_cu_memalloc_func_t gpu::Ptx::_cuda_cu_memalloc;
+gpu::Ptx::cuda_cu_memcpy_htod_func_t gpu::Ptx::_cuda_cu_memcpy_htod;
+
+// Get next java argument
+oop PTXKernelArguments::next_arg(BasicType expectedType) {
+  assert(_index < _args->length(), "out of bounds");
+  oop arg=((objArrayOop) (_args))->obj_at(_index++);
+  assert(expectedType == T_OBJECT || java_lang_boxing_object::is_instance(arg, expectedType), "arg type mismatch");
+  return arg;
+}
+
+void PTXKernelArguments::do_int()    { 
+  // If the parameter is a return value, 
+  if (is_return_type()) {
+    // Allocate device memory for T_INT return value pointer on device. Size in bytes
+    int status = gpu::Ptx::_cuda_cu_memalloc(&_return_value_ptr, T_INT_BYTE_SIZE);
+    if (status != GRAAL_CUDA_SUCCESS) {
+      tty->print_cr("[CUDA] *** Error (%d) Failed to allocate memory for return value pointer on device", status);
+      _success = false;
+      return;
+    }
+    // Push _return_value_ptr to _kernelBuffer
+    *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = _return_value_ptr;
+    _bufferOffset += sizeof(_return_value_ptr);
+  }
+  else {
+    // Get the next java argument and its value which should be a T_INT
+    oop arg = next_arg(T_INT);
+    // Copy the java argument value to kernelArgBuffer
+    jvalue intval;
+    if (java_lang_boxing_object::get_value(arg, &intval) != T_INT) {
+      tty->print_cr("[CUDA] *** Error: Unexpected argument type; expecting T_INT");
+      _success = false;
+      return;
+    }
+    *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = intval.i;
+    _bufferOffset += sizeof(intval.i);
+  }
+  return;
+}
+
+void PTXKernelArguments::do_long()    { 
+  // If the parameter is a return value, 
+  if (is_return_type()) {
+    // Allocate device memory for T_LONG return value pointer on device. Size in bytes
+    int status = gpu::Ptx::_cuda_cu_memalloc(&_return_value_ptr, T_LONG_BYTE_SIZE);
+    if (status != GRAAL_CUDA_SUCCESS) {
+      tty->print_cr("[CUDA] *** Error (%d) Failed to allocate memory for return value pointer on device", status);
+      _success = false;
+      return;
+    }
+    // Push _return_value_ptr to _kernelBuffer
+    *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = _return_value_ptr;
+    _bufferOffset += sizeof(_return_value_ptr);
+  }
+  else {
+    // Get the next java argument and its value which should be a T_LONG
+    oop arg = next_arg(T_LONG);
+    // Copy the java argument value to kernelArgBuffer
+    jvalue val;
+    if (java_lang_boxing_object::get_value(arg, &val) != T_LONG) {
+      tty->print_cr("[CUDA] *** Error: Unexpected argument type; expecting T_LONG");
+      _success = false;
+      return;
+    }
+    *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = val.j;
+    _bufferOffset += sizeof(val.j);
+  }
+  return;
+}
+
+void PTXKernelArguments::do_byte()    { 
+  // If the parameter is a return value, 
+  if (is_return_type()) {
+    // Allocate device memory for T_BYTE return value pointer on device. Size in bytes
+    int status = gpu::Ptx::_cuda_cu_memalloc(&_return_value_ptr, T_BYTE_SIZE);
+    if (status != GRAAL_CUDA_SUCCESS) {
+      tty->print_cr("[CUDA] *** Error (%d) Failed to allocate memory for return value pointer on device", status);
+      _success = false;
+      return;
+    }
+    // Push _return_value_ptr to _kernelBuffer
+    *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = _return_value_ptr;
+    _bufferOffset += sizeof(_return_value_ptr);
+  }
+  else {
+    // Get the next java argument and its value which should be a T_BYTE
+    oop arg = next_arg(T_BYTE);
+    // Copy the java argument value to kernelArgBuffer
+    jvalue val;
+    if (java_lang_boxing_object::get_value(arg, &val) != T_BYTE) {
+      tty->print_cr("[CUDA] *** Error: Unexpected argument type; expecting T_BYTE");
+      _success = false;
+      return;
+    }
+    *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = val.b;
+    _bufferOffset += sizeof(val.b);
+  }
+  return;
+}
+
+// TODO implement other do_*
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/src/gpu/ptx/vm/kernelArguments.hpp	Sat Aug 31 09:18:58 2013 +0200
@@ -0,0 +1,128 @@
+/*
+ * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved.
+ * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
+ *
+ * This code is free software; you can redistribute it and/or modify it
+ * under the terms of the GNU General Public License version 2 only, as
+ * published by the Free Software Foundation.
+ *
+ * This code is distributed in the hope that it will be useful, but WITHOUT
+ * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+ * FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+ * version 2 for more details (a copy is included in the LICENSE file that
+ * accompanied this code).
+ *
+ * You should have received a copy of the GNU General Public License version
+ * 2 along with this work; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
+ *
+ * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
+ * or visit www.oracle.com if you need additional information or have any
+ * questions.
+ *
+ */
+
+#ifndef KERNEL_ARGUMENTS_PTX_HPP
+#define KERNEL_ARGUMENTS_PTX_HPP
+
+#include "runtime/gpu.hpp"
+#include "runtime/signature.hpp"
+
+#define T_BYTE_SIZE       1
+#define T_INT_BYTE_SIZE   4
+#define T_LONG_BYTE_SIZE  8
+
+class PTXKernelArguments : public SignatureIterator {
+public:
+  // Buffer holding CUdeviceptr values that represent the kernel arguments
+  char _kernelArgBuffer[1024];
+  // Current offset into _kernelArgBuffer
+  size_t _bufferOffset;
+  gpu::Ptx::CUdeviceptr _return_value_ptr;
+private:
+  // Array of java argument oops
+  arrayOop _args;
+  // Current index into _args
+  int _index;
+  // Flag to indicate successful creation of kernel argument buffer
+  bool _success;
+  // Get next java argument
+  oop next_arg(BasicType expectedType);
+
+ public:
+  PTXKernelArguments(Symbol* signature, arrayOop args, bool is_static) : SignatureIterator(signature) {
+    this->_return_type = T_ILLEGAL;
+    _index = 0;
+    _args = args;
+    _success = true;
+    _bufferOffset = 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.");
+    }
+    // Iterate over the entire signature
+    iterate();
+    assert((_success && (_index == args->length())), "arg count mismatch with signature");
+  }
+
+  inline char* device_argument_buffer() {
+    return _kernelArgBuffer;
+  }
+
+  inline size_t device_argument_buffer_size() {
+    return _bufferOffset;
+  }
+
+  // Get the return oop value
+  oop get_return_oop();
+
+  // get device return value ptr
+  gpu::Ptx::CUdeviceptr get_return_value_ptr() {
+      return _return_value_ptr;
+  }
+
+  
+  void do_byte();
+  void do_int();
+  void do_long();
+
+  inline void do_bool()   {
+    /* TODO : To be implemented */ 
+    guarantee(false, "NYI");
+  }
+  inline void do_char()   {
+    /* TODO : To be implemented */ 
+    guarantee(false, "NYI");
+  }
+  inline void do_short()  {
+    /* TODO : To be implemented */ 
+    guarantee(false, "NYI");
+  }
+  inline void do_float()  {
+    /* TODO : To be implemented */ 
+    guarantee(false, "NYI");
+  }
+  inline void do_double() {
+    /* TODO : To be implemented */ 
+    guarantee(false, "NYI");
+  }
+
+  inline void do_object() {
+    /* TODO : To be implemented */ 
+    guarantee(false, "NYI");
+  }
+  inline void do_object(int begin, int end) {
+    /* TODO : To be implemented */ 
+    guarantee(false, "NYI");
+  }
+  inline void do_array(int begin, int end)  {
+    /* TODO : To be implemented */ 
+    guarantee(false, "NYI");
+  }
+  inline void do_void() {
+    /* TODO : To be implemented */ 
+    guarantee(false, "NYI");
+  }
+};
+
+#endif  // KERNEL_ARGUMENTS_HPP
--- a/src/os_gpu/linux_ptx/vm/gpu_linux.cpp	Sat Aug 31 09:18:42 2013 +0200
+++ b/src/os_gpu/linux_ptx/vm/gpu_linux.cpp	Sat Aug 31 09:18:58 2013 +0200
@@ -39,7 +39,6 @@
  */
 
 static unsigned int nvidia_vendor_id = 0x10de;
-static unsigned int nvidia_gk110_dev_id = 0x1005;
 
 bool gpu::Linux::probe_gpu() {
   /* 
@@ -62,7 +61,7 @@
   while (fgets(contents, sizeof(contents)-1, pci_devices)) {
     sscanf(contents, "%04x%04x%04x", &bus_num_devfn_ign, &vendor, &device);
     /* Break after finding the first CUDA device. */
-    if ((vendor == nvidia_vendor_id) && (device = nvidia_gk110_dev_id)) {
+    if (vendor == nvidia_vendor_id) {
       cuda_device_exists = true;
       if (TraceGPUInteraction) {
         tty->print_cr("Found supported nVidia CUDA device vendor : 0x%04x device 0x%04x", vendor, device);
--- a/src/share/vm/graal/graalCompilerToGPU.cpp	Sat Aug 31 09:18:42 2013 +0200
+++ b/src/share/vm/graal/graalCompilerToGPU.cpp	Sat Aug 31 09:18:58 2013 +0200
@@ -24,13 +24,11 @@
 #include "precompiled.hpp"
 
 #include "graal/graalCompiler.hpp"
-#include "runtime/javaCalls.hpp"
-#include "graal/graalCompilerToVM.hpp"
 #include "graal/graalEnv.hpp"
 #include "graal/graalJavaAccess.hpp"
 #include "runtime/gpu.hpp"
 #include "runtime/javaCalls.hpp"
-
+# include "ptx/vm/kernelArguments.hpp"
 
 // Entry to native method implementation that transitions current thread to '_thread_in_vm'.
 #define C2V_VMENTRY(result_type, name, signature) \
@@ -81,27 +79,28 @@
   nmethod* nm = (nmethod*) (address) nmethodValue;
   methodHandle mh = nm->method();
   Symbol* signature = mh->signature();
-  JavaCallArguments jca(mh->size_of_parameters());
-
-  JavaArgumentUnboxer jap(signature, &jca, (arrayOop) JNIHandles::resolve(args), mh->is_static());
-  JavaValue result(jap.get_ret_type());
-  jca.set_alternative_target(nm);
 
   // start value is the kernel
   jlong startValue = HotSpotInstalledCode::codeStart(hotspotInstalledCode);
 
-  if (!gpu::execute_kernel((address)startValue, &jca)) {
+  PTXKernelArguments ptxka(signature, (arrayOop) JNIHandles::resolve(args), mh->is_static());
+  JavaValue result(ptxka.get_ret_type());
+  if (!gpu::execute_kernel((address)startValue, ptxka, result)) {
     return NULL;
   }
 
-  if (jap.get_ret_type() == T_VOID) {
+  if (ptxka.get_ret_type() == T_VOID) {
     return NULL;
-  } else if (jap.get_ret_type() == T_OBJECT || jap.get_ret_type() == T_ARRAY) {
+  } else if (ptxka.get_ret_type() == T_OBJECT || ptxka.get_ret_type() == T_ARRAY) {
     return JNIHandles::make_local((oop) result.get_jobject());
   } else {
-    oop o = java_lang_boxing_object::create(jap.get_ret_type(), (jvalue *) result.get_value_addr(), CHECK_NULL);
+    oop o = java_lang_boxing_object::create(ptxka.get_ret_type(), (jvalue *) result.get_value_addr(), CHECK_NULL);
+    if (TraceGPUInteraction) {
+      tty->print_cr("GPU execution returned %d", result.get_jint());
+    }
     return JNIHandles::make_local(o);
   }
+
 C2V_END
 
 C2V_VMENTRY(jboolean, deviceInit, (JNIEnv *env, jobject))
--- a/src/share/vm/runtime/gpu.hpp	Sat Aug 31 09:18:42 2013 +0200
+++ b/src/share/vm/runtime/gpu.hpp	Sat Aug 31 09:18:58 2013 +0200
@@ -26,6 +26,9 @@
 #define SHARE_VM_RUNTIME_GPU_HPP
 
 #include "runtime/atomic.hpp"
+#include "oops/symbol.hpp"
+
+class PTXKernelArguments;
 
 // gpu defines the interface to the graphics processor; this includes traditional
 // GPU services such as graphics kernel load and execute.
@@ -43,7 +46,7 @@
   
   static void * generate_kernel(unsigned char *code, int code_len, const char *name);
 
-  static bool execute_kernel(address kernel, JavaCallArguments * jca);
+  static bool execute_kernel(address kernel, PTXKernelArguments & ptxka, JavaValue & ret);
 
   static void set_available(bool value) {
     _available = value;
@@ -80,6 +83,7 @@
 # include "gpu_bsd.hpp"
 #endif
 
+public:
 # include "ptx/vm/gpu_ptx.hpp"
 
 };