Mercurial > hg > truffle

annotate graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotBackend.java @ 13683:de839ec35cc7

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
schedule lambda method compilation and execution on GPU (PTX) when possible; fix a couple of bugs.
author S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
date Fri, 17 Jan 2014 16:03:13 -0500
parents 1274feb0efe6
children e8bd4f3776ee
Ignore whitespace changes - Everywhere: Within whitespace: At end of lines:
rev   line source
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
1 /*
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
2 * Copyright (c) 2012, Oracle and/or its affiliates. All rights reserved.
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
3 * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
4 *
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
5 * This code is free software; you can redistribute it and/or modify it
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
6 * under the terms of the GNU General Public License version 2 only, as
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
7 * published by the Free Software Foundation.
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
8 *
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
9 * This code is distributed in the hope that it will be useful, but WITHOUT
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
10 * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
11 * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
12 * version 2 for more details (a copy is included in the LICENSE file that
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
13 * accompanied this code).
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
14 *
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
15 * You should have received a copy of the GNU General Public License version
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
16 * 2 along with this work; if not, write to the Free Software Foundation,
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
17 * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
18 *
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
19 * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
20 * or visit www.oracle.com if you need additional information or have any
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
21 * questions.
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
22 */
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
23 package com.oracle.graal.hotspot.ptx;
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
24
13624
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
25 import static com.oracle.graal.api.code.CallingConvention.Type.*;
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
26 import static com.oracle.graal.api.code.CodeUtil.*;
13624
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
27 import static com.oracle.graal.api.meta.LocationIdentity.*;
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
28 import static com.oracle.graal.compiler.GraalCompiler.*;
13624
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
29 import static com.oracle.graal.hotspot.HotSpotForeignCallLinkage.RegisterEffect.*;
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
30 import static com.oracle.graal.hotspot.HotSpotForeignCallLinkage.Transition.*;
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
31 import static com.oracle.graal.hotspot.meta.HotSpotForeignCallsProviderImpl.*;
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
32 import static com.oracle.graal.lir.LIRValueUtil.*;
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
33
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
34 import java.util.*;
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
35
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
36 import com.oracle.graal.api.code.*;
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
37 import com.oracle.graal.api.code.CallingConvention.*;
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
38 import com.oracle.graal.api.meta.*;
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
39 import com.oracle.graal.asm.*;
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
40 import com.oracle.graal.asm.ptx.*;
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
41 import com.oracle.graal.compiler.gen.*;
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
42 import com.oracle.graal.compiler.ptx.*;
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
43 import com.oracle.graal.debug.*;
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
44 import com.oracle.graal.debug.Debug.*;
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
45 import com.oracle.graal.graph.*;
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
46 import com.oracle.graal.hotspot.*;
13638
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
47 import com.oracle.graal.hotspot.HotSpotReplacementsImpl.GraphProducer;
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
48 import com.oracle.graal.hotspot.bridge.*;
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
49 import com.oracle.graal.hotspot.meta.*;
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
50 import com.oracle.graal.lir.*;
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
51 import com.oracle.graal.lir.LIRInstruction.OperandFlag;
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
52 import com.oracle.graal.lir.LIRInstruction.OperandMode;
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
53 import com.oracle.graal.lir.LIRInstruction.ValueProcedure;
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
54 import com.oracle.graal.lir.StandardOp.LabelOp;
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
55 import com.oracle.graal.lir.asm.*;
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
56 import com.oracle.graal.lir.ptx.*;
13624
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
57 import com.oracle.graal.lir.ptx.PTXMemOp.LoadReturnAddrOp;
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
58 import com.oracle.graal.nodes.*;
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
59 import com.oracle.graal.nodes.cfg.*;
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
60 import com.oracle.graal.phases.*;
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
61 import com.oracle.graal.phases.tiers.*;
13624
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
62 import com.oracle.graal.word.*;
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
63
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
64 /**
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
65 * HotSpot PTX specific backend.
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
66 */
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
67 public class PTXHotSpotBackend extends HotSpotBackend {
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
68
13624
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
69 /**
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
70 * Descriptor for the PTX runtime method for launching a kernel. The C++ signature is:
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
71 *
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
72 * <pre>
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
73 * jlong (JavaThread* thread,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
74 * jlong kernel,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
75 * jint dimX,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
76 * jint dimY,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
77 * jint dimZ,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
78 * jlong parametersAndReturnValueBuffer,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
79 * jint parametersAndReturnValueBufferSize,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
80 * jint encodedReturnTypeSize)
13624
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
81 * </pre>
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
82 */
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
83 // @formatter:off
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
84 public static final ForeignCallDescriptor LAUNCH_KERNEL = new ForeignCallDescriptor("execute_kernel_from_vm", long.class,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
85 Word.class, // thread
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
86 long.class, // kernel
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
87 int.class, // dimX
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
88 int.class, // dimY
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
89 int.class, // dimZ
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
90 long.class, // parametersAndReturnValueBuffer
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
91 int.class, // parametersAndReturnValueBufferSize
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
92 int.class); // encodedReturnTypeSize
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
93 // @formatter:on
13624
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
94
12431
7080a96be216 rename: graalRuntime -> runtime, getGraalRuntime -> getRuntime
Doug Simon <doug.simon@oracle.com>
parents: 12429
diff changeset
95 public PTXHotSpotBackend(HotSpotGraalRuntime runtime, HotSpotProviders providers) {
7080a96be216 rename: graalRuntime -> runtime, getGraalRuntime -> getRuntime
Doug Simon <doug.simon@oracle.com>
parents: 12429
diff changeset
96 super(runtime, providers);
13683
de839ec35cc7 schedule lambda method compilation and execution on GPU (PTX) when possible; fix a couple of bugs.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 13651
diff changeset
97 CompilerToGPU compilerToGPU = getRuntime().getCompilerToGPU();
de839ec35cc7 schedule lambda method compilation and execution on GPU (PTX) when possible; fix a couple of bugs.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 13651
diff changeset
98 deviceInitialized = OmitDeviceInit || compilerToGPU.deviceInit();
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
99 }
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
100
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
101 @Override
11775
b04b94b71649 Finished PTX assembler and Register -> Variable conversion
Morris Meyer <morris.meyer@oracle.com>
parents: 11596
diff changeset
102 public boolean shouldAllocateRegisters() {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
103 return false;
11775
b04b94b71649 Finished PTX assembler and Register -> Variable conversion
Morris Meyer <morris.meyer@oracle.com>
parents: 11596
diff changeset
104 }
b04b94b71649 Finished PTX assembler and Register -> Variable conversion
Morris Meyer <morris.meyer@oracle.com>
parents: 11596
diff changeset
105
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
106 /**
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
107 * Used to omit {@linkplain CompilerToGPU#deviceInit() device initialization}.
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
108 */
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
109 private static final boolean OmitDeviceInit = Boolean.getBoolean("graal.ptx.omitDeviceInit");
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
110
11775
b04b94b71649 Finished PTX assembler and Register -> Variable conversion
Morris Meyer <morris.meyer@oracle.com>
parents: 11596
diff changeset
111 @Override
13624
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
112 public void completeInitialization() {
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
113 HotSpotHostForeignCallsProvider hostForeignCalls = (HotSpotHostForeignCallsProvider) getRuntime().getHostProviders().getForeignCalls();
13638
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
114 CompilerToGPU compilerToGPU = getRuntime().getCompilerToGPU();
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
115 if (deviceInitialized) {
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
116 long launchKernel = compilerToGPU.getLaunchKernelAddress();
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
117 hostForeignCalls.registerForeignCall(LAUNCH_KERNEL, launchKernel, NativeCall, DESTROYS_REGISTERS, NOT_LEAF, NOT_REEXECUTABLE, ANY_LOCATION);
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
118 }
13624
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
119 super.completeInitialization();
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
120 }
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
121
13638
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
122 private boolean deviceInitialized;
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
123
13624
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
124 @Override
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
125 public FrameMap newFrameMap() {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
126 return new PTXFrameMap(getCodeCache());
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
127 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
128
13638
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
129 public boolean isDeviceInitialized() {
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
130 return deviceInitialized;
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
131 }
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
132
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
133 @Override
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
134 public GraphProducer getGraphProducer() {
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
135 if (!deviceInitialized) {
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
136 // GPU could not be initialized so offload is disabled
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
137 return null;
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
138 }
13651
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
139 return new GraphProducer() {
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
140
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
141 public StructuredGraph getGraphFor(ResolvedJavaMethod method) {
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
142 if (canOffloadToGPU(method)) {
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
143 ExternalCompilationResult ptxCode = PTXHotSpotBackend.this.compileKernel(method, true);
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
144 InstalledCode installedPTXCode = PTXHotSpotBackend.this.installKernel(method, ptxCode);
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
145 return new PTXWrapperBuilder(method, installedPTXCode.getStart(), getRuntime().getHostBackend().getProviders()).getGraph();
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
146 }
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
147 return null;
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
148 }
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
149
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
150 private boolean canOffloadToGPU(ResolvedJavaMethod method) {
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
151 return method.getName().contains("lambda$main$") & method.isSynthetic();
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
152 }
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
153 };
13638
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
154 }
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
155
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
156 /**
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
157 * Compiles a given method to PTX code.
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
158 *
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
159 * @param makeBinary specifies whether a GPU binary should also be generated for the PTX code.
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
160 * If true, the returned value is guaranteed to have a non-zero
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
161 * {@linkplain ExternalCompilationResult#getEntryPoint() entry point}.
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
162 * @return the PTX code compiled from {@code method}'s bytecode
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
163 */
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
164 public ExternalCompilationResult compileKernel(ResolvedJavaMethod method, boolean makeBinary) {
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
165 StructuredGraph graph = new StructuredGraph(method);
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
166 HotSpotProviders providers = getProviders();
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
167 CallingConvention cc = getCallingConvention(providers.getCodeCache(), Type.JavaCallee, method, false);
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
168 PhaseSuite<HighTierContext> graphBuilderSuite = providers.getSuites().getDefaultGraphBuilderSuite();
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
169 Suites suites = providers.getSuites().getDefaultSuites();
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
170 ExternalCompilationResult ptxCode = compileGraph(graph, cc, method, providers, this, this.getTarget(), null, graphBuilderSuite, OptimisticOptimizations.NONE, getProfilingInfo(graph),
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
171 new SpeculationLog(), suites, true, new ExternalCompilationResult(), CompilationResultBuilderFactory.Default);
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
172 if (makeBinary) {
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
173 try (Scope ds = Debug.scope("GeneratingKernelBinary")) {
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
174 long kernel = getRuntime().getCompilerToGPU().generateKernel(ptxCode.getTargetCode(), method.getName());
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
175 ptxCode.setEntryPoint(kernel);
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
176 } catch (Throwable e) {
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
177 throw Debug.handle(e);
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
178 }
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
179 }
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
180 return ptxCode;
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
181
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
182 }
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
183
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
184 public InstalledCode installKernel(ResolvedJavaMethod method, ExternalCompilationResult ptxCode) {
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
185 assert ptxCode.getEntryPoint() != 0L;
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
186 return getProviders().getCodeCache().addExternalMethod(method, ptxCode);
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
187 }
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
188
12653
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
189 static final class RegisterAnalysis extends ValueProcedure {
12762
884bee435276 Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12653
diff changeset
190 private final SortedSet<Integer> signed32 = new TreeSet<>();
12653
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
191 private final SortedSet<Integer> signed64 = new TreeSet<>();
12762
884bee435276 Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12653
diff changeset
192
884bee435276 Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12653
diff changeset
193 // unsigned8 is only for ld, st and cbt
884bee435276 Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12653
diff changeset
194 private final SortedSet<Integer> unsigned8 = new TreeSet<>();
884bee435276 Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12653
diff changeset
195 private final SortedSet<Integer> unsigned64 = new TreeSet<>();
884bee435276 Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12653
diff changeset
196
884bee435276 Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12653
diff changeset
197 // private final SortedSet<Integer> float16 = new TreeSet<>();
12653
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
198 private final SortedSet<Integer> float32 = new TreeSet<>();
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
199 private final SortedSet<Integer> float64 = new TreeSet<>();
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
200
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
201 LIRInstruction op;
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
202
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
203 void emitDeclarations(Buffer codeBuffer) {
12762
884bee435276 Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12653
diff changeset
204 for (Integer i : unsigned8) {
884bee435276 Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12653
diff changeset
205 codeBuffer.emitString(".reg .u8 %r" + i.intValue() + ";");
884bee435276 Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12653
diff changeset
206 }
12653
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
207 for (Integer i : signed32) {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
208 codeBuffer.emitString(".reg .s32 %r" + i.intValue() + ";");
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
209 }
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
210 for (Integer i : signed64) {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
211 codeBuffer.emitString(".reg .s64 %r" + i.intValue() + ";");
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
212 }
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
213 for (Integer i : unsigned64) {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
214 codeBuffer.emitString(".reg .u64 %r" + i.intValue() + ";");
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
215 }
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
216 for (Integer i : float32) {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
217 codeBuffer.emitString(".reg .f32 %r" + i.intValue() + ";");
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
218 }
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
219 for (Integer i : float64) {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
220 codeBuffer.emitString(".reg .f64 %r" + i.intValue() + ";");
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
221 }
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
222 }
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
223
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
224 @Override
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
225 public Value doValue(Value value, OperandMode mode, EnumSet<OperandFlag> flags) {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
226 if (isVariable(value)) {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
227 Variable regVal = (Variable) value;
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
228 Kind regKind = regVal.getKind();
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
229 if ((op instanceof LoadReturnAddrOp) && (mode == OperandMode.DEF)) {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
230 unsigned64.add(regVal.index);
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
231 } else {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
232 switch (regKind) {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
233 case Int:
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
234 // If the register was used as a wider signed type
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
235 // do not add it here
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
236 if (!signed64.contains(regVal.index)) {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
237 signed32.add(regVal.index);
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
238 }
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
239 break;
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
240 case Long:
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
241 // If the register was used as a narrower signed type
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
242 // remove it from there and add it to wider type.
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
243 if (signed32.contains(regVal.index)) {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
244 signed32.remove(regVal.index);
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
245 }
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
246 signed64.add(regVal.index);
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
247 break;
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
248 case Float:
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
249 // If the register was used as a wider signed type
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
250 // do not add it here
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
251 if (!float64.contains(regVal.index)) {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
252 float32.add(regVal.index);
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
253 }
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
254 break;
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
255 case Double:
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
256 // If the register was used as a narrower signed type
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
257 // remove it from there and add it to wider type.
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
258 if (float32.contains(regVal.index)) {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
259 float32.remove(regVal.index);
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
260 }
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
261 float64.add(regVal.index);
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
262 break;
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
263 case Object:
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
264 unsigned64.add(regVal.index);
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
265 break;
12762
884bee435276 Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12653
diff changeset
266 case Byte:
884bee435276 Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12653
diff changeset
267 unsigned8.add(regVal.index);
884bee435276 Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12653
diff changeset
268 break;
12653
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
269 default:
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
270 throw GraalInternalError.shouldNotReachHere("unhandled register type " + value.toString());
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
271 }
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
272 }
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
273 }
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
274 return value;
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
275 }
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
276 }
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
277
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
278 class PTXFrameContext implements FrameContext {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
279
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
280 @Override
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
281 public void enter(CompilationResultBuilder crb) {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
282 // codeBuffer.emitString(".address_size 32"); // PTX ISA version 2.3
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
283 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
284
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
285 @Override
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
286 public void leave(CompilationResultBuilder crb) {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
287 }
13234
7e237378923d made the frameContext of a CompilationResultBuilder always non-null and added FrameContext.hasFrame() to determine if a frame is actually generated
Doug Simon <doug.simon@oracle.com>
parents: 13229
diff changeset
288
7e237378923d made the frameContext of a CompilationResultBuilder always non-null and added FrameContext.hasFrame() to determine if a frame is actually generated
Doug Simon <doug.simon@oracle.com>
parents: 13229
diff changeset
289 public boolean hasFrame() {
7e237378923d made the frameContext of a CompilationResultBuilder always non-null and added FrameContext.hasFrame() to determine if a frame is actually generated
Doug Simon <doug.simon@oracle.com>
parents: 13229
diff changeset
290 return true;
7e237378923d made the frameContext of a CompilationResultBuilder always non-null and added FrameContext.hasFrame() to determine if a frame is actually generated
Doug Simon <doug.simon@oracle.com>
parents: 13229
diff changeset
291 }
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
292 }
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
293
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
294 @Override
13229
325b4e4efb60 added CompilationResultBuilderFactory to support peep-hole instrumentation of methods as their code is emitted
Doug Simon <doug.simon@oracle.com>
parents: 13227
diff changeset
295 public CompilationResultBuilder newCompilationResultBuilder(LIRGenerator lirGen, CompilationResult compilationResult, CompilationResultBuilderFactory factory) {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
296 // Omit the frame of the method:
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
297 // - has no spill slots or other slots allocated during register allocation
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
298 // - has no callee-saved registers
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
299 // - has no incoming arguments passed on the stack
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
300 // - has no instructions with debug info
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
301 FrameMap frameMap = lirGen.frameMap;
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
302 AbstractAssembler masm = createAssembler(frameMap);
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
303 PTXFrameContext frameContext = new PTXFrameContext();
13229
325b4e4efb60 added CompilationResultBuilderFactory to support peep-hole instrumentation of methods as their code is emitted
Doug Simon <doug.simon@oracle.com>
parents: 13227
diff changeset
304 CompilationResultBuilder crb = factory.createBuilder(getCodeCache(), getForeignCalls(), frameMap, masm, frameContext, compilationResult);
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
305 crb.setFrameSize(0);
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
306 return crb;
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
307 }
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
308
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
309 @Override
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
310 protected AbstractAssembler createAssembler(FrameMap frameMap) {
12762
884bee435276 Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12653
diff changeset
311 return new PTXMacroAssembler(getTarget(), frameMap.registerConfig);
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
312 }
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
313
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
314 @Override
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
315 public LIRGenerator newLIRGenerator(StructuredGraph graph, FrameMap frameMap, CallingConvention cc, LIR lir) {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
316 return new PTXLIRGenerator(graph, getProviders(), frameMap, cc, lir);
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
317 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
318
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
319 private static void emitKernelEntry(CompilationResultBuilder crb, LIRGenerator lirGen, ResolvedJavaMethod codeCacheOwner) {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
320 // Emit PTX kernel entry text based on PTXParameterOp
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
321 // instructions in the start block. Remove the instructions
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
322 // once kernel entry text and directives are emitted to
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
323 // facilitate seemless PTX code generation subsequently.
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
324 assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method";
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
325 final String name = codeCacheOwner.getName();
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
326 Buffer codeBuffer = crb.asm.codeBuffer;
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
327
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
328 // Emit initial boiler-plate directives.
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
329 codeBuffer.emitString(".version 3.0");
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
330 codeBuffer.emitString(".target sm_30");
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
331 codeBuffer.emitString0(".entry " + name + " (");
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
332 codeBuffer.emitString("");
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
333
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
334 // Get the start block
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
335 Block startBlock = lirGen.lir.cfg.getStartBlock();
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
336 // Keep a list of ParameterOp instructions to delete from the
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
337 // list of instructions in the block.
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
338 ArrayList<LIRInstruction> deleteOps = new ArrayList<>();
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
339
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
340 // Emit .param arguments to kernel entry based on ParameterOp
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
341 // instruction.
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
342 for (LIRInstruction op : lirGen.lir.lir(startBlock)) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
343 if (op instanceof PTXParameterOp) {
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
344 op.emitCode(crb);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
345 deleteOps.add(op);
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
346 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
347 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
348
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
349 // Delete ParameterOp instructions.
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
350 for (LIRInstruction op : deleteOps) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
351 lirGen.lir.lir(startBlock).remove(op);
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
352 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
353
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
354 // Start emiting body of the PTX kernel.
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
355 codeBuffer.emitString0(") {");
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
356 codeBuffer.emitString("");
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
357 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
358
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
359 // Emit .reg space declarations
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
360 private static void emitRegisterDecl(CompilationResultBuilder crb, LIRGenerator lirGen, ResolvedJavaMethod codeCacheOwner) {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
361
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
362 assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method";
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
363
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
364 Buffer codeBuffer = crb.asm.codeBuffer;
12653
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
365 RegisterAnalysis registerAnalysis = new RegisterAnalysis();
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
366
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
367 for (Block b : lirGen.lir.codeEmittingOrder()) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
368 for (LIRInstruction op : lirGen.lir.lir(b)) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
369 if (op instanceof LabelOp) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
370 // Don't consider this as a definition
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
371 } else {
12653
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
372 registerAnalysis.op = op;
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
373 op.forEachTemp(registerAnalysis);
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
374 op.forEachOutput(registerAnalysis);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
375 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
376 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
377 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
378
12653
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
379 registerAnalysis.emitDeclarations(codeBuffer);
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
380
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
381 // emit predicate register declaration
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
382 int maxPredRegNum = ((PTXLIRGenerator) lirGen).getNextPredRegNumber();
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
383 if (maxPredRegNum > 0) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
384 codeBuffer.emitString(".reg .pred %p<" + maxPredRegNum + ">;");
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
385 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
386 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
387
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
388 @Override
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
389 public void emitCode(CompilationResultBuilder crb, LIRGenerator lirGen, ResolvedJavaMethod codeCacheOwner) {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
390 assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method";
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
391 Buffer codeBuffer = crb.asm.codeBuffer;
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
392 // Emit the prologue
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
393 emitKernelEntry(crb, lirGen, codeCacheOwner);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
394
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
395 // Emit register declarations
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
396 try {
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
397 emitRegisterDecl(crb, lirGen, codeCacheOwner);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
398 } catch (GraalInternalError e) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
399 e.printStackTrace();
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
400 // TODO : Better error handling needs to be done once
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
401 // all types of parameters are handled.
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
402 codeBuffer.setPosition(0);
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
403 codeBuffer.close(false);
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
404 return;
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
405 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
406 // Emit code for the LIR
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
407 try {
13318
da0851712519 moved emitting code for LIR and queries about whether an edge goes to its lexical successor "inside" CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13234
diff changeset
408 crb.emit(lirGen.lir);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
409 } catch (GraalInternalError e) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
410 e.printStackTrace();
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
411 // TODO : Better error handling needs to be done once
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
412 // all types of parameters are handled.
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
413 codeBuffer.setPosition(0);
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
414 codeBuffer.close(false);
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
415 return;
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
416 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
417
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
418 // Emit the epilogue
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
419 codeBuffer.emitString0("}");
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
420 codeBuffer.emitString("");
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
421 }
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
422 }