Mercurial > hg > truffle

annotate graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotBackend.java @ 16217:561070049e73

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
implement Canonicalizable.Binary in the BinaryNode hierarchy
author Lukas Stadler <lukas.stadler@oracle.com>
date Wed, 25 Jun 2014 16:54:56 +0200
parents 15771ff797b4
children 0d47af538a92
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
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
25 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
26 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
27 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
28 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
29 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
30 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
31
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
32 import java.util.*;
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
33
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
34 import com.oracle.graal.api.code.*;
13766
1a81855dae10 removed unused PTX code
Doug Simon <doug.simon@oracle.com>
parents: 13758
diff changeset
35 import com.oracle.graal.api.code.CallingConvention.Type;
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.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
37 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
38 import com.oracle.graal.asm.ptx.*;
15193
96bb07a5d667 Spit up and move GraalInternalError.
Josef Eisl <josef.eisl@jku.at>
parents: 15192
diff changeset
39 import com.oracle.graal.compiler.common.*;
15192
644dfe49c0f4 Move packages com.oracle.graal.cfg to com.oracle.graal.compiler.common.cfg.
Josef Eisl <josef.eisl@jku.at>
parents: 15180
diff changeset
40 import com.oracle.graal.compiler.common.cfg.*;
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
41 import com.oracle.graal.debug.*;
13766
1a81855dae10 removed unused PTX code
Doug Simon <doug.simon@oracle.com>
parents: 13758
diff changeset
42 import com.oracle.graal.debug.Debug.Scope;
14768
3e9a960f0da1 HSAIL: preliminary deopt support
Doug Simon <doug.simon@oracle.com>
parents: 14524
diff changeset
43 import com.oracle.graal.gpu.*;
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
44 import com.oracle.graal.hotspot.*;
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
45 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
46 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
47 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
48 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
49 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
50 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
51 import com.oracle.graal.lir.asm.*;
15292
a38d791982e1 Move LIRGenerationResult* to graal.lir.
Josef Eisl <josef.eisl@jku.at>
parents: 15249
diff changeset
52 import com.oracle.graal.lir.gen.*;
12456
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.ptx.*;
15249
54dd156e5f09 [PTX] Fix regression in generation of predicate register declaration
bharadwaj
parents: 15193
diff changeset
54 import com.oracle.graal.lir.ptx.PTXControlFlow.PTXPredicatedLIRInstruction;
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
55 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
56 import com.oracle.graal.nodes.*;
15337
c4be3c1b2d6d Use NodeLIRBuilderTool instead of NodeLIRBuilder.
Josef Eisl <josef.eisl@jku.at>
parents: 15333
diff changeset
57 import com.oracle.graal.nodes.spi.*;
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
58 import com.oracle.graal.phases.*;
13758
57d1746e3b3d refactored phase for stamping parameters as non-null that was duplicated in the GPU backends
Doug Simon <doug.simon@oracle.com>
parents: 13753
diff changeset
59 import com.oracle.graal.phases.common.*;
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.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
61 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
62
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 * 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
65 */
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 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
67
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
68 private final boolean deviceInitialized;
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
69
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
70 /**
13753
80cd5c3b8827 partially fixed passing of object parameters to PTX kernels; use a C++ object for managing resource allocation and cleanup around a PTX kernel execution
Doug Simon <doug.simon@oracle.com>
parents: 13731
diff changeset
71 * Descriptor for the PTX runtime method for calling a kernel. The C++ signature is:
14921
88dfaf6448e0 Remove LIRGenerationResult from NodeLIRBuilder.
Josef Eisl <josef.eisl@jku.at>
parents: 14888
diff changeset
72 *
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
73 * <pre>
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
74 * jlong (JavaThread* thread,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
75 * jlong kernel,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
76 * jint dimX,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
77 * jint dimY,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
78 * jint dimZ,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
79 * jlong parametersAndReturnValueBuffer,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
80 * jint parametersAndReturnValueBufferSize,
13753
80cd5c3b8827 partially fixed passing of object parameters to PTX kernels; use a C++ object for managing resource allocation and cleanup around a PTX kernel execution
Doug Simon <doug.simon@oracle.com>
parents: 13731
diff changeset
81 * jint objectParametersCount,
80cd5c3b8827 partially fixed passing of object parameters to PTX kernels; use a C++ object for managing resource allocation and cleanup around a PTX kernel execution
Doug Simon <doug.simon@oracle.com>
parents: 13731
diff changeset
82 * jlong objectParametersOffsets,
80cd5c3b8827 partially fixed passing of object parameters to PTX kernels; use a C++ object for managing resource allocation and cleanup around a PTX kernel execution
Doug Simon <doug.simon@oracle.com>
parents: 13731
diff changeset
83 * jlong pinnedObjects,
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
84 * 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
85 * </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
86 */
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
87 // @formatter:off
13753
80cd5c3b8827 partially fixed passing of object parameters to PTX kernels; use a C++ object for managing resource allocation and cleanup around a PTX kernel execution
Doug Simon <doug.simon@oracle.com>
parents: 13731
diff changeset
88 public static final ForeignCallDescriptor CALL_KERNEL = new ForeignCallDescriptor("execute_kernel_from_vm", long.class,
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
89 Word.class, // thread
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
90 long.class, // kernel
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
91 int.class, // dimX
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
92 int.class, // dimY
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
93 int.class, // dimZ
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
94 long.class, // parametersAndReturnValueBuffer
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
95 int.class, // parametersAndReturnValueBufferSize
13753
80cd5c3b8827 partially fixed passing of object parameters to PTX kernels; use a C++ object for managing resource allocation and cleanup around a PTX kernel execution
Doug Simon <doug.simon@oracle.com>
parents: 13731
diff changeset
96 int.class, // objectParameterCount
80cd5c3b8827 partially fixed passing of object parameters to PTX kernels; use a C++ object for managing resource allocation and cleanup around a PTX kernel execution
Doug Simon <doug.simon@oracle.com>
parents: 13731
diff changeset
97 long.class, // objectParameterOffsets
80cd5c3b8827 partially fixed passing of object parameters to PTX kernels; use a C++ object for managing resource allocation and cleanup around a PTX kernel execution
Doug Simon <doug.simon@oracle.com>
parents: 13731
diff changeset
98 long.class, // pinnedObjects
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
99 int.class); // encodedReturnTypeSize
15011
c8e575742f36 allow compilation with custom RegisterConfig
Lukas Stadler <lukas.stadler@oracle.com>
parents: 14921
diff changeset
100
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
101 // @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
102
12431
7080a96be216 rename: graalRuntime -> runtime, getGraalRuntime -> getRuntime
Doug Simon <doug.simon@oracle.com>
parents: 12429
diff changeset
103 public PTXHotSpotBackend(HotSpotGraalRuntime runtime, HotSpotProviders providers) {
7080a96be216 rename: graalRuntime -> runtime, getGraalRuntime -> getRuntime
Doug Simon <doug.simon@oracle.com>
parents: 12429
diff changeset
104 super(runtime, providers);
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
105 if (OmitDeviceInit) {
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
106 deviceInitialized = true;
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
107 } else {
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
108 boolean init = false;
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
109 try {
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
110 init = initialize();
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
111 } catch (UnsatisfiedLinkError e) {
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
112 }
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
113 deviceInitialized = init;
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
114 }
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
115 }
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
116
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
117 /**
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
118 * Initializes the GPU device.
14921
88dfaf6448e0 Remove LIRGenerationResult from NodeLIRBuilder.
Josef Eisl <josef.eisl@jku.at>
parents: 14888
diff changeset
119 *
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
120 * @return whether or not initialization was successful
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
121 */
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
122 private static native boolean initialize();
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
123
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
124 @Override
11775
b04b94b71649 Finished PTX assembler and Register -> Variable conversion
Morris Meyer <morris.meyer@oracle.com>
parents: 11596
diff changeset
125 public boolean shouldAllocateRegisters() {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
126 return false;
11775
b04b94b71649 Finished PTX assembler and Register -> Variable conversion
Morris Meyer <morris.meyer@oracle.com>
parents: 11596
diff changeset
127 }
b04b94b71649 Finished PTX assembler and Register -> Variable conversion
Morris Meyer <morris.meyer@oracle.com>
parents: 11596
diff changeset
128
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
129 /**
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
130 * Used to omit {@linkplain #initialize() device initialization}.
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
131 */
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
132 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
133
11775
b04b94b71649 Finished PTX assembler and Register -> Variable conversion
Morris Meyer <morris.meyer@oracle.com>
parents: 11596
diff changeset
134 @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
135 public void completeInitialization() {
13830
b00e3c4d9ada foreign call from PTX wrapper to PTX runtime requires a stub
Doug Simon <doug.simon@oracle.com>
parents: 13819
diff changeset
136 HotSpotProviders hostProviders = getRuntime().getHostProviders();
b00e3c4d9ada foreign call from PTX wrapper to PTX runtime requires a stub
Doug Simon <doug.simon@oracle.com>
parents: 13819
diff changeset
137 HotSpotHostForeignCallsProvider hostForeignCalls = (HotSpotHostForeignCallsProvider) hostProviders.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
138 if (deviceInitialized) {
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
139 long launchKernel = getLaunchKernelAddress();
13830
b00e3c4d9ada foreign call from PTX wrapper to PTX runtime requires a stub
Doug Simon <doug.simon@oracle.com>
parents: 13819
diff changeset
140 hostForeignCalls.linkForeignCall(hostProviders, CALL_KERNEL, launchKernel, false, NOT_LEAF, NOT_REEXECUTABLE, ANY_LOCATION);
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
141 }
13922
0995dcbd6dd8 Change CUDA context management to support multiple executions of a kernel. Exclude GPU offloading of lambdas from java.* library code.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 13888
diff changeset
142 /* Add a shutdown hook to destroy CUDA context(s) */
0995dcbd6dd8 Change CUDA context management to support multiple executions of a kernel. Exclude GPU offloading of lambdas from java.* library code.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 13888
diff changeset
143 Runtime.getRuntime().addShutdownHook(new Thread("PTXShutdown") {
0995dcbd6dd8 Change CUDA context management to support multiple executions of a kernel. Exclude GPU offloading of lambdas from java.* library code.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 13888
diff changeset
144 @Override
0995dcbd6dd8 Change CUDA context management to support multiple executions of a kernel. Exclude GPU offloading of lambdas from java.* library code.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 13888
diff changeset
145 public void run() {
0995dcbd6dd8 Change CUDA context management to support multiple executions of a kernel. Exclude GPU offloading of lambdas from java.* library code.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 13888
diff changeset
146 destroyContext();
0995dcbd6dd8 Change CUDA context management to support multiple executions of a kernel. Exclude GPU offloading of lambdas from java.* library code.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 13888
diff changeset
147 }
0995dcbd6dd8 Change CUDA context management to support multiple executions of a kernel. Exclude GPU offloading of lambdas from java.* library code.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 13888
diff changeset
148 });
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
149 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
150 }
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
151
13922
0995dcbd6dd8 Change CUDA context management to support multiple executions of a kernel. Exclude GPU offloading of lambdas from java.* library code.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 13888
diff changeset
152 private static native void destroyContext();
0995dcbd6dd8 Change CUDA context management to support multiple executions of a kernel. Exclude GPU offloading of lambdas from java.* library code.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 13888
diff changeset
153
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
154 /**
13888
51584f76462d pulled Ptx and Hsail classes out of gpu class namespace
Doug Simon <doug.simon@oracle.com>
parents: 13839
diff changeset
155 * Gets the address of {@code Ptx::execute_kernel_from_vm()}.
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
156 */
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
157 private static native long getLaunchKernelAddress();
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
158
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
159 @Override
15011
c8e575742f36 allow compilation with custom RegisterConfig
Lukas Stadler <lukas.stadler@oracle.com>
parents: 14921
diff changeset
160 public FrameMap newFrameMap(RegisterConfig registerConfig) {
c8e575742f36 allow compilation with custom RegisterConfig
Lukas Stadler <lukas.stadler@oracle.com>
parents: 14921
diff changeset
161 return new PTXFrameMap(getCodeCache(), registerConfig);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
162 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
163
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
164 /**
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
165 * Determines if the GPU device (or simulator) is available and initialized.
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
166 */
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
167 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
168 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
169 }
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
170
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
171 /**
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
172 * Compiles a given method to PTX code.
14921
88dfaf6448e0 Remove LIRGenerationResult from NodeLIRBuilder.
Josef Eisl <josef.eisl@jku.at>
parents: 14888
diff changeset
173 *
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
174 * @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
175 * 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
176 * {@linkplain ExternalCompilationResult#getEntryPoint() entry point}.
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
177 * @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
178 */
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
179 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
180 StructuredGraph graph = new StructuredGraph(method);
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
181 HotSpotProviders providers = getProviders();
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
182 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
183 PhaseSuite<HighTierContext> graphBuilderSuite = providers.getSuites().getDefaultGraphBuilderSuite();
13758
57d1746e3b3d refactored phase for stamping parameters as non-null that was duplicated in the GPU backends
Doug Simon <doug.simon@oracle.com>
parents: 13753
diff changeset
184 graphBuilderSuite.appendPhase(new NonNullParametersPhase());
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
185 Suites suites = providers.getSuites().getDefaultSuites();
14142
4eac66a9b87d Remove reference to graph in LIRGenerator.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 14065
diff changeset
186 ExternalCompilationResult ptxCode = compileGraph(graph, null, cc, method, providers, this, this.getTarget(), null, graphBuilderSuite, OptimisticOptimizations.NONE, getProfilingInfo(graph),
14143
8836f566b0bc Remove withScope parameter on GraalCompiler.compileGraph.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 14142
diff changeset
187 null, suites, new ExternalCompilationResult(), CompilationResultBuilderFactory.Default);
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
188 if (makeBinary) {
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
189 try (Scope ds = Debug.scope("GeneratingKernelBinary")) {
13753
80cd5c3b8827 partially fixed passing of object parameters to PTX kernels; use a C++ object for managing resource allocation and cleanup around a PTX kernel execution
Doug Simon <doug.simon@oracle.com>
parents: 13731
diff changeset
190 assert ptxCode.getTargetCode() != null;
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
191 long kernel = generateKernel(ptxCode.getTargetCode(), method.getName());
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
192 ptxCode.setEntryPoint(kernel);
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
193 } catch (Throwable e) {
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
194 throw Debug.handle(e);
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
195 }
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
196 }
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
197 return ptxCode;
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
198 }
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
199
13721
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
200 /**
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
201 * Generates a GPU binary from PTX code.
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
202 */
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
203 private static native long generateKernel(byte[] targetCode, String name);
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
204
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
205 /**
13721
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
206 * A list of the {@linkplain #installKernel(ResolvedJavaMethod, ExternalCompilationResult)
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
207 * installed} kernels. This is required so that there is a strong reference to each installed
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
208 * kernel as long as it is {@linkplain HotSpotNmethod#isValid() valid}. The list is pruned of
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
209 * invalid kernels every time a new kernel is installed.
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
210 */
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
211 private List<HotSpotNmethod> installedKernels = new LinkedList<>();
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
212
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
213 public final HotSpotNmethod installKernel(ResolvedJavaMethod method, ExternalCompilationResult ptxCode) {
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
214 assert OmitDeviceInit || ptxCode.getEntryPoint() != 0L;
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
215 HotSpotNmethod kernel = getProviders().getCodeCache().addExternalMethod(method, ptxCode);
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
216 synchronized (installedKernels) {
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
217 for (Iterator<HotSpotNmethod> i = installedKernels.iterator(); i.hasNext();) {
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
218 HotSpotNmethod entry = i.next();
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
219 if (!entry.isValid()) {
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
220 i.remove();
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
221 }
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
222 }
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
223 installedKernels.add(kernel);
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
224 }
2f944a810dd4 added list to PTXHotSpotBackend for keeping a valid installed kernel (and its associated nmethods) alive
Doug Simon <doug.simon@oracle.com>
parents: 13697
diff changeset
225 return kernel;
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
226 }
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
227
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
228 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
229 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
230 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
231
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
232 // 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
233 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
234 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
235
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
236 // 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
237 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
238 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
239
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 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
241
14032
d1c1f103d42c renamed com.oracle.graal.asm.AbstractAssembler to com.oracle.graal.asm.Assembler
twisti
parents: 14031
diff changeset
242 void emitDeclarations(Assembler asm) {
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
243 for (Integer i : unsigned8) {
14031
390c4b742890 made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents: 13922
diff changeset
244 asm.emitString(".reg .u8 %r" + i.intValue() + ";");
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
245 }
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
246 for (Integer i : signed32) {
14031
390c4b742890 made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents: 13922
diff changeset
247 asm.emitString(".reg .s32 %r" + i.intValue() + ";");
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
248 }
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 for (Integer i : signed64) {
14031
390c4b742890 made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents: 13922
diff changeset
250 asm.emitString(".reg .s64 %r" + i.intValue() + ";");
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
251 }
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 for (Integer i : unsigned64) {
14031
390c4b742890 made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents: 13922
diff changeset
253 asm.emitString(".reg .u64 %r" + i.intValue() + ";");
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
254 }
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 for (Integer i : float32) {
14031
390c4b742890 made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents: 13922
diff changeset
256 asm.emitString(".reg .f32 %r" + i.intValue() + ";");
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
257 }
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 for (Integer i : float64) {
14031
390c4b742890 made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents: 13922
diff changeset
259 asm.emitString(".reg .f64 %r" + i.intValue() + ";");
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
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 }
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
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 @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
264 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
265 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
266 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
267 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
268 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
269 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
270 } 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
271 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
272 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
273 // 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
274 // 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
275 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
276 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
277 }
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
278 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
279 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
280 // 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
281 // 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
282 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
283 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
284 }
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
285 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
286 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
287 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
288 // 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
289 // 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
290 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
291 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
292 }
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
293 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
294 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
295 // 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
296 // 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
297 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
298 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
299 }
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
300 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
301 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
302 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
303 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
304 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
305 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
306 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
307 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
308 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
309 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
310 }
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
311 }
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
312 }
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
313 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
314 }
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
315 }
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
316
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
317 class PTXFrameContext implements FrameContext {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
318
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
319 @Override
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
320 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
321 // 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
322 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
323
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
324 @Override
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
325 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
326 }
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
327
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
328 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
329 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
330 }
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
331 }
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
332
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
333 @Override
14809
73774616a6b3 Decoupled LIRGenerator and LIRGenerationResult.
Josef Eisl <josef.eisl@jku.at>
parents: 14806
diff changeset
334 public CompilationResultBuilder newCompilationResultBuilder(LIRGenerationResult lirGenRes, 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
335 // 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
336 // - 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
337 // - 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
338 // - 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
339 // - has no instructions with debug info
14809
73774616a6b3 Decoupled LIRGenerator and LIRGenerationResult.
Josef Eisl <josef.eisl@jku.at>
parents: 14806
diff changeset
340 FrameMap frameMap = lirGenRes.getFrameMap();
14032
d1c1f103d42c renamed com.oracle.graal.asm.AbstractAssembler to com.oracle.graal.asm.Assembler
twisti
parents: 14031
diff changeset
341 Assembler masm = createAssembler(frameMap);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
342 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
343 CompilationResultBuilder crb = factory.createBuilder(getCodeCache(), getForeignCalls(), frameMap, masm, frameContext, compilationResult);
15180
7a9531f50cd8 renamed com.oracle.graal.api.code.CompilationResult.frameSize to totalFrameSize
twisti
parents: 15157
diff changeset
344 crb.setTotalFrameSize(0);
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
345 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
346 }
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
347
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
348 @Override
15838
15771ff797b4 Pass the compiled method to LIR factory
Christian Wimmer <christian.wimmer@oracle.com>
parents: 15338
diff changeset
349 public LIRGenerationResult newLIRGenerationResult(LIR lir, FrameMap frameMap, ResolvedJavaMethod method, Object stub) {
14809
73774616a6b3 Decoupled LIRGenerator and LIRGenerationResult.
Josef Eisl <josef.eisl@jku.at>
parents: 14806
diff changeset
350 return new LIRGenerationResultBase(lir, frameMap);
73774616a6b3 Decoupled LIRGenerator and LIRGenerationResult.
Josef Eisl <josef.eisl@jku.at>
parents: 14806
diff changeset
351 }
73774616a6b3 Decoupled LIRGenerator and LIRGenerationResult.
Josef Eisl <josef.eisl@jku.at>
parents: 14806
diff changeset
352
73774616a6b3 Decoupled LIRGenerator and LIRGenerationResult.
Josef Eisl <josef.eisl@jku.at>
parents: 14806
diff changeset
353 @Override
14032
d1c1f103d42c renamed com.oracle.graal.asm.AbstractAssembler to com.oracle.graal.asm.Assembler
twisti
parents: 14031
diff changeset
354 protected Assembler 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
355 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
356 }
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
357
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
358 @Override
15333
06ecedffb109 Use LIRGeneratorTool in Backend.
Josef Eisl <josef.eisl@jku.at>
parents: 15292
diff changeset
359 public LIRGeneratorTool newLIRGenerator(CallingConvention cc, LIRGenerationResult lirGenRes) {
14849
97a0878202c2 Apply LIRGenerator refactoring to PTX backend.
Josef Eisl <josef.eisl@jku.at>
parents: 14809
diff changeset
360 return new PTXHotSpotLIRGenerator(getProviders(), getRuntime().getConfig(), cc, lirGenRes);
97a0878202c2 Apply LIRGenerator refactoring to PTX backend.
Josef Eisl <josef.eisl@jku.at>
parents: 14809
diff changeset
361 }
97a0878202c2 Apply LIRGenerator refactoring to PTX backend.
Josef Eisl <josef.eisl@jku.at>
parents: 14809
diff changeset
362
97a0878202c2 Apply LIRGenerator refactoring to PTX backend.
Josef Eisl <josef.eisl@jku.at>
parents: 14809
diff changeset
363 @Override
15338
5e544920ad9f Rename Backend.newNodeLIRGenerator to Backend.newNodeLIRBuilder.
Josef Eisl <josef.eisl@jku.at>
parents: 15337
diff changeset
364 public NodeLIRBuilderTool newNodeLIRBuilder(StructuredGraph graph, LIRGeneratorTool lirGen) {
14921
88dfaf6448e0 Remove LIRGenerationResult from NodeLIRBuilder.
Josef Eisl <josef.eisl@jku.at>
parents: 14888
diff changeset
365 return new PTXHotSpotNodeLIRBuilder(graph, lirGen);
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
14065
5dec26f3d4a4 Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 14032
diff changeset
368 private static void emitKernelEntry(CompilationResultBuilder crb, LIR lir, ResolvedJavaMethod codeCacheOwner) {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
369 // 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
370 // 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
371 // 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
372 // facilitate seemless PTX code generation subsequently.
14065
5dec26f3d4a4 Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 14032
diff changeset
373 assert codeCacheOwner != null : lir + " is not associated with a method";
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
374 final String name = codeCacheOwner.getName();
14032
d1c1f103d42c renamed com.oracle.graal.asm.AbstractAssembler to com.oracle.graal.asm.Assembler
twisti
parents: 14031
diff changeset
375 Assembler asm = crb.asm;
12456
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 // Emit initial boiler-plate directives.
14031
390c4b742890 made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents: 13922
diff changeset
378 asm.emitString(".version 3.0");
390c4b742890 made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents: 13922
diff changeset
379 asm.emitString(".target sm_30");
390c4b742890 made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents: 13922
diff changeset
380 asm.emitString0(".entry " + name + " (");
390c4b742890 made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents: 13922
diff changeset
381 asm.emitString("");
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
382
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
383 // Get the start block
14799
01afd501f127 Updated getStartBlock() in AbstractControlFlowGraph.
Josef Eisl <josef.eisl@jku.at>
parents: 14796
diff changeset
384 AbstractBlock<?> startBlock = lir.getControlFlowGraph().getStartBlock();
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
385 // 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
386 // 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
387 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
388
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
389 // 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
390 // instruction.
14806
a8723f1ff542 LIR renamed setter and getter functions.
Josef Eisl <josef.eisl@jku.at>
parents: 14804
diff changeset
391 for (LIRInstruction op : lir.getLIRforBlock(startBlock)) {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
392 if (op instanceof PTXParameterOp) {
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
393 op.emitCode(crb);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
394 deleteOps.add(op);
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
395 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
396 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
397
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
398 // Delete ParameterOp instructions.
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
399 for (LIRInstruction op : deleteOps) {
14806
a8723f1ff542 LIR renamed setter and getter functions.
Josef Eisl <josef.eisl@jku.at>
parents: 14804
diff changeset
400 lir.getLIRforBlock(startBlock).remove(op);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
401 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
402
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
403 // Start emiting body of the PTX kernel.
14031
390c4b742890 made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents: 13922
diff changeset
404 asm.emitString0(") {");
390c4b742890 made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents: 13922
diff changeset
405 asm.emitString("");
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
406 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
407
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
408 // Emit .reg space declarations
14065
5dec26f3d4a4 Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 14032
diff changeset
409 private static void emitRegisterDecl(CompilationResultBuilder crb, LIR lir, ResolvedJavaMethod codeCacheOwner) {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
410
14065
5dec26f3d4a4 Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 14032
diff changeset
411 assert codeCacheOwner != null : lir + " is not associated with a method";
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
412
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
413 RegisterAnalysis registerAnalysis = new RegisterAnalysis();
15249
54dd156e5f09 [PTX] Fix regression in generation of predicate register declaration
bharadwaj
parents: 15193
diff changeset
414 // Assume no predicate registers are used
54dd156e5f09 [PTX] Fix regression in generation of predicate register declaration
bharadwaj
parents: 15193
diff changeset
415 int maxPredRegNum = -1;
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
416
14786
a6595f1b55b0 Make LIR use AbstractBlock. (errors)
Josef Eisl <josef.eisl@jku.at>
parents: 14145
diff changeset
417 for (AbstractBlock<?> b : lir.codeEmittingOrder()) {
14806
a8723f1ff542 LIR renamed setter and getter functions.
Josef Eisl <josef.eisl@jku.at>
parents: 14804
diff changeset
418 for (LIRInstruction op : lir.getLIRforBlock(b)) {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
419 if (op instanceof LabelOp) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
420 // 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
421 } else {
15249
54dd156e5f09 [PTX] Fix regression in generation of predicate register declaration
bharadwaj
parents: 15193
diff changeset
422 if (op instanceof PTXPredicatedLIRInstruction) {
54dd156e5f09 [PTX] Fix regression in generation of predicate register declaration
bharadwaj
parents: 15193
diff changeset
423 // Update maximum predicate register number if op uses a larger number
54dd156e5f09 [PTX] Fix regression in generation of predicate register declaration
bharadwaj
parents: 15193
diff changeset
424 int opPredRegNum = ((PTXPredicatedLIRInstruction) op).getPredRegNum();
54dd156e5f09 [PTX] Fix regression in generation of predicate register declaration
bharadwaj
parents: 15193
diff changeset
425 maxPredRegNum = (opPredRegNum > maxPredRegNum) ? opPredRegNum : maxPredRegNum;
54dd156e5f09 [PTX] Fix regression in generation of predicate register declaration
bharadwaj
parents: 15193
diff changeset
426 }
54dd156e5f09 [PTX] Fix regression in generation of predicate register declaration
bharadwaj
parents: 15193
diff changeset
427 // Record registers used in the kernel
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
428 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
429 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
430 op.forEachOutput(registerAnalysis);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
431 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
432 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
433 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
434
15249
54dd156e5f09 [PTX] Fix regression in generation of predicate register declaration
bharadwaj
parents: 15193
diff changeset
435 // Emit register declarations
14032
d1c1f103d42c renamed com.oracle.graal.asm.AbstractAssembler to com.oracle.graal.asm.Assembler
twisti
parents: 14031
diff changeset
436 Assembler asm = crb.asm;
14031
390c4b742890 made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents: 13922
diff changeset
437 registerAnalysis.emitDeclarations(asm);
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
438
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
439 // emit predicate register declaration
15249
54dd156e5f09 [PTX] Fix regression in generation of predicate register declaration
bharadwaj
parents: 15193
diff changeset
440 if (maxPredRegNum > -1) {
54dd156e5f09 [PTX] Fix regression in generation of predicate register declaration
bharadwaj
parents: 15193
diff changeset
441 asm.emitString(".reg .pred %p<" + ++maxPredRegNum + ">;");
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
442 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
443 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
444
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
445 @Override
14065
5dec26f3d4a4 Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 14032
diff changeset
446 public void emitCode(CompilationResultBuilder crb, LIR lir, ResolvedJavaMethod codeCacheOwner) {
5dec26f3d4a4 Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 14032
diff changeset
447 assert codeCacheOwner != null : lir + " is not associated with a method";
14032
d1c1f103d42c renamed com.oracle.graal.asm.AbstractAssembler to com.oracle.graal.asm.Assembler
twisti
parents: 14031
diff changeset
448 Assembler asm = crb.asm;
14031
390c4b742890 made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents: 13922
diff changeset
449
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
450 // Emit the prologue
14065
5dec26f3d4a4 Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 14032
diff changeset
451 emitKernelEntry(crb, lir, codeCacheOwner);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
452
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
453 // Emit register declarations
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
454 try {
14065
5dec26f3d4a4 Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 14032
diff changeset
455 emitRegisterDecl(crb, lir, codeCacheOwner);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
456 } catch (GraalInternalError e) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
457 e.printStackTrace();
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
458 // 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
459 // all types of parameters are handled.
14031
390c4b742890 made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents: 13922
diff changeset
460 asm.close(false);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
461 return;
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
462 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
463 // 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
464 try {
14065
5dec26f3d4a4 Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 14032
diff changeset
465 crb.emit(lir);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
466 } catch (GraalInternalError e) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
467 e.printStackTrace();
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
468 // 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
469 // all types of parameters are handled.
14031
390c4b742890 made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents: 13922
diff changeset
470 asm.close(false);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
471 return;
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
472 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
473
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
474 // Emit the epilogue
14031
390c4b742890 made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents: 13922
diff changeset
475 asm.emitString0("}");
390c4b742890 made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents: 13922
diff changeset
476 asm.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
477 }
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
478
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
479 /**
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
480 * Gets the total number of available CUDA cores.
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
481 */
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
482 public int getAvailableProcessors() {
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
483 if (!deviceInitialized) {
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
484 return 0;
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
485 }
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
486 return getAvailableProcessors0();
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
487 }
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
488
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
489 private static native int getAvailableProcessors0();
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
490
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
491 }