OpenJDK / graal / graal-jvmci-8
changeset 12387:aff825fef0fd
Merge.
| author | Christian Humer <christian.humer@gmail.com> |
|---|---|
| date | Wed, 02 Oct 2013 13:26:31 +0200 |
| parents | 91dbb0b7dc8b 04b039d82e86 |
| children | 96c1d057a5ed |
| files | graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/DeoptimizationAction.java graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ArrayTest.java graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ControlTest.java graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64SafepointOp.java graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCSafepointOp.java graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/ArrayRangeWriteBarrier.java graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/G1ArrayRangePostWriteBarrier.java graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/G1ArrayRangePreWriteBarrier.java graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/G1PostWriteBarrier.java graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/G1PreWriteBarrier.java graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/G1ReferentFieldReadBarrier.java graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/SerialArrayRangeWriteBarrier.java graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/SerialWriteBarrier.java graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/ShortCircuitAndNode.java graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/ShortCircuitBooleanNode.java graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/WriteBarrier.java graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/extended/AbstractCallNode.java graal/com.oracle.graal.nodes/src/com/oracle/graal/nodes/extended/UnsafeArrayCastNode.java graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/InsertStateAfterPlaceholderPhase.java graal/com.oracle.graal.phases.common/src/com/oracle/graal/phases/common/SafepointInsertionPhase.java graal/com.oracle.graal.truffle/src/com/oracle/graal/truffle/phases/InlineTrivialGettersPhase.java graal/com.oracle.graal.truffle/src/com/oracle/graal/truffle/substitutions/SlowPathExceptionSubstitutions.java mx/.pylintrc src/gpu/ptx/vm/kernelArguments.cpp src/gpu/ptx/vm/kernelArguments.hpp |
| diffstat | 508 files changed, 12280 insertions(+), 7752 deletions(-) [+] |
line wrap: on
line diff
--- a/.hgignore Fri Sep 06 21:37:50 2013 +0200 +++ b/.hgignore Wed Oct 02 13:26:31 2013 +0200 @@ -1,5 +1,7 @@ ^mx/env ^mx/checkstyle-timestamps +^mx/eclipseinit.timestamp +^mx/netbeansinit.timestamp ^mx/eclipse-launches ^mx/ecj.jar ^mx/includes
--- a/README_GRAAL.txt Fri Sep 06 21:37:50 2013 +0200 +++ b/README_GRAAL.txt Wed Oct 02 13:26:31 2013 +0200 @@ -11,11 +11,30 @@ % mx build -This builds the 'product' version of HotSpot with the Graal modifications. -To build the debug or fastdebug versions: +There are a number of VM configurations supported by mx which can +be explicitly specified using the --vm option. However, you'll typically +want one of these VM configurations: - mx build debug - mx build fastdebug +1. The 'server' configuration is a standard HotSpot VM that includes the + runtime support for Graal but uses the existing compilers for normal + compilation (e.g., when the interpreter threshold is hit for a method). + Compilation with Graal is only done by explicit requests to the + Graal API. This is how Truffle uses Graal. + +2. The 'graal' configuration is a VM where all compilation is performed + by Graal and no other compilers are built into the VM binary. This + VM will bootstrap Graal itself at startup unless the -XX:-BootstrapGraal + VM option is given. + +Unless you use the --vm option with the build command, you will be presented +with a dialogue to choose one of the above VM configurations for the build +as well as have the option to make it your default for subsequent commands +that need a VM specified. + +To build the debug or fastdebug builds: + +% mx --vmbuild debug build +% mx --vmbuild fastdebug build Running Graal ------------- @@ -24,47 +43,28 @@ % mx vm ... -To select the fastdebug or debug versions of the VM: +To select the fastdebug or debug builds of the VM: -% mx --fastdebug vm ... -% mx --debug vm ... +% mx --vmbuild fastdebug vm ... +% mx --vmbuild debug vm ... -Graal has an optional bootstrap step where it compiles itself before -compiling any application code. This bootstrap step currently takes about 7 seconds -on a fast x64 machine. It's useful to disable this bootstrap step when running small -programs with the -XX:-BootstrapGraal options. For example: +Other VM Configurations +----------------------- -% mx vm -XX:-BootstrapGraal ... +In addition to the VM configurations described above, there are +VM configurations that omit all VM support for Graal: +% mx --vm server-nograal build +% mx --vm server-nograal vm -version +java version "1.7.0_25" +Java(TM) SE Runtime Environment (build 1.7.0_25-b15) +OpenJDK 64-Bit Server VM (build 25.0-b43-internal, mixed mode) -Other Build Configurations --------------------------- +% mx --vm client-nograal build +% mx --vm client-nograal vm -version +java version "1.7.0_25" +Java(TM) SE Runtime Environment (build 1.7.0_25-b15) +OpenJDK 64-Bit Cleint VM (build 25.0-b43-internal, mixed mode) -By default the build commands above create a HotSpot binary where Graal -is the only compiler. This binary is the Graal VM binary and identifies as -such with the -version option: - -% mx vm -XX:-BootstrapGraal -version -java version "1.7.0_07" -Java(TM) SE Runtime Environment (build 1.7.0_07-b10) -OpenJDK 64-Bit Graal VM (build 25.0-b10-internal, mixed mode) - -It's also possible to build and execute the standard HotSpot binaries -using the --vm option: - -% mx --vm server build -% mx --vm server vm -version -java version "1.7.0_07" -Java(TM) SE Runtime Environment (build 1.7.0_07-b10) -OpenJDK 64-Bit Server VM (build 25.0-b10-internal, mixed mode) - -These standard binaries still include the code necessary to support use of the -Graal compiler for explicit compilation requests. However, in this configuration -the Graal compiler will not service VM issued compilation requests (e.g., upon -counter overflow in the interpreter). - -To build a HotSpot binary that completely omits all VM support for Graal, -define an environment variable OMIT_GRAAL (its value does not matter) and build -with the --vm option as above (doing a clean first if necessary): - -% env OMIT_GRAAL= mx --vm server build +These configurations aim to match as closely as possible the +VM(s) included in the OpenJDK binaries one can download. \ No newline at end of file
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/ArithmeticOperation.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/ArithmeticOperation.java Wed Oct 02 13:26:31 2013 +0200 @@ -22,9 +22,13 @@ */ package com.oracle.graal.api.code; +import com.oracle.graal.api.meta.*; + /** * An {@code ArithmeticOperation} is an operation that does primitive value arithmetic without side * effect. */ public interface ArithmeticOperation { + + Constant evalConst(Constant... inputs); }
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CallingConvention.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CallingConvention.java Wed Oct 02 13:26:31 2013 +0200 @@ -151,7 +151,7 @@ private boolean verify() { for (int i = 0; i < argumentLocations.length; i++) { Value location = argumentLocations[i]; - assert isStackSlot(location) || isRegister(location); + assert isStackSlot(location) || isAllocatableValue(location); } return true; }
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CodeCacheProvider.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/CodeCacheProvider.java Wed Oct 02 13:26:31 2013 +0200 @@ -72,13 +72,6 @@ ForeignCallLinkage lookupForeignCall(ForeignCallDescriptor descriptor); /** - * Encodes a deoptimization action and a deoptimization reason in an integer value. - * - * @return the encoded value as an integer - */ - int encodeDeoptActionAndReason(DeoptimizationAction action, DeoptimizationReason reason); - - /** * Determines if a {@link DataPatch} should be created for a given * {@linkplain Constant#getPrimitiveAnnotation() annotated} primitive constant that part of a * {@link CompilationResult}. A data patch is always created for an object constant.
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/DebugInfo.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/DebugInfo.java Wed Oct 02 13:26:31 2013 +0200 @@ -44,7 +44,6 @@ private final BytecodePosition bytecodePosition; private final BitSet registerRefMap; private final BitSet frameRefMap; - private final short deoptimizationReason; private RegisterSaveLayout calleeSaveInfo; /** @@ -55,11 +54,10 @@ * @param registerRefMap the register map * @param frameRefMap the reference map for {@code frame}, which may be {@code null} */ - public DebugInfo(BytecodePosition codePos, BitSet registerRefMap, BitSet frameRefMap, short deoptimizationReason) { + public DebugInfo(BytecodePosition codePos, BitSet registerRefMap, BitSet frameRefMap) { this.bytecodePosition = codePos; this.registerRefMap = registerRefMap; this.frameRefMap = frameRefMap; - this.deoptimizationReason = deoptimizationReason; } /** @@ -127,15 +125,6 @@ } /** - * Identifies the reason in case a deoptimization happens at this program counter. - * - * @return the reason of the deoptimization - */ - public short getDeoptimizationReason() { - return deoptimizationReason; - } - - /** * Sets the map from the registers (in the caller's frame) to the slots where they are saved in * the current frame. */
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/DelegatingCodeCacheProvider.java Wed Oct 02 13:26:31 2013 +0200 @@ -0,0 +1,68 @@ +/* + * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * This code is distributed in the hope that it will be useful, but WITHOUT + * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or + * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License + * version 2 for more details (a copy is included in the LICENSE file that + * accompanied this code). + * + * You should have received a copy of the GNU General Public License version + * 2 along with this work; if not, write to the Free Software Foundation, + * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. + * + * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA + * or visit www.oracle.com if you need additional information or have any + * questions. + */ +package com.oracle.graal.api.code; + +import com.oracle.graal.api.meta.*; + +/** + * A {@link CodeCacheProvider} that delegates to another {@link CodeCacheProvider}. + */ +public class DelegatingCodeCacheProvider extends DelegatingMetaAccessProvider implements CodeCacheProvider { + + public DelegatingCodeCacheProvider(CodeCacheProvider delegate) { + super(delegate); + } + + @Override + protected CodeCacheProvider delegate() { + return (CodeCacheProvider) super.delegate(); + } + + public InstalledCode addMethod(ResolvedJavaMethod method, CompilationResult compResult) { + return delegate().addMethod(method, compResult); + } + + public String disassemble(CompilationResult compResult, InstalledCode installedCode) { + return delegate().disassemble(compResult, installedCode); + } + + public RegisterConfig lookupRegisterConfig() { + return delegate().lookupRegisterConfig(); + } + + public int getMinimumOutgoingSize() { + return delegate().getMinimumOutgoingSize(); + } + + public ForeignCallLinkage lookupForeignCall(ForeignCallDescriptor descriptor) { + return delegate().lookupForeignCall(descriptor); + } + + public boolean needsDataPatch(Constant constant) { + return delegate().needsDataPatch(constant); + } + + public TargetDescription getTarget() { + return delegate().getTarget(); + } +}
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/DeoptimizationAction.java Fri Sep 06 21:37:50 2013 +0200 +++ /dev/null Thu Jan 01 00:00:00 1970 +0000 @@ -1,70 +0,0 @@ -/* - * Copyright (c) 2012, Oracle and/or its affiliates. All rights reserved. - * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. - * - * This code is free software; you can redistribute it and/or modify it - * under the terms of the GNU General Public License version 2 only, as - * published by the Free Software Foundation. - * - * 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.code; - -/** - * Specifies the action that should be taken by the runtime in case a certain deoptimization is - * triggered. - */ -public enum DeoptimizationAction { - /** - * Do not invalidate the machine code. This is typically used when deoptimizing at a point where - * it's highly likely nothing will change the likelihood of the deoptimization happening again. - * For example, a compiled array allocation where the size is negative. - */ - None(false), - - /** - * Do not invalidate the machine code, but schedule a recompilation if this deoptimization is - * triggered too often. - */ - RecompileIfTooManyDeopts(true), - - /** - * Invalidate the machine code and reset the profiling information. - */ - InvalidateReprofile(true), - - /** - * Invalidate the machine code and immediately schedule a recompilation. This is typically used - * when deoptimizing to resolve an unresolved symbol in which case extra profiling is not - * required to determine that the deoptimization will not re-occur. - */ - InvalidateRecompile(true), - - /** - * Invalidate the machine code and stop compiling the outermost method of this compilation. - */ - InvalidateStopCompiling(true); - - private final boolean invalidatesCompilation; - - private DeoptimizationAction(boolean invalidatesCompilation) { - this.invalidatesCompilation = invalidatesCompilation; - } - - public boolean doesInvalidateCompilation() { - return invalidatesCompilation; - } - -}
--- a/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/SpeculationLog.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.api.code/src/com/oracle/graal/api/code/SpeculationLog.java Wed Oct 02 13:26:31 2013 +0200 @@ -36,11 +36,11 @@ public static final int MAX_CACHE_SIZE = 1 << 15; - private List<DeoptimizationReason> speculations = new ArrayList<>(); + private List<Object> speculations = new ArrayList<>(); private boolean[] map = new boolean[10]; - private Set<DeoptimizationReason> snapshot = new HashSet<>(); + private Set<Object> snapshot = new HashSet<>(); - public short addSpeculation(DeoptimizationReason reason) { + private short addSpeculation(Object reason) { short index = (short) speculations.indexOf(reason); if (index != -1) { // Nothing to add, reason already registered. @@ -68,7 +68,10 @@ } } - public boolean maySpeculate(DeoptimizationReason reason) { - return !snapshot.contains(reason); + public Constant maySpeculate(Object reason) { + if (snapshot.contains(reason)) { + return null; + } + return Constant.forShort(addSpeculation(reason)); } }
--- a/graal/com.oracle.graal.api.meta.test/src/com/oracle/graal/api/meta/test/TestMetaAccessProvider.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.api.meta.test/src/com/oracle/graal/api/meta/test/TestMetaAccessProvider.java Wed Oct 02 13:26:31 2013 +0200 @@ -107,15 +107,13 @@ @Test public void lookupArrayLengthTest() { for (Constant c : constants) { + Integer actual = runtime.lookupArrayLength(c); if (c.getKind() != Kind.Object || c.isNull() || !c.asObject().getClass().isArray()) { - try { - int length = runtime.lookupArrayLength(c); - fail("Expected " + IllegalArgumentException.class.getName() + " for " + c + ", not " + length); - } catch (IllegalArgumentException e) { - // pass - } + assertNull(actual); } else { - assertEquals(Array.getLength(c.asObject()), runtime.lookupArrayLength(c)); + assertNotNull(actual); + int actualInt = actual; + assertEquals(Array.getLength(c.asObject()), actualInt); } } }
--- a/graal/com.oracle.graal.api.meta.test/src/com/oracle/graal/api/meta/test/TestResolvedJavaMethod.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.api.meta.test/src/com/oracle/graal/api/meta/test/TestResolvedJavaMethod.java Wed Oct 02 13:26:31 2013 +0200 @@ -265,7 +265,9 @@ ResolvedJavaMethod method2 = runtime.lookupJavaMethod(getClass().getDeclaredMethod("nullPointerExceptionOnFirstLine", Object.class, String.class)); assertEquals(0, method1.getMaxStackSize()); // some versions of javac produce bytecode with a stacksize of 2 for this method - assertTrue(3 == method2.getMaxStackSize() || 2 == method2.getMaxStackSize()); + // JSR 292 also sometimes need one more stack slot + int method2StackSize = method2.getMaxStackSize(); + assertTrue(2 <= method2StackSize && method2StackSize <= 4); } private Method findTestMethod(Method apiMethod) {
--- a/graal/com.oracle.graal.api.meta.test/src/com/oracle/graal/api/meta/test/TestResolvedJavaType.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.api.meta.test/src/com/oracle/graal/api/meta/test/TestResolvedJavaType.java Wed Oct 02 13:26:31 2013 +0200 @@ -399,9 +399,11 @@ } for (Method m : c.getDeclaredMethods()) { if (!isStatic(m.getModifiers()) && !isPrivate(m.getModifiers())) { - Method overridden = vtable.methods.put(new NameAndSignature(m), m); - if (overridden != null) { - // println(m + " overrides " + overridden); + if (isAbstract(m.getModifiers())) { + // A subclass makes a concrete method in a superclass abstract + vtable.methods.remove(new NameAndSignature(m)); + } else { + vtable.methods.put(new NameAndSignature(m), m); } } } @@ -438,7 +440,13 @@ @Test public void resolveMethodTest() { for (Class c : classes) { - if (!c.isPrimitive() && !c.isInterface()) { + if (c.isInterface() || c.isPrimitive()) { + ResolvedJavaType type = runtime.lookupJavaType(c); + for (Method m : c.getDeclaredMethods()) { + ResolvedJavaMethod impl = type.resolveMethod(runtime.lookupJavaMethod(m)); + assertEquals(m.toString(), null, impl); + } + } else { ResolvedJavaType type = runtime.lookupJavaType(c); VTable vtable = getVTable(c); for (Method impl : vtable.methods.values()) { @@ -449,6 +457,11 @@ checkResolveMethod(type, m, i); } } + for (Method m : c.getDeclaredMethods()) { + ResolvedJavaMethod impl = type.resolveMethod(runtime.lookupJavaMethod(m)); + ResolvedJavaMethod expected = isAbstract(m.getModifiers()) ? null : impl; + assertEquals(type + " " + m.toString(), expected, impl); + } } } }
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/DelegatingMetaAccessProvider.java Wed Oct 02 13:26:31 2013 +0200 @@ -0,0 +1,93 @@ +/* + * 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.api.meta; + +import java.lang.reflect.*; + +/** + * A {@link MetaAccessProvider} that delegates to another {@link MetaAccessProvider}. + */ +public class DelegatingMetaAccessProvider implements MetaAccessProvider { + + private final MetaAccessProvider delegate; + + public DelegatingMetaAccessProvider(MetaAccessProvider delegate) { + this.delegate = delegate; + } + + protected MetaAccessProvider delegate() { + return delegate; + } + + public ResolvedJavaType lookupJavaType(Class<?> clazz) { + return delegate.lookupJavaType(clazz); + } + + public ResolvedJavaMethod lookupJavaMethod(Method reflectionMethod) { + return delegate.lookupJavaMethod(reflectionMethod); + } + + public ResolvedJavaMethod lookupJavaConstructor(Constructor reflectionConstructor) { + return delegate.lookupJavaConstructor(reflectionConstructor); + } + + public ResolvedJavaField lookupJavaField(Field reflectionField) { + return delegate.lookupJavaField(reflectionField); + } + + public ResolvedJavaType lookupJavaType(Constant constant) { + return delegate.lookupJavaType(constant); + } + + public Signature parseMethodDescriptor(String methodDescriptor) { + return delegate.parseMethodDescriptor(methodDescriptor); + } + + public boolean constantEquals(Constant x, Constant y) { + return delegate.constantEquals(x, y); + } + + public Integer lookupArrayLength(Constant array) { + return delegate.lookupArrayLength(array); + } + + public Constant readUnsafeConstant(Kind kind, Object base, long displacement, boolean compressible) { + return delegate.readUnsafeConstant(kind, base, displacement, compressible); + } + + public boolean isReexecutable(ForeignCallDescriptor descriptor) { + return delegate.isReexecutable(descriptor); + } + + public LocationIdentity[] getKilledLocations(ForeignCallDescriptor descriptor) { + return delegate.getKilledLocations(descriptor); + } + + public boolean canDeoptimize(ForeignCallDescriptor descriptor) { + return delegate.canDeoptimize(descriptor); + } + + public Constant encodeDeoptActionAndReason(DeoptimizationAction action, DeoptimizationReason reason) { + return delegate().encodeDeoptActionAndReason(action, reason); + } +}
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/DeoptimizationAction.java Wed Oct 02 13:26:31 2013 +0200 @@ -0,0 +1,70 @@ +/* + * Copyright (c) 2012, Oracle and/or its affiliates. All rights reserved. + * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. + * + * This code is free software; you can redistribute it and/or modify it + * under the terms of the GNU General Public License version 2 only, as + * published by the Free Software Foundation. + * + * 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; + +/** + * Specifies the action that should be taken by the runtime in case a certain deoptimization is + * triggered. + */ +public enum DeoptimizationAction { + /** + * Do not invalidate the machine code. This is typically used when deoptimizing at a point where + * it's highly likely nothing will change the likelihood of the deoptimization happening again. + * For example, a compiled array allocation where the size is negative. + */ + None(false), + + /** + * Do not invalidate the machine code, but schedule a recompilation if this deoptimization is + * triggered too often. + */ + RecompileIfTooManyDeopts(true), + + /** + * Invalidate the machine code and reset the profiling information. + */ + InvalidateReprofile(true), + + /** + * Invalidate the machine code and immediately schedule a recompilation. This is typically used + * when deoptimizing to resolve an unresolved symbol in which case extra profiling is not + * required to determine that the deoptimization will not re-occur. + */ + InvalidateRecompile(true), + + /** + * Invalidate the machine code and stop compiling the outermost method of this compilation. + */ + InvalidateStopCompiling(true); + + private final boolean invalidatesCompilation; + + private DeoptimizationAction(boolean invalidatesCompilation) { + this.invalidatesCompilation = invalidatesCompilation; + } + + public boolean doesInvalidateCompilation() { + return invalidatesCompilation; + } + +}
--- a/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/MetaAccessProvider.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/MetaAccessProvider.java Wed Oct 02 13:26:31 2013 +0200 @@ -79,11 +79,11 @@ boolean constantEquals(Constant x, Constant y); /** - * Returns the length of an array that is wrapped in a {@link Constant} object. - * - * @throws IllegalArgumentException if {@code array} is not an array + * Returns the length of an array that is wrapped in a {@link Constant} object. If {@code array} + * is not an array, or the array length is not available at this point, the return value is + * {@code null}. */ - int lookupArrayLength(Constant array); + Integer lookupArrayLength(Constant array); /** * Reads a value of this kind using a base address and a displacement. @@ -113,4 +113,11 @@ * Determines if deoptimization can occur during a given foreign call. */ boolean canDeoptimize(ForeignCallDescriptor descriptor); + + /** + * Encodes a deoptimization action and a deoptimization reason in an integer value. + * + * @return the encoded value as an integer + */ + Constant encodeDeoptActionAndReason(DeoptimizationAction action, DeoptimizationReason reason); }
--- a/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/ResolvedJavaField.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/ResolvedJavaField.java Wed Oct 02 13:26:31 2013 +0200 @@ -50,17 +50,21 @@ boolean isSynthetic(); /** - * Gets the constant value of this field for a given object, if available. + * Gets the constant value of this field. Note that a {@code static final} field may not be + * considered constant if its declaring class is not yet initialized or if it is a well known + * field that can be updated via other means (e.g., {@link System#setOut(java.io.PrintStream)}). * * @param receiver object from which this field's value is to be read. This value is ignored if * this field is static. - * @return the constant value of this field or {@code null} if the constant value is not - * available + * @return the constant value of this field or {@code null} if this field is not considered + * constant by the runtime */ Constant readConstantValue(Constant receiver); /** - * Gets the current value of this field for a given object, if available. + * Gets the current value of this field for a given object, if available. There is no guarantee + * that the same value will be returned by this method for a field unless the field is + * considered to be {@link #readConstantValue(Constant)} by the runtime. * * @param receiver object from which this field's value is to be read. This value is ignored if * this field is static.
--- a/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/ResolvedJavaMethod.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/ResolvedJavaMethod.java Wed Oct 02 13:26:31 2013 +0200 @@ -175,12 +175,14 @@ boolean canBeInlined(); /** - * Returns the LineNumberTable of this method. + * Returns the LineNumberTable of this method or null if this method does not have a line + * numbers table. */ LineNumberTable getLineNumberTable(); /** - * Returns the localvariable table of this method. + * Returns the local variable table of this method or null if this method does not have a local + * variable table. */ LocalVariableTable getLocalVariableTable();
--- a/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/ResolvedJavaType.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.api.meta/src/com/oracle/graal/api/meta/ResolvedJavaType.java Wed Oct 02 13:26:31 2013 +0200 @@ -112,7 +112,7 @@ int getModifiers(); /** - * Checks whether this type is initialized. If a type is initialized it implies that is was + * Checks whether this type is initialized. If a type is initialized it implies that it was * {@link #isLinked() linked} and that the static initializer has run. * * @return {@code true} if this type is initialized @@ -205,10 +205,13 @@ /** * Resolves the method implementation for virtual dispatches on objects of this dynamic type. + * This resolution process only searches "up" the class hierarchy of this type. A broader search + * that also walks "down" the hierarchy is implemented by + * {@link #findUniqueConcreteMethod(ResolvedJavaMethod)}. * * @param method the method to select the implementation of - * @return the method implementation that would be selected at runtime, or {@code null} if the - * runtime cannot resolve the method at this point in time. + * @return the concrete method that would be selected at runtime, or {@code null} if there is no + * concrete implementation of {@code method} in this type or any of its superclasses */ ResolvedJavaMethod resolveMethod(ResolvedJavaMethod method);
--- a/graal/com.oracle.graal.asm.amd64/src/com/oracle/graal/asm/amd64/AMD64Assembler.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.asm.amd64/src/com/oracle/graal/asm/amd64/AMD64Assembler.java Wed Oct 02 13:26:31 2013 +0200 @@ -747,7 +747,7 @@ emitOperandHelper(0, dst); } - private void jcc(ConditionFlag cc, int jumpTarget, boolean forceDisp32) { + public void jcc(ConditionFlag cc, int jumpTarget, boolean forceDisp32) { int shortSize = 2; int longSize = 6; long disp = jumpTarget - codeBuffer.position();
--- a/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAddress.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAddress.java Wed Oct 02 13:26:31 2013 +0200 @@ -22,6 +22,7 @@ */ package com.oracle.graal.asm.ptx; +import com.oracle.graal.lir.Variable; import com.oracle.graal.api.code.*; import com.oracle.graal.api.meta.*; @@ -31,7 +32,7 @@ */ public final class PTXAddress extends AbstractAddress { - private final Register base; + private final Variable base; private final long displacement; /** @@ -39,7 +40,7 @@ * * @param base the base register */ - public PTXAddress(Register base) { + public PTXAddress(Variable base) { this(base, 0); } @@ -50,7 +51,7 @@ * @param base the base register * @param displacement the displacement */ - public PTXAddress(Register base, long displacement) { + public PTXAddress(Variable base, long displacement) { this.base = base; this.displacement = displacement; } @@ -59,7 +60,7 @@ * @return Base register that defines the start of the address computation. If not present, is * denoted by {@link Value#ILLEGAL}. */ - public Register getBase() { + public Variable getBase() { return base; }
--- a/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXAssembler.java Wed Oct 02 13:26:31 2013 +0200 @@ -22,7 +22,17 @@ */ package com.oracle.graal.asm.ptx; -import com.oracle.graal.api.code.*; +import static com.oracle.graal.api.code.ValueUtil.*; + +import com.oracle.graal.api.code.Register; +import com.oracle.graal.api.code.RegisterConfig; +import com.oracle.graal.api.code.TargetDescription; +import com.oracle.graal.api.meta.Constant; +import com.oracle.graal.api.meta.Kind; +import com.oracle.graal.api.meta.Value; +import com.oracle.graal.nodes.calc.Condition; +import com.oracle.graal.graph.GraalInternalError; +import com.oracle.graal.lir.Variable; public class PTXAssembler extends AbstractPTXAssembler { @@ -30,535 +40,549 @@ super(target); } - public final void at() { - emitString("@%p" + " " + ""); + public enum ConditionOperator { + // @formatter:off + + // Signed integer operators + S_EQ("eq"), + S_NE("ne"), + S_LT("lt"), + S_LE("le"), + S_GT("gt"), + S_GE("ge"), + + // Unsigned integer operators + U_EQ("eq"), + U_NE("ne"), + U_LO("lo"), + U_LS("ls"), + U_HI("hi"), + U_HS("hs"), + + // Bit-size integer operators + B_EQ("eq"), + B_NE("ne"), + + // Floating-point operators + F_EQ("eq"), + F_NE("ne"), + F_LT("lt"), + F_LE("le"), + F_GT("gt"), + F_GE("ge"), + + // Floating-point operators accepting NaN + F_EQU("equ"), + F_NEU("neu"), + F_LTU("ltu"), + F_LEU("leu"), + F_GTU("gtu"), + F_GEU("geu"), + + // Floating-point operators testing for NaN + F_NUM("num"), + F_NAN("nan"); + + // @formatter:on + + private final String operator; + + private ConditionOperator(String op) { + this.operator = op; + } + + public String getOperator() { + return operator; + } } - public final void atq() { - emitString("@%q" + " " + ""); + public static class StandardFormat { + + protected Kind valueKind; + protected Variable dest; + protected Value source1; + protected Value source2; + private boolean logicInstruction = false; + + public StandardFormat(Variable dst, Value src1, Value src2) { + setDestination(dst); + setSource1(src1); + setSource2(src2); + setKind(dst.getKind()); + + // testAdd2B fails this assertion + // assert valueKind == src1.getKind(); + } + + public void setKind(Kind k) { + valueKind = k; + } + + public void setDestination(Variable var) { + assert var != null; + dest = var; + } + + public void setSource1(Value val) { + assert val != null; + source1 = val; + } + + public void setSource2(Value val) { + assert val != null; + source2 = val; + } + + public void setLogicInstruction(boolean b) { + logicInstruction = b; + } + + public String typeForKind(Kind k) { + if (logicInstruction) { + switch (k.getTypeChar()) { + case 's': + return "b16"; + case 'i': + return "b32"; + case 'j': + return "b64"; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } else { + switch (k.getTypeChar()) { + case 'z': + return "u8"; + case 'b': + return "s8"; + case 's': + return "s16"; + case 'c': + return "u16"; + case 'i': + return "s32"; + case 'f': + return "f32"; + case 'j': + return "s64"; + case 'd': + return "f64"; + case 'a': + return "u64"; + case '-': + return "u32"; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } + } + + public String emit() { + return (typeForKind(valueKind) + emitRegister(dest, true) + emitValue(source1, true) + emitValue(source2, false) + ";"); + } + + public String emitValue(Value v, boolean comma) { + assert v != null; + + if (isConstant(v)) { + return (emitConstant(v)); + } else { + return (emitRegister((Variable) v, comma)); + } + } + + public String emitRegister(Variable v, boolean comma) { + return (" %r" + v.index + (comma ? "," : "")); + } + + public String emitConstant(Value v) { + Constant constant = (Constant) v; + + switch (v.getKind().getTypeChar()) { + case 'i': + return (String.valueOf((int) constant.asLong())); + case 'f': + return (String.valueOf(constant.asFloat())); + case 'j': + return (String.valueOf(constant.asLong())); + case 'd': + return (String.valueOf(constant.asDouble())); + default: + throw GraalInternalError.shouldNotReachHere(); + } + } + } + + public static class SingleOperandFormat { + + protected Variable dest; + protected Value source; + + public SingleOperandFormat(Variable dst, Value src) { + setDestination(dst); + setSource(src); + } + + public void setDestination(Variable var) { + dest = var; + } + + public void setSource(Value var) { + source = var; + } + + public String typeForKind(Kind k) { + switch (k.getTypeChar()) { + case 'z': + return "u8"; + case 'b': + return "s8"; + case 's': + return "s16"; + case 'c': + return "u16"; + case 'i': + return "s32"; + case 'f': + return "f32"; + case 'j': + return "s64"; + case 'd': + return "f64"; + case 'a': + return "u64"; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } + + public String emit() { + return (typeForKind(dest.getKind()) + " " + emitVariable(dest) + ", " + emitValue(source) + ";"); + } + + public String emitValue(Value v) { + assert v != null; + + if (isConstant(v)) { + return (emitConstant(v)); + } else { + return (emitVariable((Variable) v)); + } + } + + public String emitConstant(Value v) { + Constant constant = (Constant) v; + + switch (v.getKind().getTypeChar()) { + case 'i': + return (String.valueOf((int) constant.asLong())); + case 'f': + return (String.valueOf(constant.asFloat())); + case 'j': + return (String.valueOf(constant.asLong())); + case 'd': + return (String.valueOf(constant.asDouble())); + default: + throw GraalInternalError.shouldNotReachHere(); + } + } + + public String emitVariable(Variable v) { + String name = v.getName(); + + if (name == null) { + return (" %r" + v.index); + } else { + return name; + } + } + } + + public static class ConversionFormat extends SingleOperandFormat { + + public ConversionFormat(Variable dst, Value src) { + super(dst, src); + } + + @Override + public String emit() { + return (typeForKind(dest.getKind()) + "." + typeForKind(source.getKind()) + " " + + emitVariable(dest) + ", " + emitValue(source) + ";"); + } + } + + public static class LoadStoreFormat extends StandardFormat { + + protected PTXStateSpace space; + + public LoadStoreFormat(PTXStateSpace space, Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + setStateSpace(space); + } + + public void setStateSpace(PTXStateSpace ss) { + space = ss; + } + + public String emitAddress(Value var, Value val) { + assert var instanceof Variable; + assert val instanceof Constant; + Constant constant = (Constant) val; + return ("[" + emitRegister((Variable) var, false) + " + " + constant.asBoxedValue() + "]"); + } + + @Override + public String emitRegister(Variable var, boolean comma) { + /* if (space == Parameter) { + return ("param" + var.index); + } else { + return ("%r" + var.index); + } */ + return ("%r" + var.index); + } + + public String emit(boolean isLoad) { + if (isLoad) { + return (space.getStateName() + "." + typeForKind(valueKind) + " " + + emitRegister(dest, false) + ", " + emitAddress(source1, source2) + ";"); + } else { + return (space.getStateName() + "." + typeForKind(valueKind) + " " + + emitAddress(source1, source2) + ", " + emitRegister(dest, false) + ";"); + } + } + } + + public static class Add extends StandardFormat { + + public Add(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + } + + public void emit(PTXAssembler asm) { + asm.emitString("add." + super.emit()); + } + } + + public static class And extends StandardFormat { + + public And(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + setLogicInstruction(true); + } + + public void emit(PTXAssembler asm) { + asm.emitString("and." + super.emit()); + } + } + + public static class Div extends StandardFormat { + + public Div(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + } + + public void emit(PTXAssembler asm) { + asm.emitString("div." + super.emit()); + } + } + + public static class Mul extends StandardFormat { + + public Mul(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + } + + public void emit(PTXAssembler asm) { + asm.emitString("mul.lo." + super.emit()); + } + } + + public static class Or extends StandardFormat { + + public Or(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + setLogicInstruction(true); + } + + public void emit(PTXAssembler asm) { + asm.emitString("or." + super.emit()); + } + } + + public static class Rem extends StandardFormat { + + public Rem(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + } + + public void emit(PTXAssembler asm) { + asm.emitString("rem." + super.emit()); + } + } + + public static class Shl extends StandardFormat { + + public Shl(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + setLogicInstruction(true); + } + + public void emit(PTXAssembler asm) { + asm.emitString("shl." + super.emit()); + } + } + + public static class Shr extends StandardFormat { + + public Shr(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + } + + public void emit(PTXAssembler asm) { + asm.emitString("shr." + super.emit()); + } + } + + public static class Sub extends StandardFormat { + + public Sub(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + } + + public void emit(PTXAssembler asm) { + asm.emitString("sub." + super.emit()); + } + } + + public static class Ushr extends StandardFormat { + + public Ushr(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + setKind(Kind.Illegal); // get around not having an Unsigned Kind + } + + public void emit(PTXAssembler asm) { + asm.emitString("shr." + super.emit()); + } + } + + public static class Xor extends StandardFormat { + + public Xor(Variable dst, Value src1, Value src2) { + super(dst, src1, src2); + setLogicInstruction(true); + } + + public void emit(PTXAssembler asm) { + asm.emitString("xor." + super.emit()); + } } // Checkstyle: stop method name check - public final void add_f32(Register d, Register a, Register b) { - emitString("add.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void add_f64(Register d, Register a, Register b) { - emitString("add.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - 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_f32(Register d, Register a, float f32) { - emitString("add.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f32 + ";" + ""); - } - - public final void add_f64(Register d, Register a, double f64) { - emitString("add.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f64 + ";" + ""); - } - - 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(String tgt, int pred) { + emitString((pred >= 0) ? "" : ("@%p" + pred + " ") + "bra" + " " + tgt + ";" + ""); } public final void bra_uni(String tgt) { emitString("bra.uni" + " " + tgt + ";" + ""); } - public final void cvt_s32_f32(Register d, Register a) { - emitString("cvt.s32.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); + public static class Cvt extends ConversionFormat { + + public Cvt(Variable dst, Variable src) { + super(dst, src); + } + + public void emit(PTXAssembler asm) { + asm.emitString("cvt." + super.emit()); + } + } + + public static class Mov extends SingleOperandFormat { + + public Mov(Variable dst, Value src) { + super(dst, src); + } + + /* + public Mov(Variable dst, AbstractAddress src) { + throw GraalInternalError.unimplemented("AbstractAddress Mov"); + } + */ + + public void emit(PTXAssembler asm) { + asm.emitString("mov." + super.emit()); + } + } + + public static class Neg extends SingleOperandFormat { + + public Neg(Variable dst, Variable src) { + super(dst, src); + } + + public void emit(PTXAssembler asm) { + asm.emitString("neg." + super.emit()); + } + } + + public static class Not extends SingleOperandFormat { + + public Not(Variable dst, Variable src) { + super(dst, src); + } + + public void emit(PTXAssembler asm) { + asm.emitString("not." + super.emit()); + } + } + + public static class Ld extends LoadStoreFormat { + + public Ld(PTXStateSpace space, Variable dst, Variable src1, Value src2) { + super(space, dst, src1, src2); + } + + public void emit(PTXAssembler asm) { + asm.emitString("ld." + super.emit(true)); + } } - public final void cvt_s64_f32(Register d, Register a) { - emitString("cvt.s64.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); - } + public static class St extends LoadStoreFormat { - public final void cvt_f64_f32(Register d, Register a) { - emitString("cvt.f64.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); - } + public St(PTXStateSpace space, Variable dst, Variable src1, Value src2) { + super(space, dst, src1, src2); + } - public final void cvt_f32_f64(Register d, Register a) { - emitString("cvt.f32.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); - } - - public final void cvt_s32_f64(Register d, Register a) { - emitString("cvt.s32.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); - } - - public final void cvt_s64_f64(Register d, Register a) { - emitString("cvt.s64.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); - } - - public final void cvt_f32_s32(Register d, Register a) { - emitString("cvt.f32.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); - } - - public final void cvt_f64_s32(Register d, Register a) { - emitString("cvt.f64.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); - } - - public final void cvt_s8_s32(Register d, Register a) { - emitString("cvt.s8.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); - } - - public final void cvt_b16_s32(Register d, Register a) { - emitString("cvt.b16.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); - } - - public final void cvt_s64_s32(Register d, Register a) { - emitString("cvt.s64.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); - } - - public final void cvt_s32_s64(Register d, Register a) { - emitString("cvt.s32.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); - } - - public final void div_f32(Register d, Register a, Register b) { - emitString("div.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void div_f64(Register d, Register a, Register b) { - emitString("div.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - 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_s32(Register d, int s32, Register b) { - emitString("div.s32" + " " + "%r" + d.encoding() + ", " + s32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void div_f32(Register d, float f32, Register b) { - emitString("div.f32" + " " + "%r" + d.encoding() + ", " + f32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void div_f64(Register d, double f64, Register b) { - emitString("div.f64" + " " + "%r" + d.encoding() + ", " + f64 + ", %r" + b.encoding() + ";" + ""); - } - - 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_f32(Register d, Register a, float f32) { - emitString("div.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f32 + ";" + ""); - } - - public final void div_f64(Register d, Register a, double f64) { - emitString("div.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f64 + ";" + ""); - } - - 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 void emit(PTXAssembler asm) { + asm.emitString("st." + super.emit(false)); + } } public final void exit() { emitString("exit;" + " " + ""); } - public final void ld_global_b8(Register d, Register a, long immOff) { - emitString("ld.global.b8" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } + public static class Param extends SingleOperandFormat { - public final void ld_global_b16(Register d, Register a, long immOff) { - emitString("ld.global.b16" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } + private boolean lastParameter; - public final void ld_global_b32(Register d, Register a, long immOff) { - emitString("ld.global.b32" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } + public Param(Variable d, boolean lastParam) { + super(d, null); + setLastParameter(lastParam); + } - public final void ld_global_b64(Register d, Register a, long immOff) { - emitString("ld.global.b64" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } + public void setLastParameter(boolean value) { + lastParameter = value; + } - public final void ld_global_u8(Register d, Register a, long immOff) { - emitString("ld.global.u8" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } + public String emitParameter(Variable v) { + return (" %r" + v.index); + } - public final void ld_global_u16(Register d, Register a, long immOff) { - emitString("ld.global.u16" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - public final void ld_global_u32(Register d, Register a, long immOff) { - emitString("ld.global.u32" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - public final void ld_global_u64(Register d, Register a, long immOff) { - emitString("ld.global.u64" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - public final void ld_global_s8(Register d, Register a, long immOff) { - emitString("ld.global.s8" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - public final void ld_global_s16(Register d, Register a, long immOff) { - emitString("ld.global.s16" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - public final void ld_global_s32(Register d, Register a, long immOff) { - emitString("ld.global.s32" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - public final void ld_global_s64(Register d, Register a, long immOff) { - emitString("ld.global.s64" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - public final void ld_global_f32(Register d, Register a, long immOff) { - emitString("ld.global.f32" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - public final void ld_global_f64(Register d, Register a, long immOff) { - emitString("ld.global.f64" + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - // Load from state space to destination register - public final void ld_from_state_space(String s, Register d, Register a, long immOff) { - emitString("ld" + s + " " + "%r" + d.encoding() + ", [" + a + " + " + immOff + "]" + ";" + ""); - } - - // Load return address from return parameter which is in .param state space - public final void ld_return_address(String s, Register d, Register a, long immOff) { - emitString("ld.param." + s + " " + "%r" + d.encoding() + ", [" + a + " + " + 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_u64(@SuppressWarnings("unused") Register d, @SuppressWarnings("unused") AbstractAddress 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_lo_f32(Register d, Register a, Register b) { - emitString("mul.lo.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void mul_lo_f64(Register d, Register a, Register b) { - emitString("mul.lo.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void mul_lo_s16(Register d, Register a, Register b) { - emitString("mul.lo.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void mul_lo_s32(Register d, Register a, Register b) { - emitString("mul.lo.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void mul_lo_s64(Register d, Register a, Register b) { - emitString("mul.lo.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void mul_lo_s16(Register d, Register a, short s16) { - emitString("mul.lo.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s16 + ";" + ""); - } - - public final void mul_lo_s32(Register d, Register a, int s32) { - emitString("mul.lo.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s32 + ";" + ""); - } - - public final void mul_lo_s64(Register d, Register a, long s64) { - emitString("mul.lo.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + s64 + ";" + ""); - } - - public final void mul_lo_f32(Register d, Register a, float f32) { - emitString("mul.lo.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f32 + ";" + ""); - } - - public final void mul_lo_f64(Register d, Register a, double f64) { - emitString("mul.lo.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f64 + ";" + ""); - } - - public final void mul_lo_u16(Register d, Register a, Register b) { - emitString("mul.lo.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void mul_lo_u32(Register d, Register a, Register b) { - emitString("mul.lo.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void mul_lo_u64(Register d, Register a, Register b) { - emitString("mul.lo.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void mul_lo_u16(Register d, Register a, short u16) { - emitString("mul.lo.u16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u16 + ";" + ""); - } - - public final void mul_lo_u32(Register d, Register a, int u32) { - emitString("mul.lo.u32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u32 + ";" + ""); - } - - public final void mul_lo_u64(Register d, Register a, long u64) { - emitString("mul.lo.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + ""); - } - - public final void neg_f32(Register d, Register a) { - emitString("neg.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); - } - - public final void neg_f64(Register d, Register a) { - emitString("neg.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); - } - - 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 not_s16(Register d, Register a) { - emitString("not.s16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); - } - - public final void not_s32(Register d, Register a) { - emitString("not.s32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); - } - - public final void not_s64(Register d, Register a) { - emitString("not.s64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ";" + ""); - } - - public final void or_b16(Register d, Register a, Register b) { - emitString("or.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void or_b32(Register d, Register a, Register b) { - emitString("or.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void or_b64(Register d, Register a, Register b) { - emitString("or.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void or_b16(Register d, Register a, short b16) { - emitString("or.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b16 + ";" + ""); - } - - public final void or_b32(Register d, Register a, int b32) { - emitString("or.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b32 + ";" + ""); - } - - public final void or_b64(Register d, Register a, long b64) { - emitString("or.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b64 + ";" + ""); - } - - public final void param_8_decl(Register d, boolean lastParam) { - emitString(".param" + " " + ".s8" + " " + d + (lastParam ? "" : ",")); - } - - public final void param_32_decl(Register d, boolean lastParam) { - emitString(".param" + " " + ".s32" + " " + d + (lastParam ? "" : ",")); - } - - public final void param_64_decl(Register d, boolean lastParam) { - emitString(".param" + " " + ".s64" + " " + d + (lastParam ? "" : ",")); + public void emit(PTXAssembler asm) { + asm.emitString(".param ." + typeForKind(dest.getKind()) + emitParameter(dest) + (lastParameter ? "" : ",")); + } } public final void popc_b32(Register d, Register a) { @@ -569,54 +593,6 @@ 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;" + " " + ""); } @@ -625,501 +601,168 @@ emitString("ret.uni;" + " " + ""); } - public final void setp_eq_f32(Register a, Register b) { - emitString("setp.eq.f32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); + public static class Setp { + + private ConditionOperator operator; + private Value first, second; + private Kind kind; + private int predicate; + + public Setp(Condition condition, Value first, Value second, int predicateRegisterNumber) { + setFirst(first); + setSecond(second); + setPredicate(predicateRegisterNumber); + setKind(); + setConditionOperator(operatorForConditon(condition)); + } + + public void setFirst(Value v) { + first = v; + } + + public void setSecond(Value v) { + second = v; + } + + public void setPredicate(int p) { + predicate = p; + } + + public void setConditionOperator(ConditionOperator co) { + operator = co; + } + + private ConditionOperator operatorForConditon(Condition condition) { + char typeChar = kind.getTypeChar(); + + switch (typeChar) { + case 'z': + case 'c': + case 'a': + // unsigned + switch (condition) { + case EQ: return ConditionOperator.U_EQ; + case NE: return ConditionOperator.U_NE; + case LT: return ConditionOperator.U_LO; + case LE: return ConditionOperator.U_LS; + case GT: return ConditionOperator.U_HI; + case GE: return ConditionOperator.U_HS; + default: + throw GraalInternalError.shouldNotReachHere(); + } + case 'b': + case 's': + case 'i': + case 'j': + // signed + switch (condition) { + case EQ: return ConditionOperator.S_EQ; + case NE: return ConditionOperator.S_NE; + case LT: return ConditionOperator.S_LT; + case LE: return ConditionOperator.S_LE; + case GT: return ConditionOperator.S_GT; + case GE: + case AE: + return ConditionOperator.S_GE; + default: + throw GraalInternalError.shouldNotReachHere(); + } + case 'f': + case 'd': + // floating point - do these need to accept NaN?? + switch (condition) { + case EQ: return ConditionOperator.F_EQ; + case NE: return ConditionOperator.F_NE; + case LT: return ConditionOperator.F_LT; + case LE: return ConditionOperator.F_LE; + case GT: return ConditionOperator.F_GT; + case GE: return ConditionOperator.F_GE; + default: + throw GraalInternalError.shouldNotReachHere(); + } + default: + throw GraalInternalError.shouldNotReachHere(); + } + } + + public void setKind() { + // assert isConstant(first) && isConstant(second) == false; + + if (isConstant(first)) { + kind = second.getKind(); + } else { + kind = first.getKind(); + } + } + + public String emitValue(Value v) { + assert v != null; + + if (isConstant(v)) { + return (", " + emitConstant(v)); + } else { + return (", " + emitVariable((Variable) v)); + } + } + + public String typeForKind(Kind k) { + switch (k.getTypeChar()) { + case 'z': + return "u8"; + case 'b': + return "s8"; + case 's': + return "s16"; + case 'c': + return "u16"; + case 'i': + return "s32"; + case 'f': + return "f32"; + case 'j': + return "s64"; + case 'd': + return "f64"; + case 'a': + return "u64"; + default: + throw GraalInternalError.shouldNotReachHere(); + } + } + + public String emitConstant(Value v) { + Constant constant = (Constant) v; + + switch (v.getKind().getTypeChar()) { + case 'i': + return (String.valueOf((int) constant.asLong())); + case 'f': + return (String.valueOf(constant.asFloat())); + case 'j': + return (String.valueOf(constant.asLong())); + case 'd': + return (String.valueOf(constant.asDouble())); + default: + throw GraalInternalError.shouldNotReachHere(); + } + } + + public String emitVariable(Variable v) { + return ("%r" + v.index); + } + + public void emit(PTXAssembler asm) { + asm.emitString("setp." + operator.getOperator() + "." + typeForKind(kind) + + " %p" + predicate + emitValue(first) + emitValue(second) + ";"); + } } - - public final void setp_ne_f32(Register a, Register b) { - emitString("setp.ne.f32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_lt_f32(Register a, Register b) { - emitString("setp.lt.f32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_le_f32(Register a, Register b) { - emitString("setp.le.f32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_gt_f32(Register a, Register b) { - emitString("setp.gt.f32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_ge_f32(Register a, Register b) { - emitString("setp.ge.f32" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_eq_f32(float f32, Register b) { - emitString("setp.eq.f32" + " " + "%p" + ", " + f32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_ne_f32(float f32, Register b) { - emitString("setp.ne.f32" + " " + "%p" + ", " + f32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_lt_f32(float f32, Register b) { - emitString("setp.lt.f32" + " " + "%p" + ", " + f32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_le_f32(float f32, Register b) { - emitString("setp.le.f32" + " " + "%p" + ", " + f32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_gt_f32(float f32, Register b) { - emitString("setp.gt.f32" + " " + "%p" + ", " + f32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_ge_f32(float f32, Register b) { - emitString("setp.ge.f32" + " " + "%p" + ", " + f32 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_eq_f64(double f64, Register b) { - emitString("setp.eq.f64" + " " + "%p" + ", " + f64 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_ne_f64(double f64, Register b) { - emitString("setp.ne.f64" + " " + "%p" + ", " + f64 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_lt_f64(double f64, Register b) { - emitString("setp.lt.f64" + " " + "%p" + ", " + f64 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_le_f64(double f64, Register b) { - emitString("setp.le.f64" + " " + "%p" + ", " + f64 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_gt_f64(double f64, Register b) { - emitString("setp.gt.f64" + " " + "%p" + ", " + f64 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_ge_f64(double f64, Register b) { - emitString("setp.ge.f64" + " " + "%p" + ", " + f64 + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_eq_s64(Register a, Register b) { - emitString("setp.eq.s64" + " " + "%p" + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void setp_eq_s64(long s64, Register b) { - emitString("setp.eq.s64" + " " + "%p" + ", " + s64 + ", %r" + b.encoding() + ";" + ""); - } - - 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() + ";" + ""); - } - - // Shift left - only types supported are .b16, .b32 and .b64 - public final void shl_b16(Register d, Register a, Register b) { - emitString("shl.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void shl_b32(Register d, Register a, Register b) { - emitString("shl.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void shl_b64(Register d, Register a, Register b) { - emitString("shl.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void shl_b16_const(Register d, Register a, int b) { - emitString("shl.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b + ";" + ""); - } - - public final void shl_b32_const(Register d, Register a, int b) { - emitString("shl.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b + ";" + ""); - } - - public final void shl_b64_const(Register d, Register a, int b) { - emitString("shl.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b + ";" + ""); - } - - // Shift Right instruction - 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, long u64) { - emitString("shr.u64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + u64 + ";" + ""); - } - - // Store in global state space - - public final void st_global_b8(Register a, long immOff, Register b) { - emitString("st.global.b8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_b16(Register a, long immOff, Register b) { - emitString("st.global.b16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_b32(Register a, long immOff, Register b) { - emitString("st.global.b32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_b64(Register a, long immOff, Register b) { - emitString("st.global.b64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_u8(Register a, long immOff, Register b) { - emitString("st.global.u8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_u16(Register a, long immOff, Register b) { - emitString("st.global.u16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_u32(Register a, long immOff, Register b) { - emitString("st.global.u32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_u64(Register a, long immOff, Register b) { - emitString("st.global.u64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_s8(Register a, long immOff, Register b) { - emitString("st.global.s8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_s16(Register a, long immOff, Register b) { - emitString("st.global.s16" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_s32(Register a, long immOff, Register b) { - emitString("st.global.s32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_s64(Register a, long immOff, Register b) { - emitString("st.global.s64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_f32(Register a, long immOff, Register b) { - emitString("st.global.f32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_f64(Register a, long immOff, Register b) { - emitString("st.global.f64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - // Store return value - public final void st_global_return_value_s8(Register a, long immOff, Register b) { - emitString("st.global.s8" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_return_value_s32(Register a, long immOff, Register b) { - emitString("st.global.s32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_return_value_s64(Register a, long immOff, Register b) { - emitString("st.global.s64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_return_value_f32(Register a, long immOff, Register b) { - emitString("st.global.f32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_return_value_f64(Register a, long immOff, Register b) { - emitString("st.global.f64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_return_value_u32(Register a, long immOff, Register b) { - emitString("st.global.u32" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - public final void st_global_return_value_u64(Register a, long immOff, Register b) { - emitString("st.global.u64" + " " + "[%r" + a.encoding() + " + " + immOff + "], %r" + b.encoding() + ";" + ""); - } - - // Subtract instruction - - public final void sub_f32(Register d, Register a, Register b) { - emitString("sub.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void sub_f64(Register d, Register a, Register b) { - emitString("sub.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %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, int s32) { - emitString("sub.s64" + " " + "%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_f32(Register d, Register a, float f32) { - emitString("sub.f32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f32 + ";" + ""); - } - - public final void sub_f64(Register d, Register a, double f64) { - emitString("sub.f64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + f64 + ";" + ""); - } - - 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_f32(Register d, float f32, Register b) { - emitString("sub.f32" + " " + "%r" + d.encoding() + ", %r" + b.encoding() + ", " + f32 + ";" + ""); - } - - public final void sub_f64(Register d, double f64, Register b) { - emitString("sub.f64" + " " + "%r" + d.encoding() + ", %r" + b.encoding() + ", " + f64 + ";" + ""); - } - - 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() + ";" + ""); - } - - public final void xor_b16(Register d, Register a, Register b) { - emitString("xor.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void xor_b32(Register d, Register a, Register b) { - emitString("xor.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void xor_b64(Register d, Register a, Register b) { - emitString("xor.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", %r" + b.encoding() + ";" + ""); - } - - public final void xor_b16(Register d, Register a, short b16) { - emitString("xor.b16" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b16 + ";" + ""); - } - - public final void xor_b32(Register d, Register a, int b32) { - emitString("xor.b32" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b32 + ";" + ""); - } - - public final void xor_b64(Register d, Register a, long b64) { - emitString("xor.b64" + " " + "%r" + d.encoding() + ", %r" + a.encoding() + ", " + b64 + ";" + ""); - } - @Override public PTXAddress makeAddress(Register base, int displacement) { - return new PTXAddress(base, displacement); + throw GraalInternalError.shouldNotReachHere(); } @Override public PTXAddress getPlaceholder() { - // TODO Auto-generated method stub return null; } }
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.asm.ptx/src/com/oracle/graal/asm/ptx/PTXStateSpace.java Wed Oct 02 13:26:31 2013 +0200 @@ -0,0 +1,49 @@ +/* + * 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; + +/** + * Represents the various PTX state spaces. + */ +public enum PTXStateSpace { + + Parameter("param"), + + Shared("shared"), + + Local("local"), + + Global("global"), + + Const("const"); + + private final String stateName; + + private PTXStateSpace(String name) { + this.stateName = name; + } + + public String getStateName() { + return stateName; + } +}
--- a/graal/com.oracle.graal.asm.sparc/src/com/oracle/graal/asm/sparc/SPARCAddress.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.asm.sparc/src/com/oracle/graal/asm/sparc/SPARCAddress.java Wed Oct 02 13:26:31 2013 +0200 @@ -121,7 +121,7 @@ throw GraalInternalError.shouldNotReachHere("address has index register"); } // TODO Should we also hide the register save area size here? - if (getBase() == sp || getBase() == fp) { + if (getBase().equals(sp) || getBase().equals(fp)) { return displacement + STACK_BIAS; } return displacement;
--- a/graal/com.oracle.graal.compiler.amd64/src/com/oracle/graal/compiler/amd64/AMD64LIRGenerator.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.amd64/src/com/oracle/graal/compiler/amd64/AMD64LIRGenerator.java Wed Oct 02 13:26:31 2013 +0200 @@ -890,6 +890,6 @@ @Override public void visitInfopointNode(InfopointNode i) { - append(new InfopointOp(stateFor(i.stateAfter(), DeoptimizationReason.None), i.reason)); + append(new InfopointOp(stateFor(i.getState()), i.reason)); } }
--- a/graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/StaticDoubleSpillTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/StaticDoubleSpillTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -25,7 +25,8 @@ import java.util.*; import com.oracle.graal.compiler.hsail.test.infra.GraalKernelTester; -import org.junit.Test; + +import org.junit.*; /** * Tests the spilling of double variables into memory. @@ -118,6 +119,7 @@ } // Marked to only run on hardware until simulator spill bug is fixed. + @Ignore @Test public void test() { testGeneratedHsail();
--- a/graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/StaticIntSpillTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/StaticIntSpillTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -25,7 +25,8 @@ import java.util.*; import com.oracle.graal.compiler.hsail.test.infra.GraalKernelTester; -import org.junit.Test; + +import org.junit.*; /** * Tests the spilling of integers into memory. @@ -87,6 +88,7 @@ } // Marked to only run on hardware until simulator spill bug is fixed. + @Ignore @Test public void test() { testGeneratedHsail();
--- a/graal/com.oracle.graal.compiler.hsail/src/com/oracle/graal/compiler/hsail/HSAILBackend.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.hsail/src/com/oracle/graal/compiler/hsail/HSAILBackend.java Wed Oct 02 13:26:31 2013 +0200 @@ -60,8 +60,16 @@ } @Override + public boolean shouldAllocateRegisters() { + return true; + } + + /** + * Use the HSAIL register set when the compilation target is HSAIL. + */ + @Override public FrameMap newFrameMap() { - return new HSAILFrameMap(runtime(), target, runtime().lookupRegisterConfig()); + return new HSAILFrameMap(runtime(), target, new HSAILRegisterConfig()); } @Override
--- a/graal/com.oracle.graal.compiler.hsail/src/com/oracle/graal/compiler/hsail/HSAILCompilationResult.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.hsail/src/com/oracle/graal/compiler/hsail/HSAILCompilationResult.java Wed Oct 02 13:26:31 2013 +0200 @@ -23,33 +23,25 @@ package com.oracle.graal.compiler.hsail; -import static com.oracle.graal.api.code.CodeUtil.*; - +import java.lang.reflect.*; import java.util.logging.*; import com.oracle.graal.api.code.*; -import com.oracle.graal.api.code.CallingConvention.*; -import com.oracle.graal.api.runtime.Graal; -import com.oracle.graal.compiler.GraalCompiler; -import com.oracle.graal.java.GraphBuilderConfiguration; -import com.oracle.graal.java.GraphBuilderPhase; -import com.oracle.graal.nodes.StructuredGraph; -import com.oracle.graal.nodes.spi.GraalCodeCacheProvider; -import com.oracle.graal.phases.OptimisticOptimizations; -import com.oracle.graal.phases.PhasePlan; +import com.oracle.graal.api.code.CallingConvention.Type; +import com.oracle.graal.api.meta.*; +import com.oracle.graal.api.runtime.*; +import com.oracle.graal.compiler.*; +import com.oracle.graal.compiler.target.*; +import com.oracle.graal.debug.*; +import com.oracle.graal.graph.*; +import com.oracle.graal.hsail.*; +import com.oracle.graal.java.*; +import com.oracle.graal.nodes.*; +import com.oracle.graal.nodes.spi.*; +import com.oracle.graal.nodes.type.*; +import com.oracle.graal.phases.*; import com.oracle.graal.phases.PhasePlan.PhasePosition; -import com.oracle.graal.api.meta.*; -import com.oracle.graal.nodes.type.*; -import com.oracle.graal.graph.GraalInternalError; -import com.oracle.graal.debug.*; -import com.oracle.graal.nodes.*; -import com.oracle.graal.phases.*; import com.oracle.graal.phases.tiers.*; -import com.oracle.graal.nodes.spi.Replacements; -import com.oracle.graal.compiler.target.Backend; -import com.oracle.graal.hsail.*; - -import java.lang.reflect.Method; /** * Class that represents a HSAIL compilation result. Includes the compiled HSAIL code. @@ -96,6 +88,33 @@ return getHSAILCompilationResult(graph); } + /** + * HSAIL doesn't have a calling convention as such. Function arguments are actually passed in + * memory but then loaded into registers in the function body. This routine makes sure that + * arguments to a kernel or function are loaded (by the kernel or function body) into registers + * of the appropriate sizes. For example, int and float parameters should appear in S registers, + * whereas double and long parameters should appear in d registers. + */ + public static CallingConvention getHSAILCallingConvention(CallingConvention.Type type, TargetDescription target, ResolvedJavaMethod method, boolean stackOnly) { + Signature sig = method.getSignature(); + JavaType retType = sig.getReturnType(null); + int sigCount = sig.getParameterCount(false); + JavaType[] argTypes; + int argIndex = 0; + if (!Modifier.isStatic(method.getModifiers())) { + argTypes = new JavaType[sigCount + 1]; + argTypes[argIndex++] = method.getDeclaringClass(); + } else { + argTypes = new JavaType[sigCount]; + } + for (int i = 0; i < sigCount; i++) { + argTypes[argIndex++] = sig.getParameterType(i, null); + } + + RegisterConfig registerConfig = new HSAILRegisterConfig(); + return registerConfig.getCallingConvention(type, retType, argTypes, target, stackOnly); + } + public static HSAILCompilationResult getHSAILCompilationResult(StructuredGraph graph) { Debug.dump(graph, "Graph"); TargetDescription target = new TargetDescription(new HSAIL(), true, 8, 0, true); @@ -105,7 +124,7 @@ phasePlan.addPhase(PhasePosition.AFTER_PARSING, graphBuilderPhase); phasePlan.addPhase(PhasePosition.AFTER_PARSING, new HSAILPhase()); new HSAILPhase().apply(graph); - CallingConvention cc = getCallingConvention(runtime, Type.JavaCallee, graph.method(), false); + CallingConvention cc = getHSAILCallingConvention(Type.JavaCallee, target, graph.method(), false); try { CompilationResult compResult = GraalCompiler.compileGraph(graph, cc, graph.method(), runtime, replacements, hsailBackend, target, null, phasePlan, OptimisticOptimizations.NONE, new SpeculationLog(), suitesProvider.getDefaultSuites(), new CompilationResult());
--- a/graal/com.oracle.graal.compiler.hsail/src/com/oracle/graal/compiler/hsail/HSAILLIRGenerator.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.hsail/src/com/oracle/graal/compiler/hsail/HSAILLIRGenerator.java Wed Oct 02 13:26:31 2013 +0200 @@ -143,12 +143,12 @@ } else { baseRegister = load(base); } - } else if (base == Value.ILLEGAL) { + } else if (base.equals(Value.ILLEGAL)) { baseRegister = Value.ILLEGAL; } else { baseRegister = asAllocatable(base); } - if (index != Value.ILLEGAL) { + if (!index.equals(Value.ILLEGAL)) { if (isConstant(index)) { finalDisp += asConstant(index).asLong() * scale; } else { @@ -160,7 +160,7 @@ } else { indexRegister = convertedIndex; } - if (baseRegister == Value.ILLEGAL) { + if (baseRegister.equals(Value.ILLEGAL)) { baseRegister = asAllocatable(indexRegister); } else { baseRegister = emitAdd(baseRegister, indexRegister); @@ -579,7 +579,7 @@ } @Override - public void emitDeoptimize(DeoptimizationAction action, DeoptimizingNode deopting) { + public void emitDeoptimize(Value actionAndReason, DeoptimizingNode deopting) { append(new ReturnOp(Value.ILLEGAL)); }
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ArrayPTXTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -0,0 +1,91 @@ +/* + * 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.test; + +import static com.oracle.graal.lir.ptx.ThreadDimension.*; + +import com.oracle.graal.lir.ptx.ParallelOver; +import com.oracle.graal.lir.ptx.Warp; + +import java.lang.reflect.Method; +import java.util.Arrays; +import org.junit.Test; + +public class ArrayPTXTest extends PTXTestBase { + + @Test + public void testArray() { + int[] array1 = { + 1, 2, 3, 4, 5, 6, 7, 8, 9 + }; + int[] array2 = { + 1, 2, 3, 4, 5, 6, 7, 8, 9 + }; + int[] array3 = { + 1, 2, 3, 4, 5, 6, 7, 8, 9 + }; + + invoke(compile("testStoreArray1I"), array1, 2); + printReport("testStoreArray1I: " + Arrays.toString(array1)); + + invoke(compile("testStoreArrayWarp0"), array2, 2); + printReport("testStoreArrayWarp0: " + Arrays.toString(array2)); + + invoke(compile("testStoreArrayWarp1I"), array3, 2); + printReport("testStoreArrayWarp1I: " + Arrays.toString(array3)); + + } + + public static void testStoreArray1I(int[] array, int i) { + array[i] = 42; + } + + public static void testStoreArrayWarp0(int[] array, + @Warp(dimension = X) int i) { + array[i] = 42; + } + + public static void testStoreArrayWarp1I(@ParallelOver(dimension = X) int[] array, + @Warp(dimension = X) int i) { + array[i] = 42; + } + + + public static void printReport(String message) { + // CheckStyle: stop system..print check + System.out.println(message); + // CheckStyle: resume system..print check + } + + public static void main(String[] args) { + ArrayPTXTest test = new ArrayPTXTest(); + for (Method m : ArrayPTXTest.class.getMethods()) { + String name = m.getName(); + if (m.getAnnotation(Test.class) == null && name.startsWith("test")) { + // CheckStyle: stop system..print check + System.out.println(name + ": \n" + new String(test.compile(name).getTargetCode())); + // CheckStyle: resume system..print check + } + } + } +}
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ArrayTest.java Fri Sep 06 21:37:50 2013 +0200 +++ /dev/null Thu Jan 01 00:00:00 1970 +0000 @@ -1,117 +0,0 @@ -/* - * 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.test; - -import java.lang.reflect.Method; - -import org.junit.*; - -public class ArrayTest extends PTXTestBase { - - @Ignore - @Test - public void testArray() { - compile("testArray1I"); - compile("testArray1J"); - compile("testArray1B"); - compile("testArray1S"); - compile("testArray1C"); - compile("testArray1F"); - compile("testArray1D"); - compile("testArray1L"); - compile("testStoreArray1I"); - compile("testStoreArray1J"); - compile("testStoreArray1B"); - compile("testStoreArray1S"); - compile("testStoreArray1F"); - compile("testStoreArray1D"); - } - - public static int testArray1I(int[] array, int i) { - return array[i]; - } - - public static long testArray1J(long[] array, int i) { - return array[i]; - } - - public static byte testArray1B(byte[] array, int i) { - return array[i]; - } - - public static short testArray1S(short[] array, int i) { - return array[i]; - } - - public static char testArray1C(char[] array, int i) { - return array[i]; - } - - public static float testArray1F(float[] array, int i) { - return array[i]; - } - - public static double testArray1D(double[] array, int i) { - return array[i]; - } - - public static Object testArray1L(Object[] array, int i) { - return array[i]; - } - - public static void testStoreArray1I(int[] array, int i, int val) { - array[i] = val; - } - - public static void testStoreArray1B(byte[] array, int i, byte val) { - array[i] = val; - } - - public static void testStoreArray1S(short[] array, int i, short val) { - array[i] = val; - } - - public static void testStoreArray1J(long[] array, int i, long val) { - array[i] = val; - } - - public static void testStoreArray1F(float[] array, int i, float val) { - array[i] = val; - } - - public static void testStoreArray1D(double[] array, int i, double val) { - array[i] = val; - } - - public static void main(String[] args) { - ArrayTest test = new ArrayTest(); - for (Method m : ArrayTest.class.getMethods()) { - String name = m.getName(); - if (m.getAnnotation(Test.class) == null && name.startsWith("test")) { - // CheckStyle: stop system..print check - System.out.println(name + ": \n" + new String(test.compile(name).getTargetCode())); - // CheckStyle: resume system..print check - } - } - } -}
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/BasicPTXTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/BasicPTXTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -24,11 +24,13 @@ import java.lang.reflect.Method; +import org.junit.Ignore; import org.junit.Test; /** * Test class for small Java methods compiled to PTX kernels. */ +@Ignore public class BasicPTXTest extends PTXTestBase { @Test
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ControlPTXTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -0,0 +1,118 @@ +/* + * 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.test; + +import org.junit.*; + +import java.lang.reflect.Method; + +public class ControlPTXTest extends PTXTestBase { + + @Ignore + @Test + public void testControl() { + compile("testLoop"); + // compile("testSwitch1I"); + // compile("testStatic"); + // compile("testCall"); + // compile("testLookupSwitch1I"); + } + + public static int testLoop(int n) { + int sum = 0; + + for (int i = 0; i < n; i++) { + sum++; + } + return sum; + } + + public static int testSwitch1I(int a) { + switch (a) { + case 1: + return 2; + case 2: + return 3; + default: + return 4; + } + } + + public static int testLookupSwitch1I(int a) { + switch (a) { + case 0: + return 1; + case 1: + return 2; + case 2: + return 3; + case 3: + return 1; + case 4: + return 2; + case 5: + return 3; + case 6: + return 1; + case 7: + return 2; + case 8: + return 3; + case 9: + return 1; + case 10: + return 2; + case 11: + return 3; + default: + return -1; + } + } + + @SuppressWarnings("unused") private static Object foo = null; + + public static boolean testStatic(Object o) { + foo = o; + return true; + } + + private static int method(int a, int b) { + return a + b; + } + + public static int testCall(@SuppressWarnings("unused") Object o, int a, int b) { + return method(a, b); + } + + public static void main(String[] args) { + ControlPTXTest test = new ControlPTXTest(); + for (Method m : ControlPTXTest.class.getMethods()) { + String name = m.getName(); + if (m.getAnnotation(Test.class) == null && name.startsWith("test")) { + // CheckStyle: stop system..print check + System.out.println(name + ": \n" + new String(test.compile(name).getTargetCode())); + // CheckStyle: resume system..print check + } + } + } +}
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/ControlTest.java Fri Sep 06 21:37:50 2013 +0200 +++ /dev/null Thu Jan 01 00:00:00 1970 +0000 @@ -1,96 +0,0 @@ -/* - * 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.test; - -import org.junit.*; - -import java.lang.reflect.Method; - -public class ControlTest extends PTXTestBase { - - @Ignore - @Test - public void testControl() { - compile("testSwitch1I"); - compile("testStatic"); - compile("testCall"); - compile("testLookupSwitch1I"); - } - - public static int testSwitch1I(int a) { - switch (a) { - case 1: - return 2; - case 2: - return 3; - default: - return 4; - } - } - - public static int testLookupSwitch1I(int a) { - switch (a) { - case 0: return 1; - case 1: return 2; - case 2: return 3; - case 3: return 1; - case 4: return 2; - case 5: return 3; - case 6: return 1; - case 7: return 2; - case 8: return 3; - case 9: return 1; - case 10: return 2; - case 11: return 3; - default: return -1; - } - } - - @SuppressWarnings("unused") - private static Object foo = null; - - public static boolean testStatic(Object o) { - foo = o; - return true; - } - - private static int method(int a, int b) { - return a + b; - } - - public static int testCall(@SuppressWarnings("unused") Object o, int a, int b) { - return method(a, b); - } - - public static void main(String[] args) { - ControlTest test = new ControlTest(); - for (Method m : ControlTest.class.getMethods()) { - String name = m.getName(); - if (m.getAnnotation(Test.class) == null && name.startsWith("test")) { - // CheckStyle: stop system..print check - System.out.println(name + ": \n" + new String(test.compile(name).getTargetCode())); - // CheckStyle: resume system..print check - } - } - } -}
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/FloatPTXTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/FloatPTXTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -29,14 +29,18 @@ import com.oracle.graal.api.code.CompilationResult; /* PTX ISA 3.1 - 8.7.3 Floating-Point Instructions */ +@Ignore public class FloatPTXTest extends PTXTestBase { + @Ignore @Test public void testAdd() { CompilationResult r = compile("testAdd2F"); if (r.getTargetCode() == null) { printReport("Compilation of testAdd2F FAILED"); } + + /* r = compile("testAdd2D"); if (r.getTargetCode() == null) { printReport("Compilation of testAdd2D FAILED"); @@ -58,6 +62,7 @@ if (r.getTargetCode() == null) { printReport("Compilation of testConstD FAILED"); } + */ } public static float testAdd2F(float a, float b) { @@ -84,6 +89,7 @@ return 32.0 + a; } + @Ignore @Test public void testSub() { CompilationResult r = compile("testSub2F"); @@ -141,6 +147,7 @@ return 32.0 - a; } + @Ignore @Test public void testMul() { CompilationResult r = compile("testMul2F"); @@ -198,6 +205,7 @@ return 32.0 * a; } + @Ignore @Test public void testDiv() { CompilationResult r = compile("testDiv2F"); @@ -255,6 +263,7 @@ return 32.0 / a; } + @Ignore @Test public void testNeg() { CompilationResult r = compile("testNeg2F"); @@ -276,6 +285,7 @@ return -a; } + @Ignore @Test public void testRem() { // need linkage to PTX remainder() @@ -360,9 +370,7 @@ FloatPTXTest test = new FloatPTXTest(); for (Method m : FloatPTXTest.class.getMethods()) { String name = m.getName(); - if (m.getAnnotation(Test.class) == null && - name.startsWith("test") && - name.startsWith("testRem") == false) { + if (m.getAnnotation(Test.class) == null && name.startsWith("test") && name.startsWith("testRem") == false) { // CheckStyle: stop system..print check System.out.println(name + ": \n" + new String(test.compile(name).getTargetCode())); // CheckStyle: resume system..print check
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/IntegerPTXTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/IntegerPTXTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -26,11 +26,19 @@ import java.lang.reflect.Method; - public class IntegerPTXTest extends PTXTestBase { @Test public void testAdd() { + /* + Integer r4 = (Integer) invoke(compile("testAdd2B"), (byte) 6, (byte) 4); + if (r4 == null) { + printReport("testAdd2B FAILED"); + } else if (r4.intValue() == testAdd2B((byte) 6, (byte) 4)) { + printReport("testAdd2B PASSED"); + } else { + printReport("testAdd2B FAILED"); + } */ Integer r4 = (Integer) invoke(compile("testAdd2I"), 18, 24); if (r4 == null) { @@ -41,16 +49,14 @@ printReport("testAdd2I FAILED"); } - Long r2 = (Long) invoke(compile("testAdd2L"), (long) 12, (long) 6); + /* Long r2 = (Long) invoke(compile("testAdd2L"), (long) 12, (long) 6); if (r2 == null) { printReport("testAdd2L FAILED"); } else if (r2.longValue() == testAdd2L(12, 6)) { printReport("testAdd2L PASSED"); } else { printReport("testAdd2L FAILED"); - } - - //invoke(compile("testAdd2B"), (byte) 6, (byte) 4); + } r4 = (Integer) invoke(compile("testAddIConst"), 5); if (r4 == null) { @@ -68,8 +74,7 @@ printReport("testAddConstI PASSED"); } else { printReport("testAddConstI FAILED"); - } - + } */ } public static int testAdd2I(int a, int b) { @@ -92,6 +97,7 @@ return 32 + a; } + @Ignore @Test public void testSub() { @@ -149,6 +155,7 @@ return 32 - a; } + @Ignore @Test public void testMul() { @@ -260,6 +267,7 @@ return 32 / a; } + @Ignore @Test public void testRem() { Integer r1 = (Integer) invoke(compile("testRem2I"), 8, 4); @@ -288,6 +296,7 @@ public static long testRem2L(long a, long b) { return a % b; } + @Ignore @Test public void testIntConversion() {
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/LogicPTXTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/LogicPTXTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -24,10 +24,11 @@ import java.lang.reflect.Method; +import org.junit.Ignore; import org.junit.Test; - /* PTX ISA 3.1 - 8.7.5 Logic and Shift Instructions */ +@Ignore public class LogicPTXTest extends PTXTestBase { @Test @@ -105,7 +106,7 @@ compile("testShiftRight2I"); compile("testShiftRight2L"); compile("testUnsignedShiftRight2I"); - compile("testUnsignedShiftRight2L"); + // compile("testUnsignedShiftRight2L"); } public static int testShiftRight2I(int a, int b) {
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/PTXTestBase.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/PTXTestBase.java Wed Oct 02 13:26:31 2013 +0200 @@ -30,12 +30,16 @@ import com.oracle.graal.api.runtime.Graal; import com.oracle.graal.compiler.GraalCompiler; import com.oracle.graal.compiler.ptx.PTXBackend; +import com.oracle.graal.compiler.ptx.PTXTargetMethodAssembler; import com.oracle.graal.compiler.test.GraalCompilerTest; import com.oracle.graal.debug.Debug; +import com.oracle.graal.hotspot.meta.HotSpotNmethod; import com.oracle.graal.hotspot.meta.HotSpotRuntime; import com.oracle.graal.hotspot.meta.HotSpotResolvedJavaMethod; +import com.oracle.graal.hotspot.ptx.PTXHotSpotRuntime; import com.oracle.graal.java.GraphBuilderConfiguration; import com.oracle.graal.java.GraphBuilderPhase; +import com.oracle.graal.lir.ptx.ParallelOver; import com.oracle.graal.nodes.StructuredGraph; import com.oracle.graal.nodes.spi.GraalCodeCacheProvider; import com.oracle.graal.phases.OptimisticOptimizations; @@ -43,6 +47,8 @@ import com.oracle.graal.phases.PhasePlan.PhasePosition; import com.oracle.graal.phases.tiers.*; import com.oracle.graal.ptx.PTX; + +import java.lang.annotation.Annotation; import java.lang.reflect.Modifier; public abstract class PTXTestBase extends GraalCompilerTest { @@ -50,32 +56,36 @@ private StructuredGraph sg; protected CompilationResult compile(String test) { - StructuredGraph graph = parse(test); - sg = graph; - Debug.dump(graph, "Graph"); - TargetDescription target = new TargetDescription(new PTX(), true, 1, 0, true); - PTXBackend ptxBackend = new PTXBackend(Graal.getRequiredCapability(GraalCodeCacheProvider.class), target); - PhasePlan phasePlan = new PhasePlan(); - GraphBuilderPhase graphBuilderPhase = new GraphBuilderPhase(runtime, GraphBuilderConfiguration.getDefault(), OptimisticOptimizations.NONE); - phasePlan.addPhase(PhasePosition.AFTER_PARSING, graphBuilderPhase); - phasePlan.addPhase(PhasePosition.AFTER_PARSING, new PTXPhase()); - new PTXPhase().apply(graph); - CallingConvention cc = getCallingConvention(runtime, Type.JavaCallee, graph.method(), false); - /* - * Use Suites.createDefaultSuites() instead of GraalCompilerTest.suites. The - * GraalCompilerTest.suites variable contains the Suites for the HotSpotRuntime. This code - * will not run on hotspot, so it should use the plain Graal default suites, without hotspot - * specific phases. - * - * Ultimately we might want to have both the kernel and the code natively compiled for GPU fallback to CPU in cases - * of ECC failure on kernel invocation. - */ - CompilationResult result = GraalCompiler.compileGraph(graph, cc, graph.method(), runtime, - graalRuntime().getReplacements(), ptxBackend, target, - null, phasePlan, - OptimisticOptimizations.NONE, new SpeculationLog(), - Suites.createDefaultSuites(), new ExternalCompilationResult()); - return result; + if (runtime instanceof PTXHotSpotRuntime) { + StructuredGraph graph = parse(test); + sg = graph; + Debug.dump(graph, "Graph"); + TargetDescription target = new TargetDescription(new PTX(), true, 1, 0, true); + PTXBackend ptxBackend = new PTXBackend(Graal.getRequiredCapability(GraalCodeCacheProvider.class), target); + PhasePlan phasePlan = new PhasePlan(); + GraphBuilderPhase graphBuilderPhase = new GraphBuilderPhase(runtime, GraphBuilderConfiguration.getDefault(), OptimisticOptimizations.NONE); + phasePlan.addPhase(PhasePosition.AFTER_PARSING, graphBuilderPhase); + phasePlan.addPhase(PhasePosition.AFTER_PARSING, new PTXPhase()); + new PTXPhase().apply(graph); + CallingConvention cc = getCallingConvention(runtime, Type.JavaCallee, graph.method(), false); + /* + * Use Suites.createDefaultSuites() instead of GraalCompilerTest.suites. The + * GraalCompilerTest.suites variable contains the Suites for the HotSpotRuntime. This code + * will not run on hotspot, so it should use the plain Graal default suites, without hotspot + * specific phases. + * + * Ultimately we might want to have both the kernel and the code natively compiled for GPU fallback to CPU in cases + * of ECC failure on kernel invocation. + */ + CompilationResult result = GraalCompiler.compileGraph(graph, cc, graph.method(), runtime, + graalRuntime().getReplacements(), ptxBackend, target, + null, phasePlan, + OptimisticOptimizations.NONE, new SpeculationLog(), + Suites.createDefaultSuites(), new ExternalCompilationResult()); + return result; + } else { + return null; + } } protected StructuredGraph getStructuredGraph() { @@ -83,6 +93,9 @@ } protected Object invoke(CompilationResult result, Object... args) { + if (result == null) { + return null; + } try { if (((ExternalCompilationResult) result).getEntryPoint() == 0) { Debug.dump(result, "[CUDA] *** Null entry point - Not launching kernel"); @@ -94,8 +107,47 @@ boolean isStatic = Modifier.isStatic(compiledMethod.getModifiers()); Object[] executeArgs = argsWithReceiver((isStatic ? null : this), args); HotSpotRuntime hsr = (HotSpotRuntime) runtime; - InstalledCode installedCode = hsr.addExternalMethod(sg.method(), result, sg); - Object r = installedCode.executeVarargs(executeArgs); + InstalledCode installedCode = hsr.addExternalMethod(compiledMethod, result, sg); + Annotation[][] params = compiledMethod.getParameterAnnotations(); + + int dimensionX = 1; + int dimensionY = 1; + int dimensionZ = 1; + + for (int p = 0; p < params.length; p++) { + Annotation[] annos = params[p]; + if (annos != null) { + for (int a = 0; a < annos.length; a++) { + Annotation aa = annos[a]; + if (args[p] instanceof int[] && aa.annotationType().equals(ParallelOver.class)) { + int[] iarray = (int[]) args[p]; + ParallelOver threadBlockDimension = (ParallelOver) aa; + switch (threadBlockDimension.dimension()) { + case X: + dimensionX = iarray.length; + break; + case Y: + dimensionY = iarray.length; + break; + case Z: + dimensionZ = iarray.length; + break; + } + } + } + } + } + Object r; + if (dimensionX != 1 || dimensionY != 1 || dimensionZ != 1) { + /* + * for now assert that the warp array block is no larger than the number of physical gpu cores. + */ + assert dimensionX * dimensionY * dimensionZ < PTXTargetMethodAssembler.getAvailableProcessors(); + + r = ((HotSpotNmethod) installedCode).executeParallel(dimensionX, dimensionY, dimensionZ, executeArgs); + } else { + r = installedCode.executeVarargs(executeArgs); + } return r; } catch (Throwable th) { th.printStackTrace();
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java Wed Oct 02 13:26:31 2013 +0200 @@ -22,7 +22,7 @@ */ package com.oracle.graal.compiler.ptx; -import static com.oracle.graal.api.code.ValueUtil.*; +import static com.oracle.graal.lir.LIRValueUtil.*; import java.util.*; @@ -40,6 +40,7 @@ import com.oracle.graal.lir.LIRInstruction.OperandFlag; import com.oracle.graal.lir.LIRInstruction.OperandMode; import com.oracle.graal.lir.LIRInstruction.ValueProcedure; +import com.oracle.graal.lir.StandardOp.LabelOp; import com.oracle.graal.graph.GraalInternalError; /** @@ -52,6 +53,11 @@ } @Override + public boolean shouldAllocateRegisters() { + return false; + } + + @Override public FrameMap newFrameMap() { return new PTXFrameMap(runtime(), target, runtime().lookupRegisterConfig()); } @@ -80,7 +86,7 @@ @Override public TargetMethodAssembler newAssembler(LIRGenerator lirGen, CompilationResult compilationResult) { - // Omit the frame if the method: + // Omit the frame of 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 @@ -89,14 +95,13 @@ AbstractAssembler masm = createAssembler(frameMap); HotSpotFrameContext frameContext = new HotSpotFrameContext(); TargetMethodAssembler tasm = new PTXTargetMethodAssembler(target, runtime(), frameMap, masm, frameContext, compilationResult); - tasm.setFrameSize(frameMap.frameSize()); + tasm.setFrameSize(0); return tasm; } - private static void emitKernelEntry(TargetMethodAssembler tasm, LIRGenerator lirGen, - ResolvedJavaMethod codeCacheOwner) { + private static void emitKernelEntry(TargetMethodAssembler tasm, LIRGenerator lirGen, ResolvedJavaMethod codeCacheOwner) { // Emit PTX kernel entry text based on PTXParameterOp - // instructions in the start block. Remove the instructions + // instructions in the start block. Remove the instructions // once kernel entry text and directives are emitted to // facilitate seemless PTX code generation subsequently. assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method"; @@ -132,8 +137,6 @@ // Start emiting body of the PTX kernel. codeBuffer.emitString0(") {"); codeBuffer.emitString(""); - - codeBuffer.emitString(".reg .u64" + " %rax;"); } // Emit .reg space declarations @@ -144,23 +147,53 @@ final SortedSet<Integer> signed32 = new TreeSet<>(); final SortedSet<Integer> signed64 = new TreeSet<>(); + final SortedSet<Integer> unsigned64 = new TreeSet<>(); + final SortedSet<Integer> float32 = new TreeSet<>(); + final SortedSet<Integer> float64 = new TreeSet<>(); ValueProcedure trackRegisterKind = new ValueProcedure() { @Override public Value doValue(Value value, OperandMode mode, EnumSet<OperandFlag> flags) { - if (isRegister(value)) { - RegisterValue regVal = (RegisterValue) value; + if (isVariable(value)) { + Variable regVal = (Variable) value; Kind regKind = regVal.getKind(); switch (regKind) { - case Int: - signed32.add(regVal.getRegister().encoding()); - break; - case Long: - signed64.add(regVal.getRegister().encoding()); - break; - default : - throw GraalInternalError.shouldNotReachHere("unhandled register type " + value.toString()); + case Int: + // If the register was used as a wider signed type + // do not add it here + if (!signed64.contains(regVal.index)) { + signed32.add(regVal.index); + } + break; + case Long: + // If the register was used as a narrower signed type + // remove it from there and add it to wider type. + if (signed32.contains(regVal.index)) { + signed32.remove(regVal.index); + } + signed64.add(regVal.index); + break; + case Float: + // If the register was used as a wider signed type + // do not add it here + if (!float64.contains(regVal.index)) { + float32.add(regVal.index); + } + break; + case Double: + // If the register was used as a narrower signed type + // remove it from there and add it to wider type. + if (float32.contains(regVal.index)) { + float32.remove(regVal.index); + } + float64.add(regVal.index); + break; + case Object: + unsigned64.add(regVal.index); + break; + default: + throw GraalInternalError.shouldNotReachHere("unhandled register type " + value.toString()); } } return value; @@ -169,7 +202,12 @@ for (Block b : lirGen.lir.codeEmittingOrder()) { for (LIRInstruction op : lirGen.lir.lir(b)) { - op.forEachOutput(trackRegisterKind); + if (op instanceof LabelOp) { + // Don't consider this as a definition + } else { + op.forEachTemp(trackRegisterKind); + op.forEachOutput(trackRegisterKind); + } } } @@ -179,6 +217,20 @@ for (Integer i : signed64) { codeBuffer.emitString(".reg .s64 %r" + i.intValue() + ";"); } + for (Integer i : unsigned64) { + codeBuffer.emitString(".reg .u64 %r" + i.intValue() + ";"); + } + for (Integer i : float32) { + codeBuffer.emitString(".reg .f32 %r" + i.intValue() + ";"); + } + for (Integer i : float64) { + codeBuffer.emitString(".reg .f64 %r" + i.intValue() + ";"); + } + // emit predicate register declaration + int maxPredRegNum = ((PTXLIRGenerator) lirGen).getNextPredRegNumber(); + if (maxPredRegNum > 0) { + codeBuffer.emitString(".reg .pred %p<" + maxPredRegNum + ">;"); + } } @Override @@ -192,6 +244,7 @@ try { emitRegisterDecl(tasm, lirGen, codeCacheOwner); } catch (GraalInternalError e) { + e.printStackTrace(); // TODO : Better error handling needs to be done once // all types of parameters are handled. codeBuffer.setPosition(0); @@ -202,6 +255,7 @@ try { lirGen.lir.emitCode(tasm); } catch (GraalInternalError e) { + e.printStackTrace(); // TODO : Better error handling needs to be done once // all types of parameters are handled. codeBuffer.setPosition(0);
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXLIRGenerator.java Wed Oct 02 13:26:31 2013 +0200 @@ -24,6 +24,7 @@ package com.oracle.graal.compiler.ptx; import static com.oracle.graal.api.code.ValueUtil.*; +import static com.oracle.graal.api.meta.Value.*; import static com.oracle.graal.lir.ptx.PTXArithmetic.*; import static com.oracle.graal.lir.ptx.PTXBitManipulationOp.IntrinsicOpcode.*; import static com.oracle.graal.lir.ptx.PTXCompare.*; @@ -32,6 +33,7 @@ import com.oracle.graal.api.meta.*; import com.oracle.graal.asm.*; import com.oracle.graal.compiler.gen.*; +import com.oracle.graal.debug.Debug; import com.oracle.graal.graph.*; import com.oracle.graal.lir.*; import com.oracle.graal.lir.StandardOp.JumpOp; @@ -44,8 +46,6 @@ import com.oracle.graal.lir.ptx.PTXArithmetic.Unary2Op; import com.oracle.graal.lir.ptx.PTXCompare.CompareOp; import com.oracle.graal.lir.ptx.PTXControlFlow.BranchOp; -import com.oracle.graal.lir.ptx.PTXControlFlow.CondMoveOp; -import com.oracle.graal.lir.ptx.PTXControlFlow.FloatCondMoveOp; import com.oracle.graal.lir.ptx.PTXControlFlow.ReturnOp; import com.oracle.graal.lir.ptx.PTXControlFlow.ReturnNoValOp; import com.oracle.graal.lir.ptx.PTXControlFlow.SequentialSwitchOp; @@ -59,13 +59,22 @@ import com.oracle.graal.lir.ptx.PTXMemOp.StoreReturnValOp; import com.oracle.graal.nodes.*; import com.oracle.graal.nodes.calc.*; +import com.oracle.graal.nodes.calc.ConvertNode.Op; import com.oracle.graal.nodes.java.*; +import java.lang.annotation.*; + + /** * This class implements the PTX specific portion of the LIR generator. */ public class PTXLIRGenerator extends LIRGenerator { + // Number of the predicate register that can be used when needed. + // This value will be recorded and incremented in the LIR instruction + // that sets a predicate register. (e.g., CompareOp) + private int nextPredRegNum; + public static final ForeignCallDescriptor ARITHMETIC_FREM = new ForeignCallDescriptor("arithmeticFrem", float.class, float.class, float.class); public static final ForeignCallDescriptor ARITHMETIC_DREM = new ForeignCallDescriptor("arithmeticDrem", double.class, double.class, double.class); @@ -73,13 +82,20 @@ @Override public LIRInstruction createMove(AllocatableValue result, Value input) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXSpillMoveFactory.createMove()"); } } public PTXLIRGenerator(StructuredGraph graph, CodeCacheProvider runtime, TargetDescription target, FrameMap frameMap, CallingConvention cc, LIR lir) { super(graph, runtime, target, frameMap, cc, lir); lir.spillMoveFactory = new PTXSpillMoveFactory(); + int callVariables = cc.getArgumentCount() + (cc.getReturn() == Value.ILLEGAL ? 0 : 1); + lir.setFirstVariableNumber(callVariables); + nextPredRegNum = 0; + } + + public int getNextPredRegNumber() { + return nextPredRegNum; } @Override @@ -107,7 +123,9 @@ if (isRegister(value)) { return asRegister(value).asValue(value.getKind().getStackKind()); } else if (isStackSlot(value)) { - return StackSlot.get(value.getKind().getStackKind(), asStackSlot(value).getRawOffset(), asStackSlot(value).getRawAddFrameSize()); + return StackSlot.get(value.getKind().getStackKind(), + asStackSlot(value).getRawOffset(), + asStackSlot(value).getRawAddFrameSize()); } else { throw GraalInternalError.shouldNotReachHere(); } @@ -119,23 +137,62 @@ public void emitPrologue() { // Need to emit .param directives based on incoming arguments and return value CallingConvention incomingArguments = cc; - int argCount = incomingArguments.getArgumentCount(); - // Additional argument for return value. - Value[] params = new Value[argCount + 1]; - for (int i = 0; i < argCount; i++) { - params[i] = toParamKind(incomingArguments.getArgument(i)); + Object returnObject = incomingArguments.getReturn(); + AllocatableValue[] params; + int argCount; + + if (returnObject == Value.ILLEGAL) { + params = incomingArguments.getArguments(); + } else { + argCount = incomingArguments.getArgumentCount(); + params = new Variable[argCount + 1]; + for (int i = 0; i < argCount; i++) { + params[i] = incomingArguments.getArgument(i); + } + params[argCount] = (Variable) returnObject; } - // Add the return value as the last parameter. - params[argCount] = incomingArguments.getReturn(); append(new PTXParameterOp(params)); for (LocalNode local : graph.getNodes(LocalNode.class)) { Value param = params[local.index()]; - assert param.getKind() == local.kind().getStackKind(); - setResult(local, emitLoadParam(param.getKind(), param, null)); + Annotation[] annos = graph.method().getParameterAnnotations()[local.index()]; + Warp warpAnnotation = null; + + if (annos != null) { + for (int a = 0; a < annos.length; a++) { + if (annos[a].annotationType().equals(Warp.class)) { + warpAnnotation = (Warp) annos[a]; + } + } + } + if (warpAnnotation != null) { + setResult(local, emitWarpParam(param.getKind(), warpAnnotation)); + } else { + setResult(local, emitLoadParam(param.getKind(), param, null)); + } } } + public Variable emitWarpParam(Kind kind, Warp annotation) { + Variable result = newVariable(kind); + Variable tid = newVariable(Kind.Char); + + switch (annotation.dimension()) { + case X: + tid.setName("%tid.x"); + break; + case Y: + tid.setName("%tid.y"); + break; + case Z: + tid.setName("%tid.y"); + break; + } + emitMove(result, tid); + + return result; + } + @Override public Variable emitMove(Value input) { Variable result = newVariable(input.getKind()); @@ -156,37 +213,45 @@ public PTXAddressValue emitAddress(Value base, long displacement, Value index, int scale) { AllocatableValue baseRegister; long finalDisp = displacement; + if (isConstant(base)) { if (asConstant(base).isNull()) { baseRegister = Value.ILLEGAL; - } else if (asConstant(base).getKind() != Kind.Object) { + } else if (asConstant(base).getKind() != Kind.Object && !runtime.needsDataPatch(asConstant(base))) { finalDisp += asConstant(base).asLong(); baseRegister = Value.ILLEGAL; } else { baseRegister = load(base); } + } else if (base.equals(Value.ILLEGAL)) { + baseRegister = Value.ILLEGAL; } else { baseRegister = asAllocatable(base); } - if (index != Value.ILLEGAL && scale != 0) { + if (!index.equals(Value.ILLEGAL)) { if (isConstant(index)) { finalDisp += asConstant(index).asLong() * scale; } else { + Value convertedIndex; Value indexRegister; + + convertedIndex = emitConvert(Op.I2L, index); if (scale != 1) { - indexRegister = emitMul(index, Constant.forInt(scale)); + if (CodeUtil.isPowerOf2(scale)) { + indexRegister = emitShl(convertedIndex, Constant.forInt(CodeUtil.log2(scale))); + } else { + indexRegister = emitMul(convertedIndex, Constant.forInt(scale)); + } } else { - indexRegister = index; + indexRegister = convertedIndex; } - - if (baseRegister == Value.ILLEGAL) { + if (baseRegister.equals(Value.ILLEGAL)) { baseRegister = asAllocatable(indexRegister); } else { - Variable newBase = newVariable(Kind.Int); - emitMove(newBase, baseRegister); - baseRegister = newBase; - baseRegister = emitAdd(baseRegister, indexRegister); + Variable longBaseRegister = newVariable(Kind.Long); + emitMove(longBaseRegister, baseRegister); + baseRegister = emitAdd(longBaseRegister, indexRegister); } } } @@ -195,6 +260,8 @@ } private PTXAddressValue asAddress(Value address) { + assert address != null; + if (address instanceof PTXAddressValue) { return (PTXAddressValue) address; } else { @@ -219,7 +286,7 @@ @Override public Variable emitAddress(StackSlot address) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitAddress()"); } @Override @@ -231,24 +298,24 @@ public void emitCompareBranch(Value left, Value right, Condition cond, boolean unorderedIsTrue, LabelRef label) { switch (left.getKind().getStackKind()) { case Int: - append(new CompareOp(ICMP, cond, left, right)); - append(new BranchOp(cond, label)); + append(new CompareOp(ICMP, cond, left, right, nextPredRegNum)); + append(new BranchOp(cond, label, nextPredRegNum++)); break; case Long: - append(new CompareOp(LCMP, cond, left, right)); - append(new BranchOp(cond, label)); + append(new CompareOp(LCMP, cond, left, right, nextPredRegNum)); + append(new BranchOp(cond, label, nextPredRegNum++)); break; case Float: - append(new CompareOp(FCMP, cond, left, right)); - append(new BranchOp(cond, label)); + append(new CompareOp(FCMP, cond, left, right, nextPredRegNum)); + append(new BranchOp(cond, label, nextPredRegNum++)); break; case Double: - append(new CompareOp(DCMP, cond, left, right)); - append(new BranchOp(cond, label)); + append(new CompareOp(DCMP, cond, left, right, nextPredRegNum)); + append(new BranchOp(cond, label, nextPredRegNum++)); break; case Object: - append(new CompareOp(ACMP, cond, left, right)); - append(new BranchOp(cond, label)); + append(new CompareOp(ACMP, cond, left, right, nextPredRegNum)); + append(new BranchOp(cond, label, nextPredRegNum++)); break; default: throw GraalInternalError.shouldNotReachHere("" + left.getKind()); @@ -257,79 +324,22 @@ @Override public void emitOverflowCheckBranch(LabelRef label, boolean negated) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitOverflowCheckBranch()"); } @Override public void emitIntegerTestBranch(Value left, Value right, boolean negated, LabelRef label) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitIntegerTestBranch()"); } @Override public Variable emitConditionalMove(Value left, Value right, Condition cond, boolean unorderedIsTrue, Value trueValue, Value falseValue) { - boolean mirrored = emitCompare(cond, left, right); - Condition finalCondition = mirrored ? cond.mirror() : cond; - - Variable result = newVariable(trueValue.getKind()); - switch (left.getKind().getStackKind()) { - case Int: - case Long: - case Object: - append(new CondMoveOp(result, finalCondition, load(trueValue), loadNonConst(falseValue))); - break; - case Float: - case Double: - append(new FloatCondMoveOp(result, finalCondition, unorderedIsTrue, load(trueValue), load(falseValue))); - break; - default: - throw GraalInternalError.shouldNotReachHere("missing: " + left.getKind()); - } - return result; + // TODO: There is no conventional conditional move instruction in PTX. + // So, this method is changed to throw NYI exception. + // To be revisited if this needs to be really implemented. + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitConditionalMove()"); } - /** - * This method emits the compare instruction, and may reorder the operands. It returns true if - * it did so. - * - * - * @param a the left operand of the comparison - * @param b the right operand of the comparison - * @return true if the left and right operands were switched, false otherwise - */ - private boolean emitCompare(Condition cond, Value a, Value b) { - Variable left; - Value right; - boolean mirrored; - if (LIRValueUtil.isVariable(b)) { - left = load(b); - right = loadNonConst(a); - mirrored = true; - } else { - left = load(a); - right = loadNonConst(b); - mirrored = false; - } - switch (left.getKind().getStackKind()) { - case Int: - append(new CompareOp(ICMP, cond, left, right)); - break; - case Long: - append(new CompareOp(LCMP, cond, left, right)); - break; - case Object: - append(new CompareOp(ACMP, cond, left, right)); - break; - case Float: - append(new CompareOp(FCMP, cond, left, right)); - break; - case Double: - append(new CompareOp(DCMP, cond, left, right)); - break; - default: - throw GraalInternalError.shouldNotReachHere(); - } - return mirrored; - } @Override public Variable emitIntegerTestMove(Value left, Value right, Value trueValue, Value falseValue) { @@ -483,12 +493,12 @@ @Override public Variable emitUDiv(Value a, Value b, DeoptimizingNode deopting) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitUDiv()"); } @Override public Variable emitURem(Value a, Value b, DeoptimizingNode deopting) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitURem()"); } @Override @@ -548,7 +558,7 @@ append(new Op2Stack(ISHL, result, a, loadNonConst(b))); break; case Long: - append(new Op1Stack(LSHL, result, loadNonConst(b))); + append(new Op2Stack(LSHL, result, a, loadNonConst(b))); break; default: throw GraalInternalError.shouldNotReachHere(); @@ -564,7 +574,7 @@ append(new Op2Stack(ISHR, result, a, loadNonConst(b))); break; case Long: - append(new Op1Stack(LSHR, result, loadNonConst(b))); + append(new Op2Stack(LSHR, result, a, loadNonConst(b))); break; default: throw GraalInternalError.shouldNotReachHere(); @@ -663,28 +673,28 @@ } @Override - public void emitDeoptimize(DeoptimizationAction action, DeoptimizingNode deopting) { + public void emitDeoptimize(Value actionAndReason, DeoptimizingNode deopting) { append(new ReturnOp(Value.ILLEGAL)); } @Override public void emitMembar(int barriers) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitMembar()"); } @Override protected void emitDirectCall(DirectCallTargetNode callTarget, Value result, Value[] parameters, Value[] temps, LIRFrameState callState) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitDirectCall()"); } @Override protected void emitIndirectCall(IndirectCallTargetNode callTarget, Value result, Value[] parameters, Value[] temps, LIRFrameState callState) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitIndirectCall()"); } @Override protected void emitForeignCall(ForeignCallLinkage callTarget, Value result, Value[] arguments, Value[] temps, LIRFrameState info) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitForeignCall()"); } @Override @@ -698,47 +708,47 @@ @Override public void emitBitScanForward(Variable result, Value value) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitBitScanForward()"); } @Override public void emitBitScanReverse(Variable result, Value value) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitBitScanReverse()"); } @Override public Value emitMathAbs(Value input) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitMathAbs()"); } @Override public Value emitMathSqrt(Value input) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitMathSqrt()"); } @Override public Value emitMathLog(Value input, boolean base10) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitMathLog()"); } @Override public Value emitMathCos(Value input) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitMathCos()"); } @Override public Value emitMathSin(Value input) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitMathSin()"); } @Override public Value emitMathTan(Value input) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitMathTan()"); } @Override public void emitByteSwap(Variable result, Value input) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitByteSwap()"); } @Override @@ -755,10 +765,10 @@ // Making a copy of the switch value is necessary because jump table destroys the input // value if (key.getKind() == Kind.Int || key.getKind() == Kind.Long) { - append(new SequentialSwitchOp(keyConstants, keyTargets, defaultTarget, key, Value.ILLEGAL)); + append(new SequentialSwitchOp(keyConstants, keyTargets, defaultTarget, key, Value.ILLEGAL, nextPredRegNum)); } else { assert key.getKind() == Kind.Object : key.getKind(); - append(new SequentialSwitchOp(keyConstants, keyTargets, defaultTarget, key, newVariable(Kind.Object))); + append(new SequentialSwitchOp(keyConstants, keyTargets, defaultTarget, key, newVariable(Kind.Object), nextPredRegNum)); } } @@ -772,39 +782,39 @@ // Making a copy of the switch value is necessary because jump table destroys the input // value Variable tmp = emitMove(key); - append(new TableSwitchOp(lowKey, defaultTarget, targets, tmp, newVariable(target.wordKind))); + append(new TableSwitchOp(lowKey, defaultTarget, targets, tmp, newVariable(target.wordKind), nextPredRegNum++)); } @Override public void visitCompareAndSwap(LoweredCompareAndSwapNode node, Value address) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.visitCompareAndSwap()"); } @Override public void visitBreakpointNode(BreakpointNode node) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.visitBreakpointNode()"); } @Override public void visitSafepointNode(SafepointNode i) { - // LIRFrameState info = state(); + // LIRFrameState info = state(i); // append(new PTXSafepointOp(info, runtime().config, this)); - throw new InternalError("NYI"); + Debug.log("visitSafePointNode unimplemented"); } @Override public void emitUnwind(Value operand) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitUnwind()"); } @Override public void emitNullCheck(ValueNode v, DeoptimizingNode deopting) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.emitNullCheck()"); } @Override public void visitInfopointNode(InfopointNode i) { - throw new InternalError("NYI"); + throw GraalInternalError.unimplemented("PTXLIRGenerator.visitInfopointNode()"); } public Variable emitLoadParam(Kind kind, Value address, DeoptimizingNode deopting) { @@ -827,6 +837,13 @@ append(new StoreReturnValOp(kind, storeAddress, input, deopting != null ? state(deopting) : null)); } + @Override + public AllocatableValue resultOperandFor(Kind kind) { + if (kind == Kind.Void) { + return ILLEGAL; + } + return (new Variable(kind, 0)); + } @Override public void visitReturn(ReturnNode x) {
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXTargetMethodAssembler.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXTargetMethodAssembler.java Wed Oct 02 13:26:31 2013 +0200 @@ -36,6 +36,12 @@ private static CompilerToGPU toGPU = HotSpotGraalRuntime.graalRuntime().getCompilerToGPU(); private static boolean validDevice = toGPU.deviceInit(); + private static final int totalProcessors = (validDevice ? toGPU.availableProcessors() : 0); + + public static int getAvailableProcessors() { + return totalProcessors; + } + // detach ?? public PTXTargetMethodAssembler(TargetDescription target, CodeCacheProvider runtime, FrameMap frameMap,
--- a/graal/com.oracle.graal.compiler.sparc/src/com/oracle/graal/compiler/sparc/SPARCLIRGenerator.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.sparc/src/com/oracle/graal/compiler/sparc/SPARCLIRGenerator.java Wed Oct 02 13:26:31 2013 +0200 @@ -837,7 +837,7 @@ } @Override - public void emitDeoptimize(DeoptimizationAction action, DeoptimizingNode deopting) { + public void emitDeoptimize(Value actionAndReason, DeoptimizingNode deopting) { append(new ReturnOp(Value.ILLEGAL)); }
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/BoxingEliminationTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/BoxingEliminationTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -300,16 +300,16 @@ final ValueNode getResult(String snippet) { processMethod(snippet); - assertEquals(1, graph.getNodes(ReturnNode.class).count()); - return graph.getNodes(ReturnNode.class).first().result(); + assertEquals(1, graph.getNodes().filter(ReturnNode.class).count()); + return graph.getNodes().filter(ReturnNode.class).first().result(); } private void processMethod(final String snippet) { graph = parse(snippet); Assumptions assumptions = new Assumptions(false); HighTierContext context = new HighTierContext(runtime(), assumptions, replacements, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL); - new InliningPhase().apply(graph, context); - new PartialEscapePhase(false).apply(graph, context); + new InliningPhase(new CanonicalizerPhase(true)).apply(graph, context); + new PartialEscapePhase(false, new CanonicalizerPhase(true)).apply(graph, context); } private void compareGraphs(final String snippet, final String referenceSnippet) { @@ -326,19 +326,19 @@ Assumptions assumptions = new Assumptions(false); HighTierContext context = new HighTierContext(runtime(), assumptions, replacements, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL); CanonicalizerPhase canonicalizer = new CanonicalizerPhase(true); - new InliningPhase().apply(graph, context); + new InliningPhase(new CanonicalizerPhase(true)).apply(graph, context); if (loopPeeling) { new LoopTransformHighPhase().apply(graph); } new DeadCodeEliminationPhase().apply(graph); canonicalizer.apply(graph, context); - new PartialEscapePhase(false).apply(graph, context); + new PartialEscapePhase(false, canonicalizer).apply(graph, context); new DeadCodeEliminationPhase().apply(graph); canonicalizer.apply(graph, context); StructuredGraph referenceGraph = parse(referenceSnippet); - new InliningPhase().apply(referenceGraph, context); + new InliningPhase(new CanonicalizerPhase(true)).apply(referenceGraph, context); new DeadCodeEliminationPhase().apply(referenceGraph); new CanonicalizerPhase(true).apply(referenceGraph, context);
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/CompareCanonicalizerTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/CompareCanonicalizerTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -30,12 +30,13 @@ import com.oracle.graal.nodes.*; import com.oracle.graal.nodes.calc.*; import com.oracle.graal.phases.common.*; +import com.oracle.graal.phases.tiers.*; public class CompareCanonicalizerTest extends GraalCompilerTest { private StructuredGraph getCanonicalizedGraph(String name) { StructuredGraph graph = parse(name); - new CanonicalizerPhase.Instance(runtime(), null, true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), null, replacements)); return graph; } @@ -53,7 +54,7 @@ assertEquals(referenceGraph, graph); } Assumptions assumptions = new Assumptions(false); - new CanonicalizerPhase.Instance(runtime(), assumptions, true).apply(referenceGraph); + new CanonicalizerPhase(true).apply(referenceGraph, new PhaseContext(runtime(), assumptions, replacements)); for (int i = 1; i < 4; i++) { StructuredGraph graph = getCanonicalizedGraph("canonicalCompare" + i); assertEquals(referenceGraph, graph);
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ConditionTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ConditionTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -45,7 +45,7 @@ Constant b = Constant.forInt(i < 100 ? a.asInt() : rand.nextInt()); boolean result1 = c1.foldCondition(a, b, null, false); boolean result2 = c2.foldCondition(a, b, null, false); - if (result1 && implies) { + if (result1) { assertTrue(result2); } } @@ -60,7 +60,7 @@ for (Condition c1 : Condition.values()) { for (Condition c2 : Condition.values()) { Condition join = c1.join(c2); - assertTrue(join == c2.join(c1)); + assertEquals(join, c2.join(c1)); if (join != null) { for (int i = 0; i < 1000; i++) { Constant a = Constant.forInt(rand.nextInt()); @@ -70,6 +70,8 @@ boolean resultJoin = join.foldCondition(a, b, null, false); if (result1 && result2) { assertTrue(resultJoin); + } else { + assertFalse(resultJoin); } } } @@ -83,7 +85,7 @@ for (Condition c1 : Condition.values()) { for (Condition c2 : Condition.values()) { Condition meet = c1.meet(c2); - assertTrue(meet == c2.meet(c1)); + assertEquals(meet, c2.meet(c1)); if (meet != null) { for (int i = 0; i < 1000; i++) { Constant a = Constant.forInt(rand.nextInt()); @@ -93,6 +95,8 @@ boolean resultMeet = meet.foldCondition(a, b, null, false); if (result1 || result2) { assertTrue(resultMeet); + } else { + assertFalse(resultMeet); } } }
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ConditionalEliminationTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ConditionalEliminationTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -32,6 +32,7 @@ import com.oracle.graal.nodes.java.*; import com.oracle.graal.nodes.java.MethodCallTargetNode.InvokeKind; import com.oracle.graal.phases.common.*; +import com.oracle.graal.phases.tiers.*; /** * Collection of tests for {@link ConditionalEliminationPhase} including those that triggered bugs @@ -141,7 +142,7 @@ StructuredGraph graph = parse("testNullnessSnippet"); new ConditionalEliminationPhase(runtime()).apply(graph); - new CanonicalizerPhase.Instance(runtime(), null, true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), null, replacements)); for (ConstantNode constant : graph.getNodes().filter(ConstantNode.class)) { assertTrue("unexpected constant: " + constant, constant.asConstant().isNull() || constant.asConstant().asInt() > 0); } @@ -163,7 +164,7 @@ @Test public void testDisjunction() { StructuredGraph graph = parse("testDisjunctionSnippet"); - new CanonicalizerPhase.Instance(runtime(), null, true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), null, replacements)); IfNode ifNode = (IfNode) graph.start().next(); InstanceOfNode instanceOf = (InstanceOfNode) ifNode.condition(); IsNullNode x = graph.unique(new IsNullNode(graph.getLocal(0))); @@ -171,9 +172,9 @@ ShortCircuitOrNode disjunction = graph.unique(new ShortCircuitOrNode(x, false, y, false, NOT_FREQUENT_PROBABILITY)); LogicNegationNode negation = graph.unique(new LogicNegationNode(disjunction)); ifNode.setCondition(negation); - new CanonicalizerPhase.Instance(runtime(), null, true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), null, replacements)); new ConditionalEliminationPhase(runtime()).apply(graph); - new CanonicalizerPhase.Instance(runtime(), null, true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), null, replacements)); for (ConstantNode constant : graph.getNodes().filter(ConstantNode.class)) { assertTrue("unexpected constant: " + constant, constant.asConstant().isNull() || constant.asConstant().asInt() > 0); } @@ -191,10 +192,10 @@ public void testInvoke() { test("testInvokeSnippet", new Integer(16)); StructuredGraph graph = parse("testInvokeSnippet"); - new CanonicalizerPhase.Instance(runtime(), null, true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), null, replacements)); new ConditionalEliminationPhase(runtime()).apply(graph); - InvokeNode invoke = graph.getNodes(InvokeNode.class).first(); + InvokeNode invoke = graph.getNodes().filter(InvokeNode.class).first(); assertEquals(InvokeKind.Special, ((MethodCallTargetNode) invoke.callTarget()).invokeKind()); } @@ -221,11 +222,43 @@ @Test public void testTypeMerging() { StructuredGraph graph = parse("testTypeMergingSnippet"); - new CanonicalizerPhase.Instance(runtime(), null, true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), null, replacements)); new ConditionalEliminationPhase(runtime()).apply(graph); - new CanonicalizerPhase.Instance(runtime(), null, true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), null, replacements)); assertEquals(0, graph.getNodes().filter(StoreFieldNode.class).count()); } + public static String testInstanceOfCheckCastSnippet(Object e) { + if (e instanceof Entry) { + return ((Entry) e).name; + } + return null; + } + + @Test + public void testInstanceOfCheckCast() { + StructuredGraph graph = parse("testInstanceOfCheckCastSnippet"); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), null, replacements)); + new ConditionalEliminationPhase(runtime()).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), null, replacements)); + + assertEquals(0, graph.getNodes().filter(CheckCastNode.class).count()); + } + + @Test + @Ignore + public void testInstanceOfCheckCastLowered() { + StructuredGraph graph = parse("testInstanceOfCheckCastSnippet"); + + CanonicalizerPhase canonicalizer = new CanonicalizerPhase(true); + PhaseContext context = new PhaseContext(runtime(), null, replacements); + + new LoweringPhase(canonicalizer).apply(graph, context); + canonicalizer.apply(graph, context); + new ConditionalEliminationPhase(runtime()).apply(graph); + canonicalizer.apply(graph, context); + + assertEquals(0, graph.getNodes().filter(GuardNode.class).count()); + } }
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/DegeneratedLoopsTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/DegeneratedLoopsTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -83,7 +83,7 @@ public void run() { StructuredGraph graph = parse(snippet); HighTierContext context = new HighTierContext(runtime(), new Assumptions(false), replacements, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL); - new InliningPhase().apply(graph, context); + new InliningPhase(new CanonicalizerPhase(true)).apply(graph, context); new CanonicalizerPhase(true).apply(graph, context); Debug.dump(graph, "Graph"); StructuredGraph referenceGraph = parse(REFERENCE_SNIPPET);
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/EliminateNestedCheckCastsTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/EliminateNestedCheckCastsTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -31,6 +31,7 @@ import com.oracle.graal.nodes.*; import com.oracle.graal.nodes.java.*; import com.oracle.graal.phases.common.*; +import com.oracle.graal.phases.tiers.*; public class EliminateNestedCheckCastsTest extends GraalCompilerTest { @@ -113,8 +114,8 @@ public StructuredGraph call() throws Exception { Debug.dump(graph, "After parsing: " + snippet); Assert.assertEquals(checkcasts, graph.getNodes().filter(CheckCastNode.class).count()); - new CanonicalizerPhase.Instance(runtime(), new Assumptions(false), true).apply(graph); - Assert.assertEquals(afterCanon, graph.getNodes(CheckCastNode.class).count()); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), new Assumptions(false), replacements)); + Assert.assertEquals(afterCanon, graph.getNodes().filter(CheckCastNode.class).count()); return graph; }
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/FinalizableSubclassTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/FinalizableSubclassTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -69,7 +69,7 @@ GraphBuilderConfiguration conf = GraphBuilderConfiguration.getSnippetDefault(); new GraphBuilderPhase(runtime, conf, OptimisticOptimizations.ALL).apply(graph); HighTierContext context = new HighTierContext(runtime(), assumptions, replacements, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL); - new InliningPhase().apply(graph, context); + new InliningPhase(new CanonicalizerPhase(true)).apply(graph, context); new CanonicalizerPhase(true).apply(graph, context); return graph; }
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/FloatingReadTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/FloatingReadTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -29,7 +29,6 @@ import com.oracle.graal.graph.*; import com.oracle.graal.nodes.*; import com.oracle.graal.nodes.extended.*; -import com.oracle.graal.nodes.spi.Lowerable.*; import com.oracle.graal.phases.common.*; import com.oracle.graal.phases.tiers.*; @@ -60,7 +59,7 @@ public void run() { StructuredGraph graph = parse(snippet); PhaseContext context = new PhaseContext(runtime(), new Assumptions(false), replacements); - new LoweringPhase(LoweringType.BEFORE_GUARDS).apply(graph, context); + new LoweringPhase(new CanonicalizerPhase(true)).apply(graph, context); new FloatingReadPhase().apply(graph); ReturnNode returnNode = null;
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/GraalCompilerTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/GraalCompilerTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -25,9 +25,11 @@ import static com.oracle.graal.api.code.CodeUtil.*; import static com.oracle.graal.phases.GraalOptions.*; +import java.io.*; import java.lang.reflect.*; import java.util.*; import java.util.concurrent.*; +import java.util.concurrent.atomic.*; import org.junit.*; import org.junit.internal.*; @@ -87,7 +89,7 @@ } @BeforeClass - public static void initializeDebgging() { + public static void initializeDebugging() { DebugEnvironment.initialize(System.out); } @@ -114,8 +116,8 @@ protected void assertConstantReturn(StructuredGraph graph, int value) { String graphString = getCanonicalGraphString(graph, false); - Assert.assertEquals("unexpected number of ReturnNodes: " + graphString, graph.getNodes(ReturnNode.class).count(), 1); - ValueNode result = graph.getNodes(ReturnNode.class).first().result(); + Assert.assertEquals("unexpected number of ReturnNodes: " + graphString, graph.getNodes().filter(ReturnNode.class).count(), 1); + ValueNode result = graph.getNodes().filter(ReturnNode.class).first().result(); Assert.assertTrue("unexpected ReturnNode result node: " + graphString, result.isConstant()); Assert.assertEquals("unexpected ReturnNode result kind: " + graphString, result.asConstant().getKind(), Kind.Int); Assert.assertEquals("unexpected ReturnNode result: " + graphString, result.asConstant().asInt(), value); @@ -169,7 +171,7 @@ return parse(getMethod(methodName)); } - private static int compilationId = 0; + private static AtomicInteger compilationId = new AtomicInteger(); /** * Compares two given objects for {@linkplain Assert#assertEquals(Object, Object) equality}. @@ -206,11 +208,45 @@ } } + @SuppressWarnings("serial") + public static class MultiCauseAssertionError extends AssertionError { + + private Throwable[] causes; + + public MultiCauseAssertionError(String message, Throwable... causes) { + super(message); + this.causes = causes; + } + + @Override + public void printStackTrace(PrintStream out) { + super.printStackTrace(out); + int num = 0; + for (Throwable cause : causes) { + if (cause != null) { + out.print("cause " + (num++)); + cause.printStackTrace(out); + } + } + } + + @Override + public void printStackTrace(PrintWriter out) { + super.printStackTrace(out); + int num = 0; + for (Throwable cause : causes) { + if (cause != null) { + out.print("cause " + (num++) + ": "); + cause.printStackTrace(out); + } + } + } + } + protected void testN(int n, final String name, final Object... args) { - final Throwable[] errors = new Throwable[n]; + final List<Throwable> errors = new ArrayList<>(n); Thread[] threads = new Thread[n]; for (int i = 0; i < n; i++) { - final int idx = i; Thread t = new Thread(i + ":" + name) { @Override @@ -218,26 +254,23 @@ try { test(name, args); } catch (Throwable e) { - errors[idx] = e; + errors.add(e); } } }; threads[i] = t; t.start(); } - int failed = 0; for (int i = 0; i < n; i++) { try { threads[i].join(); } catch (InterruptedException e) { - errors[i] = e; - } - if (errors[i] != null) { - errors[i].printStackTrace(); - failed++; + errors.add(e); } } - Assert.assertTrue(failed + " of " + n + " failed", failed == 0); + if (!errors.isEmpty()) { + throw new MultiCauseAssertionError(errors.size() + " failures", errors.toArray(new Throwable[errors.size()])); + } } protected Object referenceInvoke(Method method, Object receiver, Object... args) throws IllegalAccessException, IllegalArgumentException, InvocationTargetException { @@ -421,7 +454,7 @@ } } - final int id = compilationId++; + final int id = compilationId.incrementAndGet(); InstalledCode installedCode = Debug.scope("Compiling", new Object[]{runtime, new DebugDumpScope(String.valueOf(id), true)}, new Callable<InstalledCode>() {
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/IfCanonicalizerTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/IfCanonicalizerTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -31,6 +31,7 @@ import com.oracle.graal.graph.*; import com.oracle.graal.nodes.*; import com.oracle.graal.phases.common.*; +import com.oracle.graal.phases.tiers.*; /** * In the following tests, the usages of local variable "a" are replaced with the integer constant @@ -144,7 +145,7 @@ n.replaceFirstInput(local, constant); } Debug.dump(graph, "Graph"); - new CanonicalizerPhase.Instance(runtime(), new Assumptions(false), true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), new Assumptions(false), replacements)); for (FrameState fs : local.usages().filter(FrameState.class).snapshot()) { fs.replaceFirstInput(local, null); }
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/InfopointReasonTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/InfopointReasonTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -74,7 +74,7 @@ final Method method = getMethod("testMethod"); final StructuredGraph graph = parseDebug(method); int graphLineSPs = 0; - for (InfopointNode ipn : graph.getNodes(InfopointNode.class)) { + for (InfopointNode ipn : graph.getNodes().filter(InfopointNode.class)) { if (ipn.reason == InfopointReason.LINE_NUMBER) { ++graphLineSPs; }
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/InvokeExceptionTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/InvokeExceptionTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -67,7 +67,7 @@ } Assumptions assumptions = new Assumptions(false); HighTierContext context = new HighTierContext(runtime(), assumptions, replacements, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL); - new InliningPhase(hints).apply(graph, context); + new InliningPhase(hints, new CanonicalizerPhase(true)).apply(graph, context); new CanonicalizerPhase(true).apply(graph, context); new DeadCodeEliminationPhase().apply(graph); }
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/InvokeHintsTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/InvokeHintsTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -78,7 +78,7 @@ Assumptions assumptions = new Assumptions(false); HighTierContext context = new HighTierContext(runtime(), assumptions, replacements, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL); - new InliningPhase(hints).apply(graph, context); + new InliningPhase(hints, new CanonicalizerPhase(true)).apply(graph, context); new CanonicalizerPhase(true).apply(graph, context); new DeadCodeEliminationPhase().apply(graph); StructuredGraph referenceGraph = parse(REFERENCE_SNIPPET);
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/LockEliminationTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/LockEliminationTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -29,7 +29,6 @@ import com.oracle.graal.api.code.*; import com.oracle.graal.nodes.*; import com.oracle.graal.nodes.java.*; -import com.oracle.graal.nodes.spi.Lowerable.LoweringType; import com.oracle.graal.phases.*; import com.oracle.graal.phases.common.*; import com.oracle.graal.phases.tiers.*; @@ -62,7 +61,7 @@ test("testSynchronizedSnippet", new A(), new A()); StructuredGraph graph = getGraph("testSynchronizedSnippet"); - new CanonicalizerPhase.Instance(runtime(), null, true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), null, replacements)); new LockEliminationPhase().apply(graph); assertEquals(1, graph.getNodes().filter(MonitorEnterNode.class).count()); assertEquals(1, graph.getNodes().filter(MonitorExitNode.class).count()); @@ -80,7 +79,7 @@ test("testSynchronizedMethodSnippet", new A()); StructuredGraph graph = getGraph("testSynchronizedMethodSnippet"); - new CanonicalizerPhase.Instance(runtime(), null, true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), null, replacements)); new LockEliminationPhase().apply(graph); assertEquals(1, graph.getNodes().filter(MonitorEnterNode.class).count()); assertEquals(1, graph.getNodes().filter(MonitorExitNode.class).count()); @@ -93,10 +92,10 @@ Assumptions assumptions = new Assumptions(true); HighTierContext context = new HighTierContext(runtime(), assumptions, replacements, null, phasePlan, OptimisticOptimizations.ALL); new CanonicalizerPhase(true).apply(graph, context); - new InliningPhase().apply(graph, context); + new InliningPhase(new CanonicalizerPhase(true)).apply(graph, context); new CanonicalizerPhase(true).apply(graph, context); new DeadCodeEliminationPhase().apply(graph); - new LoweringPhase(LoweringType.BEFORE_GUARDS).apply(graph, context); + new LoweringPhase(new CanonicalizerPhase(true)).apply(graph, context); new ValueAnchorCleanupPhase().apply(graph); new LockEliminationPhase().apply(graph); return graph;
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/LoopUnswitchTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/LoopUnswitchTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -30,6 +30,7 @@ import com.oracle.graal.loop.phases.*; import com.oracle.graal.nodes.*; import com.oracle.graal.phases.common.*; +import com.oracle.graal.phases.tiers.*; public class LoopUnswitchTest extends GraalCompilerTest { @@ -133,8 +134,8 @@ } Assumptions assumptions = new Assumptions(false); - new CanonicalizerPhase.Instance(runtime(), assumptions, true).apply(graph); - new CanonicalizerPhase.Instance(runtime(), assumptions, true).apply(referenceGraph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), assumptions, replacements)); + new CanonicalizerPhase(true).apply(referenceGraph, new PhaseContext(runtime(), assumptions, replacements)); Debug.scope("Test", new DebugDumpScope("Test:" + snippet), new Runnable() { @Override
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/MemoryScheduleTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/MemoryScheduleTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -36,7 +36,6 @@ import com.oracle.graal.nodes.*; import com.oracle.graal.nodes.cfg.*; import com.oracle.graal.nodes.extended.*; -import com.oracle.graal.nodes.spi.Lowerable.LoweringType; import com.oracle.graal.nodes.util.*; import com.oracle.graal.phases.*; import com.oracle.graal.phases.common.*; @@ -222,7 +221,7 @@ public void testArrayCopy() { SchedulePhase schedule = getFinalSchedule("testArrayCopySnippet", TestMode.INLINED_WITHOUT_FRAMESTATES, MemoryScheduling.OPTIMAL); StructuredGraph graph = schedule.getCFG().getStartBlock().getBeginNode().graph(); - ReturnNode ret = graph.getNodes(ReturnNode.class).first(); + ReturnNode ret = graph.getNodes().filter(ReturnNode.class).first(); assertTrue(ret.result() instanceof FloatingReadNode); assertEquals(schedule.getCFG().blockFor(ret), schedule.getCFG().blockFor(ret.result())); assertReadWithinReturnBlock(schedule, true); @@ -502,9 +501,9 @@ HighTierContext context = new HighTierContext(runtime(), assumptions, replacements, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL); new CanonicalizerPhase(true).apply(graph, context); if (mode == TestMode.INLINED_WITHOUT_FRAMESTATES) { - new InliningPhase().apply(graph, context); + new InliningPhase(new CanonicalizerPhase(true)).apply(graph, context); } - new LoweringPhase(LoweringType.BEFORE_GUARDS).apply(graph, context); + new LoweringPhase(new CanonicalizerPhase(true)).apply(graph, context); if (mode == TestMode.WITHOUT_FRAMESTATES || mode == TestMode.INLINED_WITHOUT_FRAMESTATES) { for (Node node : graph.getNodes()) { if (node instanceof StateSplit) { @@ -523,8 +522,8 @@ MidTierContext midContext = new MidTierContext(runtime(), assumptions, replacements, runtime().getTarget(), OptimisticOptimizations.ALL); new GuardLoweringPhase().apply(graph, midContext); - new LoweringPhase(LoweringType.AFTER_GUARDS).apply(graph, midContext); - new LoweringPhase(LoweringType.AFTER_FSA).apply(graph, midContext); + new LoweringPhase(new CanonicalizerPhase(true)).apply(graph, midContext); + new LoweringPhase(new CanonicalizerPhase(true)).apply(graph, midContext); SchedulePhase schedule = new SchedulePhase(SchedulingStrategy.LATEST_OUT_OF_LOOPS, memsched); schedule.apply(graph);
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/MonitorGraphTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/MonitorGraphTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -95,7 +95,7 @@ } Assumptions assumptions = new Assumptions(false); HighTierContext context = new HighTierContext(runtime(), assumptions, replacements, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL); - new InliningPhase(hints).apply(graph, context); + new InliningPhase(hints, new CanonicalizerPhase(true)).apply(graph, context); new CanonicalizerPhase(true).apply(graph, context); new DeadCodeEliminationPhase().apply(graph); return graph;
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/PushNodesThroughPiTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/PushNodesThroughPiTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -30,7 +30,6 @@ import com.oracle.graal.nodes.*; import com.oracle.graal.nodes.calc.*; import com.oracle.graal.nodes.extended.*; -import com.oracle.graal.nodes.spi.Lowerable.LoweringType; import com.oracle.graal.nodes.type.*; import com.oracle.graal.phases.common.*; import com.oracle.graal.phases.tiers.*; @@ -94,7 +93,7 @@ StructuredGraph graph = parse(snippet); PhaseContext context = new PhaseContext(runtime(), new Assumptions(false), replacements); CanonicalizerPhase canonicalizer = new CanonicalizerPhase(true); - new LoweringPhase(LoweringType.BEFORE_GUARDS).apply(graph, context); + new LoweringPhase(canonicalizer).apply(graph, context); canonicalizer.apply(graph, context); new PushThroughPiPhase().apply(graph); canonicalizer.apply(graph, context);
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ReadAfterCheckCastTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ReadAfterCheckCastTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -29,7 +29,6 @@ import com.oracle.graal.debug.*; import com.oracle.graal.nodes.*; import com.oracle.graal.nodes.extended.*; -import com.oracle.graal.nodes.spi.Lowerable.*; import com.oracle.graal.phases.common.*; import com.oracle.graal.phases.tiers.*; @@ -85,7 +84,7 @@ public void run() { StructuredGraph graph = parse(snippet); PhaseContext context = new PhaseContext(runtime(), new Assumptions(false), replacements); - new LoweringPhase(LoweringType.BEFORE_GUARDS).apply(graph, context); + new LoweringPhase(new CanonicalizerPhase(true)).apply(graph, context); new FloatingReadPhase().apply(graph); new EliminatePartiallyRedundantGuardsPhase(true, false).apply(graph); new ReadEliminationPhase().apply(graph);
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ReassociateAndCanonicalTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ReassociateAndCanonicalTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -28,6 +28,7 @@ import com.oracle.graal.graph.*; import com.oracle.graal.nodes.*; import com.oracle.graal.phases.common.*; +import com.oracle.graal.phases.tiers.*; public class ReassociateAndCanonicalTest extends GraalCompilerTest { @@ -241,12 +242,12 @@ return (2 - rnd) - 1; } - private <T extends Node & Node.IterableNodeType> void test(String test, String ref) { + private <T extends Node & IterableNodeType> void test(String test, String ref) { StructuredGraph testGraph = parse(test); Assumptions assumptions = new Assumptions(false); - new CanonicalizerPhase.Instance(runtime(), assumptions, true).apply(testGraph); + new CanonicalizerPhase(true).apply(testGraph, new PhaseContext(runtime(), assumptions, replacements)); StructuredGraph refGraph = parse(ref); - new CanonicalizerPhase.Instance(runtime(), assumptions, true).apply(refGraph); + new CanonicalizerPhase(true).apply(refGraph, new PhaseContext(runtime(), assumptions, replacements)); assertEquals(testGraph, refGraph); } }
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ScalarTypeSystemTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ScalarTypeSystemTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -28,6 +28,7 @@ import com.oracle.graal.debug.*; import com.oracle.graal.nodes.*; import com.oracle.graal.phases.common.*; +import com.oracle.graal.phases.tiers.*; /** * In the following tests, the scalar type system of the compiler should be complete enough to see @@ -166,9 +167,9 @@ StructuredGraph graph = parse(snippet); Debug.dump(graph, "Graph"); Assumptions assumptions = new Assumptions(false); - new CanonicalizerPhase.Instance(runtime(), assumptions, true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), assumptions, replacements)); new ConditionalEliminationPhase(runtime()).apply(graph); - new CanonicalizerPhase.Instance(runtime(), assumptions, true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), assumptions, replacements)); StructuredGraph referenceGraph = parse(referenceSnippet); assertEquals(referenceGraph, graph); }
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/StampCanonicalizerTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/StampCanonicalizerTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -28,6 +28,7 @@ import com.oracle.graal.nodes.*; import com.oracle.graal.nodes.type.*; import com.oracle.graal.phases.common.*; +import com.oracle.graal.phases.tiers.*; /** * This class tests some specific patterns the stamp system should be able to canonicalize away @@ -109,7 +110,7 @@ private void testZeroReturn(String methodName) { StructuredGraph graph = parse(methodName); - new CanonicalizerPhase.Instance(runtime(), new Assumptions(false), true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), new Assumptions(false), replacements)); new DeadCodeEliminationPhase().apply(graph); assertConstantReturn(graph, 0); }
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/StraighteningTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/StraighteningTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -28,6 +28,7 @@ import com.oracle.graal.debug.*; import com.oracle.graal.nodes.*; import com.oracle.graal.phases.common.*; +import com.oracle.graal.phases.tiers.*; public class StraighteningTest extends GraalCompilerTest { @@ -89,7 +90,7 @@ // No debug scope to reduce console noise for @Test(expected = ...) tests StructuredGraph graph = parse(snippet); Debug.dump(graph, "Graph"); - new CanonicalizerPhase.Instance(runtime(), new Assumptions(false), true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), new Assumptions(false), replacements)); StructuredGraph referenceGraph = parse(REFERENCE_SNIPPET); assertEquals(referenceGraph, graph); }
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/TypeSystemTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/TypeSystemTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -35,6 +35,7 @@ import com.oracle.graal.nodes.java.*; import com.oracle.graal.phases.common.*; import com.oracle.graal.phases.schedule.*; +import com.oracle.graal.phases.tiers.*; /** * In the following tests, the scalar type system of the compiler should be complete enough to see @@ -44,7 +45,7 @@ @Test public void test1() { - test("test1Snippet", CheckCastNode.class); + testHelper("test1Snippet", CheckCastNode.class); } public static int test1Snippet(Object a) { @@ -56,7 +57,7 @@ @Test public void test2() { - test("test2Snippet", CheckCastNode.class); + testHelper("test2Snippet", CheckCastNode.class); } public static int test2Snippet(Object a) { @@ -161,7 +162,7 @@ @Test public void test6() { - test("test6Snippet", CheckCastNode.class); + testHelper("test6Snippet", CheckCastNode.class); } public static int test6Snippet(int i) throws IOException { @@ -184,13 +185,13 @@ StructuredGraph graph = parse(snippet); Debug.dump(graph, "Graph"); Assumptions assumptions = new Assumptions(false); - new CanonicalizerPhase.Instance(runtime(), assumptions, true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), assumptions, replacements)); new ConditionalEliminationPhase(runtime()).apply(graph); - new CanonicalizerPhase.Instance(runtime(), assumptions, true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), assumptions, replacements)); // a second canonicalizer is needed to process nested MaterializeNodes - new CanonicalizerPhase.Instance(runtime(), assumptions, true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), assumptions, replacements)); StructuredGraph referenceGraph = parse(referenceSnippet); - new CanonicalizerPhase.Instance(runtime(), assumptions, true).apply(referenceGraph); + new CanonicalizerPhase(true).apply(referenceGraph, new PhaseContext(runtime(), assumptions, replacements)); assertEquals(referenceGraph, graph); } @@ -236,14 +237,14 @@ } } - private <T extends Node & Node.IterableNodeType> void test(String snippet, Class<T> clazz) { + private <T extends Node> void testHelper(String snippet, Class<T> clazz) { StructuredGraph graph = parse(snippet); Debug.dump(graph, "Graph"); Assumptions assumptions = new Assumptions(false); - new CanonicalizerPhase.Instance(runtime(), assumptions, true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), assumptions, replacements)); new ConditionalEliminationPhase(runtime()).apply(graph); - new CanonicalizerPhase.Instance(runtime(), assumptions, true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), assumptions, replacements)); Debug.dump(graph, "Graph"); - Assert.assertFalse("shouldn't have nodes of type " + clazz, graph.getNodes(clazz).iterator().hasNext()); + Assert.assertFalse("shouldn't have nodes of type " + clazz, graph.getNodes().filter(clazz).iterator().hasNext()); } }
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/deopt/CompiledMethodTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/deopt/CompiledMethodTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -32,6 +32,7 @@ import com.oracle.graal.graph.*; import com.oracle.graal.nodes.*; import com.oracle.graal.phases.common.*; +import com.oracle.graal.phases.tiers.*; import com.oracle.graal.test.*; /** @@ -55,7 +56,7 @@ public void test1() { Method method = getMethod("testMethod"); final StructuredGraph graph = parse(method); - new CanonicalizerPhase.Instance(runtime(), new Assumptions(false), true).apply(graph); + new CanonicalizerPhase(true).apply(graph, new PhaseContext(runtime(), new Assumptions(false), replacements)); new DeadCodeEliminationPhase().apply(graph); for (Node node : graph.getNodes()) {
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ea/EarlyReadEliminationTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ea/EarlyReadEliminationTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -42,7 +42,7 @@ graph = parse(snippet); Assumptions assumptions = new Assumptions(false); HighTierContext context = new HighTierContext(runtime(), assumptions, replacements, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL); - new InliningPhase().apply(graph, context); - new EarlyReadEliminationPhase().apply(graph, context); + new InliningPhase(new CanonicalizerPhase(true)).apply(graph, context); + new EarlyReadEliminationPhase(new CanonicalizerPhase(true)).apply(graph, context); } }
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ea/EscapeAnalysisTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ea/EscapeAnalysisTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -234,17 +234,18 @@ Assumptions assumptions = new Assumptions(false); HighTierContext context = new HighTierContext(runtime(), assumptions, replacements, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL); - new InliningPhase().apply(graph, context); + new InliningPhase(new CanonicalizerPhase(true)).apply(graph, context); new DeadCodeEliminationPhase().apply(graph); new CanonicalizerPhase(true).apply(graph, context); - new PartialEscapePhase(iterativeEscapeAnalysis).apply(graph, context); - Assert.assertEquals(1, graph.getNodes(ReturnNode.class).count()); - ReturnNode returnNode = graph.getNodes(ReturnNode.class).first(); + new PartialEscapePhase(iterativeEscapeAnalysis, new CanonicalizerPhase(true)).apply(graph, context); + Assert.assertEquals(1, graph.getNodes().filter(ReturnNode.class).count()); + ReturnNode returnNode = graph.getNodes().filter(ReturnNode.class).first(); if (expectedConstantResult != null) { Assert.assertTrue(returnNode.result().toString(), returnNode.result().isConstant()); Assert.assertEquals(expectedConstantResult, returnNode.result().asConstant()); } - int newInstanceCount = graph.getNodes(NewInstanceNode.class).count() + graph.getNodes(NewArrayNode.class).count() + graph.getNodes(CommitAllocationNode.class).count(); + int newInstanceCount = graph.getNodes().filter(NewInstanceNode.class).count() + graph.getNodes().filter(NewArrayNode.class).count() + + graph.getNodes().filter(CommitAllocationNode.class).count(); Assert.assertEquals(0, newInstanceCount); return returnNode; }
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ea/IterativeInliningTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ea/IterativeInliningTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -82,8 +82,8 @@ final ReturnNode getReturn(String snippet) { processMethod(snippet); - assertEquals(1, graph.getNodes(ReturnNode.class).count()); - return graph.getNodes(ReturnNode.class).first(); + assertEquals(1, graph.getNodes().filter(ReturnNode.class).count()); + return graph.getNodes().filter(ReturnNode.class).first(); } private void processMethod(final String snippet) {
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ea/PEAReadEliminationTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ea/PEAReadEliminationTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -236,15 +236,15 @@ final ReturnNode getReturn(String snippet) { processMethod(snippet); - assertEquals(1, graph.getNodes(ReturnNode.class).count()); - return graph.getNodes(ReturnNode.class).first(); + assertEquals(1, graph.getNodes().filter(ReturnNode.class).count()); + return graph.getNodes().filter(ReturnNode.class).first(); } protected void processMethod(final String snippet) { graph = parse(snippet); Assumptions assumptions = new Assumptions(false); HighTierContext context = new HighTierContext(runtime(), assumptions, replacements, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL); - new InliningPhase().apply(graph, context); - new PartialEscapePhase(false, true).apply(graph, context); + new InliningPhase(new CanonicalizerPhase(true)).apply(graph, context); + new PartialEscapePhase(false, true, new CanonicalizerPhase(true)).apply(graph, context); } }
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ea/PartialEscapeAnalysisTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/ea/PartialEscapeAnalysisTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -126,7 +126,7 @@ @Test public void testCache() { - testMaterialize("testCacheSnippet", 0.5, 1); + testMaterialize("testCacheSnippet", 0.75, 1); } public static class CacheKey { @@ -167,13 +167,13 @@ final void testMaterialize(final String snippet, double expectedProbability, int expectedCount, Class<? extends Node>... invalidNodeClasses) { StructuredGraph result = processMethod(snippet); try { - Assert.assertTrue("partial escape analysis should have removed all NewInstanceNode allocations", result.getNodes(NewInstanceNode.class).isEmpty()); - Assert.assertTrue("partial escape analysis should have removed all NewArrayNode allocations", result.getNodes(NewArrayNode.class).isEmpty()); + Assert.assertTrue("partial escape analysis should have removed all NewInstanceNode allocations", result.getNodes().filter(NewInstanceNode.class).isEmpty()); + Assert.assertTrue("partial escape analysis should have removed all NewArrayNode allocations", result.getNodes().filter(NewArrayNode.class).isEmpty()); NodesToDoubles nodeProbabilities = new ComputeProbabilityClosure(result).apply(); double probabilitySum = 0; int materializeCount = 0; - for (CommitAllocationNode materialize : result.getNodes(CommitAllocationNode.class)) { + for (CommitAllocationNode materialize : result.getNodes().filter(CommitAllocationNode.class)) { probabilitySum += nodeProbabilities.get(materialize) * materialize.getVirtualObjects().size(); materializeCount += materialize.getVirtualObjects().size(); } @@ -199,10 +199,10 @@ Assumptions assumptions = new Assumptions(false); HighTierContext context = new HighTierContext(runtime(), assumptions, replacements, null, getDefaultPhasePlan(), OptimisticOptimizations.ALL); - new InliningPhase().apply(graph, context); + new InliningPhase(new CanonicalizerPhase(true)).apply(graph, context); new DeadCodeEliminationPhase().apply(graph); new CanonicalizerPhase(true).apply(graph, context); - new PartialEscapePhase(false).apply(graph, context); + new PartialEscapePhase(false, new CanonicalizerPhase(true)).apply(graph, context); for (MergeNode merge : graph.getNodes(MergeNode.class)) { merge.setStateAfter(null);
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/inlining/InliningTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/inlining/InliningTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -240,7 +240,7 @@ HighTierContext context = new HighTierContext(runtime(), assumptions, replacements, null, phasePlan, OptimisticOptimizations.ALL); Debug.dump(graph, "Graph"); new CanonicalizerPhase(true).apply(graph, context); - new InliningPhase().apply(graph, context); + new InliningPhase(new CanonicalizerPhase(true)).apply(graph, context); Debug.dump(graph, "Graph"); new CanonicalizerPhase(true).apply(graph, context); new DeadCodeEliminationPhase().apply(graph); @@ -279,7 +279,7 @@ private static int[] countMethodInfopoints(StructuredGraph graph) { int start = 0; int end = 0; - for (InfopointNode ipn : graph.getNodes(InfopointNode.class)) { + for (InfopointNode ipn : graph.getNodes().filter(InfopointNode.class)) { if (ipn.reason == InfopointReason.METHOD_START) { ++start; } else if (ipn.reason == InfopointReason.METHOD_END) {
--- a/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/GraalCompiler.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/GraalCompiler.java Wed Oct 02 13:26:31 2013 +0200 @@ -229,7 +229,7 @@ List<Block> codeEmittingOrder = ComputeBlockOrder.computeCodeEmittingOrder(blocks.length, startBlock, nodeProbabilities); List<Block> linearScanOrder = ComputeBlockOrder.computeLinearScanOrder(blocks.length, startBlock, nodeProbabilities); - LIR lir = new LIR(schedule.getCFG(), schedule.getBlockToNodesMap(), linearScanOrder, codeEmittingOrder, speculationLog); + LIR lir = new LIR(schedule.getCFG(), schedule.getBlockToNodesMap(), linearScanOrder, codeEmittingOrder); Debug.dump(lir, "After linear scan order"); return lir; @@ -238,7 +238,7 @@ } - public static LIRGenerator emitLIR(Backend backend, final TargetDescription target, final LIR lir, StructuredGraph graph, CallingConvention cc) { + public static LIRGenerator emitLIR(final Backend backend, final TargetDescription target, final LIR lir, StructuredGraph graph, CallingConvention cc) { final FrameMap frameMap = backend.newFrameMap(); final LIRGenerator lirGen = backend.newLIRGenerator(graph, frameMap, cc, lir); @@ -269,7 +269,9 @@ Debug.scope("Allocator", new Runnable() { public void run() { - new LinearScan(target, lir, lirGen, frameMap).allocate(); + if (backend.shouldAllocateRegisters()) { + new LinearScan(target, lir, lirGen, frameMap).allocate(); + } } }); return lirGen;
--- a/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/GraalDebugConfig.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/GraalDebugConfig.java Wed Oct 02 13:26:31 2013 +0200 @@ -36,8 +36,6 @@ public class GraalDebugConfig implements DebugConfig { // @formatter:off - @Option(help = "Enable scope-based debugging", name = "Debug") - public static final OptionValue<Boolean> DebugEnabled = new OptionValue<>(true); @Option(help = "Pattern for scope(s) to in which dumping is enabled (see DebugFilter and Debug.dump)") public static final OptionValue<String> Dump = new OptionValue<>(null); @Option(help = "Pattern for scope(s) to in which metering is enabled (see DebugFilter and Debug.metric)") @@ -53,7 +51,7 @@ "Partial - aggregate by partially qualified name (e.g., A.B.C.D.Counter and X.Y.Z.D.Counter will be merged to D.Counter)%n" + "Complete - aggregate by qualified name%n" + "Thread - aggregate by qualified name and thread") - public static final OptionValue<String> DebugValueSummary = new OptionValue<>("Complete"); + public static final OptionValue<String> DebugValueSummary = new OptionValue<>("Name"); @Option(help = "Send Graal IR to dump handlers on error") public static final OptionValue<Boolean> DumpOnError = new OptionValue<>(false); @Option(help = "Enable expensive assertions") @@ -68,6 +66,10 @@ }; // @formatter:on + public static boolean areDebugScopePatternsEnabled() { + return DumpOnError.getValue() || Dump.getValue() != null || Meter.getValue() != null || Time.getValue() != null || Log.getValue() != null; + } + private final DebugFilter logFilter; private final DebugFilter meterFilter; private final DebugFilter timerFilter;
--- a/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/gen/DebugInfoBuilder.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/gen/DebugInfoBuilder.java Wed Oct 02 13:26:31 2013 +0200 @@ -48,7 +48,7 @@ protected HashMap<VirtualObjectNode, VirtualObject> virtualObjects = new HashMap<>(); protected IdentityHashMap<VirtualObjectNode, EscapeObjectState> objectStates = new IdentityHashMap<>(); - public LIRFrameState build(FrameState topState, short reason, LabelRef exceptionEdge) { + public LIRFrameState build(FrameState topState, LabelRef exceptionEdge) { assert virtualObjects.size() == 0; assert objectStates.size() == 0; @@ -101,11 +101,11 @@ } objectStates.clear(); - return newLIRFrameState(reason, exceptionEdge, frame, virtualObjectsArray); + return newLIRFrameState(exceptionEdge, frame, virtualObjectsArray); } - protected LIRFrameState newLIRFrameState(short reason, LabelRef exceptionEdge, BytecodeFrame frame, VirtualObject[] virtualObjectsArray) { - return new LIRFrameState(frame, virtualObjectsArray, exceptionEdge, reason); + protected LIRFrameState newLIRFrameState(LabelRef exceptionEdge, BytecodeFrame frame, VirtualObject[] virtualObjectsArray) { + return new LIRFrameState(frame, virtualObjectsArray, exceptionEdge); } protected BytecodeFrame computeFrameForState(FrameState state) { @@ -122,7 +122,7 @@ if (state.outerFrameState() != null) { caller = computeFrameForState(state.outerFrameState()); } - assert state.bci != FrameState.UNKNOWN_BCI : "bci == " + state.bci; + assert state.bci >= FrameState.BEFORE_BCI : "bci == " + state.bci; return new BytecodeFrame(caller, state.method(), state.bci, state.rethrowException(), state.duringCall(), values, numLocals, numStack, numLocks); }
--- a/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/gen/LIRGenerator.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/gen/LIRGenerator.java Wed Oct 02 13:26:31 2013 +0200 @@ -228,26 +228,26 @@ if (!deopt.canDeoptimize()) { return null; } - return stateFor(deopt.getDeoptimizationState(), deopt.getDeoptimizationReason()); + return stateFor(deopt.getDeoptimizationState()); } public LIRFrameState stateWithExceptionEdge(DeoptimizingNode deopt, LabelRef exceptionEdge) { if (!deopt.canDeoptimize()) { return null; } - return stateForWithExceptionEdge(deopt.getDeoptimizationState(), deopt.getDeoptimizationReason(), exceptionEdge); + return stateForWithExceptionEdge(deopt.getDeoptimizationState(), exceptionEdge); } - public LIRFrameState stateFor(FrameState state, DeoptimizationReason reason) { - return stateForWithExceptionEdge(state, reason, null); + public LIRFrameState stateFor(FrameState state) { + return stateForWithExceptionEdge(state, null); } - public LIRFrameState stateForWithExceptionEdge(FrameState state, DeoptimizationReason reason, LabelRef exceptionEdge) { + public LIRFrameState stateForWithExceptionEdge(FrameState state, LabelRef exceptionEdge) { if (needOnlyOopMaps()) { - return new LIRFrameState(null, null, null, (short) -1); + return new LIRFrameState(null, null, null); } assert state != null; - return debugInfoBuilder.build(state, lir.getDeoptimizationReasons().addSpeculation(reason), exceptionEdge); + return debugInfoBuilder.build(state, exceptionEdge); } /** @@ -616,7 +616,15 @@ @Override public Variable emitForeignCall(ForeignCallLinkage linkage, DeoptimizingNode info, Value... args) { - LIRFrameState state = !linkage.canDeoptimize() ? null : stateFor(info.getDeoptimizationState(), info.getDeoptimizationReason()); + LIRFrameState state = null; + if (linkage.canDeoptimize()) { + if (info != null) { + state = stateFor(info.getDeoptimizationState()); + } else { + assert needOnlyOopMaps(); + state = new LIRFrameState(null, null, null); + } + } // move the arguments into the correct location CallingConvention linkageCc = linkage.getOutgoingCallingConvention();
--- a/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/phases/HighTier.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/phases/HighTier.java Wed Oct 02 13:26:31 2013 +0200 @@ -22,12 +22,12 @@ */ package com.oracle.graal.compiler.phases; +import static com.oracle.graal.compiler.phases.HighTier.Options.*; import static com.oracle.graal.phases.GraalOptions.*; import com.oracle.graal.api.code.*; import com.oracle.graal.api.meta.*; import com.oracle.graal.loop.phases.*; -import com.oracle.graal.nodes.spi.Lowerable.LoweringType; import com.oracle.graal.options.*; import com.oracle.graal.phases.*; import com.oracle.graal.phases.common.*; @@ -37,12 +37,15 @@ public class HighTier extends PhaseSuite<HighTierContext> { - // @formatter:off - @Option(help = "") - public static final OptionValue<Boolean> VerifyUsageWithEquals = new OptionValue<>(true); - @Option(help = "Enable inlining") - public static final OptionValue<Boolean> Inline = new OptionValue<>(true); - // @formatter:on + static class Options { + + // @formatter:off + @Option(help = "") + public static final OptionValue<Boolean> VerifyUsageWithEquals = new OptionValue<>(true); + @Option(help = "Enable inlining") + public static final OptionValue<Boolean> Inline = new OptionValue<>(true); + // @formatter:on + } public HighTier() { CanonicalizerPhase canonicalizer = new CanonicalizerPhase(!AOTCompilation.getValue()); @@ -60,12 +63,12 @@ if (IterativeInlining.getValue()) { appendPhase(new IterativeInliningPhase(canonicalizer)); } else { - appendPhase(new InliningPhase()); + appendPhase(new InliningPhase(canonicalizer)); appendPhase(new DeadCodeEliminationPhase()); if (ConditionalElimination.getValue() && OptCanonicalizer.getValue()) { appendPhase(canonicalizer); - appendPhase(new IterativeConditionalEliminationPhase()); + appendPhase(new IterativeConditionalEliminationPhase(canonicalizer)); } } } @@ -73,15 +76,15 @@ appendPhase(new CleanTypeProfileProxyPhase()); if (FullUnroll.getValue()) { - appendPhase(new LoopFullUnrollPhase(!AOTCompilation.getValue())); + appendPhase(new LoopFullUnrollPhase(canonicalizer)); } if (OptTailDuplication.getValue()) { - appendPhase(new TailDuplicationPhase()); + appendPhase(new TailDuplicationPhase(canonicalizer)); } if (PartialEscapeAnalysis.getValue()) { - appendPhase(new PartialEscapePhase(true)); + appendPhase(new PartialEscapePhase(true, canonicalizer)); } if (OptConvertDeoptsToGuards.getValue()) { @@ -98,6 +101,6 @@ appendPhase(canonicalizer); } - appendPhase(new LoweringPhase(LoweringType.BEFORE_GUARDS)); + appendPhase(new LoweringPhase(canonicalizer)); } }
--- a/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/phases/LowTier.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/phases/LowTier.java Wed Oct 02 13:26:31 2013 +0200 @@ -22,7 +22,8 @@ */ package com.oracle.graal.compiler.phases; -import com.oracle.graal.nodes.spi.Lowerable.LoweringType; +import static com.oracle.graal.phases.GraalOptions.*; + import com.oracle.graal.phases.*; import com.oracle.graal.phases.common.*; import com.oracle.graal.phases.tiers.*; @@ -30,7 +31,9 @@ public class LowTier extends PhaseSuite<LowTierContext> { public LowTier() { - appendPhase(new LoweringPhase(LoweringType.AFTER_FSA)); + CanonicalizerPhase canonicalizer = new CanonicalizerPhase(!AOTCompilation.getValue()); + + appendPhase(new LoweringPhase(canonicalizer)); appendPhase(new ExpandLogicPhase());
--- a/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/phases/MidTier.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/phases/MidTier.java Wed Oct 02 13:26:31 2013 +0200 @@ -25,7 +25,6 @@ import static com.oracle.graal.phases.GraalOptions.*; import com.oracle.graal.loop.phases.*; -import com.oracle.graal.nodes.spi.Lowerable.LoweringType; import com.oracle.graal.phases.*; import com.oracle.graal.phases.common.*; import com.oracle.graal.phases.tiers.*; @@ -47,11 +46,11 @@ appendPhase(new LockEliminationPhase()); if (OptReadElimination.getValue()) { - appendPhase(new EarlyReadEliminationPhase()); + appendPhase(new EarlyReadEliminationPhase(canonicalizer)); } if (OptFloatingReads.getValue()) { - IncrementalCanonicalizerPhase<MidTierContext> incCanonicalizer = new IncrementalCanonicalizerPhase<>(); + IncrementalCanonicalizerPhase<MidTierContext> incCanonicalizer = new IncrementalCanonicalizerPhase<>(canonicalizer); incCanonicalizer.appendPhase(new FloatingReadPhase()); appendPhase(incCanonicalizer); if (OptReadElimination.getValue()) { @@ -69,7 +68,7 @@ } if (ConditionalElimination.getValue() && OptCanonicalizer.getValue()) { - appendPhase(new IterativeConditionalEliminationPhase()); + appendPhase(new IterativeConditionalEliminationPhase(canonicalizer)); } if (OptEliminatePartiallyRedundantGuards.getValue()) { @@ -82,14 +81,16 @@ appendPhase(new LoopSafepointEliminationPhase()); - appendPhase(new SafepointInsertionPhase()); + appendPhase(new LoopSafepointInsertionPhase()); appendPhase(new GuardLoweringPhase()); - appendPhase(new LoweringPhase(LoweringType.AFTER_GUARDS)); + appendPhase(new LoweringPhase(canonicalizer)); appendPhase(new FrameStateAssignmentPhase()); + appendPhase(new DeoptimizationGroupingPhase()); + if (OptCanonicalizer.getValue()) { appendPhase(canonicalizer); }
--- a/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/target/Backend.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/target/Backend.java Wed Oct 02 13:26:31 2013 +0200 @@ -55,6 +55,8 @@ public abstract TargetMethodAssembler newAssembler(LIRGenerator lirGen, CompilationResult compilationResult); + public abstract boolean shouldAllocateRegisters(); + /** * Emits the code for a given graph. *
--- a/graal/com.oracle.graal.debug.test/src/com/oracle/graal/debug/test/DebugHistogramTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.debug.test/src/com/oracle/graal/debug/test/DebugHistogramTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -27,6 +27,7 @@ import org.junit.*; import com.oracle.graal.debug.*; +import com.oracle.graal.debug.internal.*; public class DebugHistogramTest { @@ -34,8 +35,13 @@ public void testEmptyHistogram() { DebugHistogram histogram = Debug.createHistogram("TestHistogram"); ByteArrayOutputStream outputStream = new ByteArrayOutputStream(); - histogram.print(new PrintStream(outputStream)); + + new DebugHistogramAsciiPrinter(new PrintStream(outputStream)).print(histogram); Assert.assertEquals("TestHistogram is empty.\n", outputStream.toString()); + + outputStream.reset(); + new DebugHistogramRPrinter(new PrintStream(outputStream)).print(histogram); + Assert.assertEquals("", outputStream.toString()); } @Test @@ -44,19 +50,28 @@ ByteArrayOutputStream outputStream = new ByteArrayOutputStream(); histogram.add(new Integer(1)); histogram.add(new Integer(1)); - histogram.print(new PrintStream(outputStream)); + new DebugHistogramAsciiPrinter(new PrintStream(outputStream)).print(histogram); String[] lines = outputStream.toString().split("\n"); - Assert.assertEquals(4, lines.length); - Assert.assertEquals("TestHistogram has 1 unique elements and 2 total elements:", lines[0]); - Assert.assertEquals( - "--------------------------------------------------------------------------------------------------------------------------------------------------------------------------", - lines[1]); - Assert.assertEquals( - "| 1 | 2 | ==================================================================================================== |", - lines[2]); - Assert.assertEquals( - "--------------------------------------------------------------------------------------------------------------------------------------------------------------------------", - lines[3]); + // @formatter:off + String[] expected = { + "TestHistogram has 1 unique elements and 2 total elements:", + "--------------------------------------------------------------------------------------------------------------------------------------------------------------------------", + "| 1 | 2 | ==================================================================================================== |", + "--------------------------------------------------------------------------------------------------------------------------------------------------------------------------" + }; + // @formatter:on + Assert.assertArrayEquals(expected, lines); + + outputStream.reset(); + new DebugHistogramRPrinter(new PrintStream(outputStream)).print(histogram); + lines = outputStream.toString().split("\n"); + // @formatter:off + expected = new String[] { + "TestHistogram <- c(2);", + "names(TestHistogram) <- c(\"1\");" + }; + // @formatter:on + Assert.assertArrayEquals(expected, lines); } @Test @@ -66,22 +81,29 @@ histogram.add(new Integer(1)); histogram.add(new Integer(2)); histogram.add(new Integer(2)); - histogram.print(new PrintStream(outputStream)); + new DebugHistogramAsciiPrinter(new PrintStream(outputStream)).print(histogram); String[] lines = outputStream.toString().split("\n"); - Assert.assertEquals(5, lines.length); - Assert.assertEquals("TestHistogram has 2 unique elements and 3 total elements:", lines[0]); - Assert.assertEquals( - "--------------------------------------------------------------------------------------------------------------------------------------------------------------------------", - lines[1]); - Assert.assertEquals( - "| 2 | 2 | ==================================================================================================== |", - lines[2]); - Assert.assertEquals( - "| 1 | 1 | ================================================== |", - lines[3]); - Assert.assertEquals( - "--------------------------------------------------------------------------------------------------------------------------------------------------------------------------", - lines[4]); + // @formatter:off + String[] expected = new String[] { + "TestHistogram has 2 unique elements and 3 total elements:", + "--------------------------------------------------------------------------------------------------------------------------------------------------------------------------", + "| 2 | 2 | ==================================================================================================== |", + "| 1 | 1 | ================================================== |", + "--------------------------------------------------------------------------------------------------------------------------------------------------------------------------" + }; + // @formatter:on + Assert.assertArrayEquals(expected, lines); + + outputStream.reset(); + new DebugHistogramRPrinter(new PrintStream(outputStream)).print(histogram); + lines = outputStream.toString().split("\n"); + // @formatter:off + expected = new String[] { + "TestHistogram <- c(2, 1);", + "names(TestHistogram) <- c(\"2\", \"1\");" + }; + // @formatter:on + Assert.assertArrayEquals(expected, lines); } @Test @@ -89,7 +111,7 @@ DebugHistogram histogram = Debug.createHistogram("TestHistogram"); ByteArrayOutputStream outputStream = new ByteArrayOutputStream(); histogram.add("MyCustomValue"); - histogram.print(new PrintStream(outputStream), Integer.MAX_VALUE, 10, 10); + new DebugHistogramAsciiPrinter(new PrintStream(outputStream), Integer.MAX_VALUE, 10, 10).print(histogram); String[] lines = outputStream.toString().split("\n"); Assert.assertEquals(4, lines.length); Assert.assertEquals("TestHistogram has 1 unique elements and 1 total elements:", lines[0]);
--- a/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/Debug.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/Debug.java Wed Oct 02 13:26:31 2013 +0200 @@ -22,24 +22,49 @@ */ package com.oracle.graal.debug; -import com.oracle.graal.debug.internal.*; +import static com.oracle.graal.debug.Debug.Initialization.*; import java.io.*; import java.util.*; import java.util.concurrent.*; +import com.oracle.graal.debug.internal.*; + +/** + * Scope based debugging facility. This facility is {@link #isEnabled()} if assertions are enabled + * for the {@link Debug} class or the {@value Initialization#INITIALIZER_PROPERTY_NAME} system + * property is {@code "true"} when {@link Debug} is initialized. + */ public class Debug { - private static boolean ENABLED = false; + /** + * Class to assist with initialization of {@link Debug}. + */ + public static class Initialization { - public static void enable() { - ENABLED = true; + public static final String INITIALIZER_PROPERTY_NAME = "graal.debug.enable"; + + private static boolean initialized; + + /** + * Determines if {@link Debug} has been initialized. + */ + public static boolean isDebugInitialized() { + return initialized; + } + } - public static void disable() { - ENABLED = false; + @SuppressWarnings("all") + private static boolean initialize() { + boolean assertionsEnabled = false; + assert assertionsEnabled = true; + Initialization.initialized = true; + return assertionsEnabled || Boolean.getBoolean(INITIALIZER_PROPERTY_NAME); } + private static final boolean ENABLED = initialize(); + public static boolean isEnabled() { return ENABLED; } @@ -330,7 +355,7 @@ public static DebugMetric metric(String name) { if (ENABLED) { - return new MetricImpl(name); + return new MetricImpl(name, true); } else { return VOID_METRIC; } @@ -342,6 +367,9 @@ } } + /** + * Creates an object for counting value frequencies. + */ public static DebugHistogram createHistogram(String name) { return new DebugHistogramImpl(name); } @@ -402,11 +430,19 @@ public void add(long value) { } + + public void setConditional(boolean flag) { + throw new InternalError("Cannot make void metric conditional"); + } + + public boolean isConditional() { + return false; + } }; public static DebugTimer timer(String name) { if (ENABLED) { - return new TimerImpl(name); + return new TimerImpl(name, true); } else { return VOID_TIMER; } @@ -417,5 +453,13 @@ public TimerCloseable start() { return TimerImpl.VOID_CLOSEABLE; } + + public void setConditional(boolean flag) { + throw new InternalError("Cannot make void timer conditional"); + } + + public boolean isConditional() { + return false; + } }; }
--- a/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/DebugHistogram.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/DebugHistogram.java Wed Oct 02 13:26:31 2013 +0200 @@ -22,15 +22,75 @@ */ package com.oracle.graal.debug; -import java.io.*; +import java.util.*; +/** + * Facility for recording value frequencies. + */ public interface DebugHistogram { + /** + * Gets the name specified when this objected was {@linkplain Debug#createHistogram(String) + * created}. + */ String getName(); + /** + * Increments the count for a given value. + */ void add(Object value); - void print(PrintStream os); + /** + * A value and a frequency. The ordering imposed by {@link #compareTo(CountedValue)} places + * values with higher frequencies first. + */ + public class CountedValue implements Comparable<CountedValue> { - void print(PrintStream os, int limit, int nameSize, int barSize); + private int count; + private final Object value; + + public CountedValue(int count, Object value) { + this.count = count; + this.value = value; + } + + public int compareTo(CountedValue o) { + if (count < o.count) { + return 1; + } else if (count > o.count) { + return -1; + } + return 0; + } + + @Override + public String toString() { + return count + " -> " + value; + } + + public void inc() { + count++; + } + + public int getCount() { + return count; + } + + public Object getValue() { + return value; + } + } + + /** + * Gets a list of the counted values, sorted in descending order of frequency. + */ + List<CountedValue> getValues(); + + /** + * Interface for a service that can render a visualization of a histogram. + */ + public interface Printer { + + void print(DebugHistogram histogram); + } }
--- a/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/DebugMetric.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/DebugMetric.java Wed Oct 02 13:26:31 2013 +0200 @@ -22,9 +22,32 @@ */ package com.oracle.graal.debug; +/** + * A counter for some value of interest. + */ public interface DebugMetric { + /** + * Adds 1 to this counter if metering is {@link Debug#isMeterEnabled() enabled} or this is an + * {@linkplain #isConditional() unconditional} metric. + */ void increment(); + /** + * Adds {@code value} to this counter if metering is {@link Debug#isMeterEnabled() enabled} or + * this is an {@linkplain #isConditional() unconditional} metric. + */ void add(long value); + + /** + * Sets a flag determining if this counter is only enabled if metering is + * {@link Debug#isMeterEnabled() enabled}. + */ + void setConditional(boolean flag); + + /** + * Determines if this counter is only enabled if metering is {@link Debug#isMeterEnabled() + * enabled}. + */ + boolean isConditional(); }
--- a/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/DebugTimer.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/DebugTimer.java Wed Oct 02 13:26:31 2013 +0200 @@ -24,7 +24,36 @@ import com.oracle.graal.debug.internal.*; +/** + * A timer for some activity of interest. A timer should be deployed using the try-with-resources + * pattern: + * + * <pre> + * try (TimerCloseable a = timer.start()) { + * // the code to time + * } + * </pre> + */ public interface DebugTimer { + /** + * Starts this timer if timing is {@linkplain Debug#isTimeEnabled() enabled} or this is an + * {@linkplain #isConditional() unconditional} timer. + * + * @return an object that must be closed once the activity has completed to add the elapsed time + * since this call to the total for this timer + */ TimerCloseable start(); + + /** + * Sets a flag determining if this timer is only enabled if metering is + * {@link Debug#isMeterEnabled() enabled}. + */ + void setConditional(boolean flag); + + /** + * Determines if this timer is only enabled if metering is {@link Debug#isMeterEnabled() + * enabled}. + */ + boolean isConditional(); }
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/internal/DebugHistogramAsciiPrinter.java Wed Oct 02 13:26:31 2013 +0200 @@ -0,0 +1,101 @@ +/* + * 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.debug.internal; + +import java.io.*; +import java.util.*; + +import com.oracle.graal.debug.*; +import com.oracle.graal.debug.DebugHistogram.*; + +/** + * Renders a textual representation of a histogram to a given print stream. + */ +public class DebugHistogramAsciiPrinter implements Printer { + + public static final int NumberSize = 10; + public static final int DefaultNameSize = 50; + public static final int DefaultBarSize = 100; + + private PrintStream os; + private int limit; + private int nameSize; + private int barSize; + + public DebugHistogramAsciiPrinter(PrintStream os) { + this(os, Integer.MAX_VALUE, DefaultNameSize, DefaultBarSize); + } + + /** + * @param os where to print + * @param limit limits printing to the {@code limit} most frequent values + * @param nameSize the width of the value names column + * @param barSize the width of the value frequency column + */ + public DebugHistogramAsciiPrinter(PrintStream os, int limit, int nameSize, int barSize) { + this.os = os; + this.limit = limit; + this.nameSize = nameSize; + this.barSize = barSize; + } + + public void print(DebugHistogram histogram) { + List<CountedValue> list = histogram.getValues(); + if (list.isEmpty()) { + os.printf("%s is empty.\n", histogram.getName()); + return; + } + + // Sum up the total number of elements. + int total = 0; + for (CountedValue cv : list) { + total += cv.getCount(); + } + + // Print header. + os.printf("%s has %d unique elements and %d total elements:\n", histogram.getName(), list.size(), total); + + int max = list.get(0).getCount(); + final int lineSize = nameSize + NumberSize + barSize + 10; + printLine(os, '-', lineSize); + String formatString = "| %-" + nameSize + "s | %-" + NumberSize + "d | %-" + barSize + "s |\n"; + for (int i = 0; i < list.size() && i < limit; ++i) { + CountedValue cv = list.get(i); + int value = cv.getCount(); + char[] bar = new char[(int) (((double) value / (double) max) * barSize)]; + Arrays.fill(bar, '='); + String objectString = String.valueOf(cv.getValue()); + if (objectString.length() > nameSize) { + objectString = objectString.substring(0, nameSize - 3) + "..."; + } + os.printf(formatString, objectString, value, new String(bar)); + } + printLine(os, '-', lineSize); + } + + private static void printLine(PrintStream printStream, char c, int lineSize) { + char[] charArr = new char[lineSize]; + Arrays.fill(charArr, c); + printStream.printf("%s\n", new String(charArr)); + } +}
--- a/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/internal/DebugHistogramImpl.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/internal/DebugHistogramImpl.java Wed Oct 02 13:26:31 2013 +0200 @@ -22,28 +22,25 @@ */ package com.oracle.graal.debug.internal; -import java.io.*; import java.util.*; import com.oracle.graal.debug.*; public class DebugHistogramImpl implements DebugHistogram { - public static final int NumberSize = 10; - public static final int DefaultNameSize = 50; - public static final int DefaultBarSize = 100; private final String name; - private HashMap<Object, Integer> map = new HashMap<>(); + private HashMap<Object, CountedValue> map = new HashMap<>(); public DebugHistogramImpl(String name) { this.name = name; } public void add(Object value) { - if (!map.containsKey(value)) { - map.put(value, 1); + CountedValue cv = map.get(value); + if (cv == null) { + map.put(value, new CountedValue(1, value)); } else { - map.put(value, map.get(value) + 1); + cv.inc(); } } @@ -52,59 +49,9 @@ return name; } - @Override - public void print(PrintStream os) { - print(os, Integer.MAX_VALUE, DefaultNameSize, DefaultBarSize); - } - - public void print(PrintStream os, int limit, int nameSize, int barSize) { - - List<Object> list = new ArrayList<>(map.keySet()); - if (list.size() == 0) { - // No elements in the histogram. - os.printf("%s is empty.\n", name); - return; - } - - // Sort from highest to smallest. - Collections.sort(list, new Comparator<Object>() { - - @Override - public int compare(Object o1, Object o2) { - return map.get(o2) - map.get(o1); - } - }); - - // Sum up the total number of elements. - int total = 0; - for (Object o : list) { - total += map.get(o); - } - - // Print header. - os.printf("%s has %d unique elements and %d total elements:\n", name, list.size(), total); - - int max = map.get(list.get(0)); - final int lineSize = nameSize + NumberSize + barSize + 10; - printLine(os, '-', lineSize); - String formatString = "| %-" + nameSize + "s | %-" + NumberSize + "d | %-" + barSize + "s |\n"; - for (int i = 0; i < list.size() && i < limit; ++i) { - Object o = list.get(i); - int value = map.get(o); - char[] bar = new char[(int) (((double) value / (double) max) * barSize)]; - Arrays.fill(bar, '='); - String objectString = o.toString(); - if (objectString.length() > nameSize) { - objectString = objectString.substring(0, nameSize - 3) + "..."; - } - os.printf(formatString, objectString, value, new String(bar)); - } - printLine(os, '-', lineSize); - } - - private static void printLine(PrintStream printStream, char c, int lineSize) { - char[] charArr = new char[lineSize]; - Arrays.fill(charArr, c); - printStream.printf("%s\n", new String(charArr)); + public List<CountedValue> getValues() { + ArrayList<CountedValue> res = new ArrayList<>(map.values()); + Collections.sort(res); + return res; } }
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/internal/DebugHistogramRPrinter.java Wed Oct 02 13:26:31 2013 +0200 @@ -0,0 +1,81 @@ +/* + * 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.debug.internal; + +import java.io.*; +import java.util.*; + +import com.oracle.graal.debug.*; +import com.oracle.graal.debug.DebugHistogram.CountedValue; +import com.oracle.graal.debug.DebugHistogram.Printer; + +/** + * Renders a histogram as an R script to a given print stream. The R script emitted for a histogram + * is a simple set of statements for defining a vector of named objects. + */ +public class DebugHistogramRPrinter implements Printer { + + private PrintStream os; + private int limit; + + public DebugHistogramRPrinter(PrintStream os) { + this(os, Integer.MAX_VALUE); + } + + /** + * @param os where to print + * @param limit limits printing to the {@code limit} most frequent values + */ + public DebugHistogramRPrinter(PrintStream os, int limit) { + this.os = os; + this.limit = limit; + } + + public void print(DebugHistogram histogram) { + List<CountedValue> list = histogram.getValues(); + if (list.isEmpty()) { + return; + } + + String var = histogram.getName().replace('-', '.').replace(' ', '_'); + os.print(var + " <- c("); + for (int i = 0; i < list.size() && i < limit; ++i) { + CountedValue cv = list.get(i); + if (i != 0) { + os.print(", "); + } + os.print(cv.getCount()); + } + os.println(");"); + + os.print("names(" + var + ") <- c("); + for (int i = 0; i < list.size() && i < limit; ++i) { + CountedValue cv = list.get(i); + if (i != 0) { + os.print(", "); + } + os.print("\"" + cv.getValue() + "\""); + } + os.println(");"); + } +}
--- a/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/internal/DebugScope.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/internal/DebugScope.java Wed Oct 02 13:26:31 2013 +0200 @@ -92,7 +92,6 @@ private static ThreadLocal<DebugScope> instanceTL = new ThreadLocal<>(); private static ThreadLocal<DebugConfig> configTL = new ThreadLocal<>(); private static ThreadLocal<Throwable> lastExceptionThrownTL = new ThreadLocal<>(); - private static DebugTimer scopeTime = Debug.timer("ScopeTime"); private final DebugScope parent; private IndentImpl lastUsedIndent; @@ -210,6 +209,9 @@ for (DebugDumpHandler dumpHandler : config.dumpHandlers()) { dumpHandler.dump(object, message); } + } else { + PrintStream out = System.out; + out.println("Forced dump ignored because debugging is disabled - use -G:Dump=xxx option"); } } @@ -239,7 +241,7 @@ } instanceTL.set(newChild); newChild.updateFlags(); - try (TimerCloseable a = scopeTime.start()) { + try { return executeScope(runnable, callable); } finally { newChild.context = null;
--- a/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/internal/DebugValue.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/internal/DebugValue.java Wed Oct 02 13:26:31 2013 +0200 @@ -30,10 +30,12 @@ private final String name; private int index; + private boolean conditional; - protected DebugValue(String name) { + protected DebugValue(String name, boolean conditional) { this.name = name; this.index = -1; + this.conditional = conditional; } protected long getCurrentValue() { @@ -46,6 +48,14 @@ DebugScope.getInstance().setCurrentValue(index, l); } + public void setConditional(boolean flag) { + conditional = flag; + } + + public boolean isConditional() { + return conditional; + } + private void ensureInitialized() { if (index == -1) { index = KeyRegistry.register(this);
--- a/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/internal/MetricImpl.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/internal/MetricImpl.java Wed Oct 02 13:26:31 2013 +0200 @@ -26,8 +26,8 @@ public final class MetricImpl extends DebugValue implements DebugMetric { - public MetricImpl(String name) { - super(name); + public MetricImpl(String name, boolean conditional) { + super(name, conditional); } public void increment() { @@ -35,7 +35,7 @@ } public void add(long value) { - if (Debug.isMeterEnabled()) { + if (!isConditional() || Debug.isMeterEnabled()) { super.addToCurrentValue(value); } }
--- a/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/internal/TimerCloseable.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/internal/TimerCloseable.java Wed Oct 02 13:26:31 2013 +0200 @@ -22,6 +22,12 @@ */ package com.oracle.graal.debug.internal; +import com.oracle.graal.debug.*; + +/** + * An object returned by {@link DebugTimer#start()} that when closed, stops the associated timer and + * adds the elapsed time since {@code start()} to the total for the timer. + */ public interface TimerCloseable extends AutoCloseable { void close();
--- a/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/internal/TimerImpl.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.debug/src/com/oracle/graal/debug/internal/TimerImpl.java Wed Oct 02 13:26:31 2013 +0200 @@ -37,59 +37,78 @@ } }; - private ThreadLocal<Long> valueToSubstract = new ThreadLocal<>(); + /** + * Records the most recent active timer. + */ + private static ThreadLocal<AbstractTimer> currentTimer = new ThreadLocal<>(); - public TimerImpl(String name) { - super(name); + private final DebugValue flat; + + public TimerImpl(String name, boolean conditional) { + super(name + "_Accm", conditional); + this.flat = new DebugValue(name + "_Flat", conditional) { + + @Override + public String toString(long value) { + return valueToString(value); + } + }; } @Override public TimerCloseable start() { - if (Debug.isTimeEnabled()) { + if (!isConditional() || Debug.isTimeEnabled()) { long startTime; if (threadMXBean.isCurrentThreadCpuTimeSupported()) { startTime = threadMXBean.getCurrentThreadCpuTime(); } else { startTime = System.nanoTime(); } - if (valueToSubstract.get() == null) { - valueToSubstract.set(0L); - } - long previousValueToSubstract = valueToSubstract.get(); + AbstractTimer result; if (threadMXBean.isCurrentThreadCpuTimeSupported()) { - result = new CpuTimeTimer(startTime, previousValueToSubstract); + result = new CpuTimeTimer(startTime); } else { - result = new SystemNanosTimer(startTime, previousValueToSubstract); + result = new SystemNanosTimer(startTime); } - valueToSubstract.set(0L); + currentTimer.set(result); return result; } else { return VOID_CLOSEABLE; } } + public static String valueToString(long value) { + return String.format("%d.%d ms", value / 1000000, (value / 100000) % 10); + } + @Override public String toString(long value) { - return String.format("%d.%d ms", value / 1000000, (value / 100000) % 10); + return valueToString(value); } private abstract class AbstractTimer implements TimerCloseable { + private final AbstractTimer parent; private final long startTime; - private final long previousValueToSubstract; + private long nestedTimeToSubtract; - private AbstractTimer(long startTime, long previousValueToSubstract) { + private AbstractTimer(long startTime) { + this.parent = currentTimer.get(); this.startTime = startTime; - this.previousValueToSubstract = previousValueToSubstract; } @Override public void close() { - long timeSpan = currentTime() - startTime; - long oldValueToSubstract = valueToSubstract.get(); - valueToSubstract.set(timeSpan + previousValueToSubstract); - TimerImpl.this.addToCurrentValue(timeSpan - oldValueToSubstract); + long endTime = currentTime(); + long timeSpan = endTime - startTime; + if (parent != null) { + parent.nestedTimeToSubtract += timeSpan; + } + currentTimer.set(parent); + long flatTime = timeSpan - nestedTimeToSubtract; + TimerImpl.this.addToCurrentValue(timeSpan); + flat.addToCurrentValue(flatTime); } protected abstract long currentTime(); @@ -97,8 +116,8 @@ private final class SystemNanosTimer extends AbstractTimer { - public SystemNanosTimer(long startTime, long previousValueToSubstract) { - super(startTime, previousValueToSubstract); + public SystemNanosTimer(long startTime) { + super(startTime); } @Override @@ -109,8 +128,8 @@ private final class CpuTimeTimer extends AbstractTimer { - public CpuTimeTimer(long startTime, long previousValueToSubstract) { - super(startTime, previousValueToSubstract); + public CpuTimeTimer(long startTime) { + super(startTime); } @Override
--- a/graal/com.oracle.graal.graph.test/src/com/oracle/graal/graph/test/TypedNodeIteratorTest.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.graph.test/src/com/oracle/graal/graph/test/TypedNodeIteratorTest.java Wed Oct 02 13:26:31 2013 +0200 @@ -32,7 +32,7 @@ public class TypedNodeIteratorTest { - private static class TestNode extends Node implements Node.IterableNodeType, TestNodeInterface { + private static class TestNode extends Node implements IterableNodeType, TestNodeInterface { private final String name;
--- a/graal/com.oracle.graal.graph.test/src/com/oracle/graal/graph/test/TypedNodeIteratorTest2.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.graph.test/src/com/oracle/graal/graph/test/TypedNodeIteratorTest2.java Wed Oct 02 13:26:31 2013 +0200 @@ -43,7 +43,7 @@ } } - private static class NodeB extends NodeA implements Node.IterableNodeType { + private static class NodeB extends NodeA implements IterableNodeType { public NodeB(String name) { super(name);
--- a/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/Graph.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/Graph.java Wed Oct 02 13:26:31 2013 +0200 @@ -25,7 +25,6 @@ import java.util.*; import com.oracle.graal.graph.GraphEvent.NodeEvent; -import com.oracle.graal.graph.Node.IterableNodeType; import com.oracle.graal.graph.Node.ValueNumberable; import com.oracle.graal.graph.iterators.*; @@ -57,8 +56,8 @@ private int deletedNodeCount; private GraphEventLog eventLog; - NodeChangedListener inputChanged; - NodeChangedListener usagesDroppedZero; + NodeChangedListener inputChangedListener; + NodeChangedListener usagesDroppedToZeroListener; private final HashMap<CacheEntry, Node> cachedNodes = new HashMap<>(); private static final class CacheEntry { @@ -66,6 +65,8 @@ private final Node node; public CacheEntry(Node node) { + assert node.getNodeClass().valueNumberable(); + assert node.getNodeClass().isLeafNode(); this.node = node; } @@ -82,8 +83,8 @@ if (obj instanceof CacheEntry) { CacheEntry other = (CacheEntry) obj; NodeClass nodeClass = node.getNodeClass(); - if (other.node.getNodeClass() == nodeClass) { - return nodeClass.valueNumberable() && nodeClass.valueEqual(node, other.node) && nodeClass.edgesEqual(node, other.node); + if (other.node.getClass() == node.getClass()) { + return nodeClass.valueEqual(node, other.node); } } return false; @@ -156,7 +157,7 @@ int usageModCount(Node node) { int id = extractOriginalNodeId(node); if (id >= 0 && id < nodeUsageModCounts.length) { - return nodeUsageModCounts[node.id]; + return nodeUsageModCounts[id]; } return 0; } @@ -187,8 +188,7 @@ */ public Graph copy(String newName) { Graph copy = new Graph(newName); - Map<Node, Node> emptyMap = Collections.emptyMap(); - copy.addDuplicates(getNodes(), emptyMap); + copy.addDuplicates(getNodes(), this, this.getNodeCount(), (Map<Node, Node>) null); return copy; } @@ -223,6 +223,24 @@ * @return the node which was added to the graph */ public <T extends Node> T add(T node) { + if (node.getNodeClass().valueNumberable()) { + throw new IllegalStateException("Using add for value numberable node. Consider using either unique or addWithoutUnique."); + } + return addHelper(node); + } + + public <T extends Node> T addWithoutUnique(T node) { + return addHelper(node); + } + + public <T extends Node> T addOrUnique(T node) { + if (node.getNodeClass().valueNumberable()) { + return uniqueHelper(node); + } + return add(node); + } + + private <T extends Node> T addHelper(T node) { node.initialize(this); return node; } @@ -232,24 +250,54 @@ void nodeChanged(Node node); } - public void trackInputChange(NodeChangedListener inputChangedListener) { - assert this.inputChanged == null; - this.inputChanged = inputChangedListener; + private static class ChainedNodeChangedListener implements NodeChangedListener { + + NodeChangedListener head; + NodeChangedListener next; + + ChainedNodeChangedListener(NodeChangedListener head, NodeChangedListener next) { + this.head = head; + this.next = next; + } + + public void nodeChanged(Node node) { + head.nodeChanged(node); + next.nodeChanged(node); + } + } + + public void trackInputChange(NodeChangedListener listener) { + if (inputChangedListener == null) { + inputChangedListener = listener; + } else { + inputChangedListener = new ChainedNodeChangedListener(listener, inputChangedListener); + } } public void stopTrackingInputChange() { - assert inputChanged != null; - inputChanged = null; + assert inputChangedListener != null; + if (inputChangedListener instanceof ChainedNodeChangedListener) { + inputChangedListener = ((ChainedNodeChangedListener) inputChangedListener).next; + } else { + inputChangedListener = null; + } } - public void trackUsagesDroppedZero(NodeChangedListener usagesDroppedZeroListener) { - assert this.usagesDroppedZero == null; - this.usagesDroppedZero = usagesDroppedZeroListener; + public void trackUsagesDroppedZero(NodeChangedListener listener) { + if (usagesDroppedToZeroListener == null) { + usagesDroppedToZeroListener = listener; + } else { + usagesDroppedToZeroListener = new ChainedNodeChangedListener(listener, usagesDroppedToZeroListener); + } } public void stopTrackingUsagesDroppedZero() { - assert usagesDroppedZero != null; - usagesDroppedZero = null; + assert usagesDroppedToZeroListener != null; + if (usagesDroppedToZeroListener instanceof ChainedNodeChangedListener) { + usagesDroppedToZeroListener = ((ChainedNodeChangedListener) usagesDroppedToZeroListener).next; + } else { + usagesDroppedToZeroListener = null; + } } /** @@ -261,53 +309,76 @@ * @return the node which was added to the graph or a <i>similar</i> which was already in the * graph. */ - @SuppressWarnings("unchecked") public <T extends Node & ValueNumberable> T unique(T node) { assert checkValueNumberable(node); + return uniqueHelper(node); + } - for (Node input : node.inputs()) { - if (input != null) { - for (Node usage : input.usages()) { - if (usage != node && node.getNodeClass().valueEqual(node, usage) && node.getNodeClass().edgesEqual(node, usage)) { - return (T) usage; - } - } - return add(node); + @SuppressWarnings("unchecked") + <T extends Node> T uniqueHelper(T node) { + assert node.getNodeClass().valueNumberable(); + Node other = this.findDuplicate(node); + if (other != null) { + return (T) other; + } else { + Node result = addHelper(node); + if (node.getNodeClass().isLeafNode()) { + putNodeIntoCache(result); } - } - Node cachedNode = cachedNodes.get(new CacheEntry(node)); - if (cachedNode != null && cachedNode.isAlive()) { - return (T) cachedNode; - } else { - Node result = add(node); - cachedNodes.put(new CacheEntry(node), result); return (T) result; } } + void putNodeIntoCache(Node node) { + assert node.graph() == this || node.graph() == null; + assert node.getNodeClass().valueNumberable(); + assert node.getNodeClass().isLeafNode() : node.getClass(); + cachedNodes.put(new CacheEntry(node), node); + } + + Node findNodeInCache(Node node) { + CacheEntry key = new CacheEntry(node); + Node result = cachedNodes.get(key); + if (result != null && result.isDeleted()) { + cachedNodes.remove(key); + return null; + } + return result; + } + public Node findDuplicate(Node node) { - if (node.getNodeClass().valueNumberable()) { + NodeClass nodeClass = node.getNodeClass(); + assert nodeClass.valueNumberable(); + if (nodeClass.isLeafNode()) { + Node cachedNode = findNodeInCache(node); + if (cachedNode != null) { + return cachedNode; + } else { + return null; + } + } else { + + int minCount = Integer.MAX_VALUE; + Node minCountNode = null; for (Node input : node.inputs()) { if (input != null) { - for (Node usage : input.usages()) { - if (usage != node && node.getNodeClass().valueEqual(node, usage) && node.getNodeClass().edgesEqual(node, usage)) { - return usage; - } + int estimate = input.getUsageCountUpperBound(); + if (estimate == 0) { + return null; + } else if (estimate < minCount) { + minCount = estimate; + minCountNode = input; } - return null; } } - } - CacheEntry key = new CacheEntry(node); - Node cachedNode = cachedNodes.get(key); - if (cachedNode != null) { - if (!cachedNode.isAlive()) { - cachedNodes.remove(key); + if (minCountNode != null) { + for (Node usage : minCountNode.usages()) { + if (usage != node && nodeClass == usage.getNodeClass() && nodeClass.valueEqual(node, usage) && nodeClass.edgesEqual(node, usage)) { + return usage; + } + } return null; } - return cachedNode != node ? cachedNode : null; - } else { - cachedNodes.put(key, node); return null; } } @@ -319,6 +390,10 @@ return true; } + public boolean isNew(int mark, Node node) { + return node.id >= mark; + } + /** * Gets a mark that can be used with {@link #getNewNodes(int)}. */ @@ -518,7 +593,7 @@ * {@code type}. * * @param type the type of node to return - * @return an {@link Iterable} providing all the matching nodes. + * @return an {@link Iterable} providing all the matching nodes */ public <T extends Node & IterableNodeType> NodeIterable<T> getNodes(final Class<T> type) { final NodeClass nodeClass = NodeClass.get(type); @@ -671,14 +746,14 @@ * @param replacementsMap the replacement map (can be null if no replacement is to be performed) * @return a map which associates the original nodes from {@code nodes} to their duplicates */ - public Map<Node, Node> addDuplicates(Iterable<Node> newNodes, Map<Node, Node> replacementsMap) { + public Map<Node, Node> addDuplicates(Iterable<Node> newNodes, final Graph oldGraph, int estimatedNodeCount, Map<Node, Node> replacementsMap) { DuplicationReplacement replacements; if (replacementsMap == null) { replacements = null; } else { replacements = new MapReplacement(replacementsMap); } - return addDuplicates(newNodes, replacements); + return addDuplicates(newNodes, oldGraph, estimatedNodeCount, replacements); } public interface DuplicationReplacement { @@ -702,19 +777,8 @@ } - private static final DuplicationReplacement NO_REPLACEMENT = new DuplicationReplacement() { - - @Override - public Node replacement(Node original) { - return original; - } - }; - @SuppressWarnings("all") - public Map<Node, Node> addDuplicates(Iterable<Node> newNodes, DuplicationReplacement replacements) { - if (replacements == null) { - replacements = NO_REPLACEMENT; - } - return NodeClass.addGraphDuplicate(this, newNodes, replacements); + public Map<Node, Node> addDuplicates(Iterable<Node> newNodes, final Graph oldGraph, int estimatedNodeCount, DuplicationReplacement replacements) { + return NodeClass.addGraphDuplicate(this, oldGraph, estimatedNodeCount, newNodes, replacements); } }
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/IterableNodeType.java Wed Oct 02 13:26:31 2013 +0200 @@ -0,0 +1,32 @@ +/* + * 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.graph; + +/** + * A marker for a node type supporting {@linkplain Graph#getNodes(Class) fast iteration} of its + * instances in a graph. The support for fast iteration comes with a memory cost (e.g., extra data + * structures {@link Graph}) so only node types for which fast iteration provides a compilation + * performance benefit should implement this interface. + */ +public interface IterableNodeType { +}
--- a/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/Node.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/Node.java Wed Oct 02 13:26:31 2013 +0200 @@ -104,6 +104,9 @@ /** * Determines if the stamp of the instantiated intrinsic node has its stamp set from the * return type of the annotated method. + * <p> + * When it is set to true, the stamp that is passed in to the constructor of ValueNode is + * ignored and can therefore safely be {@code null}. */ boolean setStampFromReturnType() default false; } @@ -111,9 +114,6 @@ public interface ValueNumberable { } - public interface IterableNodeType { - } - private Graph graph; int id; @@ -240,6 +240,16 @@ } } + int getUsageCountUpperBound() { + if (usage0 == null) { + return 0; + } + if (usage1 == null) { + return 1; + } + return 2 + extraUsages.length; + } + /** * Gets the list of nodes that use this node (e.g., as an input). */ @@ -423,15 +433,15 @@ assert assertTrue(result, "not found in usages, old input: %s", oldInput); } if (newInput != null) { - NodeChangedListener inputChanged = graph.inputChanged; - if (inputChanged != null) { - inputChanged.nodeChanged(this); + NodeChangedListener listener = graph.inputChangedListener; + if (listener != null) { + listener.nodeChanged(this); } newInput.addUsage(this); } else if (oldInput != null && oldInput.usages().isEmpty()) { - NodeChangedListener nodeChangedListener = graph.usagesDroppedZero; - if (nodeChangedListener != null) { - nodeChangedListener.nodeChanged(oldInput); + NodeChangedListener listener = graph.usagesDroppedToZeroListener; + if (listener != null) { + listener.nodeChanged(oldInput); } } } @@ -449,7 +459,7 @@ oldSuccessor.predecessor = null; } if (newSuccessor != null) { - assert assertTrue(newSuccessor.predecessor == null, "unexpected non-null predecessor in new successor (%s): %s", newSuccessor, newSuccessor.predecessor); + assert assertTrue(newSuccessor.predecessor == null, "unexpected non-null predecessor in new successor (%s): %s, this=%s", newSuccessor, newSuccessor.predecessor, this); newSuccessor.predecessor = this; } } @@ -485,9 +495,9 @@ boolean result = usage.getNodeClass().replaceFirstInput(usage, this, other); assert assertTrue(result, "not found in inputs, usage: %s", usage); if (other != null) { - NodeChangedListener inputChanged = graph.inputChanged; - if (inputChanged != null) { - inputChanged.nodeChanged(usage); + NodeChangedListener listener = graph.inputChangedListener; + if (listener != null) { + listener.nodeChanged(usage); } other.addUsage(usage); } @@ -532,9 +542,9 @@ for (Node input : inputs()) { removeThisFromUsages(input); if (input.usages().isEmpty()) { - NodeChangedListener nodeChangedListener = graph.usagesDroppedZero; - if (nodeChangedListener != null) { - nodeChangedListener.nodeChanged(input); + NodeChangedListener listener = graph.usagesDroppedToZeroListener; + if (listener != null) { + listener.nodeChanged(input); } } } @@ -584,15 +594,31 @@ return newNode; } - public Node clone(Graph into) { + static int count = 0; + + public final Node clone(Graph into) { + return clone(into, true); + } + + final Node clone(Graph into, boolean clearInputsAndSuccessors) { + NodeClass nodeClass = getNodeClass(); + if (nodeClass.valueNumberable() && nodeClass.isLeafNode()) { + Node otherNode = into.findNodeInCache(this); + if (otherNode != null) { + return otherNode; + } + } + Node newNode = null; try { newNode = (Node) this.clone(); } catch (CloneNotSupportedException e) { throw new GraalInternalError(e).addContext(this); } - getNodeClass().clearInputs(newNode); - getNodeClass().clearSuccessors(newNode); + if (clearInputsAndSuccessors) { + nodeClass.clearInputs(newNode); + nodeClass.clearSuccessors(newNode); + } newNode.graph = into; newNode.typeCacheNext = null; newNode.id = INITIAL_ID; @@ -601,9 +627,17 @@ newNode.usage1 = null; newNode.extraUsages = NO_NODES; newNode.predecessor = null; + + if (nodeClass.valueNumberable() && nodeClass.isLeafNode()) { + into.putNodeIntoCache(newNode); + } + newNode.afterClone(this); return newNode; } + protected void afterClone(@SuppressWarnings("unused") Node other) { + } + public boolean verify() { assertTrue(isAlive(), "cannot verify inactive nodes (id=%d)", id); assertTrue(graph() != null, "null graph"); @@ -779,6 +813,7 @@ boolean neighborsAlternate = ((flags & FormattableFlags.LEFT_JUSTIFY) == FormattableFlags.LEFT_JUSTIFY); int neighborsFlags = (neighborsAlternate ? FormattableFlags.ALTERNATE | FormattableFlags.LEFT_JUSTIFY : 0); + NodeClass nodeClass = getNodeClass(); if (width > 0) { if (this.predecessor != null) { formatter.format(" pred={"); @@ -789,10 +824,10 @@ NodeClassIterator inputIter = inputs().iterator(); while (inputIter.hasNext()) { Position position = inputIter.nextPosition(); - Node input = getNodeClass().get(this, position); + Node input = nodeClass.get(this, position); if (input != null) { formatter.format(" "); - formatter.format(getNodeClass().getName(position)); + formatter.format(nodeClass.getName(position)); formatter.format("={"); input.formatTo(formatter, neighborsFlags, width - 1, 0); formatter.format("}"); @@ -817,10 +852,10 @@ NodeClassIterator succIter = successors().iterator(); while (succIter.hasNext()) { Position position = succIter.nextPosition(); - Node successor = getNodeClass().get(this, position); + Node successor = nodeClass.get(this, position); if (successor != null) { formatter.format(" "); - formatter.format(getNodeClass().getName(position)); + formatter.format(nodeClass.getName(position)); formatter.format("={"); successor.formatTo(formatter, neighborsFlags, 0, precision - 1); formatter.format("}");
--- a/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeClass.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeClass.java Wed Oct 02 13:26:31 2013 +0200 @@ -26,12 +26,11 @@ import java.lang.reflect.*; import java.util.*; -import java.util.Map.Entry; import java.util.concurrent.*; +import com.oracle.graal.debug.*; import com.oracle.graal.graph.Graph.DuplicationReplacement; import com.oracle.graal.graph.Node.Input; -import com.oracle.graal.graph.Node.IterableNodeType; import com.oracle.graal.graph.Node.Successor; import com.oracle.graal.graph.Node.Verbosity; @@ -139,12 +138,16 @@ private final long[] successorOffsets; private final Class<?>[] dataTypes; private final boolean canGVN; + private final boolean isLeafNode; private final int startGVNNumber; private final String shortName; private final String nameTemplate; private final int iterableId; private int[] iterableIds; + private static final DebugMetric ITERABLE_NODE_TYPES = Debug.metric("IterableNodeTypes"); + private final DebugMetric nodeIterableCount; + private NodeClass(Class<?> clazz) { super(clazz); assert NODE_CLASS.isAssignableFrom(clazz); @@ -185,14 +188,15 @@ } this.nameTemplate = newNameTemplate == null ? newShortName : newNameTemplate; this.shortName = newShortName; - if (Node.IterableNodeType.class.isAssignableFrom(clazz)) { + if (IterableNodeType.class.isAssignableFrom(clazz)) { + ITERABLE_NODE_TYPES.increment(); this.iterableId = nextIterableId++; List<NodeClass> existingClasses = new LinkedList<>(); for (FieldIntrospection nodeClass : allClasses.values()) { if (clazz.isAssignableFrom(nodeClass.clazz)) { existingClasses.add((NodeClass) nodeClass); } - if (nodeClass.clazz.isAssignableFrom(clazz) && Node.IterableNodeType.class.isAssignableFrom(nodeClass.clazz)) { + if (nodeClass.clazz.isAssignableFrom(clazz) && IterableNodeType.class.isAssignableFrom(nodeClass.clazz)) { NodeClass superNodeClass = (NodeClass) nodeClass; superNodeClass.iterableIds = Arrays.copyOf(superNodeClass.iterableIds, superNodeClass.iterableIds.length + 1); superNodeClass.iterableIds[superNodeClass.iterableIds.length - 1] = this.iterableId; @@ -209,6 +213,9 @@ this.iterableId = NOT_ITERABLE; this.iterableIds = null; } + + isLeafNode = (this.inputOffsets.length == 0 && this.successorOffsets.length == 0); + nodeIterableCount = Debug.metric("NodeIterable_" + shortName); } @Override @@ -236,6 +243,7 @@ } public int[] iterableIds() { + nodeIterableCount.increment(); return iterableIds; } @@ -247,6 +255,10 @@ return canGVN; } + public boolean isLeafNode() { + return isLeafNode; + } + public static int cacheSize() { return nextIterableId; } @@ -417,6 +429,7 @@ */ public static final class NodeClassIterator implements Iterator<Node> { + private final NodeClass nodeClass; private final Node node; private final int modCount; private final int directCount; @@ -434,6 +447,7 @@ */ private NodeClassIterator(Node node, long[] offsets, int directCount) { this.node = node; + this.nodeClass = node.getNodeClass(); this.modCount = MODIFICATION_COUNTS_ENABLED ? node.modCount() : 0; this.offsets = offsets; this.directCount = directCount; @@ -500,9 +514,9 @@ public Position nextPosition() { try { if (index < directCount) { - return new Position(offsets == node.getNodeClass().inputOffsets, index, NOT_ITERABLE); + return new Position(offsets == nodeClass.inputOffsets, index, NOT_ITERABLE); } else { - return new Position(offsets == node.getNodeClass().inputOffsets, index, subIndex); + return new Position(offsets == nodeClass.inputOffsets, index, subIndex); } } finally { forward(); @@ -581,7 +595,7 @@ } public boolean valueEqual(Node a, Node b) { - if (!canGVN || a.getNodeClass() != b.getNodeClass()) { + if (!canGVN || a.getClass() != b.getClass()) { return a == b; } for (int i = 0; i < dataOffsets.length; ++i) { @@ -659,6 +673,84 @@ return fieldNames.get(pos.input ? inputOffsets[pos.index] : successorOffsets[pos.index]); } + void updateInputSuccInPlace(Node node, InplaceUpdateClosure duplicationReplacement) { + int index = 0; + while (index < directInputCount) { + Node input = getNode(node, inputOffsets[index]); + if (input != null) { + Node newInput = duplicationReplacement.replacement(input, true); + node.updateUsages(null, newInput); + putNode(node, inputOffsets[index], newInput); + } + index++; + } + + if (index < inputOffsets.length) { + updateInputLists(node, duplicationReplacement, index); + } + + index = 0; + while (index < directSuccessorCount) { + Node successor = getNode(node, successorOffsets[index]); + if (successor != null) { + Node newSucc = duplicationReplacement.replacement(successor, false); + node.updatePredecessor(null, newSucc); + putNode(node, successorOffsets[index], newSucc); + } + index++; + } + + if (index < successorOffsets.length) { + updateSuccLists(node, duplicationReplacement, index); + } + } + + private void updateInputLists(Node node, InplaceUpdateClosure duplicationReplacement, int startIndex) { + int index = startIndex; + while (index < inputOffsets.length) { + NodeList<Node> list = getNodeList(node, inputOffsets[index]); + putNodeList(node, inputOffsets[index], updateInputListCopy(list, node, duplicationReplacement)); + assert list != null : clazz; + index++; + } + } + + private void updateSuccLists(Node node, InplaceUpdateClosure duplicationReplacement, int startIndex) { + int index = startIndex; + while (index < successorOffsets.length) { + NodeList<Node> list = getNodeList(node, successorOffsets[index]); + putNodeList(node, successorOffsets[index], updateSuccListCopy(list, node, duplicationReplacement)); + assert list != null : clazz; + index++; + } + } + + private static NodeInputList<Node> updateInputListCopy(NodeList<Node> list, Node node, InplaceUpdateClosure duplicationReplacement) { + int size = list.size(); + NodeInputList<Node> result = new NodeInputList<>(node, size); + for (int i = 0; i < list.count(); ++i) { + Node oldNode = list.get(i); + if (oldNode != null) { + Node newNode = duplicationReplacement.replacement(oldNode, true); + result.set(i, newNode); + } + } + return result; + } + + private static NodeSuccessorList<Node> updateSuccListCopy(NodeList<Node> list, Node node, InplaceUpdateClosure duplicationReplacement) { + int size = list.size(); + NodeSuccessorList<Node> result = new NodeSuccessorList<>(node, size); + for (int i = 0; i < list.count(); ++i) { + Node oldNode = list.get(i); + if (oldNode != null) { + Node newNode = duplicationReplacement.replacement(oldNode, false); + result.set(i, newNode); + } + } + return result; + } + public void set(Node node, Position pos, Node x) { long offset = pos.input ? inputOffsets[pos.index] : successorOffsets[pos.index]; if (pos.subIndex == NOT_ITERABLE) { @@ -955,78 +1047,113 @@ return nameTemplate; } - static Map<Node, Node> addGraphDuplicate(Graph graph, Iterable<Node> nodes, DuplicationReplacement replacements) { - Map<Node, Node> newNodes = new IdentityHashMap<>(); - Map<Node, Node> replacementsMap = new IdentityHashMap<>(); - // create node duplicates + interface InplaceUpdateClosure { + + Node replacement(Node node, boolean isInput); + } + + static Map<Node, Node> addGraphDuplicate(final Graph graph, final Graph oldGraph, int estimatedNodeCount, Iterable<Node> nodes, final DuplicationReplacement replacements) { + final Map<Node, Node> newNodes = (estimatedNodeCount > (oldGraph.getNodeCount() + oldGraph.getDeletedNodeCount() >> 4)) ? new NodeNodeMap(oldGraph) : new IdentityHashMap<Node, Node>(); + createNodeDuplicates(graph, nodes, replacements, newNodes); + + InplaceUpdateClosure replacementClosure = new InplaceUpdateClosure() { + + public Node replacement(Node node, boolean isInput) { + Node target = newNodes.get(node); + if (target == null) { + Node replacement = node; + if (replacements != null) { + replacement = replacements.replacement(node); + } + if (replacement != node) { + target = replacement; + } else if (node.graph() == graph && isInput) { + // patch to the outer world + target = node; + } + + } + return target; + } + + }; + + // re-wire inputs + for (Node oldNode : nodes) { + Node node = newNodes.get(oldNode); + NodeClass oldNodeClass = oldNode.getNodeClass(); + NodeClass nodeClass = node.getNodeClass(); + if (replacements == null || replacements.replacement(oldNode) == oldNode) { + nodeClass.updateInputSuccInPlace(node, replacementClosure); + } else { + transferValuesDifferentNodeClass(graph, replacements, newNodes, oldNode, node, oldNodeClass, nodeClass); + } + } + + return newNodes; + } + + private static void createNodeDuplicates(final Graph graph, Iterable<Node> nodes, final DuplicationReplacement replacements, final Map<Node, Node> newNodes) { for (Node node : nodes) { if (node != null) { assert !node.isDeleted() : "trying to duplicate deleted node: " + node; - Node replacement = replacements.replacement(node); + Node replacement = node; + if (replacements != null) { + replacement = replacements.replacement(node); + } if (replacement != node) { assert replacement != null; newNodes.put(node, replacement); } else { - Node newNode = node.clone(graph); + Node newNode = node.clone(graph, false); + assert newNode.usages().count() == 0 || newNode.inputs().count() == 0; assert newNode.getClass() == node.getClass(); newNodes.put(node, newNode); } } } - // re-wire inputs - for (Entry<Node, Node> entry : newNodes.entrySet()) { - Node oldNode = entry.getKey(); - Node node = entry.getValue(); - for (NodeClassIterator iter = oldNode.inputs().iterator(); iter.hasNext();) { - Position pos = iter.nextPosition(); - if (!pos.isValidFor(node, oldNode)) { - continue; + } + + private static void transferValuesDifferentNodeClass(final Graph graph, final DuplicationReplacement replacements, final Map<Node, Node> newNodes, Node oldNode, Node node, NodeClass oldNodeClass, + NodeClass nodeClass) { + for (NodeClassIterator iter = oldNode.inputs().iterator(); iter.hasNext();) { + Position pos = iter.nextPosition(); + if (!nodeClass.isValid(pos, oldNodeClass)) { + continue; + } + Node input = oldNodeClass.get(oldNode, pos); + Node target = newNodes.get(input); + if (target == null) { + Node replacement = input; + if (replacements != null) { + replacement = replacements.replacement(input); } - Node input = oldNode.getNodeClass().get(oldNode, pos); - Node target = newNodes.get(input); - if (target == null) { - target = replacementsMap.get(input); - if (target == null) { - Node replacement = replacements.replacement(input); - if (replacement != input) { - replacementsMap.put(input, replacement); - assert isAssignable(node.getNodeClass().fieldTypes.get(node.getNodeClass().inputOffsets[pos.index]), replacement); - target = replacement; - } else if (input.graph() == graph) { // patch to the outer world - target = input; - } - } + if (replacement != input) { + assert isAssignable(nodeClass.fieldTypes.get(nodeClass.inputOffsets[pos.index]), replacement); + target = replacement; + } else if (input.graph() == graph) { // patch to the outer world + target = input; } - node.getNodeClass().set(node, pos, target); } + nodeClass.set(node, pos, target); } - // re-wire successors - for (Entry<Node, Node> entry : newNodes.entrySet()) { - Node oldNode = entry.getKey(); - Node node = entry.getValue(); - for (NodeClassIterator iter = oldNode.successors().iterator(); iter.hasNext();) { - Position pos = iter.nextPosition(); - if (!pos.isValidFor(node, oldNode)) { - continue; + for (NodeClassIterator iter = oldNode.successors().iterator(); iter.hasNext();) { + Position pos = iter.nextPosition(); + if (!nodeClass.isValid(pos, oldNodeClass)) { + continue; + } + Node succ = oldNodeClass.get(oldNode, pos); + Node target = newNodes.get(succ); + if (target == null) { + Node replacement = replacements.replacement(succ); + if (replacement != succ) { + assert isAssignable(nodeClass.fieldTypes.get(node.getNodeClass().successorOffsets[pos.index]), replacement); + target = replacement; } - Node succ = oldNode.getNodeClass().get(oldNode, pos); - Node target = newNodes.get(succ); - if (target == null) { - target = replacementsMap.get(succ); - if (target == null) { - Node replacement = replacements.replacement(succ); - if (replacement != succ) { - replacementsMap.put(succ, replacement); - assert isAssignable(node.getNodeClass().fieldTypes.get(node.getNodeClass().successorOffsets[pos.index]), replacement); - target = replacement; - } - } - } - node.getNodeClass().set(node, pos, target); } + nodeClass.set(node, pos, target); } - return newNodes; } private static boolean isAssignable(Class<?> fieldType, Node replacement) {
--- a/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeMap.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeMap.java Wed Oct 02 13:26:31 2013 +0200 @@ -26,11 +26,11 @@ import java.util.*; import java.util.Map.Entry; -public final class NodeMap<T> { +public class NodeMap<T> { private final Graph graph; private final boolean autogrow; - private Object[] values; + protected Object[] values; public NodeMap(Graph graph) { this(graph, false); @@ -54,6 +54,29 @@ return (T) values[node.id()]; } + public boolean isEmpty() { + return !entries().iterator().hasNext(); + } + + public boolean containsKey(Object key) { + if (key instanceof Node) { + Node node = (Node) key; + if (node.graph() == graph()) { + return get(node) != null; + } + } + return false; + } + + public boolean containsValue(Object value) { + for (Object o : values) { + if (o == value) { + return true; + } + } + return false; + } + public Graph graph() { return graph; }
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.graph/src/com/oracle/graal/graph/NodeNodeMap.java Wed Oct 02 13:26:31 2013 +0200 @@ -0,0 +1,75 @@ +/* + * Copyright (c) 2011, 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. + */ +package com.oracle.graal.graph; + +import java.util.*; + +public final class NodeNodeMap extends NodeMap<Node> implements Map<Node, Node> { + + public NodeNodeMap(Graph graph) { + super(graph, true); + } + + public NodeNodeMap(NodeNodeMap copyFrom) { + super(copyFrom); + } + + public Node get(Object key) { + return super.get((Node) key); + } + + public Node put(Node key, Node value) { + Node oldValue = super.get(key); + super.set(key, value); + return oldValue; + } + + public Node remove(Object key) { + throw new UnsupportedOperationException("Cannot remove keys from this map"); + } + + public void putAll(Map<? extends Node, ? extends Node> m) { + for (Entry<? extends Node, ? extends Node> entry : m.entrySet()) { + put(entry.getKey(), entry.getValue()); + } + } + + public Set<Node> keySet() { + throw new UnsupportedOperationException("Cannot get key set from this map"); + } + + public Collection<Node> values() { + ArrayList<Node> result = new ArrayList<>(this.size()); + for (int i = 0; i < values.length; ++i) { + Object v = values[i]; + if (v != null) { + result.add((Node) v); + } + } + return result; + } + + public Set<java.util.Map.Entry<Node, Node>> entrySet() { + throw new UnsupportedOperationException("Cannot get entry set for this map"); + } +}
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64DeoptimizeOp.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64DeoptimizeOp.java Wed Oct 02 13:26:31 2013 +0200 @@ -23,12 +23,8 @@ package com.oracle.graal.hotspot.amd64; import static com.oracle.graal.hotspot.HotSpotBackend.*; -import static com.oracle.graal.hotspot.HotSpotGraalRuntime.*; -import com.oracle.graal.api.code.*; -import com.oracle.graal.api.meta.*; import com.oracle.graal.asm.amd64.*; -import com.oracle.graal.hotspot.*; import com.oracle.graal.lir.*; import com.oracle.graal.lir.amd64.*; import com.oracle.graal.lir.asm.*; @@ -36,21 +32,14 @@ @Opcode("DEOPT") final class AMD64DeoptimizeOp extends AMD64LIRInstruction { - private DeoptimizationAction action; - private DeoptimizationReason reason; @State private LIRFrameState info; - AMD64DeoptimizeOp(DeoptimizationAction action, DeoptimizationReason reason, LIRFrameState info) { - this.action = action; - this.reason = reason; + AMD64DeoptimizeOp(LIRFrameState info) { this.info = info; } @Override public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) { - HotSpotGraalRuntime runtime = graalRuntime(); - Register thread = runtime.getRuntime().threadRegister(); - masm.movl(new AMD64Address(thread, runtime.getConfig().pendingDeoptimizationOffset), tasm.runtime.encodeDeoptActionAndReason(action, reason)); AMD64Call.directCall(tasm, masm, tasm.runtime.lookupForeignCall(UNCOMMON_TRAP), null, false, info); } }
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotBackend.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotBackend.java Wed Oct 02 13:26:31 2013 +0200 @@ -63,6 +63,11 @@ } @Override + public boolean shouldAllocateRegisters() { + return true; + } + + @Override public FrameMap newFrameMap() { return new AMD64FrameMap(runtime(), target, runtime().lookupRegisterConfig()); } @@ -231,11 +236,10 @@ FrameMap frameMap = tasm.frameMap; RegisterConfig regConfig = frameMap.registerConfig; HotSpotVMConfig config = runtime().config; - Label unverifiedStub = installedCodeOwner == null || isStatic(installedCodeOwner.getModifiers()) ? null : new Label(); + Label verifiedStub = new Label(); // Emit the prefix - - if (unverifiedStub != null) { + if (installedCodeOwner != null && !isStatic(installedCodeOwner.getModifiers())) { tasm.recordMark(Marks.MARK_UNVERIFIED_ENTRY); CallingConvention cc = regConfig.getCallingConvention(JavaCallee, null, new JavaType[]{runtime().lookupJavaType(Object.class)}, target, false); Register inlineCacheKlass = rax; // see definition of IC_Klass in @@ -252,11 +256,12 @@ } else { asm.cmpq(inlineCacheKlass, src); } - asm.jcc(ConditionFlag.NotEqual, unverifiedStub); + AMD64Call.directConditionalJmp(tasm, asm, runtime().lookupForeignCall(IC_MISS_HANDLER), ConditionFlag.NotEqual); } asm.align(config.codeEntryAlignment); tasm.recordMark(Marks.MARK_OSR_ENTRY); + asm.bind(verifiedStub); tasm.recordMark(Marks.MARK_VERIFIED_ENTRY); // Emit code for the LIR @@ -273,10 +278,5 @@ // it has no calls that can cause such "return" entries assert !frameMap.accessesCallerFrame() : lirGen.getGraph(); } - - if (unverifiedStub != null) { - asm.bind(unverifiedStub); - AMD64Call.directJmp(tasm, asm, runtime().lookupForeignCall(IC_MISS_HANDLER)); - } } }
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotDeoptimizeCallerOp.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotDeoptimizeCallerOp.java Wed Oct 02 13:26:31 2013 +0200 @@ -23,12 +23,8 @@ package com.oracle.graal.hotspot.amd64; import static com.oracle.graal.hotspot.HotSpotBackend.*; -import static com.oracle.graal.hotspot.HotSpotGraalRuntime.*; -import com.oracle.graal.api.code.*; -import com.oracle.graal.api.meta.*; import com.oracle.graal.asm.amd64.*; -import com.oracle.graal.hotspot.*; import com.oracle.graal.lir.*; import com.oracle.graal.lir.amd64.*; import com.oracle.graal.lir.asm.*; @@ -39,21 +35,9 @@ @Opcode("DEOPT_CALLER") final class AMD64HotSpotDeoptimizeCallerOp extends AMD64HotSpotEpilogueOp { - private final DeoptimizationAction action; - private final DeoptimizationReason reason; - - AMD64HotSpotDeoptimizeCallerOp(DeoptimizationAction action, DeoptimizationReason reason) { - this.action = action; - this.reason = reason; - } - @Override public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) { leaveFrameAndRestoreRbp(tasm, masm); - - HotSpotGraalRuntime runtime = graalRuntime(); - Register thread = runtime.getRuntime().threadRegister(); - masm.movl(new AMD64Address(thread, runtime.getConfig().pendingDeoptimizationOffset), tasm.runtime.encodeDeoptActionAndReason(action, reason)); AMD64Call.directJmp(tasm, masm, tasm.runtime.lookupForeignCall(UNCOMMON_TRAP)); } }
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotLIRGenerator.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotLIRGenerator.java Wed Oct 02 13:26:31 2013 +0200 @@ -25,6 +25,7 @@ import static com.oracle.graal.amd64.AMD64.*; import static com.oracle.graal.api.code.ValueUtil.*; import static com.oracle.graal.hotspot.HotSpotBackend.*; +import static com.oracle.graal.hotspot.HotSpotGraalRuntime.*; import java.lang.reflect.*; import java.util.*; @@ -36,6 +37,7 @@ import com.oracle.graal.asm.amd64.AMD64Address.Scale; import com.oracle.graal.compiler.amd64.*; import com.oracle.graal.compiler.gen.*; +import com.oracle.graal.debug.*; import com.oracle.graal.graph.*; import com.oracle.graal.hotspot.*; import com.oracle.graal.hotspot.amd64.AMD64HotSpotMove.CompareAndSwapCompressedOp; @@ -176,7 +178,7 @@ @Override protected void emitReturn(Value input) { - append(new AMD64HotSpotReturnOp(input)); + append(new AMD64HotSpotReturnOp(input, getStub() != null)); } @Override @@ -242,7 +244,7 @@ Variable result; if (linkage.canDeoptimize()) { - assert info != null; + assert info != null || stub != null; append(new AMD64HotSpotCRuntimeCallPrologueOp()); result = super.emitForeignCall(linkage, info, args); append(new AMD64HotSpotCRuntimeCallEpilogueOp()); @@ -287,7 +289,7 @@ @Override public void visitSafepointNode(SafepointNode i) { LIRFrameState info = state(i); - append(new AMD64SafepointOp(info, runtime().config, this)); + append(new AMD64HotSpotSafepointOp(info, runtime().config, this)); } @SuppressWarnings("hiding") @@ -361,14 +363,29 @@ append(new AMD64HotSpotUnwindOp(exceptionParameter)); } + private void moveDeoptimizationActionAndReasonToThread(Value actionAndReason) { + int pendingDeoptimizationOffset = graalRuntime().getConfig().pendingDeoptimizationOffset; + RegisterValue thread = runtime().threadRegister().asValue(HotSpotGraalRuntime.wordKind()); + AMD64AddressValue pendingDeoptAddress = new AMD64AddressValue(actionAndReason.getKind(), thread, pendingDeoptimizationOffset); + if (actionAndReason instanceof Constant && !runtime.needsDataPatch((Constant) actionAndReason)) { + Constant constantActionAndReason = (Constant) actionAndReason; + assert !runtime.needsDataPatch(constantActionAndReason); + append(new StoreConstantOp(constantActionAndReason.getKind(), pendingDeoptAddress, constantActionAndReason, null)); + } else { + append(new StoreOp(actionAndReason.getKind(), pendingDeoptAddress, load(actionAndReason), null)); + } + } + @Override - public void emitDeoptimize(DeoptimizationAction action, DeoptimizingNode deopting) { - append(new AMD64DeoptimizeOp(action, deopting.getDeoptimizationReason(), state(deopting))); + public void emitDeoptimize(Value actionAndReason, DeoptimizingNode deopting) { + moveDeoptimizationActionAndReasonToThread(actionAndReason); + append(new AMD64DeoptimizeOp(state(deopting))); } @Override public void emitDeoptimizeCaller(DeoptimizationAction action, DeoptimizationReason reason) { - append(new AMD64HotSpotDeoptimizeCallerOp(action, reason)); + moveDeoptimizationActionAndReasonToThread(runtime.encodeDeoptActionAndReason(action, reason)); + append(new AMD64HotSpotDeoptimizeCallerOp()); } @Override @@ -522,4 +539,13 @@ append(new CondMoveOp(result, Condition.EQ, load(Constant.TRUE), Constant.FALSE)); setResult(node, result); } + + @Override + public void visitInfopointNode(InfopointNode i) { + if (i.getState() != null && i.getState().bci == FrameState.AFTER_BCI) { + Debug.log("Ignoring InfopointNode for AFTER_BCI"); + } else { + super.visitInfopointNode(i); + } + } }
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotReturnOp.java Fri Sep 06 21:37:50 2013 +0200 +++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotReturnOp.java Wed Oct 02 13:26:31 2013 +0200 @@ -22,10 +22,16 @@ */ package com.oracle.graal.hotspot.amd64; +import static com.oracle.graal.hotspot.HotSpotGraalRuntime.*; import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*; +import static com.oracle.graal.phases.GraalOptions.*; +import com.oracle.graal.amd64.*; +import com.oracle.graal.api.code.*; import com.oracle.graal.api.meta.*; import com.oracle.graal.asm.amd64.*; +import com.oracle.graal.graph.*; +import com.oracle.graal.hotspot.*; import com.oracle.graal.lir.*; import com.oracle.graal.lir.asm.*; @@ -36,14 +42,31 @@ final class AMD64HotSpotReturnOp extends AMD64HotSpotEpilogueOp { @Use({REG, ILLEGAL}) protected Value value; + private final boolean isStub; - AMD64HotSpotReturnOp(Value value) { + AMD64HotSpotReturnOp(Value value, boolean isStub) { this.value = value; + this.isStub = isStub; } + private static Register findPollOnReturnScratchRegister() { + RegisterConfig config = HotSpotGraalRuntime.graalRuntime().getRuntime().lookupRegisterConfig(); + for (Register r : config.getAllocatableRegisters(Kind.Long)) { + if (r != config.getReturnRegister(Kind.Long) && r != AMD64.rbp) { + return r; + } + } + throw GraalInternalError.shouldNotReachHere(); + } + + private static final Register scratchForSafepointOnReturn = findPollOnReturnScratchRegister(); + @Override public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) { leaveFrameAndRestoreRbp(tasm, masm); + if (!isStub && (tasm.frameContext != null || !OptEliminateSafepoints.getValue())) { + AMD64HotSpotSafepointOp.emitCode(tasm, masm, graalRuntime().getConfig(), true, null, scratchForSafepointOnReturn); + } masm.ret(0); } }
--- /dev/null Thu Jan 01 00:00:00 1970 +0000 +++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotSafepointOp.java Wed Oct 02 13:26:31 2013 +0200 @@ -0,0 +1,79 @@ +/* + * 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