# HG changeset patch
# User Doug Simon <doug.simon@oracle.com>
# Date 1381964897 -7200
# Node ID f87c68d79f07e7418c6456a81c6e13741922582d
# Parent  d08accd5892573d2141820465dd93b2d5f1b2171
improved support for co-existing, multiple backends (GRAAL-363)

diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.api.runtime/src/com/oracle/graal/api/runtime/Graal.java
--- a/graal/com.oracle.graal.api.runtime/src/com/oracle/graal/api/runtime/Graal.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.api.runtime/src/com/oracle/graal/api/runtime/Graal.java	Thu Oct 17 01:08:17 2013 +0200
@@ -69,5 +69,10 @@
         public <T> T getCapability(Class<T> clazz) {
             return null;
         }
+
+        @Override
+        public <T> T getCapability(Class<T> clazz, String selector) {
+            return null;
+        }
     }
 }
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.api.runtime/src/com/oracle/graal/api/runtime/GraalRuntime.java
--- a/graal/com.oracle.graal.api.runtime/src/com/oracle/graal/api/runtime/GraalRuntime.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.api.runtime/src/com/oracle/graal/api/runtime/GraalRuntime.java	Thu Oct 17 01:08:17 2013 +0200
@@ -27,4 +27,14 @@
     String getName();
 
     <T> T getCapability(Class<T> clazz);
+
+    /**
+     * Requests a capability represented by a given class.
+     * 
+     * @param selector a value the runtime will use to refine the capability returned (if any). An
+     *            example may be the name of an architecture when asking for a capability
+     *            representing a compiler backend.
+     */
+    <T> T getCapability(Class<T> clazz, String selector);
+
 }
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/BasicHSAILTest.java
--- a/graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/BasicHSAILTest.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.compiler.hsail.test/src/com/oracle/graal/compiler/hsail/test/BasicHSAILTest.java	Thu Oct 17 01:08:17 2013 +0200
@@ -24,15 +24,19 @@
 
 import org.junit.*;
 
-import com.oracle.graal.compiler.test.GraalCompilerTest;
+import com.oracle.graal.compiler.test.*;
 import com.oracle.graal.hotspot.hsail.*;
-import com.oracle.graal.nodes.StructuredGraph;
+import com.oracle.graal.nodes.*;
 
 /**
  * Test class for small Java methods compiled to HSAIL kernels.
  */
 public class BasicHSAILTest extends GraalCompilerTest {
 
+    public BasicHSAILTest() {
+        super("HSAIL");
+    }
+
     public void testAdd() {
         test("testAddSnippet");
     }
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/PTXTestBase.java
--- a/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/PTXTestBase.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.compiler.ptx.test/src/com/oracle/graal/compiler/ptx/test/PTXTestBase.java	Thu Oct 17 01:08:17 2013 +0200
@@ -31,6 +31,7 @@
 import com.oracle.graal.api.code.CallingConvention.Type;
 import com.oracle.graal.compiler.*;
 import com.oracle.graal.compiler.ptx.*;
+import com.oracle.graal.compiler.target.*;
 import com.oracle.graal.compiler.test.*;
 import com.oracle.graal.debug.*;
 import com.oracle.graal.hotspot.meta.*;
@@ -41,7 +42,6 @@
 import com.oracle.graal.phases.*;
 import com.oracle.graal.phases.PhasePlan.PhasePosition;
 import com.oracle.graal.phases.tiers.*;
-import com.oracle.graal.ptx.*;
 
 public abstract class PTXTestBase extends GraalCompilerTest {
 
@@ -54,13 +54,17 @@
 
     }
 
+    public PTXTestBase() {
+        super("PTX");
+    }
+
     protected CompilationResult compile(String test) {
-        if (getCodeCache() instanceof PTXHotSpotCodeCacheProvider) {
+        if (getBackend() instanceof PTXHotSpotBackend) {
             StructuredGraph graph = parse(test);
             sg = graph;
             Debug.dump(graph, "Graph");
-            TargetDescription target = new TargetDescription(new PTX(), true, 1, 0, true);
-            PTXBackend ptxBackend = new PTXBackend(getProviders());
+            Backend ptxBackend = getBackend();
+            TargetDescription target = ptxBackend.getTarget();
             PhasePlan phasePlan = new PhasePlan();
             GraphBuilderPhase graphBuilderPhase = new GraphBuilderPhase(getMetaAccess(), getForeignCalls(), GraphBuilderConfiguration.getDefault(), OptimisticOptimizations.NONE);
             phasePlan.addPhase(PhasePosition.AFTER_PARSING, graphBuilderPhase);
@@ -76,8 +80,9 @@
              * 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.
              */
+            Suites suites = Suites.createDefaultSuites();
             CompilationResult result = GraalCompiler.compileGraph(graph, cc, graph.method(), getProviders(), ptxBackend, target, null, phasePlan, OptimisticOptimizations.NONE, new SpeculationLog(),
-                            Suites.createDefaultSuites(), new ExternalCompilationResult());
+                            suites, new ExternalCompilationResult());
             return result;
         } else {
             return null;
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java
--- a/graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java	Thu Oct 17 01:05:13 2013 +0200
+++ /dev/null	Thu Jan 01 00:00:00 1970 +0000
@@ -1,273 +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;
-
-import static com.oracle.graal.lir.LIRValueUtil.*;
-
-import java.util.*;
-
-import com.oracle.graal.api.code.*;
-import com.oracle.graal.api.meta.*;
-import com.oracle.graal.asm.*;
-import com.oracle.graal.asm.ptx.*;
-import com.oracle.graal.compiler.gen.*;
-import com.oracle.graal.compiler.target.*;
-import com.oracle.graal.graph.*;
-import com.oracle.graal.lir.*;
-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.lir.asm.*;
-import com.oracle.graal.lir.ptx.*;
-import com.oracle.graal.nodes.*;
-import com.oracle.graal.nodes.cfg.*;
-import com.oracle.graal.phases.util.*;
-
-/**
- * PTX specific backend.
- */
-public class PTXBackend extends Backend {
-
-    public PTXBackend(Providers providers) {
-        super(providers);
-    }
-
-    @Override
-    public boolean shouldAllocateRegisters() {
-        return false;
-    }
-
-    @Override
-    public FrameMap newFrameMap() {
-        return new PTXFrameMap(getCodeCache());
-    }
-
-    @Override
-    public LIRGenerator newLIRGenerator(StructuredGraph graph, FrameMap frameMap, CallingConvention cc, LIR lir) {
-        return new PTXLIRGenerator(graph, getProviders(), frameMap, cc, lir);
-    }
-
-    class PTXFrameContext implements FrameContext {
-
-        @Override
-        public void enter(TargetMethodAssembler tasm) {
-            // codeBuffer.emitString(".address_size 32"); // PTX ISA version 2.3
-        }
-
-        @Override
-        public void leave(TargetMethodAssembler tasm) {
-        }
-    }
-
-    @Override
-    protected AbstractAssembler createAssembler(FrameMap frameMap) {
-        return new PTXAssembler(getTarget(), frameMap.registerConfig);
-    }
-
-    @Override
-    public TargetMethodAssembler newAssembler(LIRGenerator lirGen, CompilationResult compilationResult) {
-        // 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
-        // - has no instructions with debug info
-        FrameMap frameMap = lirGen.frameMap;
-        AbstractAssembler masm = createAssembler(frameMap);
-        PTXFrameContext frameContext = new PTXFrameContext();
-        TargetMethodAssembler tasm = new PTXTargetMethodAssembler(getCodeCache(), getForeignCalls(), frameMap, masm, frameContext, compilationResult);
-        tasm.setFrameSize(0);
-        return tasm;
-    }
-
-    private static void emitKernelEntry(TargetMethodAssembler tasm, LIRGenerator lirGen, ResolvedJavaMethod codeCacheOwner) {
-        // Emit PTX kernel entry text based on PTXParameterOp
-        // instructions in the start block. Remove the instructions
-        // once kernel entry text and directives are emitted to
-        // facilitate seemless PTX code generation subsequently.
-        assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method";
-        final String name = codeCacheOwner.getName();
-        Buffer codeBuffer = tasm.asm.codeBuffer;
-
-        // Emit initial boiler-plate directives.
-        codeBuffer.emitString(".version 3.0");
-        codeBuffer.emitString(".target sm_30");
-        codeBuffer.emitString0(".entry " + name + " (");
-        codeBuffer.emitString("");
-
-        // Get the start block
-        Block startBlock = lirGen.lir.cfg.getStartBlock();
-        // Keep a list of ParameterOp instructions to delete from the
-        // list of instructions in the block.
-        ArrayList<LIRInstruction> deleteOps = new ArrayList<>();
-
-        // Emit .param arguments to kernel entry based on ParameterOp
-        // instruction.
-        for (LIRInstruction op : lirGen.lir.lir(startBlock)) {
-            if (op instanceof PTXParameterOp) {
-                op.emitCode(tasm);
-                deleteOps.add(op);
-            }
-        }
-
-        // Delete ParameterOp instructions.
-        for (LIRInstruction op : deleteOps) {
-            lirGen.lir.lir(startBlock).remove(op);
-        }
-
-        // Start emiting body of the PTX kernel.
-        codeBuffer.emitString0(") {");
-        codeBuffer.emitString("");
-    }
-
-    // Emit .reg space declarations
-    private static void emitRegisterDecl(TargetMethodAssembler tasm, LIRGenerator lirGen, ResolvedJavaMethod codeCacheOwner) {
-
-        assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method";
-
-        Buffer codeBuffer = tasm.asm.codeBuffer;
-
-        final SortedSet<Integer> signed32 = new TreeSet<>();
-        final SortedSet<Integer> signed64 = new TreeSet<>();
-        final SortedSet<Integer> unsigned64 = new TreeSet<>();
-        final SortedSet<Integer> float32 = new TreeSet<>();
-        final SortedSet<Integer> float64 = new TreeSet<>();
-
-        ValueProcedure trackRegisterKind = new ValueProcedure() {
-
-            @Override
-            public Value doValue(Value value, OperandMode mode, EnumSet<OperandFlag> flags) {
-                if (isVariable(value)) {
-                    Variable regVal = (Variable) value;
-                    Kind regKind = regVal.getKind();
-                    switch (regKind) {
-                        case Int:
-                            // If the register was used as a wider signed type
-                            // do not add it here
-                            if (!signed64.contains(regVal.index)) {
-                                signed32.add(regVal.index);
-                            }
-                            break;
-                        case Long:
-                            // If the register was used as a narrower signed type
-                            // remove it from there and add it to wider type.
-                            if (signed32.contains(regVal.index)) {
-                                signed32.remove(regVal.index);
-                            }
-                            signed64.add(regVal.index);
-                            break;
-                        case Float:
-                            // If the register was used as a wider signed type
-                            // do not add it here
-                            if (!float64.contains(regVal.index)) {
-                                float32.add(regVal.index);
-                            }
-                            break;
-                        case Double:
-                            // If the register was used as a narrower signed type
-                            // remove it from there and add it to wider type.
-                            if (float32.contains(regVal.index)) {
-                                float32.remove(regVal.index);
-                            }
-                            float64.add(regVal.index);
-                            break;
-                        case Object:
-                            unsigned64.add(regVal.index);
-                            break;
-                        default:
-                            throw GraalInternalError.shouldNotReachHere("unhandled register type " + value.toString());
-                    }
-                }
-                return value;
-            }
-        };
-
-        for (Block b : lirGen.lir.codeEmittingOrder()) {
-            for (LIRInstruction op : lirGen.lir.lir(b)) {
-                if (op instanceof LabelOp) {
-                    // Don't consider this as a definition
-                } else {
-                    op.forEachTemp(trackRegisterKind);
-                    op.forEachOutput(trackRegisterKind);
-                }
-            }
-        }
-
-        for (Integer i : signed32) {
-            codeBuffer.emitString(".reg .s32 %r" + i.intValue() + ";");
-        }
-        for (Integer i : signed64) {
-            codeBuffer.emitString(".reg .s64 %r" + i.intValue() + ";");
-        }
-        for (Integer i : unsigned64) {
-            codeBuffer.emitString(".reg .u64 %r" + i.intValue() + ";");
-        }
-        for (Integer i : float32) {
-            codeBuffer.emitString(".reg .f32 %r" + i.intValue() + ";");
-        }
-        for (Integer i : float64) {
-            codeBuffer.emitString(".reg .f64 %r" + i.intValue() + ";");
-        }
-        // emit predicate register declaration
-        int maxPredRegNum = ((PTXLIRGenerator) lirGen).getNextPredRegNumber();
-        if (maxPredRegNum > 0) {
-            codeBuffer.emitString(".reg .pred %p<" + maxPredRegNum + ">;");
-        }
-        codeBuffer.emitString(".reg .pred %r;");  // used for setp bool
-    }
-
-    @Override
-    public void emitCode(TargetMethodAssembler tasm, LIRGenerator lirGen, ResolvedJavaMethod codeCacheOwner) {
-        assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method";
-        Buffer codeBuffer = tasm.asm.codeBuffer;
-        // Emit the prologue
-        emitKernelEntry(tasm, lirGen, codeCacheOwner);
-
-        // Emit register declarations
-        try {
-            emitRegisterDecl(tasm, lirGen, codeCacheOwner);
-        } catch (GraalInternalError e) {
-            e.printStackTrace();
-            // TODO : Better error handling needs to be done once
-            // all types of parameters are handled.
-            codeBuffer.setPosition(0);
-            codeBuffer.close(false);
-            return;
-        }
-        // Emit code for the LIR
-        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);
-            codeBuffer.close(false);
-            return;
-        }
-
-        // Emit the epilogue
-        codeBuffer.emitString0("}");
-        codeBuffer.emitString("");
-    }
-}
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/GraalCompilerTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/GraalCompilerTest.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/GraalCompilerTest.java	Thu Oct 17 01:08:17 2013 +0200
@@ -78,12 +78,30 @@
 public abstract class GraalCompilerTest extends GraalTest {
 
     private final Providers providers;
-    protected final Backend backend;
-    protected final Suites suites;
+    private final Backend backend;
+    private final Suites suites;
 
     public GraalCompilerTest() {
-        this.providers = GraalCompiler.getGraalProviders();
         this.backend = Graal.getRequiredCapability(Backend.class);
+        this.providers = getBackend().getProviders();
+        this.suites = Graal.getRequiredCapability(SuitesProvider.class).createSuites();
+    }
+
+    /**
+     * Set up a test for a non-default backend. The test should check (via {@link #getBackend()} )
+     * whether the desired backend is available.
+     * 
+     * @param arch the name of the desired backend architecture
+     */
+    public GraalCompilerTest(String arch) {
+        Backend b = Graal.getRuntime().getCapability(Backend.class, arch);
+        if (b != null) {
+            this.backend = b;
+        } else {
+            // Fall back to the default/host backend
+            this.backend = Graal.getRuntime().getCapability(Backend.class);
+        }
+        this.providers = backend.getProviders();
         this.suites = Graal.getRequiredCapability(SuitesProvider.class).createSuites();
     }
 
@@ -157,6 +175,14 @@
         return result.toString();
     }
 
+    protected Backend getBackend() {
+        return backend;
+    }
+
+    protected Suites getSuites() {
+        return suites;
+    }
+
     protected Providers getProviders() {
         return providers;
     }
@@ -487,8 +513,8 @@
                 GraphBuilderPhase graphBuilderPhase = new GraphBuilderPhase(getMetaAccess(), getForeignCalls(), GraphBuilderConfiguration.getDefault(), OptimisticOptimizations.ALL);
                 phasePlan.addPhase(PhasePosition.AFTER_PARSING, graphBuilderPhase);
                 CallingConvention cc = getCallingConvention(getCodeCache(), Type.JavaCallee, graph.method(), false);
-                final CompilationResult compResult = GraalCompiler.compileGraph(graph, cc, method, getProviders(), backend, getCodeCache().getTarget(), null, phasePlan, OptimisticOptimizations.ALL,
-                                new SpeculationLog(), suites, new CompilationResult());
+                final CompilationResult compResult = GraalCompiler.compileGraph(graph, cc, method, getProviders(), getBackend(), getCodeCache().getTarget(), null, phasePlan,
+                                OptimisticOptimizations.ALL, new SpeculationLog(), getSuites(), new CompilationResult());
                 if (printCompilation) {
                     TTY.println(String.format("@%-6d Graal %-70s %-45s %-50s | %4dms %5dB", id, "", "", "", System.currentTimeMillis() - start, compResult.getTargetCodeSize()));
                 }
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/InfopointReasonTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/InfopointReasonTest.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/InfopointReasonTest.java	Thu Oct 17 01:08:17 2013 +0200
@@ -59,8 +59,8 @@
         final Method method = getMethod("testMethod");
         final StructuredGraph graph = parse(method);
         CallingConvention cc = getCallingConvention(getCodeCache(), Type.JavaCallee, graph.method(), false);
-        final CompilationResult cr = GraalCompiler.compileGraph(graph, cc, graph.method(), getProviders(), backend, getCodeCache().getTarget(), null, getDefaultPhasePlan(),
-                        OptimisticOptimizations.ALL, new SpeculationLog(), suites, new CompilationResult());
+        final CompilationResult cr = GraalCompiler.compileGraph(graph, cc, graph.method(), getProviders(), getBackend(), getCodeCache().getTarget(), null, getDefaultPhasePlan(),
+                        OptimisticOptimizations.ALL, new SpeculationLog(), getSuites(), new CompilationResult());
         for (Infopoint sp : cr.getInfopoints()) {
             assertNotNull(sp.reason);
             if (sp instanceof Call) {
@@ -81,8 +81,8 @@
         }
         assertTrue(graphLineSPs > 0);
         CallingConvention cc = getCallingConvention(getCodeCache(), Type.JavaCallee, graph.method(), false);
-        final CompilationResult cr = GraalCompiler.compileGraph(graph, cc, graph.method(), getProviders(), backend, getCodeCache().getTarget(), null, getDefaultPhasePlan(true),
-                        OptimisticOptimizations.ALL, new SpeculationLog(), suites, new CompilationResult());
+        final CompilationResult cr = GraalCompiler.compileGraph(graph, cc, graph.method(), getProviders(), getBackend(), getCodeCache().getTarget(), null, getDefaultPhasePlan(true),
+                        OptimisticOptimizations.ALL, new SpeculationLog(), getSuites(), new CompilationResult());
         int lineSPs = 0;
         for (Infopoint sp : cr.getInfopoints()) {
             assertNotNull(sp.reason);
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/backend/AllocatorTest.java
--- a/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/backend/AllocatorTest.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.compiler.test/src/com/oracle/graal/compiler/test/backend/AllocatorTest.java	Thu Oct 17 01:08:17 2013 +0200
@@ -119,7 +119,7 @@
 
             @Override
             public LIR call() {
-                return GraalCompiler.emitHIR(getProviders(), backend.getTarget(), graph, assumptions, null, phasePlan, OptimisticOptimizations.NONE, new SpeculationLog(), suites);
+                return GraalCompiler.emitHIR(getProviders(), getBackend().getTarget(), graph, assumptions, null, phasePlan, OptimisticOptimizations.NONE, new SpeculationLog(), getSuites());
             }
         });
 
@@ -128,7 +128,7 @@
             @Override
             public RegisterStats call() {
                 CallingConvention cc = getCallingConvention(getCodeCache(), Type.JavaCallee, graph.method(), false);
-                GraalCompiler.emitLIR(backend, backend.getTarget(), lir, graph, cc);
+                GraalCompiler.emitLIR(getBackend(), getBackend().getTarget(), lir, graph, cc);
                 return new RegisterStats(lir);
             }
         });
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/target/Backend.java
--- a/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/target/Backend.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.compiler/src/com/oracle/graal/compiler/target/Backend.java	Thu Oct 17 01:08:17 2013 +0200
@@ -32,7 +32,7 @@
 import com.oracle.graal.phases.util.*;
 
 /**
- * The {@code Backend} class represents a compiler backend for Graal.
+ * Represents a compiler backend for Graal.
  */
 public abstract class Backend {
 
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotBackendFactory.java
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotBackendFactory.java	Thu Oct 17 01:08:17 2013 +0200
@@ -0,0 +1,141 @@
+/*
+ * 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.hotspot.amd64;
+
+import java.util.*;
+
+import com.oracle.graal.amd64.*;
+import com.oracle.graal.api.code.*;
+import com.oracle.graal.api.meta.*;
+import com.oracle.graal.api.runtime.*;
+import com.oracle.graal.hotspot.*;
+import com.oracle.graal.hotspot.meta.*;
+import com.oracle.graal.phases.util.*;
+
+@ServiceProvider(HotSpotBackendFactory.class)
+public class AMD64HotSpotBackendFactory implements HotSpotBackendFactory {
+
+    protected Architecture createArchitecture(HotSpotVMConfig config) {
+        return new AMD64(config.useSSE, config.useAVX);
+    }
+
+    protected TargetDescription createTarget(HotSpotVMConfig config) {
+        final int stackFrameAlignment = 16;
+        final int implicitNullCheckLimit = 4096;
+        final boolean inlineObjects = true;
+        return new TargetDescription(createArchitecture(config), true, stackFrameAlignment, implicitNullCheckLimit, inlineObjects);
+    }
+
+    @Override
+    public HotSpotBackend createBackend(HotSpotGraalRuntime runtime, HotSpotBackend host) {
+        assert host == null;
+        TargetDescription target = createTarget(runtime.getConfig());
+
+        HotSpotMetaAccessProvider metaAccess = new HotSpotMetaAccessProvider(runtime);
+        HotSpotCodeCacheProvider codeCache = new AMD64HotSpotCodeCacheProvider(runtime, target);
+        HotSpotConstantReflectionProvider constantReflection = new HotSpotConstantReflectionProvider(runtime);
+        Value[] nativeABICallerSaveRegisters = createNativeABICallerSaveRegisters(runtime.getConfig(), codeCache.getRegisterConfig());
+        HotSpotForeignCallsProvider foreignCalls = new AMD64HotSpotForeignCallsProvider(runtime, metaAccess, codeCache, nativeABICallerSaveRegisters);
+        HotSpotLoweringProvider lowerer = new AMD64HotSpotLoweringProvider(runtime, metaAccess, foreignCalls);
+        // Replacements cannot have speculative optimizations since they have
+        // to be valid for the entire run of the VM.
+        Assumptions assumptions = new Assumptions(false);
+        Providers p = new Providers(metaAccess, codeCache, constantReflection, foreignCalls, lowerer, null);
+        HotSpotReplacementsImpl replacements = new HotSpotReplacementsImpl(p, runtime.getConfig(), assumptions);
+        HotSpotDisassemblerProvider disassembler = new HotSpotDisassemblerProvider(runtime);
+        HotSpotSuitesProvider suites = new HotSpotSuitesProvider(runtime);
+        HotSpotRegisters registers = new HotSpotRegisters(AMD64.r15, AMD64.r12, AMD64.rsp);
+        HotSpotProviders providers = new HotSpotProviders(metaAccess, codeCache, constantReflection, foreignCalls, lowerer, replacements, disassembler, suites, registers);
+
+        return new AMD64HotSpotBackend(runtime, providers);
+    }
+
+    private static Value[] createNativeABICallerSaveRegisters(HotSpotVMConfig config, RegisterConfig regConfig) {
+        List<Register> callerSave = new ArrayList<>(Arrays.asList(regConfig.getAllocatableRegisters()));
+        if (config.windowsOs) {
+            // http://msdn.microsoft.com/en-us/library/9z1stfyw.aspx
+            callerSave.remove(AMD64.rdi);
+            callerSave.remove(AMD64.rsi);
+            callerSave.remove(AMD64.rbx);
+            callerSave.remove(AMD64.rbp);
+            callerSave.remove(AMD64.rsp);
+            callerSave.remove(AMD64.r12);
+            callerSave.remove(AMD64.r13);
+            callerSave.remove(AMD64.r14);
+            callerSave.remove(AMD64.r15);
+            callerSave.remove(AMD64.xmm6);
+            callerSave.remove(AMD64.xmm7);
+            callerSave.remove(AMD64.xmm8);
+            callerSave.remove(AMD64.xmm9);
+            callerSave.remove(AMD64.xmm10);
+            callerSave.remove(AMD64.xmm11);
+            callerSave.remove(AMD64.xmm12);
+            callerSave.remove(AMD64.xmm13);
+            callerSave.remove(AMD64.xmm14);
+            callerSave.remove(AMD64.xmm15);
+        } else {
+            /*
+             * System V Application Binary Interface, AMD64 Architecture Processor Supplement
+             * 
+             * Draft Version 0.96
+             * 
+             * http://www.uclibc.org/docs/psABI-x86_64.pdf
+             * 
+             * 3.2.1
+             * 
+             * ...
+             * 
+             * This subsection discusses usage of each register. Registers %rbp, %rbx and %r12
+             * through %r15 "belong" to the calling function and the called function is required to
+             * preserve their values. In other words, a called function must preserve these
+             * registers' values for its caller. Remaining registers "belong" to the called
+             * function. If a calling function wants to preserve such a register value across a
+             * function call, it must save the value in its local stack frame.
+             */
+            callerSave.remove(AMD64.rbp);
+            callerSave.remove(AMD64.rbx);
+            callerSave.remove(AMD64.r12);
+            callerSave.remove(AMD64.r13);
+            callerSave.remove(AMD64.r14);
+            callerSave.remove(AMD64.r15);
+        }
+        Value[] nativeABICallerSaveRegisters = new Value[callerSave.size()];
+        for (int i = 0; i < callerSave.size(); i++) {
+            nativeABICallerSaveRegisters[i] = callerSave.get(i).asValue();
+        }
+        return nativeABICallerSaveRegisters;
+    }
+
+    public String getArchitecture() {
+        return "AMD64";
+    }
+
+    public String getGraalRuntimeName() {
+        return "basic";
+    }
+
+    @Override
+    public String toString() {
+        return getGraalRuntimeName() + ":" + getArchitecture();
+    }
+}
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotCRuntimeCallEpilogueOp.java
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotCRuntimeCallEpilogueOp.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotCRuntimeCallEpilogueOp.java	Thu Oct 17 01:08:17 2013 +0200
@@ -22,11 +22,8 @@
  */
 package com.oracle.graal.hotspot.amd64;
 
