Mercurial > hg > truffle
changeset 12662:c6cc96cc6a1f
Merge.
| author | Christian Humer <christian.humer@gmail.com> |
|---|---|
| date | Mon, 04 Nov 2013 12:20:17 +0100 |
| parents | 0497d6702cff (current diff) 2c4aa758ee18 (diff) |
| children | 7f507f082daa |
| files | graal/com.oracle.truffle.api/src/com/oracle/truffle/api/frame/FrameSlotImpl.java mxtool/mx.py |
| diffstat | 101 files changed, 2454 insertions(+), 833 deletions(-) [+] |
line wrap: on
line diff
--- a/.hgignore Mon Nov 04 12:18:58 2013 +0100 +++ b/.hgignore Mon Nov 04 12:20:17 2013 +0100 @@ -3,6 +3,7 @@ ^mx/eclipseinit.timestamp ^mx/netbeansinit.timestamp ^mx/eclipse-launches +^mx/currentAnnotationProcessors ^mx/ecj.jar ^mx/includes ^build/ @@ -16,7 +17,6 @@ \.checkstyle$ \.classpath \.factorypath -\.currentAnnotationProcessors \.externalToolBuilders \.project \.settings/
--- a/graal/com.oracle.graal.asm.hsail/src/com/oracle/graal/asm/hsail/HSAILAssembler.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.asm.hsail/src/com/oracle/graal/asm/hsail/HSAILAssembler.java Mon Nov 04 12:20:17 2013 +0100 @@ -225,7 +225,7 @@ public static final String getArgTypeBitwiseLogical(Value src) { String length = getArgType(src); - String prefix = "b" + (length.endsWith("64") ? "64" : "32"); + String prefix = "_b" + (length.endsWith("64") ? "64" : "32"); return prefix; } @@ -268,10 +268,6 @@ emitString(prefix + getArgType(dest) + "_" + getArgType(src) + " " + HSAIL.mapRegister(dest) + ", " + HSAIL.mapRegister(src) + ";"); } - public void emitArg1(String mnemonic, Value dest, Value src) { - emitString(mnemonic + "_" + getArgType(src) + " " + HSAIL.mapRegister(dest) + ", " + mapRegOrConstToString(src) + ";" + ""); - } - public static String mapAddress(HSAILAddress addr) { return "[$d" + addr.getBase().encoding() + " + " + addr.getDisplacement() + "]"; } @@ -304,40 +300,84 @@ } - public final void emit(String mnemonic, Value dest, Value src0, Value src1) { + /** + * Emits an instruction. + * + * @param mnemonic the instruction mnemonic + * @param dest the destination operand + * @param sources the source operands + */ + public final void emit(String mnemonic, Value dest, Value... sources) { String prefix = getArgType(dest); - emit(mnemonic + "_" + prefix, dest, "", src0, src1); - } - - public final void emitForceUnsigned(String mnemonic, Value dest, Value src0, Value src1) { - String prefix = getArgTypeForceUnsigned(dest); - emit(mnemonic + "_" + prefix, dest, "", src0, src1); + emitTextFormattedInstruction(mnemonic + "_" + prefix, dest, sources); } - private void emit(String instr, Value dest, String controlRegString, Value src0, Value src1) { - assert (!isConstant(dest)); - emitString(String.format("%s %s, %s%s, %s;", instr, HSAIL.mapRegister(dest), controlRegString, mapRegOrConstToString(src0), mapRegOrConstToString(src1))); + /** + * Emits an unsigned instruction. + * + * @param mnemonic the instruction mnemonic + * @param dest the destination argument + * @param sources the source arguments + * + */ + public final void emitForceUnsigned(String mnemonic, Value dest, Value... sources) { + String prefix = getArgTypeForceUnsigned(dest); + emitTextFormattedInstruction(mnemonic + "_" + prefix, dest, sources); } - private void emit(String instr, Value dest, Value src0, Value src1, Value src2) { - assert (!isConstant(dest)); - emitString(String.format("%s %s, %s, %s, %s;", instr, HSAIL.mapRegister(dest), mapRegOrConstToString(src0), mapRegOrConstToString(src1), mapRegOrConstToString(src2))); - } - - public final void cmovCommon(Value dest, Value trueReg, Value falseReg, int width) { - String instr = (width == 32 ? "cmov_b32" : "cmov_b64"); - emit(instr, dest, "$c0, ", trueReg, falseReg); + /** + * Emits an instruction for a bitwise logical operation. + * + * @param mnemonic the instruction mnemonic + * @param dest the destination + * @param sources the source operands + */ + public final void emitForceBitwise(String mnemonic, Value dest, Value... sources) { + String prefix = getArgTypeBitwiseLogical(dest); + emitTextFormattedInstruction(mnemonic + prefix, dest, sources); } /** + * Central helper routine that emits a text formatted HSAIL instruction via call to + * AbstractAssembler.emitString. All the emit routines in the assembler end up calling this one. * - * Emit code to generate a 32-bit or 64-bit bitwise logical operation. Used to generate bitwise - * AND, OR and XOR. + * @param instr the full instruction mnenomics including any prefixes + * @param dest the destination operand + * @param sources the source operand */ - public final void emitBitwiseLogical(String mnemonic, Value dest, Value src0, Value src1) { - // Bitwise ops don't use a control register so the fourth arg is empty. - String prefix = getArgTypeBitwiseLogical(dest); - emit(mnemonic + "_" + prefix, dest, "", src0, src1); + private void emitTextFormattedInstruction(String instr, Value dest, Value... sources) { + /** + * Destination can't be a constant and no instruction has > 3 source operands. + */ + assert (!isConstant(dest) && sources.length <= 3); + switch (sources.length) { + case 3: + // Emit an instruction with three source operands. + emitString(String.format("%s %s, %s, %s, %s;", instr, HSAIL.mapRegister(dest), mapRegOrConstToString(sources[0]), mapRegOrConstToString(sources[1]), mapRegOrConstToString(sources[2]))); + break; + case 2: + // Emit an instruction with two source operands. + emitString(String.format("%s %s, %s, %s;", instr, HSAIL.mapRegister(dest), mapRegOrConstToString(sources[0]), mapRegOrConstToString(sources[1]))); + break; + default: + // Emit an instruction with one source operand. + emitString(String.format("%s %s, %s;", instr, HSAIL.mapRegister(dest), mapRegOrConstToString(sources[0]))); + break; + } + } + + /** + * Emits a conditional move instruction. + * + * @param dest the destination operand storing result of the move + * @param trueReg the register that should be copied to dest if the condition is true + * @param falseReg the register that should be copied to dest if the condition is false + * @param width the width of the instruction (32 or 64 bits) + */ + public final void emitConditionalMove(Value dest, Value trueReg, Value falseReg, int width) { + assert (!isConstant(dest)); + String instr = (width == 32 ? "cmov_b32" : "cmov_b64"); + emitString(String.format("%s %s, %s%s, %s;", instr, HSAIL.mapRegister(dest), "$c0, ", mapRegOrConstToString(trueReg), mapRegOrConstToString(falseReg))); } /** @@ -351,12 +391,12 @@ // only use add if result is not starting as null (unsigned compare) emitCompare(result, Constant.forLong(0), "eq", false, true); emit("add", result, result, Constant.forLong(narrowOopBase)); - cmovCommon(result, Constant.forLong(0), result, 64); + emitConditionalMove(result, Constant.forLong(0), result, 64); } else { // only use mad if result is not starting as null (unsigned compare) emitCompare(result, Constant.forLong(0), "eq", false, true); - emit("mad_u64 ", result, result, Constant.forInt(1 << narrowOopShift), Constant.forLong(narrowOopBase)); - cmovCommon(result, Constant.forLong(0), result, 64); + emitTextFormattedInstruction("mad_u64 ", result, result, Constant.forInt(1 << narrowOopShift), Constant.forLong(narrowOopBase)); + emitConditionalMove(result, Constant.forLong(0), result, 64); } } @@ -369,13 +409,18 @@ // only use sub if result is not starting as null (unsigned compare) emitCompare(result, Constant.forLong(0), "eq", false, true); emit("sub", result, result, Constant.forLong(narrowOopBase)); - cmovCommon(result, Constant.forLong(0), result, 64); + emitConditionalMove(result, Constant.forLong(0), result, 64); } if (narrowOopShift != 0) { emit("shr", result, result, Constant.forInt(narrowOopShift)); } } + /** + * Emits a comment. Useful for debugging purposes. + * + * @param comment + */ public void emitComment(String comment) { emitString(comment); }
--- a/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Mon Nov 04 12:20:17 2013 +0100 @@ -105,6 +105,7 @@ protected Value source1; protected Value source2; private boolean logicInstruction = false; + private boolean ldRetAddrInstruction = false; public StandardFormat(Variable dst, Value src1, Value src2) { setDestination(dst); @@ -139,8 +140,18 @@ logicInstruction = b; } + public void setLdRetAddrInstruction(boolean b) { + ldRetAddrInstruction = b; + } + public String typeForKind(Kind k) { - if (logicInstruction) { + if (ldRetAddrInstruction) { + if (System.getProperty("os.arch").compareTo("amd64") == 0) { + return "u64"; + } else { + return "u32"; + } + } else if (logicInstruction) { switch (k.getTypeChar()) { case 's': return "b16"; @@ -658,16 +669,16 @@ } public static class Param extends SingleOperandFormat { - - private boolean lastParameter; + // Last parameter holds the return parameter. + private boolean returnParameter; public Param(Variable d, boolean lastParam) { super(d, null); - setLastParameter(lastParam); + setReturnParameter(lastParam); } - public void setLastParameter(boolean value) { - lastParameter = value; + public void setReturnParameter(boolean value) { + returnParameter = value; } public String emitParameter(Variable v) { @@ -675,30 +686,38 @@ } public void emit(PTXAssembler asm) { - asm.emitString(".param ." + paramForKind(dest.getKind()) + emitParameter(dest) + (lastParameter ? "" : ",")); + asm.emitString(".param ." + paramForKind(dest.getKind()) + emitParameter(dest) + (returnParameter ? "" : ",")); } public String paramForKind(Kind k) { - switch (k.getTypeChar()) { - case 'z': - case 'f': - return "s32"; - case 'b': - return "s8"; - case 's': - return "s16"; - case 'c': - return "u16"; - case 'i': - return "s32"; - case 'j': - return "s64"; - case 'd': - return "f64"; - case 'a': + if (returnParameter) { + if (System.getProperty("os.arch").compareTo("amd64") == 0) { return "u64"; - default: - throw GraalInternalError.shouldNotReachHere(); + } else { + return "u32"; + } + } else { + switch (k.getTypeChar()) { + case 'z': + case 'f': + return "s32"; + case 'b': + return "s8"; + case 's': + return "s16"; + case 'c': + return "u16"; + case 'i': + return "s32"; + case 'j': + return "s64"; + case 'd': + return "f64"; + case 'a': + return "u64"; + default: + throw GraalInternalError.shouldNotReachHere(); + } } }
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/DoubleNegTest.java Mon Nov 04 12:20:17 2013 +0100 @@ -0,0 +1,81 @@ +/* + * Copyright (c) 2009, 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ + +package com.oracle.graal.compiler.hsail.test; + +import org.junit.Test; +import com.oracle.graal.compiler.hsail.test.infra.*; + +/** + * Tests negation of double values. Generates a neg_f64 instruction. + */ +public class DoubleNegTest extends GraalKernelTester { + + static final int num = 20; + // Output array storing the results of negation operations. + @Result protected double[] outArray = new double[num]; + + /** + * The static "kernel" method we will be testing. This method performs a negation operation on + * an element of an input array and writes the result to the corresponding index of an output + * array. By convention the gid is the last parameter. + * + * @param out the output array. + * @param ina the input array. + * @param gid the parameter used to index into the input and output arrays. + */ + public static void run(double[] out, double[] ina, int gid) { + out[gid] = -(ina[gid]); + } + + /** + * Tests the HSAIL code generated for this unit test by comparing the result of executing this + * code with the result of executing a sequential Java version of this unit test. + */ + @Test + public void test() { + super.testGeneratedHsail(); + } + + /** + * Initializes the input and output arrays passed to the run routine. + * + * @param in the input array. + */ + void setupArrays(double[] in) { + for (int i = 0; i < num; i++) { + in[i] = i < num / 2 ? i : -i; + outArray[i] = 0; + } + } + + /** + * Dispatches the HSAIL kernel for this test case. + */ + @Override + public void runTest() { + double[] inArray = new double[num]; + setupArrays(inArray); + dispatchMethodKernel(num, outArray, inArray); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/FloatNegTest.java Mon Nov 04 12:20:17 2013 +0100 @@ -0,0 +1,81 @@ +/* + * Copyright (c) 2009, 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ + +package com.oracle.graal.compiler.hsail.test; + +import org.junit.Test; +import com.oracle.graal.compiler.hsail.test.infra.*; + +/** + * Tests negation of float values. Generates a neg_f32 instruction. + */ +public class FloatNegTest extends GraalKernelTester { + + static final int num = 20; + // Output array storing the results of negation operations. + @Result protected float[] outArray = new float[num]; + + /** + * The static "kernel" method we will be testing. This method performs a negation operation on + * an element of an input array and writes the result to the corresponding index of an output + * array. By convention the gid is the last parameter. + * + * @param out the output array. + * @param ina the input array. + * @param gid the parameter used to index into the input and output arrays. + */ + public static void run(float[] out, float[] ina, int gid) { + out[gid] = -(ina[gid]); + } + + /** + * Tests the HSAIL code generated for this unit test by comparing the result of executing this + * code with the result of executing a sequential Java version of this unit test. + */ + @Test + public void test() { + super.testGeneratedHsail(); + } + + /** + * Initializes the input and output arrays passed to the run routine. + * + * @param in the input array. + */ + void setupArrays(float[] in) { + for (int i = 0; i < num; i++) { + in[i] = i < num / 2 ? i : -i; + outArray[i] = 0; + } + } + + /** + * Dispatches the HSAIL kernel for this test case. + */ + @Override + public void runTest() { + float[] inArray = new float[num]; + setupArrays(inArray); + dispatchMethodKernel(num, outArray, inArray); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/IntBitwiseNotTest.java Mon Nov 04 12:20:17 2013 +0100 @@ -0,0 +1,81 @@ +/* + * Copyright (c) 2009, 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ + +package com.oracle.graal.compiler.hsail.test; + +import org.junit.Test; +import com.oracle.graal.compiler.hsail.test.infra.*; + +/** + * Tests bitwise not of an integer. + */ +public class IntBitwiseNotTest extends GraalKernelTester { + + static final int num = 20; + // Output array storing the results of the operations. + @Result protected int[] outArray = new int[num]; + + /** + * The static "kernel" method we will be testing. This method performs a bitwise not operation + * on an element of an input array and writes the result to the corresponding index of an output + * array. By convention the gid is the last parameter. + * + * @param out the output array + * @param ina the input array + * @param gid the parameter used to index into the input and output arrays + */ + public static void run(int[] out, int[] ina, int gid) { + out[gid] = ~(ina[gid]); + } + + /** + * Tests the HSAIL code generated for this unit test by comparing the result of executing this + * code with the result of executing a sequential Java version of this unit test. + */ + @Test + public void test() { + super.testGeneratedHsail(); + } + + /** + * Initializes the input and output arrays passed to the run routine. + * + * @param in the input array + */ + void setupArrays(int[] in) { + for (int i = 0; i < num; i++) { + in[i] = i < num / 2 ? i : -i; + outArray[i] = 0; + } + } + + /** + * Dispatches the HSAIL kernel for this test case. + */ + @Override + public void runTest() { + int[] inArray = new int[num]; + setupArrays(inArray); + dispatchMethodKernel(num, outArray, inArray); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/IntBitwiseShiftLeftTest.java Mon Nov 04 12:20:17 2013 +0100 @@ -0,0 +1,89 @@ +/* + * Copyright (c) 2009, 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ + +package com.oracle.graal.compiler.hsail.test; + +import org.junit.Test; + +import com.oracle.graal.compiler.hsail.test.infra.*; + +/** + * Tests bitwise shift left of int values. Generates an shl_s32 instruction. + */ +public class IntBitwiseShiftLeftTest extends GraalKernelTester { + + static final int num = 100; + // Output array containing the results of shift operations. + @Result protected int[] outArray = new int[num]; + + /** + * The static "kernel" method we will be testing. This method performs a bitwise shift left + * operation on each element of an input array and writes each result to the corresponding index + * of an output array. By convention the gid is the last parameter. + * + * @param out the output array + * @param ina the input array + * @param shiftAmount an array of values used for the shift magnitude + * @param gid the parameter used to index into the input and output arrays + */ + public static void run(int[] out, int[] ina, int[] shiftAmount, int gid) { + out[gid] = ina[gid] << shiftAmount[gid]; + } + + /** + * Tests the HSAIL code generated for this unit test by comparing the result of executing this + * code with the result of executing a sequential Java version of this unit test. + */ + @Test + public void test() { + super.testGeneratedHsail(); + } + + /** + * Initializes the arrays passed to the run routine. + * + * We do this in such a way that the input arrays contain a mix of negative and positive values. + * As a result, the work items will exercise all the different combinations for the sign of the + * value being shifted and the sign of the shift magnitude. + * + * @param in the input array + */ + void setupArrays(int[] in, int[] shiftAmount) { + for (int i = 0; i < num; i++) { + in[i] = i < num / 2 ? i : -i; + shiftAmount[i] = (i & 1) == 0 ? i : -i; + outArray[i] = 0; + } + } + + /** + * Dispatches the HSAIL kernel for this test case. + */ + @Override + public void runTest() { + int[] inArray = new int[num]; + int[] shiftAmount = new int[num]; + setupArrays(inArray, shiftAmount); + dispatchMethodKernel(num, outArray, inArray, shiftAmount); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/IntBitwiseShiftRightTest.java Mon Nov 04 12:20:17 2013 +0100 @@ -0,0 +1,89 @@ +/* + * Copyright (c) 2009, 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ + +package com.oracle.graal.compiler.hsail.test; + +import org.junit.Test; + +import com.oracle.graal.compiler.hsail.test.infra.*; + +/** + * Tests bitwise shift right of int values. Generates an shr_s32 instruction. + */ +public class IntBitwiseShiftRightTest extends GraalKernelTester { + + static final int num = 100; + // Output array containing the results of shift operations. + @Result protected int[] outArray = new int[num]; + + /** + * The static "kernel" method we will be testing. This method performs a bitwise shift left + * operation on each element of an input array and writes each result to the corresponding index + * of an output array. By convention the gid is the last parameter. + * + * @param out the output array + * @param ina the input array + * @param shiftAmount an array of values used for the shift magnitude + * @param gid the parameter used to index into the input and output arrays + */ + public static void run(int[] out, int[] ina, int[] shiftAmount, int gid) { + out[gid] = ina[gid] >> shiftAmount[gid]; + } + + /** + * Tests the HSAIL code generated for this unit test by comparing the result of executing this + * code with the result of executing a sequential Java version of this unit test. + */ + @Test + public void test() { + super.testGeneratedHsail(); + } + + /** + * Initializes the arrays passed to the run routine. + * + * We do this in such a way that the input arrays contain a mix of negative and positive values. + * As a result, the work items will exercise all the different combinations for the sign of the + * value being shifted and the sign of the shift magnitude. + * + * @param in the input array + */ + void setupArrays(int[] in, int[] shiftAmount) { + for (int i = 0; i < num; i++) { + in[i] = i < num / 2 ? i : -i; + shiftAmount[i] = (i & 1) == 0 ? i : -i; + outArray[i] = 0; + } + } + + /** + * Dispatches the HSAIL kernel for this test case. + */ + @Override + public void runTest() { + int[] inArray = new int[num]; + int[] shiftAmount = new int[num]; + setupArrays(inArray, shiftAmount); + dispatchMethodKernel(num, outArray, inArray, shiftAmount); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/IntBitwiseShiftRightUnsignedTest.java Mon Nov 04 12:20:17 2013 +0100 @@ -0,0 +1,88 @@ +/* + * Copyright (c) 2009, 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ + +package com.oracle.graal.compiler.hsail.test; + +import org.junit.Test; +import com.oracle.graal.compiler.hsail.test.infra.*; + +/** + * Tests unsigned bitwise shift right of int values. Generates an shr_u32 instruction. + */ +public class IntBitwiseShiftRightUnsignedTest extends GraalKernelTester { + + static final int num = 20; + // Output array containing the results of shift operations. + @Result protected int[] outArray = new int[num]; + + /** + * The static "kernel" method we will be testing. This method performs a bitwise unsigned shift + * right operation on an element of an input array and writes the result to the corresponding + * index of an output array. By convention the gid is the last parameter. + * + * @param out the output array + * @param ina the input array + * @param shiftAmount array containing values used for the shift magnitude + * @param gid the parameter used to index into the input and output arrays + */ + public static void run(int[] out, int[] ina, int[] shiftAmount, int gid) { + out[gid] = ina[gid] >>> shiftAmount[gid]; + } + + /** + * Tests the HSAIL code generated for this unit test by comparing the result of executing this + * code with the result of executing a sequential Java version of this unit test. + */ + @Test + public void test() { + super.testGeneratedHsail(); + } + + /** + * Initializes the arrays passed to the run routine. + * + * We do this in such a way that the input arrays contain a mix of negative and positive values. + * As a result, the work groups will exercise all the different combinations for the sign of the + * value being shifted and the sign of the shift magnitude. + * + * @param in the input array. + */ + void setupArrays(int[] in, int[] shiftAmount) { + for (int i = 0; i < num; i++) { + in[i] = i < num / 2 ? i : -i; + shiftAmount[i] = (i & 1) == 0 ? i : -i; + outArray[i] = 0; + } + } + + /** + * Dispatches the HSAIL kernel for this test case. + */ + @Override + public void runTest() { + int[] inArray = new int[num]; + int[] shiftAmount = new int[num]; + setupArrays(inArray, shiftAmount); + dispatchMethodKernel(num, outArray, inArray, shiftAmount); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/IntNegTest.java Mon Nov 04 12:20:17 2013 +0100 @@ -0,0 +1,81 @@ +/* + * Copyright (c) 2009, 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ + +package com.oracle.graal.compiler.hsail.test; + +import org.junit.Test; +import com.oracle.graal.compiler.hsail.test.infra.*; + +/** + * Tests integer negation. Generates a neg_s32 instruction. + */ +public class IntNegTest extends GraalKernelTester { + + static final int num = 20; + // Output array storing the results of negation operations. + @Result protected int[] outArray = new int[num]; + + /** + * The static "kernel" method we will be testing. This method performs a negation operation on + * an element of an input array and writes the result to the corresponding index of an output + * array. By convention the gid is the last parameter. + * + * @param out the output array. + * @param ina the input array. + * @param gid the parameter used to index into the input and output arrays. + */ + public static void run(int[] out, int[] ina, int gid) { + out[gid] = -(ina[gid]); + } + + /** + * Tests the HSAIL code generated for this unit test by comparing the result of executing this + * code with the result of executing a sequential Java version of this unit test. + */ + @Test + public void test() { + super.testGeneratedHsail(); + } + + /** + * Initializes the input and output arrays passed to the run routine. + * + * @param in the input array. + */ + void setupArrays(int[] in) { + for (int i = 0; i < num; i++) { + in[i] = i < num / 2 ? i : -i; + outArray[i] = 0; + } + } + + /** + * Dispatches the HSAIL kernel for this test case. + */ + @Override + public void runTest() { + int[] inArray = new int[num]; + setupArrays(inArray); + dispatchMethodKernel(num, outArray, inArray); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/LongBitwiseNotTest.java Mon Nov 04 12:20:17 2013 +0100 @@ -0,0 +1,81 @@ +/* + * Copyright (c) 2009, 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ + +package com.oracle.graal.compiler.hsail.test; + +import org.junit.Test; +import com.oracle.graal.compiler.hsail.test.infra.*; + +/** + * Tests bitwise not of a long. + */ +public class LongBitwiseNotTest extends GraalKernelTester { + + static final int num = 20; + // Output array storing the results of the operations. + @Result protected long[] outArray = new long[num]; + + /** + * The static "kernel" method we will be testing. This method performs a bitwise not operation + * on an element of an input array and writes the result to the corresponding index of an output + * array. By convention the gid is the last parameter. + * + * @param out the output array + * @param ina the input array + * @param gid the parameter used to index into the input and output arrays + */ + public static void run(long[] out, long[] ina, int gid) { + out[gid] = ~(ina[gid]); + } + + /** + * Tests the HSAIL code generated for this unit test by comparing the result of executing this + * code with the result of executing a sequential Java version of this unit test. + */ + @Test + public void test() { + super.testGeneratedHsail(); + } + + /** + * Initializes the input and output arrays passed to the run routine. + * + * @param in the input array + */ + void setupArrays(long[] in) { + for (int i = 0; i < num; i++) { + in[i] = i < num / 2 ? i : -i; + outArray[i] = 0; + } + } + + /** + * Dispatches the HSAIL kernel for this test case. + */ + @Override + public void runTest() { + long[] inArray = new long[num]; + setupArrays(inArray); + dispatchMethodKernel(num, outArray, inArray); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/LongBitwiseShiftLeftTest.java Mon Nov 04 12:20:17 2013 +0100 @@ -0,0 +1,94 @@ +/* + * Copyright (c) 2009, 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ + +package com.oracle.graal.compiler.hsail.test; + +import org.junit.Test; +import com.oracle.graal.compiler.hsail.test.infra.*; + +/** + * Tests bitwise left shift of long values. Generates an shl_s64 instruction. + */ +public class LongBitwiseShiftLeftTest extends GraalKernelTester { + + static final int num = 100; + // Output array containing the results of shift operations. + @Result protected long[] outArray = new long[num]; + + /** + * The static "kernel" method we will be testing. This method performs a bitwise shift left + * operation on an element of an input array and writes the result to the corresponding index of + * an output array. By convention the gid is the last parameter. + * + * @param out the output array + * @param ina the input array + * @param shiftAmount an array of values used for the shift magnitude + * @param gid the parameter used to index into the input and output arrays + */ + public static void run(long[] out, long[] ina, int[] shiftAmount, int gid) { + out[gid] = ina[gid] << shiftAmount[gid]; + } + + /** + * Tests the HSAIL code generated for this unit test by comparing the result of executing this + * code with the result of executing a sequential Java version of this unit test. + */ + @Test + public void test() { + super.testGeneratedHsail(); + } + + /** + * Initializes the arrays passed to the run routine. + * + * We do this in such a way that the input arrays contain a mix of negative and positive values. + * As a result, the work items will exercise all the different combinations for the sign of the + * value being shifted and the sign of the shift magnitude. + * + * @param in the input array. + */ + void setupArrays(long[] in, int[] shiftAmount) { + for (int i = 0; i < num; i++) { + /** + * Fill lower half of in[] with positive numbers and upper half with negative numbers. + */ + in[i] = i < num / 2 ? i : -i; + /** + * Fill shiftAmount[] so that even elements are positive and odd elements are negative. + */ + shiftAmount[i] = (i & 1) == 0 ? i : -i; + outArray[i] = 0; + } + } + + /** + * Dispatches the HSAIL kernel for this test case. + */ + @Override + public void runTest() { + long[] inArray = new long[num]; + int[] shiftAmount = new int[num]; + setupArrays(inArray, shiftAmount); + dispatchMethodKernel(num, outArray, inArray, shiftAmount); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/LongBitwiseShiftRightTest.java Mon Nov 04 12:20:17 2013 +0100 @@ -0,0 +1,94 @@ +/* + * Copyright (c) 2009, 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ + +package com.oracle.graal.compiler.hsail.test; + +import org.junit.Test; +import com.oracle.graal.compiler.hsail.test.infra.*; + +/** + * Tests bitwise right shift of long values. Generates an shr_s64 instruction. + */ +public class LongBitwiseShiftRightTest extends GraalKernelTester { + + static final int num = 100; + // Output array containing the results of shift operations. + @Result protected long[] outArray = new long[num]; + + /** + * The static "kernel" method we will be testing. This method performs a bitwise shift righj + * operation on an element of an input array and writes the result to the corresponding index of + * an output array. By convention the gid is the last parameter. + * + * @param out the output array + * @param ina the input array + * @param shiftAmount an array of values used for the shift magnitude + * @param gid the parameter used to index into the input and output arrays + */ + public static void run(long[] out, long[] ina, int[] shiftAmount, int gid) { + out[gid] = ina[gid] >> shiftAmount[gid]; + } + + /** + * Tests the HSAIL code generated for this unit test by comparing the result of executing this + * code with the result of executing a sequential Java version of this unit test. + */ + @Test + public void test() { + super.testGeneratedHsail(); + } + + /** + * Initializes the arrays passed to the run routine. + * + * We do this in such a way that the input arrays contain a mix of negative and positive values. + * As a result, the work items will exercise all the different combinations for the sign of the + * value being shifted and the sign of the shift magnitude. + * + * @param in the input array + */ + void setupArrays(long[] in, int[] shiftAmount) { + for (int i = 0; i < num; i++) { + /** + * Fill lower half of in[] with positive numbers and upper half with negative numbers. + */ + in[i] = i < num / 2 ? i : -i; + /** + * Fill shiftAmount[] so that even elements are positive and odd elements are negative. + */ + shiftAmount[i] = (i & 1) == 0 ? i : -i; + outArray[i] = 0; + } + } + + /** + * Dispatches the HSAIL kernel for this test case. + */ + @Override + public void runTest() { + long[] inArray = new long[num]; + int[] shiftAmount = new int[num]; + setupArrays(inArray, shiftAmount); + dispatchMethodKernel(num, outArray, inArray, shiftAmount); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/LongBitwiseShiftRightUnsignedTest.java Mon Nov 04 12:20:17 2013 +0100 @@ -0,0 +1,88 @@ +/* + * Copyright (c) 2009, 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.graal.compiler.hsail.test; + +import org.junit.Test; +import com.oracle.graal.compiler.hsail.test.infra.*; + +/** + * Tests unsigned bitwise right shift of long values. Generates an shr_u64 instruction. + */ +public class LongBitwiseShiftRightUnsignedTest extends GraalKernelTester { + + static final int num = 20; + // Output array containing the results of shift operations. + @Result protected long[] outArray = new long[num]; + + /** + * The static "kernel" method we will be testing. This method performs an unsigned bitwise shift + * right operation on an element of an input array and writes the result to the corresponding + * index of an output array. By convention the gid is the last parameter. + * + * @param out the output array. + * @param ina the input array. + * @param shiftAmount array containing values used for the shift magnitude + * @param gid the parameter used to index into the input and output arrays + */ + public static void run(long[] out, long[] ina, int[] shiftAmount, int gid) { + out[gid] = ina[gid] >>> shiftAmount[gid]; + } + + /** + * Tests the HSAIL code generated for this unit test by comparing the result of executing this + * code with the result of executing a sequential Java version of this unit test. + */ + @Test + public void test() { + super.testGeneratedHsail(); + } + + /** + * Initializes the arrays passed to the run routine. + * + * We do this in such a way that the input arrays contain a mix of negative and positive values. + * As a result, the work groups will exercise all the different combinations for the sign of the + * value being shifted and the sign of the shift magnitude. + * + * @param in the input array. + */ + void setupArrays(long[] in, int[] shiftAmount) { + for (int i = 0; i < num; i++) { + in[i] = i < num / 2 ? i : -i; + shiftAmount[i] = (i & 1) == 0 ? i : -i; + outArray[i] = 0; + } + } + + /** + * Dispatches the HSAIL kernel for this test case. + */ + @Override + public void runTest() { + long[] inArray = new long[num]; + int[] shiftAmount = new int[num]; + setupArrays(inArray, shiftAmount); + dispatchMethodKernel(num, outArray, inArray, shiftAmount); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/LongNegTest.java Mon Nov 04 12:20:17 2013 +0100 @@ -0,0 +1,81 @@ +/* + * Copyright (c) 2009, 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ + +package com.oracle.graal.compiler.hsail.test; + +import org.junit.Test; +import com.oracle.graal.compiler.hsail.test.infra.*; + +/** + * Tests negation of long values. Generates a neg_s64 instruction. + */ +public class LongNegTest extends GraalKernelTester { + + static final int num = 20; + // Output array storing the results of negation operations. + @Result protected long[] outArray = new long[num]; + + /** + * The static "kernel" method we will be testing. This method performs a negation operation on + * an element of an input array and writes the result to the corresponding index of an output + * array. By convention the gid is the last parameter. + * + * @param out the output array. + * @param ina the input array. + * @param gid the parameter used to index into the input and output arrays. + */ + public static void run(long[] out, long[] ina, int gid) { + out[gid] = -(ina[gid]); + } + + /** + * Tests the HSAIL code generated for this unit test by comparing the result of executing this + * code with the result of executing a sequential Java version of this unit test. + */ + @Test + public void test() { + super.testGeneratedHsail(); + } + + /** + * Initializes the input and output arrays passed to the run routine. + * + * @param in the input array. + */ + void setupArrays(long[] in) { + for (int i = 0; i < num; i++) { + in[i] = i < num / 2 ? i : -i; + outArray[i] = 0; + } + } + + /** + * Dispatches the HSAIL kernel for this test case. + */ + @Override + public void runTest() { + long[] inArray = new long[num]; + setupArrays(inArray); + dispatchMethodKernel(num, outArray, inArray); + } +}
--- a/graal/com.oracle.graal.compiler.hsail/src/com/oracle/graal/compiler/hsail/HSAILLIRGenerator.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.compiler.hsail/src/com/oracle/graal/compiler/hsail/HSAILLIRGenerator.java Mon Nov 04 12:20:17 2013 +0100 @@ -246,13 +246,30 @@ throw GraalInternalError.unimplemented(); } + /** + * Generates the LIR instruction for a negation operation. + * + * @param input the value that is being negated + * @return Variable that represents the result of the negation + */ @Override public Variable emitNegate(Value input) { Variable result = newVariable(input.getKind()); switch (input.getKind()) { case Int: + // Note: The Int case also handles the negation of shorts, bytes, and chars because + // Java treats these types as ints at the bytecode level. append(new Op1Stack(INEG, result, input)); break; + case Long: + append(new Op1Stack(LNEG, result, input)); + break; + case Double: + append(new Op1Stack(DNEG, result, input)); + break; + case Float: + append(new Op1Stack(FNEG, result, input)); + break; default: throw GraalInternalError.shouldNotReachHere(); } @@ -260,13 +277,24 @@ } + /** + * Generates the LIR instruction for a bitwise NOT operation. + * + * @param input the source operand + * @return Variable that represents the result of the operation + */ @Override public Variable emitNot(Value input) { Variable result = newVariable(input.getKind()); switch (input.getKind()) { case Int: + // Note: The Int case also covers other primitive integral types smaller than an int + // (char, byte, short) because Java treats these types as ints. append(new Op1Stack(INOT, result, input)); break; + case Long: + append(new Op1Stack(LNOT, result, input)); + break; default: throw GraalInternalError.shouldNotReachHere(); } @@ -491,11 +519,20 @@ return result; } + /** + * Generates the LIR instruction for a shift left operation. + * + * @param a The value that is being shifted + * @param b The shift amount + * @return Variable that represents the result of the operation + */ @Override public Variable emitShl(Value a, Value b) { Variable result = newVariable(a.getKind()); switch (a.getKind()) { case Int: + // Note: The Int case also covers the shifting of bytes, shorts and chars because + // Java treats these types as ints at the bytecode level. append(new ShiftOp(ISHL, result, a, b)); break; case Long: @@ -507,11 +544,38 @@ return result; } + /** + * Generates the LIR instruction for a shift right operation. + * + * @param a The value that is being shifted + * @param b The shift amount + * @return Variable that represents the result of the operation + */ @Override public Variable emitShr(Value a, Value b) { - throw GraalInternalError.unimplemented(); + Variable result = newVariable(a.getKind()); + switch (a.getKind()) { + case Int: + // Note: The Int case also covers the shifting of bytes, shorts and chars because + // Java treats these types as ints at the bytecode level. + append(new ShiftOp(ISHR, result, a, b)); + break; + case Long: + append(new ShiftOp(LSHR, result, a, b)); + break; + default: + GraalInternalError.shouldNotReachHere(); + } + return result; } + /** + * Generates the LIR instruction for an unsigned shift right operation. + * + * @param a The value that is being shifted + * @param b The shift amount + * @return Variable that represents the result of the operation + */ @Override public Variable emitUShr(Value a, Value b) { Variable result = newVariable(a.getKind()); @@ -575,7 +639,6 @@ case F2L: append(new Op1Stack(F2L, result, input)); break; - default: throw GraalInternalError.shouldNotReachHere(); }
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ArrayPTXTest.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ArrayPTXTest.java Mon Nov 04 12:20:17 2013 +0100 @@ -40,7 +40,11 @@ int[] array3 = {1, 2, 3, 4, 5, 6, 7, 8, 9}; invoke(compile("testStoreArray1I"), array1, 2); - printReport("testStoreArray1I: " + Arrays.toString(array1)); + if (array1[2] == 42) { + printReport("testStoreArray1I: " + Arrays.toString(array1) + " PASSED"); + } else { + printReport("testStoreArray1I: " + Arrays.toString(array1) + " FAILED"); + } invoke(compile("testStoreArrayWarp0"), array2, 2); printReport("testStoreArrayWarp0: " + Arrays.toString(array2));
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ConditionalEliminationTest.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ConditionalEliminationTest.java Mon Nov 04 12:20:17 2013 +0100 @@ -144,8 +144,8 @@ StructuredGraph graph = parse("testNullnessSnippet"); new ConditionalEliminationPhase(getMetaAccess()).apply(graph); new CanonicalizerPhase(true).apply(graph, new PhaseContext(getProviders(), null)); - for (ConstantNode constant : graph.getNodes().filter(ConstantNode.class)) { - if (ConstantNodeRecordsUsages || !constant.gatherUsages().isEmpty()) { + for (ConstantNode constant : getConstantNodes(graph)) { + if (ConstantNodeRecordsUsages || !constant.gatherUsages(graph).isEmpty()) { assertTrue("unexpected constant: " + constant, constant.asConstant().isNull() || constant.asConstant().asInt() > 0); } } @@ -178,8 +178,8 @@ new CanonicalizerPhase(true).apply(graph, new PhaseContext(getProviders(), null)); new ConditionalEliminationPhase(getMetaAccess()).apply(graph); new CanonicalizerPhase(true).apply(graph, new PhaseContext(getProviders(), null)); - for (ConstantNode constant : graph.getNodes().filter(ConstantNode.class)) { - if (ConstantNodeRecordsUsages || !constant.gatherUsages().isEmpty()) { + for (ConstantNode constant : getConstantNodes(graph)) { + if (ConstantNodeRecordsUsages || !constant.gatherUsages(graph).isEmpty()) { assertTrue("unexpected constant: " + constant, constant.asConstant().isNull() || constant.asConstant().asInt() > 0); } }
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/GraalCompilerTest.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/GraalCompilerTest.java Mon Nov 04 12:20:17 2013 +0100 @@ -119,9 +119,9 @@ protected int countUnusedConstants(StructuredGraph graph) { int total = 0; - for (ConstantNode node : graph.getNodes().filter(ConstantNode.class)) { + for (ConstantNode node : getConstantNodes(graph)) { if (!ConstantNodeRecordsUsages) { - if (node.gatherUsages().isEmpty()) { + if (node.gatherUsages(graph).isEmpty()) { total++; } } else { @@ -520,7 +520,6 @@ if (cached.isValid()) { return cached; } - } }
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/deopt/CompiledMethodTest.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/deopt/CompiledMethodTest.java Mon Nov 04 12:20:17 2013 +0100 @@ -29,7 +29,6 @@ import com.oracle.graal.api.code.*; import com.oracle.graal.api.meta.*; import com.oracle.graal.compiler.test.*; -import com.oracle.graal.graph.*; import com.oracle.graal.nodes.*; import com.oracle.graal.phases.common.*; import com.oracle.graal.phases.tiers.*; @@ -59,12 +58,9 @@ new CanonicalizerPhase(true).apply(graph, new PhaseContext(getProviders(), new Assumptions(false))); new DeadCodeEliminationPhase().apply(graph); - for (Node node : graph.getNodes()) { - if (node instanceof ConstantNode) { - ConstantNode constant = (ConstantNode) node; - if (constant.kind() == Kind.Object && " ".equals(constant.value.asObject())) { - constant.replace(ConstantNode.forObject("-", getMetaAccess(), graph)); - } + for (ConstantNode node : ConstantNode.getConstantNodes(graph)) { + if (node.kind() == Kind.Object && " ".equals(node.getValue().asObject())) { + node.replace(graph, ConstantNode.forObject("-", getMetaAccess(), graph)); } }
--- a/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/gen/DebugInfoBuilder.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/gen/DebugInfoBuilder.java Mon Nov 04 12:20:17 2013 +0100 @@ -187,7 +187,7 @@ } } else if (value instanceof ConstantNode) { Debug.metric("StateConstants").increment(); - return ((ConstantNode) value).value; + return ((ConstantNode) value).getValue(); } else if (value != null) { Debug.metric("StateVariables").increment();
--- a/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/gen/LIRGenerator.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/gen/LIRGenerator.java Mon Nov 04 12:20:17 2013 +0100 @@ -159,7 +159,7 @@ if (nodeOperands == null) { return null; } - Value operand = nodeOperands.get(node); + Value operand = !node.isExternal() ? nodeOperands.get(node) : null; if (operand == null) { return getConstantOperand(node); } @@ -171,7 +171,7 @@ Constant value = node.asConstant(); if (value != null) { if (canInlineConstant(value)) { - return setResult(node, value); + return !node.isExternal() ? setResult(node, value) : value; } else { Variable loadedValue; if (constantsLoadedInCurrentBlock == null) {
--- a/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/Graph.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/Graph.java Mon Nov 04 12:20:17 2013 +0100 @@ -235,7 +235,7 @@ public <T extends Node> T addOrUnique(T node) { if (node.getNodeClass().valueNumberable()) { - return uniqueHelper(node); + return uniqueHelper(node, true); } return add(node); } @@ -301,27 +301,38 @@ } /** - * Adds a new node to the graph, if a <i>similar</i> node already exists in the graph, the - * provided node will not be added to the graph but the <i>similar</i> node will be returned - * instead. + * Looks for a node <i>similar</i> to {@code node} and returns it if found. Otherwise + * {@code node} is added to this graph and returned. * - * @param node - * @return the node which was added to the graph or a <i>similar</i> which was already in the - * graph. + * @return a node similar to {@code node} if one exists, otherwise {@code node} */ public <T extends Node & ValueNumberable> T unique(T node) { - assert checkValueNumberable(node); - return uniqueHelper(node); + return uniqueHelper(node, true); + } + + /** + * Looks for a node <i>similar</i> to {@code node}. If not found, {@code node} is added to the + * cache in this graph used to canonicalize nodes. + * <p> + * Note that node must implement {@link ValueNumberable} and must be an + * {@linkplain Node#isExternal() external} node. + * + * @return a node similar to {@code node} if one exists, otherwise {@code node} + */ + public <T extends Node> T uniqueWithoutAdd(T node) { + assert node.isExternal() : node; + assert node instanceof ValueNumberable : node; + return uniqueHelper(node, false); } @SuppressWarnings("unchecked") - <T extends Node> T uniqueHelper(T node) { + <T extends Node> T uniqueHelper(T node, boolean addIfMissing) { assert node.getNodeClass().valueNumberable(); Node other = this.findDuplicate(node); if (other != null) { return (T) other; } else { - Node result = addHelper(node); + Node result = addIfMissing ? addHelper(node) : node; if (node.getNodeClass().isLeafNode()) { putNodeIntoCache(result); } @@ -330,7 +341,7 @@ } void putNodeIntoCache(Node node) { - assert node.graph() == this || node.graph() == null; + assert node.isExternal() || node.graph() == this || node.graph() == null; assert node.getNodeClass().valueNumberable(); assert node.getNodeClass().isLeafNode() : node.getClass(); cachedNodes.put(new CacheEntry(node), node); @@ -383,13 +394,6 @@ } } - private static boolean checkValueNumberable(Node node) { - if (!node.getNodeClass().valueNumberable()) { - throw new VerificationError("node is not valueNumberable").addContext(node); - } - return true; - } - public boolean isNew(int mark, Node node) { return node.id >= mark; } @@ -650,6 +654,7 @@ } void register(Node node) { + assert !node.isExternal(); assert node.id() == Node.INITIAL_ID; int id = nodes.size(); nodes.add(id, node);
--- a/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/Node.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/Node.java Mon Nov 04 12:20:17 2013 +0100 @@ -156,7 +156,12 @@ return id; } + /** + * Gets the graph context of this node. This must not be called for {@linkplain #isExternal() + * external} nodes. + */ public Graph graph() { + assert !isExternal() : "external node has no graph: " + this; return graph; } @@ -283,6 +288,15 @@ } /** + * Determines if this node has a {@linkplain #graph() graph} context or is external to any + * graph. The {@link #graph()} method must only be called on nodes for which this method returns + * true. + */ + public boolean isExternal() { + return false; + } + + /** * Finds the index of the last non-null entry in a node array. The search assumes that all * non-null entries precede the first null entry in the array. * @@ -517,7 +531,7 @@ assert assertFalse(other == this, "cannot replace a node with itself"); assert assertFalse(isDeleted(), "cannot replace deleted node"); assert assertTrue(other == null || !other.isDeleted(), "cannot replace with deleted node %s", other); - assert assertTrue(other == null || other.graph() == graph, "cannot replace with node in different graph: %s", other == null ? null : other.graph()); + assert assertTrue(other == null || other.isExternal() || other.graph() == graph, "cannot replace with node in different graph: %s", other == null || other.isExternal() ? null : other.graph()); return true; } @@ -574,7 +588,7 @@ assert assertFalse(isDeleted(), "cannot clear inputs of deleted node"); for (Node input : inputs()) { - if (input.recordsUsages()) { + if (input.recordsUsages() && !input.isExternal()) { removeThisFromUsages(input); if (input.usages().isEmpty()) { NodeChangedListener listener = graph.usagesDroppedToZeroListener; @@ -623,6 +637,7 @@ } public final Node copyWithInputs() { + assert !isExternal(); Node newNode = clone(graph); NodeClass clazz = getNodeClass(); clazz.copyInputs(this, newNode); @@ -661,6 +676,7 @@ } final Node clone(Graph into, boolean clearInputsAndSuccessors) { + assert !isExternal(); NodeClass nodeClass = getNodeClass(); if (nodeClass.valueNumberable() && nodeClass.isLeafNode()) { Node otherNode = into.findNodeInCache(this); @@ -695,6 +711,23 @@ return newNode; } + /** + * Determines if a given node is {@linkplain Graph#uniqueWithoutAdd(Node) unique} within a given + * graph if the node is non-null and {@linkplain #isExternal() external}. + * + * @param node node to check + * @param graph graph context to use + * @return true if node is null, not external or unique within {@code graph} otherwise raises a + * {@link VerificationError} + */ + public static boolean verifyUniqueIfExternal(Node node, Graph graph) { + if (node != null && node.isExternal()) { + Node cached = graph.findNodeInCache(node); + node.assertTrue(cached == node, "external node does not match canonical node %s", cached); + } + return true; + } + protected void afterClone(@SuppressWarnings("unused") Node other) { } @@ -703,7 +736,8 @@ assertTrue(graph() != null, "null graph"); for (Node input : inputs()) { assertTrue(!input.recordsUsages() || input.usages().contains(this), "missing usage in input %s", input); - assertTrue(input.graph() == graph(), "mismatching graph in input %s", input); + assert verifyUniqueIfExternal(input, graph()); + assertTrue(input.isExternal() || input.graph() == graph(), "mismatching graph in input %s", input); } for (Node successor : successors()) { assertTrue(successor.predecessor() == this, "missing predecessor in %s (actual: %s)", successor, successor.predecessor()); @@ -757,17 +791,18 @@ } /** - * hashCode and equals should always rely on object identity alone, thus hashCode and equals are - * final. + * Nodes always use an {@linkplain System#identityHashCode(Object) identity} hash code. For this + * reason, {@linkplain #isExternal() external} nodes should still be {@link Graph#unique unique} + * within the context of a graph. */ @Override public final int hashCode() { - return super.hashCode(); + return System.identityHashCode(this); } /** - * hashCode and equals should always rely on object identity alone, thus hashCode and equals are - * final. + * Equality tests must rely solely on identity. For this reason, {@linkplain #isExternal() + * external} nodes should still be {@link Graph#unique unique} within the context of a graph. */ @Override public final boolean equals(Object obj) {
--- a/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeClass.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeClass.java Mon Nov 04 12:20:17 2013 +0100 @@ -30,9 +30,7 @@ import com.oracle.graal.debug.*; import com.oracle.graal.graph.Graph.DuplicationReplacement; -import com.oracle.graal.graph.Node.Input; -import com.oracle.graal.graph.Node.Successor; -import com.oracle.graal.graph.Node.Verbosity; +import com.oracle.graal.graph.Node.*; import com.oracle.graal.graph.spi.*; /** @@ -835,6 +833,7 @@ if (input != null) { Node newInput = duplicationReplacement.replacement(input, true); node.updateUsages(null, newInput); + assert Node.verifyUniqueIfExternal(newInput, node.graph()); putNode(node, inputOffsets[index], newInput); } index++; @@ -887,6 +886,7 @@ Node oldNode = list.get(i); if (oldNode != null) { Node newNode = duplicationReplacement.replacement(oldNode, true); + assert Node.verifyUniqueIfExternal(newNode, node.graph()); result.set(i, newNode); } } @@ -971,6 +971,7 @@ } public boolean replaceFirstInput(Node node, Node old, Node other) { + assert Node.verifyUniqueIfExternal(other, node.graph()); int index = 0; while (index < directInputCount) { Node input = getNode(node, inputOffsets[index]); @@ -1268,6 +1269,9 @@ InplaceUpdateClosure replacementClosure = new InplaceUpdateClosure() { public Node replacement(Node node, boolean isInput) { + if (node.isExternal() && node instanceof ValueNumberable) { + return graph.uniqueWithoutAdd(node); + } Node target = newNodes.get(node); if (target == null) { Node replacement = node;
--- a/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeMap.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeMap.java Mon Nov 04 12:20:17 2013 +0100 @@ -160,4 +160,27 @@ } }; } + + @Override + public String toString() { + Iterator<Entry<Node, T>> i = entries().iterator(); + if (!i.hasNext()) { + return "{}"; + } + + StringBuilder sb = new StringBuilder(); + sb.append('{'); + while (true) { + Entry<Node, T> e = i.next(); + Node key = e.getKey(); + T value = e.getValue(); + sb.append(key); + sb.append('='); + sb.append(value); + if (!i.hasNext()) { + return sb.append('}').toString(); + } + sb.append(',').append(' '); + } + } }
--- a/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeWorkList.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeWorkList.java Mon Nov 04 12:20:17 2013 +0100 @@ -60,7 +60,7 @@ public void addAll(Iterable<? extends Node> nodes) { for (Node node : nodes) { - if (node.isAlive()) { + if (node.isAlive() && !node.isExternal()) { this.add(node); } }
--- a/graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotBackend.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotBackend.java Mon Nov 04 12:20:17 2013 +0100 @@ -44,6 +44,7 @@ import com.oracle.graal.lir.ptx.*; import com.oracle.graal.nodes.*; import com.oracle.graal.nodes.cfg.*; +import com.oracle.graal.lir.ptx.PTXMemOp.LoadReturnAddrOp; /** * HotSpot PTX specific backend. @@ -64,6 +65,84 @@ return new PTXFrameMap(getCodeCache()); } + static final class RegisterAnalysis extends ValueProcedure { + private final SortedSet<Integer> unsigned64 = new TreeSet<>(); + private final SortedSet<Integer> signed64 = new TreeSet<>(); + private final SortedSet<Integer> float32 = new TreeSet<>(); + private final SortedSet<Integer> signed32 = new TreeSet<>(); + private final SortedSet<Integer> float64 = new TreeSet<>(); + + LIRInstruction op; + + void emitDeclarations(Buffer codeBuffer) { + for (Integer i : signed32) { + codeBuffer.emitString(".reg .s32 %r" + i.intValue() + ";"); + } + for (Integer i : signed64) { + codeBuffer.emitString(".reg .s64 %r" + i.intValue() + ";"); + } + for (Integer i : unsigned64) { + codeBuffer.emitString(".reg .u64 %r" + i.intValue() + ";"); + } + for (Integer i : float32) { + codeBuffer.emitString(".reg .f32 %r" + i.intValue() + ";"); + } + for (Integer i : float64) { + codeBuffer.emitString(".reg .f64 %r" + i.intValue() + ";"); + } + } + + @Override + public Value doValue(Value value, OperandMode mode, EnumSet<OperandFlag> flags) { + if (isVariable(value)) { + Variable regVal = (Variable) value; + Kind regKind = regVal.getKind(); + if ((op instanceof LoadReturnAddrOp) && (mode == OperandMode.DEF)) { + unsigned64.add(regVal.index); + } else { + switch (regKind) { + case Int: + // If the register was used as a wider signed type + // do not add it here + if (!signed64.contains(regVal.index)) { + signed32.add(regVal.index); + } + break; + case Long: + // If the register was used as a narrower signed type + // remove it from there and add it to wider type. + if (signed32.contains(regVal.index)) { + signed32.remove(regVal.index); + } + signed64.add(regVal.index); + break; + case Float: + // If the register was used as a wider signed type + // do not add it here + if (!float64.contains(regVal.index)) { + float32.add(regVal.index); + } + break; + case Double: + // If the register was used as a narrower signed type + // remove it from there and add it to wider type. + if (float32.contains(regVal.index)) { + float32.remove(regVal.index); + } + float64.add(regVal.index); + break; + case Object: + unsigned64.add(regVal.index); + break; + default: + throw GraalInternalError.shouldNotReachHere("unhandled register type " + value.toString()); + } + } + } + return value; + } + } + class PTXFrameContext implements FrameContext { @Override @@ -147,94 +226,27 @@ assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method"; Buffer codeBuffer = tasm.asm.codeBuffer; - - final SortedSet<Integer> signed32 = new TreeSet<>(); - final SortedSet<Integer> signed64 = new TreeSet<>(); - final SortedSet<Integer> unsigned64 = new TreeSet<>(); - final SortedSet<Integer> float32 = new TreeSet<>(); - final SortedSet<Integer> float64 = new TreeSet<>(); - - ValueProcedure trackRegisterKind = new ValueProcedure() { - - @Override - public Value doValue(Value value, OperandMode mode, EnumSet<OperandFlag> flags) { - if (isVariable(value)) { - Variable regVal = (Variable) value; - Kind regKind = regVal.getKind(); - switch (regKind) { - case Int: - // If the register was used as a wider signed type - // do not add it here - if (!signed64.contains(regVal.index)) { - signed32.add(regVal.index); - } - break; - case Long: - // If the register was used as a narrower signed type - // remove it from there and add it to wider type. - if (signed32.contains(regVal.index)) { - signed32.remove(regVal.index); - } - signed64.add(regVal.index); - break; - case Float: - // If the register was used as a wider signed type - // do not add it here - if (!float64.contains(regVal.index)) { - float32.add(regVal.index); - } - break; - case Double: - // If the register was used as a narrower signed type - // remove it from there and add it to wider type. - if (float32.contains(regVal.index)) { - float32.remove(regVal.index); - } - float64.add(regVal.index); - break; - case Object: - unsigned64.add(regVal.index); - break; - default: - throw GraalInternalError.shouldNotReachHere("unhandled register type " + value.toString()); - } - } - return value; - } - }; + RegisterAnalysis registerAnalysis = new RegisterAnalysis(); for (Block b : lirGen.lir.codeEmittingOrder()) { for (LIRInstruction op : lirGen.lir.lir(b)) { if (op instanceof LabelOp) { // Don't consider this as a definition } else { - op.forEachTemp(trackRegisterKind); - op.forEachOutput(trackRegisterKind); + registerAnalysis.op = op; + op.forEachTemp(registerAnalysis); + op.forEachOutput(registerAnalysis); } } } - for (Integer i : signed32) { - codeBuffer.emitString(".reg .s32 %r" + i.intValue() + ";"); - } - for (Integer i : signed64) { - codeBuffer.emitString(".reg .s64 %r" + i.intValue() + ";"); - } - for (Integer i : unsigned64) { - codeBuffer.emitString(".reg .u64 %r" + i.intValue() + ";"); - } - for (Integer i : float32) { - codeBuffer.emitString(".reg .f32 %r" + i.intValue() + ";"); - } - for (Integer i : float64) { - codeBuffer.emitString(".reg .f64 %r" + i.intValue() + ";"); - } + registerAnalysis.emitDeclarations(codeBuffer); + // emit predicate register declaration int maxPredRegNum = ((PTXLIRGenerator) lirGen).getNextPredRegNumber(); if (maxPredRegNum > 0) { codeBuffer.emitString(".reg .pred %p<" + maxPredRegNum + ">;"); } - codeBuffer.emitString(".reg .pred %r;"); // used for setp bool } @Override
--- a/graal/com.oracle.graal.hotspot.test/src/com/oracle/graal/hotspot/test/AheadOfTimeCompilationTest.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.hotspot.test/src/com/oracle/graal/hotspot/test/AheadOfTimeCompilationTest.java Mon Nov 04 12:20:17 2013 +0100 @@ -23,6 +23,7 @@ package com.oracle.graal.hotspot.test; import static com.oracle.graal.api.code.CodeUtil.*; +import static com.oracle.graal.nodes.ConstantNode.*; import static com.oracle.graal.phases.GraalOptions.*; import org.junit.*; @@ -64,8 +65,8 @@ @Test public void testStaticFinalObjectAOT() { StructuredGraph result = compile("getStaticFinalObject", true); - assertEquals(1, result.getNodes().filter(ConstantNode.class).count()); - assertEquals(getCodeCache().getTarget().wordKind, result.getNodes().filter(ConstantNode.class).first().kind()); + assertEquals(1, getConstantNodes(result).count()); + assertEquals(getCodeCache().getTarget().wordKind, getConstantNodes(result).first().kind()); assertEquals(2, result.getNodes(FloatingReadNode.class).count()); assertEquals(0, result.getNodes().filter(ReadNode.class).count()); } @@ -73,8 +74,8 @@ @Test public void testStaticFinalObject() { StructuredGraph result = compile("getStaticFinalObject", false); - assertEquals(1, result.getNodes().filter(ConstantNode.class).count()); - assertEquals(Kind.Object, result.getNodes().filter(ConstantNode.class).first().kind()); + assertEquals(1, getConstantNodes(result).count()); + assertEquals(Kind.Object, getConstantNodes(result).first().kind()); assertEquals(0, result.getNodes(FloatingReadNode.class).count()); assertEquals(0, result.getNodes().filter(ReadNode.class).count()); } @@ -87,7 +88,7 @@ public void testClassObjectAOT() { StructuredGraph result = compile("getClassObject", true); - NodeIterable<ConstantNode> filter = result.getNodes().filter(ConstantNode.class); + NodeIterable<ConstantNode> filter = getConstantNodes(result); assertEquals(1, filter.count()); HotSpotResolvedObjectType type = (HotSpotResolvedObjectType) getMetaAccess().lookupJavaType(AheadOfTimeCompilationTest.class); assertEquals(type.klass(), filter.first().asConstant()); @@ -100,7 +101,7 @@ public void testClassObject() { StructuredGraph result = compile("getClassObject", false); - NodeIterable<ConstantNode> filter = result.getNodes().filter(ConstantNode.class); + NodeIterable<ConstantNode> filter = getConstantNodes(result); assertEquals(1, filter.count()); Object mirror = filter.first().asConstant().asObject(); assertEquals(Class.class, mirror.getClass()); @@ -117,7 +118,7 @@ @Test public void testPrimitiveClassObjectAOT() { StructuredGraph result = compile("getPrimitiveClassObject", true); - NodeIterable<ConstantNode> filter = result.getNodes().filter(ConstantNode.class); + NodeIterable<ConstantNode> filter = getConstantNodes(result); assertEquals(1, filter.count()); assertEquals(getCodeCache().getTarget().wordKind, filter.first().kind()); @@ -128,7 +129,7 @@ @Test public void testPrimitiveClassObject() { StructuredGraph result = compile("getPrimitiveClassObject", false); - NodeIterable<ConstantNode> filter = result.getNodes().filter(ConstantNode.class); + NodeIterable<ConstantNode> filter = getConstantNodes(result); assertEquals(1, filter.count()); Object mirror = filter.first().asConstant().asObject(); assertEquals(Class.class, mirror.getClass()); @@ -156,7 +157,7 @@ private void testStringObjectCommon(boolean compileAOT) { StructuredGraph result = compile("getStringObject", compileAOT); - NodeIterable<ConstantNode> filter = result.getNodes().filter(ConstantNode.class); + NodeIterable<ConstantNode> filter = getConstantNodes(result); assertEquals(1, filter.count()); Object mirror = filter.first().asConstant().asObject(); assertEquals(String.class, mirror.getClass()); @@ -177,8 +178,8 @@ assertEquals(2, result.getNodes(FloatingReadNode.class).count()); assertEquals(1, result.getNodes(PiNode.class).count()); - assertEquals(1, result.getNodes().filter(ConstantNode.class).count()); - ConstantNode constant = result.getNodes().filter(ConstantNode.class).first(); + assertEquals(1, getConstantNodes(result).count()); + ConstantNode constant = getConstantNodes(result).first(); assertEquals(Kind.Long, constant.kind()); assertEquals(((HotSpotResolvedObjectType) getMetaAccess().lookupJavaType(Boolean.class)).klass(), constant.asConstant()); } @@ -188,8 +189,8 @@ StructuredGraph result = compile("getBoxedBoolean", false); assertEquals(0, result.getNodes(FloatingReadNode.class).count()); assertEquals(0, result.getNodes(PiNode.class).count()); - assertEquals(1, result.getNodes().filter(ConstantNode.class).count()); - ConstantNode constant = result.getNodes().filter(ConstantNode.class).first(); + assertEquals(1, getConstantNodes(result).count()); + ConstantNode constant = getConstantNodes(result).first(); assertEquals(Kind.Object, constant.kind()); assertEquals(Boolean.TRUE, constant.asConstant().asObject()); }
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/VMToCompiler.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/VMToCompiler.java Mon Nov 04 12:20:17 2013 +0100 @@ -51,6 +51,8 @@ void bootstrap() throws Throwable; + void compileTheWorld() throws Throwable; + PrintStream log(); JavaMethod createUnresolvedJavaMethod(String name, String signature, JavaType holder);
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/VMToCompilerImpl.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/VMToCompilerImpl.java Mon Nov 04 12:20:17 2013 +0100 @@ -28,7 +28,6 @@ import static com.oracle.graal.hotspot.CompilationTask.*; import static com.oracle.graal.hotspot.HotSpotGraalRuntime.*; import static com.oracle.graal.java.GraphBuilderPhase.*; -import static com.oracle.graal.phases.GraalOptions.*; import static com.oracle.graal.phases.common.InliningUtil.*; import java.io.*; @@ -40,7 +39,7 @@ import com.oracle.graal.api.code.*; import com.oracle.graal.api.meta.*; import com.oracle.graal.compiler.*; -import com.oracle.graal.compiler.CompilerThreadFactory.*; +import com.oracle.graal.compiler.CompilerThreadFactory.DebugConfigAccess; import com.oracle.graal.debug.*; import com.oracle.graal.debug.internal.*; import com.oracle.graal.graph.*; @@ -334,11 +333,11 @@ } System.gc(); phaseTransition("bootstrap2"); + } - if (CompileTheWorld.getValue() != null) { - new CompileTheWorld().compile(); - System.exit(0); - } + public void compileTheWorld() throws Throwable { + new CompileTheWorld().compile(); + System.exit(0); } private MetricRateInPhase parsedBytecodesPerSecond;
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotHostLoweringProvider.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotHostLoweringProvider.java Mon Nov 04 12:20:17 2013 +0100 @@ -695,13 +695,14 @@ base /= scale; assert NumUtil.isInt(base); + StructuredGraph graph = location.graph(); if (index == null) { - return ConstantNode.forInt((int) base, location.graph()); + return ConstantNode.forInt((int) base, graph); } else { if (base == 0) { return index; } else { - return IntegerArithmeticNode.add(ConstantNode.forInt((int) base, location.graph()), index); + return IntegerArithmeticNode.add(graph, ConstantNode.forInt((int) base, graph), index); } } } @@ -709,7 +710,7 @@ private GuardingNode createBoundsCheck(AccessIndexedNode n, LoweringTool tool) { StructuredGraph g = n.graph(); ValueNode array = n.array(); - ValueNode arrayLength = readArrayLength(array, tool.getConstantReflection()); + ValueNode arrayLength = readArrayLength(n.graph(), array, tool.getConstantReflection()); if (arrayLength == null) { Stamp stamp = StampFactory.positiveInt(); ReadNode readArrayLength = g.add(new ReadNode(array, ConstantLocationNode.create(FINAL_LOCATION, Kind.Int, runtime.getConfig().arrayLengthOffset, g), stamp, BarrierType.NONE, false));
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedJavaMethod.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedJavaMethod.java Mon Nov 04 12:20:17 2013 +0100 @@ -74,11 +74,11 @@ */ public static HotSpotResolvedObjectType getHolder(long metaspaceMethod) { HotSpotVMConfig config = runtime().getConfig(); - long constMethod = unsafe.getLong(metaspaceMethod + config.methodConstMethodOffset); + long constMethod = unsafe.getAddress(metaspaceMethod + config.methodConstMethodOffset); assert constMethod != 0; - long constantPool = unsafe.getLong(constMethod + config.constMethodConstantsOffset); + long constantPool = unsafe.getAddress(constMethod + config.constMethodConstantsOffset); assert constantPool != 0; - long holder = unsafe.getLong(constantPool + config.constantPoolHolderOffset); + long holder = unsafe.getAddress(constantPool + config.constantPoolHolderOffset); assert holder != 0; return (HotSpotResolvedObjectType) HotSpotResolvedObjectType.fromMetaspaceKlass(holder); } @@ -231,7 +231,7 @@ return 0; } HotSpotVMConfig config = runtime().getConfig(); - long metaspaceConstMethod = unsafe.getLong(metaspaceMethod + config.methodConstMethodOffset); + long metaspaceConstMethod = unsafe.getAddress(metaspaceMethod + config.methodConstMethodOffset); return unsafe.getShort(metaspaceConstMethod + config.methodMaxLocalsOffset) & 0xFFFF; } @@ -242,7 +242,7 @@ return 0; } HotSpotVMConfig config = runtime().getConfig(); - long metaspaceConstMethod = unsafe.getLong(metaspaceMethod + config.methodConstMethodOffset); + long metaspaceConstMethod = unsafe.getAddress(metaspaceMethod + config.methodConstMethodOffset); return config.extraStackEntries + (unsafe.getShort(metaspaceConstMethod + config.constMethodMaxStackOffset) & 0xFFFF); }
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/phases/AheadOfTimeVerificationPhase.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/phases/AheadOfTimeVerificationPhase.java Mon Nov 04 12:20:17 2013 +0100 @@ -22,6 +22,8 @@ */ package com.oracle.graal.hotspot.phases; +import static com.oracle.graal.nodes.ConstantNode.*; + import com.oracle.graal.api.meta.*; import com.oracle.graal.nodes.*; import com.oracle.graal.phases.*; @@ -38,8 +40,8 @@ @Override protected boolean verify(StructuredGraph graph, PhaseContext context) { - for (ConstantNode node : graph.getNodes().filter(ConstantNode.class)) { - if (node.recordsUsages() || !node.gatherUsages().isEmpty()) { + for (ConstantNode node : getConstantNodes(graph)) { + if (node.recordsUsages() || !node.gatherUsages(graph).isEmpty()) { assert !isObject(node) || isNullReference(node) || isInternedString(node) : "illegal object constant: " + node; } }
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/phases/LoadJavaMirrorWithKlassPhase.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/phases/LoadJavaMirrorWithKlassPhase.java Mon Nov 04 12:20:17 2013 +0100 @@ -23,6 +23,7 @@ package com.oracle.graal.hotspot.phases; import static com.oracle.graal.api.meta.LocationIdentity.*; +import static com.oracle.graal.nodes.ConstantNode.*; import com.oracle.graal.api.meta.*; import com.oracle.graal.hotspot.meta.*; @@ -52,23 +53,30 @@ this.classMirrorOffset = classMirrorOffset; } + private FloatingReadNode getClassConstantReplacement(StructuredGraph graph, PhaseContext context, Constant constant) { + if (constant.getKind() == Kind.Object && constant.asObject() instanceof Class<?>) { + MetaAccessProvider metaAccess = context.getMetaAccess(); + ResolvedJavaType type = metaAccess.lookupJavaType((Class<?>) constant.asObject()); + assert type instanceof HotSpotResolvedObjectType; + + Constant klass = ((HotSpotResolvedObjectType) type).klass(); + ConstantNode klassNode = ConstantNode.forConstant(klass, metaAccess, graph); + + Stamp stamp = StampFactory.exactNonNull(metaAccess.lookupJavaType(Class.class)); + LocationNode location = graph.unique(ConstantLocationNode.create(FINAL_LOCATION, stamp.kind(), classMirrorOffset, graph)); + FloatingReadNode freadNode = graph.unique(new FloatingReadNode(klassNode, location, null, stamp)); + return freadNode; + } + return null; + } + @Override protected void run(StructuredGraph graph, PhaseContext context) { - for (ConstantNode node : graph.getNodes().filter(ConstantNode.class)) { + for (ConstantNode node : getConstantNodes(graph)) { Constant constant = node.asConstant(); - if (constant.getKind() == Kind.Object && constant.asObject() instanceof Class<?>) { - MetaAccessProvider metaAccess = context.getMetaAccess(); - ResolvedJavaType type = metaAccess.lookupJavaType((Class<?>) constant.asObject()); - assert type instanceof HotSpotResolvedObjectType; - - Constant klass = ((HotSpotResolvedObjectType) type).klass(); - ConstantNode klassNode = ConstantNode.forConstant(klass, metaAccess, graph); - - Stamp stamp = StampFactory.exactNonNull(metaAccess.lookupJavaType(Class.class)); - LocationNode location = graph.unique(ConstantLocationNode.create(FINAL_LOCATION, stamp.kind(), classMirrorOffset, graph)); - FloatingReadNode freadNode = graph.unique(new FloatingReadNode(klassNode, location, null, stamp)); - - node.replace(freadNode); + FloatingReadNode freadNode = getClassConstantReplacement(graph, context, constant); + if (freadNode != null) { + node.replace(graph, freadNode); } } }
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/replacements/InstanceOfSnippets.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/replacements/InstanceOfSnippets.java Mon Nov 04 12:20:17 2013 +0100 @@ -216,7 +216,7 @@ Arguments args; - StructuredGraph graph = hub.graph(); + StructuredGraph graph = instanceOf.graph(); if (hintInfo.hintHitProbability >= hintHitProbabilityThresholdForDeoptimizingSnippet()) { Hints hints = createHints(hintInfo, providers.getMetaAccess(), false, graph); args = new Arguments(instanceofWithProfile, graph.getGuardsStage());
--- a/graal/com.oracle.graal.lir.hsail/src/com/oracle/graal/lir/hsail/HSAILArithmetic.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.lir.hsail/src/com/oracle/graal/lir/hsail/HSAILArithmetic.java Mon Nov 04 12:20:17 2013 +0100 @@ -35,85 +35,89 @@ * Defines arithmetic instruction nodes. */ public enum HSAILArithmetic { - IADD, - OADD, - ISUB, - FSUB, - DSUB, - IMUL, - FMUL, - DMUL, - IDIV, + CALL, FDIV, - DDIV, - IUADD, - IUSUB, - IUMUL, - IUDIV, - LADD, + FREM, + D2F, + D2I, + D2L, DADD, - FADD, - LSUB, - LMUL, - LDIV, - LUADD, - LUSUB, - LUMUL, - LUDIV, - IMAX, - LMAX, - IUMAX, - LUMAX, - IMIN, - LMIN, - IUMIN, - LUMIN, - IREM, - LREM, - FREM, + DDIV, + DMUL, + DNEG, DREM, - IUREM, - LUREM, - ICARRY, - LCARRY, - IUCARRY, - LUCARRY, - IAND, - LAND, - INEG, - INOT, - I2B, - I2S, - I2L, - I2C, + DSUB, F2D, F2I, F2L, - D2F, - I2F, + FADD, + FMUL, + FNEG, + FSUB, + I2B, + I2C, I2D, - D2I, + I2F, + I2L, + I2S, + IADD, + IAND, + ICARRY, + IDIV, + IMAX, + IMIN, + IMUL, + INEG, + INOT, + IOR, + IREM, + ISHL, + ISHR, + ISUB, + IUADD, + IUCARRY, + IUDIV, + IUMAX, + IUMIN, + IUMUL, + IUREM, + IUSHR, + IUSUB, + IXOR, + L2D, L2F, - D2L, + L2I, + LADD, + LAND, + LCARRY, + LDIV, + LMAX, + LMIN, + LMUL, + LNEG, + LNOT, + LOR, + LREM, + LSHL, + LSHR, + LSUB, + LUADD, + LUCARRY, + LUDIV, + LUMAX, + LUMIN, + LUMUL, + LUREM, + LUSHR, + LUSUB, + LXOR, MOV_F2I, MOV_D2L, - L2D, MOV_I2F, MOV_L2D, - ISHL, - LSHL, - ISHR, - LSHR, - IUSHR, - LUSHR, + OADD, SQRT, - UNDEF, - CALL, - L2I, - IXOR, - LXOR, - IOR, - LOR; + UNDEF; public static class Op1Stack extends HSAILLIRInstruction { @Opcode private final HSAILArithmetic opcode; @@ -313,7 +317,7 @@ masm.emitConvertIntToByte(dst, src); break; case SQRT: - masm.emitArg1("sqrt", dst, src); + masm.emit("sqrt", dst, src); break; case UNDEF: masm.undefined("undefined node"); @@ -322,7 +326,16 @@ masm.undefined("undefined node CALL"); break; case INOT: - masm.emitArg1("not", dst, src); + case LNOT: + // Emit the HSAIL instruction for a bitwise not. + masm.emitForceBitwise("not", dst, src); + break; + case INEG: + case LNEG: + case FNEG: + case DNEG: + // Emit the HSAIL instruction for a negate operation. + masm.emit("neg", dst, src); break; default: throw GraalInternalError.shouldNotReachHere(); @@ -402,15 +415,15 @@ break; case IAND: case LAND: - masm.emitBitwiseLogical("and", dst, src1, src2); + masm.emitForceBitwise("and", dst, src1, src2); break; case IXOR: case LXOR: - masm.emitBitwiseLogical("xor", dst, src1, src2); + masm.emitForceBitwise("xor", dst, src1, src2); break; case IOR: case LOR: - masm.emitBitwiseLogical("or", dst, src1, src2); + masm.emitForceBitwise("or", dst, src1, src2); break; case IREM: masm.emit("rem", dst, src1, src2);
--- a/graal/com.oracle.graal.lir.hsail/src/com/oracle/graal/lir/hsail/HSAILControlFlow.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.lir.hsail/src/com/oracle/graal/lir/hsail/HSAILControlFlow.java Mon Nov 04 12:20:17 2013 +0100 @@ -265,6 +265,6 @@ default: throw GraalInternalError.shouldNotReachHere(); } - masm.cmovCommon(result, trueValue, falseValue, width); + masm.emitConditionalMove(result, trueValue, falseValue, width); } }
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMemOp.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMemOp.java Mon Nov 04 12:20:17 2013 +0100 @@ -166,7 +166,9 @@ case Long: case Float: case Double: - new Ld(Parameter, result, addr.getBase(), Constant.forLong(addr.getDisplacement())).emit(masm); + Ld ldIns = new Ld(Parameter, result, addr.getBase(), Constant.forLong(addr.getDisplacement())); + ldIns.setLdRetAddrInstruction(true); + ldIns.emit(masm); break; default: throw GraalInternalError.shouldNotReachHere();
--- a/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/BasicInductionVariable.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/BasicInductionVariable.java Mon Nov 04 12:20:17 2013 +0100 @@ -45,6 +45,11 @@ } @Override + public StructuredGraph graph() { + return phi.graph(); + } + + @Override public Direction direction() { Stamp stamp = rawStride.stamp(); if (stamp instanceof IntegerStamp) { @@ -83,7 +88,7 @@ return rawStride; } if (op instanceof IntegerSubNode) { - return rawStride.graph().unique(new NegateNode(rawStride)); + return graph().unique(new NegateNode(rawStride)); } throw GraalInternalError.shouldNotReachHere(); } @@ -117,7 +122,7 @@ @Override public ValueNode extremumNode(boolean assumePositiveTripCount, Kind kind) { Kind fromKind = phi.kind(); - Graph graph = phi.graph(); + StructuredGraph graph = graph(); ValueNode stride = strideNode(); ValueNode maxTripCount = loop.counted().maxTripCountNode(assumePositiveTripCount); ValueNode initNode = this.initNode(); @@ -127,13 +132,13 @@ maxTripCount = graph.unique(new ConvertNode(convertOp, maxTripCount)); initNode = graph.unique(new ConvertNode(convertOp, initNode)); } - return IntegerArithmeticNode.add(IntegerArithmeticNode.mul(stride, IntegerArithmeticNode.sub(maxTripCount, ConstantNode.forIntegerKind(kind, 1, graph))), initNode); + return IntegerArithmeticNode.add(graph, IntegerArithmeticNode.mul(graph, stride, IntegerArithmeticNode.sub(graph, maxTripCount, ConstantNode.forIntegerKind(kind, 1, graph))), initNode); } @Override public ValueNode exitValueNode() { ValueNode maxTripCount = loop.counted().maxTripCountNode(false); - return IntegerArithmeticNode.add(IntegerArithmeticNode.mul(strideNode(), maxTripCount), initNode()); + return IntegerArithmeticNode.add(graph(), IntegerArithmeticNode.mul(graph(), strideNode(), maxTripCount), initNode()); } @Override
--- a/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/CountedLoopInfo.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/CountedLoopInfo.java Mon Nov 04 12:20:17 2013 +0100 @@ -25,7 +25,6 @@ import static com.oracle.graal.nodes.calc.IntegerArithmeticNode.*; import com.oracle.graal.api.meta.*; -import com.oracle.graal.graph.*; import com.oracle.graal.loop.InductionVariable.Direction; import com.oracle.graal.nodes.*; import com.oracle.graal.nodes.calc.*; @@ -54,12 +53,12 @@ public ValueNode maxTripCountNode(boolean assumePositive) { StructuredGraph graph = iv.valueNode().graph(); Kind kind = iv.valueNode().kind(); - IntegerArithmeticNode range = IntegerArithmeticNode.sub(end, iv.initNode()); + IntegerArithmeticNode range = IntegerArithmeticNode.sub(graph, end, iv.initNode()); if (oneOff) { if (iv.direction() == Direction.Up) { - range = IntegerArithmeticNode.add(range, ConstantNode.forIntegerKind(kind, 1, graph)); + range = IntegerArithmeticNode.add(graph, range, ConstantNode.forIntegerKind(kind, 1, graph)); } else { - range = IntegerArithmeticNode.sub(range, ConstantNode.forIntegerKind(kind, 1, graph)); + range = IntegerArithmeticNode.sub(graph, range, ConstantNode.forIntegerKind(kind, 1, graph)); } } IntegerDivNode div = graph.add(new IntegerDivNode(kind, range, iv.strideNode())); @@ -139,20 +138,20 @@ return overflowGuard; } Kind kind = iv.valueNode().kind(); - Graph graph = iv.valueNode().graph(); + StructuredGraph graph = iv.valueNode().graph(); CompareNode cond; // we use a negated guard with a < condition to achieve a >= ConstantNode one = ConstantNode.forIntegerKind(kind, 1, graph); if (iv.direction() == Direction.Up) { - IntegerArithmeticNode v1 = sub(ConstantNode.forIntegerKind(kind, kind.getMaxValue(), graph), sub(iv.strideNode(), one)); + IntegerArithmeticNode v1 = sub(graph, ConstantNode.forIntegerKind(kind, kind.getMaxValue(), graph), sub(graph, iv.strideNode(), one)); if (oneOff) { - v1 = sub(v1, one); + v1 = sub(graph, v1, one); } cond = graph.unique(new IntegerLessThanNode(v1, end)); } else { assert iv.direction() == Direction.Down; - IntegerArithmeticNode v1 = add(ConstantNode.forIntegerKind(kind, kind.getMinValue(), graph), sub(one, iv.strideNode())); + IntegerArithmeticNode v1 = add(graph, ConstantNode.forIntegerKind(kind, kind.getMinValue(), graph), sub(graph, one, iv.strideNode())); if (oneOff) { - v1 = add(v1, one); + v1 = add(graph, v1, one); } cond = graph.unique(new IntegerLessThanNode(end, v1)); }
--- a/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/DerivedOffsetInductionVariable.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/DerivedOffsetInductionVariable.java Mon Nov 04 12:20:17 2013 +0100 @@ -41,6 +41,11 @@ } @Override + public StructuredGraph graph() { + return base.graph(); + } + + @Override public Direction direction() { return base.direction(); } @@ -81,7 +86,7 @@ @Override public ValueNode strideNode() { if (value instanceof IntegerSubNode && base.valueNode() == value.y()) { - return value.graph().unique(new NegateNode(base.strideNode())); + return graph().unique(new NegateNode(base.strideNode())); } return base.strideNode(); } @@ -123,14 +128,14 @@ private ValueNode op(ValueNode b, ValueNode o) { if (value instanceof IntegerAddNode) { - return IntegerArithmeticNode.add(b, o); + return IntegerArithmeticNode.add(graph(), b, o); } if (value instanceof IntegerSubNode) { if (base.valueNode() == value.x()) { - return IntegerArithmeticNode.sub(b, o); + return IntegerArithmeticNode.sub(graph(), b, o); } else { assert base.valueNode() == value.y(); - return IntegerArithmeticNode.sub(o, b); + return IntegerArithmeticNode.sub(graph(), o, b); } } throw GraalInternalError.shouldNotReachHere();
--- a/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/DerivedScaledInductionVariable.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/DerivedScaledInductionVariable.java Mon Nov 04 12:20:17 2013 +0100 @@ -48,6 +48,11 @@ } @Override + public StructuredGraph graph() { + return base.graph(); + } + + @Override public Direction direction() { Stamp stamp = scale.stamp(); if (stamp instanceof IntegerStamp) { @@ -68,12 +73,12 @@ @Override public ValueNode initNode() { - return IntegerArithmeticNode.mul(base.initNode(), scale); + return IntegerArithmeticNode.mul(graph(), base.initNode(), scale); } @Override public ValueNode strideNode() { - return IntegerArithmeticNode.mul(base.strideNode(), scale); + return IntegerArithmeticNode.mul(graph(), base.strideNode(), scale); } @Override @@ -98,12 +103,12 @@ @Override public ValueNode extremumNode(boolean assumePositiveTripCount, Kind kind) { - return IntegerArithmeticNode.mul(base.extremumNode(assumePositiveTripCount, kind), ConvertNode.convert(kind, scale)); + return IntegerArithmeticNode.mul(graph(), base.extremumNode(assumePositiveTripCount, kind), ConvertNode.convert(kind, scale)); } @Override public ValueNode exitValueNode() { - return IntegerArithmeticNode.mul(base.exitValueNode(), scale); + return IntegerArithmeticNode.mul(graph(), base.exitValueNode(), scale); } @Override
--- a/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/InductionVariable.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/InductionVariable.java Mon Nov 04 12:20:17 2013 +0100 @@ -46,6 +46,8 @@ } } + public abstract StructuredGraph graph(); + protected final LoopEx loop; public InductionVariable(LoopEx loop) {
--- a/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/LoopEx.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/LoopEx.java Mon Nov 04 12:20:17 2013 +0100 @@ -77,7 +77,7 @@ } public boolean isOutsideLoop(Node n) { - return !whole().contains(n); + return n.isExternal() || !whole().contains(n); } public LoopBeginNode loopBegin() {
--- a/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/LoopFragment.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.loop/src/com/oracle/graal/loop/LoopFragment.java Mon Nov 04 12:20:17 2013 +0100 @@ -72,7 +72,10 @@ @SuppressWarnings("unchecked") public <New extends Node, Old extends New> New getDuplicatedNode(Old n) { assert isDuplicate(); - return (New) duplicationMap.get(n); + if (!n.isExternal()) { + return (New) duplicationMap.get(n); + } + return n; } protected <New extends Node, Old extends New> void putDuplicatedNode(Old oldNode, New newNode) {
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/ConstantNode.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/ConstantNode.java Mon Nov 04 12:20:17 2013 +0100 @@ -25,23 +25,45 @@ import java.util.*; import com.oracle.graal.api.meta.*; +import com.oracle.graal.debug.*; import com.oracle.graal.graph.*; +import com.oracle.graal.graph.iterators.*; import com.oracle.graal.nodes.calc.*; import com.oracle.graal.nodes.spi.*; import com.oracle.graal.nodes.type.*; /** - * The {@code ConstantNode} represents a constant such as an integer value, long, float, object - * reference, address, etc. + * The {@code ConstantNode} represents a {@link Constant constant}. */ @NodeInfo(shortName = "Const", nameTemplate = "Const({p#rawvalue})") -public class ConstantNode extends FloatingNode implements LIRLowerable { +public final class ConstantNode extends FloatingNode implements LIRLowerable { - public final Constant value; + private static final DebugMetric ConstantNodes = Debug.metric("ConstantNodes"); + + private final Constant value; protected ConstantNode(Constant value) { super(StampFactory.forConstant(value)); this.value = value; + ConstantNodes.increment(); + } + + /** + * Constructs a new ConstantNode representing the specified constant. + * + * @param value the constant + */ + protected ConstantNode(Constant value, MetaAccessProvider metaAccess) { + super(StampFactory.forConstant(value, metaAccess)); + this.value = value; + ConstantNodes.increment(); + } + + /** + * @return the constant value represented by this node + */ + public Constant getValue() { + return value; } /** @@ -55,24 +77,24 @@ return ConstantNodeRecordsUsages; } - /** - * Constructs a new ConstantNode representing the specified constant. - * - * @param value the constant - */ - protected ConstantNode(Constant value, MetaAccessProvider metaAccess) { - super(StampFactory.forConstant(value, metaAccess)); - this.value = value; + @Override + public boolean isAlive() { + return true; + } + + @Override + public boolean isExternal() { + return true; } /** * Computes the usages of this node by iterating over all the nodes in the graph, searching for * those that have this node as an input. */ - public List<Node> gatherUsages() { + public List<Node> gatherUsages(StructuredGraph graph) { assert !ConstantNodeRecordsUsages; List<Node> usages = new ArrayList<>(); - for (Node node : graph().getNodes()) { + for (Node node : graph.getNodes()) { for (Node input : node.inputs()) { if (input == this) { usages.add(node); @@ -82,20 +104,45 @@ return usages; } - public void replace(Node replacement) { + /** + * Gathers all the {@link ConstantNode}s that are inputs to the {@linkplain Graph#getNodes() + * live nodes} in a given graph. This is an expensive operation that should only be used in + * test/verification/AOT code. + */ + public static NodeIterable<ConstantNode> getConstantNodes(StructuredGraph graph) { + Map<ConstantNode, ConstantNode> result = new HashMap<>(); + for (Node node : graph.getNodes()) { + for (Node input : node.inputs()) { + if (input instanceof ConstantNode) { + result.put((ConstantNode) input, (ConstantNode) input); + } + } + } + return new ConstantNodeList(result.keySet()); + } + + /** + * Replaces this node at its usages with another node. If {@value #ConstantNodeRecordsUsages} is + * false, this is an expensive operation that should only be used in test/verification/AOT code. + */ + public void replace(StructuredGraph graph, Node replacement) { if (!recordsUsages()) { - List<Node> usages = gatherUsages(); + List<Node> usages = gatherUsages(graph); for (Node usage : usages) { usage.replaceFirstInput(this, replacement); } - graph().removeFloating(this); + if (!isExternal()) { + graph.removeFloating(this); + } } else { + assert graph == graph(); graph().replaceFloating(this, replacement); } } @Override public void generate(LIRGeneratorTool gen) { + assert ConstantNodeRecordsUsages : "LIR generator should generate constants per-usage"; if (gen.canInlineConstant(value) || onlyUsedInVirtualState()) { gen.setResult(this, value); } else { @@ -118,9 +165,9 @@ if (constant.getKind().getStackKind() == Kind.Int && constant.getKind() != Kind.Int) { return forInt(constant.asInt(), graph); } else if (constant.getKind() == Kind.Object) { - return graph.unique(new ConstantNode(constant, metaAccess)); + return unique(graph, new ConstantNode(constant, metaAccess)); } else { - return graph.unique(new ConstantNode(constant)); + return unique(graph, new ConstantNode(constant)); } } @@ -136,100 +183,95 @@ * Returns a node for a double constant. * * @param d the double value for which to create the instruction - * @param graph * @return a node for a double constant */ public static ConstantNode forDouble(double d, Graph graph) { - return graph.unique(new ConstantNode(Constant.forDouble(d))); + return unique(graph, new ConstantNode(Constant.forDouble(d))); } /** * Returns a node for a float constant. * * @param f the float value for which to create the instruction - * @param graph * @return a node for a float constant */ public static ConstantNode forFloat(float f, Graph graph) { - return graph.unique(new ConstantNode(Constant.forFloat(f))); + return unique(graph, new ConstantNode(Constant.forFloat(f))); } /** * Returns a node for an long constant. * * @param i the long value for which to create the instruction - * @param graph * @return a node for an long constant */ public static ConstantNode forLong(long i, Graph graph) { - return graph.unique(new ConstantNode(Constant.forLong(i))); + return unique(graph, new ConstantNode(Constant.forLong(i))); } /** * Returns a node for an integer constant. * * @param i the integer value for which to create the instruction - * @param graph * @return a node for an integer constant */ public static ConstantNode forInt(int i, Graph graph) { - return graph.unique(new ConstantNode(Constant.forInt(i))); + return unique(graph, new ConstantNode(Constant.forInt(i))); } /** * Returns a node for a boolean constant. * * @param i the boolean value for which to create the instruction - * @param graph * @return a node representing the boolean */ public static ConstantNode forBoolean(boolean i, Graph graph) { - return graph.unique(new ConstantNode(Constant.forInt(i ? 1 : 0))); + return unique(graph, new ConstantNode(Constant.forInt(i ? 1 : 0))); } /** * Returns a node for a byte constant. * * @param i the byte value for which to create the instruction - * @param graph * @return a node representing the byte */ public static ConstantNode forByte(byte i, Graph graph) { - return graph.unique(new ConstantNode(Constant.forInt(i))); + return unique(graph, new ConstantNode(Constant.forInt(i))); } /** * Returns a node for a char constant. * * @param i the char value for which to create the instruction - * @param graph * @return a node representing the char */ public static ConstantNode forChar(char i, Graph graph) { - return graph.unique(new ConstantNode(Constant.forInt(i))); + return unique(graph, new ConstantNode(Constant.forInt(i))); } /** * Returns a node for a short constant. * * @param i the short value for which to create the instruction - * @param graph * @return a node representing the short */ public static ConstantNode forShort(short i, Graph graph) { - return graph.unique(new ConstantNode(Constant.forInt(i))); + return unique(graph, new ConstantNode(Constant.forInt(i))); } /** * Returns a node for an object constant. * * @param o the object value for which to create the instruction - * @param graph * @return a node representing the object */ public static ConstantNode forObject(Object o, MetaAccessProvider metaAccess, Graph graph) { assert !(o instanceof Constant) : "wrapping a Constant into a Constant"; - return graph.unique(new ConstantNode(Constant.forObject(o), metaAccess)); + return unique(graph, new ConstantNode(Constant.forObject(o), metaAccess)); + } + + private static ConstantNode unique(Graph graph, ConstantNode node) { + return graph.uniqueWithoutAdd(node); } public static ConstantNode forIntegerKind(Kind kind, long value, Graph graph) { @@ -292,4 +334,15 @@ return super.toString(verbosity); } } + + static class ConstantNodeList extends NodeList<ConstantNode> { + + public ConstantNodeList(Collection<ConstantNode> nodes) { + super(nodes.toArray(new ConstantNode[nodes.size()])); + } + + @Override + protected void update(ConstantNode oldNode, ConstantNode newNode) { + } + } }
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/ValueNode.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/ValueNode.java Mon Nov 04 12:20:17 2013 +0100 @@ -110,7 +110,7 @@ * @return {@code true} if this value represents the null constant */ public final boolean isNullConstant() { - return this instanceof ConstantNode && ((ConstantNode) this).value.isNull(); + return this instanceof ConstantNode && ((ConstantNode) this).getValue().isNull(); } /** @@ -121,7 +121,7 @@ */ public final Constant asConstant() { if (this instanceof ConstantNode) { - return ((ConstantNode) this).value; + return ((ConstantNode) this).getValue(); } return null; }
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/BinaryNode.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/BinaryNode.java Mon Nov 04 12:20:17 2013 +0100 @@ -83,7 +83,7 @@ } } - public static BinaryNode add(ValueNode x, ValueNode y) { + public static BinaryNode add(StructuredGraph graph, ValueNode x, ValueNode y) { assert x.kind() == y.kind(); switch (x.kind()) { case Byte: @@ -91,7 +91,7 @@ case Short: case Int: case Long: - return IntegerArithmeticNode.add(x, y); + return IntegerArithmeticNode.add(graph, x, y); case Float: case Double: return x.graph().unique(new FloatAddNode(x.kind(), x, y, false)); @@ -100,7 +100,7 @@ } } - public static BinaryNode sub(ValueNode x, ValueNode y) { + public static BinaryNode sub(StructuredGraph graph, ValueNode x, ValueNode y) { assert x.kind() == y.kind(); switch (x.kind()) { case Byte: @@ -108,7 +108,7 @@ case Short: case Int: case Long: - return IntegerArithmeticNode.sub(x, y); + return IntegerArithmeticNode.sub(graph, x, y); case Float: case Double: return x.graph().unique(new FloatSubNode(x.kind(), x, y, false)); @@ -117,7 +117,7 @@ } } - public static BinaryNode mul(ValueNode x, ValueNode y) { + public static BinaryNode mul(StructuredGraph graph, ValueNode x, ValueNode y) { assert x.kind() == y.kind(); switch (x.kind()) { case Byte: @@ -125,7 +125,7 @@ case Short: case Int: case Long: - return IntegerArithmeticNode.mul(x, y); + return IntegerArithmeticNode.mul(graph, x, y); case Float: case Double: return x.graph().unique(new FloatMulNode(x.kind(), x, y, false)); @@ -168,9 +168,10 @@ //@formatter:on /** * Tries to re-associate values which satisfy the criterion. For example with a constantness - * criterion : (a + 2) + 1 => a + (1 + 2)<br> - * This method accepts only reassociable operations (see - * {@linkplain #canTryReassociate(BinaryNode)}) such as +, -, *, &, | and ^ + * criterion: {@code (a + 2) + 1 => a + (1 + 2)} + * <p> + * This method accepts only {@linkplain #canTryReassociate(BinaryNode) reassociable} operations + * such as +, -, *, &, | and ^ */ public static BinaryNode reassociate(BinaryNode node, NodePredicate criterion) { assert canTryReassociate(node); @@ -217,28 +218,29 @@ ValueNode a = match2.getOtherValue(other); if (node instanceof IntegerAddNode || node instanceof IntegerSubNode) { BinaryNode associated; + StructuredGraph graph = node.graph(); if (invertM1) { - associated = IntegerArithmeticNode.sub(m2, m1); + associated = IntegerArithmeticNode.sub(graph, m2, m1); } else if (invertM2) { - associated = IntegerArithmeticNode.sub(m1, m2); + associated = IntegerArithmeticNode.sub(graph, m1, m2); } else { - associated = IntegerArithmeticNode.add(m1, m2); + associated = IntegerArithmeticNode.add(graph, m1, m2); } if (invertA) { - return IntegerArithmeticNode.sub(associated, a); + return IntegerArithmeticNode.sub(graph, associated, a); } if (aSub) { - return IntegerArithmeticNode.sub(a, associated); + return IntegerArithmeticNode.sub(graph, a, associated); } - return IntegerArithmeticNode.add(a, associated); + return IntegerArithmeticNode.add(graph, a, associated); } else if (node instanceof IntegerMulNode) { - return IntegerArithmeticNode.mul(a, IntegerAddNode.mul(m1, m2)); + return IntegerArithmeticNode.mul(node.graph(), a, IntegerAddNode.mul(node.graph(), m1, m2)); } else if (node instanceof AndNode) { - return BitLogicNode.and(a, BitLogicNode.and(m1, m2)); + return BitLogicNode.and(node.graph(), a, BitLogicNode.and(node.graph(), m1, m2)); } else if (node instanceof OrNode) { - return BitLogicNode.or(a, BitLogicNode.or(m1, m2)); + return BitLogicNode.or(node.graph(), a, BitLogicNode.or(node.graph(), m1, m2)); } else if (node instanceof XorNode) { - return BitLogicNode.xor(a, BitLogicNode.xor(m1, m2)); + return BitLogicNode.xor(node.graph(), a, BitLogicNode.xor(node.graph(), m1, m2)); } else { throw GraalInternalError.shouldNotReachHere(); }
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/BitLogicNode.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/BitLogicNode.java Mon Nov 04 12:20:17 2013 +0100 @@ -23,7 +23,6 @@ package com.oracle.graal.nodes.calc; import com.oracle.graal.api.meta.*; -import com.oracle.graal.graph.*; import com.oracle.graal.nodes.*; import com.oracle.graal.nodes.spi.*; @@ -43,9 +42,8 @@ assert kind == Kind.Int || kind == Kind.Long; } - public static BitLogicNode and(ValueNode v1, ValueNode v2) { - assert v1.kind() == v2.kind() && v1.graph() == v2.graph(); - Graph graph = v1.graph(); + public static BitLogicNode and(StructuredGraph graph, ValueNode v1, ValueNode v2) { + assert v1.kind() == v2.kind(); switch (v1.kind()) { case Int: return graph.unique(new AndNode(Kind.Int, v1, v2)); @@ -56,9 +54,8 @@ } } - public static BitLogicNode or(ValueNode v1, ValueNode v2) { - assert v1.kind() == v2.kind() && v1.graph() == v2.graph(); - Graph graph = v1.graph(); + public static BitLogicNode or(StructuredGraph graph, ValueNode v1, ValueNode v2) { + assert v1.kind() == v2.kind(); switch (v1.kind()) { case Int: return graph.unique(new OrNode(Kind.Int, v1, v2)); @@ -69,9 +66,8 @@ } } - public static BitLogicNode xor(ValueNode v1, ValueNode v2) { - assert v1.kind() == v2.kind() && v1.graph() == v2.graph(); - Graph graph = v1.graph(); + public static BitLogicNode xor(StructuredGraph graph, ValueNode v1, ValueNode v2) { + assert v1.kind() == v2.kind(); switch (v1.kind()) { case Int: return graph.unique(new XorNode(Kind.Int, v1, v2));
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/CompareNode.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/CompareNode.java Mon Nov 04 12:20:17 2013 +0100 @@ -144,7 +144,7 @@ return this; } - public static CompareNode createCompareNode(Condition condition, ValueNode x, ValueNode y) { + public static CompareNode createCompareNode(StructuredGraph graph, Condition condition, ValueNode x, ValueNode y) { assert x.kind() == y.kind(); assert condition.isCanonical() : "condition is not canonical: " + condition; assert x.kind() != Kind.Double && x.kind() != Kind.Float; @@ -166,6 +166,6 @@ comparison = new IntegerBelowThanNode(x, y); } - return x.graph().unique(comparison); + return graph.unique(comparison); } }
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/ConditionalNode.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/ConditionalNode.java Mon Nov 04 12:20:17 2013 +0100 @@ -113,8 +113,8 @@ generator.emitConditional(this); } - private ConditionalNode(Condition condition, ValueNode x, ValueNode y) { - this(createCompareNode(condition, x, y)); + private ConditionalNode(@InjectedNodeParameter StructuredGraph graph, Condition condition, ValueNode x, ValueNode y) { + this(createCompareNode(graph, condition, x, y)); } private ConditionalNode(ValueNode type, ValueNode object) {
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/IntegerAddNode.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/IntegerAddNode.java Mon Nov 04 12:20:17 2013 +0100 @@ -80,17 +80,17 @@ } if (c < 0) { if (kind() == Kind.Int) { - return IntegerArithmeticNode.sub(x(), ConstantNode.forInt((int) -c, graph())); + return IntegerArithmeticNode.sub(graph(), x(), ConstantNode.forInt((int) -c, graph())); } else { assert kind() == Kind.Long; - return IntegerArithmeticNode.sub(x(), ConstantNode.forLong(-c, graph())); + return IntegerArithmeticNode.sub(graph(), x(), ConstantNode.forLong(-c, graph())); } } } if (x() instanceof NegateNode) { - return IntegerArithmeticNode.sub(y(), ((NegateNode) x()).x()); + return IntegerArithmeticNode.sub(graph(), y(), ((NegateNode) x()).x()); } else if (y() instanceof NegateNode) { - return IntegerArithmeticNode.sub(x(), ((NegateNode) y()).x()); + return IntegerArithmeticNode.sub(graph(), x(), ((NegateNode) y()).x()); } return this; }
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/IntegerArithmeticNode.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/IntegerArithmeticNode.java Mon Nov 04 12:20:17 2013 +0100 @@ -23,7 +23,6 @@ package com.oracle.graal.nodes.calc; import com.oracle.graal.api.meta.*; -import com.oracle.graal.graph.*; import com.oracle.graal.nodes.*; import com.oracle.graal.nodes.spi.*; @@ -34,9 +33,8 @@ assert kind == Kind.Int || kind == Kind.Long; } - public static IntegerAddNode add(ValueNode v1, ValueNode v2) { - assert v1.kind() == v2.kind() && v1.graph() == v2.graph(); - Graph graph = v1.graph(); + public static IntegerAddNode add(StructuredGraph graph, ValueNode v1, ValueNode v2) { + assert v1.kind() == v2.kind(); switch (v1.kind()) { case Int: return graph.unique(new IntegerAddNode(Kind.Int, v1, v2)); @@ -47,9 +45,8 @@ } } - public static IntegerMulNode mul(ValueNode v1, ValueNode v2) { - assert v1.kind() == v2.kind() && v1.graph() == v2.graph(); - Graph graph = v1.graph(); + public static IntegerMulNode mul(StructuredGraph graph, ValueNode v1, ValueNode v2) { + assert v1.kind() == v2.kind(); switch (v1.kind()) { case Int: return graph.unique(new IntegerMulNode(Kind.Int, v1, v2)); @@ -60,9 +57,8 @@ } } - public static IntegerSubNode sub(ValueNode v1, ValueNode v2) { - assert v1.kind() == v2.kind() && v1.graph() == v2.graph(); - Graph graph = v1.graph(); + public static IntegerSubNode sub(StructuredGraph graph, ValueNode v1, ValueNode v2) { + assert v1.kind() == v2.kind(); switch (v1.kind()) { case Int: return graph.unique(new IntegerSubNode(Kind.Int, v1, v2));
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/IntegerBelowThanNode.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/IntegerBelowThanNode.java Mon Nov 04 12:20:17 2013 +0100 @@ -60,7 +60,7 @@ } else { if (x().isConstant() && x().asConstant().asLong() == 0) { // 0 |<| y is the same as 0 != y - return graph().unique(new LogicNegationNode(CompareNode.createCompareNode(Condition.EQ, x(), y()))); + return graph().unique(new LogicNegationNode(CompareNode.createCompareNode(graph(), Condition.EQ, x(), y()))); } if (x().stamp() instanceof IntegerStamp && y().stamp() instanceof IntegerStamp) {
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/IntegerDivNode.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/IntegerDivNode.java Mon Nov 04 12:20:17 2013 +0100 @@ -79,7 +79,7 @@ } RightShiftNode sign = graph().unique(new RightShiftNode(kind(), x(), ConstantNode.forInt(bits - 1, graph()))); UnsignedRightShiftNode round = graph().unique(new UnsignedRightShiftNode(kind(), sign, ConstantNode.forInt(bits - log2, graph()))); - dividend = IntegerArithmeticNode.add(dividend, round); + dividend = IntegerArithmeticNode.add(graph(), dividend, round); } RightShiftNode shift = graph().unique(new RightShiftNode(kind(), dividend, ConstantNode.forInt(log2, graph()))); if (c < 0) {
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/IntegerSubNode.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/IntegerSubNode.java Mon Nov 04 12:20:17 2013 +0100 @@ -99,10 +99,10 @@ } if (c < 0) { if (kind() == Kind.Int) { - return IntegerArithmeticNode.add(x(), ConstantNode.forInt((int) -c, graph())); + return IntegerArithmeticNode.add(graph(), x(), ConstantNode.forInt((int) -c, graph())); } else { assert kind() == Kind.Long; - return IntegerArithmeticNode.add(x(), ConstantNode.forLong(-c, graph())); + return IntegerArithmeticNode.add(graph(), x(), ConstantNode.forLong(-c, graph())); } } } else if (x().isConstant()) { @@ -113,7 +113,7 @@ return BinaryNode.reassociate(this, ValueNode.isConstantPredicate()); } if (y() instanceof NegateNode) { - return IntegerArithmeticNode.add(x(), ((NegateNode) y()).x()); + return IntegerArithmeticNode.add(graph(), x(), ((NegateNode) y()).x()); } return this; }
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/NegateNode.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/NegateNode.java Mon Nov 04 12:20:17 2013 +0100 @@ -82,7 +82,7 @@ } if (x() instanceof IntegerSubNode) { IntegerSubNode sub = (IntegerSubNode) x; - return IntegerArithmeticNode.sub(sub.y(), sub.x()); + return IntegerArithmeticNode.sub(graph(), sub.y(), sub.x()); } return this; }
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/extended/AddLocationNode.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/extended/AddLocationNode.java Mon Nov 04 12:20:17 2013 +0100 @@ -82,7 +82,7 @@ IndexedLocationNode yIdx = (IndexedLocationNode) y; if (xIdx.getIndexScaling() == yIdx.getIndexScaling()) { long displacement = xIdx.getDisplacement() + yIdx.getDisplacement(); - ValueNode index = IntegerArithmeticNode.add(xIdx.getIndex(), yIdx.getIndex()); + ValueNode index = IntegerArithmeticNode.add(graph(), xIdx.getIndex(), yIdx.getIndex()); return IndexedLocationNode.create(getLocationIdentity(), getValueKind(), displacement, index, graph(), xIdx.getIndexScaling()); } }
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/java/ArrayLengthNode.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/java/ArrayLengthNode.java Mon Nov 04 12:20:17 2013 +0100 @@ -47,7 +47,7 @@ @Override public Node canonical(CanonicalizerTool tool) { - ValueNode length = readArrayLength(array(), tool.getConstantReflection()); + ValueNode length = readArrayLength(graph(), array(), tool.getConstantReflection()); if (length != null) { return length; } @@ -57,10 +57,12 @@ /** * Gets the length of an array if possible. * + * @param graph TODO * @param array an array + * * @return a node representing the length of {@code array} or null if it is not available */ - public static ValueNode readArrayLength(ValueNode array, ConstantReflectionProvider constantReflection) { + public static ValueNode readArrayLength(StructuredGraph graph, ValueNode array, ConstantReflectionProvider constantReflection) { if (array instanceof ArrayLengthProvider) { ValueNode length = ((ArrayLengthProvider) array).length(); if (length != null) { @@ -72,7 +74,7 @@ if (constantValue != null && constantValue.isNonNull()) { Integer constantLength = constantReflection.lookupArrayLength(constantValue); if (constantLength != null) { - return ConstantNode.forInt(constantLength, array.graph()); + return ConstantNode.forInt(constantLength, graph); } } }
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/type/StampFactory.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/type/StampFactory.java Mon Nov 04 12:20:17 2013 +0100 @@ -160,6 +160,12 @@ return forFloat(kind, value.asDouble(), value.asDouble(), !Double.isNaN(value.asDouble())); case Illegal: return illegal(Kind.Illegal); + case Object: + if (value.isNull()) { + return alwaysNull(); + } else { + return objectNonNull(); + } default: throw new GraalInternalError("unexpected kind: %s", kind); }
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/util/GraphUtil.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/util/GraphUtil.java Mon Nov 04 12:20:17 2013 +0100 @@ -123,7 +123,7 @@ } public static void killWithUnusedFloatingInputs(Node node) { - if (node.recordsUsages()) { + if (node.recordsUsages() && !node.isExternal()) { List<Node> floatingInputs = node.inputs().filter(isFloatingNode()).snapshot(); node.safeDelete();
--- a/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/ConditionalEliminationPhase.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/ConditionalEliminationPhase.java Mon Nov 04 12:20:17 2013 +0100 @@ -510,7 +510,7 @@ PiNode piNode; if (isNull) { ConstantNode nullObject = ConstantNode.forObject(null, metaAccess, graph); - piNode = graph.unique(new PiNode(nullObject, StampFactory.forConstant(nullObject.value, metaAccess), replacementAnchor.asNode())); + piNode = graph.unique(new PiNode(nullObject, StampFactory.forConstant(nullObject.getValue(), metaAccess), replacementAnchor.asNode())); } else { piNode = graph.unique(new PiNode(object, StampFactory.declared(type, nonNull), replacementAnchor.asNode())); }
--- a/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/DeadCodeEliminationPhase.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/DeadCodeEliminationPhase.java Mon Nov 04 12:20:17 2013 +0100 @@ -112,13 +112,17 @@ for (Node node : graph.getNodes()) { if (flood.isMarked(node)) { for (Node input : node.inputs()) { - flood.add(input); + if (!input.isExternal()) { + flood.add(input); + } } } } for (Node current : flood) { for (Node input : current.inputs()) { - flood.add(input); + if (!input.isExternal()) { + flood.add(input); + } } } }
--- a/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/InliningUtil.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/InliningUtil.java Mon Nov 04 12:20:17 2013 +0100 @@ -39,6 +39,7 @@ import com.oracle.graal.debug.*; import com.oracle.graal.graph.*; import com.oracle.graal.graph.Graph.DuplicationReplacement; +import com.oracle.graal.graph.Node.ValueNumberable; import com.oracle.graal.graph.Node.Verbosity; import com.oracle.graal.nodes.*; import com.oracle.graal.nodes.calc.*; @@ -487,7 +488,7 @@ ConstantNode typeHub = ConstantNode.forConstant(type.getEncoding(Representation.ObjectHub), metaAccess, graph); LoadHubNode receiverHub = graph.unique(new LoadHubNode(nonNullReceiver, typeHub.kind(), null)); - CompareNode typeCheck = CompareNode.createCompareNode(Condition.EQ, receiverHub, typeHub); + CompareNode typeCheck = CompareNode.createCompareNode(graph, Condition.EQ, receiverHub, typeHub); FixedGuardNode guard = graph.add(new FixedGuardNode(typeCheck, DeoptimizationReason.TypeCheckedInliningViolated, DeoptimizationAction.InvalidateReprofile)); assert invoke.predecessor() != null; @@ -803,7 +804,7 @@ FixedNode lastSucc = successors[concretes.size()]; for (int i = concretes.size() - 1; i >= 0; --i) { LoadMethodNode method = graph.add(new LoadMethodNode(concretes.get(i), hub, constantMethods[i].kind())); - CompareNode methodCheck = CompareNode.createCompareNode(Condition.EQ, method, constantMethods[i]); + CompareNode methodCheck = CompareNode.createCompareNode(graph, Condition.EQ, method, constantMethods[i]); IfNode ifNode = graph.add(new IfNode(methodCheck, successors[i], lastSucc, probability[i])); method.setNext(ifNode); lastSucc = method; @@ -1433,7 +1434,12 @@ if (returnNode.result() instanceof LocalNode) { returnValue = localReplacement.replacement(returnNode.result()); } else if (returnNode.result() != null) { - returnValue = duplicates.get(returnNode.result()); + returnValue = returnNode.result(); + if (!returnValue.isExternal()) { + returnValue = duplicates.get(returnValue); + } else if (returnValue instanceof ValueNumberable) { + returnValue = graph.uniqueWithoutAdd(returnValue); + } } invoke.asNode().replaceAtUsages(returnValue); Node returnDuplicate = duplicates.get(returnNode);
--- a/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/LoweringPhase.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/LoweringPhase.java Mon Nov 04 12:20:17 2013 +0100 @@ -98,7 +98,7 @@ // Short cut creation of null check guard if the object is known to be non-null. return null; } - StructuredGraph graph = object.graph(); + StructuredGraph graph = guardedNode.asNode().graph(); if (graph.getGuardsStage().ordinal() > GuardsStage.FLOATING_GUARDS.ordinal()) { NullCheckNode nullCheck = graph.add(new NullCheckNode(object)); graph.addBeforeFixed((FixedNode) guardedNode, nullCheck);
--- a/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/TailDuplicationPhase.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/TailDuplicationPhase.java Mon Nov 04 12:20:17 2013 +0100 @@ -364,7 +364,7 @@ // stop iterating: fixed nodes within the given set are traversal roots // anyway, and all other // fixed nodes are known to be outside. - } else if (!aboveBound.isMarked(node)) { + } else if (!node.isExternal() && !aboveBound.isMarked(node)) { worklist.add(node); aboveBound.mark(node); } @@ -377,7 +377,9 @@ while (!worklist.isEmpty()) { Node current = worklist.remove(); for (Node input : current.inputs()) { - aboveClosure.apply(current, input); + if (!input.isExternal()) { + aboveClosure.apply(current, input); + } } }
--- a/graal/com.oracle.graal.phases/src/com/oracle/graal/phases/schedule/SchedulePhase.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.phases/src/com/oracle/graal/phases/schedule/SchedulePhase.java Mon Nov 04 12:20:17 2013 +0100 @@ -672,6 +672,10 @@ * Determines the earliest block in which the given node can be scheduled. */ private Block earliestBlock(Node node) { + if (node.isExternal()) { + return cfg.getStartBlock(); + } + Block earliest = cfg.getNodeToBlock().get(node); if (earliest != null) { return earliest; @@ -987,7 +991,7 @@ for (Node input : state.inputs()) { if (input instanceof VirtualState) { addUnscheduledToLatestSorting(b, (VirtualState) input, sortedInstructions, visited, reads, beforeLastLocation); - } else { + } else if (!input.isExternal()) { addToLatestSorting(b, (ScheduledNode) input, sortedInstructions, visited, reads, beforeLastLocation); } } @@ -1004,8 +1008,9 @@ if (input instanceof FrameState) { assert state == null; state = (FrameState) input; - } else { + } else if (!input.isExternal()) { addToLatestSorting(b, (ScheduledNode) input, sortedInstructions, visited, reads, beforeLastLocation); + } } List<FloatingNode> inputs = phantomInputs.get(i);
--- a/graal/com.oracle.graal.phases/src/com/oracle/graal/phases/verify/VerifyUsageWithEquals.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.phases/src/com/oracle/graal/phases/verify/VerifyUsageWithEquals.java Mon Nov 04 12:20:17 2013 +0100 @@ -72,7 +72,7 @@ for (ObjectEqualsNode cn : graph.getNodes().filter(ObjectEqualsNode.class)) { if (!isEqualsMethod(graph)) { // bail out if we compare an object of type klass with == or != (except null checks) - assert !(checkUsage(cn.x(), cn.y(), context.getMetaAccess()) && checkUsage(cn.y(), cn.x(), context.getMetaAccess())) : "Verifcation of " + klass.getName() + + assert !(checkUsage(cn.x(), cn.y(), context.getMetaAccess()) && checkUsage(cn.y(), cn.x(), context.getMetaAccess())) : "Verification of " + klass.getName() + " usage failed: Comparing " + cn.x() + " and " + cn.y() + " in " + graph.method() + " must use .equals() for object equality, not '==' or '!='"; } }
--- a/graal/com.oracle.graal.printer/src/com/oracle/graal/printer/BinaryGraphPrinter.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.printer/src/com/oracle/graal/printer/BinaryGraphPrinter.java Mon Nov 04 12:20:17 2013 +0100 @@ -137,8 +137,9 @@ ControlFlowGraph cfg = schedule == null ? null : schedule.getCFG(); BlockMap<List<ScheduledNode>> blockToNodes = schedule == null ? null : schedule.getBlockToNodesMap(); Block[] blocks = cfg == null ? null : cfg.getBlocks(); - writeNodes(graph); - writeBlocks(blocks, blockToNodes); + Map<Node, Integer> externalNodes = new HashMap<>(); + writeNodes(graph, externalNodes); + writeBlocks(blocks, blockToNodes, externalNodes); } private void flush() throws IOException { @@ -388,7 +389,17 @@ } @SuppressWarnings("deprecation") - private void writeNodes(Graph graph) throws IOException { + private static int getNodeId(Node node, Map<Node, Integer> externalNodes) { + if (!node.isExternal()) { + return node.getId(); + } else { + Integer id = externalNodes.get(node); + assert id != null; + return id; + } + } + + private void writeNodes(Graph graph, Map<Node, Integer> externalNodes) throws IOException { NodesToDoubles probabilities = null; if (PrintGraphProbabilities.getValue()) { try { @@ -397,14 +408,26 @@ } } Map<Object, Object> props = new HashMap<>(); - writeInt(graph.getNodeCount()); + int firstExternalNodeId = graph.getNodeCount() + graph.getDeletedNodeCount(); + for (Node node : graph.getNodes()) { + for (Node input : node.inputs()) { + if (input.isExternal()) { + if (!externalNodes.containsKey(input)) { + externalNodes.put(input, externalNodes.size() + firstExternalNodeId); + } + } + } + } + + writeInt(graph.getNodeCount() + externalNodes.size()); + for (Node node : graph.getNodes()) { NodeClass nodeClass = node.getNodeClass(); node.getDebugProperties(props); if (probabilities != null && node instanceof FixedNode && probabilities.contains((FixedNode) node)) { props.put("probability", probabilities.get((FixedNode) node)); } - writeInt(node.getId()); + writeInt(getNodeId(node, externalNodes)); writePoolObject(nodeClass); writeByte(node.predecessor() == null ? 0 : 1); // properties @@ -415,20 +438,35 @@ writePropertyObject(entry.getValue()); } // inputs - writeEdges(node, nodeClass.getFirstLevelInputPositions()); + writeEdges(node, nodeClass.getFirstLevelInputPositions(), externalNodes); // successors - writeEdges(node, nodeClass.getFirstLevelSuccessorPositions()); + writeEdges(node, nodeClass.getFirstLevelSuccessorPositions(), externalNodes); props.clear(); } + for (Node node : externalNodes.keySet()) { + NodeClass nodeClass = node.getNodeClass(); + node.getDebugProperties(props); + writeInt(getNodeId(node, externalNodes)); + writePoolObject(nodeClass); + writeByte(0); + // properties + writeShort((char) props.size()); + for (Entry<Object, Object> entry : props.entrySet()) { + String key = entry.getKey().toString(); + writePoolObject(key); + writePropertyObject(entry.getValue()); + } + props.clear(); + } } - private void writeEdges(Node node, Collection<Position> positions) throws IOException { + private void writeEdges(Node node, Collection<Position> positions, Map<Node, Integer> externalNodes) throws IOException { NodeClass nodeClass = node.getNodeClass(); for (Position pos : positions) { if (pos.subIndex == NodeClass.NOT_ITERABLE) { Node edge = nodeClass.get(node, pos); - writeNodeRef(edge); + writeNodeRef(edge, externalNodes); } else { NodeList<?> list = nodeClass.getNodeList(node, pos); if (list == null) { @@ -438,24 +476,22 @@ assert listSize == ((char) listSize); writeShort((char) listSize); for (Node edge : list) { - writeNodeRef(edge); + writeNodeRef(edge, externalNodes); } } } } } - @SuppressWarnings("deprecation") - private void writeNodeRef(Node edge) throws IOException { + private void writeNodeRef(Node edge, Map<Node, Integer> externalNodes) throws IOException { if (edge != null) { - writeInt(edge.getId()); + writeInt(getNodeId(edge, externalNodes)); } else { writeInt(-1); } } - @SuppressWarnings("deprecation") - private void writeBlocks(Block[] blocks, BlockMap<List<ScheduledNode>> blockToNodes) throws IOException { + private void writeBlocks(Block[] blocks, BlockMap<List<ScheduledNode>> blockToNodes, Map<Node, Integer> externalNodes) throws IOException { if (blocks != null) { writeInt(blocks.length); for (Block block : blocks) { @@ -463,7 +499,7 @@ writeInt(block.getId()); writeInt(nodes.size()); for (Node node : nodes) { - writeInt(node.getId()); + writeInt(getNodeId(node, externalNodes)); } writeInt(block.getSuccessors().size()); for (Block sux : block.getSuccessors()) {
--- a/graal/com.oracle.graal.replacements/src/com/oracle/graal/replacements/InstanceOfSnippetsTemplates.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.replacements/src/com/oracle/graal/replacements/InstanceOfSnippetsTemplates.java Mon Nov 04 12:20:17 2013 +0100 @@ -141,7 +141,7 @@ } if (condition == null || condition.y() != testValue) { // Re-use previously generated condition if the trueValue for the test is the same - condition = createCompareNode(Condition.EQ, result, testValue); + condition = createCompareNode(result.graph(), Condition.EQ, result, testValue); } return condition; } @@ -152,13 +152,13 @@ * @param t the true value for the materialization * @param f the false value for the materialization */ - ValueNode asMaterialization(ValueNode t, ValueNode f) { + ValueNode asMaterialization(StructuredGraph graph, ValueNode t, ValueNode f) { assert isInitialized(); if (t == this.trueValue && f == this.falseValue) { // Can simply use the phi result if the same materialized values are expected. return result; } else { - return t.graph().unique(new ConditionalNode(asCondition(trueValue), t, f)); + return graph.unique(new ConditionalNode(asCondition(trueValue), t, f)); } } } @@ -230,7 +230,7 @@ @Override public void replaceUsingInstantiation() { - ValueNode newValue = instantiation.asMaterialization(trueValue, falseValue); + ValueNode newValue = instantiation.asMaterialization(usage.graph(), trueValue, falseValue); usage.replaceAtUsages(newValue); usage.clearInputs(); assert usage.usages().isEmpty();
--- a/graal/com.oracle.graal.replacements/src/com/oracle/graal/replacements/NodeIntrinsificationPhase.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.replacements/src/com/oracle/graal/replacements/NodeIntrinsificationPhase.java Mon Nov 04 12:20:17 2013 +0100 @@ -72,6 +72,7 @@ ResolvedJavaMethod target = methodCallTargetNode.targetMethod(); NodeIntrinsic intrinsic = target.getAnnotation(Node.NodeIntrinsic.class); ResolvedJavaType declaringClass = target.getDeclaringClass(); + StructuredGraph graph = methodCallTargetNode.graph(); if (intrinsic != null) { assert target.getAnnotation(Fold.class) == null; assert Modifier.isStatic(target.getModifiers()) : "node intrinsic must be static: " + target; @@ -86,10 +87,10 @@ // Create the new node instance. ResolvedJavaType c = getNodeClass(target, intrinsic); - Node newInstance = createNodeInstance(c, parameterTypes, methodCallTargetNode.invoke().asNode().stamp(), intrinsic.setStampFromReturnType(), nodeConstructorArguments); + Node newInstance = createNodeInstance(graph, c, parameterTypes, methodCallTargetNode.invoke().asNode().stamp(), intrinsic.setStampFromReturnType(), nodeConstructorArguments); // Replace the invoke with the new node. - newInstance = methodCallTargetNode.graph().addOrUnique(newInstance); + newInstance = graph.addOrUnique(newInstance); methodCallTargetNode.invoke().intrinsify(newInstance); // Clean up checkcast instructions inserted by javac if the return type is generic. @@ -194,12 +195,13 @@ return result; } - private Node createNodeInstance(ResolvedJavaType nodeClass, ResolvedJavaType[] parameterTypes, Stamp invokeStamp, boolean setStampFromReturnType, Constant[] nodeConstructorArguments) { + private Node createNodeInstance(StructuredGraph graph, ResolvedJavaType nodeClass, ResolvedJavaType[] parameterTypes, Stamp invokeStamp, boolean setStampFromReturnType, + Constant[] nodeConstructorArguments) { ResolvedJavaMethod constructor = null; Constant[] arguments = null; for (ResolvedJavaMethod c : nodeClass.getDeclaredConstructors()) { - Constant[] match = match(c, parameterTypes, nodeConstructorArguments); + Constant[] match = match(graph, c, parameterTypes, nodeConstructorArguments); if (match != null) { if (constructor == null) { @@ -246,7 +248,7 @@ return false; } - private Constant[] match(ResolvedJavaMethod c, ResolvedJavaType[] parameterTypes, Constant[] nodeConstructorArguments) { + private Constant[] match(StructuredGraph graph, ResolvedJavaMethod c, ResolvedJavaType[] parameterTypes, Constant[] nodeConstructorArguments) { Constant[] arguments = null; Constant[] injected = null; @@ -257,6 +259,8 @@ injected = injected == null ? new Constant[1] : Arrays.copyOf(injected, injected.length + 1); if (signature[i].equals(metaAccess.lookupJavaType(MetaAccessProvider.class))) { injected[injected.length - 1] = Constant.forObject(metaAccess); + } else if (signature[i].equals(metaAccess.lookupJavaType(StructuredGraph.class))) { + injected[injected.length - 1] = Constant.forObject(graph); } else if (signature[i].equals(metaAccess.lookupJavaType(ForeignCallsProvider.class))) { injected[injected.length - 1] = Constant.forObject(providers.getForeignCalls()); } else {
--- a/graal/com.oracle.graal.truffle/src/com/oracle/graal/truffle/FrameWithoutBoxing.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.truffle/src/com/oracle/graal/truffle/FrameWithoutBoxing.java Mon Nov 04 12:20:17 2013 +0100 @@ -251,7 +251,7 @@ resize(); } byte tag = this.getTags()[slotIndex]; - if (accessKind == FrameSlotKind.Object ? tag != 0 : tag != accessKind.ordinal()) { + if (tag != accessKind.ordinal()) { CompilerDirectives.transferToInterpreter(); if (slot.getKind() == accessKind || tag == 0) { descriptor.getTypeConversion().updateFrameSlot(this, slot, getValue(slot)); @@ -284,7 +284,7 @@ } else if (tag == FrameSlotKind.Float.ordinal()) { return getFloatUnsafe(slot); } else { - assert tag == FrameSlotKind.Object.ordinal() || tag == FrameSlotKind.Illegal.ordinal(); + assert tag == FrameSlotKind.Object.ordinal(); return getObjectUnsafe(slot); } } @@ -300,13 +300,47 @@ } } - @Override - public boolean isInitialized(FrameSlot slot) { + private byte getTag(FrameSlot slot) { int slotIndex = slot.getIndex(); if (slotIndex >= getTags().length) { CompilerDirectives.transferToInterpreter(); resize(); } - return getTags()[slotIndex] != 0; + return getTags()[slotIndex]; + } + + @Override + public boolean isObject(FrameSlot slot) { + return getTag(slot) == FrameSlotKind.Object.ordinal(); + } + + @Override + public boolean isByte(FrameSlot slot) { + return getTag(slot) == FrameSlotKind.Byte.ordinal(); + } + + @Override + public boolean isBoolean(FrameSlot slot) { + return getTag(slot) == FrameSlotKind.Boolean.ordinal(); + } + + @Override + public boolean isInt(FrameSlot slot) { + return getTag(slot) == FrameSlotKind.Int.ordinal(); + } + + @Override + public boolean isLong(FrameSlot slot) { + return getTag(slot) == FrameSlotKind.Long.ordinal(); + } + + @Override + public boolean isFloat(FrameSlot slot) { + return getTag(slot) == FrameSlotKind.Float.ordinal(); + } + + @Override + public boolean isDouble(FrameSlot slot) { + return getTag(slot) == FrameSlotKind.Double.ordinal(); } }
--- a/graal/com.oracle.graal.truffle/src/com/oracle/graal/truffle/GraalTruffleRuntime.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.truffle/src/com/oracle/graal/truffle/GraalTruffleRuntime.java Mon Nov 04 12:20:17 2013 +0100 @@ -93,7 +93,7 @@ @Override public MaterializedFrame createMaterializedFrame(Arguments arguments) { - return createMaterializedFrame(arguments); + return createMaterializedFrame(arguments, new FrameDescriptor()); } @Override
--- a/graal/com.oracle.graal.truffle/src/com/oracle/graal/truffle/nodes/typesystem/CustomizedUnsafeLoadMacroNode.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.truffle/src/com/oracle/graal/truffle/nodes/typesystem/CustomizedUnsafeLoadMacroNode.java Mon Nov 04 12:20:17 2013 +0100 @@ -61,10 +61,9 @@ } else { locationIdentity = ObjectLocationIdentity.create(locationIdentityObject); } - Node result = graph().add( - new UnsafeLoadNode(objectArgument, offsetArgument, this.getTargetMethod().getSignature().getReturnKind(), locationIdentity, CompareNode.createCompareNode(Condition.EQ, - conditionArgument, ConstantNode.forBoolean(true, graph())))); - + CompareNode compare = CompareNode.createCompareNode(graph(), Condition.EQ, conditionArgument, ConstantNode.forBoolean(true, graph())); + Kind returnKind = this.getTargetMethod().getSignature().getReturnKind(); + Node result = graph().add(new UnsafeLoadNode(objectArgument, offsetArgument, returnKind, locationIdentity, compare)); return result; } return this;
--- a/graal/com.oracle.graal.truffle/src/com/oracle/graal/truffle/nodes/typesystem/UnsafeTypeCastMacroNode.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.truffle/src/com/oracle/graal/truffle/nodes/typesystem/UnsafeTypeCastMacroNode.java Mon Nov 04 12:20:17 2013 +0100 @@ -58,7 +58,7 @@ return objectArgument; } ResolvedJavaType lookupJavaType = tool.getMetaAccess().lookupJavaType(c); - ConditionAnchorNode valueAnchorNode = graph().add(new ConditionAnchorNode(CompareNode.createCompareNode(Condition.EQ, conditionArgument, ConstantNode.forBoolean(true, graph())))); + ConditionAnchorNode valueAnchorNode = graph().add(new ConditionAnchorNode(CompareNode.createCompareNode(graph(), Condition.EQ, conditionArgument, ConstantNode.forBoolean(true, graph())))); UnsafeCastNode piCast = graph().unique(new UnsafeCastNode(objectArgument, StampFactory.declaredNonNull(lookupJavaType), valueAnchorNode)); this.replaceAtUsages(piCast); return valueAnchorNode;
--- a/graal/com.oracle.graal.virtual/src/com/oracle/graal/virtual/phases/ea/EffectsClosure.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.virtual/src/com/oracle/graal/virtual/phases/ea/EffectsClosure.java Mon Nov 04 12:20:17 2013 +0100 @@ -44,6 +44,7 @@ protected final NodeMap<ValueNode> aliases; protected final BlockMap<GraphEffectList> blockEffects; private final IdentityHashMap<Loop, GraphEffectList> loopMergeEffects = new IdentityHashMap<>(); + private final IdentityHashMap<LoopBeginNode, BlockT> loopEntryStates = new IdentityHashMap<>(); private boolean changed; @@ -122,6 +123,10 @@ FixedWithNextNode lastFixedNode = null; for (Node node : schedule.getBlockToNodesMap().get(block)) { aliases.set(node, null); + if (node instanceof LoopExitNode) { + LoopExitNode loopExit = (LoopExitNode) node; + processLoopExit(loopExit, loopEntryStates.get(loopExit.loopBegin()), state, blockEffects.get(block)); + } changed |= processNode(node, state, effects, lastFixedNode); if (node instanceof FixedWithNextNode) { lastFixedNode = (FixedWithNextNode) node; @@ -169,10 +174,9 @@ loopMergeEffects.put(loop, mergeProcessor.afterMergeEffects); assert info.exitStates.size() == loop.exits.size(); + loopEntryStates.put(loop.loopBegin(), loopEntryState); for (int i = 0; i < loop.exits.size(); i++) { - BlockT exitState = info.exitStates.get(i); - assert exitState != null : "no loop exit state at " + loop.exits.get(i) + " / " + loop.header; - processLoopExit((LoopExitNode) loop.exits.get(i).getBeginNode(), loopEntryState, exitState, blockEffects.get(loop.exits.get(i))); + assert info.exitStates.get(i) != null : "no loop exit state at " + loop.exits.get(i) + " / " + loop.header; } return info.exitStates; @@ -227,7 +231,7 @@ public ValueNode getScalarAlias(ValueNode node) { assert !(node instanceof VirtualObjectNode); - if (node == null || !node.isAlive() || aliases.isNew(node)) { + if (node == null || !node.isAlive() || node.isExternal() || aliases.isNew(node)) { return node; } ValueNode result = aliases.get(node);
--- a/graal/com.oracle.graal.virtual/src/com/oracle/graal/virtual/phases/ea/PartialEscapeClosure.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.virtual/src/com/oracle/graal/virtual/phases/ea/PartialEscapeClosure.java Mon Nov 04 12:20:17 2013 +0100 @@ -245,7 +245,7 @@ if (obj.isVirtual()) { for (int i = 0; i < obj.getEntries().length; i++) { ValueNode value = obj.getEntry(i); - if (!(value instanceof VirtualObjectNode)) { + if (!(value instanceof VirtualObjectNode || value.isConstant())) { if (exitNode.loopBegin().isPhiAtMerge(value) || initialObj == null || !initialObj.isVirtual() || initialObj.getEntry(i) != value) { ProxyNode proxy = new ProxyNode(value, exitNode, PhiType.Value, null); obj.setEntry(i, proxy); @@ -616,7 +616,7 @@ } public ObjectState getObjectState(PartialEscapeBlockState<?> state, ValueNode value) { - if (value == null) { + if (value == null || value.isExternal()) { return null; } if (value.isAlive() && !aliases.isNew(value)) {
--- a/graal/com.oracle.graal.virtual/src/com/oracle/graal/virtual/phases/ea/VirtualUtil.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.virtual/src/com/oracle/graal/virtual/phases/ea/VirtualUtil.java Mon Nov 04 12:20:17 2013 +0100 @@ -69,18 +69,22 @@ for (Node node : graph.getNodes()) { if (flood.isMarked(node)) { for (Node input : node.inputs()) { - flood.add(input); - if (!path.containsKey(input)) { - path.put(input, node); + if (!input.isExternal()) { + flood.add(input); + if (!path.containsKey(input)) { + path.put(input, node); + } } } } } for (Node current : flood) { for (Node input : current.inputs()) { - flood.add(input); - if (!path.containsKey(input)) { - path.put(input, current); + if (!input.isExternal()) { + flood.add(input); + if (!path.containsKey(input)) { + path.put(input, current); + } } } }
--- a/graal/com.oracle.graal.word/src/com/oracle/graal/word/phases/WordTypeRewriterPhase.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.graal.word/src/com/oracle/graal/word/phases/WordTypeRewriterPhase.java Mon Nov 04 12:20:17 2013 +0100 @@ -114,8 +114,8 @@ if (isWord(node)) { if (node.isConstant()) { ConstantNode oldConstant = (ConstantNode) node; - assert oldConstant.value.getKind() == Kind.Object; - WordBase value = (WordBase) oldConstant.value.asObject(); + assert oldConstant.getValue().getKind() == Kind.Object; + WordBase value = (WordBase) oldConstant.getValue().asObject(); ConstantNode newConstant = ConstantNode.forIntegerKind(wordKind, value.rawValue(), node.graph()); graph.replaceFloating(oldConstant, newConstant);
--- a/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/frame/Frame.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/frame/Frame.java Mon Nov 04 12:20:17 2013 +0100 @@ -191,11 +191,37 @@ MaterializedFrame materialize(); /** - * To check whether the given {@link FrameSlot} has been initialized or not. An initialized slot - * has previously been read or modified. - * - * @param slot the slot - * @return true if the slot is uninitialized. + * Check whether the given {@link FrameSlot} is of type object. + */ + boolean isObject(FrameSlot slot); + + /** + * Check whether the given {@link FrameSlot} is of type byte. + */ + boolean isByte(FrameSlot slot); + + /** + * Check whether the given {@link FrameSlot} is of type boolean. + */ + boolean isBoolean(FrameSlot slot); + + /** + * Check whether the given {@link FrameSlot} is of type int. */ - boolean isInitialized(FrameSlot slot); + boolean isInt(FrameSlot slot); + + /** + * Check whether the given {@link FrameSlot} is of type long. + */ + boolean isLong(FrameSlot slot); + + /** + * Check whether the given {@link FrameSlot} is of type float. + */ + boolean isFloat(FrameSlot slot); + + /** + * Check whether the given {@link FrameSlot} is of type double. + */ + boolean isDouble(FrameSlot slot); }
--- a/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/frame/FrameDescriptor.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/frame/FrameDescriptor.java Mon Nov 04 12:20:17 2013 +0100 @@ -36,8 +36,8 @@ public final class FrameDescriptor implements Cloneable { private final FrameTypeConversion typeConversion; - private final ArrayList<FrameSlotImpl> slots; - private final HashMap<Object, FrameSlotImpl> identifierToSlotMap; + private final ArrayList<FrameSlot> slots; + private final HashMap<Object, FrameSlot> identifierToSlotMap; private Assumption version; private HashMap<Object, Assumption> identifierToNotInFrameAssumptionMap; @@ -58,7 +58,7 @@ public FrameSlot addFrameSlot(Object identifier, FrameSlotKind kind) { assert !identifierToSlotMap.containsKey(identifier); - FrameSlotImpl slot = new FrameSlotImpl(this, identifier, slots.size(), kind); + FrameSlot slot = new FrameSlot(this, identifier, slots.size(), kind); slots.add(slot); identifierToSlotMap.put(identifier, slot); updateVersion();
--- a/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/frame/FrameSlot.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/frame/FrameSlot.java Mon Nov 04 12:20:17 2013 +0100 @@ -24,18 +24,51 @@ */ package com.oracle.truffle.api.frame; +import com.oracle.truffle.api.*; + /** * A slot in a frame that can store a value of a given type. */ -public interface FrameSlot extends Cloneable { +public final class FrameSlot implements Cloneable { + + private final FrameDescriptor descriptor; + private final Object identifier; + private final int index; + @com.oracle.truffle.api.CompilerDirectives.CompilationFinal private FrameSlotKind kind; - Object getIdentifier(); + public FrameSlot(FrameDescriptor descriptor, Object identifier, int index, FrameSlotKind kind) { + this.descriptor = descriptor; + this.identifier = identifier; + this.index = index; + this.kind = kind; + } - int getIndex(); + public Object getIdentifier() { + return identifier; + } + + public int getIndex() { + return index; + } - FrameSlotKind getKind(); + public FrameSlotKind getKind() { + return kind; + } - void setKind(FrameSlotKind kind); + public void setKind(final FrameSlotKind kind) { + if (this.kind != kind) { + CompilerDirectives.transferToInterpreter(); + this.kind = kind; + this.descriptor.updateVersion(); + } + } - FrameDescriptor getFrameDescriptor(); + @Override + public String toString() { + return "[" + index + "," + identifier + "," + kind + "]"; + } + + public FrameDescriptor getFrameDescriptor() { + return this.descriptor; + } }
--- a/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/frame/FrameSlotImpl.java Mon Nov 04 12:18:58 2013 +0100 +++ /dev/null Thu Jan 01 00:00:00 1970 +0000 @@ -1,69 +0,0 @@ -/* - * Copyright (c) 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.frame; - -public final class FrameSlotImpl implements FrameSlot { - - private final FrameDescriptor descriptor; - private final Object identifier; - private final int index; - @com.oracle.truffle.api.CompilerDirectives.CompilationFinal private FrameSlotKind kind; - - public FrameSlotImpl(FrameDescriptor descriptor, Object identifier, int index, FrameSlotKind kind) { - this.descriptor = descriptor; - this.identifier = identifier; - this.index = index; - this.kind = kind; - } - - public Object getIdentifier() { - return identifier; - } - - public int getIndex() { - return index; - } - - public FrameSlotKind getKind() { - return kind; - } - - public void setKind(final FrameSlotKind kind) { - if (this.kind != kind) { - this.kind = kind; - this.descriptor.updateVersion(); - } - } - - @Override - public String toString() { - return "[" + index + "," + identifier + "," + kind + "]"; - } - - @Override - public FrameDescriptor getFrameDescriptor() { - return this.descriptor; - } -}
--- a/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/frame/FrameUtil.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/frame/FrameUtil.java Mon Nov 04 12:20:17 2013 +0100 @@ -25,88 +25,115 @@ package com.oracle.truffle.api.frame; public final class FrameUtil { - /** - * Write access to a local variable of type {@link Object}. - * - * Sets the frame slot type to {@link Object} if it isn't already. + * Read a frame slot that is guaranteed to be of the desired kind (either previously checked by + * a guard or statically known). * - * @param slot the slot of the local variable - * @param value the new value of the local variable + * @param frameSlot the slot of the variable + * @throws IllegalStateException if the slot kind does not match + * @see Frame#getObject(FrameSlot) */ - public static void setObjectSafe(Frame frame, FrameSlot slot, Object value) { - frame.setObject(slot, value); + public static Object getObjectSafe(Frame frame, FrameSlot frameSlot) { + try { + return frame.getObject(frameSlot); + } catch (FrameSlotTypeException e) { + throw new IllegalStateException(); + } } /** - * Write access to a local variable of type {@code byte}. - * - * Sets the frame slot type to {@code byte} if it isn't already. + * Read a frame slot that is guaranteed to be of the desired kind (either previously checked by + * a guard or statically known). * - * @param slot the slot of the local variable - * @param value the new value of the local variable + * @param frameSlot the slot of the variable + * @throws IllegalStateException if the slot kind does not match + * @see Frame#getByte(FrameSlot) */ - public static void setByteSafe(Frame frame, FrameSlot slot, byte value) { - frame.setByte(slot, value); + public static byte getByteSafe(Frame frame, FrameSlot frameSlot) { + try { + return frame.getByte(frameSlot); + } catch (FrameSlotTypeException e) { + throw new IllegalStateException(); + } } /** - * Write access to a local variable of type {@code boolean}. - * - * Sets the frame slot type to {@code boolean} if it isn't already. + * Read a frame slot that is guaranteed to be of the desired kind (either previously checked by + * a guard or statically known). * - * @param slot the slot of the local variable - * @param value the new value of the local variable + * @param frameSlot the slot of the variable + * @throws IllegalStateException if the slot kind does not match + * @see Frame#getBoolean(FrameSlot) */ - public static void setBooleanSafe(Frame frame, FrameSlot slot, boolean value) { - frame.setBoolean(slot, value); + public static boolean getBooleanSafe(Frame frame, FrameSlot frameSlot) { + try { + return frame.getBoolean(frameSlot); + } catch (FrameSlotTypeException e) { + throw new IllegalStateException(); + } } /** - * Write access to a local variable of type {@code int}. - * - * Sets the frame slot type to {@code int} if it isn't already. + * Read a frame slot that is guaranteed to be of the desired kind (either previously checked by + * a guard or statically known). * - * @param slot the slot of the local variable - * @param value the new value of the local variable + * @param frameSlot the slot of the variable + * @throws IllegalStateException if the slot kind does not match + * @see Frame#getInt(FrameSlot) */ - public static void setIntSafe(Frame frame, FrameSlot slot, int value) { - frame.setInt(slot, value); + public static int getIntSafe(Frame frame, FrameSlot frameSlot) { + try { + return frame.getInt(frameSlot); + } catch (FrameSlotTypeException e) { + throw new IllegalStateException(); + } } /** - * Write access to a local variable of type {@code long}. - * - * Sets the frame slot type to {@code long} if it isn't already. + * Read a frame slot that is guaranteed to be of the desired kind (either previously checked by + * a guard or statically known). * - * @param slot the slot of the local variable - * @param value the new value of the local variable + * @param frameSlot the slot of the variable + * @throws IllegalStateException if the slot kind does not match + * @see Frame#getLong(FrameSlot) */ - public static void setLongSafe(Frame frame, FrameSlot slot, long value) { - frame.setLong(slot, value); + public static long getLongSafe(Frame frame, FrameSlot frameSlot) { + try { + return frame.getLong(frameSlot); + } catch (FrameSlotTypeException e) { + throw new IllegalStateException(); + } } /** - * Write access to a local variable of type {@code float}. - * - * Sets the frame slot type to {@code float} if it isn't already. + * Read a frame slot that is guaranteed to be of the desired kind (either previously checked by + * a guard or statically known). * - * @param slot the slot of the local variable - * @param value the new value of the local variable + * @param frameSlot the slot of the variable + * @throws IllegalStateException if the slot kind does not match + * @see Frame#getDouble(FrameSlot) */ - public static void setFloatSafe(Frame frame, FrameSlot slot, float value) { - frame.setFloat(slot, value); + public static double getDoubleSafe(Frame frame, FrameSlot frameSlot) { + try { + return frame.getDouble(frameSlot); + } catch (FrameSlotTypeException e) { + throw new IllegalStateException(); + } } /** - * Write access to a local variable of type {@code double}. - * - * Sets the frame slot type to {@code double} if it isn't already. + * Read a frame slot that is guaranteed to be of the desired kind (either previously checked by + * a guard or statically known). * - * @param slot the slot of the local variable - * @param value the new value of the local variable + * @param frameSlot the slot of the variable + * @throws IllegalStateException if the slot kind does not match + * @see Frame#getFloat(FrameSlot) */ - public static void setDoubleSafe(Frame frame, FrameSlot slot, double value) { - frame.setDouble(slot, value); + public static float getFloatSafe(Frame frame, FrameSlot frameSlot) { + try { + return frame.getFloat(frameSlot); + } catch (FrameSlotTypeException e) { + throw new IllegalStateException(); + } } }
--- a/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/frame/NativeFrame.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/frame/NativeFrame.java Mon Nov 04 12:20:17 2013 +0100 @@ -48,77 +48,77 @@ @Override public Object getObject(FrameSlot slot) { - throw new UnsupportedOperationException("native frame"); + throw unsupportedInNativeFrame(); } @Override public void setObject(FrameSlot slot, Object value) { - throw new UnsupportedOperationException("native frame"); + throw unsupportedInNativeFrame(); } @Override public byte getByte(FrameSlot slot) { - throw new UnsupportedOperationException("native frame"); + throw unsupportedInNativeFrame(); } @Override public void setByte(FrameSlot slot, byte value) { - throw new UnsupportedOperationException("native frame"); + throw unsupportedInNativeFrame(); } @Override public boolean getBoolean(FrameSlot slot) { - throw new UnsupportedOperationException("native frame"); + throw unsupportedInNativeFrame(); } @Override public void setBoolean(FrameSlot slot, boolean value) { - throw new UnsupportedOperationException("native frame"); + throw unsupportedInNativeFrame(); } @Override public int getInt(FrameSlot slot) { - throw new UnsupportedOperationException("native frame"); + throw unsupportedInNativeFrame(); } @Override public void setInt(FrameSlot slot, int value) { - throw new UnsupportedOperationException("native frame"); + throw unsupportedInNativeFrame(); } @Override public long getLong(FrameSlot slot) { - throw new UnsupportedOperationException("native frame"); + throw unsupportedInNativeFrame(); } @Override public void setLong(FrameSlot slot, long value) { - throw new UnsupportedOperationException("native frame"); + throw unsupportedInNativeFrame(); } @Override public float getFloat(FrameSlot slot) { - throw new UnsupportedOperationException("native frame"); + throw unsupportedInNativeFrame(); } @Override public void setFloat(FrameSlot slot, float value) { - throw new UnsupportedOperationException("native frame"); + throw unsupportedInNativeFrame(); } @Override public double getDouble(FrameSlot slot) { - throw new UnsupportedOperationException("native frame"); + throw unsupportedInNativeFrame(); } @Override public void setDouble(FrameSlot slot, double value) { - throw new UnsupportedOperationException("native frame"); + throw unsupportedInNativeFrame(); } @Override public Object getValue(FrameSlot slot) { - throw new UnsupportedOperationException("native frame"); + throw unsupportedInNativeFrame(); } @Override @@ -133,7 +133,7 @@ @Override public MaterializedFrame materialize() { - throw new UnsupportedOperationException("native frame"); + throw unsupportedInNativeFrame(); } @Override @@ -143,11 +143,45 @@ @Override public FrameDescriptor getFrameDescriptor() { - throw new UnsupportedOperationException("native frame"); + throw unsupportedInNativeFrame(); + } + + @Override + public boolean isObject(FrameSlot slot) { + throw unsupportedInNativeFrame(); + } + + @Override + public boolean isByte(FrameSlot slot) { + throw unsupportedInNativeFrame(); + } + + @Override + public boolean isBoolean(FrameSlot slot) { + throw unsupportedInNativeFrame(); } @Override - public boolean isInitialized(FrameSlot slot) { + public boolean isInt(FrameSlot slot) { + throw unsupportedInNativeFrame(); + } + + @Override + public boolean isLong(FrameSlot slot) { + throw unsupportedInNativeFrame(); + } + + @Override + public boolean isFloat(FrameSlot slot) { + throw unsupportedInNativeFrame(); + } + + @Override + public boolean isDouble(FrameSlot slot) { + throw unsupportedInNativeFrame(); + } + + private static UnsupportedOperationException unsupportedInNativeFrame() { throw new UnsupportedOperationException("native frame"); } }
--- a/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/impl/DefaultMaterializedFrame.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/impl/DefaultMaterializedFrame.java Mon Nov 04 12:20:17 2013 +0100 @@ -136,7 +136,37 @@ } @Override - public boolean isInitialized(FrameSlot slot) { - return wrapped.isInitialized(slot); + public boolean isObject(FrameSlot slot) { + return wrapped.isObject(slot); + } + + @Override + public boolean isByte(FrameSlot slot) { + return wrapped.isByte(slot); + } + + @Override + public boolean isBoolean(FrameSlot slot) { + return wrapped.isBoolean(slot); + } + + @Override + public boolean isInt(FrameSlot slot) { + return wrapped.isInt(slot); + } + + @Override + public boolean isLong(FrameSlot slot) { + return wrapped.isLong(slot); + } + + @Override + public boolean isFloat(FrameSlot slot) { + return wrapped.isFloat(slot); + } + + @Override + public boolean isDouble(FrameSlot slot) { + return wrapped.isDouble(slot); } }
--- a/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/impl/DefaultVirtualFrame.java Mon Nov 04 12:18:58 2013 +0100 +++ b/graal/com.oracle.truffle.api/src/com/oracle/truffle/api/impl/DefaultVirtualFrame.java Mon Nov 04 12:20:17 2013 +0100 @@ -200,12 +200,46 @@ } } - @Override - public boolean isInitialized(FrameSlot slot) { + private byte getTag(FrameSlot slot) { int slotIndex = slot.getIndex(); if (slotIndex >= tags.length) { resize(); } - return tags[slotIndex] != 0; + return tags[slotIndex]; + } + + @Override + public boolean isObject(FrameSlot slot) { + return getTag(slot) == FrameSlotKind.Object.ordinal(); + } + + @Override + public boolean isByte(FrameSlot slot) { + return getTag(slot) == FrameSlotKind.Byte.ordinal(); + } + + @Override + public boolean isBoolean(FrameSlot slot) { + return getTag(slot) == FrameSlotKind.Boolean.ordinal(); + } + + @Override + public boolean isInt(FrameSlot slot) { + return getTag(slot) == FrameSlotKind.Int.ordinal(); + } + + @Override + public boolean isLong(FrameSlot slot) { + return getTag(slot) == FrameSlotKind.Long.ordinal(); + } + + @Override + public boolean isFloat(FrameSlot slot) { + return getTag(slot) == FrameSlotKind.Float.ordinal(); + } + + @Override + public boolean isDouble(FrameSlot slot) { + return getTag(slot) == FrameSlotKind.Double.ordinal(); } }
--- a/mx/commands.py Mon Nov 04 12:18:58 2013 +0100 +++ b/mx/commands.py Mon Nov 04 12:20:17 2013 +0100 @@ -821,7 +821,7 @@ def harness(projectscp, vmArgs): if not exists(javaClass) or getmtime(javaClass) < getmtime(javaSource): subprocess.check_call([mx.java().javac, '-cp', projectscp, '-d', mxdir, javaSource]) - if not isGraalEnabled(_get_vm()): + if _get_vm() != 'graal': prefixArgs = ['-esa', '-ea'] else: prefixArgs = ['-XX:-BootstrapGraal', '-esa', '-ea']
--- a/mx/sanitycheck.py Mon Nov 04 12:18:58 2013 +0100 +++ b/mx/sanitycheck.py Mon Nov 04 12:20:17 2013 +0100 @@ -277,8 +277,8 @@ args = ['-XX:+CompileTheWorld', '-Xbootclasspath/p:' + rtjar] - if commands.isGraalEnabled(vm): - args += ['-XX:+BootstrapGraal', '-G:-Debug'] + if vm == 'graal': + args += ['-XX:+BootstrapGraal'] if mode >= CTWMode.NoInline: if not commands.isGraalEnabled(vm): args.append('-XX:-Inline')
--- a/mxtool/mx.py Mon Nov 04 12:18:58 2013 +0100 +++ b/mxtool/mx.py Mon Nov 04 12:20:17 2013 +0100 @@ -328,15 +328,16 @@ def update_current_annotation_processors_file(self): aps = self.annotation_processors() outOfDate = False - currentApsFile = join(self.dir, '.currentAnnotationProcessors') - if exists(currentApsFile): + currentApsFile = join(self.suite.mxDir, 'currentAnnotationProcessors', self.name) + currentApsFileExists = exists(currentApsFile) + if currentApsFileExists: with open(currentApsFile) as fp: currentAps = [l.strip() for l in fp.readlines()] if currentAps != aps: outOfDate = True - else: - outOfDate = True - if outOfDate: + if outOfDate or not currentApsFileExists: + if not exists(dirname(currentApsFile)): + os.mkdir(dirname(currentApsFile)) with open(currentApsFile, 'w') as fp: for ap in aps: print >> fp, ap @@ -786,7 +787,7 @@ abort('Missing "suite=<name>" in ' + projectsFile) def _commands_name(self): - return 'mx_' + self.name.replace('-','_') + return 'mx_' + self.name.replace('-', '_') def _find_commands(self, name): commandsPath = join(self.mxDir, name + '.py') @@ -1988,7 +1989,7 @@ continue # Ensure that the output directories are clean - #prepareOutputDirs(p, True) + # prepareOutputDirs(p, True) built.add(p.name)
--- a/src/gpu/ptx/vm/gpu_ptx.cpp Mon Nov 04 12:18:58 2013 +0100 +++ b/src/gpu/ptx/vm/gpu_ptx.cpp Mon Nov 04 12:20:17 2013 +0100 @@ -49,7 +49,21 @@ 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; +gpu::Ptx::cuda_cu_mem_host_register_func_t gpu::Ptx::_cuda_cu_mem_host_register; +gpu::Ptx::cuda_cu_mem_host_get_device_pointer_func_t gpu::Ptx::_cuda_cu_mem_host_get_device_pointer; +gpu::Ptx::cuda_cu_mem_host_unregister_func_t gpu::Ptx::_cuda_cu_mem_host_unregister; +#define STRINGIFY(x) #x + +#define LOOKUP_CUDA_FUNCTION(name, alias) \ + _##alias = \ + CAST_TO_FN_PTR(alias##_func_t, os::dll_lookup(handle, STRINGIFY(name))); \ + if (_##alias == NULL) { \ + tty->print_cr("[CUDA] ***** Error: Failed to lookup %s", STRINGIFY(name)); \ + return 0; \ + } \ + +#define LOOKUP_CUDA_V2_FUNCTION(name, alias) LOOKUP_CUDA_FUNCTION(name##_v2, alias) /* * see http://en.wikipedia.org/wiki/CUDA#Supported_GPUs @@ -199,7 +213,7 @@ tty->print_cr("[CUDA] Failed to get GRAAL_CU_DEVICE_ATTRIBUTE_WARP_SIZE: %d", _cu_device); return 0; } - + status = _cuda_cu_device_get_attribute(&async_engines, GRAAL_CU_DEVICE_ATTRIBUTE_ASYNC_ENGINE_COUNT, _cu_device); @@ -234,7 +248,7 @@ tty->print_cr("[CUDA] Max threads per block: %d warp size: %d", max_threads_per_block, warp_size); } return (total); - + } void *gpu::Ptx::generate_kernel(unsigned char *code, int code_len, const char *name) { @@ -262,7 +276,7 @@ jit_option_values[2] = (void *)(size_t)jit_register_count; /* Create CUDA context to compile and execute the kernel */ - int status = _cuda_cu_ctx_create(&_device_context, 0, _cu_device); + int status = _cuda_cu_ctx_create(&_device_context, GRAAL_CU_CTX_MAP_HOST, _cu_device); if (status != GRAAL_CUDA_SUCCESS) { tty->print_cr("[CUDA] Failed to create CUDA context for device(%d): %d", _cu_device, status); @@ -443,9 +457,6 @@ tty->print_cr("[CUDA] TODO *** Unhandled return type: %d", return_type); } - // Copy all reference arguments from device to host memory. - ptxka.copyRefArgsFromDtoH(); - // Free device memory allocated for result status = gpu::Ptx::_cuda_cu_memfree(ptxka._dev_return_value); if (status != GRAAL_CUDA_SUCCESS) { @@ -487,40 +498,36 @@ void *handle = os::dll_load(cuda_library_name, buffer, STD_BUFFER_SIZE); free(buffer); if (handle != NULL) { - _cuda_cu_init = - 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_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_ctx_set_current = - CAST_TO_FN_PTR(cuda_cu_ctx_set_current_func_t, os::dll_lookup(handle, "cuCtxSetCurrent")); - _cuda_cu_device_get_count = - CAST_TO_FN_PTR(cuda_cu_device_get_count_func_t, os::dll_lookup(handle, "cuDeviceGetCount")); - _cuda_cu_device_get_name = - CAST_TO_FN_PTR(cuda_cu_device_get_name_func_t, os::dll_lookup(handle, "cuDeviceGetName")); - _cuda_cu_device_get = - CAST_TO_FN_PTR(cuda_cu_device_get_func_t, os::dll_lookup(handle, "cuDeviceGet")); - _cuda_cu_device_compute_capability = - CAST_TO_FN_PTR(cuda_cu_device_compute_capability_func_t, os::dll_lookup(handle, "cuDeviceComputeCapability")); - _cuda_cu_device_get_attribute = - CAST_TO_FN_PTR(cuda_cu_device_get_attribute_func_t, os::dll_lookup(handle, "cuDeviceGetAttribute")); - _cuda_cu_module_get_function = - CAST_TO_FN_PTR(cuda_cu_module_get_function_func_t, os::dll_lookup(handle, "cuModuleGetFunction")); - _cuda_cu_module_load_data_ex = - 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")); + LOOKUP_CUDA_FUNCTION(cuInit, cuda_cu_init); + LOOKUP_CUDA_FUNCTION(cuCtxSynchronize, cuda_cu_ctx_synchronize); + LOOKUP_CUDA_FUNCTION(cuCtxSetCurrent, cuda_cu_ctx_set_current); + LOOKUP_CUDA_FUNCTION(cuDeviceGetCount, cuda_cu_device_get_count); + LOOKUP_CUDA_FUNCTION(cuDeviceGetName, cuda_cu_device_get_name); + LOOKUP_CUDA_FUNCTION(cuDeviceGet, cuda_cu_device_get); + LOOKUP_CUDA_FUNCTION(cuDeviceComputeCapability, cuda_cu_device_compute_capability); + LOOKUP_CUDA_FUNCTION(cuDeviceGetAttribute, cuda_cu_device_get_attribute); + LOOKUP_CUDA_FUNCTION(cuModuleGetFunction, cuda_cu_module_get_function); + LOOKUP_CUDA_FUNCTION(cuModuleLoadDataEx, cuda_cu_module_load_data_ex); + LOOKUP_CUDA_FUNCTION(cuLaunchKernel, cuda_cu_launch_kernel); + LOOKUP_CUDA_FUNCTION(cuMemHostRegister, cuda_cu_mem_host_register); + LOOKUP_CUDA_FUNCTION(cuMemHostUnregister, cuda_cu_mem_host_unregister); +#if defined(__x86_64) || defined(AMD64) || defined(_M_AMD64) + LOOKUP_CUDA_V2_FUNCTION(cuCtxCreate, cuda_cu_ctx_create); + LOOKUP_CUDA_V2_FUNCTION(cuCtxDestroy, cuda_cu_ctx_destroy); + LOOKUP_CUDA_V2_FUNCTION(cuMemAlloc, cuda_cu_memalloc); + LOOKUP_CUDA_V2_FUNCTION(cuMemFree, cuda_cu_memfree); + LOOKUP_CUDA_V2_FUNCTION(cuMemcpyHtoD, cuda_cu_memcpy_htod); + LOOKUP_CUDA_V2_FUNCTION(cuMemcpyDtoH, cuda_cu_memcpy_dtoh); + LOOKUP_CUDA_V2_FUNCTION(cuMemHostGetDevicePointer, cuda_cu_mem_host_get_device_pointer); +#else + LOOKUP_CUDA_FUNCTION(cuCtxCreate, cuda_cu_ctx_create); + LOOKUP_CUDA_FUNCTION(cuCtxDestroy, cuda_cu_ctx_destroy); + LOOKUP_CUDA_FUNCTION(cuMemAlloc, cuda_cu_memalloc); + LOOKUP_CUDA_FUNCTION(cuMemFree, cuda_cu_memfree); + LOOKUP_CUDA_FUNCTION(cuMemcpyHtoD, cuda_cu_memcpy_htod); + LOOKUP_CUDA_FUNCTION(cuMemcpyDtoH, cuda_cu_memcpy_dtoh); + LOOKUP_CUDA_FUNCTION(cuMemHostGetDevicePointer, cuda_cu_mem_host_get_device_pointer); +#endif if (TraceGPUInteraction) { tty->print_cr("[CUDA] Success: library linkage");
--- a/src/gpu/ptx/vm/gpu_ptx.hpp Mon Nov 04 12:18:58 2013 +0100 +++ b/src/gpu/ptx/vm/gpu_ptx.hpp Mon Nov 04 12:20:17 2013 +0100 @@ -46,6 +46,13 @@ #define GRAAL_CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES 4 #define GRAAL_CUDA_ERROR_NO_BINARY_FOR_GPU 209 +/* + * Flags for cuMemHostRegister + */ + +#define GRAAL_CU_MEMHOSTREGISTER_PORTABLE 1 +#define GRAAL_CU_MEMHOSTREGISTER_DEVICEMAP 2 + /** * End of array terminator for the extra parameter to * ::cuLaunchKernel @@ -73,6 +80,12 @@ */ #define GRAAL_CU_LAUNCH_PARAM_BUFFER_SIZE ((void*) 0x02) +/* + * Context creation flags + */ + +#define GRAAL_CU_CTX_MAP_HOST 0x08 + class Ptx { friend class gpu; @@ -90,9 +103,11 @@ typedef unsigned int CUdeviceptr; #endif +typedef int CUdevice; /**< CUDA device */ + 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_create_func_t)(void*, unsigned int, CUdevice); typedef int (*cuda_cu_ctx_destroy_func_t)(void*); typedef int (*cuda_cu_ctx_synchronize_func_t)(void); typedef int (*cuda_cu_ctx_set_current_func_t)(void*); @@ -107,10 +122,13 @@ 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*, size_t); + typedef int (*cuda_cu_memalloc_func_t)(gpu::Ptx::CUdeviceptr*, size_t); typedef int (*cuda_cu_memfree_func_t)(gpu::Ptx::CUdeviceptr); typedef int (*cuda_cu_memcpy_htod_func_t)(gpu::Ptx::CUdeviceptr, const void*, unsigned int); typedef int (*cuda_cu_memcpy_dtoh_func_t)(const void*, gpu::Ptx::CUdeviceptr, unsigned int); + typedef int (*cuda_cu_mem_host_register_func_t)(void*, size_t, unsigned int); + typedef int (*cuda_cu_mem_host_get_device_pointer_func_t)(gpu::Ptx::CUdeviceptr*, void*, unsigned int); + typedef int (*cuda_cu_mem_host_unregister_func_t)(void*); public: static cuda_cu_init_func_t _cuda_cu_init; @@ -130,6 +148,9 @@ static cuda_cu_memcpy_htod_func_t _cuda_cu_memcpy_htod; static cuda_cu_memcpy_dtoh_func_t _cuda_cu_memcpy_dtoh; static cuda_cu_ctx_set_current_func_t _cuda_cu_ctx_set_current; + static cuda_cu_mem_host_register_func_t _cuda_cu_mem_host_register; + static cuda_cu_mem_host_get_device_pointer_func_t _cuda_cu_mem_host_get_device_pointer; + static cuda_cu_mem_host_unregister_func_t _cuda_cu_mem_host_unregister; protected: static void* _device_context;
--- a/src/gpu/ptx/vm/ptxKernelArguments.cpp Mon Nov 04 12:18:58 2013 +0100 +++ b/src/gpu/ptx/vm/ptxKernelArguments.cpp Mon Nov 04 12:20:17 2013 +0100 @@ -38,20 +38,32 @@ return arg; } +/* + * Pad kernel argument buffer to naturally align for given size. + */ +void PTXKernelArguments::pad_kernel_argument_buffer(size_t dataSz) { + while ((_bufferOffset % dataSz) != 0) { + *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = (char) 0; + _bufferOffset += sizeof(char); + } + return; +} void PTXKernelArguments::do_int() { // If the parameter is a return value, if (is_return_type()) { - if (is_kernel_arg_setup()) { - // Allocate device memory for T_INT return value pointer on device. Size in bytes - int status = gpu::Ptx::_cuda_cu_memalloc(&_dev_return_value, 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 _dev_return_value to _kernelBuffer - *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = _dev_return_value; + // Allocate device memory for T_INT return value pointer on device. Size in bytes + int status = gpu::Ptx::_cuda_cu_memalloc(&_dev_return_value, 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; } + + // Kernel arguments are expected to be naturally aligned. + // Insert padding into kernel argument buffer, if needed. + pad_kernel_argument_buffer(sizeof(_dev_return_value)); + // Push _dev_return_value to _kernelBuffer + *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = _dev_return_value; _bufferOffset += sizeof(_dev_return_value); } else { // Get the next java argument and its value which should be a T_INT @@ -63,9 +75,13 @@ _success = false; return; } - if (is_kernel_arg_setup()) { - *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = intval.i; - } + + // Kernel arguments are expected to be naturally aligned. + // Insert padding into kernel argument buffer, if needed. + pad_kernel_argument_buffer(sizeof(intval.i)); + // Push _dev_return_value to _kernelBuffer + *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = intval.i; + // Advance _bufferOffset _bufferOffset += sizeof(intval.i); } @@ -75,17 +91,18 @@ void PTXKernelArguments::do_float() { // If the parameter is a return value, if (is_return_type()) { - if (is_kernel_arg_setup()) { - // Allocate device memory for T_INT return value pointer on device. Size in bytes - int status = gpu::Ptx::_cuda_cu_memalloc(&_dev_return_value, T_FLOAT_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 _dev_return_value to _kernelBuffer - *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = _dev_return_value; + // Allocate device memory for T_FLOAT return value pointer on device. Size in bytes + int status = gpu::Ptx::_cuda_cu_memalloc(&_dev_return_value, T_FLOAT_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; } + // Kernel arguments are expected to be naturally aligned. + // Insert padding into kernel argument buffer, if needed. + pad_kernel_argument_buffer(sizeof(_dev_return_value)); + // Push _dev_return_value to _kernelBuffer + *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = _dev_return_value; // Advance _bufferOffset _bufferOffset += sizeof(_dev_return_value); } else { @@ -98,9 +115,11 @@ _success = false; return; } - if (is_kernel_arg_setup()) { - *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = (gpu::Ptx::CUdeviceptr) floatval.f; - } + // Kernel arguments are expected to be naturally aligned. + // Insert padding into kernel argument buffer, if needed. + pad_kernel_argument_buffer(sizeof(floatval.f)); + *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = (gpu::Ptx::CUdeviceptr) floatval.f; + // Advance _bufferOffset _bufferOffset += sizeof(floatval.f); } @@ -111,18 +130,19 @@ // If the parameter is a return value, jvalue doubleval; if (is_return_type()) { - if (is_kernel_arg_setup()) { - // Allocate device memory for T_INT return value pointer on device. Size in bytes - int status = gpu::Ptx::_cuda_cu_memalloc(&_dev_return_value, T_DOUBLE_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 _dev_return_value to _kernelBuffer - *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = _dev_return_value; + // Allocate device memory for T_DOUBLE return value pointer on device. Size in bytes + int status = gpu::Ptx::_cuda_cu_memalloc(&_dev_return_value, T_DOUBLE_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; } - // Advance _bufferOffset + // Kernel arguments are expected to be naturally aligned. + // Insert padding into kernel argument buffer, if needed. + pad_kernel_argument_buffer(sizeof(_dev_return_value)); + // Push _dev_return_value to _kernelBuffer + *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = _dev_return_value; + // Advance _bufferOffset. _bufferOffset += sizeof(doubleval.d); } else { // Get the next java argument and its value which should be a T_INT @@ -133,11 +153,16 @@ _success = false; return; } - if (is_kernel_arg_setup()) { - *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = (gpu::Ptx::CUdeviceptr) doubleval.d; - } + // Kernel arguments are expected to be naturally aligned. + // Insert padding into kernel argument buffer, if needed. + pad_kernel_argument_buffer(sizeof(doubleval.d)); + *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = (gpu::Ptx::CUdeviceptr) doubleval.d; + // Advance _bufferOffset _bufferOffset += sizeof(doubleval.d); + // For a 64-bit host, since size of double is 8, there is no need + // to pad the kernel argument buffer to ensure 8-byte alignment of + // the next potential argument to be pushed. } return; } @@ -145,17 +170,18 @@ void PTXKernelArguments::do_long() { // If the parameter is a return value, if (is_return_type()) { - if (is_kernel_arg_setup()) { - // Allocate device memory for T_LONG return value pointer on device. Size in bytes - int status = gpu::Ptx::_cuda_cu_memalloc(&_dev_return_value, 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 _dev_return_value to _kernelBuffer - *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = _dev_return_value; + // Allocate device memory for T_LONG return value pointer on device. Size in bytes + int status = gpu::Ptx::_cuda_cu_memalloc(&_dev_return_value, 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; } + // Kernel arguments are expected to be naturally aligned. + // Insert padding into kernel argument buffer, if needed. + pad_kernel_argument_buffer(sizeof(_dev_return_value)); + // Push _dev_return_value to _kernelBuffer + *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = _dev_return_value; // Advance _bufferOffset _bufferOffset += sizeof(_dev_return_value); } else { @@ -168,11 +194,16 @@ _success = false; return; } - if (is_kernel_arg_setup()) { - *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = val.j; - } + // Kernel arguments are expected to be naturally aligned. + // Insert padding into kernel argument buffer, if needed. + pad_kernel_argument_buffer(sizeof(val.j)); + *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = val.j; + // Advance _bufferOffset _bufferOffset += sizeof(val.j); + // For a 64-bit host, since size of long is 8, there is no need + // to pad the kernel argument buffer to ensure 8-byte alignment of + // the next potential argument to be pushed. } return; } @@ -180,17 +211,19 @@ void PTXKernelArguments::do_byte() { // If the parameter is a return value, if (is_return_type()) { - if (is_kernel_arg_setup()) { - // Allocate device memory for T_BYTE return value pointer on device. Size in bytes - int status = gpu::Ptx::_cuda_cu_memalloc(&_dev_return_value, 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 _dev_return_value to _kernelBuffer - *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = _dev_return_value; + // Allocate device memory for T_BYTE return value pointer on device. Size in bytes + int status = gpu::Ptx::_cuda_cu_memalloc(&_dev_return_value, 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; } + // Kernel arguments are expected to be naturally aligned. + // Insert padding into kernel argument buffer, if needed. + pad_kernel_argument_buffer(sizeof(_dev_return_value)); + // Push _dev_return_value to _kernelBuffer + *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = _dev_return_value; + // Advance _bufferOffset _bufferOffset += sizeof(_dev_return_value); } else { @@ -203,11 +236,16 @@ _success = false; return; } - if (is_kernel_arg_setup()) { - *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = val.b; - } + // Kernel arguments are expected to be naturally aligned. + // Insert padding into kernel argument buffer, if needed. + pad_kernel_argument_buffer(sizeof(val.b)); + *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = val.b; + // Advance _bufferOffset _bufferOffset += sizeof(val.b); + // For a 64-bit host, since size of T_BYTE is 8, there is no need + // to pad the kernel argument buffer to ensure 8-byte alignment of + // the next potential argument to be pushed. } return; } @@ -215,32 +253,34 @@ void PTXKernelArguments::do_bool() { // If the parameter is a return value, if (is_return_type()) { - if (is_kernel_arg_setup()) { - // Allocate device memory for T_BYTE return value pointer on device. Size in bytes - int status = gpu::Ptx::_cuda_cu_memalloc(&_dev_return_value, T_BOOLEAN_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 _dev_return_value to _kernelBuffer - *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = _dev_return_value; + // Allocate device memory for T_BYTE return value pointer on device. Size in bytes + int status = gpu::Ptx::_cuda_cu_memalloc(&_dev_return_value, T_BOOLEAN_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; } - // Advance _bufferOffset + // Kernel arguments are expected to be naturally aligned. + // Insert padding into kernel argument buffer, if needed. + pad_kernel_argument_buffer(sizeof(_dev_return_value)); + // Push _dev_return_value to _kernelBuffer + *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = _dev_return_value; _bufferOffset += sizeof(_dev_return_value); } else { - // Get the next java argument and its value which should be a T_BYTE - oop arg = next_arg(T_BYTE); + // Get the next java argument and its value which should be a T_BOOLEAN + oop arg = next_arg(T_BOOLEAN); // Copy the java argument value to kernelArgBuffer jvalue val; if (java_lang_boxing_object::get_value(arg, &val) != T_BOOLEAN) { - tty->print_cr("[CUDA] *** Error: Unexpected argument type; expecting T_BYTE"); + tty->print_cr("[CUDA] *** Error: Unexpected argument type; expecting T_BOOLEAN"); _success = false; return; } - if (is_kernel_arg_setup()) { - *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = val.z; - } + // Kernel arguments are expected to be naturally aligned. + // Insert padding into kernel argument buffer, if needed. + pad_kernel_argument_buffer(sizeof(val.z)); + *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = val.z; + // Advance _bufferOffset _bufferOffset += sizeof(val.z); } @@ -257,35 +297,28 @@ gpu::Ptx::CUdeviceptr arrayArgOnDev; int status; - if (is_kernel_arg_setup()) { - // Allocate device memory for array argument on device. Size in bytes - status = gpu::Ptx::_cuda_cu_memalloc(&arrayArgOnDev, argSize); - if (status != GRAAL_CUDA_SUCCESS) { - tty->print_cr("[CUDA] *** Error (%d) Failed to allocate memory for array argument on device", - status); - _success = false; - return; - } - // Copy array argument to device - status = gpu::Ptx::_cuda_cu_memcpy_htod(arrayArgOnDev, arg, argSize); - if (status != GRAAL_CUDA_SUCCESS) { - tty->print_cr("[CUDA] *** Error (%d) Failed to copy array argument content to device memory", - status); - _success = false; - return; - } + // Register host memory for use by the device. Size in bytes + status = gpu::Ptx::_cuda_cu_mem_host_register(arg, argSize, GRAAL_CU_MEMHOSTREGISTER_DEVICEMAP); + if (status != GRAAL_CUDA_SUCCESS) { + tty->print_cr("[CUDA] *** Error (%d) Failed to register host memory for array argument on device", + status); + _success = false; + return; + } + // Get device pointer + status = gpu::Ptx::_cuda_cu_mem_host_get_device_pointer(&arrayArgOnDev, arg, 0); + if (status != GRAAL_CUDA_SUCCESS) { + tty->print_cr("[CUDA] *** Error (%d) Failed to get device pointer of mapped pinned memory of array argument.", + status); + _success = false; + return; + } - // Push device array argument to _kernelBuffer - *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = arrayArgOnDev; - } else { - arrayArgOnDev = *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]); - status = gpu::Ptx::_cuda_cu_memcpy_dtoh(arg, arrayArgOnDev, argSize); - if (status != GRAAL_CUDA_SUCCESS) { - tty->print_cr("[CUDA] *** Error (%d) Failed to copy array argument to host", status); - _success = false; - return; - } - } + // Kernel arguments are expected to be naturally aligned. + // Insert padding into kernel argument buffer, if needed. + pad_kernel_argument_buffer(sizeof(arrayArgOnDev)); + // Push device array argument to _kernelBuffer + *((gpu::Ptx::CUdeviceptr*) &_kernelArgBuffer[_bufferOffset]) = arrayArgOnDev; // Advance _bufferOffset _bufferOffset += sizeof(arrayArgOnDev);
--- a/src/gpu/ptx/vm/ptxKernelArguments.hpp Mon Nov 04 12:18:58 2013 +0100 +++ b/src/gpu/ptx/vm/ptxKernelArguments.hpp Mon Nov 04 12:20:17 2013 +0100 @@ -45,10 +45,6 @@ // Device pointer holding return value gpu::Ptx::CUdeviceptr _dev_return_value; - // Indicates if signature iteration is being done during kernel - // setup i.e., java arguments are being copied to device pointers. - bool _kernelArgSetup; - private: // Array of java argument oops arrayOop _args; @@ -68,8 +64,6 @@ _success = true; _bufferOffset = 0; _dev_return_value = 0; - _kernelArgSetup = true; - //_dev_call_by_reference_args_index = 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."); @@ -87,17 +81,6 @@ return _bufferOffset; } - void copyRefArgsFromDtoH() { - _kernelArgSetup = false; - _bufferOffset = 0; - _index = 0; - iterate(); - } - - inline bool is_kernel_arg_setup() { - return _kernelArgSetup; - } - // Get the return oop value oop get_return_oop(); @@ -106,6 +89,10 @@ return _dev_return_value; } + /* + * Pad kernel argument buffer to naturally align for given size. + */ + void pad_kernel_argument_buffer(size_t); void do_byte(); void do_bool();
--- a/src/share/vm/classfile/vmSymbols.hpp Mon Nov 04 12:18:58 2013 +0100 +++ b/src/share/vm/classfile/vmSymbols.hpp Mon Nov 04 12:20:17 2013 +0100 @@ -352,6 +352,7 @@ template(com_oracle_graal_truffle_GraalTruffleRuntime, "com/oracle/graal/truffle/GraalTruffleRuntime") \ template(startCompiler_name, "startCompiler") \ template(bootstrap_name, "bootstrap") \ + template(compileTheWorld_name, "compileTheWorld") \ template(shutdownCompiler_name, "shutdownCompiler") \ template(compileMethod_name, "compileMethod") \ template(compileMethod_signature, "(JLcom/oracle/graal/hotspot/meta/HotSpotResolvedObjectType;IZ)V") \
--- a/src/share/vm/graal/graalCompiler.cpp Mon Nov 04 12:18:58 2013 +0100 +++ b/src/share/vm/graal/graalCompiler.cpp Mon Nov 04 12:20:17 2013 +0100 @@ -98,15 +98,23 @@ VMToCompiler::finalizeOptions(CITime); if (UseCompiler) { - VMToCompiler::startCompiler(BootstrapGraal); + bool bootstrap = GRAALVM_ONLY(BootstrapGraal) NOT_GRAALVM(false); + VMToCompiler::startCompiler(bootstrap); _initialized = true; - if (BootstrapGraal) { - // We turn off CompileTheWorld and complete the VM startup so that - // Graal can be compiled by C1/C2 when we do a CTW. - NOT_PRODUCT(CompileTheWorld = false); - CompilationPolicy::completed_vm_startup(); + CompilationPolicy::completed_vm_startup(); + if (bootstrap) { VMToCompiler::bootstrap(); } + + +#ifndef PRODUCT + if (CompileTheWorld) { + // We turn off CompileTheWorld so that Graal can + // be compiled by C1/C2 when Graal does a CTW. + CompileTheWorld = false; + VMToCompiler::compileTheWorld(); + } +#endif } } }
--- a/src/share/vm/graal/graalCompilerToGPU.cpp Mon Nov 04 12:18:58 2013 +0100 +++ b/src/share/vm/graal/graalCompilerToGPU.cpp Mon Nov 04 12:20:17 2013 +0100 @@ -99,16 +99,19 @@ if (TraceGPUInteraction) { switch (ptxka.get_ret_type()) { case T_INT: - tty->print_cr("GPU execution returned %d", result.get_jint()); + tty->print_cr("GPU execution returned (int) %d", result.get_jint()); + break; + case T_LONG: + tty->print_cr("GPU execution returned (long) %ld", result.get_jlong()); break; case T_FLOAT: - tty->print_cr("GPU execution returned %f", result.get_jfloat()); + tty->print_cr("GPU execution returned (float) %f", result.get_jfloat()); break; case T_DOUBLE: - tty->print_cr("GPU execution returned %f", result.get_jdouble()); + tty->print_cr("GPU execution returned (double) %f", result.get_jdouble()); break; default: - tty->print_cr("GPU returned unhandled"); + tty->print_cr("**** Value returned by GPU not yet handled"); break; } }
--- a/src/share/vm/graal/graalGlobals.hpp Mon Nov 04 12:18:58 2013 +0100 +++ b/src/share/vm/graal/graalGlobals.hpp Mon Nov 04 12:20:17 2013 +0100 @@ -49,8 +49,8 @@ product(bool, DebugGraal, true, \ "Enable JVMTI for the compiler thread") \ \ - product(bool, BootstrapGraal, GRAALVM_ONLY(true) NOT_GRAALVM(false), \ - "Bootstrap Graal before running Java main method") \ + GRAALVM_ONLY(product(bool, BootstrapGraal, true, \ + "Bootstrap Graal before running Java main method")) \ \ product(intx, TraceGraal, 0, \ "Trace level for Graal") \
--- a/src/share/vm/graal/graalVMToCompiler.cpp Mon Nov 04 12:18:58 2013 +0100 +++ b/src/share/vm/graal/graalVMToCompiler.cpp Mon Nov 04 12:20:17 2013 +0100 @@ -162,7 +162,16 @@ JavaCallArguments args; args.push_oop(instance()); JavaCalls::call_interface(&result, vmToCompilerKlass(), vmSymbols::bootstrap_name(), vmSymbols::void_method_signature(), &args, THREAD); - check_pending_exception("Error while calling boostrap"); + check_pending_exception("Error while calling bootstrap"); +} + +void VMToCompiler::compileTheWorld() { + JavaThread* THREAD = JavaThread::current(); + JavaValue result(T_VOID); + JavaCallArguments args; + args.push_oop(instance()); + JavaCalls::call_interface(&result, vmToCompilerKlass(), vmSymbols::compileTheWorld_name(), vmSymbols::void_method_signature(), &args, THREAD); + check_pending_exception("Error while calling compileTheWorld"); } oop VMToCompiler::createJavaField(Handle holder, Handle name, Handle type, int index, int flags, jboolean internal, TRAPS) {
--- a/src/share/vm/graal/graalVMToCompiler.hpp Mon Nov 04 12:18:58 2013 +0100 +++ b/src/share/vm/graal/graalVMToCompiler.hpp Mon Nov 04 12:20:17 2013 +0100 @@ -72,6 +72,9 @@ // public abstract void bootstrap(); static void bootstrap(); + // public abstract void compileTheWorld(); + static void compileTheWorld(); + // public abstract JavaField createJavaField(JavaType holder, String name, JavaType type, int flags, int offset); static oop createJavaField(Handle holder, Handle name, Handle type, int index, int flags, jboolean internal, TRAPS);