Mercurial > hg > truffle
changeset 7849:268d3e74191e
Merge.
| author | Christian Humer <christian.humer@gmail.com> |
|---|---|
| date | Tue, 19 Feb 2013 17:41:10 +0100 |
| parents | 698cd036a1ca (current diff) 8959b331ef3e (diff) |
| children | 85891f9c2197 |
| files | src/share/vm/graal/graalVmIds.cpp src/share/vm/graal/graalVmIds.hpp |
| diffstat | 72 files changed, 2943 insertions(+), 383 deletions(-) [+] |
line wrap: on
line diff
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CodeCacheProvider.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CodeCacheProvider.java Tue Feb 19 17:41:10 2013 +0100 @@ -81,11 +81,6 @@ int getMinimumOutgoingSize(); /** - * Performs any runtime-specific conversion on the object used to describe the target of a call. - */ - Object lookupCallTarget(Object callTarget); - - /** * Gets the signature and linkage information for a runtime call. */ RuntimeCallTarget lookupRuntimeCall(Descriptor descriptor);
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CompilationResult.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CompilationResult.java Tue Feb 19 17:41:10 2013 +0100 @@ -94,7 +94,7 @@ /** * The target of the call. */ - public final Object target; + public final InvokeTarget target; /** * The size of the call instruction. @@ -108,7 +108,7 @@ */ public final boolean direct; - public Call(Object target, int pcOffset, int size, boolean direct, DebugInfo debugInfo) { + public Call(InvokeTarget target, int pcOffset, int size, boolean direct, DebugInfo debugInfo) { super(pcOffset, debugInfo); this.size = size; this.target = target; @@ -438,11 +438,11 @@ * * @param codePos the position of the call in the code array * @param size the size of the call instruction - * @param target the {@link CodeCacheProvider#lookupCallTarget(Object) target} being called + * @param target the being called * @param debugInfo the debug info for the call * @param direct specifies if this is a {@linkplain Call#direct direct} call */ - public void recordCall(int codePos, int size, Object target, DebugInfo debugInfo, boolean direct) { + public void recordCall(int codePos, int size, InvokeTarget target, DebugInfo debugInfo, boolean direct) { final Call call = new Call(target, codePos, size, direct, debugInfo); addSafepoint(call); }
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/Register.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/Register.java Tue Feb 19 17:41:10 2013 +0100 @@ -65,6 +65,13 @@ public final int encoding; /** + * The assembler calls this method to get the register's encoding. + */ + public int encoding() { + return encoding; + } + + /** * The size of the stack slot used to spill the value of this register. */ public final int spillSlotSize;
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/RuntimeCallTarget.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/RuntimeCallTarget.java Tue Feb 19 17:41:10 2013 +0100 @@ -24,11 +24,13 @@ import java.util.*; +import com.oracle.graal.api.meta.*; + /** * The name, signature and calling convention of a call from compiled code to the runtime. The * target of such a call may be a leaf stub or a call into the runtime code proper. */ -public interface RuntimeCallTarget { +public interface RuntimeCallTarget extends InvokeTarget { /** * The name and signature of a runtime call.
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/StackSlot.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/StackSlot.java Tue Feb 19 17:41:10 2013 +0100 @@ -165,7 +165,7 @@ private static StackSlot[][] makeCache(int cachePerKindSize, int sign, boolean addFrameSize) { StackSlot[][] cache = new StackSlot[Kind.values().length][]; - for (Kind kind : new Kind[]{Illegal, Int, Long, Float, Double, Object, Jsr}) { + for (Kind kind : new Kind[]{Illegal, Int, Long, Float, Double, Object}) { StackSlot[] slots = new StackSlot[cachePerKindSize]; for (int i = 0; i < cachePerKindSize; i++) { slots[i] = new StackSlot(kind, sign * i * CACHE_GRANULARITY, addFrameSize);
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/TargetDescription.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/TargetDescription.java Tue Feb 19 17:41:10 2013 +0100 @@ -139,8 +139,6 @@ return 8; case Object: return wordSize; - case Jsr: - return 4; default: return 0; }
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/ValueUtil.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/ValueUtil.java Tue Feb 19 17:41:10 2013 +0100 @@ -89,7 +89,7 @@ } public static Register asIntReg(Value value) { - assert value.getKind() == Kind.Int || value.getKind() == Kind.Jsr; + assert value.getKind() == Kind.Int; return asRegister(value); }
--- a/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/Constant.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/Constant.java Tue Feb 19 17:41:10 2013 +0100 @@ -134,8 +134,6 @@ return (short) primitive; case Char: return (char) primitive; - case Jsr: - return (int) primitive; case Int: return (int) primitive; case Long: @@ -163,12 +161,12 @@ /** * Returns the primitive int value this constant represents. The constant must have a - * {@link Kind#getStackKind()} of {@link Kind#Int}, or kind {@link Kind#Jsr}. + * {@link Kind#getStackKind()} of {@link Kind#Int}. * * @return the constant value */ public int asInt() { - assert getKind().getStackKind() == Kind.Int || getKind() == Kind.Jsr; + assert getKind().getStackKind() == Kind.Int; return (int) primitive; } @@ -185,13 +183,12 @@ /** * Returns the primitive long value this constant represents. The constant must have kind - * {@link Kind#Long}, a {@link Kind#getStackKind()} of {@link Kind#Int}, or kind - * {@link Kind#Jsr}. + * {@link Kind#Long}, a {@link Kind#getStackKind()} of {@link Kind#Int}. * * @return the constant value */ public long asLong() { - assert getKind() == Kind.Long || getKind().getStackKind() == Kind.Int || getKind() == Kind.Jsr; + assert getKind() == Kind.Long || getKind().getStackKind() == Kind.Int; return primitive; } @@ -365,16 +362,6 @@ } /** - * Creates a boxed address (jsr/ret address) constant. - * - * @param i the address value to box - * @return a boxed copy of {@code value} - */ - public static Constant forJsr(int i) { - return new Constant(Kind.Jsr, null, i); - } - - /** * Creates a boxed object constant. * * @param o the object value to box
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/InvokeTarget.java Tue Feb 19 17:41:10 2013 +0100 @@ -0,0 +1,29 @@ +/* + * Copyright (c) 2013, 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.api.meta; + +/** + * Represents the resolved target of an invocation. + */ +public interface InvokeTarget { +}
--- a/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/Kind.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/Kind.java Tue Feb 19 17:41:10 2013 +0100 @@ -60,9 +60,6 @@ /** The void float kind. */ Void('v', "void", false, java.lang.Void.TYPE, java.lang.Void.class), - /** Denote a bytecode address in a {@code JSR} bytecode. */ - Jsr('r', "jsr", false, null, null), - /** The non-type. */ Illegal('-', "illegal", false, null, null); @@ -318,7 +315,6 @@ return java.lang.Character.MIN_VALUE; case Short: return java.lang.Short.MIN_VALUE; - case Jsr: case Int: return java.lang.Integer.MIN_VALUE; case Long: @@ -343,7 +339,6 @@ return java.lang.Character.MAX_VALUE; case Short: return java.lang.Short.MAX_VALUE; - case Jsr: case Int: return java.lang.Integer.MAX_VALUE; case Long: @@ -367,7 +362,6 @@ case Char: case Short: return 16; - case Jsr: case Int: return 32; case Long:
--- a/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/ResolvedJavaMethod.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/ResolvedJavaMethod.java Tue Feb 19 17:41:10 2013 +0100 @@ -1,5 +1,5 @@ /* - * Copyright (c) 2009, 2012, Oracle and/or its affiliates. All rights reserved. + * 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 @@ -31,7 +31,7 @@ * Represents a resolved Java method. Methods, like fields and types, are resolved through * {@link ConstantPool constant pools}. */ -public interface ResolvedJavaMethod extends JavaMethod { +public interface ResolvedJavaMethod extends JavaMethod, InvokeTarget { /** * Returns the bytecodes of this method, if the method has code. The returned byte array does
--- a/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/ResolvedJavaType.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/ResolvedJavaType.java Tue Feb 19 17:41:10 2013 +0100 @@ -1,5 +1,5 @@ /* - * Copyright (c) 2009, 2012, Oracle and/or its affiliates. All rights reserved. + * 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 @@ -253,4 +253,9 @@ * Returns name of source file of this type. */ String getSourceFileName(); + + /** + * Returns the class file path - if available - of this type, or {@code null}. + */ + String getClassFilePath(); }
--- a/graal/com.oracle.graal.asm.amd64/src/com/oracle/graal/asm/amd64/AMD64Assembler.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.asm.amd64/src/com/oracle/graal/asm/amd64/AMD64Assembler.java Tue Feb 19 17:41:10 2013 +0100 @@ -3071,9 +3071,4 @@ public void fstp(int i) { emitx87(0xDD, 0xD8, i); } - - @Override - public void bangStack(int disp) { - movq(new Address(target.wordKind, AMD64.RSP, -disp), AMD64.rax); - } }
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/AbstractPTXAssembler.java Tue Feb 19 17:41:10 2013 +0100 @@ -0,0 +1,63 @@ +/* + * 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 com.oracle.graal.api.code.*; +import com.oracle.graal.asm.*; + +/** + * The platform-dependent base class for the PTX assembler. + */ +public abstract class AbstractPTXAssembler extends AbstractAssembler { + + public AbstractPTXAssembler(TargetDescription target) { + super(target); + } + + public static final String UNBOUND_TARGET = "L" + Integer.MAX_VALUE; + + @Override + public final void bind(Label l) { + super.bind(l); + emitString0("L" + l.toString() + ":\n"); + } + + @Override + public void align(int modulus) { + // TODO Auto-generated method stub + } + + @Override + public void jmp(Label l) { + // TODO Auto-generated method stub + } + + @Override + protected void patchJumpTarget(int branch, int jumpTarget) { + final int spaces = UNBOUND_TARGET.length(); + String targetString = String.format("L%-" + spaces + "s", jumpTarget + ";"); + int offset = "\tbra ".length(); // XXX we need a better way to figure this out + codeBuffer.emitString(targetString, branch + offset); + } + +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAsmOptions.java Tue Feb 19 17:41:10 2013 +0100 @@ -0,0 +1,27 @@ +/* + * 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; + +public class PTXAsmOptions { + // Nothing for now +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Tue Feb 19 17:41:10 2013 +0100 @@ -0,0 +1,741 @@ +/* + * 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 com.oracle.graal.api.code.*; + +public class PTXAssembler extends AbstractPTXAssembler { + + @SuppressWarnings("unused") + public PTXAssembler(TargetDescription target, RegisterConfig registerConfig) { + super(target); + } + + public final void at() { + emitString("@%p" + " " + ""); + } + + public final void add_s16(Register d, Register a, Register b) { + emitString("add.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void add_s32(Register d, Register a, Register b) { + emitString("add.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void add_s64(Register d, Register a, Register b) { + emitString("add.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void add_s16(Register d, Register a, short s16) { + emitString("add.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + ""); + } + + public final void add_s32(Register d, Register a, int s32) { + emitString("add.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + ""); + } + + public final void add_s64(Register d, Register a, long s64) { + emitString("add.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + ""); + } + + public final void add_u16(Register d, Register a, Register b) { + emitString("add.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void add_u32(Register d, Register a, Register b) { + emitString("add.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void add_u64(Register d, Register a, Register b) { + emitString("add.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void add_u16(Register d, Register a, short u16) { + emitString("add.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + ""); + } + + public final void add_u32(Register d, Register a, int u32) { + emitString("add.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); + } + + public final void add_u64(Register d, Register a, long u64) { + emitString("add.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + ""); + } + + public final void add_sat_s32(Register d, Register a, Register b) { + emitString("add.sat.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void add_sat_s32(Register d, Register a, int s32) { + emitString("add.sat.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + ""); + } + + public final void and_b16(Register d, Register a, Register b) { + emitString("and.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void and_b32(Register d, Register a, Register b) { + emitString("and.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void and_b64(Register d, Register a, Register b) { + emitString("and.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void and_b16(Register d, Register a, short b16) { + emitString("and.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b16 + ";" + ""); + } + + public final void and_b32(Register d, Register a, int b32) { + emitString("and.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b32 + ";" + ""); + } + + public final void and_b64(Register d, Register a, long b64) { + emitString("and.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b64 + ";" + ""); + } + + public final void bra(String tgt) { + emitString("bra" + " " + tgt + ";" + ""); + } + + public final void bra_uni(String tgt) { + emitString("bra.uni" + " " + tgt + ";" + ""); + } + + public final void div_s16(Register d, Register a, Register b) { + emitString("div.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void div_s32(Register d, Register a, Register b) { + emitString("div.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void div_s64(Register d, Register a, Register b) { + emitString("div.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void div_s16(Register d, Register a, short s16) { + emitString("div.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + ""); + } + + public final void div_s32(Register d, Register a, int s32) { + emitString("div.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + ""); + } + + public final void div_s64(Register d, Register a, long s64) { + emitString("div.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + ""); + } + + public final void div_u16(Register d, Register a, Register b) { + emitString("div.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void div_u32(Register d, Register a, Register b) { + emitString("div.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void div_u64(Register d, Register a, Register b) { + emitString("div.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void div_u16(Register d, Register a, short u16) { + emitString("div.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + ""); + } + + public final void div_u32(Register d, Register a, int u32) { + emitString("div.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); + } + + public final void div_u64(Register d, Register a, long u64) { + emitString("div.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + ""); + } + + public final void exit() { + emitString("exit;" + " " + ""); + } + + public final void ld_global_b8(Register d, Register a, int immOff) { + emitString("ld.global.b8" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); + } + + public final void ld_global_b16(Register d, Register a, int immOff) { + emitString("ld.global.b16" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); + } + + public final void ld_global_b32(Register d, Register a, int immOff) { + emitString("ld.global.b32" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); + } + + public final void ld_global_b64(Register d, Register a, int immOff) { + emitString("ld.global.b64" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); + } + + public final void ld_global_u8(Register d, Register a, int immOff) { + emitString("ld.global.u8" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); + } + + public final void ld_global_u16(Register d, Register a, int immOff) { + emitString("ld.global.u16" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); + } + + public final void ld_global_u32(Register d, Register a, int immOff) { + emitString("ld.global.u32" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); + } + + public final void ld_global_u64(Register d, Register a, int immOff) { + emitString("ld.global.u64" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); + } + + public final void ld_global_s8(Register d, Register a, int immOff) { + emitString("ld.global.s8" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); + } + + public final void ld_global_s16(Register d, Register a, int immOff) { + emitString("ld.global.s16" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); + } + + public final void ld_global_s32(Register d, Register a, int immOff) { + emitString("ld.global.s32" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); + } + + public final void ld_global_s64(Register d, Register a, int immOff) { + emitString("ld.global.s64" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); + } + + public final void ld_global_f32(Register d, Register a, int immOff) { + emitString("ld.global.f32" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); + } + + public final void ld_global_f64(Register d, Register a, int immOff) { + emitString("ld.global.f64" + " " + "%r" + d.encoding() + ", [%r" + a.encoding() + " + " + immOff + "]" + ";" + ""); + } + + public final void mov_b16(Register d, Register a) { + emitString("mov.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); + } + + public final void mov_b32(Register d, Register a) { + emitString("mov.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); + } + + public final void mov_b64(Register d, Register a) { + emitString("mov.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); + } + + public final void mov_u16(Register d, Register a) { + emitString("mov.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); + } + + public final void mov_u32(Register d, Register a) { + emitString("mov.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); + } + + public final void mov_u64(Register d, Register a) { + emitString("mov.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); + } + + public final void mov_s16(Register d, Register a) { + emitString("mov.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); + } + + public final void mov_s32(Register d, Register a) { + emitString("mov.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); + } + + public final void mov_s64(Register d, Register a) { + emitString("mov.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); + } + + public final void mov_f32(Register d, Register a) { + emitString("mov.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); + } + + public final void mov_f64(Register d, Register a) { + emitString("mov.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); + } + + public final void mov_b16(Register d, short b16) { + emitString("mov.b16" + " " + "%r" + d.encoding() + ", " + b16 + ";" + ""); + } + + public final void mov_b32(Register d, int b32) { + emitString("mov.b32" + " " + "%r" + d.encoding() + ", " + b32 + ";" + ""); + } + + public final void mov_b64(Register d, long b64) { + emitString("mov.b64" + " " + "%r" + d.encoding() + ", " + b64 + ";" + ""); + } + + public final void mov_u16(Register d, short u16) { + emitString("mov.u16" + " " + "%r" + d.encoding() + ", " + u16 + ";" + ""); + } + + public final void mov_u32(Register d, int u32) { + emitString("mov.u32" + " " + "%r" + d.encoding() + ", " + u32 + ";" + ""); + } + + public final void mov_u64(Register d, long u64) { + emitString("mov.u64" + " " + "%r" + d.encoding() + ", " + u64 + ";" + ""); + } + + public final void mov_s16(Register d, short s16) { + emitString("mov.s16" + " " + "%r" + d.encoding() + ", " + s16 + ";" + ""); + } + + public final void mov_s32(Register d, int s32) { + emitString("mov.s32" + " " + "%r" + d.encoding() + ", " + s32 + ";" + ""); + } + + public final void mov_s64(Register d, long s64) { + emitString("mov.s64" + " " + "%r" + d.encoding() + ", " + s64 + ";" + ""); + } + + public final void mov_f32(Register d, float f32) { + emitString("mov.f32" + " " + "%r" + d.encoding() + ", " + f32 + ";" + ""); + } + + public final void mov_f64(Register d, double f64) { + emitString("mov.f64" + " " + "%r" + d.encoding() + ", " + f64 + ";" + ""); + } + + public final void mul_s16(Register d, Register a, Register b) { + emitString("mul.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void mul_s32(Register d, Register a, Register b) { + emitString("mul.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void mul_s64(Register d, Register a, Register b) { + emitString("mul.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void mul_s16(Register d, Register a, short s16) { + emitString("mul.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + ""); + } + + public final void mul_s32(Register d, Register a, int s32) { + emitString("mul.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + ""); + } + + public final void mul_s64(Register d, Register a, long s64) { + emitString("mul.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + ""); + } + + public final void mul_u16(Register d, Register a, Register b) { + emitString("mul.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void mul_u32(Register d, Register a, Register b) { + emitString("mul.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void mul_u64(Register d, Register a, Register b) { + emitString("mul.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void mul_u16(Register d, Register a, short u16) { + emitString("mul.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + ""); + } + + public final void mul_u32(Register d, Register a, int u32) { + emitString("mul.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); + } + + public final void mul_u64(Register d, Register a, long u64) { + emitString("mul.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + ""); + } + + public final void neg_s16(Register d, Register a) { + emitString("neg.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); + } + + public final void neg_s32(Register d, Register a) { + emitString("neg.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); + } + + public final void neg_s64(Register d, Register a) { + emitString("neg.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); + } + + public final void popc_b32(Register d, Register a) { + emitString("popc.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); + } + + public final void popc_b64(Register d, Register a) { + emitString("popc.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); + } + + public final void rem_s16(Register d, Register a, Register b) { + emitString("rem.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void rem_s32(Register d, Register a, Register b) { + emitString("rem.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void rem_s64(Register d, Register a, Register b) { + emitString("rem.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void rem_s16(Register d, Register a, short s16) { + emitString("rem.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + ""); + } + + public final void rem_s32(Register d, Register a, int s32) { + emitString("rem.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + ""); + } + + public final void rem_s64(Register d, Register a, long s64) { + emitString("rem.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + ""); + } + + public final void rem_u16(Register d, Register a, Register b) { + emitString("rem.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void rem_u32(Register d, Register a, Register b) { + emitString("rem.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void rem_u64(Register d, Register a, Register b) { + emitString("rem.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void rem_u16(Register d, Register a, short u16) { + emitString("rem.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + ""); + } + + public final void rem_u32(Register d, Register a, int u32) { + emitString("rem.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); + } + + public final void rem_u64(Register d, Register a, long u64) { + emitString("rem.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + ""); + } + + public final void ret() { + emitString("ret;" + " " + ""); + } + + public final void ret_uni() { + emitString("ret.uni;" + " " + ""); + } + + public final void setp_eq_s32(Register a, Register b) { + emitString("setp.eq.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_ne_s32(Register a, Register b) { + emitString("setp.ne.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_lt_s32(Register a, Register b) { + emitString("setp.lt.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_le_s32(Register a, Register b) { + emitString("setp.le.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_gt_s32(Register a, Register b) { + emitString("setp.gt.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_ge_s32(Register a, Register b) { + emitString("setp.ge.s32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_eq_s32(Register a, int s32) { + emitString("setp.eq.s32" + " " + "%p" + ", %r" + a.encoding() + ", " + s32 + ";" + ""); + } + + public final void setp_ne_s32(Register a, int s32) { + emitString("setp.ne.s32" + " " + "%p" + ", %r" + a.encoding() + ", " + s32 + ";" + ""); + } + + public final void setp_lt_s32(Register a, int s32) { + emitString("setp.lt.s32" + " " + "%p" + ", %r" + a.encoding() + ", " + s32 + ";" + ""); + } + + public final void setp_le_s32(Register a, int s32) { + emitString("setp.le.s32" + " " + "%p" + ", %r" + a.encoding() + ", " + s32 + ";" + ""); + } + + public final void setp_gt_s32(Register a, int s32) { + emitString("setp.gt.s32" + " " + "%p" + ", %r" + a.encoding() + ", " + s32 + ";" + ""); + } + + public final void setp_ge_s32(Register a, int s32) { + emitString("setp.ge.s32" + " " + "%p" + ", %r" + a.encoding() + ", " + s32 + ";" + ""); + } + + public final void setp_eq_s32(int s32, Register b) { + emitString("setp.eq.s32" + " " + "%p" + ", " + s32 + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_ne_s32(int s32, Register b) { + emitString("setp.ne.s32" + " " + "%p" + ", " + s32 + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_lt_s32(int s32, Register b) { + emitString("setp.lt.s32" + " " + "%p" + ", " + s32 + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_le_s32(int s32, Register b) { + emitString("setp.le.s32" + " " + "%p" + ", " + s32 + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_gt_s32(int s32, Register b) { + emitString("setp.gt.s32" + " " + "%p" + ", " + s32 + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_ge_s32(int s32, Register b) { + emitString("setp.ge.s32" + " " + "%p" + ", " + s32 + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_eq_u32(Register a, Register b) { + emitString("setp.eq.u32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_ne_u32(Register a, Register b) { + emitString("setp.ne.u32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_lt_u32(Register a, Register b) { + emitString("setp.lt.u32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_le_u32(Register a, Register b) { + emitString("setp.le.u32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_gt_u32(Register a, Register b) { + emitString("setp.gt.u32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_ge_u32(Register a, Register b) { + emitString("setp.ge.u32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_eq_u32(Register a, int u32) { + emitString("setp.eq.u32" + " " + "%p" + ", %r" + a.encoding() + ", " + u32 + ";" + ""); + } + + public final void setp_ne_u32(Register a, int u32) { + emitString("setp.ne.u32" + " " + "%p" + ", %r" + a.encoding() + ", " + u32 + ";" + ""); + } + + public final void setp_lt_u32(Register a, int u32) { + emitString("setp.lt.u32" + " " + "%p" + ", %r" + a.encoding() + ", " + u32 + ";" + ""); + } + + public final void setp_le_u32(Register a, int u32) { + emitString("setp.le.u32" + " " + "%p" + ", %r" + a.encoding() + ", " + u32 + ";" + ""); + } + + public final void setp_gt_u32(Register a, int u32) { + emitString("setp.gt.u32" + " " + "%p" + ", %r" + a.encoding() + ", " + u32 + ";" + ""); + } + + public final void setp_ge_u32(Register a, int u32) { + emitString("setp.ge.u32" + " " + "%p" + ", %r" + a.encoding() + ", " + u32 + ";" + ""); + } + + public final void setp_eq_u32(int u32, Register b) { + emitString("setp.eq.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_ne_u32(int u32, Register b) { + emitString("setp.ne.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_lt_u32(int u32, Register b) { + emitString("setp.lt.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_le_u32(int u32, Register b) { + emitString("setp.le.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_gt_u32(int u32, Register b) { + emitString("setp.gt.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + ""); + } + + public final void setp_ge_u32(int u32, Register b) { + emitString("setp.ge.u32" + " " + "%p" + ", " + u32 + ", %r" + b.encoding() + ";" + ""); + } + + public final void shr_s16(Register d, Register a, Register b) { + emitString("shr.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void shr_s32(Register d, Register a, Register b) { + emitString("shr.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void shr_s64(Register d, Register a, Register b) { + emitString("shr.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void shr_s16(Register d, Register a, int u32) { + emitString("shr.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); + } + + public final void shr_s32(Register d, Register a, int u32) { + emitString("shr.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); + } + + public final void shr_s64(Register d, Register a, int u32) { + emitString("shr.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); + } + + public final void shr_u16(Register d, Register a, Register b) { + emitString("shr.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void shr_u32(Register d, Register a, Register b) { + emitString("shr.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void shr_u64(Register d, Register a, Register b) { + emitString("shr.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void shr_u16(Register d, Register a, int u32) { + emitString("shr.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); + } + + public final void shr_u32(Register d, Register a, int u32) { + emitString("shr.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); + } + + public final void shr_u64(Register d, Register a, int u32) { + emitString("shr.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); + } + + public final void st_global_b8(Register a, int immOff, Register b) { + emitString("st.global.b8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); + } + + public final void st_global_b16(Register a, int immOff, Register b) { + emitString("st.global.b16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); + } + + public final void st_global_b32(Register a, int immOff, Register b) { + emitString("st.global.b32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); + } + + public final void st_global_b64(Register a, int immOff, Register b) { + emitString("st.global.b64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); + } + + public final void st_global_u8(Register a, int immOff, Register b) { + emitString("st.global.u8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); + } + + public final void st_global_u16(Register a, int immOff, Register b) { + emitString("st.global.u16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); + } + + public final void st_global_u32(Register a, int immOff, Register b) { + emitString("st.global.u32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); + } + + public final void st_global_u64(Register a, int immOff, Register b) { + emitString("st.global.u64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); + } + + public final void st_global_s8(Register a, int immOff, Register b) { + emitString("st.global.s8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); + } + + public final void st_global_s16(Register a, int immOff, Register b) { + emitString("st.global.s16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); + } + + public final void st_global_s32(Register a, int immOff, Register b) { + emitString("st.global.s32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); + } + + public final void st_global_s64(Register a, int immOff, Register b) { + emitString("st.global.s64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); + } + + public final void st_global_f32(Register a, int immOff, Register b) { + emitString("st.global.f32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); + } + + public final void st_global_f64(Register a, int immOff, Register b) { + emitString("st.global.f64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); + } + + public final void sub_s16(Register d, Register a, Register b) { + emitString("sub.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void sub_s32(Register d, Register a, Register b) { + emitString("sub.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void sub_s64(Register d, Register a, Register b) { + emitString("sub.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void sub_s16(Register d, Register a, short s16) { + emitString("sub.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + ""); + } + + public final void sub_s32(Register d, Register a, int s32) { + emitString("sub.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + ""); + } + + public final void sub_s64(Register d, Register a, long s64) { + emitString("sub.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + ""); + } + + public final void sub_s16(Register d, short s16, Register b) { + emitString("sub.s16" + " " + "%r" + d.encoding() + ", " + s16 + ", %r" + b.encoding() + ";" + ""); + } + + public final void sub_s32(Register d, int s32, Register b) { + emitString("sub.s32" + " " + "%r" + d.encoding() + ", " + s32 + ", %r" + b.encoding() + ";" + ""); + } + + public final void sub_s64(Register d, long s64, Register b) { + emitString("sub.s64" + " " + "%r" + d.encoding() + ", " + s64 + ", %r" + b.encoding() + ";" + ""); + } + + public final void sub_sat_s32(Register d, Register a, Register b) { + emitString("sub.sat.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + } + + public final void sub_sat_s32(Register d, Register a, int s32) { + emitString("sub.sat.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + ""); + } + + public final void sub_sat_s32(Register d, int s32, Register b) { + emitString("sub.sat.s32" + " " + "%r" + d.encoding() + ", " + s32 + ", %r" + b.encoding() + ";" + ""); + } +}
--- a/graal/com.oracle.graal.asm.sparc/src/com/oracle/graal/asm/sparc/SPARCAssembler.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.asm.sparc/src/com/oracle/graal/asm/sparc/SPARCAssembler.java Tue Feb 19 17:41:10 2013 +0100 @@ -52,9 +52,4 @@ protected void patchJumpTarget(int branch, int jumpTarget) { // SPARC: Implement patching of jump target. } - - @Override - public void bangStack(int disp) { - // SPARC: Implement stack banging. - } }
--- a/graal/com.oracle.graal.asm/src/com/oracle/graal/asm/AbstractAssembler.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.asm/src/com/oracle/graal/asm/AbstractAssembler.java Tue Feb 19 17:41:10 2013 +0100 @@ -44,7 +44,7 @@ } } - public final void bind(Label l) { + public void bind(Label l) { assert !l.isBound() : "can bind label only once"; l.bind(codeBuffer.position()); l.patchInstructions(this); @@ -56,14 +56,6 @@ protected abstract void patchJumpTarget(int branch, int jumpTarget); - /** - * Emits instruction(s) that access an address specified by a given displacement from the stack - * pointer in the direction that the stack grows (which is down on most architectures). - * - * @param disp the displacement from the stack pointer at which the stack should be accessed - */ - public abstract void bangStack(int disp); - protected final void emitByte(int x) { codeBuffer.emitByte(x); } @@ -79,4 +71,16 @@ protected final void emitLong(long x) { codeBuffer.emitLong(x); } + + /** + * Some GPU architectures have a text based encoding. + */ + protected final void emitString(String x) { + codeBuffer.emitString(x); + } + + // XXX for pretty-printing + protected final void emitString0(String x) { + codeBuffer.emitString0(x); + } }
--- a/graal/com.oracle.graal.asm/src/com/oracle/graal/asm/Buffer.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.asm/src/com/oracle/graal/asm/Buffer.java Tue Feb 19 17:41:10 2013 +0100 @@ -37,10 +37,6 @@ data = new byte[AsmOptions.InitialCodeBufferSize]; } - public void reset() { - position = 0; - } - public int position() { return position; } @@ -106,6 +102,27 @@ position = emitLong(b, position); } + private static final String NEWLINE = System.getProperty("line.separator"); + + public void emitString(String s) { + position = emitString("\t", position); // XXX REMOVE ME pretty-printing + position = emitString(s, position); + position = emitString(NEWLINE, position); + } + + // XXX for pretty-printing + public void emitString0(String s) { + emitBytes(s.getBytes(), 0, s.length()); + } + + public int emitBytes(byte[] arr, int pos) { + final int len = arr.length; + final int newPos = pos + len; + ensureSize(newPos); + System.arraycopy(arr, 0, data, pos, len); + return newPos; + } + public int emitByte(int b, int pos) { assert NumUtil.isUByte(b); int newPos = pos + 1; @@ -120,6 +137,10 @@ public abstract int emitLong(long b, int pos); + public int emitString(String s, int pos) { + return emitBytes(s.getBytes(), pos); + } + public int getByte(int pos) { return data[pos] & 0xff; }
--- a/graal/com.oracle.graal.compiler.amd64/src/com/oracle/graal/compiler/amd64/AMD64LIRGenerator.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.compiler.amd64/src/com/oracle/graal/compiler/amd64/AMD64LIRGenerator.java Tue Feb 19 17:41:10 2013 +0100 @@ -314,7 +314,6 @@ mirrored = false; } switch (left.getKind().getStackKind()) { - case Jsr: case Int: append(new CompareOp(ICMP, left, right)); break;
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java Tue Feb 19 17:41:10 2013 +0100 @@ -0,0 +1,115 @@ +/* + * 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.compiler.ptx; + +import com.oracle.graal.api.code.*; +import com.oracle.graal.api.meta.*; +import com.oracle.graal.asm.*; +import com.oracle.graal.asm.ptx.*; +import com.oracle.graal.compiler.gen.*; +import com.oracle.graal.compiler.target.*; +import com.oracle.graal.lir.*; +import com.oracle.graal.lir.asm.*; +import com.oracle.graal.nodes.*; +import com.oracle.graal.phases.*; + +/** + * PTX specific backend. + */ +public class PTXBackend extends Backend { + + public PTXBackend(CodeCacheProvider runtime, TargetDescription target) { + super(runtime, target); + } + + @Override + public LIRGenerator newLIRGenerator(StructuredGraph graph, FrameMap frameMap, ResolvedJavaMethod method, LIR lir) { + return new PTXLIRGenerator(graph, runtime(), target, frameMap, method, lir); + } + + class HotSpotFrameContext implements FrameContext { + + @Override + public void enter(TargetMethodAssembler tasm) { + Buffer codeBuffer = tasm.asm.codeBuffer; + codeBuffer.emitString(".version 1.4"); + codeBuffer.emitString(".target sm_10"); + // codeBuffer.emitString(".address_size 32"); // PTX ISA version 2.3 + } + + @Override + public void leave(TargetMethodAssembler tasm) { + } + } + + @Override + public TargetMethodAssembler newAssembler(FrameMap frameMap, LIR lir) { + // Omit the frame if the method: + // - has no spill slots or other slots allocated during register allocation + // - has no callee-saved registers + // - has no incoming arguments passed on the stack + // - has no instructions with debug info + boolean omitFrame = GraalOptions.CanOmitFrame && frameMap.frameSize() == frameMap.initialFrameSize && frameMap.registerConfig.getCalleeSaveLayout().registers.length == 0 && + !lir.hasArgInCallerFrame() && !lir.hasDebugInfo(); + + AbstractAssembler masm = new PTXAssembler(target, frameMap.registerConfig); + HotSpotFrameContext frameContext = omitFrame ? null : new HotSpotFrameContext(); + TargetMethodAssembler tasm = new TargetMethodAssembler(target, runtime(), frameMap, masm, frameContext, lir.stubs); + tasm.setFrameSize(frameMap.frameSize()); + tasm.compilationResult.setCustomStackAreaOffset(frameMap.offsetToCustomArea()); + return tasm; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, ResolvedJavaMethod method, LIR lir) { + // Emit the prologue + final String name = method.getName(); + Buffer codeBuffer = tasm.asm.codeBuffer; + codeBuffer.emitString0(".entry " + name + " ("); + codeBuffer.emitString(""); + + Signature signature = method.getSignature(); + for (int i = 0; i < signature.getParameterCount(false); i++) { + System.err.println(i + ": " + signature.getParameterKind(i)); + String param = ".param .u32 param" + i; + codeBuffer.emitString(param); + } + + codeBuffer.emitString0(") {"); + codeBuffer.emitString(""); + + // XXX For now declare one predicate and all registers + codeBuffer.emitString(" .reg .pred %p;"); + codeBuffer.emitString(" .reg .u32 %r<16>;"); + + // Emit code for the LIR + lir.emitCode(tasm); + + // Emit the epilogue + codeBuffer.emitString0("}"); + codeBuffer.emitString(""); + + byte[] data = codeBuffer.copyData(0, codeBuffer.position()); + System.err.println(new String(data)); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXDeoptimizationStub.java Tue Feb 19 17:41:10 2013 +0100 @@ -0,0 +1,61 @@ +/* + * 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.compiler.ptx; + +import com.oracle.graal.api.code.*; +import com.oracle.graal.api.code.RuntimeCallTarget.*; +import com.oracle.graal.api.meta.*; +import com.oracle.graal.asm.*; +import com.oracle.graal.asm.ptx.*; +import com.oracle.graal.lir.*; +import com.oracle.graal.lir.ptx.*; +import com.oracle.graal.lir.asm.*; + +public class PTXDeoptimizationStub extends PTXCode { + + public static final Descriptor DEOPTIMIZE = new Descriptor("deoptimize", true, void.class); + public static final Descriptor SET_DEOPT_INFO = new Descriptor("setDeoptInfo", true, void.class, Object.class); + + public final Label label = new Label(); + public final LIRFrameState info; + public final DeoptimizationAction action; + public final DeoptimizationReason reason; + public final Object deoptInfo; + + public PTXDeoptimizationStub(DeoptimizationAction action, DeoptimizationReason reason, LIRFrameState info, Object deoptInfo) { + this.action = action; + this.reason = reason; + this.info = info; + this.deoptInfo = deoptInfo; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + masm.exit(); + } + + @Override + public String description() { + return "deopt stub[reason=" + reason + ", action=" + action + "]"; + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java Tue Feb 19 17:41:10 2013 +0100 @@ -0,0 +1,485 @@ +/* + * 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.compiler.ptx; + +import static com.oracle.graal.api.code.ValueUtil.*; +import static com.oracle.graal.lir.ptx.PTXArithmetic.*; +import static com.oracle.graal.lir.ptx.PTXCompare.*; +import static com.oracle.graal.lir.ptx.PTXBitManipulationOp.IntrinsicOpcode.*; + +import com.oracle.graal.api.code.*; +import com.oracle.graal.api.meta.*; +import com.oracle.graal.asm.*; +import com.oracle.graal.compiler.gen.*; +import com.oracle.graal.compiler.target.*; +import com.oracle.graal.graph.*; +import com.oracle.graal.lir.*; +import com.oracle.graal.lir.StandardOp.JumpOp; +import com.oracle.graal.lir.ptx.PTXBitManipulationOp; +import com.oracle.graal.lir.ptx.PTXCompare.CompareOp; +import com.oracle.graal.lir.ptx.PTXControlFlow.BranchOp; +import com.oracle.graal.lir.ptx.PTXControlFlow.ReturnOp; +import com.oracle.graal.lir.ptx.PTXMove.LoadOp; +import com.oracle.graal.lir.ptx.PTXMove.MoveFromRegOp; +import com.oracle.graal.lir.ptx.PTXMove.MoveToRegOp; +import com.oracle.graal.lir.ptx.PTXMove.StoreOp; +import com.oracle.graal.nodes.*; +import com.oracle.graal.nodes.calc.Condition; +import com.oracle.graal.nodes.calc.ConvertNode; +import com.oracle.graal.nodes.extended.IndexedLocationNode; +import com.oracle.graal.nodes.extended.LocationNode; +import com.oracle.graal.nodes.java.*; + +/** + * This class implements the PTX specific portion of the LIR generator. + */ +public class PTXLIRGenerator extends LIRGenerator { + + public static class PTXSpillMoveFactory implements LIR.SpillMoveFactory { + + @Override + public LIRInstruction createMove(Value result, Value input) { + throw new InternalError("NYI"); + } + } + + public PTXLIRGenerator(StructuredGraph graph, CodeCacheProvider runtime, TargetDescription target, FrameMap frameMap, ResolvedJavaMethod method, LIR lir) { + super(graph, runtime, target, frameMap, method, lir); + lir.spillMoveFactory = new PTXSpillMoveFactory(); + } + + @Override + protected void emitNode(ValueNode node) { + if (node instanceof LIRGenLowerable) { + ((LIRGenLowerable) node).generate(this); + } else { + super.emitNode(node); + } + } + + @Override + public boolean canStoreConstant(Constant c) { + // Operand b must be in the .reg state space. + return false; + } + + @Override + public boolean canInlineConstant(Constant c) { + switch (c.getKind()) { + case Long: + return NumUtil.isInt(c.asLong()) && !runtime.needsDataPatch(c); + case Object: + return c.isNull(); + default: + return true; + } + } + + @Override + public Address makeAddress(LocationNode location, ValueNode object) { + Value base = operand(object); + Value index = Value.ILLEGAL; + int scale = 1; + int displacement = location.displacement(); + + if (isConstant(base)) { + if (asConstant(base).isNull()) { + base = Value.ILLEGAL; + } else if (asConstant(base).getKind() != Kind.Object) { + long newDisplacement = displacement + asConstant(base).asLong(); + if (NumUtil.isInt(newDisplacement)) { + assert !runtime.needsDataPatch(asConstant(base)); + displacement = (int) newDisplacement; + base = Value.ILLEGAL; + } else { + Value newBase = newVariable(Kind.Long); + emitMove(base, newBase); + base = newBase; + } + } + } + + if (location instanceof IndexedLocationNode) { + IndexedLocationNode indexedLoc = (IndexedLocationNode) location; + + index = operand(indexedLoc.index()); + if (indexedLoc.indexScalingEnabled()) { + scale = target().sizeInBytes(location.getValueKind()); + } + if (isConstant(index)) { + long newDisplacement = displacement + asConstant(index).asLong() * scale; + // only use the constant index if the resulting displacement fits into a 32 bit + // offset + if (NumUtil.isInt(newDisplacement)) { + displacement = (int) newDisplacement; + index = Value.ILLEGAL; + } else { + // create a temporary variable for the index, the pointer load cannot handle a + // constant index + Value newIndex = newVariable(Kind.Long); + emitMove(index, newIndex); + index = newIndex; + } + } + } + + return new Address(location.getValueKind(), base, index, Address.Scale.fromInt(scale), displacement); + } + + @Override + public Variable emitMove(Value input) { + Variable result = newVariable(input.getKind()); + emitMove(input, result); + return result; + } + + @Override + public void emitMove(Value src, Value dst) { + if (isRegister(src) || isStackSlot(dst)) { + append(new MoveFromRegOp(dst, src)); + } else { + append(new MoveToRegOp(dst, src)); + } + } + + @Override + public Variable emitLoad(Value loadAddress, boolean canTrap) { + Variable result = newVariable(loadAddress.getKind()); + append(new LoadOp(result, loadAddress, canTrap ? state() : null)); + return result; + } + + @Override + public void emitStore(Value storeAddress, Value inputVal, boolean canTrap) { + Value input = loadForStore(inputVal, storeAddress.getKind()); + append(new StoreOp(storeAddress, input, canTrap ? state() : null)); + } + + @Override + public Variable emitLea(Value address) { + throw new InternalError("NYI"); + } + + @Override + public void emitJump(LabelRef label, LIRFrameState info) { + append(new JumpOp(label, info)); + } + + @Override + public void emitCompareBranch(Value left, Value right, Condition cond, boolean unorderedIsTrue, LabelRef label, LIRFrameState info) { + switch (left.getKind().getStackKind()) { + case Int: + append(new CompareOp(ICMP, cond, left, right)); + append(new BranchOp(cond, label, info)); + break; + case Object: + append(new CompareOp(ACMP, cond, left, right)); + append(new BranchOp(cond, label, info)); + break; + default: + throw GraalInternalError.shouldNotReachHere("" + left.getKind()); + } + } + + @Override + public void emitIntegerTestBranch(Value left, Value right, boolean negated, LabelRef label, LIRFrameState info) { + throw new InternalError("NYI"); + } + + @Override + public Variable emitConditionalMove(Value left, Value right, Condition cond, boolean unorderedIsTrue, Value trueValue, Value falseValue) { + throw new InternalError("NYI"); + } + + @Override + public Variable emitIntegerTestMove(Value left, Value right, Value trueValue, Value falseValue) { + throw new InternalError("NYI"); + } + + @Override + public Variable emitNegate(Value input) { + Variable result = newVariable(input.getKind()); + switch (input.getKind()) { + case Int: + append(new Op1Stack(INEG, result, input)); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + return result; + } + + @Override + public Variable emitAdd(Value a, Value b) { + Variable result = newVariable(a.getKind()); + switch (a.getKind()) { + case Int: + append(new Op2Stack(IADD, result, a, loadNonConst(b))); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + return result; + } + + @Override + public Variable emitSub(Value a, Value b) { + Variable result = newVariable(a.getKind()); + switch (a.getKind()) { + case Int: + append(new Op2Stack(ISUB, result, a, loadNonConst(b))); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + return result; + } + + @Override + public Variable emitMul(Value a, Value b) { + Variable result = newVariable(a.getKind()); + switch (a.getKind()) { + case Int: + append(new Op2Reg(IMUL, result, a, loadNonConst(b))); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + return result; + } + + @Override + protected boolean peephole(ValueNode valueNode) { + // No peephole optimizations for now + return false; + } + + @Override + public Value emitDiv(Value a, Value b) { + throw new InternalError("NYI"); + } + + @Override + public Value emitRem(Value a, Value b) { + throw new InternalError("NYI"); + } + + @Override + public Variable emitUDiv(Value a, Value b) { + throw new InternalError("NYI"); + } + + @Override + public Variable emitURem(Value a, Value b) { + throw new InternalError("NYI"); + } + + @Override + public Variable emitAnd(Value a, Value b) { + Variable result = newVariable(a.getKind()); + switch (a.getKind()) { + case Int: + append(new Op2Stack(IAND, result, a, loadNonConst(b))); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + return result; + } + + @Override + public Variable emitOr(Value a, Value b) { + throw new InternalError("NYI"); + } + + @Override + public Variable emitXor(Value a, Value b) { + throw new InternalError("NYI"); + } + + @Override + public Variable emitShl(Value a, Value b) { + throw new InternalError("NYI"); + } + + @Override + public Variable emitShr(Value a, Value b) { + throw new InternalError("NYI"); + } + + @Override + public Variable emitUShr(Value a, Value b) { + Variable result = newVariable(a.getKind()); + switch (a.getKind()) { + case Int: + append(new ShiftOp(IUSHR, result, a, b)); + break; + default: + GraalInternalError.shouldNotReachHere(); + } + return result; + } + + @Override + public Variable emitConvert(ConvertNode.Op opcode, Value inputVal) { + throw new InternalError("NYI"); + } + + @Override + public void emitDeoptimizeOnOverflow(DeoptimizationAction action, DeoptimizationReason reason, Object deoptInfo) { + throw new InternalError("NYI"); + } + + @Override + public void emitDeoptimize(DeoptimizationAction action, DeoptimizationReason reason, Object deoptInfo) { + throw new InternalError("NYI"); + } + + @Override + public void emitMembar(int barriers) { + throw new InternalError("NYI"); + } + + @Override + protected void emitDirectCall(DirectCallTargetNode callTarget, Value result, Value[] parameters, Value[] temps, LIRFrameState callState) { + throw new InternalError("NYI"); + } + + @Override + protected void emitIndirectCall(IndirectCallTargetNode callTarget, Value result, Value[] parameters, Value[] temps, LIRFrameState callState) { + throw new InternalError("NYI"); + } + + @Override + protected void emitCall(RuntimeCallTarget callTarget, Value result, Value[] arguments, Value[] temps, Value targetAddress, LIRFrameState info) { + throw new InternalError("NYI"); + } + + @Override + public void emitBitCount(Variable result, Value value) { + if (value.getKind().getStackKind() == Kind.Int) { + append(new PTXBitManipulationOp(IPOPCNT, result, value)); + } else { + append(new PTXBitManipulationOp(LPOPCNT, result, value)); + } + } + + @Override + public void emitBitScanForward(Variable result, Value value) { + throw new InternalError("NYI"); + } + + @Override + public void emitBitScanReverse(Variable result, Value value) { + throw new InternalError("NYI"); + } + + @Override + public void emitMathAbs(Variable result, Variable input) { + throw new InternalError("NYI"); + } + + @Override + public void emitMathSqrt(Variable result, Variable input) { + throw new InternalError("NYI"); + } + + @Override + public void emitMathLog(Variable result, Variable input, boolean base10) { + throw new InternalError("NYI"); + } + + @Override + public void emitMathCos(Variable result, Variable input) { + throw new InternalError("NYI"); + } + + @Override + public void emitMathSin(Variable result, Variable input) { + throw new InternalError("NYI"); + } + + @Override + public void emitMathTan(Variable result, Variable input) { + throw new InternalError("NYI"); + } + + @Override + public void emitByteSwap(Variable result, Value input) { + throw new InternalError("NYI"); + } + + @Override + protected void emitReturn(Value input) { + append(new ReturnOp(input)); + } + + @Override + protected void emitSequentialSwitch(Constant[] keyConstants, LabelRef[] keyTargets, LabelRef defaultTarget, Value key) { + throw new InternalError("NYI"); + } + + @Override + protected void emitSwitchRanges(int[] lowKeys, int[] highKeys, LabelRef[] targets, LabelRef defaultTarget, Value key) { + throw new InternalError("NYI"); + } + + @Override + protected void emitTableSwitch(int lowKey, LabelRef defaultTarget, LabelRef[] targets, Value key) { + throw new InternalError("NYI"); + } + + @Override + protected LabelRef createDeoptStub(DeoptimizationAction action, DeoptimizationReason reason, LIRFrameState info, Object deoptInfo) { + assert info.topFrame.getBCI() >= 0 : "invalid bci for deopt framestate"; + PTXDeoptimizationStub stub = new PTXDeoptimizationStub(action, reason, info, deoptInfo); + lir.stubs.add(stub); + return LabelRef.forLabel(stub.label); + } + + @Override + protected void emitNullCheckGuard(ValueNode object) { + throw new InternalError("NYI"); + } + + @Override + public void visitCompareAndSwap(CompareAndSwapNode node) { + throw new InternalError("NYI"); + } + + @Override + public void visitBreakpointNode(BreakpointNode node) { + throw new InternalError("NYI"); + } + + @Override + public void visitExceptionObject(ExceptionObjectNode i) { + throw new InternalError("NYI"); + } + + @Override + public void visitSafepointNode(SafepointNode i) { + throw new InternalError("NYI"); + } +}
--- a/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/gen/LIRGenerator.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/gen/LIRGenerator.java Tue Feb 19 17:41:10 2013 +0100 @@ -151,7 +151,6 @@ public Variable newVariable(Kind kind) { Kind stackKind = kind.getStackKind(); switch (stackKind) { - case Jsr: case Int: case Long: case Object:
--- a/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/target/Backend.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/target/Backend.java Tue Feb 19 17:41:10 2013 +0100 @@ -28,7 +28,6 @@ import com.oracle.graal.lir.*; import com.oracle.graal.lir.asm.*; import com.oracle.graal.nodes.*; -import com.oracle.graal.phases.*; /** * The {@code Backend} class represents a compiler backend for Graal. @@ -56,30 +55,6 @@ public abstract TargetMethodAssembler newAssembler(FrameMap frameMap, LIR lir); /** - * Emits code to do stack overflow checking. - * - * @param afterFrameInit specifies if the stack pointer has already been adjusted to allocate - * the current frame - */ - protected static void emitStackOverflowCheck(TargetMethodAssembler tasm, boolean afterFrameInit) { - if (GraalOptions.StackShadowPages > 0) { - int frameSize = tasm.frameMap.frameSize(); - if (frameSize > 0) { - int lastFramePage = frameSize / tasm.target.pageSize; - // emit multiple stack bangs for methods with frames larger than a page - for (int i = 0; i <= lastFramePage; i++) { - int disp = (i + GraalOptions.StackShadowPages) * tasm.target.pageSize; - if (afterFrameInit) { - disp -= frameSize; - } - tasm.blockComment("[stack overflow check]"); - tasm.asm.bangStack(disp); - } - } - } - } - - /** * Emits the code for a given method. This includes any architecture/runtime specific * prefix/suffix. A prefix typically contains the code for setting up the frame, spilling * callee-save registers, stack overflow checking, handling multiple entry points etc. A suffix
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64DirectCallOp.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64DirectCallOp.java Tue Feb 19 17:41:10 2013 +0100 @@ -72,7 +72,7 @@ private final InvokeKind invokeKind; - AMD64DirectCallOp(Object target, Value result, Value[] parameters, Value[] temps, LIRFrameState state, InvokeKind invokeKind, LIR lir) { + AMD64DirectCallOp(InvokeTarget target, Value result, Value[] parameters, Value[] temps, LIRFrameState state, InvokeKind invokeKind, LIR lir) { super(target, result, parameters, temps, state); this.invokeKind = invokeKind;
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotBackend.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotBackend.java Tue Feb 19 17:41:10 2013 +0100 @@ -31,6 +31,7 @@ import com.oracle.graal.amd64.*; import com.oracle.graal.api.code.*; +import com.oracle.graal.api.code.RuntimeCallTarget.*; import com.oracle.graal.api.meta.*; import com.oracle.graal.asm.*; import com.oracle.graal.asm.amd64.AMD64Assembler.ConditionFlag; @@ -55,6 +56,10 @@ */ public class AMD64HotSpotBackend extends HotSpotBackend { + public static final Descriptor EXCEPTION_HANDLER = new Descriptor("exceptionHandler", true, void.class); + public static final Descriptor DEOPT_HANDLER = new Descriptor("deoptHandler", true, void.class); + public static final Descriptor IC_MISS_HANDLER = new Descriptor("icMissHandler", true, void.class); + public AMD64HotSpotBackend(HotSpotRuntime runtime, TargetDescription target) { super(runtime, target); } @@ -162,6 +167,32 @@ } } + /** + * Emits code to do stack overflow checking. + * + * @param afterFrameInit specifies if the stack pointer has already been adjusted to allocate + * the current frame + */ + protected static void emitStackOverflowCheck(TargetMethodAssembler tasm, boolean afterFrameInit) { + if (GraalOptions.StackShadowPages > 0) { + + AMD64MacroAssembler asm = (AMD64MacroAssembler) tasm.asm; + int frameSize = tasm.frameMap.frameSize(); + if (frameSize > 0) { + int lastFramePage = frameSize / tasm.target.pageSize; + // emit multiple stack bangs for methods with frames larger than a page + for (int i = 0; i <= lastFramePage; i++) { + int disp = (i + GraalOptions.StackShadowPages) * tasm.target.pageSize; + if (afterFrameInit) { + disp -= frameSize; + } + tasm.blockComment("[stack overflow check]"); + asm.movq(new Address(asm.target.wordKind, AMD64.RSP, -disp), AMD64.rax); + } + } + } + } + class HotSpotFrameContext implements FrameContext { @Override @@ -196,7 +227,7 @@ RegisterConfig regConfig = tasm.frameMap.registerConfig; if (csl != null && csl.size != 0) { - tasm.targetMethod.setRegisterRestoreEpilogueOffset(asm.codeBuffer.position()); + tasm.compilationResult.setRegisterRestoreEpilogueOffset(asm.codeBuffer.position()); // saved all registers, restore all registers int frameToCSA = tasm.frameMap.offsetToCalleeSaveArea(); asm.restore(csl, frameToCSA); @@ -242,7 +273,7 @@ HotSpotFrameContext frameContext = omitFrame ? null : new HotSpotFrameContext(); TargetMethodAssembler tasm = new TargetMethodAssembler(target, runtime(), frameMap, masm, frameContext, lir.stubs); tasm.setFrameSize(frameMap.frameSize()); - tasm.targetMethod.setCustomStackAreaOffset(frameMap.offsetToCustomArea()); + tasm.compilationResult.setCustomStackAreaOffset(frameMap.offsetToCustomArea()); return tasm; } @@ -279,11 +310,11 @@ boolean frameOmitted = tasm.frameContext == null; if (!frameOmitted) { tasm.recordMark(Marks.MARK_EXCEPTION_HANDLER_ENTRY); - AMD64Call.directCall(tasm, asm, config.handleExceptionStub, null); + AMD64Call.directCall(tasm, asm, runtime().lookupRuntimeCall(EXCEPTION_HANDLER), null); AMD64Call.shouldNotReachHere(tasm, asm); tasm.recordMark(Marks.MARK_DEOPT_HANDLER_ENTRY); - AMD64Call.directCall(tasm, asm, config.handleDeoptStub, null); + AMD64Call.directCall(tasm, asm, runtime().lookupRuntimeCall(DEOPT_HANDLER), null); AMD64Call.shouldNotReachHere(tasm, asm); } else { // No need to emit the stubs for entries back into the method since @@ -293,7 +324,7 @@ if (unverifiedStub != null) { asm.bind(unverifiedStub); - AMD64Call.directJmp(tasm, asm, config.inlineCacheMissStub); + AMD64Call.directJmp(tasm, asm, runtime().lookupRuntimeCall(IC_MISS_HANDLER)); } for (int i = 0; i < GraalOptions.MethodEndBreakpointGuards; ++i) { @@ -301,4 +332,5 @@ } } + }
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotRuntime.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotRuntime.java Tue Feb 19 17:41:10 2013 +0100 @@ -177,6 +177,17 @@ /* arg3: r */ word, /* arg4: inLength */ Kind.Int)); + addRuntimeCall(AMD64HotSpotBackend.EXCEPTION_HANDLER, config.handleExceptionStub, + /* temps */ null, + /* ret */ ret(Kind.Void)); + + addRuntimeCall(AMD64HotSpotBackend.DEOPT_HANDLER, config.handleDeoptStub, + /* temps */ null, + /* ret */ ret(Kind.Void)); + + addRuntimeCall(AMD64HotSpotBackend.IC_MISS_HANDLER, config.inlineCacheMissStub, + /* temps */ null, + /* ret */ ret(Kind.Void)); // @formatter:on }
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64IndirectCallOp.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64IndirectCallOp.java Tue Feb 19 17:41:10 2013 +0100 @@ -52,7 +52,7 @@ @Use({REG}) protected Value metaspaceMethod; - AMD64IndirectCallOp(Object targetMethod, Value result, Value[] parameters, Value[] temps, Value metaspaceMethod, Value targetAddress, LIRFrameState state) { + AMD64IndirectCallOp(InvokeTarget targetMethod, Value result, Value[] parameters, Value[] temps, Value metaspaceMethod, Value targetAddress, LIRFrameState state) { super(targetMethod, result, parameters, temps, targetAddress, state); this.metaspaceMethod = metaspaceMethod; }
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/HotSpotGraalRuntime.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/HotSpotGraalRuntime.java Tue Feb 19 17:41:10 2013 +0100 @@ -211,8 +211,6 @@ break; case Int: return impl.typeInt; - case Jsr: - break; case Long: return impl.typeLong; case Object:
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/HotSpotRuntimeCallTarget.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/HotSpotRuntimeCallTarget.java Tue Feb 19 17:41:10 2013 +0100 @@ -30,7 +30,7 @@ /** * The details required to link a HotSpot runtime or stub call. */ -public class HotSpotRuntimeCallTarget implements RuntimeCallTarget { +public class HotSpotRuntimeCallTarget implements RuntimeCallTarget, InvokeTarget { /** * The descriptor of the stub. This is for informational purposes only.
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedObjectType.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedObjectType.java Tue Feb 19 17:41:10 2013 +0100 @@ -1,5 +1,5 @@ /* - * Copyright (c) 2011, Oracle and/or its affiliates. All rights reserved. + * Copyright (c) 2011, 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 @@ -29,6 +29,7 @@ import java.lang.annotation.*; import java.lang.reflect.*; +import java.net.*; import java.util.*; import com.oracle.graal.api.meta.*; @@ -497,4 +498,16 @@ } return null; } + + @Override + public String getClassFilePath() { + Class<?> cls = mirror(); + String name = cls.getName(); + int dot = name.lastIndexOf('.'); + if (dot != -1) { + name = name.substring(dot + 1); + } + URL classFilePath = cls.getResource(name + ".class"); + return classFilePath == null ? null : classFilePath.getPath(); + } }
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedPrimitiveType.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedPrimitiveType.java Tue Feb 19 17:41:10 2013 +0100 @@ -1,5 +1,5 @@ /* - * Copyright (c) 2011, Oracle and/or its affiliates. All rights reserved. + * Copyright (c) 2011, 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 @@ -189,4 +189,9 @@ public Class<?> mirror() { return javaMirror; } + + @Override + public String getClassFilePath() { + return null; + } }
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotRuntime.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotRuntime.java Tue Feb 19 17:41:10 2013 +0100 @@ -776,13 +776,6 @@ return HotSpotResolvedObjectType.fromClass(clazz); } - public Object lookupCallTarget(Object callTarget) { - if (callTarget instanceof HotSpotRuntimeCallTarget) { - return ((HotSpotRuntimeCallTarget) callTarget).getAddress(); - } - return callTarget; - } - /** * Gets the stub corresponding to a given method. *
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/HotSpotDirectCallTargetNode.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/HotSpotDirectCallTargetNode.java Tue Feb 19 17:41:10 2013 +0100 @@ -34,7 +34,7 @@ private final InvokeKind invokeKind; - public HotSpotDirectCallTargetNode(List<ValueNode> arguments, Stamp returnStamp, JavaType[] signature, Object target, Type callType, InvokeKind invokeKind) { + public HotSpotDirectCallTargetNode(List<ValueNode> arguments, Stamp returnStamp, JavaType[] signature, InvokeTarget target, Type callType, InvokeKind invokeKind) { super(arguments, returnStamp, signature, target, callType); this.invokeKind = invokeKind; }
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/HotSpotIndirectCallTargetNode.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/HotSpotIndirectCallTargetNode.java Tue Feb 19 17:41:10 2013 +0100 @@ -33,7 +33,7 @@ @Input private ValueNode metaspaceMethod; - public HotSpotIndirectCallTargetNode(ValueNode metaspaceMethod, ValueNode computedAddress, List<ValueNode> arguments, Stamp returnStamp, JavaType[] signature, Object target, Type callType) { + public HotSpotIndirectCallTargetNode(ValueNode metaspaceMethod, ValueNode computedAddress, List<ValueNode> arguments, Stamp returnStamp, JavaType[] signature, InvokeTarget target, Type callType) { super(computedAddress, arguments, returnStamp, signature, target, callType); this.metaspaceMethod = metaspaceMethod; }
--- a/graal/com.oracle.graal.java/src/com/oracle/graal/java/FrameStateBuilder.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.java/src/com/oracle/graal/java/FrameStateBuilder.java Tue Feb 19 17:41:10 2013 +0100 @@ -479,15 +479,6 @@ } /** - * Pushes a value onto the stack and checks that it is a JSR return address. - * - * @param x the instruction to push onto the stack - */ - public void jpush(ValueNode x) { - xpush(assertJsr(x)); - } - - /** * Pushes a value onto the stack and checks that it is a long. * * @param x the instruction to push onto the stack @@ -566,15 +557,6 @@ } /** - * Pops a value off of the stack and checks that it is a JSR return address. - * - * @return x the instruction popped off the stack - */ - public ValueNode jpop() { - return assertJsr(xpop()); - } - - /** * Pops a value off of the stack and checks that it is a long. * * @return x the instruction popped off the stack
--- a/graal/com.oracle.graal.java/src/com/oracle/graal/java/GraphBuilderPhase.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.java/src/com/oracle/graal/java/GraphBuilderPhase.java Tue Feb 19 17:41:10 2013 +0100 @@ -1103,7 +1103,7 @@ if (successor.jsrScope.nextReturnAddress() != stream().nextBCI()) { throw new JsrNotSupportedBailout("unstructured control flow (internal limitation)"); } - frameState.push(Kind.Jsr, ConstantNode.forJsr(stream().nextBCI(), currentGraph)); + frameState.push(Kind.Int, ConstantNode.forInt(stream().nextBCI(), currentGraph)); appendGoto(createTarget(successor, frameState)); } @@ -1112,7 +1112,7 @@ ValueNode local = frameState.loadLocal(localIndex); JsrScope scope = currentBlock.jsrScope; int retAddress = scope.nextReturnAddress(); - append(currentGraph.add(new FixedGuardNode(currentGraph.unique(new IntegerEqualsNode(local, ConstantNode.forJsr(retAddress, currentGraph))), DeoptimizationReason.JavaSubroutineMismatch, + append(currentGraph.add(new FixedGuardNode(currentGraph.unique(new IntegerEqualsNode(local, ConstantNode.forInt(retAddress, currentGraph))), DeoptimizationReason.JavaSubroutineMismatch, DeoptimizationAction.InvalidateReprofile))); if (!successor.jsrScope.equals(scope.pop())) { throw new JsrNotSupportedBailout("unstructured control flow (ret leaves more than one scope)");
--- a/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Call.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Call.java Tue Feb 19 17:41:10 2013 +0100 @@ -45,9 +45,9 @@ @Temp protected Value[] temps; @State protected LIRFrameState state; - protected final Object callTarget; + protected final InvokeTarget callTarget; - public DirectCallOp(Object callTarget, Value result, Value[] parameters, Value[] temps, LIRFrameState state) { + public DirectCallOp(InvokeTarget callTarget, Value result, Value[] parameters, Value[] temps, LIRFrameState state) { this.callTarget = callTarget; this.result = result; this.parameters = parameters; @@ -79,9 +79,9 @@ @Temp protected Value[] temps; @State protected LIRFrameState state; - protected final Object callTarget; + protected final InvokeTarget callTarget; - public IndirectCallOp(Object callTarget, Value result, Value[] parameters, Value[] temps, Value targetAddress, LIRFrameState state) { + public IndirectCallOp(InvokeTarget callTarget, Value result, Value[] parameters, Value[] temps, Value targetAddress, LIRFrameState state) { this.callTarget = callTarget; this.result = result; this.parameters = parameters; @@ -103,7 +103,7 @@ } } - public static void directCall(TargetMethodAssembler tasm, AMD64MacroAssembler masm, Object callTarget, LIRFrameState info) { + public static void directCall(TargetMethodAssembler tasm, AMD64MacroAssembler masm, InvokeTarget callTarget, LIRFrameState info) { int before = masm.codeBuffer.position(); if (callTarget instanceof RuntimeCallTarget) { long maxOffset = ((RuntimeCallTarget) callTarget).getMaxCallTargetOffset(); @@ -121,24 +121,24 @@ masm.call(); } int after = masm.codeBuffer.position(); - tasm.recordDirectCall(before, after, tasm.runtime.lookupCallTarget(callTarget), info); + tasm.recordDirectCall(before, after, callTarget, info); tasm.recordExceptionHandlers(after, info); masm.ensureUniquePC(); } - public static void directJmp(TargetMethodAssembler tasm, AMD64MacroAssembler masm, Object target) { + public static void directJmp(TargetMethodAssembler tasm, AMD64MacroAssembler masm, InvokeTarget target) { int before = masm.codeBuffer.position(); masm.jmp(0, true); int after = masm.codeBuffer.position(); - tasm.recordDirectCall(before, after, tasm.runtime.lookupCallTarget(target), null); + tasm.recordDirectCall(before, after, target, null); masm.ensureUniquePC(); } - public static void indirectCall(TargetMethodAssembler tasm, AMD64MacroAssembler masm, Register dst, Object callTarget, LIRFrameState info) { + public static void indirectCall(TargetMethodAssembler tasm, AMD64MacroAssembler masm, Register dst, InvokeTarget callTarget, LIRFrameState info) { int before = masm.codeBuffer.position(); masm.call(dst); int after = masm.codeBuffer.position(); - tasm.recordIndirectCall(before, after, tasm.runtime.lookupCallTarget(callTarget), info); + tasm.recordIndirectCall(before, after, callTarget, info); tasm.recordExceptionHandlers(after, info); masm.ensureUniquePC(); }
--- a/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Compare.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Compare.java Tue Feb 19 17:41:10 2013 +0100 @@ -54,7 +54,6 @@ protected void verify() { super.verify(); assert (name().startsWith("I") && x.getKind() == Kind.Int && y.getKind().getStackKind() == Kind.Int) - || (name().startsWith("I") && x.getKind() == Kind.Jsr && y.getKind() == Kind.Jsr) || (name().startsWith("L") && x.getKind() == Kind.Long && y.getKind() == Kind.Long) || (name().startsWith("A") && x.getKind() == Kind.Object && y.getKind() == Kind.Object) || (name().startsWith("F") && x.getKind() == Kind.Float && y.getKind() == Kind.Float)
--- a/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64ControlFlow.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64ControlFlow.java Tue Feb 19 17:41:10 2013 +0100 @@ -374,7 +374,7 @@ } JumpTable jt = new JumpTable(jumpTablePos, lowKey, highKey, 4); - tasm.targetMethod.addAnnotation(jt); + tasm.compilationResult.addAnnotation(jt); } private static void floatJcc(AMD64MacroAssembler masm, ConditionFlag condition, boolean unorderedIsTrue, Label label) {
--- a/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Move.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.lir.amd64/src/com/oracle/graal/lir/amd64/AMD64Move.java Tue Feb 19 17:41:10 2013 +0100 @@ -256,7 +256,6 @@ return; } switch (input.getKind()) { - case Jsr: case Int: masm.movl(asRegister(result), asRegister(input)); break; case Long: masm.movq(asRegister(result), asRegister(input)); break; case Float: masm.movflt(asFloatReg(result), asFloatReg(input)); break; @@ -268,7 +267,6 @@ private static void reg2stack(TargetMethodAssembler tasm, AMD64MacroAssembler masm, Value result, Value input) { switch (input.getKind()) { - case Jsr: case Int: masm.movl(tasm.asAddress(result), asRegister(input)); break; case Long: masm.movq(tasm.asAddress(result), asRegister(input)); break; case Float: masm.movflt(tasm.asAddress(result), asFloatReg(input)); break; @@ -280,7 +278,6 @@ private static void stack2reg(TargetMethodAssembler tasm, AMD64MacroAssembler masm, Value result, Value input) { switch (input.getKind()) { - case Jsr: case Int: masm.movl(asRegister(result), tasm.asAddress(input)); break; case Long: masm.movq(asRegister(result), tasm.asAddress(input)); break; case Float: masm.movflt(asFloatReg(result), tasm.asAddress(input)); break; @@ -295,7 +292,6 @@ // in all cases. For example, an object constant can be loaded to a long register when unsafe casts occurred (e.g., // for a write barrier where arithmetic operations are then performed on the pointer). switch (input.getKind().getStackKind()) { - case Jsr: case Int: if (tasm.runtime.needsDataPatch(input)) { tasm.recordDataReferenceInCode(input, 0, true); @@ -354,7 +350,6 @@ private static void const2stack(TargetMethodAssembler tasm, AMD64MacroAssembler masm, Value result, Constant input) { assert !tasm.runtime.needsDataPatch(input); switch (input.getKind().getStackKind()) { - case Jsr: case Int: masm.movl(tasm.asAddress(result), input.asInt()); break; case Long: masm.movlong(tasm.asAddress(result), input.asLong()); break; case Float: masm.movl(tasm.asAddress(result), floatToRawIntBits(input.asFloat())); break; @@ -415,7 +410,6 @@ case Byte: masm.movb(storeAddr, c.asInt() & 0xFF); break; case Char: case Short: masm.movw(storeAddr, c.asInt() & 0xFFFF); break; - case Jsr: case Int: masm.movl(storeAddr, c.asInt()); break; case Long: if (NumUtil.isInt(c.asLong())) {
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXArithmetic.java Tue Feb 19 17:41:10 2013 +0100 @@ -0,0 +1,297 @@ +/* + * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ +package com.oracle.graal.lir.ptx; + +import static com.oracle.graal.api.code.ValueUtil.*; +import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; + +import com.oracle.graal.api.code.*; +import com.oracle.graal.api.meta.*; +import com.oracle.graal.asm.ptx.*; +import com.oracle.graal.graph.*; +import com.oracle.graal.lir.*; +import com.oracle.graal.lir.asm.*; + +// @formatter:off +public enum PTXArithmetic { + IADD, ISUB, IMUL, IDIV, IDIVREM, IREM, IUDIV, IUREM, IAND, IOR, IXOR, ISHL, ISHR, IUSHR, + LADD, LSUB, LMUL, LDIV, LDIVREM, LREM, LUDIV, LUREM, LAND, LOR, LXOR, LSHL, LSHR, LUSHR, + FADD, FSUB, FMUL, FDIV, FAND, FOR, FXOR, + DADD, DSUB, DMUL, DDIV, DAND, DOR, DXOR, + INEG, LNEG, + I2L, L2I, I2B, I2C, I2S, + F2D, D2F, + I2F, I2D, F2I, D2I, + L2F, L2D, F2L, D2L, + MOV_I2F, MOV_L2D, MOV_F2I, MOV_D2L; + + + public static class Op1Reg extends PTXLIRInstruction { + @Opcode private final PTXArithmetic opcode; + @Def({REG, HINT}) protected Value result; + @Use({REG}) protected Value x; + + public Op1Reg(PTXArithmetic opcode, Value result, Value x) { + this.opcode = opcode; + this.result = result; + this.x = x; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + emit(tasm, masm, opcode, result, x, null); + } + } + + public static class Op1Stack extends PTXLIRInstruction { + @Opcode private final PTXArithmetic opcode; + @Def({REG, HINT}) protected Value result; + @Use({REG, STACK, CONST}) protected Value x; + + public Op1Stack(PTXArithmetic opcode, Value result, Value x) { + this.opcode = opcode; + this.result = result; + this.x = x; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + emit(tasm, masm, opcode, result, x, null); + } + } + + public static class Op2Stack extends PTXLIRInstruction { + @Opcode private final PTXArithmetic opcode; + @Def({REG, HINT}) protected Value result; + @Use({REG, STACK, CONST}) protected Value x; + @Alive({REG, STACK, CONST}) protected Value y; + + public Op2Stack(PTXArithmetic opcode, Value result, Value x, Value y) { + this.opcode = opcode; + this.result = result; + this.x = x; + this.y = y; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + emit(tasm, masm, opcode, result, x, y, null); + } + + @Override + public void verify() { + super.verify(); + verifyKind(opcode, result, x, y); + } + } + + public static class Op2Reg extends PTXLIRInstruction { + @Opcode private final PTXArithmetic opcode; + @Def({REG, HINT}) protected Value result; + @Use({REG, STACK, CONST}) protected Value x; + @Alive({REG, CONST}) protected Value y; + + public Op2Reg(PTXArithmetic opcode, Value result, Value x, Value y) { + this.opcode = opcode; + this.result = result; + this.x = x; + this.y = y; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + emit(tasm, masm, opcode, result, x, y, null); + } + + @Override + public void verify() { + super.verify(); + verifyKind(opcode, result, x, y); + } + } + + public static class Op2RegCommutative extends PTXLIRInstruction { + @Opcode private final PTXArithmetic opcode; + @Def({REG, HINT}) protected Value result; + @Use({REG, STACK, CONST}) protected Value x; + @Use({REG, CONST}) protected Value y; + + public Op2RegCommutative(PTXArithmetic opcode, Value result, Value x, Value y) { + this.opcode = opcode; + this.result = result; + this.x = x; + this.y = y; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + if (sameRegister(result, y)) { + emit(tasm, masm, opcode, result, x, null); + } else { + PTXMove.move(tasm, masm, result, x); + emit(tasm, masm, opcode, result, y, null); + } + } + + @Override + protected void verify() { + super.verify(); + verifyKind(opcode, result, x, y); + } + } + + public static class ShiftOp extends PTXLIRInstruction { + @Opcode private final PTXArithmetic opcode; + @Def({REG, HINT}) protected Value result; + @Use({REG, STACK, CONST}) protected Value x; + @Alive({REG, CONST}) protected Value y; + + public ShiftOp(PTXArithmetic opcode, Value result, Value x, Value y) { + this.opcode = opcode; + this.result = result; + this.x = x; + this.y = y; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + emit(tasm, masm, opcode, result, x, y, null); + } + + @Override + public void verify() { + super.verify(); + verifyKind(opcode, result, x, x); + assert y.getKind().getStackKind() == Kind.Int; + } + } + + public static class DivOp extends PTXLIRInstruction { + @Opcode private final PTXArithmetic opcode; + @Def protected Value result; + @Use protected Value x; + @Alive protected Value y; + @State protected LIRFrameState state; + + public DivOp(PTXArithmetic opcode, Value result, Value x, Value y, LIRFrameState state) { + this.opcode = opcode; + this.result = result; + this.x = x; + this.y = y; + this.state = state; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + emit(tasm, masm, opcode, result, y, state); + } + + @Override + protected void verify() { + super.verify(); + verifyKind(opcode, result, x, y); + } + } + + + @SuppressWarnings("unused") + protected static void emit(TargetMethodAssembler tasm, PTXAssembler masm, PTXArithmetic opcode, Value result) { + switch (opcode) { + default: throw GraalInternalError.shouldNotReachHere(); + } + } + + public static void emit(TargetMethodAssembler tasm, PTXAssembler masm, PTXArithmetic opcode, Value dst, Value src, LIRFrameState info) { + int exceptionOffset = -1; + if (isRegister(src)) { + Register a = asIntReg(src); + Register d = asIntReg(dst); + switch (opcode) { + case INEG: masm.neg_s32(d, a); break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } else if (isConstant(src)) { + switch (opcode) { + case ISUB: masm.sub_s32(asIntReg(dst), asIntReg(dst), tasm.asIntConst(src)); break; + case IAND: masm.and_b32(asIntReg(dst), asIntReg(dst), tasm.asIntConst(src)); break; + default: throw GraalInternalError.shouldNotReachHere(); + } + } else { + switch (opcode) { + default: throw GraalInternalError.shouldNotReachHere(); + } + } + + if (info != null) { + assert exceptionOffset != -1; + tasm.recordImplicitException(exceptionOffset, info); + } + } + + public static void emit(TargetMethodAssembler tasm, PTXAssembler masm, PTXArithmetic opcode, Value dst, Value src1, Value src2, LIRFrameState info) { + int exceptionOffset = -1; + if (isConstant(src1)) { + int a = tasm.asIntConst(src1); + Register b = asIntReg(src2); + Register d = asIntReg(dst); + switch (opcode) { + case ISUB: masm.sub_s32(d, a, b); break; + case IAND: masm.and_b32(d, b, a); break; + default: throw GraalInternalError.shouldNotReachHere(); + } + } else if (isConstant(src2)) { + Register a = asIntReg(src1); + int b = tasm.asIntConst(src2); + Register d = asIntReg(dst); + switch (opcode) { + case IADD: masm.add_s32(d, a, b); break; + case IAND: masm.and_b32(d, a, b); break; + case IUSHR: masm.shr_u32(d, a, b); break; + default: throw GraalInternalError.shouldNotReachHere(); + } + } else { + Register a = asIntReg(src1); + Register b = asIntReg(src2); + Register d = asIntReg(dst); + switch (opcode) { + case IADD: masm.add_s32(d, a, b); break; + case ISUB: masm.sub_s32(d, a, b); break; + case IMUL: masm.mul_s32(d, a, b); break; + default: throw GraalInternalError.shouldNotReachHere(); + } + } + + if (info != null) { + assert exceptionOffset != -1; + tasm.recordImplicitException(exceptionOffset, info); + } + } + + private static void verifyKind(PTXArithmetic opcode, Value result, Value x, Value y) { + assert (opcode.name().startsWith("I") && result.getKind() == Kind.Int && x.getKind().getStackKind() == Kind.Int && y.getKind().getStackKind() == Kind.Int) + || (opcode.name().startsWith("L") && result.getKind() == Kind.Long && x.getKind() == Kind.Long && y.getKind() == Kind.Long) + || (opcode.name().startsWith("F") && result.getKind() == Kind.Float && x.getKind() == Kind.Float && y.getKind() == Kind.Float) + || (opcode.name().startsWith("D") && result.getKind() == Kind.Double && x.getKind() == Kind.Double && y.getKind() == Kind.Double); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXBitManipulationOp.java Tue Feb 19 17:41:10 2013 +0100 @@ -0,0 +1,62 @@ +/* + * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ +package com.oracle.graal.lir.ptx; + +import com.oracle.graal.api.code.*; +import com.oracle.graal.api.meta.*; +import com.oracle.graal.asm.ptx.*; +import com.oracle.graal.graph.*; +import com.oracle.graal.lir.asm.*; + +public class PTXBitManipulationOp extends PTXLIRInstruction { + + public enum IntrinsicOpcode { + IPOPCNT, LPOPCNT, IBSR, LBSR, BSF; + } + + @Opcode private final IntrinsicOpcode opcode; + @Def protected Value result; + @Use({OperandFlag.REG, OperandFlag.ADDR}) protected Value input; + + public PTXBitManipulationOp(IntrinsicOpcode opcode, Value result, Value input) { + this.opcode = opcode; + this.result = result; + this.input = input; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + Register dst = ValueUtil.asIntReg(result); + Register src = ValueUtil.asRegister(input); + switch (opcode) { + case IPOPCNT: + masm.popc_b32(dst, src); + break; + case LPOPCNT: + masm.popc_b64(dst, src); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXCode.java Tue Feb 19 17:41:10 2013 +0100 @@ -0,0 +1,40 @@ +/* + * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ +package com.oracle.graal.lir.ptx; + +import com.oracle.graal.asm.ptx.*; +import com.oracle.graal.lir.*; +import com.oracle.graal.lir.asm.*; + +/** + * Convenience class to provide PTXAssembler for the {@link #emitCode} method. + */ +public abstract class PTXCode implements LIR.Code { + + @Override + public final void emitCode(TargetMethodAssembler tasm) { + emitCode(tasm, (PTXAssembler) tasm.asm); + } + + public abstract void emitCode(TargetMethodAssembler tasm, PTXAssembler masm); +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXCompare.java Tue Feb 19 17:41:10 2013 +0100 @@ -0,0 +1,226 @@ +/* + * Copyright (c) 2011, 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.lir.ptx; + +import static com.oracle.graal.api.code.ValueUtil.*; +import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; + +import com.oracle.graal.api.code.*; +import com.oracle.graal.api.meta.*; +import com.oracle.graal.asm.ptx.*; +import com.oracle.graal.graph.*; +import com.oracle.graal.lir.asm.*; +import com.oracle.graal.nodes.calc.*; + +public enum PTXCompare { + ICMP, LCMP, ACMP, FCMP, DCMP; + + public static class CompareOp extends PTXLIRInstruction { + + @Opcode private final PTXCompare opcode; + @Use({REG}) protected Value x; + @Use({REG, STACK, CONST}) protected Value y; + private final Condition condition; + + public CompareOp(PTXCompare opcode, Condition condition, Value x, Value y) { + this.opcode = opcode; + this.condition = condition; + this.x = x; + this.y = y; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + emit(tasm, masm, opcode, condition, x, y); + } + + @Override + protected void verify() { + super.verify(); + assert (name().startsWith("I") && x.getKind() == Kind.Int && y.getKind().getStackKind() == Kind.Int) || (name().startsWith("L") && x.getKind() == Kind.Long && y.getKind() == Kind.Long) || + (name().startsWith("A") && x.getKind() == Kind.Object && y.getKind() == Kind.Object) || + (name().startsWith("F") && x.getKind() == Kind.Float && y.getKind() == Kind.Float) || (name().startsWith("D") && x.getKind() == Kind.Double && y.getKind() == Kind.Double); + } + } + + public static void emit(TargetMethodAssembler tasm, PTXAssembler masm, PTXCompare opcode, Condition condition, Value x, Value y) { + if (isConstant(x)) { + int a = tasm.asIntConst(x); + Register b = asIntReg(y); + switch (opcode) { + case ICMP: + emitCompareConstReg(masm, condition, a, b); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } else if (isConstant(y)) { + Register a = asIntReg(x); + int b = tasm.asIntConst(y); + switch (opcode) { + case ICMP: + emitCompareRegConst(masm, condition, a, b); + break; + case ACMP: + if (((Constant) y).isNull()) { + switch (condition) { + case EQ: + masm.setp_eq_s32(a, b); + break; + case NE: + masm.setp_ne_s32(a, b); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } else { + throw GraalInternalError.shouldNotReachHere("Only null object constants are allowed in comparisons"); + } + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } else { + Register a = asIntReg(x); + Register b = asIntReg(y); + switch (opcode) { + case ICMP: + emitCompareRegReg(masm, condition, a, b); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } + } + + private static void emitCompareConstReg(PTXAssembler masm, Condition condition, int a, Register b) { + switch (condition) { + case EQ: + masm.setp_eq_s32(a, b); + break; + case NE: + masm.setp_ne_s32(a, b); + break; + case LT: + masm.setp_lt_s32(a, b); + break; + case LE: + masm.setp_le_s32(a, b); + break; + case GT: + masm.setp_gt_s32(a, b); + break; + case GE: + masm.setp_ge_s32(a, b); + break; + case AT: + masm.setp_gt_u32(a, b); + break; + case AE: + masm.setp_ge_u32(a, b); + break; + case BT: + masm.setp_lt_u32(a, b); + break; + case BE: + masm.setp_le_u32(a, b); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } + + private static void emitCompareRegConst(PTXAssembler masm, Condition condition, Register a, int b) { + switch (condition) { + case EQ: + masm.setp_eq_s32(a, b); + break; + case NE: + masm.setp_ne_s32(a, b); + break; + case LT: + masm.setp_lt_s32(a, b); + break; + case LE: + masm.setp_le_s32(a, b); + break; + case GT: + masm.setp_gt_s32(a, b); + break; + case GE: + masm.setp_ge_s32(a, b); + break; + case AT: + masm.setp_gt_u32(a, b); + break; + case AE: + masm.setp_ge_u32(a, b); + break; + case BT: + masm.setp_lt_u32(a, b); + break; + case BE: + masm.setp_le_u32(a, b); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } + + private static void emitCompareRegReg(PTXAssembler masm, Condition condition, Register a, Register b) { + switch (condition) { + case EQ: + masm.setp_eq_s32(a, b); + break; + case NE: + masm.setp_ne_s32(a, b); + break; + case LT: + masm.setp_lt_s32(a, b); + break; + case LE: + masm.setp_le_s32(a, b); + break; + case GT: + masm.setp_gt_s32(a, b); + break; + case GE: + masm.setp_ge_s32(a, b); + break; + case AT: + masm.setp_gt_u32(a, b); + break; + case AE: + masm.setp_ge_u32(a, b); + break; + case BT: + masm.setp_lt_u32(a, b); + break; + case BE: + masm.setp_le_u32(a, b); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXControlFlow.java Tue Feb 19 17:41:10 2013 +0100 @@ -0,0 +1,85 @@ +/* + * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ +package com.oracle.graal.lir.ptx; + +import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; + +import com.oracle.graal.api.meta.*; +import com.oracle.graal.asm.*; +import com.oracle.graal.asm.ptx.*; +import com.oracle.graal.lir.*; +import com.oracle.graal.lir.asm.*; +import com.oracle.graal.nodes.calc.*; + +public class PTXControlFlow { + + public static class ReturnOp extends PTXLIRInstruction { + + @Use({REG, ILLEGAL}) protected Value x; + + public ReturnOp(Value x) { + this.x = x; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + if (tasm.frameContext != null) { + tasm.frameContext.leave(tasm); + } + masm.exit(); + } + } + + public static class BranchOp extends PTXLIRInstruction implements StandardOp.BranchOp { + + protected Condition condition; + protected LabelRef destination; + @State protected LIRFrameState state; + + public BranchOp(Condition condition, LabelRef destination, LIRFrameState state) { + this.condition = condition; + this.destination = destination; + this.state = state; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + masm.at(); + Label l = destination.label(); + l.addPatchAt(tasm.asm.codeBuffer.position()); + String target = l.isBound() ? "L" + l.toString() : AbstractPTXAssembler.UNBOUND_TARGET; + masm.bra(target); + } + + @Override + public LabelRef destination() { + return destination; + } + + @Override + public void negate(LabelRef newDestination) { + destination = newDestination; + condition = condition.negate(); + } + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXLIRInstruction.java Tue Feb 19 17:41:10 2013 +0100 @@ -0,0 +1,40 @@ +/* + * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ +package com.oracle.graal.lir.ptx; + +import com.oracle.graal.asm.ptx.*; +import com.oracle.graal.lir.*; +import com.oracle.graal.lir.asm.*; + +/** + * Convenience class to provide PTXAssembler for the {@link #emitCode} method. + */ +public abstract class PTXLIRInstruction extends LIRInstruction { + + @Override + public final void emitCode(TargetMethodAssembler tasm) { + emitCode(tasm, (PTXAssembler) tasm.asm); + } + + public abstract void emitCode(TargetMethodAssembler tasm, PTXAssembler masm); +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.lir.ptx/src/com/oracle/graal/lir/ptx/PTXMove.java Tue Feb 19 17:41:10 2013 +0100 @@ -0,0 +1,284 @@ +/* + * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ +package com.oracle.graal.lir.ptx; + +import static com.oracle.graal.api.code.ValueUtil.*; +import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; + +import com.oracle.graal.api.code.*; +import com.oracle.graal.api.meta.*; +import com.oracle.graal.asm.ptx.*; +import com.oracle.graal.graph.*; +import com.oracle.graal.lir.*; +import com.oracle.graal.lir.LIRInstruction.Opcode; +import com.oracle.graal.lir.StandardOp.MoveOp; +import com.oracle.graal.lir.asm.*; + +public class PTXMove { + + @Opcode("MOVE") + public static class SpillMoveOp extends PTXLIRInstruction implements MoveOp { + + @Def({REG, STACK}) protected Value result; + @Use({REG, STACK, CONST}) protected Value input; + + public SpillMoveOp(Value result, Value input) { + this.result = result; + this.input = input; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + move(tasm, masm, getResult(), getInput()); + } + + @Override + public Value getInput() { + return input; + } + + @Override + public Value getResult() { + return result; + } + } + + @Opcode("MOVE") + public static class MoveToRegOp extends PTXLIRInstruction implements MoveOp { + + @Def({REG, HINT}) protected Value result; + @Use({REG, STACK, CONST}) protected Value input; + + public MoveToRegOp(Value result, Value input) { + this.result = result; + this.input = input; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + move(tasm, masm, getResult(), getInput()); + } + + @Override + public Value getInput() { + return input; + } + + @Override + public Value getResult() { + return result; + } + } + + @Opcode("MOVE") + public static class MoveFromRegOp extends PTXLIRInstruction implements MoveOp { + + @Def({REG, STACK}) protected Value result; + @Use({REG, CONST, HINT}) protected Value input; + + public MoveFromRegOp(Value result, Value input) { + this.result = result; + this.input = input; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + move(tasm, masm, getResult(), getInput()); + } + + @Override + public Value getInput() { + return input; + } + + @Override + public Value getResult() { + return result; + } + } + + public static class LoadOp extends PTXLIRInstruction { + + @Def({REG}) protected Value result; + @Use({ADDR}) protected Value address; + @State protected LIRFrameState state; + + public LoadOp(Value result, Value address, LIRFrameState state) { + this.result = result; + this.address = address; + this.state = state; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + load(tasm, masm, result, (Address) address, state); + } + } + + public static class StoreOp extends PTXLIRInstruction { + + @Use({ADDR}) protected Value address; + @Use({REG, CONST}) protected Value input; + @State protected LIRFrameState state; + + public StoreOp(Value address, Value input, LIRFrameState state) { + this.address = address; + this.input = input; + this.state = state; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + store(tasm, masm, (Address) address, input, state); + } + } + + public static class LeaOp extends PTXLIRInstruction { + + @Def({REG}) protected Value result; + @Use({ADDR, STACK, UNINITIALIZED}) protected Value address; + + public LeaOp(Value result, Value address) { + this.result = result; + this.address = address; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + throw new InternalError("NYI"); + } + } + + @Opcode("CAS") + public static class CompareAndSwapOp extends PTXLIRInstruction { + + @Def protected Value result; + @Use({ADDR}) protected Value address; + @Use protected Value cmpValue; + @Use protected Value newValue; + + public CompareAndSwapOp(Value result, Address address, Value cmpValue, Value newValue) { + this.result = result; + this.address = address; + this.cmpValue = cmpValue; + this.newValue = newValue; + } + + @Override + public void emitCode(TargetMethodAssembler tasm, PTXAssembler masm) { + compareAndSwap(tasm, masm, result, (Address) address, cmpValue, newValue); + } + } + + public static void move(TargetMethodAssembler tasm, PTXAssembler masm, Value result, Value input) { + if (isRegister(input)) { + if (isRegister(result)) { + reg2reg(masm, result, input); + } else { + throw GraalInternalError.shouldNotReachHere(); + } + } else if (isConstant(input)) { + if (isRegister(result)) { + const2reg(tasm, masm, result, (Constant) input); + } else { + throw GraalInternalError.shouldNotReachHere(); + } + } else { + throw GraalInternalError.shouldNotReachHere(); + } + } + + private static void reg2reg(PTXAssembler masm, Value result, Value input) { + if (asRegister(input).equals(asRegister(result))) { + return; + } + switch (input.getKind()) { + case Int: + masm.mov_s32(asRegister(result), asRegister(input)); + break; + case Object: + masm.mov_u64(asRegister(result), asRegister(input)); + break; + default: + throw GraalInternalError.shouldNotReachHere("kind=" + result.getKind()); + } + } + + private static void const2reg(TargetMethodAssembler tasm, PTXAssembler masm, Value result, Constant input) { + switch (input.getKind().getStackKind()) { + case Int: + if (tasm.runtime.needsDataPatch(input)) { + tasm.recordDataReferenceInCode(input, 0, true); + } + masm.mov_s32(asRegister(result), input.asInt()); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } + + @SuppressWarnings("unused") + public static void load(TargetMethodAssembler tasm, PTXAssembler masm, Value result, Address loadAddr, LIRFrameState info) { + Register a = asRegister(loadAddr.getBase()); + int immOff = loadAddr.getDisplacement(); + switch (loadAddr.getKind()) { + case Int: + masm.ld_global_s32(asRegister(result), a, immOff); + break; + case Object: + masm.ld_global_u32(asRegister(result), a, immOff); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } + + @SuppressWarnings("unused") + public static void store(TargetMethodAssembler tasm, PTXAssembler masm, Address storeAddr, Value input, LIRFrameState info) { + Register a = asRegister(storeAddr.getBase()); + int immOff = storeAddr.getDisplacement(); + if (isRegister(input)) { + switch (storeAddr.getKind()) { + case Int: + masm.st_global_s32(a, immOff, asRegister(input)); + break; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } else if (isConstant(input)) { + Constant c = (Constant) input; + switch (storeAddr.getKind()) { + default: + throw GraalInternalError.shouldNotReachHere(); + } + + } else { + throw GraalInternalError.shouldNotReachHere(); + } + } + + @SuppressWarnings("unused") + protected static void compareAndSwap(TargetMethodAssembler tasm, PTXAssembler masm, Value result, Address address, Value cmpValue, Value newValue) { + throw new InternalError("NYI"); + } +}
--- a/graal/com.oracle.graal.lir/src/com/oracle/graal/lir/asm/TargetMethodAssembler.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.lir/src/com/oracle/graal/lir/asm/TargetMethodAssembler.java Tue Feb 19 17:41:10 2013 +0100 @@ -48,7 +48,7 @@ } public final AbstractAssembler asm; - public final CompilationResult targetMethod; + public final CompilationResult compilationResult; public final TargetDescription target; public final CodeCacheProvider runtime; public final FrameMap frameMap; @@ -73,52 +73,52 @@ this.frameMap = frameMap; this.stubs = stubs; this.asm = asm; - this.targetMethod = new CompilationResult(); + this.compilationResult = new CompilationResult(); this.frameContext = frameContext; // 0 is a valid pc for safepoints in template methods this.lastSafepointPos = -1; } public void setFrameSize(int frameSize) { - targetMethod.setFrameSize(frameSize); + compilationResult.setFrameSize(frameSize); } private static final CompilationResult.Mark[] NO_REFS = {}; public CompilationResult.Mark recordMark(Object id) { - return targetMethod.recordMark(asm.codeBuffer.position(), id, NO_REFS); + return compilationResult.recordMark(asm.codeBuffer.position(), id, NO_REFS); } public CompilationResult.Mark recordMark(Object id, CompilationResult.Mark... references) { - return targetMethod.recordMark(asm.codeBuffer.position(), id, references); + return compilationResult.recordMark(asm.codeBuffer.position(), id, references); } public void blockComment(String s) { - targetMethod.addAnnotation(new CompilationResult.CodeComment(asm.codeBuffer.position(), s)); + compilationResult.addAnnotation(new CompilationResult.CodeComment(asm.codeBuffer.position(), s)); } public CompilationResult finishTargetMethod(Object name, boolean isStub) { // Install code, data and frame size - targetMethod.setTargetCode(asm.codeBuffer.close(false), asm.codeBuffer.position()); + compilationResult.setTargetCode(asm.codeBuffer.close(false), asm.codeBuffer.position()); // Record exception handlers if they exist if (exceptionInfoList != null) { for (ExceptionInfo ei : exceptionInfoList) { int codeOffset = ei.codeOffset; - targetMethod.recordExceptionHandler(codeOffset, ei.exceptionEdge.label().position()); + compilationResult.recordExceptionHandler(codeOffset, ei.exceptionEdge.label().position()); } } // Set the info on callee-saved registers - targetMethod.setCalleeSaveLayout(frameMap.registerConfig.getCalleeSaveLayout()); + compilationResult.setCalleeSaveLayout(frameMap.registerConfig.getCalleeSaveLayout()); Debug.metric("TargetMethods").increment(); - Debug.metric("CodeBytesEmitted").add(targetMethod.getTargetCodeSize()); - Debug.metric("SafepointsEmitted").add(targetMethod.getSafepoints().size()); - Debug.metric("DataPatches").add(targetMethod.getDataReferences().size()); - Debug.metric("ExceptionHandlersEmitted").add(targetMethod.getExceptionHandlers().size()); + Debug.metric("CodeBytesEmitted").add(compilationResult.getTargetCodeSize()); + Debug.metric("SafepointsEmitted").add(compilationResult.getSafepoints().size()); + Debug.metric("DataPatches").add(compilationResult.getDataReferences().size()); + Debug.metric("ExceptionHandlersEmitted").add(compilationResult.getExceptionHandlers().size()); Debug.log("Finished target method %s, isStub %b", name, isStub); - return targetMethod; + return compilationResult; } public void recordExceptionHandlers(int pcOffset, LIRFrameState info) { @@ -137,23 +137,23 @@ if (info != null) { assert lastSafepointPos < pcOffset : lastSafepointPos + "<" + pcOffset; lastSafepointPos = pcOffset; - targetMethod.recordSafepoint(pcOffset, info.debugInfo()); + compilationResult.recordSafepoint(pcOffset, info.debugInfo()); assert info.exceptionEdge == null; } } - public void recordDirectCall(int posBefore, int posAfter, Object callTarget, LIRFrameState info) { + public void recordDirectCall(int posBefore, int posAfter, InvokeTarget callTarget, LIRFrameState info) { DebugInfo debugInfo = info != null ? info.debugInfo() : null; assert lastSafepointPos < posAfter; lastSafepointPos = posAfter; - targetMethod.recordCall(posBefore, posAfter - posBefore, callTarget, debugInfo, true); + compilationResult.recordCall(posBefore, posAfter - posBefore, callTarget, debugInfo, true); } - public void recordIndirectCall(int posBefore, int posAfter, Object callTarget, LIRFrameState info) { + public void recordIndirectCall(int posBefore, int posAfter, InvokeTarget callTarget, LIRFrameState info) { DebugInfo debugInfo = info != null ? info.debugInfo() : null; assert lastSafepointPos < posAfter; lastSafepointPos = posAfter; - targetMethod.recordCall(posBefore, posAfter - posBefore, callTarget, debugInfo, false); + compilationResult.recordCall(posBefore, posAfter - posBefore, callTarget, debugInfo, false); } public void recordSafepoint(int pos, LIRFrameState info) { @@ -161,14 +161,14 @@ DebugInfo debugInfo = info.debugInfo(); assert lastSafepointPos < pos; lastSafepointPos = pos; - targetMethod.recordSafepoint(pos, debugInfo); + compilationResult.recordSafepoint(pos, debugInfo); } public Address recordDataReferenceInCode(Constant data, int alignment, boolean inlined) { assert data != null; int pos = asm.codeBuffer.position(); Debug.log("Data reference in code: pos = %d, data = %s", pos, data.toString()); - targetMethod.recordDataReference(pos, data, alignment, inlined); + compilationResult.recordDataReference(pos, data, alignment, inlined); return Address.Placeholder; } @@ -181,7 +181,7 @@ * including long constants that fit into the 32-bit range. */ public int asIntConst(Value value) { - assert (value.getKind().getStackKind() == Kind.Int || value.getKind() == Kind.Jsr || value.getKind() == Kind.Long) && isConstant(value); + assert (value.getKind().getStackKind() == Kind.Int || value.getKind() == Kind.Long) && isConstant(value); Constant constant = (Constant) value; assert !runtime.needsDataPatch(constant) : constant + " should be in a DataPatch"; long c = constant.asLong();
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/AbstractCallTargetNode.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/AbstractCallTargetNode.java Tue Feb 19 17:41:10 2013 +0100 @@ -32,10 +32,10 @@ private final Stamp returnStamp; private final JavaType[] signature; - private final Object target; + private final InvokeTarget target; private final CallingConvention.Type callType; - public AbstractCallTargetNode(List<ValueNode> arguments, Stamp returnStamp, JavaType[] signature, Object target, CallingConvention.Type callType) { + public AbstractCallTargetNode(List<ValueNode> arguments, Stamp returnStamp, JavaType[] signature, InvokeTarget target, CallingConvention.Type callType) { super(arguments); this.returnStamp = returnStamp; this.signature = signature; @@ -52,7 +52,7 @@ return signature; } - public Object target() { + public InvokeTarget target() { return target; }
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/ConstantNode.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/ConstantNode.java Tue Feb 19 17:41:10 2013 +0100 @@ -165,17 +165,6 @@ } /** - * Returns a node for an address (jsr/ret address) constant. - * - * @param i the address value for which to create the instruction - * @param graph - * @return a node representing the address - */ - public static ConstantNode forJsr(int i, Graph graph) { - return graph.unique(new ConstantNode(Constant.forJsr(i))); - } - - /** * Returns a node for an object constant. * * @param o the object value for which to create the instruction
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/DirectCallTargetNode.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/DirectCallTargetNode.java Tue Feb 19 17:41:10 2013 +0100 @@ -30,7 +30,7 @@ public class DirectCallTargetNode extends AbstractCallTargetNode { - public DirectCallTargetNode(List<ValueNode> arguments, Stamp returnStamp, JavaType[] signature, Object target, CallingConvention.Type callType) { + public DirectCallTargetNode(List<ValueNode> arguments, Stamp returnStamp, JavaType[] signature, InvokeTarget target, CallingConvention.Type callType) { super(arguments, returnStamp, signature, target, callType); }
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/IndirectCallTargetNode.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/IndirectCallTargetNode.java Tue Feb 19 17:41:10 2013 +0100 @@ -32,7 +32,7 @@ @Input protected ValueNode computedAddress; - public IndirectCallTargetNode(ValueNode computedAddress, List<ValueNode> arguments, Stamp returnStamp, JavaType[] signature, Object target, CallingConvention.Type callType) { + public IndirectCallTargetNode(ValueNode computedAddress, List<ValueNode> arguments, Stamp returnStamp, JavaType[] signature, InvokeTarget target, CallingConvention.Type callType) { super(arguments, returnStamp, signature, target, callType); this.computedAddress = computedAddress; }
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/ValueNodeUtil.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/ValueNodeUtil.java Tue Feb 19 17:41:10 2013 +0100 @@ -31,7 +31,7 @@ public class ValueNodeUtil { public static ValueNode assertKind(Kind kind, ValueNode x) { - assert x != null && ((x.kind() == kind) || (x.kind() == Kind.Jsr && kind == Kind.Object)) : "kind=" + kind + ", value=" + x + ((x == null) ? "" : ", value.kind=" + x.kind()); + assert x != null && x.kind() == kind : "kind=" + kind + ", value=" + x + ((x == null) ? "" : ", value.kind=" + x.kind()); return x; } @@ -48,11 +48,6 @@ return x; } - public static ValueNode assertJsr(ValueNode x) { - assert x != null && (x.kind() == Kind.Jsr); - return x; - } - public static ValueNode assertInt(ValueNode x) { assert x != null && (x.kind() == Kind.Int); return x;
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/Condition.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/calc/Condition.java Tue Feb 19 17:41:10 2013 +0100 @@ -342,7 +342,6 @@ case Byte: case Char: case Short: - case Jsr: case Int: { int x = lt.asInt(); int y = rt.asInt();
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/type/IntegerStamp.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/type/IntegerStamp.java Tue Feb 19 17:41:10 2013 +0100 @@ -199,7 +199,6 @@ return 0xffffL; case Short: return 0xffffL; - case Jsr: case Int: return 0xffffffffL; case Long:
--- a/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/type/StampFactory.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/type/StampFactory.java Tue Feb 19 17:41:10 2013 +0100 @@ -57,8 +57,6 @@ setCache(Kind.Float, new FloatStamp(Kind.Float)); setCache(Kind.Double, new FloatStamp(Kind.Double)); - setCache(Kind.Jsr, new IntegerStamp(Kind.Jsr)); - setCache(Kind.Object, objectStamp); setCache(Kind.Void, voidStamp); }
--- a/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/InliningUtil.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/InliningUtil.java Tue Feb 19 17:41:10 2013 +0100 @@ -966,8 +966,6 @@ return logNotInlinedMethodAndReturnFalse(invoke, method, "the method's class is not initialized"); } else if (!method.canBeInlined()) { return logNotInlinedMethodAndReturnFalse(invoke, method, "it is marked non-inlinable"); - } else if (computeInliningLevel(invoke) > GraalOptions.MaximumInlineLevel) { - return logNotInlinedMethodAndReturnFalse(invoke, method, "it exceeds the maximum inlining depth"); } else if (computeRecursiveInliningLevel(invoke.stateAfter(), method) > GraalOptions.MaximumRecursiveInlining) { return logNotInlinedMethodAndReturnFalse(invoke, method, "it exceeds the maximum recursive inlining depth"); } else if (new OptimisticOptimizations(method).lessOptimisticThan(optimisticOpts)) {
--- a/graal/com.oracle.graal.phases/src/com/oracle/graal/phases/GraalOptions.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.phases/src/com/oracle/graal/phases/GraalOptions.java Tue Feb 19 17:41:10 2013 +0100 @@ -43,7 +43,6 @@ static boolean InlineMonomorphicCalls = true; static boolean InlinePolymorphicCalls = true; static boolean InlineMegamorphicCalls = ____; - public static int MaximumInlineLevel = 30; public static int MaximumDesiredSize = 5000; public static int MaximumRecursiveInlining = 1; public static boolean LimitInlinedProbability = ____;
--- a/graal/com.oracle.graal.phases/src/com/oracle/graal/phases/OptimisticOptimizations.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.phases/src/com/oracle/graal/phases/OptimisticOptimizations.java Tue Feb 19 17:41:10 2013 +0100 @@ -56,7 +56,6 @@ * TODO (chaeubl): see GRAAL-75 (remove when we are sure that optimistic optimizations * are not disabled unnecessarily */ - TTY.println("WARN: deactivated optimistic optimization %s for %s", optimization.name(), MetaUtil.format("%H.%n(%p)", method)); disabledOptimisticOptsMetric.increment(); } }
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.ptx/src/com/oracle/graal/ptx/PTX.java Tue Feb 19 17:41:10 2013 +0100 @@ -0,0 +1,119 @@ +/* + * 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.ptx; + +import static com.oracle.graal.api.code.MemoryBarriers.*; +import static com.oracle.graal.api.code.Register.RegisterFlag.*; + +import java.nio.*; + +import com.oracle.graal.api.code.*; +import com.oracle.graal.api.code.Register.*; + +/** + * Represents the PTX architecture. + */ +public class PTX extends Architecture { + + // @formatter:off + + /* + * Register State Space + * + * Registers (.reg state space) are fast storage locations. The number of + * registers is limited, and will vary from platform to platform. When the + * limit is exceeded, register variables will be spilled to memory, causing + * changes in performance. For each architecture, there is a recommended + * maximum number of registers to use (see the "CUDA Programming Guide" for + * details). + */ + + // General purpose registers + public static final Register r0 = new Register(0, 0, 8, "r0", CPU, RegisterFlag.Byte); + public static final Register r1 = new Register(1, 1, 8, "r1", CPU, RegisterFlag.Byte); + public static final Register r2 = new Register(2, 2, 8, "r2", CPU, RegisterFlag.Byte); + public static final Register r3 = new Register(3, 3, 8, "r3", CPU, RegisterFlag.Byte); + public static final Register r4 = new Register(4, 4, 8, "r4", CPU, RegisterFlag.Byte); + public static final Register r5 = new Register(5, 5, 8, "r5", CPU, RegisterFlag.Byte); + public static final Register r6 = new Register(6, 6, 8, "r6", CPU, RegisterFlag.Byte); + public static final Register r7 = new Register(7, 7, 8, "r7", CPU, RegisterFlag.Byte); + + public static final Register r8 = new Register(8, 8, 8, "r8", CPU, RegisterFlag.Byte); + public static final Register r9 = new Register(9, 9, 8, "r9", CPU, RegisterFlag.Byte); + public static final Register r10 = new Register(10, 10, 8, "r10", CPU, RegisterFlag.Byte); + public static final Register r11 = new Register(11, 11, 8, "r11", CPU, RegisterFlag.Byte); + public static final Register r12 = new Register(12, 12, 8, "r12", CPU, RegisterFlag.Byte); + public static final Register r13 = new Register(13, 13, 8, "r13", CPU, RegisterFlag.Byte); + public static final Register r14 = new Register(14, 14, 8, "r14", CPU, RegisterFlag.Byte); + public static final Register r15 = new Register(15, 15, 8, "r15", CPU, RegisterFlag.Byte); + + public static final Register[] gprRegisters = { + r0, r1, r2, r3, r4, r5, r6, r7, + r8, r9, r10, r11, r12, r13, r14, r15 + }; + + // Floating point registers + public static final Register f0 = new Register(16, 0, 8, "f0", FPU); + public static final Register f1 = new Register(17, 1, 8, "f1", FPU); + public static final Register f2 = new Register(18, 2, 8, "f2", FPU); + public static final Register f3 = new Register(19, 3, 8, "f3", FPU); + public static final Register f4 = new Register(20, 4, 8, "f4", FPU); + public static final Register f5 = new Register(21, 5, 8, "f5", FPU); + public static final Register f6 = new Register(22, 6, 8, "f6", FPU); + public static final Register f7 = new Register(23, 7, 8, "f7", FPU); + + public static final Register f8 = new Register(24, 8, 8, "f8", FPU); + public static final Register f9 = new Register(25, 9, 8, "f9", FPU); + public static final Register f10 = new Register(26, 10, 8, "f10", FPU); + public static final Register f11 = new Register(27, 11, 8, "f11", FPU); + public static final Register f12 = new Register(28, 12, 8, "f12", FPU); + public static final Register f13 = new Register(29, 13, 8, "f13", FPU); + public static final Register f14 = new Register(30, 14, 8, "f14", FPU); + public static final Register f15 = new Register(31, 15, 8, "f15", FPU); + + public static final Register[] fpuRegisters = { + f0, f1, f2, f3, f4, f5, f6, f7, + f8, f9, f10, f11, f12, f13, f14, f15 + }; + + public static final Register[] allRegisters = { + // GPR + r0, r1, r2, r3, r4, r5, r6, r7, + r8, r9, r10, r11, r12, r13, r14, r15, + // FPU + f0, f1, f2, f3, f4, f5, f6, f7, + f8, f9, f10, f11, f12, f13, f14, f15 + }; + + public PTX() { + super("PTX", + 8, + ByteOrder.LITTLE_ENDIAN, + allRegisters, + LOAD_STORE | STORE_STORE, + 0, + r15.encoding + 1, + 8); + } + // @formatter:on +}
--- a/graal/com.oracle.graal.snippets.test/src/com/oracle/graal/snippets/InstanceOfTest.java Tue Feb 19 17:27:02 2013 +0100 +++ b/graal/com.oracle.graal.snippets.test/src/com/oracle/graal/snippets/InstanceOfTest.java Tue Feb 19 17:41:10 2013 +0100 @@ -280,7 +280,7 @@ @Test public void test10() { Mark[] noMarks = {}; - Call callAt63 = new Call("ignore", 63, 5, true, null); + Call callAt63 = new Call(null, 63, 5, true, null); Mark markAt63 = new Mark(63, "1", noMarks); test("compareSites", callAt63, callAt63); test("compareSites", callAt63, markAt63);
--- a/mx/projects Tue Feb 19 17:27:02 2013 +0100 +++ b/mx/projects Tue Feb 19 17:41:10 2013 +0100 @@ -62,6 +62,13 @@ project@com.oracle.graal.amd64@checkstyle=com.oracle.graal.graph project@com.oracle.graal.amd64@javaCompliance=1.7 +# graal.ptx +project@com.oracle.graal.ptx@subDir=graal +project@com.oracle.graal.ptx@sourceDirs=src +project@com.oracle.graal.ptx@dependencies=com.oracle.graal.api.code +project@com.oracle.graal.ptx@checkstyle=com.oracle.graal.graph +project@com.oracle.graal.ptx@javaCompliance=1.7 + # graal.sparc project@com.oracle.graal.sparc@subDir=graal project@com.oracle.graal.sparc@sourceDirs=src @@ -137,6 +144,13 @@ project@com.oracle.graal.lir.amd64@checkstyle=com.oracle.graal.graph project@com.oracle.graal.lir.amd64@javaCompliance=1.7 +# graal.lir.ptx +project@com.oracle.graal.lir.ptx@subDir=graal +project@com.oracle.graal.lir.ptx@sourceDirs=src +project@com.oracle.graal.lir.ptx@dependencies=com.oracle.graal.lir,com.oracle.graal.asm.ptx +project@com.oracle.graal.lir.ptx@checkstyle=com.oracle.graal.graph +project@com.oracle.graal.lir.ptx@javaCompliance=1.7 + # graal.lir.sparc project@com.oracle.graal.lir.sparc@subDir=graal project@com.oracle.graal.lir.sparc@sourceDirs=src @@ -221,6 +235,13 @@ project@com.oracle.graal.compiler.amd64@checkstyle=com.oracle.graal.graph project@com.oracle.graal.compiler.amd64@javaCompliance=1.7 +# graal.compiler.ptx +project@com.oracle.graal.compiler.ptx@subDir=graal +project@com.oracle.graal.compiler.ptx@sourceDirs=src +project@com.oracle.graal.compiler.ptx@dependencies=com.oracle.graal.compiler,com.oracle.graal.lir.ptx +project@com.oracle.graal.compiler.ptx@checkstyle=com.oracle.graal.graph +project@com.oracle.graal.compiler.ptx@javaCompliance=1.7 + # graal.compiler.sparc project@com.oracle.graal.compiler.sparc@subDir=graal project@com.oracle.graal.compiler.sparc@sourceDirs=src @@ -297,6 +318,13 @@ project@com.oracle.graal.asm.amd64.test@checkstyle=com.oracle.graal.graph project@com.oracle.graal.asm.amd64.test@javaCompliance=1.7 +# graal.asm.ptx +project@com.oracle.graal.asm.ptx@subDir=graal +project@com.oracle.graal.asm.ptx@sourceDirs=src +project@com.oracle.graal.asm.ptx@dependencies=com.oracle.graal.asm +project@com.oracle.graal.asm.ptx@checkstyle=com.oracle.graal.graph +project@com.oracle.graal.asm.ptx@javaCompliance=1.7 + # graal.asm.sparc project@com.oracle.graal.asm.sparc@subDir=graal project@com.oracle.graal.asm.sparc@sourceDirs=src
--- a/src/share/vm/classfile/systemDictionary.hpp Tue Feb 19 17:27:02 2013 +0100 +++ b/src/share/vm/classfile/systemDictionary.hpp Tue Feb 19 17:41:10 2013 +0100 @@ -186,6 +186,7 @@ do_klass(GraalBitMap_klass, java_util_BitSet, Opt) \ /* graal.hotspot */ \ do_klass(HotSpotCompilationResult_klass, com_oracle_graal_hotspot_HotSpotCompilationResult, Opt) \ + do_klass(HotSpotRuntimeCallTarget_klass, com_oracle_graal_hotspot_HotSpotRuntimeCallTarget, Opt) \ do_klass(HotSpotCodeInfo_klass, com_oracle_graal_hotspot_meta_HotSpotCodeInfo, Opt) \ do_klass(HotSpotInstalledCode_klass, com_oracle_graal_hotspot_meta_HotSpotInstalledCode, Opt) \ do_klass(HotSpotJavaType_klass, com_oracle_graal_hotspot_meta_HotSpotJavaType, Opt) \
--- a/src/share/vm/classfile/vmSymbols.hpp Tue Feb 19 17:27:02 2013 +0100 +++ b/src/share/vm/classfile/vmSymbols.hpp Tue Feb 19 17:41:10 2013 +0100 @@ -295,6 +295,7 @@ template(com_oracle_graal_hotspot_HotSpotKlassOop, "com/oracle/graal/hotspot/HotSpotKlassOop") \ template(com_oracle_graal_hotspot_HotSpotOptions, "com/oracle/graal/hotspot/HotSpotOptions") \ template(com_oracle_graal_hotspot_HotSpotCompilationResult, "com/oracle/graal/hotspot/HotSpotCompilationResult") \ + template(com_oracle_graal_hotspot_HotSpotRuntimeCallTarget, "com/oracle/graal/hotspot/HotSpotRuntimeCallTarget") \ template(com_oracle_graal_hotspot_bridge_VMToCompiler, "com/oracle/graal/hotspot/bridge/VMToCompiler") \ template(com_oracle_graal_hotspot_meta_HotSpotCodeInfo, "com/oracle/graal/hotspot/meta/HotSpotCodeInfo") \ template(com_oracle_graal_hotspot_meta_HotSpotInstalledCode, "com/oracle/graal/hotspot/meta/HotSpotInstalledCode") \
--- a/src/share/vm/graal/graalCodeInstaller.cpp Tue Feb 19 17:27:02 2013 +0100 +++ b/src/share/vm/graal/graalCodeInstaller.cpp Tue Feb 19 17:41:10 2013 +0100 @@ -29,7 +29,6 @@ #include "graal/graalCodeInstaller.hpp" #include "graal/graalJavaAccess.hpp" #include "graal/graalCompilerToVM.hpp" -#include "graal/graalVmIds.hpp" #include "graal/graalRuntime.hpp" #include "asm/register.hpp" #include "classfile/vmSymbols.hpp" @@ -209,7 +208,7 @@ return new ConstantOopWriteValue(JNIHandles::make_local(obj)); } } else if (type == T_ADDRESS) { - return new ConstantLongValue(prim); + ShouldNotReachHere(); } tty->print("%i", type); } else if (value->is_a(VirtualObject::klass())) { @@ -362,7 +361,7 @@ const char* cname = java_lang_String::as_utf8_string(_name); blob = BufferBlob::create(strdup(cname), &buffer); // this is leaking strings... but only a limited number of stubs will be created IF_TRACE_graal_3 Disassembler::decode((CodeBlob*) blob); - id = VmIds::addStub(blob->code_begin()); + id = (jlong)blob->code_begin(); } void CodeInstaller::initialize_fields(oop comp_result, methodHandle method) { @@ -606,7 +605,7 @@ oop hotspot_method = NULL; // JavaMethod oop global_stub = NULL; - if (target_klass->is_subclass_of(SystemDictionary::Long_klass())) { + if (target_klass->is_subclass_of(SystemDictionary::HotSpotRuntimeCallTarget_klass())) { global_stub = target; } else { hotspot_method = target; @@ -659,21 +658,20 @@ } if (global_stub != NULL) { - assert(java_lang_boxing_object::is_instance(global_stub, T_LONG), "global_stub needs to be of type Long"); - + jlong global_stub_destination = HotSpotRuntimeCallTarget::address(global_stub); if (inst->is_call()) { // NOTE: for call without a mov, the offset must fit a 32-bit immediate // see also CompilerToVM.getMaxCallTargetOffset() NativeCall* call = nativeCall_at((address) (inst)); - call->set_destination(VmIds::getStub(global_stub)); + call->set_destination((address) global_stub_destination); _instructions->relocate(call->instruction_address(), runtime_call_Relocation::spec(), Assembler::call32_operand); } else if (inst->is_mov_literal64()) { NativeMovConstReg* mov = nativeMovConstReg_at((address) (inst)); - mov->set_data((intptr_t) VmIds::getStub(global_stub)); + mov->set_data((intptr_t) global_stub_destination); _instructions->relocate(mov->instruction_address(), runtime_call_Relocation::spec(), Assembler::imm_operand); } else { NativeJump* jump = nativeJump_at((address) (inst)); - jump->set_jump_destination(VmIds::getStub(global_stub)); + jump->set_jump_destination((address) global_stub_destination); _instructions->relocate((address)inst, runtime_call_Relocation::spec(), Assembler::call32_operand); } TRACE_graal_3("relocating (stub) at %p", inst);
--- a/src/share/vm/graal/graalCompiler.cpp Tue Feb 19 17:27:02 2013 +0100 +++ b/src/share/vm/graal/graalCompiler.cpp Tue Feb 19 17:41:10 2013 +0100 @@ -28,7 +28,6 @@ #include "graal/graalJavaAccess.hpp" #include "graal/graalVMToCompiler.hpp" #include "graal/graalCompilerToVM.hpp" -#include "graal/graalVmIds.hpp" #include "graal/graalEnv.hpp" #include "graal/graalRuntime.hpp" #include "runtime/arguments.hpp" @@ -188,7 +187,7 @@ } Handle GraalCompiler::get_JavaType(Symbol* klass_name, TRAPS) { - return VMToCompiler::createUnresolvedJavaType(VmIds::toString<Handle>(klass_name, THREAD), THREAD); + return VMToCompiler::createUnresolvedJavaType(java_lang_String::create_from_symbol(klass_name, THREAD), THREAD); } Handle GraalCompiler::get_JavaTypeFromSignature(Symbol* signature, KlassHandle loading_klass, TRAPS) { @@ -255,12 +254,12 @@ } Handle GraalCompiler::get_JavaType(KlassHandle klass, TRAPS) { - Handle name = VmIds::toString<Handle>(klass->name(), THREAD); + Handle name = java_lang_String::create_from_symbol(klass->name(), THREAD); return createHotSpotResolvedObjectType(klass, name, CHECK_NULL); } Handle GraalCompiler::get_JavaField(int offset, int flags, Symbol* field_name, Handle field_holder, Handle field_type, TRAPS) { - Handle name = VmIds::toString<Handle>(field_name, CHECK_NULL); + Handle name = java_lang_String::create_from_symbol(field_name, CHECK_NULL); return VMToCompiler::createJavaField(field_holder, name, field_type, offset, flags, false, CHECK_NULL); }
--- a/src/share/vm/graal/graalCompilerToVM.cpp Tue Feb 19 17:27:02 2013 +0100 +++ b/src/share/vm/graal/graalCompilerToVM.cpp Tue Feb 19 17:41:10 2013 +0100 @@ -37,7 +37,6 @@ #include "graal/graalJavaAccess.hpp" #include "graal/graalCodeInstaller.hpp" #include "graal/graalVMToCompiler.hpp" -#include "graal/graalVmIds.hpp" Method* getMethodFromHotSpotMethod(oop hotspot_method) { @@ -120,7 +119,7 @@ C2V_VMENTRY(jstring, getSignature, (JNIEnv *env, jobject, jlong metaspace_method)) Method* method = asMethod(metaspace_method); assert(method != NULL && method->signature() != NULL, "signature required"); - return VmIds::toString<jstring>(method->signature(), THREAD); + return (jstring)JNIHandles::make_local(java_lang_String::create_from_symbol(method->signature(), THREAD)()); C2V_END C2V_VMENTRY(jobjectArray, initializeExceptionHandlers, (JNIEnv *, jobject, jlong metaspace_method, jobjectArray java_handlers)) @@ -271,7 +270,7 @@ C2V_VMENTRY(void, initializeMethod,(JNIEnv *, jobject, jlong metaspace_method, jobject hotspot_method)) methodHandle method = asMethod(metaspace_method); - Handle name = VmIds::toString<Handle>(method->name(), CHECK); + Handle name = java_lang_String::create_from_symbol(method->name(), CHECK); InstanceKlass::cast(HotSpotResolvedJavaMethod::klass())->initialize(CHECK); HotSpotResolvedJavaMethod::set_name(hotspot_method, name()); HotSpotResolvedJavaMethod::set_codeSize(hotspot_method, method->code_size()); @@ -323,8 +322,8 @@ C2V_VMENTRY(jobject, lookupType, (JNIEnv *env, jobject, jstring jname, jobject accessingClass, jboolean eagerResolve)) ResourceMark rm; - Symbol* nameSymbol = VmIds::toSymbol(jname); Handle name = JNIHandles::resolve(jname); + Symbol* nameSymbol = java_lang_String::as_symbol(name, THREAD); assert(nameSymbol != NULL, "name to symbol creation failed"); oop result = NULL; @@ -432,8 +431,8 @@ return JNIHandles::make_local(THREAD, VMToCompiler::createResolvedJavaMethod(holder, method(), THREAD)); } else { // Get the method's name and signature. - Handle name = VmIds::toString<Handle>(cp->name_ref_at(index), CHECK_NULL); - Handle signature = VmIds::toString<Handle>(cp->signature_ref_at(index), CHECK_NULL); + Handle name = java_lang_String::create_from_symbol(cp->name_ref_at(index), CHECK_NULL); + Handle signature = java_lang_String::create_from_symbol(cp->signature_ref_at(index), CHECK_NULL); int holder_index = cp->klass_ref_index_at(index); Handle type = GraalCompiler::get_JavaType(cp, holder_index, cp->pool_holder(), CHECK_NULL); return JNIHandles::make_local(THREAD, VMToCompiler::createUnresolvedJavaMethod(name, signature, type, THREAD)); @@ -514,8 +513,8 @@ assert(JNIHandles::resolve(resolved_type) != NULL, ""); Klass* klass = java_lang_Class::as_Klass(HotSpotResolvedObjectType::javaMirror(resolved_type)); - Symbol* name_symbol = VmIds::toSymbol(name); - Symbol* signature_symbol = VmIds::toSymbol(signature); + Symbol* name_symbol = java_lang_String::as_symbol(JNIHandles::resolve(name), THREAD); + Symbol* signature_symbol = java_lang_String::as_symbol(JNIHandles::resolve(signature), THREAD); methodHandle method = klass->lookup_method(name_symbol, signature_symbol); if (method.is_null()) { if (TraceGraal >= 3) { @@ -551,7 +550,7 @@ Handle type = GraalCompiler::get_JavaTypeFromSignature(fs.signature(), k, Thread::current()); int flags = fs.access_flags().as_int(); bool internal = fs.access_flags().is_internal(); - Handle name = VmIds::toString<Handle>(fs.name(), Thread::current()); + Handle name = java_lang_String::create_from_symbol(fs.name(), Thread::current()); Handle field = VMToCompiler::createJavaField(JNIHandles::resolve(klass), name, type, fs.offset(), flags, internal, Thread::current()); fields.append(field()); } @@ -598,6 +597,7 @@ #define set_boolean(name, value) do { env->SetBooleanField(config, getFieldID(env, config, name, "Z"), value); } while (0) #define set_int(name, value) do { env->SetIntField(config, getFieldID(env, config, name, "I"), value); } while (0) #define set_long(name, value) do { env->SetLongField(config, getFieldID(env, config, name, "J"), value); } while (0) +#define set_stub(name, value) do { set_long(name, (jlong) value); } while (0) #define set_object(name, value) do { env->SetObjectField(config, getFieldID(env, config, name, "Ljava/lang/Object;"), value); } while (0) #define set_int_array(name, value) do { env->SetObjectField(config, getFieldID(env, config, name, "[I"), value); } while (0) @@ -707,41 +707,41 @@ set_int("layoutHelperHeaderSizeMask", Klass::_lh_header_size_mask); set_int("layoutHelperOffset", in_bytes(Klass::layout_helper_offset())); - set_long("debugStub", VmIds::addStub((address)warning)); - set_long("instanceofStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_slow_subtype_check_id))); - set_long("newInstanceStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_new_instance_id))); - set_long("newArrayStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_new_array_id))); - set_long("newMultiArrayStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_new_multi_array_id))); - set_long("identityHashCodeStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_identity_hash_code_id))); - set_long("threadIsInterruptedStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_thread_is_interrupted_id))); - set_long("inlineCacheMissStub", VmIds::addStub(SharedRuntime::get_ic_miss_stub())); - set_long("handleExceptionStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_handle_exception_nofpu_id))); - set_long("handleDeoptStub", VmIds::addStub(SharedRuntime::deopt_blob()->unpack())); - set_long("monitorEnterStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_monitorenter_id))); - set_long("monitorExitStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_monitorexit_id))); - set_long("verifyOopStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_verify_oop_id))); - set_long("vmErrorStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_vm_error_id))); - set_long("deoptimizeStub", VmIds::addStub(SharedRuntime::deopt_blob()->uncommon_trap())); - set_long("unwindExceptionStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_unwind_exception_call_id))); - set_long("osrMigrationEndStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_OSR_migration_end_id))); - set_long("registerFinalizerStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_register_finalizer_id))); - set_long("setDeoptInfoStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_set_deopt_info_id))); - set_long("createNullPointerExceptionStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_create_null_pointer_exception_id))); - set_long("createOutOfBoundsExceptionStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_create_out_of_bounds_exception_id))); - set_long("javaTimeMillisStub", VmIds::addStub(CAST_FROM_FN_PTR(address, os::javaTimeMillis))); - set_long("javaTimeNanosStub", VmIds::addStub(CAST_FROM_FN_PTR(address, os::javaTimeNanos))); - set_long("arithmeticFremStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_arithmetic_frem_id))); - set_long("arithmeticDremStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_arithmetic_drem_id))); - set_long("arithmeticSinStub", VmIds::addStub(CAST_FROM_FN_PTR(address, SharedRuntime::dsin))); - set_long("arithmeticCosStub", VmIds::addStub(CAST_FROM_FN_PTR(address, SharedRuntime::dcos))); - set_long("arithmeticTanStub", VmIds::addStub(CAST_FROM_FN_PTR(address, SharedRuntime::dtan))); - set_long("logPrimitiveStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_log_primitive_id))); - set_long("logObjectStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_log_object_id))); - set_long("logPrintfStub", VmIds::addStub(GraalRuntime::entry_for(GraalRuntime::graal_log_printf_id))); - set_long("aescryptEncryptBlockStub", VmIds::addStub(StubRoutines::aescrypt_encryptBlock())); - set_long("aescryptDecryptBlockStub", VmIds::addStub(StubRoutines::aescrypt_decryptBlock())); - set_long("cipherBlockChainingEncryptAESCryptStub", VmIds::addStub(StubRoutines::cipherBlockChaining_encryptAESCrypt())); - set_long("cipherBlockChainingDecryptAESCryptStub", VmIds::addStub(StubRoutines::cipherBlockChaining_decryptAESCrypt())); + set_stub("debugStub", (address)warning); + set_stub("instanceofStub", GraalRuntime::entry_for(GraalRuntime::graal_slow_subtype_check_id)); + set_stub("newInstanceStub", GraalRuntime::entry_for(GraalRuntime::graal_new_instance_id)); + set_stub("newArrayStub", GraalRuntime::entry_for(GraalRuntime::graal_new_array_id)); + set_stub("newMultiArrayStub", GraalRuntime::entry_for(GraalRuntime::graal_new_multi_array_id)); + set_stub("identityHashCodeStub", GraalRuntime::entry_for(GraalRuntime::graal_identity_hash_code_id)); + set_stub("threadIsInterruptedStub", GraalRuntime::entry_for(GraalRuntime::graal_thread_is_interrupted_id)); + set_stub("inlineCacheMissStub", SharedRuntime::get_ic_miss_stub()); + set_stub("handleExceptionStub", GraalRuntime::entry_for(GraalRuntime::graal_handle_exception_nofpu_id)); + set_stub("handleDeoptStub", SharedRuntime::deopt_blob()->unpack()); + set_stub("monitorEnterStub", GraalRuntime::entry_for(GraalRuntime::graal_monitorenter_id)); + set_stub("monitorExitStub", GraalRuntime::entry_for(GraalRuntime::graal_monitorexit_id)); + set_stub("verifyOopStub", GraalRuntime::entry_for(GraalRuntime::graal_verify_oop_id)); + set_stub("vmErrorStub", GraalRuntime::entry_for(GraalRuntime::graal_vm_error_id)); + set_stub("deoptimizeStub", SharedRuntime::deopt_blob()->uncommon_trap()); + set_stub("unwindExceptionStub", GraalRuntime::entry_for(GraalRuntime::graal_unwind_exception_call_id)); + set_stub("osrMigrationEndStub", GraalRuntime::entry_for(GraalRuntime::graal_OSR_migration_end_id)); + set_stub("registerFinalizerStub", GraalRuntime::entry_for(GraalRuntime::graal_register_finalizer_id)); + set_stub("setDeoptInfoStub", GraalRuntime::entry_for(GraalRuntime::graal_set_deopt_info_id)); + set_stub("createNullPointerExceptionStub", GraalRuntime::entry_for(GraalRuntime::graal_create_null_pointer_exception_id)); + set_stub("createOutOfBoundsExceptionStub", GraalRuntime::entry_for(GraalRuntime::graal_create_out_of_bounds_exception_id)); + set_stub("javaTimeMillisStub", CAST_FROM_FN_PTR(address, os::javaTimeMillis)); + set_stub("javaTimeNanosStub", CAST_FROM_FN_PTR(address, os::javaTimeNanos)); + set_stub("arithmeticFremStub", GraalRuntime::entry_for(GraalRuntime::graal_arithmetic_frem_id)); + set_stub("arithmeticDremStub", GraalRuntime::entry_for(GraalRuntime::graal_arithmetic_drem_id)); + set_stub("arithmeticSinStub", CAST_FROM_FN_PTR(address, SharedRuntime::dsin)); + set_stub("arithmeticCosStub", CAST_FROM_FN_PTR(address, SharedRuntime::dcos)); + set_stub("arithmeticTanStub", CAST_FROM_FN_PTR(address, SharedRuntime::dtan)); + set_stub("logPrimitiveStub", GraalRuntime::entry_for(GraalRuntime::graal_log_primitive_id)); + set_stub("logObjectStub", GraalRuntime::entry_for(GraalRuntime::graal_log_object_id)); + set_stub("logPrintfStub", GraalRuntime::entry_for(GraalRuntime::graal_log_printf_id)); + set_stub("aescryptEncryptBlockStub", StubRoutines::aescrypt_encryptBlock()); + set_stub("aescryptDecryptBlockStub", StubRoutines::aescrypt_decryptBlock()); + set_stub("cipherBlockChainingEncryptAESCryptStub", StubRoutines::cipherBlockChaining_encryptAESCrypt()); + set_stub("cipherBlockChainingDecryptAESCryptStub", StubRoutines::cipherBlockChaining_decryptAESCrypt()); set_int("deoptReasonNone", Deoptimization::Reason_none); set_int("deoptReasonNullCheck", Deoptimization::Reason_null_check);
--- a/src/share/vm/graal/graalJavaAccess.hpp Tue Feb 19 17:27:02 2013 +0100 +++ b/src/share/vm/graal/graalJavaAccess.hpp Tue Feb 19 17:41:10 2013 +0100 @@ -88,6 +88,9 @@ oop_field(HotSpotCompilationResult, sites, "[Lcom/oracle/graal/api/code/CompilationResult$Site;") \ oop_field(HotSpotCompilationResult, exceptionHandlers, "[Lcom/oracle/graal/api/code/CompilationResult$ExceptionHandler;") \ end_class \ + start_class(HotSpotRuntimeCallTarget) \ + long_field(HotSpotRuntimeCallTarget, address) \ + end_class \ start_class(ExceptionHandler) \ int_field(ExceptionHandler, startBCI) \ int_field(ExceptionHandler, endBCI) \ @@ -122,7 +125,7 @@ int_field(CompilationResult_Site, pcOffset) \ end_class \ start_class(CompilationResult_Call) \ - oop_field(CompilationResult_Call, target, "Ljava/lang/Object;") \ + oop_field(CompilationResult_Call, target, "Lcom/oracle/graal/api/meta/InvokeTarget;") \ oop_field(CompilationResult_Call, debugInfo, "Lcom/oracle/graal/api/code/DebugInfo;") \ end_class \ start_class(CompilationResult_DataPatch) \
--- a/src/share/vm/graal/graalVmIds.cpp Tue Feb 19 17:27:02 2013 +0100 +++ /dev/null Thu Jan 01 00:00:00 1970 +0000 @@ -1,37 +0,0 @@ -/* - * Copyright (c) 2011, Oracle and/or its affiliates. All rights reserved. - * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. - * - * This code is free software; you can redistribute it and/or modify it - * under the terms of the GNU General Public License version 2 only, as - * published by the Free Software Foundation. - * - * This code is distributed in the hope that it will be useful, but WITHOUT - * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or - * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License - * version 2 for more details (a copy is included in the LICENSE file that - * accompanied this code). - * - * You should have received a copy of the GNU General Public License version - * 2 along with this work; if not, write to the Free Software Foundation, - * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. - * - * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA - * or visit www.oracle.com if you need additional information or have any - * questions. - */ - -#include "precompiled.hpp" -#include "graal/graalVmIds.hpp" -#include "ci/ciUtilities.hpp" - -// VmIds implementation - -jlong VmIds::addStub(address stub) { - return (jlong)stub; -} - -address VmIds::getStub(jlong id) { - return (address)id; -} -
--- a/src/share/vm/graal/graalVmIds.hpp Tue Feb 19 17:27:02 2013 +0100 +++ /dev/null Thu Jan 01 00:00:00 1970 +0000 @@ -1,86 +0,0 @@ -/* - * Copyright (c) 2011, Oracle and/or its affiliates. All rights reserved. - * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. - * - * This code is free software; you can redistribute it and/or modify it - * under the terms of the GNU General Public License version 2 only, as - * published by the Free Software Foundation. - * - * This code is distributed in the hope that it will be useful, but WITHOUT - * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or - * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License - * version 2 for more details (a copy is included in the LICENSE file that - * accompanied this code). - * - * You should have received a copy of the GNU General Public License version - * 2 along with this work; if not, write to the Free Software Foundation, - * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. - * - * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA - * or visit www.oracle.com if you need additional information or have any - * questions. - */ - -#ifndef SHARE_VM_GRAAL_GRAAL_VM_IDS_HPP -#define SHARE_VM_GRAAL_GRAAL_VM_IDS_HPP - -#include "memory/allocation.hpp" -#include "utilities/growableArray.hpp" -#include "oops/oop.hpp" -#include "runtime/handles.hpp" -#include "runtime/thread.hpp" -#include "classfile/javaClasses.hpp" -#include "runtime/jniHandles.hpp" - -class VmIds : public AllStatic { - -public: - // Adds a stub address, and returns the corresponding vmId (which is of type STUB) - static jlong addStub(address stub); - - // Returns the stub address with the given vmId - static address getStub(jlong id); - - // Returns the stub address with the given vmId taken from a java.lang.Long - static address getStub(oop id); - - // Helper function to convert a symbol to a java.lang.String object - template <typename T> static T toString(Symbol* symbol, TRAPS); - - // Helper function to convert a java.lang.String object to a symbol (this will return NULL if the symbol doesn't exist in the system) - static Symbol* toSymbol(jstring string); - - // Helper function to get the contents of a java.lang.Long - static jlong getBoxedLong(oop obj); -}; - -inline address VmIds::getStub(oop obj) { - return getStub(getBoxedLong(obj)); -} - -template <> inline Handle VmIds::toString<Handle>(Symbol* symbol, TRAPS) { - return java_lang_String::create_from_symbol(symbol, THREAD); -} - -template <> inline oop VmIds::toString<oop>(Symbol* symbol, TRAPS) { - return toString<Handle>(symbol, THREAD)(); -} - -template <> inline jstring VmIds::toString<jstring>(Symbol* symbol, TRAPS) { - return (jstring)JNIHandles::make_local(toString<oop>(symbol, THREAD)); -} - -template <> inline jobject VmIds::toString<jobject>(Symbol* symbol, TRAPS) { - return JNIHandles::make_local(toString<oop>(symbol, THREAD)); -} - -inline Symbol* VmIds::toSymbol(jstring string) { - return java_lang_String::as_symbol(JNIHandles::resolve(string), Thread::current()); -} - -inline jlong VmIds::getBoxedLong(oop obj) { - assert(obj->is_oop(true), "cannot unbox null or non-oop"); - return obj->long_field(java_lang_boxing_object::value_offset_in_bytes(T_LONG)); -} - -#endif // SHARE_VM_GRAAL_GRAAL_VM_IDS_HPP