-import static com.oracle.graal.hotspot.HotSpotGraalRuntime.*;
-
 import com.oracle.graal.api.code.*;
 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.*;
@@ -34,13 +31,20 @@
 @Opcode("CRUNTIME_CALL_EPILOGUE")
 final class AMD64HotSpotCRuntimeCallEpilogueOp extends AMD64LIRInstruction {
 
+    private final int threadLastJavaSpOffset;
+    private final int threadLastJavaFpOffset;
+    private final Register thread;
+
+    public AMD64HotSpotCRuntimeCallEpilogueOp(int threadLastJavaSpOffset, int threadLastJavaFpOffset, Register thread) {
+        this.threadLastJavaSpOffset = threadLastJavaSpOffset;
+        this.threadLastJavaFpOffset = threadLastJavaFpOffset;
+        this.thread = thread;
+    }
+
     @Override
     public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) {
         // reset last Java frame:
-        HotSpotVMConfig config = runtime().getConfig();
-        Register thread = runtime().getProviders().getRegisters().getThreadRegister();
-
-        masm.movslq(new AMD64Address(thread, config.threadLastJavaSpOffset), 0);
-        masm.movslq(new AMD64Address(thread, config.threadLastJavaFpOffset), 0);
+        masm.movslq(new AMD64Address(thread, threadLastJavaSpOffset), 0);
+        masm.movslq(new AMD64Address(thread, threadLastJavaFpOffset), 0);
     }
 }
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotCRuntimeCallPrologueOp.java
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotCRuntimeCallPrologueOp.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotCRuntimeCallPrologueOp.java	Thu Oct 17 01:08:17 2013 +0200
@@ -23,7 +23,6 @@
 package com.oracle.graal.hotspot.amd64;
 
 import static com.oracle.graal.amd64.AMD64.*;
-import static com.oracle.graal.hotspot.HotSpotGraalRuntime.*;
 
 import com.oracle.graal.api.code.*;
 import com.oracle.graal.asm.amd64.*;
@@ -34,11 +33,17 @@
 @Opcode
 final class AMD64HotSpotCRuntimeCallPrologueOp extends AMD64LIRInstruction {
 
+    private final int threadLastJavaSpOffset;
+    private final Register thread;
+
+    public AMD64HotSpotCRuntimeCallPrologueOp(int threadLastJavaSpOffset, Register thread) {
+        this.threadLastJavaSpOffset = threadLastJavaSpOffset;
+        this.thread = thread;
+    }
+
     @Override
     public void emitCode(TargetMethodAssembler tasm, AMD64MacroAssembler masm) {
-
         // save last Java frame
-        Register thread = runtime().getProviders().getRegisters().getThreadRegister();
-        masm.movq(new AMD64Address(thread, runtime().getConfig().threadLastJavaSpOffset), rsp);
+        masm.movq(new AMD64Address(thread, threadLastJavaSpOffset), rsp);
     }
 }
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotCodeCacheProvider.java
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotCodeCacheProvider.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotCodeCacheProvider.java	Thu Oct 17 01:08:17 2013 +0200
@@ -28,13 +28,12 @@
 
 public class AMD64HotSpotCodeCacheProvider extends HotSpotCodeCacheProvider {
 
-    public AMD64HotSpotCodeCacheProvider(HotSpotGraalRuntime runtime) {
-        super(runtime);
-
+    public AMD64HotSpotCodeCacheProvider(HotSpotGraalRuntime runtime, TargetDescription target) {
+        super(runtime, target);
     }
 
     @Override
     protected RegisterConfig createRegisterConfig() {
-        return new AMD64HotSpotRegisterConfig(runtime.getTarget().arch, runtime.getConfig());
+        return new AMD64HotSpotRegisterConfig(getTarget().arch, runtime.getConfig());
     }
 }
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotForeignCallsProvider.java
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotForeignCallsProvider.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotForeignCallsProvider.java	Thu Oct 17 01:08:17 2013 +0200
@@ -39,10 +39,13 @@
 import com.oracle.graal.hotspot.*;
 import com.oracle.graal.hotspot.meta.*;
 
-public class AMD64HotSpotForeignCallsProvider extends HotSpotForeignCallsProvider {
+public class AMD64HotSpotForeignCallsProvider extends HotSpotHostForeignCallsProvider {
+
+    private final Value[] nativeABICallerSaveRegisters;
 
-    public AMD64HotSpotForeignCallsProvider(HotSpotGraalRuntime runtime) {
-        super(runtime);
+    public AMD64HotSpotForeignCallsProvider(HotSpotGraalRuntime runtime, MetaAccessProvider metaAccess, CodeCacheProvider codeCache, Value[] nativeABICallerSaveRegisters) {
+        super(runtime, metaAccess, codeCache);
+        this.nativeABICallerSaveRegisters = nativeABICallerSaveRegisters;
     }
 
     @Override
@@ -68,4 +71,10 @@
 
         super.initialize(providers);
     }
+
+    @Override
+    public Value[] getNativeABICallerSaveRegisters() {
+        return nativeABICallerSaveRegisters;
+    }
+
 }
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotGraalRuntime.java
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotGraalRuntime.java	Thu Oct 17 01:05:13 2013 +0200
+++ /dev/null	Thu Jan 01 00:00:00 1970 +0000
@@ -1,155 +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.hotspot.amd64;
-
-import java.util.*;
-
-import com.oracle.graal.amd64.*;
-import com.oracle.graal.api.code.*;
-import com.oracle.graal.api.meta.*;
-import com.oracle.graal.hotspot.*;
-import com.oracle.graal.hotspot.meta.*;
-import com.oracle.graal.phases.util.*;
-
-/**
- * AMD64 specific implementation of {@link HotSpotGraalRuntime}.
- */
-public class AMD64HotSpotGraalRuntime extends HotSpotGraalRuntime {
-
-    private Value[] nativeABICallerSaveRegisters;
-
-    protected AMD64HotSpotGraalRuntime() {
-    }
-
-    /**
-     * Called from C++ code to retrieve the singleton instance, creating it first if necessary.
-     */
-    public static HotSpotGraalRuntime makeInstance() {
-        HotSpotGraalRuntime runtime = runtime();
-        if (runtime == null) {
-            HotSpotGraalRuntimeFactory factory = findFactory("AMD64");
-            if (factory != null) {
-                runtime = factory.createRuntime();
-            } else {
-                runtime = new AMD64HotSpotGraalRuntime();
-            }
-            runtime.completeInitialization();
-        }
-        return runtime;
-    }
-
-    protected Architecture createArchitecture() {
-        return new AMD64(config.useSSE, config.useAVX);
-    }
-
-    @Override
-    protected HotSpotProviders createProviders() {
-        HotSpotMetaAccessProvider metaAccess = new HotSpotMetaAccessProvider(this);
-        HotSpotCodeCacheProvider codeCache = new AMD64HotSpotCodeCacheProvider(this);
-        HotSpotConstantReflectionProvider constantReflection = new HotSpotConstantReflectionProvider(this);
-        HotSpotForeignCallsProvider foreignCalls = new AMD64HotSpotForeignCallsProvider(this);
-        HotSpotLoweringProvider lowerer = new AMD64HotSpotLoweringProvider(this, metaAccess, foreignCalls);
-        // Replacements cannot have speculative optimizations since they have
-        // to be valid for the entire run of the VM.
-        Assumptions assumptions = new Assumptions(false);
-        Providers p = new Providers(metaAccess, codeCache, constantReflection, foreignCalls, lowerer, null);
-        HotSpotReplacementsImpl replacements = new HotSpotReplacementsImpl(p, getConfig(), assumptions);
-        HotSpotDisassemblerProvider disassembler = new HotSpotDisassemblerProvider(this);
-        HotSpotSuitesProvider suites = new HotSpotSuitesProvider(this);
-        HotSpotRegisters registers = new HotSpotRegisters(AMD64.r15, AMD64.r12, AMD64.rsp);
-        return new HotSpotProviders(metaAccess, codeCache, constantReflection, foreignCalls, lowerer, replacements, disassembler, suites, registers);
-    }
-
-    @Override
-    protected TargetDescription createTarget() {
-        final int stackFrameAlignment = 16;
-        final int implicitNullCheckLimit = 4096;
-        final boolean inlineObjects = true;
-        return new TargetDescription(createArchitecture(), true, stackFrameAlignment, implicitNullCheckLimit, inlineObjects);
-    }
-
-    @Override
-    protected HotSpotBackend createBackend() {
-        return new AMD64HotSpotBackend(this, getProviders());
-    }
-
-    @Override
-    protected Value[] getNativeABICallerSaveRegisters() {
-        if (nativeABICallerSaveRegisters == null) {
-            List<Register> callerSave = new ArrayList<>(Arrays.asList(getProviders().getCodeCache().getRegisterConfig().getAllocatableRegisters()));
-            if (getConfig().windowsOs) {
-                // http://msdn.microsoft.com/en-us/library/9z1stfyw.aspx
-                callerSave.remove(AMD64.rdi);
-                callerSave.remove(AMD64.rsi);
-                callerSave.remove(AMD64.rbx);
-                callerSave.remove(AMD64.rbp);
-                callerSave.remove(AMD64.rsp);
-                callerSave.remove(AMD64.r12);
-                callerSave.remove(AMD64.r13);
-                callerSave.remove(AMD64.r14);
-                callerSave.remove(AMD64.r15);
-                callerSave.remove(AMD64.xmm6);
-                callerSave.remove(AMD64.xmm7);
-                callerSave.remove(AMD64.xmm8);
-                callerSave.remove(AMD64.xmm9);
-                callerSave.remove(AMD64.xmm10);
-                callerSave.remove(AMD64.xmm11);
-                callerSave.remove(AMD64.xmm12);
-                callerSave.remove(AMD64.xmm13);
-                callerSave.remove(AMD64.xmm14);
-                callerSave.remove(AMD64.xmm15);
-            } else {
-                /*
-                 * System V Application Binary Interface, AMD64 Architecture Processor Supplement
-                 * 
-                 * Draft Version 0.96
-                 * 
-                 * http://www.uclibc.org/docs/psABI-x86_64.pdf
-                 * 
-                 * 3.2.1
-                 * 
-                 * ...
-                 * 
-                 * This subsection discusses usage of each register. Registers %rbp, %rbx and %r12
-                 * through %r15 "belong" to the calling function and the called function is required
-                 * to preserve their values. In other words, a called function must preserve these
-                 * registers' values for its caller. Remaining registers "belong" to the called
-                 * function. If a calling function wants to preserve such a register value across a
-                 * function call, it must save the value in its local stack frame.
-                 */
-                callerSave.remove(AMD64.rbp);
-                callerSave.remove(AMD64.rbx);
-                callerSave.remove(AMD64.r12);
-                callerSave.remove(AMD64.r13);
-                callerSave.remove(AMD64.r14);
-                callerSave.remove(AMD64.r15);
-            }
-            nativeABICallerSaveRegisters = new Value[callerSave.size()];
-            for (int i = 0; i < callerSave.size(); i++) {
-                nativeABICallerSaveRegisters[i] = callerSave.get(i).asValue();
-            }
-        }
-
-        return nativeABICallerSaveRegisters;
-    }
-}
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotGraalRuntimeFactory.java
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotGraalRuntimeFactory.java	Thu Oct 17 01:05:13 2013 +0200
+++ /dev/null	Thu Jan 01 00:00:00 1970 +0000
@@ -1,45 +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.hotspot.amd64;
-
-import com.oracle.graal.api.runtime.*;
-import com.oracle.graal.hotspot.*;
-
-@ServiceProvider(HotSpotGraalRuntimeFactory.class)
-public class AMD64HotSpotGraalRuntimeFactory implements HotSpotGraalRuntimeFactory {
-
-    @Override
-    public HotSpotGraalRuntime createRuntime() {
-        return new AMD64HotSpotGraalRuntime();
-    }
-
-    @Override
-    public String getArchitecture() {
-        return "AMD64";
-    }
-
-    @Override
-    public String getName() {
-        return "basic";
-    }
-}
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotJumpToExceptionHandlerInCallerOp.java
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotJumpToExceptionHandlerInCallerOp.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotJumpToExceptionHandlerInCallerOp.java	Thu Oct 17 01:08:17 2013 +0200
@@ -24,7 +24,6 @@
 
 import static com.oracle.graal.amd64.AMD64.*;
 import static com.oracle.graal.api.code.ValueUtil.*;
-import static com.oracle.graal.hotspot.HotSpotGraalRuntime.*;
 import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*;
 
 import com.oracle.graal.api.code.*;
@@ -44,11 +43,15 @@
     @Use(REG) AllocatableValue handlerInCallerPc;
     @Use(REG) AllocatableValue exception;
     @Use(REG) AllocatableValue exceptionPc;
+    private final Register thread;
+    private final int isMethodHandleReturnOffset;
 
-    AMD64HotSpotJumpToExceptionHandlerInCallerOp(AllocatableValue handlerInCallerPc, AllocatableValue exception, AllocatableValue exceptionPc) {
+    AMD64HotSpotJumpToExceptionHandlerInCallerOp(AllocatableValue handlerInCallerPc, AllocatableValue exception, AllocatableValue exceptionPc, int isMethodHandleReturnOffset, Register thread) {
         this.handlerInCallerPc = handlerInCallerPc;
         this.exception = exception;
         this.exceptionPc = exceptionPc;
+        this.isMethodHandleReturnOffset = isMethodHandleReturnOffset;
+        this.thread = thread;
     }
 
     @Override
@@ -59,8 +62,6 @@
         masm.incrementq(rsp, 8);
 
         // Restore rsp from rbp if the exception PC is a method handle call site.
-        Register thread = runtime().getProviders().getRegisters().getThreadRegister();
-        int isMethodHandleReturnOffset = runtime().getConfig().threadIsMethodHandleReturnOffset;
         AMD64Address dst = new AMD64Address(thread, isMethodHandleReturnOffset);
         masm.cmpl(dst, 0);
         masm.cmovq(ConditionFlag.NotEqual, rsp, rbp);
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotLIRGenerator.java
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotLIRGenerator.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotLIRGenerator.java	Thu Oct 17 01:08:17 2013 +0200
@@ -180,9 +180,24 @@
         }
     }
 
+    private Register findPollOnReturnScratchRegister() {
+        RegisterConfig regConfig = getProviders().getCodeCache().getRegisterConfig();
+        for (Register r : regConfig.getAllocatableRegisters(Kind.Long)) {
+            if (r != regConfig.getReturnRegister(Kind.Long) && r != AMD64.rbp) {
+                return r;
+            }
+        }
+        throw GraalInternalError.shouldNotReachHere();
+    }
+
+    private Register pollOnReturnScratchRegister;
+
     @Override
     protected void emitReturn(Value input) {
-        append(new AMD64HotSpotReturnOp(input, getStub() != null));
+        if (pollOnReturnScratchRegister == null) {
+            pollOnReturnScratchRegister = findPollOnReturnScratchRegister();
+        }
+        append(new AMD64HotSpotReturnOp(input, getStub() != null, pollOnReturnScratchRegister));
     }
 
     @Override
@@ -249,9 +264,10 @@
 
         if (linkage.canDeoptimize()) {
             assert info != null || stub != null;
-            append(new AMD64HotSpotCRuntimeCallPrologueOp());
+            Register thread = getProviders().getRegisters().getThreadRegister();
+            append(new AMD64HotSpotCRuntimeCallPrologueOp(config.threadLastJavaSpOffset, thread));
             result = super.emitForeignCall(linkage, info, args);
-            append(new AMD64HotSpotCRuntimeCallEpilogueOp());
+            append(new AMD64HotSpotCRuntimeCallEpilogueOp(config.threadLastJavaSpOffset, config.threadLastJavaFpOffset, thread));
         } else {
             result = super.emitForeignCall(linkage, null, args);
         }
@@ -408,7 +424,8 @@
         RegisterValue exceptionPcFixed = (RegisterValue) outgoingCc.getArgument(1);
         emitMove(exceptionFixed, operand(exception));
         emitMove(exceptionPcFixed, operand(exceptionPc));
-        AMD64HotSpotJumpToExceptionHandlerInCallerOp op = new AMD64HotSpotJumpToExceptionHandlerInCallerOp(handler, exceptionFixed, exceptionPcFixed);
+        Register thread = getProviders().getRegisters().getThreadRegister();
+        AMD64HotSpotJumpToExceptionHandlerInCallerOp op = new AMD64HotSpotJumpToExceptionHandlerInCallerOp(handler, exceptionFixed, exceptionPcFixed, config.threadIsMethodHandleReturnOffset, thread);
         append(op);
     }
 
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotLoweringProvider.java
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotLoweringProvider.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotLoweringProvider.java	Thu Oct 17 01:08:17 2013 +0200
@@ -41,7 +41,7 @@
 
     @Override
     public void initialize() {
-        HotSpotProviders providers = runtime.getProviders();
+        HotSpotProviders providers = runtime.getHostProviders();
         convertSnippets = new AMD64ConvertSnippets.Templates(providers, runtime.getTarget());
         super.initialize();
     }
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotReturnOp.java
--- a/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotReturnOp.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.amd64/src/com/oracle/graal/hotspot/amd64/AMD64HotSpotReturnOp.java	Thu Oct 17 01:08:17 2013 +0200
@@ -26,12 +26,9 @@
 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.*;
 
@@ -43,24 +40,14 @@
 
     @Use({REG, ILLEGAL}) protected Value value;
     private final boolean isStub;
+    private final Register scratchForSafepointOnReturn;
 
-    AMD64HotSpotReturnOp(Value value, boolean isStub) {
+    AMD64HotSpotReturnOp(Value value, boolean isStub, Register scratchForSafepointOnReturn) {
         this.value = value;
         this.isStub = isStub;
+        this.scratchForSafepointOnReturn = scratchForSafepointOnReturn;
     }
 
-    private static Register findPollOnReturnScratchRegister() {
-        RegisterConfig config = HotSpotGraalRuntime.runtime().getProviders().getCodeCache().getRegisterConfig();
-        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);
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.hsail/src/com/oracle/graal/hotspot/hsail/ForEachToGraal.java
--- a/graal/com.oracle.graal.hotspot.hsail/src/com/oracle/graal/hotspot/hsail/ForEachToGraal.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.hsail/src/com/oracle/graal/hotspot/hsail/ForEachToGraal.java	Thu Oct 17 01:08:17 2013 +0200
@@ -58,7 +58,7 @@
             }
         }
         HotSpotGraalRuntime runtime = HotSpotGraalRuntime.runtime();
