Mercurial > hg > truffle
changeset 12762:884bee435276
Implement support for passing byte arguments in Java methods to be compiled to PTX.
Fix incorrect .reg and .param declarations in PTX kernels.
Refactor and clean up PTXAssembler to PTXMacroAssembler.
line wrap: on
line diff
--- a/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Tue Nov 12 16:21:56 2013 +0100 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Tue Nov 12 13:54:05 2013 -0500 @@ -100,21 +100,17 @@ public static class StandardFormat { + // Type of destination value protected Kind valueKind; protected Variable dest; protected Value source1; protected Value source2; - private boolean logicInstruction = false; - private boolean ldRetAddrInstruction = false; public StandardFormat(Variable dst, Value src1, Value src2) { setDestination(dst); setSource1(src1); setSource2(src2); setKind(dst.getKind()); - - // testAdd2B fails this assertion - // assert valueKind == src1.getKind(); } public void setKind(Kind k) { @@ -124,6 +120,7 @@ public void setDestination(Variable var) { assert var != null; dest = var; + setKind(var.getKind()); } public void setSource1(Value val) { @@ -136,57 +133,33 @@ source2 = val; } - public void setLogicInstruction(boolean b) { - logicInstruction = b; - } - - public void setLdRetAddrInstruction(boolean b) { - ldRetAddrInstruction = b; - } - public String typeForKind(Kind k) { - if (ldRetAddrInstruction) { - if (System.getProperty("os.arch").compareTo("amd64") == 0) { + switch (k.getTypeChar()) { + // Boolean + case 'z': + return "u8"; + // Byte + case 'b': + return "b8"; + // Short + case 's': + return "s16"; + case 'c': + return "u16"; + case 'i': + return "s32"; + case 'f': + return "f32"; + case 'j': + return "s64"; + case 'd': + return "f64"; + case 'a': return "u64"; - } else { + case '-': return "u32"; - } - } else if (logicInstruction) { - switch (k.getTypeChar()) { - case 's': - return "b16"; - case 'i': - return "b32"; - case 'j': - return "b64"; - default: - throw GraalInternalError.shouldNotReachHere(); - } - } else { - switch (k.getTypeChar()) { - case 'z': - return "u8"; - case 'b': - return "s8"; - case 's': - return "s16"; - case 'c': - return "u16"; - case 'i': - return "s32"; - case 'f': - return "f32"; - case 'j': - return "s64"; - case 'd': - return "f64"; - case 'a': - return "u64"; - case '-': - return "u32"; - default: - throw GraalInternalError.shouldNotReachHere(); - } + default: + throw GraalInternalError.shouldNotReachHere(); } } @@ -236,6 +209,32 @@ } } + public static class LogicInstructionFormat extends StandardFormat { + public LogicInstructionFormat(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + } + + @Override + public String emit() { + String kindStr; + switch (valueKind.getTypeChar()) { + case 's': + kindStr = "b16"; + break; + case 'i': + kindStr = "b32"; + break; + case 'j': + kindStr = "b64"; + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + + return (kindStr + emitRegister(dest, true) + emitValue(source1, true) + emitValue(source2, false) + ";"); + } + } + public static class SingleOperandFormat { protected Variable dest; @@ -397,132 +396,6 @@ } } - public static class Add extends StandardFormat { - - public Add(Variable dst, Value src1, Value src2) { - super(dst, src1, src2); - } - - public void emit(PTXAssembler asm) { - asm.emitString("add." + super.emit()); - } - } - - public static class And extends StandardFormat { - - public And(Variable dst, Value src1, Value src2) { - super(dst, src1, src2); - setLogicInstruction(true); - } - - public void emit(PTXAssembler asm) { - asm.emitString("and." + super.emit()); - } - } - - public static class Div extends StandardFormat { - - public Div(Variable dst, Value src1, Value src2) { - super(dst, src1, src2); - } - - public void emit(PTXAssembler asm) { - asm.emitString("div." + super.emit()); - } - } - - public static class Mul extends StandardFormat { - - public Mul(Variable dst, Value src1, Value src2) { - super(dst, src1, src2); - } - - public void emit(PTXAssembler asm) { - asm.emitString("mul.lo." + super.emit()); - } - } - - public static class Or extends StandardFormat { - - public Or(Variable dst, Value src1, Value src2) { - super(dst, src1, src2); - setLogicInstruction(true); - } - - public void emit(PTXAssembler asm) { - asm.emitString("or." + super.emit()); - } - } - - public static class Rem extends StandardFormat { - - public Rem(Variable dst, Value src1, Value src2) { - super(dst, src1, src2); - } - - public void emit(PTXAssembler asm) { - asm.emitString("rem." + super.emit()); - } - } - - public static class Shl extends StandardFormat { - - public Shl(Variable dst, Value src1, Value src2) { - super(dst, src1, src2); - setLogicInstruction(true); - } - - public void emit(PTXAssembler asm) { - asm.emitString("shl." + super.emit()); - } - } - - public static class Shr extends StandardFormat { - - public Shr(Variable dst, Value src1, Value src2) { - super(dst, src1, src2); - } - - public void emit(PTXAssembler asm) { - asm.emitString("shr." + super.emit()); - } - } - - public static class Sub extends StandardFormat { - - public Sub(Variable dst, Value src1, Value src2) { - super(dst, src1, src2); - } - - public void emit(PTXAssembler asm) { - asm.emitString("sub." + super.emit()); - } - } - - public static class Ushr extends StandardFormat { - - public Ushr(Variable dst, Value src1, Value src2) { - super(dst, src1, src2); - setKind(Kind.Illegal); // get around not having an Unsigned Kind - } - - public void emit(PTXAssembler asm) { - asm.emitString("shr." + super.emit()); - } - } - - public static class Xor extends StandardFormat { - - public Xor(Variable dst, Value src1, Value src2) { - super(dst, src1, src2); - setLogicInstruction(true); - } - - public void emit(PTXAssembler asm) { - asm.emitString("xor." + super.emit()); - } - } - // Checkstyle: stop method name check public final void bra(String tgt, int pred) { assert pred >= 0; @@ -541,93 +414,6 @@ emitString("bra.uni" + " " + tgt + ";" + ""); } - public static class Cvt extends ConversionFormat { - - public Cvt(Variable dst, Variable src, Kind dstKind, Kind srcKind) { - super(dst, src, dstKind, srcKind); - } - - public void emit(PTXAssembler asm) { - if (dest.getKind() == Kind.Float || dest.getKind() == Kind.Double) { - // round-to-zero - might not be right - asm.emitString("cvt.rz." + super.emit()); - } else { - asm.emitString("cvt." + super.emit()); - } - } - } - - public static class Mov extends SingleOperandFormat { - - private int predicateRegisterNumber = -1; - - public Mov(Variable dst, Value src) { - super(dst, src); - } - - public Mov(Variable dst, Value src, int predicate) { - super(dst, src); - this.predicateRegisterNumber = predicate; - } - - /* - * public Mov(Variable dst, AbstractAddress src) { throw - * GraalInternalError.unimplemented("AbstractAddress Mov"); } - */ - - public void emit(PTXAssembler asm) { - if (predicateRegisterNumber >= 0) { - asm.emitString("@%p" + String.valueOf(predicateRegisterNumber) + " mov." + super.emit()); - } else { - asm.emitString("mov." + super.emit()); - } - } - } - - public static class Neg extends SingleOperandFormat { - - public Neg(Variable dst, Variable src) { - super(dst, src); - } - - public void emit(PTXAssembler asm) { - asm.emitString("neg." + super.emit()); - } - } - - public static class Not extends BinarySingleOperandFormat { - - public Not(Variable dst, Variable src) { - super(dst, src); - } - - public void emit(PTXAssembler asm) { - asm.emitString("not." + super.emit()); - } - } - - public static class Ld extends LoadStoreFormat { - - public Ld(PTXStateSpace space, Variable dst, Variable src1, Value src2) { - super(space, dst, src1, src2); - } - - public void emit(PTXAssembler asm) { - asm.emitString("ld." + super.emit(true)); - } - } - - public static class St extends LoadStoreFormat { - - public St(PTXStateSpace space, Variable dst, Variable src1, Value src2) { - super(space, dst, src1, src2); - } - - public void emit(PTXAssembler asm) { - asm.emitString("st." + super.emit(false)); - } - } - public final void exit() { emitString("exit;" + " " + ""); } @@ -673,61 +459,6 @@ } } - public static class Param extends SingleOperandFormat { - // Last parameter holds the return parameter. - private boolean returnParameter; - - public Param(Variable d, boolean lastParam) { - super(d, null); - setReturnParameter(lastParam); - } - - public void setReturnParameter(boolean value) { - returnParameter = value; - } - - public String emitParameter(Variable v) { - return (" param" + v.index); - } - - public void emit(PTXAssembler asm) { - asm.emitString(".param ." + paramForKind(dest.getKind()) + emitParameter(dest) + (returnParameter ? "" : ",")); - } - - public String paramForKind(Kind k) { - if (returnParameter) { - if (System.getProperty("os.arch").compareTo("amd64") == 0) { - return "u64"; - } 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(); - } - } - } - - } - public final void popc_b32(Register d, Register a) { emitString("popc.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); }
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXMacroAssembler.java Tue Nov 12 13:54:05 2013 -0500 @@ -0,0 +1,322 @@ +/* + * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ +package com.oracle.graal.asm.ptx; + +import sun.misc.*; + +import com.oracle.graal.api.code.*; +import com.oracle.graal.lir.*; +import com.oracle.graal.api.meta.*; +import com.oracle.graal.graph.GraalInternalError; + +public class PTXMacroAssembler extends PTXAssembler { + + public PTXMacroAssembler(TargetDescription target, RegisterConfig registerConfig) { + super(target, registerConfig); + } + + public static class LoadAddr extends LoadStoreFormat { + + public LoadAddr(PTXStateSpace space, Variable dst, Variable src1, Value src2) { + super(space, dst, src1, src2); + } + + public void emit(PTXMacroAssembler masm) { + String ldAddrStr = "ld." + space.getStateName(); + if (Unsafe.ADDRESS_SIZE == 8) { + ldAddrStr = ldAddrStr + ".u64"; + } else { + ldAddrStr = ldAddrStr + ".u32"; + } + masm.emitString(ldAddrStr + " " + emitRegister(dest, false) + ", " + emitAddress(source1, source2) + ";"); + } + } + + public static class Ld extends LoadStoreFormat { + + public Ld(PTXStateSpace space, Variable dst, Variable src1, Value src2) { + super(space, dst, src1, src2); + } + + public void emit(PTXMacroAssembler asm) { + asm.emitString("ld." + super.emit(true)); + } + } + + public static class St extends LoadStoreFormat { + + public St(PTXStateSpace space, Variable dst, Variable src1, Value src2) { + super(space, dst, src1, src2); + } + + public void emit(PTXMacroAssembler asm) { + asm.emitString("st." + super.emit(false)); + } + } + + public static class LoadParam extends Ld { + // Type of the operation is dependent on src1's type. + public LoadParam(PTXStateSpace space, Variable dst, Variable src1, Value src2) { + super(space, dst, src1, src2); + setKind(src1.getKind()); + } + } + + public static class Add extends StandardFormat { + + public Add(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + } + + public void emit(PTXMacroAssembler asm) { + asm.emitString("add." + super.emit()); + } + } + + public static class And extends LogicInstructionFormat { + + public And(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + } + + public void emit(PTXMacroAssembler asm) { + asm.emitString("and." + super.emit()); + } + } + + public static class Div extends StandardFormat { + + public Div(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + } + + public void emit(PTXMacroAssembler asm) { + asm.emitString("div." + super.emit()); + } + } + + public static class Mul extends StandardFormat { + + public Mul(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + } + + public void emit(PTXMacroAssembler asm) { + asm.emitString("mul.lo." + super.emit()); + } + } + + public static class Or extends LogicInstructionFormat { + + public Or(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + } + + public void emit(PTXMacroAssembler asm) { + asm.emitString("or." + super.emit()); + } + } + + public static class Rem extends StandardFormat { + + public Rem(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + } + + public void emit(PTXMacroAssembler asm) { + asm.emitString("rem." + super.emit()); + } + } + + public static class Shl extends LogicInstructionFormat { + + public Shl(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + } + + public void emit(PTXMacroAssembler asm) { + asm.emitString("shl." + super.emit()); + } + } + + public static class Shr extends StandardFormat { + + public Shr(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + } + + public void emit(PTXMacroAssembler asm) { + asm.emitString("shr." + super.emit()); + } + } + + public static class Sub extends StandardFormat { + + public Sub(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + } + + public void emit(PTXMacroAssembler asm) { + asm.emitString("sub." + super.emit()); + } + } + + public static class Ushr extends StandardFormat { + + public Ushr(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + setKind(Kind.Illegal); // get around not having an Unsigned Kind + } + + public void emit(PTXMacroAssembler asm) { + asm.emitString("shr." + super.emit()); + } + } + + public static class Xor extends LogicInstructionFormat { + + public Xor(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + } + + public void emit(PTXMacroAssembler asm) { + asm.emitString("xor." + super.emit()); + } + } + + public static class Cvt extends ConversionFormat { + + public Cvt(Variable dst, Variable src, Kind dstKind, Kind srcKind) { + super(dst, src, dstKind, srcKind); + } + + public void emit(PTXMacroAssembler asm) { + if (dest.getKind() == Kind.Float || dest.getKind() == Kind.Double) { + // round-to-zero - might not be right + asm.emitString("cvt.rz." + super.emit()); + } else { + asm.emitString("cvt." + super.emit()); + } + } + } + + public static class Mov extends SingleOperandFormat { + + private int predicateRegisterNumber = -1; + + public Mov(Variable dst, Value src) { + super(dst, src); + } + + public Mov(Variable dst, Value src, int predicate) { + super(dst, src); + this.predicateRegisterNumber = predicate; + } + + /* + * public Mov(Variable dst, AbstractAddress src) { throw + * GraalInternalError.unimplemented("AbstractAddress Mov"); } + */ + + public void emit(PTXMacroAssembler asm) { + if (predicateRegisterNumber >= 0) { + asm.emitString("@%p" + String.valueOf(predicateRegisterNumber) + " mov." + super.emit()); + } else { + asm.emitString("mov." + super.emit()); + } + } + } + + public static class Neg extends SingleOperandFormat { + + public Neg(Variable dst, Variable src) { + super(dst, src); + } + + public void emit(PTXMacroAssembler asm) { + asm.emitString("neg." + super.emit()); + } + } + + public static class Not extends BinarySingleOperandFormat { + + public Not(Variable dst, Variable src) { + super(dst, src); + } + + public void emit(PTXMacroAssembler asm) { + asm.emitString("not." + super.emit()); + } + } + + public static class Param extends SingleOperandFormat { + boolean isReturnParameter; + + public Param(Variable d, boolean isRet) { + super(d, null); + isReturnParameter = isRet; + } + + public String emitParameter(Variable v) { + return (" param" + v.index); + } + + public void emit(PTXMacroAssembler asm, boolean isLastParam) { + asm.emitString(".param ." + paramForKind(dest.getKind()) + emitParameter(dest) + (isLastParam ? "" : ",")); + } + + public String paramForKind(Kind k) { + if (isReturnParameter) { + if (Unsafe.ADDRESS_SIZE == 8) { + return "u64"; + } else { + return "u32"; + } + } else { + switch (k.getTypeChar()) { + case 'z': + case 'f': + return "s32"; + case 'b': + return "b8"; + 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(); + } + } + } + + } + +}
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ControlPTXTest.java Tue Nov 12 16:21:56 2013 +0100 +++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ControlPTXTest.java Tue Nov 12 13:54:05 2013 -0500 @@ -30,36 +30,78 @@ @Test public void testControl() { - Integer ret = (Integer) invoke(compile("testLoop"), 42); + String testName = "testLoop"; + int argVal1 = 42; + int argVal2; + Integer ret = (Integer) invoke(compile(testName), argVal1); + int jres = testLoop(argVal1); if (ret != null) { - printReport("testLoop: " + ret); + if (ret.intValue() == jres) { + printReport(testName + ": PASSED"); + } else { + printReport(testName + ": FAILED " + "Expected " + jres + " Got " + ret); + } } else { - printReport("testLoop: no VALUE"); + printReport(testName + ": FAILED (null returned)"); } - ret = (Integer) invoke(compile("testSwitchDefault1I"), 3); + + testName = "testSwitchDefault1I"; + argVal1 = 3; + ret = (Integer) invoke(compile(testName), argVal1); + jres = testSwitchDefault1I(argVal1); if (ret != null) { - printReport("testSwitchDefault1I: " + ret); + if (ret.intValue() == jres) { + printReport(testName + ": PASSED"); + } else { + printReport(testName + ": FAILED " + "Expected " + jres + " Got " + ret); + } } else { - printReport("testSwitchDefault1I: no VALUE"); + printReport(testName + ": FAILED (null returned)"); } - ret = (Integer) invoke(compile("testSwitch1I"), 2); + + testName = "testSwitch1I"; + argVal1 = 2; + ret = (Integer) invoke(compile(testName), argVal1); + jres = testSwitch1I(argVal1); if (ret != null) { - printReport("testSwitch1I: " + ret); - } else { - printReport("testSwitch1I: no VALUE"); - } - ret = (Integer) invoke(compile("testIfElse1I"), 222); - if (ret != null) { - printReport("testIfElse1I: " + ret); + if (ret.intValue() == jres) { + printReport(testName + ": PASSED"); + } else { + printReport(testName + ": FAILED " + "Expected " + jres + " Got " + ret); + } } else { - printReport("testIfElse1I: no VALUE"); + printReport(testName + ": FAILED (null returned)"); } - ret = (Integer) invoke(compile("testIfElse2I"), 19, 64); + + testName = "testIfElse1I"; + argVal1 = 222; + ret = (Integer) invoke(compile(testName), argVal1); + jres = testIfElse1I(argVal1); if (ret != null) { - printReport("testIfElse2I: " + (char) ret.intValue()); + if (ret.intValue() == jres) { + printReport(testName + ": PASSED"); + } else { + printReport(testName + ": FAILED " + "Expected " + jres + " Got " + ret); + } } else { - printReport("testIfElse2I: no VALUE"); + printReport(testName + ": FAILED (null returned)"); } + + testName = "testIfElse2I"; + argVal1 = 19; + argVal2 = 64; + ret = (Integer) invoke(compile(testName), argVal1, argVal2); + jres = testIfElse2I(argVal1, argVal2); + if (ret != null) { + if (ret.intValue() == jres) { + printReport(testName + ": PASSED"); + } else { + printReport(testName + ": FAILED " + "Expected " + jres + " Got " + ret); + } + } else { + printReport(testName + ": FAILED " + "Expected " + jres + " Got " + ret); + } + Boolean bret = (Boolean) invoke(compile("testIntegerTestBranch2I"), 0xff00, 0x00ff); if (bret != null) { printReport("testIntegerTestBranch2I: " + bret);
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/IntegerPTXTest.java Tue Nov 12 16:21:56 2013 +0100 +++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/IntegerPTXTest.java Tue Nov 12 13:54:05 2013 -0500 @@ -32,17 +32,17 @@ public void testAdd() { // @formatter:off - /* Integer r4 = (Integer) invoke(compile("testAdd2B"), (byte) 6, (byte) 4); + Integer r4 = (Integer) invoke(compile("testAdd2B"), (byte) 6, (byte) 4); if (r4 == null) { printReport("testAdd2B FAILED"); } else if (r4.intValue() == testAdd2B((byte) 6, (byte) 4)) { printReport("testAdd2B PASSED"); } else { printReport("testAdd2B FAILED"); - } */ + } // @formatter:on - Integer r4 = (Integer) invoke(compile("testAdd2I"), 18, 24); + r4 = (Integer) invoke(compile("testAdd2I"), 18, 24); if (r4 == null) { printReport("testAdd2I FAILED"); } else if (r4.intValue() == testAdd2I(18, 24)) { @@ -267,7 +267,6 @@ return 32 / a; } - @Ignore @Test public void testRem() { Integer r1 = (Integer) invoke(compile("testRem2I"), 8, 4);
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java Tue Nov 12 16:21:56 2013 +0100 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java Tue Nov 12 13:54:05 2013 -0500 @@ -136,11 +136,12 @@ // Need to emit .param directives based on incoming arguments and return value CallingConvention incomingArguments = cc; Object returnObject = incomingArguments.getReturn(); - AllocatableValue[] params; - int argCount; + AllocatableValue[] params = incomingArguments.getArguments(); + int argCount = incomingArguments.getArgumentCount(); if (returnObject == Value.ILLEGAL) { params = incomingArguments.getArguments(); + append(new PTXParameterOp(params, false)); } else { argCount = incomingArguments.getArgumentCount(); params = new Variable[argCount + 1]; @@ -148,9 +149,9 @@ params[i] = incomingArguments.getArgument(i); } params[argCount] = (Variable) returnObject; + append(new PTXParameterOp(params, true)); } - append(new PTXParameterOp(params)); for (LocalNode local : graph.getNodes(LocalNode.class)) { Value param = params[local.index()]; Annotation[] annos = graph.method().getParameterAnnotations()[local.index()];
--- a/graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotBackend.java Tue Nov 12 16:21:56 2013 +0100 +++ b/graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotBackend.java Tue Nov 12 13:54:05 2013 -0500 @@ -66,15 +66,23 @@ } static final class RegisterAnalysis extends ValueProcedure { - private final SortedSet<Integer> unsigned64 = new TreeSet<>(); + private final SortedSet<Integer> signed32 = new TreeSet<>(); private final SortedSet<Integer> signed64 = new TreeSet<>(); + + // unsigned8 is only for ld, st and cbt + private final SortedSet<Integer> unsigned8 = new TreeSet<>(); + private final SortedSet<Integer> unsigned64 = new TreeSet<>(); + + // private final SortedSet<Integer> float16 = 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 : unsigned8) { + codeBuffer.emitString(".reg .u8 %r" + i.intValue() + ";"); + } for (Integer i : signed32) { codeBuffer.emitString(".reg .s32 %r" + i.intValue() + ";"); } @@ -134,6 +142,9 @@ case Object: unsigned64.add(regVal.index); break; + case Byte: + unsigned8.add(regVal.index); + break; default: throw GraalInternalError.shouldNotReachHere("unhandled register type " + value.toString()); } @@ -172,7 +183,7 @@ @Override protected AbstractAssembler createAssembler(FrameMap frameMap) { - return new PTXAssembler(getTarget(), frameMap.registerConfig); + return new PTXMacroAssembler(getTarget(), frameMap.registerConfig); } @Override
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXArithmetic.java Tue Nov 12 16:21:56 2013 +0100 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXArithmetic.java Tue Nov 12 13:54:05 2013 -0500 @@ -22,7 +22,7 @@ */ package com.oracle.graal.lir.ptx; -import static com.oracle.graal.asm.ptx.PTXAssembler.*; +import static com.oracle.graal.asm.ptx.PTXMacroAssembler.*; import static com.oracle.graal.api.code.ValueUtil.*; import static com.oracle.graal.lir.LIRValueUtil.*; import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; @@ -56,7 +56,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { Variable dst = (Variable) result; Variable src = (Variable) x; if (from == Kind.Long && to == Kind.Int) { @@ -81,7 +81,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { emit(tasm, masm, opcode, result, x, null); } } @@ -100,7 +100,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { emit(tasm, masm, opcode, result, x, y, null); } @@ -125,7 +125,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { emit(tasm, masm, opcode, result, x, y, null); } @@ -150,7 +150,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { if (sameRegister(result, y)) { emit(tasm, masm, opcode, result, x, null); } else { @@ -180,7 +180,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { emit(tasm, masm, opcode, result, x, y, null); } @@ -208,7 +208,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { emit(tasm, masm, opcode, result, y, state); } @@ -219,7 +219,7 @@ } } - public static void emit(TargetMethodAssembler tasm, PTXAssembler masm, PTXArithmetic opcode, Value dst, Value src, LIRFrameState info) { + public static void emit(TargetMethodAssembler tasm, PTXMacroAssembler masm, PTXArithmetic opcode, Value dst, Value src, LIRFrameState info) { int exceptionOffset = -1; Variable dest = (Variable) dst; @@ -271,7 +271,7 @@ } } - public static void emit(TargetMethodAssembler tasm, PTXAssembler masm, PTXArithmetic opcode, + public static void emit(TargetMethodAssembler tasm, PTXMacroAssembler masm, PTXArithmetic opcode, Value dst, Value src1, Value src2, LIRFrameState info) { int exceptionOffset = -1; Variable dest = (Variable) dst;
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXBitManipulationOp.java Tue Nov 12 16:21:56 2013 +0100 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXBitManipulationOp.java Tue Nov 12 13:54:05 2013 -0500 @@ -46,7 +46,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { Register dst = ValueUtil.asIntReg(result); Register src = ValueUtil.asRegister(input); switch (opcode) {
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXCompare.java Tue Nov 12 16:21:56 2013 +0100 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXCompare.java Tue Nov 12 13:54:05 2013 -0500 @@ -54,7 +54,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { emit(masm, opcode, condition, x, y, predRegNum); }
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java Tue Nov 12 16:21:56 2013 +0100 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java Tue Nov 12 13:54:05 2013 -0500 @@ -23,6 +23,7 @@ package com.oracle.graal.lir.ptx; import static com.oracle.graal.asm.ptx.PTXAssembler.*; +import static com.oracle.graal.asm.ptx.PTXMacroAssembler.*; import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; import static com.oracle.graal.lir.LIRValueUtil.*; import static com.oracle.graal.nodes.calc.Condition.*; @@ -48,7 +49,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { if (tasm.frameContext != null) { tasm.frameContext.leave(tasm); } @@ -62,7 +63,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { if (tasm.frameContext != null) { tasm.frameContext.leave(tasm); } @@ -83,7 +84,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { masm.bra(masm.nameOf(destination.label()), predRegNum); } @@ -116,7 +117,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { cmove(tasm, masm, result, false, condition, false, trueValue, falseValue, predicate); } } @@ -140,12 +141,12 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { cmove(tasm, masm, result, true, condition, unorderedIsTrue, trueValue, falseValue, predicate); } } - private static void cmove(TargetMethodAssembler tasm, PTXAssembler asm, Value result, boolean isFloat, Condition condition, boolean unorderedIsTrue, Value trueValue, Value falseValue, + private static void cmove(TargetMethodAssembler tasm, PTXMacroAssembler asm, Value result, boolean isFloat, Condition condition, boolean unorderedIsTrue, Value trueValue, Value falseValue, int predicateRegister) { // check that we don't overwrite an input operand before it is used. assert !result.equals(trueValue); @@ -177,7 +178,7 @@ } } - private static void cmove(PTXAssembler asm, Value result, Value other, int predicateRegister) { + private static void cmove(PTXMacroAssembler asm, Value result, Value other, int predicateRegister) { if (isVariable(other)) { assert !asVariable(other).equals(asVariable(result)) : "other already overwritten by previous move"; @@ -217,7 +218,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { Kind keyKind = key.getKind(); if (keyKind == Kind.Int || keyKind == Kind.Long) { @@ -275,7 +276,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { tableswitch(tasm, masm, lowKey, defaultTarget, targets, index, scratch, predRegNum); } }
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXLIRInstruction.java Tue Nov 12 16:21:56 2013 +0100 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXLIRInstruction.java Tue Nov 12 13:54:05 2013 -0500 @@ -33,8 +33,8 @@ @Override public final void emitCode(TargetMethodAssembler tasm) { - emitCode(tasm, (PTXAssembler) tasm.asm); + emitCode(tasm, (PTXMacroAssembler) tasm.asm); } - public abstract void emitCode(TargetMethodAssembler tasm, PTXAssembler masm); + public abstract void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm); }
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMemOp.java Tue Nov 12 16:21:56 2013 +0100 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMemOp.java Tue Nov 12 13:54:05 2013 -0500 @@ -22,7 +22,7 @@ */ package com.oracle.graal.lir.ptx; -import static com.oracle.graal.asm.ptx.PTXAssembler.*; +import static com.oracle.graal.asm.ptx.PTXMacroAssembler.*; import static com.oracle.graal.asm.ptx.PTXStateSpace.*; import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; @@ -51,7 +51,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { PTXAddress addr = address.toAddress(); switch (kind) { case Byte: @@ -87,7 +87,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { PTXAddress addr = address.toAddress(); switch (kind) { case Byte: @@ -122,7 +122,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { PTXAddress addr = address.toAddress(); switch (kind) { case Byte: @@ -133,7 +133,7 @@ case Float: case Double: case Object: - new Ld(Parameter, result, addr.getBase(), Constant.forLong(addr.getDisplacement())).emit(masm); + new LoadParam(Parameter, result, addr.getBase(), Constant.forLong(addr.getDisplacement())).emit(masm); break; default: throw GraalInternalError.shouldNotReachHere(); @@ -159,16 +159,14 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { PTXAddress addr = address.toAddress(); switch (kind) { case Int: case Long: case Float: case Double: - Ld ldIns = new Ld(Parameter, result, addr.getBase(), Constant.forLong(addr.getDisplacement())); - ldIns.setLdRetAddrInstruction(true); - ldIns.emit(masm); + new LoadAddr(Parameter, result, addr.getBase(), Constant.forLong(addr.getDisplacement())).emit(masm); break; default: throw GraalInternalError.shouldNotReachHere(); @@ -193,7 +191,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { PTXAddress addr = address.toAddress(); switch (kind) {
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java Tue Nov 12 16:21:56 2013 +0100 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java Tue Nov 12 13:54:05 2013 -0500 @@ -22,7 +22,7 @@ */ package com.oracle.graal.lir.ptx; -import static com.oracle.graal.asm.ptx.PTXAssembler.*; +import static com.oracle.graal.asm.ptx.PTXMacroAssembler.*; import static com.oracle.graal.api.code.ValueUtil.*; import static com.oracle.graal.lir.LIRValueUtil.*; import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; @@ -49,7 +49,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { move(tasm, masm, getResult(), getInput()); } @@ -76,7 +76,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { move(tasm, masm, getResult(), getInput()); } @@ -103,7 +103,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { move(tasm, masm, getResult(), getInput()); } @@ -129,7 +129,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { throw new InternalError("NYI"); } } @@ -145,7 +145,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { throw new InternalError("NYI"); } } @@ -166,12 +166,12 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { compareAndSwap(tasm, masm, result, address, cmpValue, newValue); } } - public static void move(TargetMethodAssembler tasm, PTXAssembler masm, Value result, Value input) { + public static void move(TargetMethodAssembler tasm, PTXMacroAssembler masm, Value result, Value input) { if (isVariable(input)) { if (isVariable(result)) { reg2reg(masm, result, input); @@ -189,7 +189,7 @@ } } - private static void reg2reg(PTXAssembler masm, Value result, Value input) { + private static void reg2reg(PTXMacroAssembler masm, Value result, Value input) { Variable dest = (Variable) result; Variable source = (Variable) input; @@ -209,7 +209,7 @@ } } - private static void const2reg(TargetMethodAssembler tasm, PTXAssembler masm, Value result, Constant input) { + private static void const2reg(TargetMethodAssembler tasm, PTXMacroAssembler masm, Value result, Constant input) { Variable dest = (Variable) result; switch (input.getKind().getStackKind()) {
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXParameterOp.java Tue Nov 12 16:21:56 2013 +0100 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXParameterOp.java Tue Nov 12 13:54:05 2013 -0500 @@ -23,7 +23,7 @@ package com.oracle.graal.lir.ptx; -import static com.oracle.graal.asm.ptx.PTXAssembler.*; +import static com.oracle.graal.asm.ptx.PTXMacroAssembler.*; import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; import com.oracle.graal.api.meta.*; @@ -34,18 +34,23 @@ public class PTXParameterOp extends LIRInstruction { @Def({REG}) protected Value[] params; + // True if the parameter list has return argument as the last + // item of the array params. + private boolean hasReturnParam; - public PTXParameterOp(Value[] params) { + public PTXParameterOp(Value[] params, boolean hasReturn) { this.params = params; + hasReturnParam = hasReturn; } @Override public void emitCode(TargetMethodAssembler tasm) { - PTXAssembler masm = (PTXAssembler) tasm.asm; + PTXMacroAssembler masm = (PTXMacroAssembler) tasm.asm; // Emit parameter directives for arguments int argCount = params.length; for (int i = 0; i < argCount; i++) { - new Param((Variable) params[i], (i == (argCount - 1))).emit(masm); + boolean isReturnParam = (hasReturnParam && (i == (argCount - 1))); + new Param((Variable) params[i], isReturnParam).emit(masm, (i == (argCount - 1))); } } }
--- a/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXTestOp.java Tue Nov 12 16:21:56 2013 +0100 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXTestOp.java Tue Nov 12 13:54:05 2013 -0500 @@ -44,7 +44,7 @@ } @Override - public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + public void emitCode(TargetMethodAssembler tasm, PTXMacroAssembler masm) { emit(masm, x, y, predicate); }