# HG changeset patch # User Doug Simon # Date 1377933538 -7200 # Node ID 96e4e5333a257e7e35ffb1084643b38cd6b51416 # Parent 94779c895aadc825035dd85375bd1097a0330c7e# Parent 49bb1bc983c66bc3d7eb5c9ccb27d6d44555a3c1 Merge. diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CodeCacheProvider.java --- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CodeCacheProvider.java 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(); } diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/Constant.java --- 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()); + } + } } diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java --- a/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java 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() + ";" + ""); } diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ArrayTest.java --- 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"); diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/BasicPTXTest.java --- 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; } diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ControlTest.java --- 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"); diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/FloatPTXTest.java --- 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()) { diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/IntegerPTXTest.java --- 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()) { diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/PTXTestBase.java --- 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; } } } diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java --- 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 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 signed32 = new TreeSet<>(); + final SortedSet signed64 = new TreeSet<>(); + + ValueProcedure trackRegisterKind = new ValueProcedure() { + + @Override + public Value doValue(Value value, OperandMode mode, EnumSet 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("}"); diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java --- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java 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(); + } } diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXTargetMethodAssembler.java --- 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); } diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotBackend.java --- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotBackend.java 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); diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotLIRGenerator.java --- 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)); } diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotMove.java --- /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); + } +} diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotRuntime.java --- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotRuntime.java 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 { diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Move.java --- a/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Move.java 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); - } } diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java --- 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; diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMemOp.java --- /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()); + } + } + } +} diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java --- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java 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; diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXParameterOp.java --- /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()); + } + } + } +} diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/ConstantNode.java --- 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 diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/PhiNode.java --- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/PhiNode.java 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 diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/BinaryNode.java --- 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; } diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.api.dsl.test/src/com/oracle/truffle/api/dsl/test/CompilerErrorTest.java --- /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 { + + } + +} diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.api.dsl.test/src/com/oracle/truffle/api/dsl/test/SpecializationGroupingTest.java --- 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; diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.api.dsl.test/src/com/oracle/truffle/api/dsl/test/TestHelper.java --- 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 List> permutations(List list) { + return permutations(new ArrayList(), list, new ArrayList>()); + } + + static Object[][] permutations(Object... list) { + List> permutations = permutations(Arrays.asList(list)); + + Object[][] a = new Object[permutations.size()][]; + int index = 0; + for (List p : permutations) { + a[index] = p.toArray(new Object[p.size()]); + index++; + } + + return a; + } + + static List> permutations(List prefix, List suffix, List> output) { + if (suffix.size() == 1) { + ArrayList newElement = new ArrayList<>(prefix); + newElement.addAll(suffix); + output.add(newElement); + return output; + } + + for (int i = 0; i < suffix.size(); i++) { + List newPrefix = new ArrayList<>(prefix); + newPrefix.add(suffix.get(i)); + List 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 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> permuts = permutations(Arrays.asList(testValues)); + for (List list : permuts) { + TestRootNode root = createRoot(factory); + for (Object object : list) { + assertValue(root, result, object); + } + } + } + + static void assertValue(TestRootNode root, Object result, Object testValues) { + if (testValues instanceof Object[]) { + assertEquals(result, executeWith(root, (Object[]) testValues)); + } else { + assertEquals(result, executeWith(root, testValues)); + } + } + } diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.api.dsl.test/src/com/oracle/truffle/api/dsl/test/TypeSystemErrorsTest.java --- /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 { + + } + +} diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.api.dsl/src/com/oracle/truffle/api/dsl/ExpectError.java --- /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(); + +} diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.api.dsl/src/com/oracle/truffle/api/dsl/ImplicitCast.java --- /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 { + +} diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/AbstractParser.java --- 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()); diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/TruffleProcessor.java --- 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") diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/TruffleTypes.java --- 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 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() { diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/Utils.java --- 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 mirrors, Class annotationClass) { TypeElement expectedAnnotationType = processingEnv.getElementUtils().getTypeElement(annotationClass.getCanonicalName()); + return findAnnotationMirror(mirrors, expectedAnnotationType); + } + + public static AnnotationMirror findAnnotationMirror(List mirrors, TypeElement expectedAnnotationType) { for (AnnotationMirror mirror : mirrors) { DeclaredType annotationType = mirror.getAnnotationType(); TypeElement actualAnnotationType = (TypeElement) annotationType.asElement(); diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/CreateCastParser.java --- 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 childNames = Utils.getAnnotationValueList(String.class, mirror, "value"); CreateCastData cast = new CreateCastData(method, childNames); diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/ExecutableTypeMethodParser.java --- 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 nodeTypeMirrors(NodeData nodeData) { - // executable types not yet available - if (nodeData.getTypeSystem() == null) { - return Collections.emptyList(); - } - List types = new ArrayList<>(nodeData.getTypeSystem().getPrimitiveTypeMirrors()); - types.add(nodeData.getTypeSystem().getVoidType().getPrimitiveType()); + List 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); } diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/GenericParser.java --- 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; } diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/NodeMethodParser.java --- 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))); diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/NodeParser.java --- 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 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 parseChildren(List elements, final List typeHierarchy) { Set 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 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) { diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/ShortCircuitParser.java --- 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)) { diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/SpecializationGroup.java --- 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 typeGuards; private final List guards; + private final NodeData node; private final SpecializationData specialization; private final List 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 children, List assumptionMatches, List typeGuardsMatches, List 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 elseConnectedGuards = previous.getElseConnectableGuards(); + SpecializationGroup previous = this.getPreviousGroup(); + if (previous == null) { + return null; + } + List 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 iterator = guardMatches.iterator(); iterator.hasNext();) { GuardData guardMatch = iterator.next(); - List 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; diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/SpecializationListenerParser.java --- 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); } diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/node/SpecializationMethodParser.java --- 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"); diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/template/MessageContainer.java --- 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()); + public final void emitMessages(ProcessorContext context, TypeElement baseElement, Log log) { + emitMessagesImpl(context, baseElement, log, new HashSet(), null); } - private void emitMessagesImpl(TypeElement baseElement, Log log, Set visitedSinks) { + private void emitMessagesImpl(ProcessorContext context, TypeElement baseElement, Log log, Set visitedSinks, List verifiedMessages) { + List childMessages; + if (verifiedMessages == null) { + childMessages = collectMessagesWithElementChildren(new HashSet(), 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 collectMessagesWithElementChildren(Set visitedSinks, Element e) { + if (visitedSinks.contains(this)) { + return Collections.emptyList(); + } + visitedSinks.add(this); + + List 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 msgs) { + TypeElement expectError = context.getTruffleTypes().getExpectError(); + if (expectError != null) { + Element element = getMessageElement(); + AnnotationMirror mirror = Utils.findAnnotationMirror(element.getAnnotationMirrors(), expectError); + if (mirror != null) { + List 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 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) { diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/template/Template.java --- 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 extensionElements; - public Template(TypeElement templateType, String templateMethodName, AnnotationMirror annotation) { this.templateType = templateType; this.templateMethodName = templateMethodName; @@ -68,14 +65,6 @@ return annotation; } - public List getExtensionElements() { - return extensionElements; - } - - public void setExtensionElements(List extensionMethods) { - this.extensionElements = extensionMethods; - } - @Override public String toString() { return getClass().getSimpleName() + "[" + Utils.getSimpleName(getTemplateType()) + "]"; diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/template/TemplateMethodParser.java --- 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. emptyList())); + E invalidMethod = create(new TemplateMethod(id, template, methodSpecification, method, annotation, returnTypeMirror, Collections. emptyList()), true); String expectedReturnType = returnTypeSpec.toSignatureString(true); String actualReturnType = Utils.getSimpleName(method.getReturnType()); @@ -163,7 +164,7 @@ List parameters = parseParameters(methodSpecification, parameterTypes); if (parameters == null) { if (isEmitErrors()) { - E invalidMethod = create(new TemplateMethod(id, template, methodSpecification, method, annotation, returnTypeMirror, Collections. emptyList())); + E invalidMethod = create(new TemplateMethod(id, template, methodSpecification, method, annotation, returnTypeMirror, Collections. 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) { diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/GuardData.java --- 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(); } } diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/GuardParser.java --- 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 diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/ImplicitCastData.java --- /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); + } + +} diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/ImplicitCastParser.java --- /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 { + + public ImplicitCastParser(ProcessorContext context, TypeSystemData typeSystem) { + super(context, typeSystem); + } + + @Override + public Class getAnnotationType() { + return ImplicitCast.class; + } + + @Override + public MethodSpec createSpecification(ExecutableElement method, AnnotationMirror mirror) { + List 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); + } + +} diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/TypeCastParser.java --- 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 diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/TypeCheckParser.java --- 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"); diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/TypeSystemData.java --- 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 primitiveTypeMirrors = new ArrayList<>(); private List boxedTypeMirrors = new ArrayList<>(); + private List casts; + private List checks; + private TypeMirror genericType; private TypeData voidType; @@ -58,6 +61,14 @@ } } + public void setCasts(List casts) { + this.casts = casts; + } + + public void setChecks(List 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; } diff -r 94779c895aad -r 96e4e5333a25 graal/com.oracle.truffle.dsl.processor/src/com/oracle/truffle/dsl/processor/typesystem/TypeSystemParser.java --- 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 { - public static final List> ANNOTATIONS = Arrays.asList(TypeSystem.class); + public static final List> ANNOTATIONS = Arrays.asList(TypeSystem.class, ExpectError.class); public TypeSystemParser(ProcessorContext c) { super(c); @@ -84,11 +84,17 @@ verifyExclusiveMethodAnnotation(typeSystem, TypeCast.class, TypeCheck.class); List elements = new ArrayList<>(context.getEnvironment().getElementUtils().getAllMembers(templateType)); - + List implicitCasts = new ImplicitCastParser(context, typeSystem).parse(elements); List casts = new TypeCastParser(context, typeSystem).parse(elements); List 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; } diff -r 94779c895aad -r 96e4e5333a25 src/gpu/ptx/vm/gpu_ptx.cpp --- 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, ¶mBufferSz, + 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"); } diff -r 94779c895aad -r 96e4e5333a25 src/gpu/ptx/vm/gpu_ptx.hpp --- 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 diff -r 94779c895aad -r 96e4e5333a25 src/gpu/ptx/vm/kernelArguments.cpp --- /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_* diff -r 94779c895aad -r 96e4e5333a25 src/gpu/ptx/vm/kernelArguments.hpp --- /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 diff -r 94779c895aad -r 96e4e5333a25 src/os_gpu/linux_ptx/vm/gpu_linux.cpp --- 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); diff -r 94779c895aad -r 96e4e5333a25 src/share/vm/graal/graalCompilerToGPU.cpp --- 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)) diff -r 94779c895aad -r 96e4e5333a25 src/share/vm/runtime/gpu.hpp --- 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" };