-        HotSpotProviders providers = runtime.getProviders();
+        HotSpotProviders providers = runtime.getHostProviders();
         MetaAccessProvider metaAccess = providers.getMetaAccess();
         ResolvedJavaMethod method = metaAccess.lookupJavaMethod(acceptMethod);
         StructuredGraph graph = new StructuredGraph(method);
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.hsail/src/com/oracle/graal/hotspot/hsail/HSAILCompilationResult.java
--- a/graal/com.oracle.graal.hotspot.hsail/src/com/oracle/graal/hotspot/hsail/HSAILCompilationResult.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.hsail/src/com/oracle/graal/hotspot/hsail/HSAILCompilationResult.java	Thu Oct 17 01:08:17 2013 +0200
@@ -23,6 +23,8 @@
 
 package com.oracle.graal.hotspot.hsail;
 
+import static com.oracle.graal.hotspot.HotSpotGraalRuntime.*;
+
 import java.lang.reflect.*;
 import java.util.logging.*;
 
@@ -31,9 +33,9 @@
 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.hotspot.*;
 import com.oracle.graal.hotspot.meta.*;
 import com.oracle.graal.java.*;
 import com.oracle.graal.nodes.*;
@@ -73,17 +75,26 @@
         consoleHandler.setLevel(logLevel);
     }
 
-    private static final HotSpotGraalRuntime runtime = new HSAILHotSpotGraalRuntime();
+    private static final HSAILHotSpotBackend backend;
+    static {
+        // Look for installed HSAIL backend
+        HSAILHotSpotBackend b = (HSAILHotSpotBackend) Graal.getRuntime().getCapability(Backend.class, "HSAIL");
+        if (b == null) {
+            // Fall back to a new instance
+            b = new HSAILHotSpotBackendFactory().createBackend(runtime(), runtime().getHostBackend());
+        }
+        backend = b;
+    }
 
     public static HSAILCompilationResult getHSAILCompilationResult(Method meth) {
-        HotSpotMetaAccessProvider metaAccess = runtime.getProviders().getMetaAccess();
+        HotSpotMetaAccessProvider metaAccess = backend.getProviders().getMetaAccess();
         ResolvedJavaMethod javaMethod = metaAccess.lookupJavaMethod(meth);
         return getHSAILCompilationResult(javaMethod);
     }
 
     public static HSAILCompilationResult getHSAILCompilationResult(ResolvedJavaMethod javaMethod) {
-        HotSpotMetaAccessProvider metaAccess = runtime.getProviders().getMetaAccess();
-        ForeignCallsProvider foreignCalls = runtime.getProviders().getForeignCalls();
+        HotSpotMetaAccessProvider metaAccess = backend.getProviders().getMetaAccess();
+        ForeignCallsProvider foreignCalls = backend.getProviders().getForeignCalls();
         StructuredGraph graph = new StructuredGraph(javaMethod);
         new GraphBuilderPhase(metaAccess, foreignCalls, GraphBuilderConfiguration.getEagerDefault(), OptimisticOptimizations.ALL).apply(graph);
         return getHSAILCompilationResult(graph);
@@ -112,15 +123,14 @@
             argTypes[argIndex++] = sig.getParameterType(i, null);
         }
 
-        RegisterConfig registerConfig = runtime.getProviders().getCodeCache().getRegisterConfig();
+        RegisterConfig registerConfig = backend.getProviders().getCodeCache().getRegisterConfig();
         return registerConfig.getCallingConvention(type, retType, argTypes, target, stackOnly);
     }
 
     public static HSAILCompilationResult getHSAILCompilationResult(StructuredGraph graph) {
         Debug.dump(graph, "Graph");
-        Providers providers = runtime.getProviders();
+        Providers providers = backend.getProviders();
         TargetDescription target = providers.getCodeCache().getTarget();
-        HSAILHotSpotBackend hsailBackend = (HSAILHotSpotBackend) runtime.getBackend();
         PhasePlan phasePlan = new PhasePlan();
         GraphBuilderPhase graphBuilderPhase = new GraphBuilderPhase(providers.getMetaAccess(), providers.getForeignCalls(), GraphBuilderConfiguration.getDefault(), OptimisticOptimizations.NONE);
         phasePlan.addPhase(PhasePosition.AFTER_PARSING, graphBuilderPhase);
@@ -129,11 +139,11 @@
         CallingConvention cc = getHSAILCallingConvention(Type.JavaCallee, target, graph.method(), false);
         SuitesProvider suitesProvider = Graal.getRequiredCapability(SuitesProvider.class);
         try {
-            CompilationResult compResult = GraalCompiler.compileGraph(graph, cc, graph.method(), providers, hsailBackend, target, null, phasePlan, OptimisticOptimizations.NONE, new SpeculationLog(),
+            CompilationResult compResult = GraalCompiler.compileGraph(graph, cc, graph.method(), providers, backend, target, null, phasePlan, OptimisticOptimizations.NONE, new SpeculationLog(),
                             suitesProvider.getDefaultSuites(), new CompilationResult());
             return new HSAILCompilationResult(compResult);
         } catch (GraalInternalError e) {
-            String partialCode = hsailBackend.getPartialCodeString();
+            String partialCode = backend.getPartialCodeString();
             if (partialCode != null && !partialCode.equals("")) {
                 logger.fine("-------------------\nPartial Code Generation:\n--------------------");
                 logger.fine(partialCode);
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.hsail/src/com/oracle/graal/hotspot/hsail/HSAILHotSpotBackendFactory.java
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.graal.hotspot.hsail/src/com/oracle/graal/hotspot/hsail/HSAILHotSpotBackendFactory.java	Thu Oct 17 01:08:17 2013 +0200
@@ -0,0 +1,73 @@
+/*
+ * 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.hotspot.hsail;
+
+import com.oracle.graal.api.code.*;
+import com.oracle.graal.api.meta.*;
+import com.oracle.graal.api.runtime.*;
+import com.oracle.graal.hotspot.*;
+import com.oracle.graal.hotspot.meta.*;
+import com.oracle.graal.hsail.*;
+import com.oracle.graal.nodes.spi.*;
+
+@ServiceProvider(HotSpotBackendFactory.class)
+public class HSAILHotSpotBackendFactory implements HotSpotBackendFactory {
+
+    @Override
+    public HSAILHotSpotBackend createBackend(HotSpotGraalRuntime runtime, HotSpotBackend hostBackend) {
+        HotSpotProviders host = hostBackend.getProviders();
+
+        HotSpotMetaAccessProvider metaAccess = host.getMetaAccess();
+        HSAILHotSpotCodeCacheProvider codeCache = new HSAILHotSpotCodeCacheProvider(runtime, createTarget());
+        ConstantReflectionProvider constantReflection = host.getConstantReflection();
+        HotSpotForeignCallsProvider foreignCalls = new HSAILHotSpotForeignCallsProvider(host.getForeignCalls());
+        LoweringProvider lowerer = new HSAILHotSpotLoweringProvider(host.getLowerer());
+        Replacements replacements = host.getReplacements();
+        HotSpotDisassemblerProvider disassembler = host.getDisassembler();
+        HotSpotSuitesProvider suites = host.getSuites();
+        HotSpotRegisters registers = new HotSpotRegisters(Register.None, Register.None, Register.None);
+        HotSpotProviders providers = new HotSpotProviders(metaAccess, codeCache, constantReflection, foreignCalls, lowerer, replacements, disassembler, suites, registers);
+
+        return new HSAILHotSpotBackend(runtime, providers);
+    }
+
+    protected TargetDescription createTarget() {
+        final int stackFrameAlignment = 8;
+        final int implicitNullCheckLimit = 0;
+        final boolean inlineObjects = true;
+        return new TargetDescription(new HSAIL(), true, stackFrameAlignment, implicitNullCheckLimit, inlineObjects);
+    }
+
+    public String getArchitecture() {
+        return "HSAIL";
+    }
+
+    public String getGraalRuntimeName() {
+        return "basic";
+    }
+
+    @Override
+    public String toString() {
+        return getGraalRuntimeName() + ":" + getArchitecture();
+    }
+}
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.hsail/src/com/oracle/graal/hotspot/hsail/HSAILHotSpotCodeCacheProvider.java
--- a/graal/com.oracle.graal.hotspot.hsail/src/com/oracle/graal/hotspot/hsail/HSAILHotSpotCodeCacheProvider.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.hsail/src/com/oracle/graal/hotspot/hsail/HSAILHotSpotCodeCacheProvider.java	Thu Oct 17 01:08:17 2013 +0200
@@ -29,8 +29,8 @@
 
 public class HSAILHotSpotCodeCacheProvider extends HotSpotCodeCacheProvider {
 
-    public HSAILHotSpotCodeCacheProvider(HotSpotGraalRuntime runtime) {
-        super(runtime);
+    public HSAILHotSpotCodeCacheProvider(HotSpotGraalRuntime runtime, TargetDescription target) {
+        super(runtime, target);
 
     }
 
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.hsail/src/com/oracle/graal/hotspot/hsail/HSAILHotSpotForeignCallsProvider.java
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.graal.hotspot.hsail/src/com/oracle/graal/hotspot/hsail/HSAILHotSpotForeignCallsProvider.java	Thu Oct 17 01:08:17 2013 +0200
@@ -0,0 +1,60 @@
+/*
+ * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved.
+ * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
+ *
+ * This code is free software; you can redistribute it and/or modify it
+ * under the terms of the GNU General Public License version 2 only, as
+ * published by the Free Software Foundation.
+ *
+ * This code is distributed in the hope that it will be useful, but WITHOUT
+ * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+ * FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+ * version 2 for more details (a copy is included in the LICENSE file that
+ * accompanied this code).
+ *
+ * You should have received a copy of the GNU General Public License version
+ * 2 along with this work; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
+ *
+ * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
+ * or visit www.oracle.com if you need additional information or have any
+ * questions.
+ */
+package com.oracle.graal.hotspot.hsail;
+
+import com.oracle.graal.api.code.*;
+import com.oracle.graal.api.meta.*;
+import com.oracle.graal.graph.*;
+import com.oracle.graal.hotspot.meta.*;
+
+public class HSAILHotSpotForeignCallsProvider implements HotSpotForeignCallsProvider {
+
+    private final ForeignCallsProvider host;
+
+    public HSAILHotSpotForeignCallsProvider(ForeignCallsProvider host) {
+        this.host = host;
+    }
+
+    public boolean isReexecutable(ForeignCallDescriptor descriptor) {
+        return host.isReexecutable(descriptor);
+    }
+
+    public LocationIdentity[] getKilledLocations(ForeignCallDescriptor descriptor) {
+        return host.getKilledLocations(descriptor);
+    }
+
+    public boolean canDeoptimize(ForeignCallDescriptor descriptor) {
+        return host.canDeoptimize(descriptor);
+    }
+
+    public ForeignCallLinkage lookupForeignCall(ForeignCallDescriptor descriptor) {
+        return host.lookupForeignCall(descriptor);
+    }
+
+    public Value[] getNativeABICallerSaveRegisters() {
+        throw GraalInternalError.unimplemented();
+    }
+
+    public void initialize(HotSpotProviders providers) {
+    }
+}
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.hsail/src/com/oracle/graal/hotspot/hsail/HSAILHotSpotGraalRuntime.java
--- a/graal/com.oracle.graal.hotspot.hsail/src/com/oracle/graal/hotspot/hsail/HSAILHotSpotGraalRuntime.java	Thu Oct 17 01:05:13 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.hotspot.hsail;
-
-import com.oracle.graal.api.code.*;
-import com.oracle.graal.api.meta.*;
-import com.oracle.graal.hotspot.*;
-import com.oracle.graal.hotspot.meta.*;
-import com.oracle.graal.hsail.*;
-import com.oracle.graal.nodes.spi.*;
-
-/**
- * HSAIL specific implementation of {@link HotSpotGraalRuntime}.
- */
-public class HSAILHotSpotGraalRuntime extends HotSpotGraalRuntime {
-
-    @Override
-    protected HotSpotProviders createProviders() {
-        HotSpotProviders host = HotSpotGraalRuntime.runtime().getProviders();
-
-        HotSpotMetaAccessProvider metaAccess = host.getMetaAccess();
-        HSAILHotSpotCodeCacheProvider codeCache = new HSAILHotSpotCodeCacheProvider(this);
-        ConstantReflectionProvider constantReflection = host.getConstantReflection();
-        HotSpotForeignCallsProvider foreignCalls = host.getForeignCalls();
-        LoweringProvider lowerer = new HSAILHotSpotLoweringProvider(host.getLowerer());
-        Replacements replacements = host.getReplacements();
-        HotSpotDisassemblerProvider disassembler = host.getDisassembler();
-        HotSpotSuitesProvider suites = host.getSuites();
-        HotSpotRegisters registers = new HotSpotRegisters(Register.None, Register.None, Register.None);
-        return new HotSpotProviders(metaAccess, codeCache, constantReflection, foreignCalls, lowerer, replacements, disassembler, suites, registers);
-    }
-
-    @Override
-    protected TargetDescription createTarget() {
-        final int stackFrameAlignment = 8;
-        final int implicitNullCheckLimit = 0;
-        final boolean inlineObjects = true;
-        return new TargetDescription(new HSAIL(), true, stackFrameAlignment, implicitNullCheckLimit, inlineObjects);
-    }
-
-    @Override
-    protected HotSpotBackend createBackend() {
-        return new HSAILHotSpotBackend(this, getProviders());
-    }
-
-    @Override
-    protected Value[] getNativeABICallerSaveRegisters() {
-        throw new InternalError("NYI");
-    }
-}
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotBackend.java
--- a/graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotBackend.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotBackend.java	Thu Oct 17 01:08:17 2013 +0200
@@ -22,15 +22,28 @@
  */
 package com.oracle.graal.hotspot.ptx;
 
+import static com.oracle.graal.lir.LIRValueUtil.*;
+
+import java.util.*;
+
 import com.oracle.graal.api.code.*;
 import com.oracle.graal.api.meta.*;
 import com.oracle.graal.asm.*;
+import com.oracle.graal.asm.ptx.*;
 import com.oracle.graal.compiler.gen.*;
+import com.oracle.graal.compiler.ptx.*;
+import com.oracle.graal.graph.*;
 import com.oracle.graal.hotspot.*;
 import com.oracle.graal.hotspot.meta.*;
 import com.oracle.graal.lir.*;
+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.lir.asm.*;
+import com.oracle.graal.lir.ptx.*;
 import com.oracle.graal.nodes.*;
+import com.oracle.graal.nodes.cfg.*;
 
 /**
  * HotSpot PTX specific backend.
@@ -43,31 +56,219 @@
 
     @Override
     public boolean shouldAllocateRegisters() {
-        throw new InternalError("NYI");
+        return false;
     }
 
     @Override
     public FrameMap newFrameMap() {
-        throw new InternalError("NYI");
+        return new PTXFrameMap(getCodeCache());
+    }
+
+    class PTXFrameContext implements FrameContext {
+
+        @Override
+        public void enter(TargetMethodAssembler tasm) {
+            // codeBuffer.emitString(".address_size 32"); // PTX ISA version 2.3
+        }
+
+        @Override
+        public void leave(TargetMethodAssembler tasm) {
+        }
     }
 
     @Override
     public TargetMethodAssembler newAssembler(LIRGenerator lirGen, CompilationResult compilationResult) {
-        throw new InternalError("NYI");
-    }
-
-    @Override
-    public void emitCode(TargetMethodAssembler tasm, LIRGenerator lirGen, ResolvedJavaMethod installedCodeOwner) {
-        throw new InternalError("NYI");
+        // 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
+        // - has no instructions with debug info
+        FrameMap frameMap = lirGen.frameMap;
+        AbstractAssembler masm = createAssembler(frameMap);
+        PTXFrameContext frameContext = new PTXFrameContext();
+        TargetMethodAssembler tasm = new PTXTargetMethodAssembler(getCodeCache(), getForeignCalls(), frameMap, masm, frameContext, compilationResult);
+        tasm.setFrameSize(0);
+        return tasm;
     }
 
     @Override
     protected AbstractAssembler createAssembler(FrameMap frameMap) {
-        throw new InternalError("NYI");
+        return new PTXAssembler(getTarget(), frameMap.registerConfig);
     }
 
     @Override
     public LIRGenerator newLIRGenerator(StructuredGraph graph, FrameMap frameMap, CallingConvention cc, LIR lir) {
-        throw new InternalError("NYI");
+        return new PTXLIRGenerator(graph, getProviders(), frameMap, cc, lir);
+    }
+
+    private static void emitKernelEntry(TargetMethodAssembler tasm, LIRGenerator lirGen, ResolvedJavaMethod codeCacheOwner) {
+        // Emit PTX kernel entry text based on PTXParameterOp
+        // instructions in the start block. Remove the instructions
+        // once kernel entry text and directives are emitted to
+        // facilitate seemless PTX code generation subsequently.
+        assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method";
+        final String name = codeCacheOwner.getName();
+        Buffer codeBuffer = tasm.asm.codeBuffer;
+
+        // Emit initial boiler-plate directives.
+        codeBuffer.emitString(".version 3.0");
+        codeBuffer.emitString(".target sm_30");
+        codeBuffer.emitString0(".entry " + name + " (");
+        codeBuffer.emitString("");
+
+        // Get the start block
+        Block startBlock = lirGen.lir.cfg.getStartBlock();
+        // Keep a list of ParameterOp instructions to delete from the
+        // list of instructions in the block.
+        ArrayList<LIRInstruction> deleteOps = new ArrayList<>();
+
+        // Emit .param arguments to kernel entry based on ParameterOp
+        // instruction.
+        for (LIRInstruction op : lirGen.lir.lir(startBlock)) {
+            if (op instanceof PTXParameterOp) {
+                op.emitCode(tasm);
+                deleteOps.add(op);
+            }
+        }
+
+        // Delete ParameterOp instructions.
+        for (LIRInstruction op : deleteOps) {
+            lirGen.lir.lir(startBlock).remove(op);
+        }
+
+        // Start emiting body of the PTX kernel.
+        codeBuffer.emitString0(") {");
+        codeBuffer.emitString("");
+    }
+
+    // Emit .reg space declarations
+    private static void emitRegisterDecl(TargetMethodAssembler tasm, LIRGenerator lirGen, ResolvedJavaMethod codeCacheOwner) {
+
+        assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method";
+
+        Buffer codeBuffer = tasm.asm.codeBuffer;
+
+        final SortedSet<Integer> signed32 = new TreeSet<>();
+        final SortedSet<Integer> signed64 = new TreeSet<>();
+        final SortedSet<Integer> unsigned64 = new TreeSet<>();
+        final SortedSet<Integer> float32 = new TreeSet<>();
+        final SortedSet<Integer> float64 = new TreeSet<>();
+
+        ValueProcedure trackRegisterKind = new ValueProcedure() {
+
+            @Override
+            public Value doValue(Value value, OperandMode mode, EnumSet<OperandFlag> flags) {
+                if (isVariable(value)) {
+                    Variable regVal = (Variable) value;
+                    Kind regKind = regVal.getKind();
+                    switch (regKind) {
+                        case Int:
+                            // If the register was used as a wider signed type
+                            // do not add it here
+                            if (!signed64.contains(regVal.index)) {
+                                signed32.add(regVal.index);
+                            }
+                            break;
+                        case Long:
+                            // If the register was used as a narrower signed type
+                            // remove it from there and add it to wider type.
+                            if (signed32.contains(regVal.index)) {
+                                signed32.remove(regVal.index);
+                            }
+                            signed64.add(regVal.index);
+                            break;
+                        case Float:
+                            // If the register was used as a wider signed type
+                            // do not add it here
+                            if (!float64.contains(regVal.index)) {
+                                float32.add(regVal.index);
+                            }
+                            break;
+                        case Double:
+                            // If the register was used as a narrower signed type
+                            // remove it from there and add it to wider type.
+                            if (float32.contains(regVal.index)) {
+                                float32.remove(regVal.index);
+                            }
+                            float64.add(regVal.index);
+                            break;
+                        case Object:
+                            unsigned64.add(regVal.index);
+                            break;
+                        default:
+                            throw GraalInternalError.shouldNotReachHere("unhandled register type " + value.toString());
+                    }
+                }
+                return value;
+            }
+        };
+
+        for (Block b : lirGen.lir.codeEmittingOrder()) {
+            for (LIRInstruction op : lirGen.lir.lir(b)) {
+                if (op instanceof LabelOp) {
+                    // Don't consider this as a definition
+                } else {
+                    op.forEachTemp(trackRegisterKind);
+                    op.forEachOutput(trackRegisterKind);
+                }
+            }
+        }
+
+        for (Integer i : signed32) {
+            codeBuffer.emitString(".reg .s32 %r" + i.intValue() + ";");
+        }
+        for (Integer i : signed64) {
+            codeBuffer.emitString(".reg .s64 %r" + i.intValue() + ";");
+        }
+        for (Integer i : unsigned64) {
+            codeBuffer.emitString(".reg .u64 %r" + i.intValue() + ";");
+        }
+        for (Integer i : float32) {
+            codeBuffer.emitString(".reg .f32 %r" + i.intValue() + ";");
+        }
+        for (Integer i : float64) {
+            codeBuffer.emitString(".reg .f64 %r" + i.intValue() + ";");
+        }
+        // emit predicate register declaration
+        int maxPredRegNum = ((PTXLIRGenerator) lirGen).getNextPredRegNumber();
+        if (maxPredRegNum > 0) {
+            codeBuffer.emitString(".reg .pred %p<" + maxPredRegNum + ">;");
+        }
+        codeBuffer.emitString(".reg .pred %r;");  // used for setp bool
+    }
+
+    @Override
+    public void emitCode(TargetMethodAssembler tasm, LIRGenerator lirGen, ResolvedJavaMethod codeCacheOwner) {
+        assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method";
+        Buffer codeBuffer = tasm.asm.codeBuffer;
+        // Emit the prologue
+        emitKernelEntry(tasm, lirGen, codeCacheOwner);
+
+        // Emit register declarations
+        try {
+            emitRegisterDecl(tasm, lirGen, codeCacheOwner);
+        } catch (GraalInternalError e) {
+            e.printStackTrace();
+            // TODO : Better error handling needs to be done once
+            // all types of parameters are handled.
+            codeBuffer.setPosition(0);
+            codeBuffer.close(false);
+            return;
+        }
+        // Emit code for the LIR
+        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);
+            codeBuffer.close(false);
+            return;
+        }
+
+        // Emit the epilogue
+        codeBuffer.emitString0("}");
+        codeBuffer.emitString("");
     }
 }
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotBackendFactory.java
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotBackendFactory.java	Thu Oct 17 01:08:17 2013 +0200
@@ -0,0 +1,75 @@
+/*
+ * 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.hotspot.ptx;
+
+import com.oracle.graal.api.code.*;
+import com.oracle.graal.api.meta.*;
+import com.oracle.graal.api.runtime.*;
+import com.oracle.graal.hotspot.*;
+import com.oracle.graal.hotspot.meta.*;
+import com.oracle.graal.nodes.spi.*;
+import com.oracle.graal.ptx.*;
+
+@ServiceProvider(HotSpotBackendFactory.class)
+public class PTXHotSpotBackendFactory implements HotSpotBackendFactory {
+
+    public HotSpotBackend createBackend(HotSpotGraalRuntime runtime, HotSpotBackend hostBackend) {
+        HotSpotProviders host = hostBackend.getProviders();
+
+        HotSpotMetaAccessProvider metaAccess = host.getMetaAccess();
+        PTXHotSpotCodeCacheProvider codeCache = new PTXHotSpotCodeCacheProvider(runtime, createTarget());
+        ConstantReflectionProvider constantReflection = host.getConstantReflection();
+        HotSpotForeignCallsProvider foreignCalls = new PTXHotSpotForeignCallsProvider();
+        HotSpotLoweringProvider lowerer = new PTXHotSpotLoweringProvider(runtime, metaAccess, foreignCalls);
+        Replacements replacements = host.getReplacements();
+        HotSpotDisassemblerProvider disassembler = host.getDisassembler();
+        HotSpotSuitesProvider suites = host.getSuites();
+        HotSpotRegistersProvider registers = new HotSpotRegisters(PTX.tid, Register.None, Register.None);
+        HotSpotProviders providers = new HotSpotProviders(metaAccess, codeCache, constantReflection, foreignCalls, lowerer, replacements, disassembler, suites, registers);
+        return new PTXHotSpotBackend(runtime, providers);
+    }
+
+    protected Architecture createArchitecture() {
+        return new PTX();
+    }
+
+    protected TargetDescription createTarget() {
+        final int stackFrameAlignment = 1;
+        final int implicitNullCheckLimit = 0;
+        final boolean inlineObjects = true;
+        return new TargetDescription(createArchitecture(), true, stackFrameAlignment, implicitNullCheckLimit, inlineObjects);
+    }
+
+    public String getArchitecture() {
+        return "PTX";
+    }
+
+    public String getGraalRuntimeName() {
+        return "basic";
+    }
+
+    @Override
+    public String toString() {
+        return getGraalRuntimeName() + ":" + getArchitecture();
+    }
+}
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotCodeCacheProvider.java
--- a/graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotCodeCacheProvider.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotCodeCacheProvider.java	Thu Oct 17 01:08:17 2013 +0200
@@ -28,8 +28,8 @@
 
 public class PTXHotSpotCodeCacheProvider extends HotSpotCodeCacheProvider {
 
-    public PTXHotSpotCodeCacheProvider(HotSpotGraalRuntime runtime) {
-        super(runtime);
+    public PTXHotSpotCodeCacheProvider(HotSpotGraalRuntime runtime, TargetDescription target) {
+        super(runtime, target);
 
     }
 
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotForeignCallsProvider.java
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotForeignCallsProvider.java	Thu Oct 17 01:08:17 2013 +0200
@@ -0,0 +1,55 @@
+/*
+ * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved.
+ * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
+ *
+ * This code is free software; you can redistribute it and/or modify it
+ * under the terms of the GNU General Public License version 2 only, as
+ * published by the Free Software Foundation.
+ *
+ * This code is distributed in the hope that it will be useful, but WITHOUT
+ * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+ * FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+ * version 2 for more details (a copy is included in the LICENSE file that
+ * accompanied this code).
+ *
+ * You should have received a copy of the GNU General Public License version
+ * 2 along with this work; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
+ *
+ * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
+ * or visit www.oracle.com if you need additional information or have any
+ * questions.
+ */
+package com.oracle.graal.hotspot.ptx;
+
+import com.oracle.graal.api.code.*;
+import com.oracle.graal.api.meta.*;
+import com.oracle.graal.graph.*;
+import com.oracle.graal.hotspot.meta.*;
+
+public class PTXHotSpotForeignCallsProvider implements HotSpotForeignCallsProvider {
+
+    public boolean isReexecutable(ForeignCallDescriptor descriptor) {
+        throw GraalInternalError.unimplemented();
+    }
+
+    public LocationIdentity[] getKilledLocations(ForeignCallDescriptor descriptor) {
+        throw GraalInternalError.unimplemented();
+    }
+
+    public boolean canDeoptimize(ForeignCallDescriptor descriptor) {
+        throw GraalInternalError.unimplemented();
+    }
+
+    public ForeignCallLinkage lookupForeignCall(ForeignCallDescriptor descriptor) {
+        throw GraalInternalError.unimplemented();
+    }
+
+    public Value[] getNativeABICallerSaveRegisters() {
+        throw GraalInternalError.unimplemented();
+    }
+
+    public void initialize(HotSpotProviders providers) {
+        throw GraalInternalError.unimplemented();
+    }
+}
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotGraalRuntime.java
--- a/graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotGraalRuntime.java	Thu Oct 17 01:05:13 2013 +0200
+++ /dev/null	Thu Jan 01 00:00:00 1970 +0000
@@ -1,96 +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.hotspot.ptx;
-
-import com.oracle.graal.api.code.*;
-import com.oracle.graal.api.meta.*;
-import com.oracle.graal.hotspot.*;
-import com.oracle.graal.hotspot.meta.*;
-import com.oracle.graal.phases.util.*;
-import com.oracle.graal.ptx.*;
-
-/**
- * PTX specific implementation of {@link HotSpotGraalRuntime}.
- */
-public class PTXHotSpotGraalRuntime extends HotSpotGraalRuntime {
-
-    protected PTXHotSpotGraalRuntime() {
-    }
-
-    /**
-     * Called from C++ code to retrieve the singleton instance, creating it first if necessary.
-     */
-    public static HotSpotGraalRuntime makeInstance() {
-        HotSpotGraalRuntime runtime = runtime();
-        if (runtime == null) {
-            HotSpotGraalRuntimeFactory factory = findFactory("PTX");
-            if (factory != null) {
-                runtime = factory.createRuntime();
-            } else {
-                runtime = new PTXHotSpotGraalRuntime();
-            }
-            runtime.completeInitialization();
-        }
-        return runtime;
-    }
-
-    @Override
-    protected HotSpotProviders createProviders() {
-        HotSpotMetaAccessProvider metaAccess = new HotSpotMetaAccessProvider(this);
-        PTXHotSpotCodeCacheProvider codeCache = new PTXHotSpotCodeCacheProvider(this);
-        HotSpotConstantReflectionProvider constantReflection = new HotSpotConstantReflectionProvider(this);
-        HotSpotForeignCallsProvider foreignCalls = new HotSpotForeignCallsProvider(this);
-        HotSpotLoweringProvider lowerer = new PTXHotSpotLoweringProvider(this, metaAccess, foreignCalls);
-        // Replacements cannot have speculative optimizations since they have
-        // to be valid for the entire run of the VM.
-        Assumptions assumptions = new Assumptions(false);
-        Providers p = new Providers(metaAccess, codeCache, constantReflection, foreignCalls, lowerer, null);
-        HotSpotReplacementsImpl replacements = new HotSpotReplacementsImpl(p, getConfig(), assumptions);
-        HotSpotDisassemblerProvider disassembler = new HotSpotDisassemblerProvider(this);
-        HotSpotSuitesProvider suites = new HotSpotSuitesProvider(this);
-        HotSpotRegistersProvider registers = new HotSpotRegisters(PTX.tid, Register.None, Register.None);
-        return new HotSpotProviders(metaAccess, codeCache, constantReflection, foreignCalls, lowerer, replacements, disassembler, suites, registers);
-    }
-
-    protected Architecture createArchitecture() {
-        return new PTX();
-    }
-
-    @Override
-    protected TargetDescription createTarget() {
-        final int stackFrameAlignment = 16;
-        final int implicitNullCheckLimit = 4096;
-        final boolean inlineObjects = true;
-        return new TargetDescription(createArchitecture(), true, stackFrameAlignment, implicitNullCheckLimit, inlineObjects);
-    }
-
-    @Override
-    protected HotSpotBackend createBackend() {
-        return new PTXHotSpotBackend(this, getProviders());
-    }
-
-    @Override
-    protected Value[] getNativeABICallerSaveRegisters() {
-        throw new InternalError("NYI");
-    }
-}
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotBackendFactory.java
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotBackendFactory.java	Thu Oct 17 01:08:17 2013 +0200
@@ -0,0 +1,88 @@
+/*
+ * 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.hotspot.sparc;
+
+import com.oracle.graal.api.code.*;
+import com.oracle.graal.api.meta.*;
+import com.oracle.graal.api.runtime.*;
+import com.oracle.graal.graph.*;
+import com.oracle.graal.hotspot.*;
+import com.oracle.graal.hotspot.meta.*;
+import com.oracle.graal.phases.util.*;
+import com.oracle.graal.sparc.*;
+
+@ServiceProvider(HotSpotBackendFactory.class)
+public class SPARCHotSpotBackendFactory implements HotSpotBackendFactory {
+
+    protected static Architecture createArchitecture() {
+        return new SPARC();
+    }
+
+    protected TargetDescription createTarget() {
+        final int stackFrameAlignment = 16;
+        final int implicitNullCheckLimit = 4096;
+        final boolean inlineObjects = true;
+        return new TargetDescription(createArchitecture(), true, stackFrameAlignment, implicitNullCheckLimit, inlineObjects);
+    }
+
+    public HotSpotBackend createBackend(HotSpotGraalRuntime runtime, HotSpotBackend host) {
+        assert host == null;
+        TargetDescription target = createTarget();
+
+        HotSpotMetaAccessProvider metaAccess = new HotSpotMetaAccessProvider(runtime);
+        HotSpotCodeCacheProvider codeCache = new SPARCHotSpotCodeCacheProvider(runtime, target);
+        HotSpotConstantReflectionProvider constantReflection = new HotSpotConstantReflectionProvider(runtime);
+        Value[] nativeABICallerSaveRegisters = createNativeABICallerSaveRegisters(runtime.getConfig(), codeCache.getRegisterConfig());
+        HotSpotForeignCallsProvider foreignCalls = new SPARCHotSpotForeignCallsProvider(runtime, metaAccess, codeCache, nativeABICallerSaveRegisters);
+        HotSpotLoweringProvider lowerer = new SPARCHotSpotLoweringProvider(runtime, metaAccess, foreignCalls);
+        // Replacements cannot have speculative optimizations since they have
+        // to be valid for the entire run of the VM.
+        Assumptions assumptions = new Assumptions(false);
+        Providers p = new Providers(metaAccess, codeCache, constantReflection, foreignCalls, lowerer, null);
+        HotSpotReplacementsImpl replacements = new HotSpotReplacementsImpl(p, runtime.getConfig(), assumptions);
+        HotSpotDisassemblerProvider disassembler = new HotSpotDisassemblerProvider(runtime);
+        HotSpotSuitesProvider suites = new HotSpotSuitesProvider(runtime);
+        HotSpotRegisters registers = new HotSpotRegisters(Register.None, Register.None, Register.None); // FIXME
+        HotSpotProviders providers = new HotSpotProviders(metaAccess, codeCache, constantReflection, foreignCalls, lowerer, replacements, disassembler, suites, registers);
+
+        return new SPARCHotSpotBackend(runtime, providers);
+    }
+
+    @SuppressWarnings("unused")
+    private static Value[] createNativeABICallerSaveRegisters(HotSpotVMConfig config, RegisterConfig regConfig) {
+        throw GraalInternalError.unimplemented();
+    }
+
+    public String getArchitecture() {
+        return "SPARC";
+    }
+
+    public String getGraalRuntimeName() {
+        return "basic";
+    }
+
+    @Override
+    public String toString() {
+        return getGraalRuntimeName() + ":" + getArchitecture();
+    }
+}
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotCRuntimeCallEpilogueOp.java
--- a/graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotCRuntimeCallEpilogueOp.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotCRuntimeCallEpilogueOp.java	Thu Oct 17 01:08:17 2013 +0200
@@ -23,30 +23,40 @@
 package com.oracle.graal.hotspot.sparc;
 
 import static com.oracle.graal.sparc.SPARC.*;
-import static com.oracle.graal.asm.sparc.SPARCMacroAssembler.*;
-import static com.oracle.graal.hotspot.HotSpotGraalRuntime.*;
 
 import com.oracle.graal.api.code.*;
 import com.oracle.graal.asm.sparc.*;
-import com.oracle.graal.hotspot.*;
+import com.oracle.graal.asm.sparc.SPARCAssembler.Stw;
+import com.oracle.graal.asm.sparc.SPARCAssembler.Stx;
+import com.oracle.graal.asm.sparc.SPARCMacroAssembler.Mov;
 import com.oracle.graal.lir.*;
+import com.oracle.graal.lir.asm.*;
 import com.oracle.graal.lir.sparc.*;
-import com.oracle.graal.lir.asm.*;
 
 @Opcode("CRUNTIME_CALL_EPILOGUE")
 final class SPARCHotSpotCRuntimeCallEpilogueOp extends SPARCLIRInstruction {
 
+    private final int threadLastJavaSpOffset;
+    private final int threadLastJavaPcOffset;
+    private final int threadJavaFrameAnchorFlagsOffset;
+    private final Register thread;
+
+    public SPARCHotSpotCRuntimeCallEpilogueOp(int threadLastJavaSpOffset, int threadLastJavaPcOffset, int threadJavaFrameAnchorFlagsOffset, Register thread) {
+        this.threadLastJavaSpOffset = threadLastJavaSpOffset;
+        this.threadLastJavaPcOffset = threadLastJavaPcOffset;
+        this.threadJavaFrameAnchorFlagsOffset = threadJavaFrameAnchorFlagsOffset;
+        this.thread = thread;
+    }
+
     @Override
     public void emitCode(TargetMethodAssembler tasm, SPARCMacroAssembler masm) {
-        Register thread = runtime().getProviders().getRegisters().getThreadRegister();
-        HotSpotVMConfig config = runtime().getConfig();
 
         // Restore the thread register when coming back from the runtime.
         new Mov(l7, thread).emit(masm);
 
         // Reset last Java frame, last Java PC and flags.
-        new Stx(g0, new SPARCAddress(thread, config.threadLastJavaSpOffset)).emit(masm);
-        new Stx(g0, new SPARCAddress(thread, config.threadLastJavaPcOffset)).emit(masm);
-        new Stw(g0, new SPARCAddress(thread, config.threadJavaFrameAnchorFlagsOffset)).emit(masm);
+        new Stx(g0, new SPARCAddress(thread, threadLastJavaSpOffset)).emit(masm);
+        new Stx(g0, new SPARCAddress(thread, threadLastJavaPcOffset)).emit(masm);
+        new Stw(g0, new SPARCAddress(thread, threadJavaFrameAnchorFlagsOffset)).emit(masm);
     }
 }
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotCRuntimeCallPrologueOp.java
--- a/graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotCRuntimeCallPrologueOp.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotCRuntimeCallPrologueOp.java	Thu Oct 17 01:08:17 2013 +0200
@@ -23,30 +23,35 @@
 package com.oracle.graal.hotspot.sparc;
 
 import static com.oracle.graal.sparc.SPARC.*;
-import static com.oracle.graal.asm.sparc.SPARCMacroAssembler.*;
-import static com.oracle.graal.hotspot.HotSpotGraalRuntime.*;
 
 import com.oracle.graal.api.code.*;
 import com.oracle.graal.asm.sparc.*;
-import com.oracle.graal.hotspot.*;
-import com.oracle.graal.hotspot.meta.*;
+import com.oracle.graal.asm.sparc.SPARCAssembler.Add;
+import com.oracle.graal.asm.sparc.SPARCAssembler.Stx;
+import com.oracle.graal.asm.sparc.SPARCMacroAssembler.Mov;
 import com.oracle.graal.lir.*;
+import com.oracle.graal.lir.asm.*;
 import com.oracle.graal.lir.sparc.*;
-import com.oracle.graal.lir.asm.*;
 
 @Opcode("CRUNTIME_CALL_PROLOGUE")
 final class SPARCHotSpotCRuntimeCallPrologueOp extends SPARCLIRInstruction {
 
+    private final int threadLastJavaSpOffset;
+    private final Register thread;
+    private final Register stackPointer;
+
+    public SPARCHotSpotCRuntimeCallPrologueOp(int threadLastJavaSpOffset, Register thread, Register stackPointer) {
+        this.threadLastJavaSpOffset = threadLastJavaSpOffset;
+        this.thread = thread;
+        this.stackPointer = stackPointer;
+    }
+
     @Override
     public void emitCode(TargetMethodAssembler tasm, SPARCMacroAssembler masm) {
-        HotSpotRegistersProvider registers = runtime().getProviders().getRegisters();
-        HotSpotVMConfig config = runtime().getConfig();
-        Register thread = registers.getThreadRegister();
-        Register stackPointer = registers.getStackPointerRegister();
 
         // Save last Java frame.
         new Add(stackPointer, new SPARCAddress(stackPointer, 0).getDisplacement(), g4).emit(masm);
-        new Stx(g4, new SPARCAddress(thread, config.threadLastJavaSpOffset)).emit(masm);
+        new Stx(g4, new SPARCAddress(thread, threadLastJavaSpOffset)).emit(masm);
 
         // Save the thread register when calling out to the runtime.
         new Mov(thread, l7).emit(masm);
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotCodeCacheProvider.java
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotCodeCacheProvider.java	Thu Oct 17 01:08:17 2013 +0200
@@ -0,0 +1,41 @@
+/*
+ * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved.
+ * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
+ *
+ * This code is free software; you can redistribute it and/or modify it
+ * under the terms of the GNU General Public License version 2 only, as
+ * published by the Free Software Foundation.
+ *
+ * This code is distributed in the hope that it will be useful, but WITHOUT
+ * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+ * FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+ * version 2 for more details (a copy is included in the LICENSE file that
+ * accompanied this code).
+ *
+ * You should have received a copy of the GNU General Public License version
+ * 2 along with this work; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
+ *
+ * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
+ * or visit www.oracle.com if you need additional information or have any
+ * questions.
+ */
+package com.oracle.graal.hotspot.sparc;
+
+import com.oracle.graal.api.code.*;
+import com.oracle.graal.graph.*;
+import com.oracle.graal.hotspot.*;
+import com.oracle.graal.hotspot.meta.*;
+
+public class SPARCHotSpotCodeCacheProvider extends HotSpotCodeCacheProvider {
+
+    public SPARCHotSpotCodeCacheProvider(HotSpotGraalRuntime runtime, TargetDescription target) {
+        super(runtime, target);
+    }
+
+    @Override
+    protected RegisterConfig createRegisterConfig() {
+        throw GraalInternalError.unimplemented();
+    }
+
+}
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotForeignCallsProvider.java
--- a/graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotForeignCallsProvider.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotForeignCallsProvider.java	Thu Oct 17 01:08:17 2013 +0200
@@ -35,10 +35,13 @@
 import com.oracle.graal.hotspot.*;
 import com.oracle.graal.hotspot.meta.*;
 
-public class SPARCHotSpotForeignCallsProvider extends HotSpotForeignCallsProvider {
+public class SPARCHotSpotForeignCallsProvider extends HotSpotHostForeignCallsProvider {
+
+    private final Value[] nativeABICallerSaveRegisters;
 
-    public SPARCHotSpotForeignCallsProvider(HotSpotGraalRuntime runtime) {
-        super(runtime);
+    public SPARCHotSpotForeignCallsProvider(HotSpotGraalRuntime runtime, MetaAccessProvider metaAccess, CodeCacheProvider codeCache, Value[] nativeABICallerSaveRegisters) {
+        super(runtime, metaAccess, codeCache);
+        this.nativeABICallerSaveRegisters = nativeABICallerSaveRegisters;
     }
 
     @Override
@@ -57,4 +60,9 @@
         register(new HotSpotForeignCallLinkage(EXCEPTION_HANDLER, 0L, PRESERVES_REGISTERS, LEAF, outgoingExceptionCc, incomingExceptionCc, NOT_REEXECUTABLE, ANY_LOCATION));
         register(new HotSpotForeignCallLinkage(EXCEPTION_HANDLER_IN_CALLER, JUMP_ADDRESS, PRESERVES_REGISTERS, LEAF, outgoingExceptionCc, incomingExceptionCc, NOT_REEXECUTABLE, ANY_LOCATION));
     }
+
+    @Override
+    public Value[] getNativeABICallerSaveRegisters() {
+        return nativeABICallerSaveRegisters;
+    }
 }
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotGraalRuntime.java
--- a/graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotGraalRuntime.java	Thu Oct 17 01:05:13 2013 +0200
+++ /dev/null	Thu Jan 01 00:00:00 1970 +0000
@@ -1,79 +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.hotspot.sparc;
-
-import com.oracle.graal.api.code.*;
-import com.oracle.graal.api.meta.*;
-import com.oracle.graal.graph.*;
-import com.oracle.graal.hotspot.*;
-import com.oracle.graal.hotspot.meta.*;
-import com.oracle.graal.sparc.*;
-
-/**
- * SPARC specific implementation of {@link HotSpotGraalRuntime}.
- */
-final class SPARCHotSpotGraalRuntime extends HotSpotGraalRuntime {
-
-    private SPARCHotSpotGraalRuntime() {
-    }
-
-    /**
-     * Called from C++ code to retrieve the singleton instance, creating it first if necessary.
-     */
-    public static HotSpotGraalRuntime makeInstance() {
-        HotSpotGraalRuntime runtime = runtime();
-        if (runtime == null) {
-            runtime = new SPARCHotSpotGraalRuntime();
-            runtime.completeInitialization();
-        }
-        return runtime;
-    }
-
-    @Override
-    protected HotSpotProviders createProviders() {
-        // TODO Auto-generated method stub
-        return null;
-    }
-
-    protected static Architecture createArchitecture() {
-        return new SPARC();
-    }
-
-    @Override
-    protected TargetDescription createTarget() {
-        final int stackFrameAlignment = 16;
-        final int implicitNullCheckLimit = 4096;
-        final boolean inlineObjects = true;
-        return new TargetDescription(createArchitecture(), true, stackFrameAlignment, implicitNullCheckLimit, inlineObjects);
-    }
-
-    @Override
-    protected HotSpotBackend createBackend() {
-        return new SPARCHotSpotBackend(this, getProviders());
-    }
-
-    @Override
-    protected Value[] getNativeABICallerSaveRegisters() {
-        throw GraalInternalError.unimplemented();
-    }
-}
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotJumpToExceptionHandlerInCallerOp.java
--- a/graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotJumpToExceptionHandlerInCallerOp.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotJumpToExceptionHandlerInCallerOp.java	Thu Oct 17 01:08:17 2013 +0200
@@ -22,15 +22,19 @@
  */
 package com.oracle.graal.hotspot.sparc;
 
-import static com.oracle.graal.asm.sparc.SPARCMacroAssembler.*;
+import static com.oracle.graal.api.code.ValueUtil.*;
+import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*;
 import static com.oracle.graal.sparc.SPARC.*;
-import static com.oracle.graal.api.code.ValueUtil.*;
-import static com.oracle.graal.hotspot.HotSpotGraalRuntime.*;
-import static com.oracle.graal.lir.LIRInstruction.OperandFlag.*;
 
 import com.oracle.graal.api.code.*;
 import com.oracle.graal.api.meta.*;
 import com.oracle.graal.asm.sparc.*;
+import com.oracle.graal.asm.sparc.SPARCAssembler.CC;
+import com.oracle.graal.asm.sparc.SPARCAssembler.ConditionFlag;
+import com.oracle.graal.asm.sparc.SPARCAssembler.Jmpl;
+import com.oracle.graal.asm.sparc.SPARCAssembler.Lduw;
+import com.oracle.graal.asm.sparc.SPARCAssembler.Movcc;
+import com.oracle.graal.asm.sparc.SPARCMacroAssembler.Cmp;
 import com.oracle.graal.lir.*;
 import com.oracle.graal.lir.asm.*;
 
@@ -44,11 +48,15 @@
     @Use(REG) AllocatableValue handlerInCallerPc;
     @Use(REG) AllocatableValue exception;
     @Use(REG) AllocatableValue exceptionPc;
+    private final Register thread;
+    private final int isMethodHandleReturnOffset;
 
-    SPARCHotSpotJumpToExceptionHandlerInCallerOp(AllocatableValue handlerInCallerPc, AllocatableValue exception, AllocatableValue exceptionPc) {
+    SPARCHotSpotJumpToExceptionHandlerInCallerOp(AllocatableValue handlerInCallerPc, AllocatableValue exception, AllocatableValue exceptionPc, int isMethodHandleReturnOffset, Register thread) {
         this.handlerInCallerPc = handlerInCallerPc;
         this.exception = exception;
         this.exceptionPc = exceptionPc;
+        this.isMethodHandleReturnOffset = isMethodHandleReturnOffset;
+        this.thread = thread;
     }
 
     @Override
@@ -56,8 +64,6 @@
         leaveFrame(tasm);
 
         // Restore SP from L7 if the exception PC is a method handle call site.
-        Register thread = runtime().getProviders().getRegisters().getThreadRegister();
-        int isMethodHandleReturnOffset = runtime().getConfig().threadIsMethodHandleReturnOffset;
         SPARCAddress dst = new SPARCAddress(thread, isMethodHandleReturnOffset);
         new Lduw(dst, o7).emit(masm);
         new Cmp(o7, o7).emit(masm);
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotLIRGenerator.java
--- a/graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotLIRGenerator.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.sparc/src/com/oracle/graal/hotspot/sparc/SPARCHotSpotLIRGenerator.java	Thu Oct 17 01:08:17 2013 +0200
@@ -100,9 +100,12 @@
 
         if (linkage.canDeoptimize()) {
             assert info != null;
-            append(new SPARCHotSpotCRuntimeCallPrologueOp());
+            HotSpotRegistersProvider registers = getProviders().getRegisters();
+            Register thread = registers.getThreadRegister();
+            Register stackPointer = registers.getStackPointerRegister();
+            append(new SPARCHotSpotCRuntimeCallPrologueOp(config.threadLastJavaSpOffset, thread, stackPointer));
             result = super.emitForeignCall(linkage, info, args);
-            append(new SPARCHotSpotCRuntimeCallEpilogueOp());
+            append(new SPARCHotSpotCRuntimeCallEpilogueOp(config.threadLastJavaSpOffset, config.threadLastJavaPcOffset, config.threadJavaFrameAnchorFlagsOffset, thread));
         } else {
             result = super.emitForeignCall(linkage, null, args);
         }
@@ -223,7 +226,8 @@
         RegisterValue exceptionPcFixed = (RegisterValue) linkageCc.getArgument(1);
         emitMove(exceptionFixed, operand(exception));
         emitMove(exceptionPcFixed, operand(exceptionPc));
-        SPARCHotSpotJumpToExceptionHandlerInCallerOp op = new SPARCHotSpotJumpToExceptionHandlerInCallerOp(handler, exceptionFixed, exceptionPcFixed);
+        Register thread = getProviders().getRegisters().getThreadRegister();
+        SPARCHotSpotJumpToExceptionHandlerInCallerOp op = new SPARCHotSpotJumpToExceptionHandlerInCallerOp(handler, exceptionFixed, exceptionPcFixed, config.threadIsMethodHandleReturnOffset, thread);
         append(op);
     }
 
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot.test/src/com/oracle/graal/hotspot/test/AheadOfTimeCompilationTest.java
--- a/graal/com.oracle.graal.hotspot.test/src/com/oracle/graal/hotspot/test/AheadOfTimeCompilationTest.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot.test/src/com/oracle/graal/hotspot/test/AheadOfTimeCompilationTest.java	Thu Oct 17 01:08:17 2013 +0200
@@ -205,7 +205,7 @@
         CallingConvention cc = getCallingConvention(getCodeCache(), Type.JavaCallee, graph.method(), false);
         // create suites everytime, as we modify options for the compiler
         final Suites suitesLocal = Graal.getRequiredCapability(SuitesProvider.class).createSuites();
-        final CompilationResult compResult = GraalCompiler.compileGraph(graph, cc, method, getProviders(), backend, getCodeCache().getTarget(), null, phasePlan, OptimisticOptimizations.ALL,
+        final CompilationResult compResult = GraalCompiler.compileGraph(graph, cc, method, getProviders(), getBackend(), getCodeCache().getTarget(), null, phasePlan, OptimisticOptimizations.ALL,
                         new SpeculationLog(), suitesLocal, new CompilationResult());
         addMethod(method, compResult);
 
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/CompilationTask.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/CompilationTask.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/CompilationTask.java	Thu Oct 17 01:08:17 2013 +0200
@@ -42,7 +42,6 @@
 import com.oracle.graal.phases.*;
 import com.oracle.graal.phases.common.*;
 import com.oracle.graal.phases.tiers.*;
-import com.oracle.graal.phases.util.*;
 
 public final class CompilationTask implements Runnable {
 
@@ -58,9 +57,8 @@
         Queued, Running
     }
 
-    private final HotSpotGraalRuntime runtime;
+    private final HotSpotBackend backend;
     private final PhasePlan plan;
-    private final SuitesProvider suitesProvider;
     private final OptimisticOptimizations optimisticOpts;
     private final HotSpotResolvedJavaMethod method;
     private final int entryBCI;
@@ -70,14 +68,13 @@
     private StructuredGraph graph;
 
     public static CompilationTask create(HotSpotGraalRuntime runtime, PhasePlan plan, OptimisticOptimizations optimisticOpts, HotSpotResolvedJavaMethod method, int entryBCI, int id) {
-        return new CompilationTask(runtime, plan, optimisticOpts, method, entryBCI, id);
+        return new CompilationTask(runtime.getHostBackend(), plan, optimisticOpts, method, entryBCI, id);
     }
 
-    private CompilationTask(HotSpotGraalRuntime runtime, PhasePlan plan, OptimisticOptimizations optimisticOpts, HotSpotResolvedJavaMethod method, int entryBCI, int id) {
+    private CompilationTask(HotSpotBackend backend, PhasePlan plan, OptimisticOptimizations optimisticOpts, HotSpotResolvedJavaMethod method, int entryBCI, int id) {
         assert id >= 0;
-        this.runtime = runtime;
+        this.backend = backend;
         this.plan = plan;
-        this.suitesProvider = runtime.getCapability(SuitesProvider.class);
         this.method = method;
         this.optimisticOpts = optimisticOpts;
         this.entryBCI = entryBCI;
@@ -143,8 +140,8 @@
 
                     @Override
                     public CompilationResult call() throws Exception {
-                        runtime.evictDeoptedGraphs();
-                        Providers providers = runtime.getProviders();
+                        backend.getRuntime().evictDeoptedGraphs();
+                        HotSpotProviders providers = backend.getProviders();
                         Replacements replacements = providers.getReplacements();
                         graph = replacements.getMethodSubstitution(method);
                         if (graph == null || entryBCI != INVOCATION_ENTRY_BCI) {
@@ -155,8 +152,9 @@
                         }
                         InliningUtil.InlinedBytecodes.add(method.getCodeSize());
                         CallingConvention cc = getCallingConvention(providers.getCodeCache(), Type.JavaCallee, graph.method(), false);
-                        return GraalCompiler.compileGraph(graph, cc, method, providers, runtime.getBackend(), runtime.getTarget(), runtime.getCache(), plan, optimisticOpts,
-                                        method.getSpeculationLog(), suitesProvider.getDefaultSuites(), new CompilationResult());
+                        Suites suites = providers.getSuites().getDefaultSuites();
+                        return GraalCompiler.compileGraph(graph, cc, method, providers, backend, backend.getTarget(), backend.getRuntime().getCache(), plan, optimisticOpts,
+                                        method.getSpeculationLog(), suites, new CompilationResult());
                     }
                 });
             } finally {
@@ -212,7 +210,7 @@
     }
 
     private void installMethod(final CompilationResult compResult) {
-        final HotSpotCodeCacheProvider codeCache = runtime.getProviders().getCodeCache();
+        final HotSpotCodeCacheProvider codeCache = backend.getProviders().getCodeCache();
         Debug.scope("CodeInstall", new Object[]{new DebugDumpScope(String.valueOf(id), true), codeCache, method}, new Runnable() {
 
             @Override
@@ -222,7 +220,7 @@
                     Debug.dump(new Object[]{compResult, installedCode}, "After code installation");
                 }
                 if (Debug.isLogEnabled()) {
-                    Debug.log("%s", runtime.getProviders().getDisassembler().disassemble(installedCode));
+                    Debug.log("%s", backend.getProviders().getDisassembler().disassemble(installedCode));
                 }
             }
 
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/CompileTheWorld.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/CompileTheWorld.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/CompileTheWorld.java	Thu Oct 17 01:08:17 2013 +0200
@@ -182,7 +182,7 @@
                     }
 
                     // Are we compiling this class?
-                    HotSpotMetaAccessProvider metaAccess = runtime.getProviders().getMetaAccess();
+                    HotSpotMetaAccessProvider metaAccess = runtime.getHostProviders().getMetaAccess();
                     if (classFileCounter >= startAt) {
                         TTY.println("CompileTheWorld (%d) : %s", classFileCounter, className);
 
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/HotSpotBackendFactory.java
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/HotSpotBackendFactory.java	Thu Oct 17 01:08:17 2013 +0200
@@ -0,0 +1,42 @@
+/*
+ * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved.
+ * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
+ *
+ * This code is free software; you can redistribute it and/or modify it
+ * under the terms of the GNU General Public License version 2 only, as
+ * published by the Free Software Foundation.
+ *
+ * This code is distributed in the hope that it will be useful, but WITHOUT
+ * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+ * FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+ * version 2 for more details (a copy is included in the LICENSE file that
+ * accompanied this code).
+ *
+ * You should have received a copy of the GNU General Public License version
+ * 2 along with this work; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
+ *
+ * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
+ * or visit www.oracle.com if you need additional information or have any
+ * questions.
+ */
+package com.oracle.graal.hotspot;
+
+import com.oracle.graal.hotspot.HotSpotGraalRuntime.Options;
+
+public interface HotSpotBackendFactory {
+
+    HotSpotBackend createBackend(HotSpotGraalRuntime runtime, HotSpotBackend host);
+
+    /**
+     * Gets the CPU architecture of this backend.
+     */
+    String getArchitecture();
+
+    /**
+     * Gets the name of the {@link Options#GraalRuntime GraalRuntime} in which the backend created
+     * by this factory should be used.
+     */
+    String getGraalRuntimeName();
+
+}
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/HotSpotForeignCallLinkage.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/HotSpotForeignCallLinkage.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/HotSpotForeignCallLinkage.java	Thu Oct 17 01:08:17 2013 +0200
@@ -31,6 +31,7 @@
 import com.oracle.graal.api.code.CallingConvention.Type;
 import com.oracle.graal.api.meta.*;
 import com.oracle.graal.compiler.target.*;
+import com.oracle.graal.hotspot.meta.*;
 import com.oracle.graal.hotspot.stubs.*;
 import com.oracle.graal.word.*;
 
@@ -119,13 +120,13 @@
      *            re-executed.
      * @param killedLocations the memory locations killed by the call
      */
-    public static HotSpotForeignCallLinkage create(ForeignCallDescriptor descriptor, long address, RegisterEffect effect, Type outgoingCcType, Type incomingCcType, Transition transition,
-                    boolean reexecutable, LocationIdentity... killedLocations) {
-        CallingConvention outgoingCc = createCallingConvention(descriptor, outgoingCcType);
-        CallingConvention incomingCc = incomingCcType == null ? null : createCallingConvention(descriptor, incomingCcType);
+    public static HotSpotForeignCallLinkage create(MetaAccessProvider metaAccess, CodeCacheProvider codeCache, HotSpotForeignCallsProvider foreignCalls, ForeignCallDescriptor descriptor,
+                    long address, RegisterEffect effect, Type outgoingCcType, Type incomingCcType, Transition transition, boolean reexecutable, LocationIdentity... killedLocations) {
+        CallingConvention outgoingCc = createCallingConvention(metaAccess, codeCache, descriptor, outgoingCcType);
+        CallingConvention incomingCc = incomingCcType == null ? null : createCallingConvention(metaAccess, codeCache, descriptor, incomingCcType);
         HotSpotForeignCallLinkage linkage = new HotSpotForeignCallLinkage(descriptor, address, effect, transition, outgoingCc, incomingCc, reexecutable, killedLocations);
         if (outgoingCcType == Type.NativeCall) {
-            linkage.temporaries = runtime().getNativeABICallerSaveRegisters();
+            linkage.temporaries = foreignCalls.getNativeABICallerSaveRegisters();
         }
         return linkage;
     }
@@ -133,17 +134,16 @@
     /**
      * Gets a calling convention for a given descriptor and call type.
      */
-    public static CallingConvention createCallingConvention(ForeignCallDescriptor descriptor, Type ccType) {
+    public static CallingConvention createCallingConvention(MetaAccessProvider metaAccess, CodeCacheProvider codeCache, ForeignCallDescriptor descriptor, Type ccType) {
         assert ccType != null;
-        MetaAccessProvider metaAccess = runtime().getProviders().getMetaAccess();
         Class<?>[] argumentTypes = descriptor.getArgumentTypes();
         JavaType[] parameterTypes = new JavaType[argumentTypes.length];
         for (int i = 0; i < parameterTypes.length; ++i) {
             parameterTypes[i] = asJavaType(argumentTypes[i], metaAccess);
         }
-        TargetDescription target = runtime().getTarget();
+        TargetDescription target = codeCache.getTarget();
         JavaType returnType = asJavaType(descriptor.getResultType(), metaAccess);
-        RegisterConfig regConfig = runtime().getProviders().getCodeCache().getRegisterConfig();
+        RegisterConfig regConfig = codeCache.getRegisterConfig();
         return regConfig.getCallingConvention(ccType, returnType, parameterTypes, target, false);
     }
 
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/HotSpotGraalRuntime.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/HotSpotGraalRuntime.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/HotSpotGraalRuntime.java	Thu Oct 17 01:08:17 2013 +0200
@@ -45,17 +45,19 @@
 
 /**
  * Singleton class holding the instance of the {@link GraalRuntime}.
- * 
- * The platform specific subclass is created by a call from the C++ HotSpot code.
  */
-public abstract class HotSpotGraalRuntime implements GraalRuntime {
+public final class HotSpotGraalRuntime implements GraalRuntime {
 
-    private static final HotSpotGraalRuntime instance = (HotSpotGraalRuntime) Graal.getRuntime();
+    private static final HotSpotGraalRuntime instance = new HotSpotGraalRuntime();
+    static {
+        instance.completeInitialization();
+    }
 
     /**
      * Gets the singleton {@link HotSpotGraalRuntime} object.
      */
     public static HotSpotGraalRuntime runtime() {
+        assert instance != null;
         return instance;
     }
 
@@ -86,7 +88,7 @@
         this.compilerToVm = toVM;
     }
 
-    // Options must not be directly declared in HotSpotGraalRuntime - see VerifyHotSpotOptionsPhase
+    // Options must not be directly declared in HotSpotGraalRuntime - see VerifyOptionsPhase
     static class Options {
 
         // @formatter:off
@@ -95,19 +97,19 @@
         // @formatter:on
     }
 
-    protected static HotSpotGraalRuntimeFactory findFactory(String architecture) {
-        HotSpotGraalRuntimeFactory basic = null;
-        HotSpotGraalRuntimeFactory selected = null;
-        HotSpotGraalRuntimeFactory nonBasic = null;
+    private static HotSpotBackendFactory findFactory(String architecture) {
+        HotSpotBackendFactory basic = null;
+        HotSpotBackendFactory selected = null;
+        HotSpotBackendFactory nonBasic = null;
         int nonBasicCount = 0;
 
-        for (HotSpotGraalRuntimeFactory factory : ServiceLoader.loadInstalled(HotSpotGraalRuntimeFactory.class)) {
-            if (factory.getArchitecture().equals(architecture)) {
-                if (factory.getName().equals(GraalRuntime.getValue())) {
+        for (HotSpotBackendFactory factory : ServiceLoader.loadInstalled(HotSpotBackendFactory.class)) {
+            if (factory.getArchitecture().equalsIgnoreCase(architecture)) {
+                if (factory.getGraalRuntimeName().equals(GraalRuntime.getValue())) {
                     assert selected == null;
                     selected = factory;
                 }
-                if (factory.getName().equals("basic")) {
+                if (factory.getGraalRuntimeName().equals("basic")) {
                     assert basic == null;
                     basic = factory;
                 } else {
@@ -174,21 +176,19 @@
     protected/* final */CompilerToGPU compilerToGpu;
     protected/* final */VMToCompiler vmToCompiler;
 
-    protected final HotSpotProviders providers;
-
-    protected final TargetDescription target;
+    protected final HotSpotProviders hostProviders;
 
     private HotSpotRuntimeInterpreterInterface runtimeInterpreterInterface;
     private volatile HotSpotGraphCache cache;
 
     protected final HotSpotVMConfig config;
-    private final HotSpotBackend backend;
+    private final HotSpotBackend hostBackend;
 
-    protected HotSpotGraalRuntime() {
+    private final Map<String, HotSpotBackend> backends = new HashMap<>();
+
+    private HotSpotGraalRuntime() {
         CompilerToVM toVM = new CompilerToVMImpl();
         CompilerToGPU toGPU = new CompilerToGPUImpl();
-
-        // initialize VmToCompiler
         VMToCompiler toCompiler = new VMToCompilerImpl(this);
 
         compilerToVm = toVM;
@@ -213,18 +213,43 @@
             printConfig(config);
         }
 
-        target = createTarget();
-        providers = createProviders();
-        assert wordKind == null || wordKind.equals(target.wordKind);
-        wordKind = target.wordKind;
+        String hostArchitecture = System.getProperty("os.arch");
+        hostBackend = findFactory(hostArchitecture).createBackend(this, null);
+        hostProviders = hostBackend.getProviders();
+        backends.put(hostArchitecture, hostBackend);
 
-        backend = createBackend();
+        String[] gpuArchitectures = getGPUArchitectures();
+        for (String arch : gpuArchitectures) {
+            HotSpotBackendFactory factory = findFactory(arch);
+            if (factory == null) {
+                throw new GraalInternalError("No backend available for specified GPU architecture \"%s\"", arch);
+            }
+            backends.put(factory.getArchitecture(), factory.createBackend(this, hostBackend));
+        }
+
+        assert wordKind == null || wordKind.equals(hostBackend.getTarget().wordKind);
+        wordKind = hostBackend.getTarget().wordKind;
+
         GraalOptions.StackShadowPages.setValue(config.stackShadowPages);
         if (GraalOptions.CacheGraphs.getValue()) {
             cache = new HotSpotGraphCache();
         }
     }
 
+    /**
+     * Gets the supported GPUs. This method first looks for a comma separated list of names in the
+     * "graal.gpu.isalist" system property. If this property is not set, then the GPU native support
+     * code is queried.
+     */
+    private String[] getGPUArchitectures() {
+        String gpuList = System.getProperty("graal.gpu.isalist");
+        if (gpuList != null) {
+            String[] gpus = gpuList.split(",");
+            return gpus;
+        }
+        return compilerToGpu.getAvailableGPUArchitectures();
+    }
+
     private static void printConfig(HotSpotVMConfig config) {
         Field[] fields = config.getClass().getDeclaredFields();
         Map<String, Field> sortedFields = new TreeMap<>();
@@ -240,23 +265,12 @@
         }
     }
 
-    protected abstract HotSpotProviders createProviders();
-
-    protected abstract TargetDescription createTarget();
-
-    protected abstract HotSpotBackend createBackend();
-
-    /**
-     * Gets the registers that must be saved across a foreign call into the runtime.
-     */
-    protected abstract Value[] getNativeABICallerSaveRegisters();
-
     public HotSpotVMConfig getConfig() {
         return config;
     }
 
     public TargetDescription getTarget() {
-        return target;
+        return hostBackend.getTarget();
     }
 
     public HotSpotGraphCache getCache() {
@@ -309,13 +323,13 @@
 
     public HotSpotRuntimeInterpreterInterface getRuntimeInterpreterInterface() {
         if (runtimeInterpreterInterface == null) {
-            runtimeInterpreterInterface = new HotSpotRuntimeInterpreterInterface(providers.getMetaAccess());
+            runtimeInterpreterInterface = new HotSpotRuntimeInterpreterInterface(hostProviders.getMetaAccess());
         }
         return runtimeInterpreterInterface;
     }
 
-    public HotSpotProviders getProviders() {
-        return providers;
+    public HotSpotProviders getHostProviders() {
+        return hostProviders;
     }
 
     public void evictDeoptedGraphs() {
@@ -337,8 +351,8 @@
     }
 
     @SuppressWarnings("unchecked")
-    @Override
-    public <T> T getCapability(Class<T> clazz) {
+    public static <T> T getCapability(HotSpotBackend backend, Class<T> clazz) {
+        HotSpotProviders providers = backend.getProviders();
         if (clazz == LoweringProvider.class) {
             return (T) providers.getLowerer();
         }
@@ -370,13 +384,28 @@
             return (T) providers.getRegisters();
         }
         if (clazz == Backend.class) {
-            return (T) getBackend();
+            return (T) backend;
         }
         return null;
     }
 
-    public HotSpotBackend getBackend() {
-        return backend;
+    @Override
+    public <T> T getCapability(Class<T> clazz) {
+        return getCapability(clazz, null);
+    }
+
+    @Override
+    public <T> T getCapability(Class<T> clazz, String selector) {
+        HotSpotBackend backend = selector == null ? hostBackend : backends.get(selector);
+        return backend == null ? null : getCapability(backend, clazz);
+    }
+
+    public HotSpotBackend getHostBackend() {
+        return hostBackend;
+    }
+
+    public Map<String, HotSpotBackend> getBackends() {
+        return Collections.unmodifiableMap(backends);
     }
 
     /**
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/HotSpotGraalRuntimeFactory.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/HotSpotGraalRuntimeFactory.java	Thu Oct 17 01:05:13 2013 +0200
+++ /dev/null	Thu Jan 01 00:00:00 1970 +0000
@@ -1,32 +0,0 @@
-/*
- * Copyright (c) 2011, 2012, Oracle and/or its affiliates. All rights reserved.
- * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
- *
- * This code is free software; you can redistribute it and/or modify it
- * under the terms of the GNU General Public License version 2 only, as
- * published by the Free Software Foundation.
- *
- * This code is distributed in the hope that it will be useful, but WITHOUT
- * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
- * FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
- * version 2 for more details (a copy is included in the LICENSE file that
- * accompanied this code).
- *
- * You should have received a copy of the GNU General Public License version
- * 2 along with this work; if not, write to the Free Software Foundation,
- * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
- *
- * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
- * or visit www.oracle.com if you need additional information or have any
- * questions.
- */
-package com.oracle.graal.hotspot;
-
-public interface HotSpotGraalRuntimeFactory {
-
-    HotSpotGraalRuntime createRuntime();
-
-    String getArchitecture();
-
-    String getName();
-}
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/CompilerToGPU.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/CompilerToGPU.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/CompilerToGPU.java	Thu Oct 17 01:08:17 2013 +0200
@@ -48,8 +48,12 @@
     int availableProcessors();
 
     /**
-     * Attempts to generate and return a bound function to the
-     * loaded method kernel on the GPU.
+     * Gets the architecture names of the available GPUs.
+     */
+    String[] getAvailableGPUArchitectures();
+
+    /**
+     * Attempts to generate and return a bound function to the loaded method kernel on the GPU.
      * 
      * @param code the text or binary values for a method kernel
      * @return the value of the bound kernel in GPU space.
@@ -58,7 +62,5 @@
 
     Object executeExternalMethodVarargs(Object[] args, HotSpotInstalledCode hotspotInstalledCode) throws InvalidInstalledCodeException;
 
-
-    Object executeParallelMethodVarargs(int dimX, int dimY, int dimZ,
-                                        Object[] args, HotSpotInstalledCode hotspotInstalledCode) throws InvalidInstalledCodeException;
+    Object executeParallelMethodVarargs(int dimX, int dimY, int dimZ, Object[] args, HotSpotInstalledCode hotspotInstalledCode) throws InvalidInstalledCodeException;
 }
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/CompilerToGPUImpl.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/CompilerToGPUImpl.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/CompilerToGPUImpl.java	Thu Oct 17 01:08:17 2013 +0200
@@ -26,7 +26,6 @@
 import com.oracle.graal.api.code.InvalidInstalledCodeException;
 import com.oracle.graal.hotspot.meta.HotSpotInstalledCode;
 
-
 /**
  * Entries into the HotSpot GPU interface from Java code.
  */
@@ -40,10 +39,9 @@
 
     public native int availableProcessors();
 
-    public native Object executeExternalMethodVarargs(Object[] args,
-                                                      HotSpotInstalledCode hotspotInstalledCode) throws InvalidInstalledCodeException;
+    public native String[] getAvailableGPUArchitectures();
 
-    public native Object executeParallelMethodVarargs(int dimX, int dimY, int dimZ,
-                                                      Object[] args,
-                                                      HotSpotInstalledCode hotspotInstalledCode) throws InvalidInstalledCodeException;
+    public native Object executeExternalMethodVarargs(Object[] args, HotSpotInstalledCode hotspotInstalledCode) throws InvalidInstalledCodeException;
+
+    public native Object executeParallelMethodVarargs(int dimX, int dimY, int dimZ, Object[] args, HotSpotInstalledCode hotspotInstalledCode) throws InvalidInstalledCodeException;
 }
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/VMToCompilerImpl.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/VMToCompilerImpl.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/bridge/VMToCompilerImpl.java	Thu Oct 17 01:08:17 2013 +0200
@@ -180,7 +180,7 @@
             }
         }
 
-        final HotSpotProviders providers = runtime.getProviders();
+        final HotSpotProviders providers = runtime.getHostProviders();
         final MetaAccessProvider metaAccess = providers.getMetaAccess();
         assert VerifyOptionsPhase.checkOptions(metaAccess, providers.getForeignCalls());
 
@@ -190,6 +190,7 @@
 
                 @Override
                 public void run() {
+
                     final Replacements replacements = providers.getReplacements();
                     ServiceLoader<ReplacementsProvider> serviceLoader = ServiceLoader.loadInstalled(ReplacementsProvider.class);
                     TargetDescription target = providers.getCodeCache().getTarget();
@@ -350,7 +351,7 @@
     private MetricRateInPhase inlinedBytecodesPerSecond;
 
     private void enqueue(Method m) throws Throwable {
-        JavaMethod javaMethod = runtime.getProviders().getMetaAccess().lookupJavaMethod(m);
+        JavaMethod javaMethod = runtime.getHostProviders().getMetaAccess().lookupJavaMethod(m);
         assert !Modifier.isAbstract(((HotSpotResolvedJavaMethod) javaMethod).getModifiers()) && !Modifier.isNative(((HotSpotResolvedJavaMethod) javaMethod).getModifiers()) : javaMethod;
         compileMethod((HotSpotResolvedJavaMethod) javaMethod, StructuredGraph.INVOCATION_ENTRY_BCI, false);
     }
@@ -692,8 +693,8 @@
 
     public PhasePlan createPhasePlan(OptimisticOptimizations optimisticOpts, boolean onStackReplacement) {
         PhasePlan phasePlan = new PhasePlan();
-        MetaAccessProvider metaAccess = runtime.getProviders().getMetaAccess();
-        ForeignCallsProvider foreignCalls = runtime.getProviders().getForeignCalls();
+        MetaAccessProvider metaAccess = runtime.getHostProviders().getMetaAccess();
+        ForeignCallsProvider foreignCalls = runtime.getHostProviders().getForeignCalls();
         phasePlan.addPhase(PhasePosition.AFTER_PARSING, new GraphBuilderPhase(metaAccess, foreignCalls, GraphBuilderConfiguration.getDefault(), optimisticOpts));
         if (onStackReplacement) {
             phasePlan.addPhase(PhasePosition.AFTER_PARSING, new OnStackReplacementPhase());
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotCodeCacheProvider.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotCodeCacheProvider.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotCodeCacheProvider.java	Thu Oct 17 01:08:17 2013 +0200
@@ -44,10 +44,12 @@
 public abstract class HotSpotCodeCacheProvider implements CodeCacheProvider {
 
     protected final HotSpotGraalRuntime runtime;
+    protected final TargetDescription target;
     protected final RegisterConfig regConfig;
 
-    public HotSpotCodeCacheProvider(HotSpotGraalRuntime runtime) {
+    public HotSpotCodeCacheProvider(HotSpotGraalRuntime runtime, TargetDescription target) {
         this.runtime = runtime;
+        this.target = target;
         regConfig = createRegisterConfig();
     }
 
@@ -57,7 +59,6 @@
     public String disassemble(CompilationResult compResult, InstalledCode installedCode) {
         byte[] code = installedCode == null ? Arrays.copyOf(compResult.getTargetCode(), compResult.getTargetCodeSize()) : installedCode.getCode();
         long start = installedCode == null ? 0L : installedCode.getStart();
-        TargetDescription target = runtime.getTarget();
         HexCodeFile hcf = new HexCodeFile(code, start, target.arch.getName(), target.wordSize * 8);
         if (compResult != null) {
             HexCodeFile.addAnnotations(hcf, compResult.getAnnotations());
@@ -191,7 +192,7 @@
 
     @Override
     public TargetDescription getTarget() {
-        return runtime.getTarget();
+        return target;
     }
 
     public String disassemble(InstalledCode code) {
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotForeignCallsProvider.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotForeignCallsProvider.java	Thu Oct 17 01:05:13 2013 +0200
+++ /dev/null	Thu Jan 01 00:00:00 1970 +0000
@@ -1,223 +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.hotspot.meta;
-
-import static com.oracle.graal.api.code.CallingConvention.Type.*;
-import static com.oracle.graal.api.meta.LocationIdentity.*;
-import static com.oracle.graal.hotspot.HotSpotBackend.*;
-import static com.oracle.graal.hotspot.HotSpotForeignCallLinkage.RegisterEffect.*;
-import static com.oracle.graal.hotspot.HotSpotForeignCallLinkage.Transition.*;
-import static com.oracle.graal.hotspot.nodes.MonitorExitStubCall.*;
-import static com.oracle.graal.hotspot.nodes.NewArrayStubCall.*;
-import static com.oracle.graal.hotspot.nodes.NewInstanceStubCall.*;
-import static com.oracle.graal.hotspot.nodes.NewMultiArrayStubCall.*;
-import static com.oracle.graal.hotspot.nodes.VMErrorNode.*;
-import static com.oracle.graal.hotspot.replacements.HotSpotReplacementsUtil.*;
-import static com.oracle.graal.hotspot.replacements.MonitorSnippets.*;
-import static com.oracle.graal.hotspot.replacements.NewObjectSnippets.*;
-import static com.oracle.graal.hotspot.replacements.SystemSubstitutions.*;
-import static com.oracle.graal.hotspot.replacements.ThreadSubstitutions.*;
-import static com.oracle.graal.hotspot.replacements.WriteBarrierSnippets.*;
-import static com.oracle.graal.hotspot.stubs.ExceptionHandlerStub.*;
-import static com.oracle.graal.hotspot.stubs.NewArrayStub.*;
-import static com.oracle.graal.hotspot.stubs.NewInstanceStub.*;
-import static com.oracle.graal.hotspot.stubs.StubUtil.*;
-import static com.oracle.graal.hotspot.stubs.UnwindExceptionToCallerStub.*;
-import static com.oracle.graal.java.GraphBuilderPhase.RuntimeCalls.*;
-import static com.oracle.graal.nodes.java.RegisterFinalizerNode.*;
-import static com.oracle.graal.replacements.Log.*;
-import static com.oracle.graal.replacements.MathSubstitutionsX86.*;
-
-import java.util.*;
-
-import com.oracle.graal.api.code.*;
-import com.oracle.graal.api.meta.*;
-import com.oracle.graal.hotspot.*;
-import com.oracle.graal.hotspot.HotSpotForeignCallLinkage.RegisterEffect;
-import com.oracle.graal.hotspot.HotSpotForeignCallLinkage.Transition;
-import com.oracle.graal.hotspot.stubs.*;
-import com.oracle.graal.word.*;
-
-/**
- * HotSpot implementation of {@link ForeignCallsProvider}.
- */
-public class HotSpotForeignCallsProvider implements ForeignCallsProvider {
-
-    public static final ForeignCallDescriptor OSR_MIGRATION_END = new ForeignCallDescriptor("OSR_migration_end", void.class, long.class);
-    public static final ForeignCallDescriptor IDENTITY_HASHCODE = new ForeignCallDescriptor("identity_hashcode", int.class, Object.class);
-    public static final ForeignCallDescriptor VERIFY_OOP = new ForeignCallDescriptor("verify_oop", Object.class, Object.class);
-    public static final ForeignCallDescriptor LOAD_AND_CLEAR_EXCEPTION = new ForeignCallDescriptor("load_and_clear_exception", Object.class, Word.class);
-
-    protected final HotSpotGraalRuntime runtime;
-
-    private final Map<ForeignCallDescriptor, HotSpotForeignCallLinkage> foreignCalls = new HashMap<>();
-
-    public HotSpotForeignCallsProvider(HotSpotGraalRuntime runtime) {
-        this.runtime = runtime;
-    }
-
-    /**
-     * Registers the linkage for a foreign call.
-     */
-    protected HotSpotForeignCallLinkage register(HotSpotForeignCallLinkage linkage) {
-        assert !foreignCalls.containsKey(linkage.getDescriptor()) : "already registered linkage for " + linkage.getDescriptor();
-        foreignCalls.put(linkage.getDescriptor(), linkage);
-        return linkage;
-    }
-
-    /**
-     * Creates and registers the details for linking a foreign call to a {@link Stub}.
-     * 
-     * @param descriptor the signature of the call to the stub
-     * @param reexecutable specifies if the stub call can be re-executed without (meaningful) side
-     *            effects. Deoptimization will not return to a point before a stub call that cannot
-     *            be re-executed.
-     * @param transition specifies if this is a {@linkplain Transition#LEAF leaf} call
-     * @param killedLocations the memory locations killed by the stub call
-     */
-    protected HotSpotForeignCallLinkage registerStubCall(ForeignCallDescriptor descriptor, boolean reexecutable, Transition transition, LocationIdentity... killedLocations) {
-        return register(HotSpotForeignCallLinkage.create(descriptor, 0L, PRESERVES_REGISTERS, JavaCall, JavaCallee, transition, reexecutable, killedLocations));
-    }
-
-    /**
-     * Creates and registers the linkage for a foreign call.
-     * 
-     * @param descriptor the signature of the foreign call
-     * @param address the address of the code to call
-     * @param outgoingCcType outgoing (caller) calling convention type
-     * @param effect specifies if the call destroys or preserves all registers (apart from
-     *            temporaries which are always destroyed)
-     * @param transition specifies if this is a {@linkplain Transition#LEAF leaf} call
-     * @param reexecutable specifies if the foreign call can be re-executed without (meaningful)
-     *            side effects. Deoptimization will not return to a point before a foreign call that
-     *            cannot be re-executed.
-     * @param killedLocations the memory locations killed by the foreign call
-     */
-    protected HotSpotForeignCallLinkage registerForeignCall(ForeignCallDescriptor descriptor, long address, CallingConvention.Type outgoingCcType, RegisterEffect effect, Transition transition,
-                    boolean reexecutable, LocationIdentity... killedLocations) {
-        Class<?> resultType = descriptor.getResultType();
-        assert transition == LEAF || resultType.isPrimitive() || Word.class.isAssignableFrom(resultType) : "non-leaf foreign calls must return objects in thread local storage: " + descriptor;
-        return register(HotSpotForeignCallLinkage.create(descriptor, address, effect, outgoingCcType, null, transition, reexecutable, killedLocations));
-    }
-
-    private static void link(Stub stub) {
-        stub.getLinkage().setCompiledStub(stub);
-    }
-
-    /**
-     * Creates a {@linkplain ForeignCallStub stub} for a foreign call.
-     * 
-     * @param descriptor the signature of the call to the stub
-     * @param address the address of the foreign code to call
-     * @param prependThread true if the JavaThread value for the current thread is to be prepended
-     *            to the arguments for the call to {@code address}
-     * @param transition specifies if this is a {@linkplain Transition#LEAF leaf} call
-     * @param reexecutable specifies if the foreign call can be re-executed without (meaningful)
-     *            side effects. Deoptimization will not return to a point before a foreign call that
-     *            cannot be re-executed.
-     * @param killedLocations the memory locations killed by the foreign call
-     */
-    private void linkForeignCall(HotSpotProviders providers, ForeignCallDescriptor descriptor, long address, boolean prependThread, Transition transition, boolean reexecutable,
-                    LocationIdentity... killedLocations) {
-        ForeignCallStub stub = new ForeignCallStub(providers, address, descriptor, prependThread, transition, reexecutable, killedLocations);
-        HotSpotForeignCallLinkage linkage = stub.getLinkage();
-        HotSpotForeignCallLinkage targetLinkage = stub.getTargetLinkage();
-        linkage.setCompiledStub(stub);
-        register(linkage);
-        register(targetLinkage);
-    }
-
-    public static final boolean PREPEND_THREAD = true;
-    public static final boolean DONT_PREPEND_THREAD = !PREPEND_THREAD;
-
-    public static final boolean REEXECUTABLE = true;
-    public static final boolean NOT_REEXECUTABLE = !REEXECUTABLE;
-
-    public static final LocationIdentity[] NO_LOCATIONS = {};
-
-    public void initialize(HotSpotProviders providers) {
-        HotSpotVMConfig c = runtime.getConfig();
-        TargetDescription target = providers.getCodeCache().getTarget();
-
-        registerForeignCall(UNCOMMON_TRAP, c.uncommonTrapStub, NativeCall, PRESERVES_REGISTERS, LEAF, REEXECUTABLE, NO_LOCATIONS);
-        registerForeignCall(DEOPT_HANDLER, c.handleDeoptStub, NativeCall, PRESERVES_REGISTERS, LEAF, REEXECUTABLE, NO_LOCATIONS);
-        registerForeignCall(IC_MISS_HANDLER, c.inlineCacheMissStub, NativeCall, PRESERVES_REGISTERS, LEAF, REEXECUTABLE, NO_LOCATIONS);
-
-        registerForeignCall(JAVA_TIME_MILLIS, c.javaTimeMillisAddress, NativeCall, DESTROYS_REGISTERS, LEAF, REEXECUTABLE, NO_LOCATIONS);
-        registerForeignCall(JAVA_TIME_NANOS, c.javaTimeNanosAddress, NativeCall, DESTROYS_REGISTERS, LEAF, REEXECUTABLE, NO_LOCATIONS);
-        registerForeignCall(ARITHMETIC_SIN, c.arithmeticSinAddress, NativeCall, DESTROYS_REGISTERS, LEAF, REEXECUTABLE, NO_LOCATIONS);
-        registerForeignCall(ARITHMETIC_COS, c.arithmeticCosAddress, NativeCall, DESTROYS_REGISTERS, LEAF, REEXECUTABLE, NO_LOCATIONS);
-        registerForeignCall(ARITHMETIC_TAN, c.arithmeticTanAddress, NativeCall, DESTROYS_REGISTERS, LEAF, REEXECUTABLE, NO_LOCATIONS);
-        registerForeignCall(LOAD_AND_CLEAR_EXCEPTION, c.loadAndClearExceptionAddress, NativeCall, DESTROYS_REGISTERS, LEAF, NOT_REEXECUTABLE, ANY_LOCATION);
-
-        registerForeignCall(EXCEPTION_HANDLER_FOR_PC, c.exceptionHandlerForPcAddress, NativeCall, DESTROYS_REGISTERS, NOT_LEAF, REEXECUTABLE, ANY_LOCATION);
-        registerForeignCall(EXCEPTION_HANDLER_FOR_RETURN_ADDRESS, c.exceptionHandlerForReturnAddressAddress, NativeCall, DESTROYS_REGISTERS, NOT_LEAF, REEXECUTABLE, ANY_LOCATION);
-        registerForeignCall(NEW_ARRAY_C, c.newArrayAddress, NativeCall, DESTROYS_REGISTERS, NOT_LEAF, REEXECUTABLE, ANY_LOCATION);
-        registerForeignCall(NEW_INSTANCE_C, c.newInstanceAddress, NativeCall, DESTROYS_REGISTERS, NOT_LEAF, REEXECUTABLE, ANY_LOCATION);
-        registerForeignCall(VM_MESSAGE_C, c.vmMessageAddress, NativeCall, DESTROYS_REGISTERS, NOT_LEAF, REEXECUTABLE, NO_LOCATIONS);
-
-        link(new NewInstanceStub(providers, target, registerStubCall(NEW_INSTANCE, REEXECUTABLE, NOT_LEAF, ANY_LOCATION)));
-        link(new NewArrayStub(providers, target, registerStubCall(NEW_ARRAY, REEXECUTABLE, NOT_LEAF, INIT_LOCATION)));
-        link(new ExceptionHandlerStub(providers, target, foreignCalls.get(EXCEPTION_HANDLER)));
-        link(new UnwindExceptionToCallerStub(providers, target, registerStubCall(UNWIND_EXCEPTION_TO_CALLER, NOT_REEXECUTABLE, NOT_LEAF, ANY_LOCATION)));
-        link(new VerifyOopStub(providers, target, registerStubCall(VERIFY_OOP, REEXECUTABLE, LEAF, NO_LOCATIONS)));
-
-        linkForeignCall(providers, IDENTITY_HASHCODE, c.identityHashCodeAddress, PREPEND_THREAD, NOT_LEAF, NOT_REEXECUTABLE, MARK_WORD_LOCATION);
-        linkForeignCall(providers, REGISTER_FINALIZER, c.registerFinalizerAddress, PREPEND_THREAD, NOT_LEAF, NOT_REEXECUTABLE, ANY_LOCATION);
-        linkForeignCall(providers, CREATE_NULL_POINTER_EXCEPTION, c.createNullPointerExceptionAddress, PREPEND_THREAD, NOT_LEAF, REEXECUTABLE, ANY_LOCATION);
-        linkForeignCall(providers, CREATE_OUT_OF_BOUNDS_EXCEPTION, c.createOutOfBoundsExceptionAddress, PREPEND_THREAD, NOT_LEAF, REEXECUTABLE, ANY_LOCATION);
-        linkForeignCall(providers, MONITORENTER, c.monitorenterAddress, PREPEND_THREAD, NOT_LEAF, NOT_REEXECUTABLE, ANY_LOCATION);
-        linkForeignCall(providers, MONITOREXIT, c.monitorexitAddress, PREPEND_THREAD, NOT_LEAF, NOT_REEXECUTABLE, ANY_LOCATION);
-        linkForeignCall(providers, NEW_MULTI_ARRAY, c.newMultiArrayAddress, PREPEND_THREAD, NOT_LEAF, REEXECUTABLE, INIT_LOCATION);
-        linkForeignCall(providers, DYNAMIC_NEW_ARRAY, c.dynamicNewArrayAddress, PREPEND_THREAD, NOT_LEAF, REEXECUTABLE, INIT_LOCATION);
-        linkForeignCall(providers, LOG_PRINTF, c.logPrintfAddress, PREPEND_THREAD, LEAF, REEXECUTABLE, NO_LOCATIONS);
-        linkForeignCall(providers, LOG_OBJECT, c.logObjectAddress, PREPEND_THREAD, NOT_LEAF, REEXECUTABLE, NO_LOCATIONS);
-        linkForeignCall(providers, LOG_PRIMITIVE, c.logPrimitiveAddress, PREPEND_THREAD, NOT_LEAF, REEXECUTABLE, NO_LOCATIONS);
-        linkForeignCall(providers, THREAD_IS_INTERRUPTED, c.threadIsInterruptedAddress, PREPEND_THREAD, NOT_LEAF, NOT_REEXECUTABLE, ANY_LOCATION);
-        linkForeignCall(providers, VM_ERROR, c.vmErrorAddress, PREPEND_THREAD, LEAF, REEXECUTABLE, NO_LOCATIONS);
-        linkForeignCall(providers, OSR_MIGRATION_END, c.osrMigrationEndAddress, DONT_PREPEND_THREAD, LEAF, NOT_REEXECUTABLE, NO_LOCATIONS);
-        linkForeignCall(providers, G1WBPRECALL, c.writeBarrierPreAddress, PREPEND_THREAD, LEAF, REEXECUTABLE, NO_LOCATIONS);
-        linkForeignCall(providers, G1WBPOSTCALL, c.writeBarrierPostAddress, PREPEND_THREAD, LEAF, REEXECUTABLE, NO_LOCATIONS);
-        linkForeignCall(providers, VALIDATE_OBJECT, c.validateObject, PREPEND_THREAD, LEAF, REEXECUTABLE, NO_LOCATIONS);
-    }
-
-    public HotSpotForeignCallLinkage lookupForeignCall(ForeignCallDescriptor descriptor) {
-        HotSpotForeignCallLinkage callTarget = foreignCalls.get(descriptor);
-        assert foreignCalls != null : descriptor;
-        callTarget.finalizeAddress(runtime.getBackend());
-        return callTarget;
-    }
-
-    @Override
-    public boolean isReexecutable(ForeignCallDescriptor descriptor) {
-        return foreignCalls.get(descriptor).isReexecutable();
-    }
-
-    public boolean canDeoptimize(ForeignCallDescriptor descriptor) {
-        return foreignCalls.get(descriptor).canDeoptimize();
-    }
-
-    public LocationIdentity[] getKilledLocations(ForeignCallDescriptor descriptor) {
-        return foreignCalls.get(descriptor).getKilledLocations();
-    }
-}
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotHostForeignCallsProvider.java
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotHostForeignCallsProvider.java	Thu Oct 17 01:08:17 2013 +0200
@@ -0,0 +1,232 @@
+/*
+ * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved.
+ * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
+ *
+ * This code is free software; you can redistribute it and/or modify it
+ * under the terms of the GNU General Public License version 2 only, as
+ * published by the Free Software Foundation.
+ *
+ * This code is distributed in the hope that it will be useful, but WITHOUT
+ * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
+ * FITNESS FOR A PARTICULAR PURPOSE.  See the GNU General Public License
+ * version 2 for more details (a copy is included in the LICENSE file that
+ * accompanied this code).
+ *
+ * You should have received a copy of the GNU General Public License version
+ * 2 along with this work; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
+ *
+ * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
+ * or visit www.oracle.com if you need additional information or have any
+ * questions.
+ */
+package com.oracle.graal.hotspot.meta;
+
+import static com.oracle.graal.api.code.CallingConvention.Type.*;
+import static com.oracle.graal.api.meta.LocationIdentity.*;
+import static com.oracle.graal.hotspot.HotSpotBackend.*;
+import static com.oracle.graal.hotspot.HotSpotForeignCallLinkage.RegisterEffect.*;
+import static com.oracle.graal.hotspot.HotSpotForeignCallLinkage.Transition.*;
+import static com.oracle.graal.hotspot.nodes.MonitorExitStubCall.*;
+import static com.oracle.graal.hotspot.nodes.NewArrayStubCall.*;
+import static com.oracle.graal.hotspot.nodes.NewInstanceStubCall.*;
+import static com.oracle.graal.hotspot.nodes.NewMultiArrayStubCall.*;
+import static com.oracle.graal.hotspot.nodes.VMErrorNode.*;
+import static com.oracle.graal.hotspot.replacements.HotSpotReplacementsUtil.*;
+import static com.oracle.graal.hotspot.replacements.MonitorSnippets.*;
+import static com.oracle.graal.hotspot.replacements.NewObjectSnippets.*;
+import static com.oracle.graal.hotspot.replacements.SystemSubstitutions.*;
+import static com.oracle.graal.hotspot.replacements.ThreadSubstitutions.*;
+import static com.oracle.graal.hotspot.replacements.WriteBarrierSnippets.*;
+import static com.oracle.graal.hotspot.stubs.ExceptionHandlerStub.*;
+import static com.oracle.graal.hotspot.stubs.NewArrayStub.*;
+import static com.oracle.graal.hotspot.stubs.NewInstanceStub.*;
+import static com.oracle.graal.hotspot.stubs.StubUtil.*;
+import static com.oracle.graal.hotspot.stubs.UnwindExceptionToCallerStub.*;
+import static com.oracle.graal.java.GraphBuilderPhase.RuntimeCalls.*;
+import static com.oracle.graal.nodes.java.RegisterFinalizerNode.*;
+import static com.oracle.graal.replacements.Log.*;
+import static com.oracle.graal.replacements.MathSubstitutionsX86.*;
+
+import java.util.*;
+
+import com.oracle.graal.api.code.*;
+import com.oracle.graal.api.meta.*;
+import com.oracle.graal.hotspot.*;
+import com.oracle.graal.hotspot.HotSpotForeignCallLinkage.RegisterEffect;
+import com.oracle.graal.hotspot.HotSpotForeignCallLinkage.Transition;
+import com.oracle.graal.hotspot.stubs.*;
+import com.oracle.graal.word.*;
+
+/**
+ * HotSpot implementation of {@link ForeignCallsProvider}.
+ */
+public abstract class HotSpotHostForeignCallsProvider implements HotSpotForeignCallsProvider {
+
+    public static final ForeignCallDescriptor OSR_MIGRATION_END = new ForeignCallDescriptor("OSR_migration_end", void.class, long.class);
+    public static final ForeignCallDescriptor IDENTITY_HASHCODE = new ForeignCallDescriptor("identity_hashcode", int.class, Object.class);
+    public static final ForeignCallDescriptor VERIFY_OOP = new ForeignCallDescriptor("verify_oop", Object.class, Object.class);
+    public static final ForeignCallDescriptor LOAD_AND_CLEAR_EXCEPTION = new ForeignCallDescriptor("load_and_clear_exception", Object.class, Word.class);
+
+    protected final HotSpotGraalRuntime runtime;
+
+    private final Map<ForeignCallDescriptor, HotSpotForeignCallLinkage> foreignCalls = new HashMap<>();
+    private final MetaAccessProvider metaAccess;
+    private final CodeCacheProvider codeCache;
+
+    public HotSpotHostForeignCallsProvider(HotSpotGraalRuntime runtime, MetaAccessProvider metaAccess, CodeCacheProvider codeCache) {
+        this.runtime = runtime;
+        this.metaAccess = metaAccess;
+        this.codeCache = codeCache;
+    }
+
+    /**
+     * Registers the linkage for a foreign call.
+     */
+    protected HotSpotForeignCallLinkage register(HotSpotForeignCallLinkage linkage) {
+        assert !foreignCalls.containsKey(linkage.getDescriptor()) : "already registered linkage for " + linkage.getDescriptor();
+        foreignCalls.put(linkage.getDescriptor(), linkage);
+        return linkage;
+    }
+
+    /**
+     * Creates and registers the details for linking a foreign call to a {@link Stub}.
+     * 
+     * @param descriptor the signature of the call to the stub
+     * @param reexecutable specifies if the stub call can be re-executed without (meaningful) side
+     *            effects. Deoptimization will not return to a point before a stub call that cannot
+     *            be re-executed.
+     * @param transition specifies if this is a {@linkplain Transition#LEAF leaf} call
+     * @param killedLocations the memory locations killed by the stub call
+     */
+    protected HotSpotForeignCallLinkage registerStubCall(ForeignCallDescriptor descriptor, boolean reexecutable, Transition transition, LocationIdentity... killedLocations) {
+        return register(HotSpotForeignCallLinkage.create(metaAccess, codeCache, this, descriptor, 0L, PRESERVES_REGISTERS, JavaCall, JavaCallee, transition, reexecutable, killedLocations));
+    }
+
+    /**
+     * Creates and registers the linkage for a foreign call.
+     * 
+     * @param descriptor the signature of the foreign call
+     * @param address the address of the code to call
+     * @param outgoingCcType outgoing (caller) calling convention type
+     * @param effect specifies if the call destroys or preserves all registers (apart from
+     *            temporaries which are always destroyed)
+     * @param transition specifies if this is a {@linkplain Transition#LEAF leaf} call
+     * @param reexecutable specifies if the foreign call can be re-executed without (meaningful)
+     *            side effects. Deoptimization will not return to a point before a foreign call that
+     *            cannot be re-executed.
+     * @param killedLocations the memory locations killed by the foreign call
+     */
+    protected HotSpotForeignCallLinkage registerForeignCall(ForeignCallDescriptor descriptor, long address, CallingConvention.Type outgoingCcType, RegisterEffect effect, Transition transition,
+                    boolean reexecutable, LocationIdentity... killedLocations) {
+        Class<?> resultType = descriptor.getResultType();
+        assert transition == LEAF || resultType.isPrimitive() || Word.class.isAssignableFrom(resultType) : "non-leaf foreign calls must return objects in thread local storage: " + descriptor;
+        return register(HotSpotForeignCallLinkage.create(metaAccess, codeCache, this, descriptor, address, effect, outgoingCcType, null, transition, reexecutable, killedLocations));
+    }
+
+    private static void link(Stub stub) {
+        stub.getLinkage().setCompiledStub(stub);
+    }
+
+    /**
+     * Creates a {@linkplain ForeignCallStub stub} for a foreign call.
+     * 
+     * @param descriptor the signature of the call to the stub
+     * @param address the address of the foreign code to call
+     * @param prependThread true if the JavaThread value for the current thread is to be prepended
+     *            to the arguments for the call to {@code address}
+     * @param transition specifies if this is a {@linkplain Transition#LEAF leaf} call
+     * @param reexecutable specifies if the foreign call can be re-executed without (meaningful)
+     *            side effects. Deoptimization will not return to a point before a foreign call that
+     *            cannot be re-executed.
+     * @param killedLocations the memory locations killed by the foreign call
+     */
+    private void linkForeignCall(HotSpotProviders providers, ForeignCallDescriptor descriptor, long address, boolean prependThread, Transition transition, boolean reexecutable,
+                    LocationIdentity... killedLocations) {
+        ForeignCallStub stub = new ForeignCallStub(providers, address, descriptor, prependThread, transition, reexecutable, killedLocations);
+        HotSpotForeignCallLinkage linkage = stub.getLinkage();
+        HotSpotForeignCallLinkage targetLinkage = stub.getTargetLinkage();
+        linkage.setCompiledStub(stub);
+        register(linkage);
+        register(targetLinkage);
+    }
+
+    public static final boolean PREPEND_THREAD = true;
+    public static final boolean DONT_PREPEND_THREAD = !PREPEND_THREAD;
+
+    public static final boolean REEXECUTABLE = true;
+    public static final boolean NOT_REEXECUTABLE = !REEXECUTABLE;
+
+    public static final LocationIdentity[] NO_LOCATIONS = {};
+
+    public void initialize(HotSpotProviders providers) {
+        HotSpotVMConfig c = runtime.getConfig();
+        TargetDescription target = providers.getCodeCache().getTarget();
+
+        registerForeignCall(UNCOMMON_TRAP, c.uncommonTrapStub, NativeCall, PRESERVES_REGISTERS, LEAF, REEXECUTABLE, NO_LOCATIONS);
+        registerForeignCall(DEOPT_HANDLER, c.handleDeoptStub, NativeCall, PRESERVES_REGISTERS, LEAF, REEXECUTABLE, NO_LOCATIONS);
+        registerForeignCall(IC_MISS_HANDLER, c.inlineCacheMissStub, NativeCall, PRESERVES_REGISTERS, LEAF, REEXECUTABLE, NO_LOCATIONS);
+
+        registerForeignCall(JAVA_TIME_MILLIS, c.javaTimeMillisAddress, NativeCall, DESTROYS_REGISTERS, LEAF, REEXECUTABLE, NO_LOCATIONS);
+        registerForeignCall(JAVA_TIME_NANOS, c.javaTimeNanosAddress, NativeCall, DESTROYS_REGISTERS, LEAF, REEXECUTABLE, NO_LOCATIONS);
+        registerForeignCall(ARITHMETIC_SIN, c.arithmeticSinAddress, NativeCall, DESTROYS_REGISTERS, LEAF, REEXECUTABLE, NO_LOCATIONS);
+        registerForeignCall(ARITHMETIC_COS, c.arithmeticCosAddress, NativeCall, DESTROYS_REGISTERS, LEAF, REEXECUTABLE, NO_LOCATIONS);
+        registerForeignCall(ARITHMETIC_TAN, c.arithmeticTanAddress, NativeCall, DESTROYS_REGISTERS, LEAF, REEXECUTABLE, NO_LOCATIONS);
+        registerForeignCall(LOAD_AND_CLEAR_EXCEPTION, c.loadAndClearExceptionAddress, NativeCall, DESTROYS_REGISTERS, LEAF, NOT_REEXECUTABLE, ANY_LOCATION);
+
+        registerForeignCall(EXCEPTION_HANDLER_FOR_PC, c.exceptionHandlerForPcAddress, NativeCall, DESTROYS_REGISTERS, NOT_LEAF, REEXECUTABLE, ANY_LOCATION);
+        registerForeignCall(EXCEPTION_HANDLER_FOR_RETURN_ADDRESS, c.exceptionHandlerForReturnAddressAddress, NativeCall, DESTROYS_REGISTERS, NOT_LEAF, REEXECUTABLE, ANY_LOCATION);
+        registerForeignCall(NEW_ARRAY_C, c.newArrayAddress, NativeCall, DESTROYS_REGISTERS, NOT_LEAF, REEXECUTABLE, ANY_LOCATION);
+        registerForeignCall(NEW_INSTANCE_C, c.newInstanceAddress, NativeCall, DESTROYS_REGISTERS, NOT_LEAF, REEXECUTABLE, ANY_LOCATION);
+        registerForeignCall(VM_MESSAGE_C, c.vmMessageAddress, NativeCall, DESTROYS_REGISTERS, NOT_LEAF, REEXECUTABLE, NO_LOCATIONS);
+
+        link(new NewInstanceStub(providers, target, registerStubCall(NEW_INSTANCE, REEXECUTABLE, NOT_LEAF, ANY_LOCATION)));
+        link(new NewArrayStub(providers, target, registerStubCall(NEW_ARRAY, REEXECUTABLE, NOT_LEAF, INIT_LOCATION)));
+        link(new ExceptionHandlerStub(providers, target, foreignCalls.get(EXCEPTION_HANDLER)));
+        link(new UnwindExceptionToCallerStub(providers, target, registerStubCall(UNWIND_EXCEPTION_TO_CALLER, NOT_REEXECUTABLE, NOT_LEAF, ANY_LOCATION)));
+        link(new VerifyOopStub(providers, target, registerStubCall(VERIFY_OOP, REEXECUTABLE, LEAF, NO_LOCATIONS)));
+
+        linkForeignCall(providers, IDENTITY_HASHCODE, c.identityHashCodeAddress, PREPEND_THREAD, NOT_LEAF, NOT_REEXECUTABLE, MARK_WORD_LOCATION);
+        linkForeignCall(providers, REGISTER_FINALIZER, c.registerFinalizerAddress, PREPEND_THREAD, NOT_LEAF, NOT_REEXECUTABLE, ANY_LOCATION);
+        linkForeignCall(providers, CREATE_NULL_POINTER_EXCEPTION, c.createNullPointerExceptionAddress, PREPEND_THREAD, NOT_LEAF, REEXECUTABLE, ANY_LOCATION);
+        linkForeignCall(providers, CREATE_OUT_OF_BOUNDS_EXCEPTION, c.createOutOfBoundsExceptionAddress, PREPEND_THREAD, NOT_LEAF, REEXECUTABLE, ANY_LOCATION);
+        linkForeignCall(providers, MONITORENTER, c.monitorenterAddress, PREPEND_THREAD, NOT_LEAF, NOT_REEXECUTABLE, ANY_LOCATION);
+        linkForeignCall(providers, MONITOREXIT, c.monitorexitAddress, PREPEND_THREAD, NOT_LEAF, NOT_REEXECUTABLE, ANY_LOCATION);
+        linkForeignCall(providers, NEW_MULTI_ARRAY, c.newMultiArrayAddress, PREPEND_THREAD, NOT_LEAF, REEXECUTABLE, INIT_LOCATION);
+        linkForeignCall(providers, DYNAMIC_NEW_ARRAY, c.dynamicNewArrayAddress, PREPEND_THREAD, NOT_LEAF, REEXECUTABLE, INIT_LOCATION);
+        linkForeignCall(providers, LOG_PRINTF, c.logPrintfAddress, PREPEND_THREAD, LEAF, REEXECUTABLE, NO_LOCATIONS);
+        linkForeignCall(providers, LOG_OBJECT, c.logObjectAddress, PREPEND_THREAD, NOT_LEAF, REEXECUTABLE, NO_LOCATIONS);
+        linkForeignCall(providers, LOG_PRIMITIVE, c.logPrimitiveAddress, PREPEND_THREAD, NOT_LEAF, REEXECUTABLE, NO_LOCATIONS);
+        linkForeignCall(providers, THREAD_IS_INTERRUPTED, c.threadIsInterruptedAddress, PREPEND_THREAD, NOT_LEAF, NOT_REEXECUTABLE, ANY_LOCATION);
+        linkForeignCall(providers, VM_ERROR, c.vmErrorAddress, PREPEND_THREAD, LEAF, REEXECUTABLE, NO_LOCATIONS);
+        linkForeignCall(providers, OSR_MIGRATION_END, c.osrMigrationEndAddress, DONT_PREPEND_THREAD, LEAF, NOT_REEXECUTABLE, NO_LOCATIONS);
+        linkForeignCall(providers, G1WBPRECALL, c.writeBarrierPreAddress, PREPEND_THREAD, LEAF, REEXECUTABLE, NO_LOCATIONS);
+        linkForeignCall(providers, G1WBPOSTCALL, c.writeBarrierPostAddress, PREPEND_THREAD, LEAF, REEXECUTABLE, NO_LOCATIONS);
+        linkForeignCall(providers, VALIDATE_OBJECT, c.validateObject, PREPEND_THREAD, LEAF, REEXECUTABLE, NO_LOCATIONS);
+    }
+
+    public HotSpotForeignCallLinkage lookupForeignCall(ForeignCallDescriptor descriptor) {
+        HotSpotForeignCallLinkage callTarget = foreignCalls.get(descriptor);
+        assert foreignCalls != null : descriptor;
+        callTarget.finalizeAddress(runtime.getHostBackend());
+        return callTarget;
+    }
+
+    @Override
+    public boolean isReexecutable(ForeignCallDescriptor descriptor) {
+        return foreignCalls.get(descriptor).isReexecutable();
+    }
+
+    public boolean canDeoptimize(ForeignCallDescriptor descriptor) {
+        return foreignCalls.get(descriptor).canDeoptimize();
+    }
+
+    public LocationIdentity[] getKilledLocations(ForeignCallDescriptor descriptor) {
+        return foreignCalls.get(descriptor).getKilledLocations();
+    }
+
+    /**
+     * Gets the registers that must be saved across a foreign call into the runtime.
+     */
+    public abstract Value[] getNativeABICallerSaveRegisters();
+}
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotLoweringProvider.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotLoweringProvider.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotLoweringProvider.java	Thu Oct 17 01:08:17 2013 +0200
@@ -27,7 +27,7 @@
 import static com.oracle.graal.api.meta.DeoptimizationReason.*;
 import static com.oracle.graal.api.meta.LocationIdentity.*;
 import static com.oracle.graal.hotspot.HotSpotGraalRuntime.*;
-import static com.oracle.graal.hotspot.meta.HotSpotForeignCallsProvider.*;
+import static com.oracle.graal.hotspot.meta.HotSpotHostForeignCallsProvider.*;
 import static com.oracle.graal.hotspot.replacements.HotSpotReplacementsUtil.*;
 import static com.oracle.graal.hotspot.replacements.NewObjectSnippets.*;
 import static com.oracle.graal.nodes.java.ArrayLengthNode.*;
@@ -80,7 +80,7 @@
 
     public void initialize() {
         HotSpotVMConfig c = runtime.getConfig();
-        HotSpotProviders providers = runtime.getProviders();
+        HotSpotProviders providers = runtime.getHostProviders();
         Replacements r = providers.getReplacements();
 
         r.registerSubstitutions(ObjectSubstitutions.class);
@@ -413,7 +413,7 @@
             }
         } else if (n instanceof DynamicCounterNode) {
             if (graph.getGuardsStage() == StructuredGraph.GuardsStage.AFTER_FSA) {
-                BenchmarkCounters.lower((DynamicCounterNode) n, runtime.getProviders().getRegisters(), runtime.getConfig(), wordKind);
+                BenchmarkCounters.lower((DynamicCounterNode) n, runtime.getHostProviders().getRegisters(), runtime.getConfig(), wordKind);
             }
         } else if (n instanceof CheckCastDynamicNode) {
             checkcastDynamicSnippets.lower((CheckCastDynamicNode) n);
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotMethodData.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotMethodData.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotMethodData.java	Thu Oct 17 01:08:17 2013 +0200
@@ -78,7 +78,7 @@
     }
 
     public int getDeoptimizationCount(DeoptimizationReason reason) {
-        int reasonIndex = runtime().getProviders().getMetaAccess().convertDeoptReason(reason);
+        int reasonIndex = runtime().getHostProviders().getMetaAccess().convertDeoptReason(reason);
         return unsafe.getByte(metaspaceMethodData + config.methodDataOopTrapHistoryOffset + reasonIndex) & 0xFF;
     }
 
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedJavaField.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedJavaField.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedJavaField.java	Thu Oct 17 01:08:17 2013 +0200
@@ -92,7 +92,7 @@
      * only called for snippets or replacements.
      */
     private static boolean isCalledForSnippets() {
-        MetaAccessProvider metaAccess = runtime().getProviders().getMetaAccess();
+        MetaAccessProvider metaAccess = runtime().getHostProviders().getMetaAccess();
         ResolvedJavaMethod makeGraphMethod = null;
         ResolvedJavaMethod initMethod = null;
         try {
@@ -119,7 +119,7 @@
     private static final Set<ResolvedJavaField> notEmbeddable = new HashSet<>();
 
     private static void addResolvedToSet(Field field) {
-        MetaAccessProvider metaAccess = runtime().getProviders().getMetaAccess();
+        MetaAccessProvider metaAccess = runtime().getHostProviders().getMetaAccess();
         notEmbeddable.add(metaAccess.lookupJavaField(field));
     }
 
@@ -216,12 +216,12 @@
         if (receiver == null) {
             assert Modifier.isStatic(flags);
             if (holder.isInitialized()) {
-                return runtime().getProviders().getConstantReflection().readUnsafeConstant(getKind(), holder.mirror(), offset, getKind() == Kind.Object);
+                return runtime().getHostProviders().getConstantReflection().readUnsafeConstant(getKind(), holder.mirror(), offset, getKind() == Kind.Object);
             }
             return null;
         } else {
             assert !Modifier.isStatic(flags);
-            return runtime().getProviders().getConstantReflection().readUnsafeConstant(getKind(), receiver.asObject(), offset, getKind() == Kind.Object);
+            return runtime().getHostProviders().getConstantReflection().readUnsafeConstant(getKind(), receiver.asObject(), offset, getKind() == Kind.Object);
         }
     }
 
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedJavaMethod.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedJavaMethod.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedJavaMethod.java	Thu Oct 17 01:08:17 2013 +0200
@@ -343,7 +343,7 @@
         // Cannot use toJava() as it ignores the return type
         HotSpotSignature sig = getSignature();
         JavaType[] sigTypes = MetaUtil.signatureToTypes(sig, null);
-        MetaAccessProvider metaAccess = runtime().getProviders().getMetaAccess();
+        MetaAccessProvider metaAccess = runtime().getHostProviders().getMetaAccess();
         for (Method method : holder.mirror().getDeclaredMethods()) {
             if (method.getName().equals(name)) {
                 if (metaAccess.lookupJavaType(method.getReturnType()).equals(sig.getReturnType(holder))) {
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedObjectType.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedObjectType.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/meta/HotSpotResolvedObjectType.java	Thu Oct 17 01:08:17 2013 +0200
@@ -544,7 +544,7 @@
         Constructor[] constructors = javaMirror.getDeclaredConstructors();
         ResolvedJavaMethod[] result = new ResolvedJavaMethod[constructors.length];
         for (int i = 0; i < constructors.length; i++) {
-            result[i] = runtime().getProviders().getMetaAccess().lookupJavaConstructor(constructors[i]);
+            result[i] = runtime().getHostProviders().getMetaAccess().lookupJavaConstructor(constructors[i]);
             assert result[i].isConstructor();
         }
         return result;
@@ -555,7 +555,7 @@
         Method[] methods = javaMirror.getDeclaredMethods();
         ResolvedJavaMethod[] result = new ResolvedJavaMethod[methods.length];
         for (int i = 0; i < methods.length; i++) {
-            result[i] = runtime().getProviders().getMetaAccess().lookupJavaMethod(methods[i]);
+            result[i] = runtime().getHostProviders().getMetaAccess().lookupJavaMethod(methods[i]);
             assert !result[i].isConstructor();
         }
         return result;
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/CurrentJavaThreadNode.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/CurrentJavaThreadNode.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/nodes/CurrentJavaThreadNode.java	Thu Oct 17 01:08:17 2013 +0200
@@ -41,7 +41,7 @@
 
     @Override
     public void generate(LIRGeneratorTool gen) {
-        Register rawThread = runtime().getProviders().getRegisters().getThreadRegister();
+        Register rawThread = runtime().getHostProviders().getRegisters().getThreadRegister();
         gen.setResult(this, rawThread.asValue(this.kind()));
     }
 
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/replacements/HotSpotReplacementsUtil.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/replacements/HotSpotReplacementsUtil.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/replacements/HotSpotReplacementsUtil.java	Thu Oct 17 01:08:17 2013 +0200
@@ -24,7 +24,7 @@
 
 import static com.oracle.graal.graph.UnsafeAccess.*;
 import static com.oracle.graal.hotspot.HotSpotGraalRuntime.*;
-import static com.oracle.graal.hotspot.meta.HotSpotForeignCallsProvider.*;
+import static com.oracle.graal.hotspot.meta.HotSpotHostForeignCallsProvider.*;
 import static com.oracle.graal.nodes.extended.BranchProbabilityNode.*;
 import sun.misc.*;
 
@@ -203,12 +203,12 @@
 
     @Fold
     public static Register threadRegister() {
-        return runtime().getProviders().getRegisters().getThreadRegister();
+        return runtime().getHostProviders().getRegisters().getThreadRegister();
     }
 
     @Fold
     public static Register stackPointerRegister() {
-        return runtime().getProviders().getRegisters().getStackPointerRegister();
+        return runtime().getHostProviders().getRegisters().getStackPointerRegister();
     }
 
     @Fold
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/replacements/LoadExceptionObjectSnippets.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/replacements/LoadExceptionObjectSnippets.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/replacements/LoadExceptionObjectSnippets.java	Thu Oct 17 01:08:17 2013 +0200
@@ -22,7 +22,7 @@
  */
 package com.oracle.graal.hotspot.replacements;
 
-import static com.oracle.graal.hotspot.meta.HotSpotForeignCallsProvider.*;
+import static com.oracle.graal.hotspot.meta.HotSpotHostForeignCallsProvider.*;
 import static com.oracle.graal.hotspot.replacements.HotSpotReplacementsUtil.*;
 import static com.oracle.graal.nodes.PiNode.*;
 import static com.oracle.graal.replacements.SnippetTemplate.*;
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/stubs/ForeignCallStub.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/stubs/ForeignCallStub.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/stubs/ForeignCallStub.java	Thu Oct 17 01:08:17 2013 +0200
@@ -84,11 +84,13 @@
      */
     public ForeignCallStub(HotSpotProviders providers, long address, ForeignCallDescriptor descriptor, boolean prependThread, Transition transition, boolean reexecutable,
                     LocationIdentity... killedLocations) {
-        super(providers, HotSpotForeignCallLinkage.create(descriptor, 0L, PRESERVES_REGISTERS, JavaCall, JavaCallee, transition, reexecutable, killedLocations));
+        super(providers, HotSpotForeignCallLinkage.create(providers.getMetaAccess(), providers.getCodeCache(), providers.getForeignCalls(), descriptor, 0L, PRESERVES_REGISTERS, JavaCall, JavaCallee,
+                        transition, reexecutable, killedLocations));
         this.prependThread = prependThread;
         Class[] targetParameterTypes = createTargetParameters(descriptor);
         ForeignCallDescriptor targetSig = new ForeignCallDescriptor(descriptor.getName() + ":C", descriptor.getResultType(), targetParameterTypes);
-        target = HotSpotForeignCallLinkage.create(targetSig, address, DESTROYS_REGISTERS, NativeCall, NativeCall, transition, reexecutable, killedLocations);
+        target = HotSpotForeignCallLinkage.create(providers.getMetaAccess(), providers.getCodeCache(), providers.getForeignCalls(), targetSig, address, DESTROYS_REGISTERS, NativeCall, NativeCall,
+                        transition, reexecutable, killedLocations);
     }
 
     /**
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/stubs/VerifyOopStub.java
--- a/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/stubs/VerifyOopStub.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.hotspot/src/com/oracle/graal/hotspot/stubs/VerifyOopStub.java	Thu Oct 17 01:08:17 2013 +0200
@@ -30,7 +30,7 @@
 import com.oracle.graal.replacements.*;
 
 /**
- * Stub called via {@link HotSpotForeignCallsProvider#VERIFY_OOP}.
+ * Stub called via {@link HotSpotHostForeignCallsProvider#VERIFY_OOP}.
  */
 public class VerifyOopStub extends SnippetStub {
 
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.replacements.test/src/com/oracle/graal/replacements/test/CompiledExceptionHandlerTest.java
--- a/graal/com.oracle.graal.replacements.test/src/com/oracle/graal/replacements/test/CompiledExceptionHandlerTest.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.replacements.test/src/com/oracle/graal/replacements/test/CompiledExceptionHandlerTest.java	Thu Oct 17 01:08:17 2013 +0200
@@ -37,7 +37,7 @@
 public class CompiledExceptionHandlerTest extends GraalCompilerTest {
 
     public CompiledExceptionHandlerTest() {
-        suites.getHighTier().findPhase(AbstractInliningPhase.class).remove();
+        getSuites().getHighTier().findPhase(AbstractInliningPhase.class).remove();
     }
 
     @Override
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.replacements.test/src/com/oracle/graal/replacements/test/DeoptimizeOnExceptionTest.java
--- a/graal/com.oracle.graal.replacements.test/src/com/oracle/graal/replacements/test/DeoptimizeOnExceptionTest.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.replacements.test/src/com/oracle/graal/replacements/test/DeoptimizeOnExceptionTest.java	Thu Oct 17 01:08:17 2013 +0200
@@ -33,7 +33,7 @@
 public class DeoptimizeOnExceptionTest extends GraalCompilerTest {
 
     public DeoptimizeOnExceptionTest() {
-        suites.getHighTier().findPhase(AbstractInliningPhase.class).remove();
+        getSuites().getHighTier().findPhase(AbstractInliningPhase.class).remove();
     }
 
     private static void raiseException(String m1, String m2, String m3, String m4, String m5) {
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.replacements.test/src/com/oracle/graal/replacements/test/InstanceOfTest.java
--- a/graal/com.oracle.graal.replacements.test/src/com/oracle/graal/replacements/test/InstanceOfTest.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.replacements.test/src/com/oracle/graal/replacements/test/InstanceOfTest.java	Thu Oct 17 01:08:17 2013 +0200
@@ -42,7 +42,7 @@
 public class InstanceOfTest extends TypeCheckTest {
 
     public InstanceOfTest() {
-        suites.getHighTier().findPhase(AbstractInliningPhase.class).remove();
+        getSuites().getHighTier().findPhase(AbstractInliningPhase.class).remove();
     }
 
     @Override
diff -r d08accd58925 -r f87c68d79f07 graal/com.oracle.graal.replacements.test/src/com/oracle/graal/replacements/test/InvokeTest.java
--- a/graal/com.oracle.graal.replacements.test/src/com/oracle/graal/replacements/test/InvokeTest.java	Thu Oct 17 01:05:13 2013 +0200
+++ b/graal/com.oracle.graal.replacements.test/src/com/oracle/graal/replacements/test/InvokeTest.java	Thu Oct 17 01:08:17 2013 +0200
@@ -33,7 +33,7 @@
 public class InvokeTest extends GraalCompilerTest {
 
     public InvokeTest() {
-        suites.getHighTier().findPhase(AbstractInliningPhase.class).remove();
+        getSuites().getHighTier().findPhase(AbstractInliningPhase.class).remove();
     }
 
     public interface I {
diff -r d08accd58925 -r f87c68d79f07 mx/projects
--- a/mx/projects	Thu Oct 17 01:05:13 2013 +0200
+++ b/mx/projects	Thu Oct 17 01:08:17 2013 +0200
@@ -28,7 +28,13 @@
 library@OKRA@urls=http://cr.openjdk.java.net/~tdeneau/okra-1.2.jar
 
 distribution@GRAAL@path=graal.jar
-distribution@GRAAL@dependencies=com.oracle.graal.hotspot.amd64,com.oracle.graal.hotspot.ptx,com.oracle.graal.truffle,com.oracle.graal.hotspot.sparc,com.oracle.graal.hotspot,com.oracle.graal.compiler.hsail
+distribution@GRAAL@dependencies=\
+    com.oracle.graal.hotspot.amd64,\
+    com.oracle.graal.hotspot.ptx,\
+    com.oracle.graal.truffle,\
+    com.oracle.graal.hotspot.sparc,\
+    com.oracle.graal.hotspot,\
+    com.oracle.graal.hotspot.hsail
 
 # graal.api.runtime
 project@com.oracle.graal.api.runtime@subDir=graal
@@ -137,7 +143,7 @@
 # graal.hotspot.ptx
 project@com.oracle.graal.hotspot.ptx@subDir=graal
 project@com.oracle.graal.hotspot.ptx@sourceDirs=src
-project@com.oracle.graal.hotspot.ptx@dependencies=com.oracle.graal.hotspot,com.oracle.graal.ptx
+project@com.oracle.graal.hotspot.ptx@dependencies=com.oracle.graal.hotspot,com.oracle.graal.ptx,com.oracle.graal.compiler.ptx
 project@com.oracle.graal.hotspot.ptx@checkstyle=com.oracle.graal.graph
 project@com.oracle.graal.hotspot.ptx@annotationProcessors=com.oracle.graal.service.processor
 project@com.oracle.graal.hotspot.ptx@javaCompliance=1.7
diff -r d08accd58925 -r f87c68d79f07 src/share/vm/classfile/vmSymbols.hpp
--- a/src/share/vm/classfile/vmSymbols.hpp	Thu Oct 17 01:05:13 2013 +0200
+++ b/src/share/vm/classfile/vmSymbols.hpp	Thu Oct 17 01:08:17 2013 +0200
@@ -314,9 +314,6 @@
   template(com_oracle_graal_hotspot_meta_HotSpotMonitorValue,        "com/oracle/graal/hotspot/meta/HotSpotMonitorValue")             \
   template(com_oracle_graal_hotspot_debug_LocalImpl,                 "com/oracle/graal/hotspot/debug/LocalImpl")                      \
   template(com_oracle_graal_hotspot_CompilerThread,                  "com/oracle/graal/hotspot/CompilerThread")                       \
-  template(com_oracle_graal_hotspot_ptx_PTXHotSpotGraalRuntime,      "com/oracle/graal/hotspot/ptx/PTXHotSpotGraalRuntime")\
-  AMD64_ONLY(template(com_oracle_graal_hotspot_amd64_AMD64HotSpotGraalRuntime,"com/oracle/graal/hotspot/amd64/AMD64HotSpotGraalRuntime"))\
-  SPARC_ONLY(template(com_oracle_graal_hotspot_sparc_SPARCHotSpotGraalRuntime,"com/oracle/graal/hotspot/sparc/SPARCHotSpotGraalRuntime"))\
   /* graal.api.meta */                                                                                                                \
   template(com_oracle_graal_api_meta_Constant,                       "com/oracle/graal/api/meta/Constant")                            \
   template(com_oracle_graal_api_meta_ConstantPool,                   "com/oracle/graal/api/meta/ConstantPool")                        \
@@ -387,11 +384,11 @@
   template(createConstantObject_signature,        "(Ljava/lang/Object;)Lcom/oracle/graal/api/meta/Constant;")                         \
   template(getVMToCompiler_name,                  "getVMToCompiler")                                                                  \
   template(getVMToCompiler_signature,             "()Lcom/oracle/graal/hotspot/bridge/VMToCompiler;")                                 \
-  template(getInstance_name,                      "getInstance")                                                                      \
-  template(getTruffleRuntimeInstance_signature,   "()Lcom/oracle/graal/truffle/GraalTruffleRuntime;")                                 \
+  template(runtime_name,                          "runtime")                                                                          \
+  template(runtime_signature,                     "()Lcom/oracle/graal/hotspot/HotSpotGraalRuntime;")                                 \
   template(makeInstance_name,                     "makeInstance")                                                                     \
+  template(makeInstance_signature,                "()Lcom/oracle/graal/truffle/GraalTruffleRuntime;")                                 \
   template(initialize_name,                       "initialize")                                                                       \
-  template(getInstance_signature,                 "()Lcom/oracle/graal/hotspot/HotSpotGraalRuntime;")                                 \
   template(forObject_name,                        "forObject")                                                                        \
   template(callbackInternal_name,                 "callbackInternal")                                                                 \
   template(callback_signature,                    "(Ljava/lang/Object;)Ljava/lang/Object;")                                           \
diff -r d08accd58925 -r f87c68d79f07 src/share/vm/graal/graalCompilerToGPU.cpp
--- a/src/share/vm/graal/graalCompilerToGPU.cpp	Thu Oct 17 01:05:13 2013 +0200
+++ b/src/share/vm/graal/graalCompilerToGPU.cpp	Thu Oct 17 01:08:17 2013 +0200
@@ -23,6 +23,7 @@
 
 #include "precompiled.hpp"
 
+#include "memory/oopFactory.hpp"
 #include "graal/graalCompiler.hpp"
 #include "graal/graalEnv.hpp"
 #include "graal/graalJavaAccess.hpp"
@@ -189,6 +190,28 @@
   return gpu::available_processors();
 C2V_END
 
+static objArrayHandle newSingletonStringArray(const char* value, TRAPS) {
+  objArrayOop res = oopFactory::new_objArray(SystemDictionary::String_klass(), 1, CHECK_NULL);
+  objArrayHandle res_h = objArrayHandle(THREAD, res);
+  Handle valueString = java_lang_String::create_from_str(value, CHECK_NULL);
+  res_h->obj_at_put(0, valueString());
+  return res_h;
+}
+
+C2V_VMENTRY(jobject, getAvailableGPUArchitectures, (JNIEnv *env, jobject))
+  if (UseGPU) {
+    if (gpu::is_available() && gpu::has_gpu_linkage()) {
+      if (gpu::get_target_il_type() == gpu::PTX) {
+        return JNIHandles::make_local(newSingletonStringArray("PTX", THREAD)());
+      }
+      if (gpu::get_target_il_type() == gpu::HSAIL) {
+        return JNIHandles::make_local(newSingletonStringArray("HSAIL", THREAD)());
+      }
+    }
+  }
+  return JNIHandles::make_local(oopFactory::new_objArray(SystemDictionary::String_klass(), 0, THREAD));
+C2V_END
+
 C2V_VMENTRY(jboolean, deviceDetach, (JNIEnv *env, jobject))
 return true;
 C2V_END
@@ -205,7 +228,7 @@
 #define CONSTANT_POOL         "Lcom/oracle/graal/api/meta/ConstantPool;"
 #define CONSTANT              "Lcom/oracle/graal/api/meta/Constant;"
 #define KIND                  "Lcom/oracle/graal/api/meta/Kind;"
-#define LOCAL                  "Lcom/oracle/graal/api/meta/Local;"
+#define LOCAL                 "Lcom/oracle/graal/api/meta/Local;"
 #define RUNTIME_CALL          "Lcom/oracle/graal/api/code/RuntimeCall;"
 #define EXCEPTION_HANDLERS    "[Lcom/oracle/graal/api/meta/ExceptionHandler;"
 #define REFLECT_METHOD        "Ljava/lang/reflect/Method;"
@@ -234,6 +257,7 @@
   {CC"deviceInit",                    CC"()Z",                                    FN_PTR(deviceInit)},
   {CC"deviceDetach",                  CC"()Z",                                    FN_PTR(deviceDetach)},
   {CC"availableProcessors",           CC"()I",                                    FN_PTR(availableProcessors)},
+  {CC"getAvailableGPUArchitectures",  CC"()["STRING,                              FN_PTR(getAvailableGPUArchitectures)},
   {CC"executeExternalMethodVarargs",  CC"(["OBJECT HS_INSTALLED_CODE")"OBJECT,    FN_PTR(executeExternalMethodVarargs)},
   {CC"executeParallelMethodVarargs",  CC"(III["OBJECT HS_INSTALLED_CODE")"OBJECT, FN_PTR(executeParallelMethodVarargs)},
 };
diff -r d08accd58925 -r f87c68d79f07 src/share/vm/graal/graalVMToCompiler.cpp
--- a/src/share/vm/graal/graalVMToCompiler.cpp	Thu Oct 17 01:05:13 2013 +0200
+++ b/src/share/vm/graal/graalVMToCompiler.cpp	Thu Oct 17 01:08:17 2013 +0200
@@ -53,32 +53,16 @@
   KlassHandle klass = loadClass(name);
 
   JavaValue result(T_OBJECT);
-  JavaCalls::call_static(&result, klass, vmSymbols::makeInstance_name(), vmSymbols::getTruffleRuntimeInstance_signature(), Thread::current());
+  JavaCalls::call_static(&result, klass, vmSymbols::makeInstance_name(), vmSymbols::makeInstance_signature(), Thread::current());
   check_pending_exception("Couldn't initialize GraalTruffleRuntime");
   return Handle((oop) result.get_jobject());
 }
 
 Handle VMToCompiler::graalRuntime() {
   if (JNIHandles::resolve(_graalRuntimePermObject) == NULL) {
-#ifdef AMD64
-    Symbol* name = NULL;
-    // Set name to PTXHotSpotRuntime if nVidia GPU was detected.
-    if (UseGPU && (gpu::get_target_il_type() == gpu::PTX) &&
-        gpu::is_available() && gpu::has_gpu_linkage()) {
-      name = vmSymbols::com_oracle_graal_hotspot_ptx_PTXHotSpotGraalRuntime();
-    }
-    
-    if (name == NULL) {
-      name = vmSymbols::com_oracle_graal_hotspot_amd64_AMD64HotSpotGraalRuntime();
-    }
-#endif
-#ifdef SPARC
-    Symbol* name = vmSymbols::com_oracle_graal_hotspot_sparc_SPARCHotSpotGraalRuntime();
-#endif
-    KlassHandle klass = loadClass(name);
-
+    KlassHandle klass = loadClass(vmSymbols::com_oracle_graal_hotspot_HotSpotGraalRuntime());
     JavaValue result(T_OBJECT);
-    JavaCalls::call_static(&result, klass, vmSymbols::makeInstance_name(), vmSymbols::getInstance_signature(), Thread::current());
+    JavaCalls::call_static(&result, klass, vmSymbols::runtime_name(), vmSymbols::runtime_signature(), Thread::current());
     check_pending_exception("Couldn't initialize HotSpotGraalRuntime");
     _graalRuntimePermObject = JNIHandles::make_global((oop) result.get_jobject());
   }