Mercurial > hg > truffle

annotate graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotBackend.java @ 13922:0995dcbd6dd8

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
Change CUDA context management to support multiple executions of a kernel. Exclude GPU offloading of lambdas from java.* library code.
author S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
date Mon, 10 Feb 2014 14:38:42 -0500
parents 51584f76462d
children 390c4b742890
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.*;
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
39 import com.oracle.graal.compiler.gen.*;
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
40 import com.oracle.graal.compiler.ptx.*;
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
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;
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
43 import com.oracle.graal.graph.*;
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
44 import com.oracle.graal.hotspot.*;
13638
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
45 import com.oracle.graal.hotspot.HotSpotReplacementsImpl.GraphProducer;
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
46 import com.oracle.graal.hotspot.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
47 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
48 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
49 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
50 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
51 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
52 import com.oracle.graal.lir.asm.*;
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
53 import com.oracle.graal.lir.ptx.*;
13624
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
54 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
55 import com.oracle.graal.nodes.*;
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
56 import com.oracle.graal.nodes.cfg.*;
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
57 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
58 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
59 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
60 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
61
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 * 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
64 */
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
65 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
66
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
67 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
68
13624
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
69 /**
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
70 * Descriptor for the PTX runtime method for calling a kernel. The C++ signature is:
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
71 *
220ed109bf77 initial code for calling PTX kernel code from Java with parameter marshaling and return value unmarshaling performed by a wrapper specified via manual graph construction
Doug Simon <doug.simon@oracle.com>
parents: 13318
diff changeset
72 * <pre>
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
73 * jlong (JavaThread* thread,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
74 * jlong kernel,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
75 * jint dimX,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
76 * jint dimY,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
77 * jint dimZ,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
78 * jlong parametersAndReturnValueBuffer,
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
79 * jint parametersAndReturnValueBufferSize,
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
80 * 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
81 * 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
82 * jlong pinnedObjects,
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
83 * 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
84 * </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
85 */
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
86 // @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
87 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
88 Word.class, // thread
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
89 long.class, // kernel
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
90 int.class, // dimX
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
91 int.class, // dimY
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
92 int.class, // dimZ
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
93 long.class, // parametersAndReturnValueBuffer
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
94 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
95 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
96 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
97 long.class, // pinnedObjects
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
98 int.class); // encodedReturnTypeSize
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
99 // @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
100
12431
7080a96be216 rename: graalRuntime -> runtime, getGraalRuntime -> getRuntime
Doug Simon <doug.simon@oracle.com>
parents: 12429
diff changeset
101 public PTXHotSpotBackend(HotSpotGraalRuntime runtime, HotSpotProviders providers) {
7080a96be216 rename: graalRuntime -> runtime, getGraalRuntime -> getRuntime
Doug Simon <doug.simon@oracle.com>
parents: 12429
diff changeset
102 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
103 if (OmitDeviceInit) {
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
104 deviceInitialized = true;
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
105 } else {
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
106 boolean init = false;
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
107 try {
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
108 init = initialize();
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
109 } catch (UnsatisfiedLinkError e) {
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
110 }
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
111 deviceInitialized = init;
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
112 }
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
113 }
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
114
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
115 /**
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
116 * Initializes the GPU device.
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 * @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
119 */
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
120 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
121
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
122 @Override
11775
b04b94b71649 Finished PTX assembler and Register -> Variable conversion
Morris Meyer <morris.meyer@oracle.com>
parents: 11596
diff changeset
123 public boolean shouldAllocateRegisters() {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
124 return false;
11775
b04b94b71649 Finished PTX assembler and Register -> Variable conversion
Morris Meyer <morris.meyer@oracle.com>
parents: 11596
diff changeset
125 }
b04b94b71649 Finished PTX assembler and Register -> Variable conversion
Morris Meyer <morris.meyer@oracle.com>
parents: 11596
diff changeset
126
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
127 /**
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
128 * 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
129 */
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
130 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
131
11775
b04b94b71649 Finished PTX assembler and Register -> Variable conversion
Morris Meyer <morris.meyer@oracle.com>
parents: 11596
diff changeset
132 @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
133 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
134 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
135 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
136 if (deviceInitialized) {
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
137 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
138 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
139 }
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
140 /* 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
141 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
142 @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
143 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
144 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
145 }
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 });
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
147 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
148 }
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
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
150 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
151
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
152 /**
13888
51584f76462d pulled Ptx and Hsail classes out of gpu class namespace
Doug Simon <doug.simon@oracle.com>
parents: 13839
diff changeset
153 * 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
154 */
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
155 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
156
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
157 @Override
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
158 public FrameMap newFrameMap() {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
159 return new PTXFrameMap(getCodeCache());
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
160 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
161
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
162 /**
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
163 * 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
164 */
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
165 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
166 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
167 }
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
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 @Override
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
170 public GraphProducer getGraphProducer() {
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
171 if (!deviceInitialized) {
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
172 // GPU could not be initialized so offload is disabled
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
173 return null;
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
174 }
13651
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
175 return new GraphProducer() {
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
176
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
177 public StructuredGraph getGraphFor(ResolvedJavaMethod method) {
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
178 if (canOffloadToGPU(method)) {
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
179 ExternalCompilationResult ptxCode = PTXHotSpotBackend.this.compileKernel(method, true);
13697
e8bd4f3776ee changed return type of HotSpotCodeCacheProvider.addExternalMethod to be HotSpotNmethod
Doug Simon <doug.simon@oracle.com>
parents: 13683
diff changeset
180 HotSpotNmethod installedPTXCode = PTXHotSpotBackend.this.installKernel(method, ptxCode);
e8bd4f3776ee changed return type of HotSpotCodeCacheProvider.addExternalMethod to be HotSpotNmethod
Doug Simon <doug.simon@oracle.com>
parents: 13683
diff changeset
181 return new PTXWrapperBuilder(method, installedPTXCode, getRuntime().getHostBackend().getProviders()).getGraph();
13651
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
182 }
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
183 return null;
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
184 }
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
185
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
186 private boolean canOffloadToGPU(ResolvedJavaMethod method) {
13778
0be9a42f28e7 PTXHotSpotBackend$1.getGraphFor() checks GPUOffload VM option (JBS:GRAAL-6)
Doug Simon <doug.simon@oracle.com>
parents: 13773
diff changeset
187 HotSpotVMConfig config = getRuntime().getConfig();
0be9a42f28e7 PTXHotSpotBackend$1.getGraphFor() checks GPUOffload VM option (JBS:GRAAL-6)
Doug Simon <doug.simon@oracle.com>
parents: 13773
diff changeset
188 return config.gpuOffload && method.getName().contains("lambda$") & method.isSynthetic();
13651
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
189 }
1274feb0efe6 renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13649
diff changeset
190 };
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
191 }
1dabd01a73bd PTX backend can alter compilation pipeline to offload selected code to the GPU
Doug Simon <doug.simon@oracle.com>
parents: 13624
diff changeset
192
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
193 /**
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
194 * Compiles a given method to PTX code.
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
195 *
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
196 * @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
197 * 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
198 * {@linkplain ExternalCompilationResult#getEntryPoint() entry point}.
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
199 * @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
200 */
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
201 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
202 StructuredGraph graph = new StructuredGraph(method);
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
203 HotSpotProviders providers = getProviders();
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
204 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
205 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
206 graphBuilderSuite.appendPhase(new NonNullParametersPhase());
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
207 Suites suites = providers.getSuites().getDefaultSuites();
13773
e7c2351ed7d5 Throw-away speculation logs should not be used. Just use null instead
Gilles Duboscq <duboscq@ssw.jku.at>
parents: 13766
diff changeset
208 ExternalCompilationResult ptxCode = compileGraph(graph, cc, method, providers, this, this.getTarget(), null, graphBuilderSuite, OptimisticOptimizations.NONE, getProfilingInfo(graph), null,
e7c2351ed7d5 Throw-away speculation logs should not be used. Just use null instead
Gilles Duboscq <duboscq@ssw.jku.at>
parents: 13766
diff changeset
209 suites, true, new ExternalCompilationResult(), CompilationResultBuilderFactory.Default);
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
210 if (makeBinary) {
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
211 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
212 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
213 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
214 ptxCode.setEntryPoint(kernel);
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
215 } catch (Throwable e) {
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
216 throw Debug.handle(e);
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
217 }
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
218 }
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
219 return ptxCode;
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
220 }
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
221
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
222 /**
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
223 * 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
224 */
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
225 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
226
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
227 /**
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
228 * 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
229 * 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
230 * 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
231 * 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
232 */
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
233 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
234
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
235 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
236 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
237 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
238 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
239 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
240 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
241 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
242 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
243 }
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
244 }
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
245 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
246 }
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
247 return kernel;
13649
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
248 }
36d4faef2c56 moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents: 13638
diff changeset
249
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
250 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
251 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
252 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
253
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
254 // 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
255 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
256 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
257
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
258 // 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
259 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
260 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
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 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
263
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 void emitDeclarations(Buffer codeBuffer) {
12762
884bee435276 Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12653
diff changeset
265 for (Integer i : unsigned8) {
884bee435276 Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12653
diff changeset
266 codeBuffer.emitString(".reg .u8 %r" + i.intValue() + ";");
884bee435276 Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12653
diff changeset
267 }
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
268 for (Integer i : signed32) {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
269 codeBuffer.emitString(".reg .s32 %r" + i.intValue() + ";");
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
270 }
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 for (Integer i : signed64) {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
272 codeBuffer.emitString(".reg .s64 %r" + i.intValue() + ";");
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
273 }
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
274 for (Integer i : unsigned64) {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
275 codeBuffer.emitString(".reg .u64 %r" + i.intValue() + ";");
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
276 }
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
277 for (Integer i : float32) {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
278 codeBuffer.emitString(".reg .f32 %r" + i.intValue() + ";");
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
279 }
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 for (Integer i : float64) {
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
281 codeBuffer.emitString(".reg .f64 %r" + i.intValue() + ";");
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
282 }
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 }
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 @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
286 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
287 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
288 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
289 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
290 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
291 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
292 } 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
293 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
294 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
295 // 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
296 // 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
297 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
298 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
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 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
301 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
302 // 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
303 // 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
304 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
305 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
306 }
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
307 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
308 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
309 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
310 // 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
311 // 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
312 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
313 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
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 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
316 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
317 // 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
318 // 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
319 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
320 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
321 }
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
322 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
323 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
324 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
325 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
326 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
327 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
328 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
329 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
330 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
331 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
332 }
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
333 }
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
334 }
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
335 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
336 }
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
337 }
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
338
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
339 class PTXFrameContext implements FrameContext {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
340
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
341 @Override
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
342 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
343 // 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
344 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
345
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
346 @Override
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
347 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
348 }
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
349
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
350 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
351 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
352 }
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
353 }
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
354
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
355 @Override
13229
325b4e4efb60 added CompilationResultBuilderFactory to support peep-hole instrumentation of methods as their code is emitted
Doug Simon <doug.simon@oracle.com>
parents: 13227
diff changeset
356 public CompilationResultBuilder newCompilationResultBuilder(LIRGenerator lirGen, CompilationResult compilationResult, CompilationResultBuilderFactory factory) {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
357 // 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
358 // - 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
359 // - 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
360 // - 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
361 // - has no instructions with debug info
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
362 FrameMap frameMap = lirGen.frameMap;
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
363 AbstractAssembler masm = createAssembler(frameMap);
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
364 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
365 CompilationResultBuilder crb = factory.createBuilder(getCodeCache(), getForeignCalls(), frameMap, masm, frameContext, compilationResult);
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
366 crb.setFrameSize(0);
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
367 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
368 }
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
369
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
370 @Override
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
371 protected AbstractAssembler createAssembler(FrameMap frameMap) {
12762
884bee435276 Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12653
diff changeset
372 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
373 }
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
374
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
375 @Override
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
376 public LIRGenerator newLIRGenerator(StructuredGraph graph, FrameMap frameMap, CallingConvention cc, LIR lir) {
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
377 return new PTXHotSpotLIRGenerator(graph, getProviders(), getRuntime().getConfig(), frameMap, cc, lir);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
378 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
379
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
380 private static void emitKernelEntry(CompilationResultBuilder crb, LIRGenerator lirGen, ResolvedJavaMethod codeCacheOwner) {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
381 // 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
382 // 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
383 // 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
384 // facilitate seemless PTX code generation subsequently.
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
385 assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method";
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
386 final String name = codeCacheOwner.getName();
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
387 Buffer codeBuffer = crb.asm.codeBuffer;
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
388
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
389 // Emit initial boiler-plate directives.
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
390 codeBuffer.emitString(".version 3.0");
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
391 codeBuffer.emitString(".target sm_30");
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
392 codeBuffer.emitString0(".entry " + name + " (");
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
393 codeBuffer.emitString("");
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
394
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
395 // Get the start block
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
396 Block startBlock = lirGen.lir.cfg.getStartBlock();
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
397 // 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
398 // 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
399 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
400
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
401 // 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
402 // instruction.
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
403 for (LIRInstruction op : lirGen.lir.lir(startBlock)) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
404 if (op instanceof PTXParameterOp) {
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
405 op.emitCode(crb);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
406 deleteOps.add(op);
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 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
409
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
410 // Delete ParameterOp instructions.
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
411 for (LIRInstruction op : deleteOps) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
412 lirGen.lir.lir(startBlock).remove(op);
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
413 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
414
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
415 // Start emiting body of the PTX kernel.
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
416 codeBuffer.emitString0(") {");
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
417 codeBuffer.emitString("");
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
418 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
419
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
420 // Emit .reg space declarations
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
421 private static void emitRegisterDecl(CompilationResultBuilder crb, LIRGenerator lirGen, ResolvedJavaMethod codeCacheOwner) {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
422
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
423 assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method";
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
424
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
425 Buffer codeBuffer = crb.asm.codeBuffer;
12653
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
426 RegisterAnalysis registerAnalysis = new RegisterAnalysis();
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
427
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
428 for (Block b : lirGen.lir.codeEmittingOrder()) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
429 for (LIRInstruction op : lirGen.lir.lir(b)) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
430 if (op instanceof LabelOp) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
431 // 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
432 } else {
12653
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
433 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
434 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
435 op.forEachOutput(registerAnalysis);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
436 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
437 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
438 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
439
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
440 registerAnalysis.emitDeclarations(codeBuffer);
1a7e7011a341 * PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents: 12456
diff changeset
441
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
442 // emit predicate register declaration
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
443 int maxPredRegNum = ((PTXLIRGenerator) lirGen).getNextPredRegNumber();
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
444 if (maxPredRegNum > 0) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
445 codeBuffer.emitString(".reg .pred %p<" + maxPredRegNum + ">;");
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
446 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
447 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
448
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
449 @Override
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
450 public void emitCode(CompilationResultBuilder crb, LIRGenerator lirGen, ResolvedJavaMethod codeCacheOwner) {
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
451 assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method";
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
452 Buffer codeBuffer = crb.asm.codeBuffer;
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
453 // Emit the prologue
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
454 emitKernelEntry(crb, lirGen, codeCacheOwner);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
455
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
456 // Emit register declarations
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
457 try {
13227
1a66453f73db renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13225
diff changeset
458 emitRegisterDecl(crb, lirGen, codeCacheOwner);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
459 } catch (GraalInternalError e) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
460 e.printStackTrace();
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
461 // 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
462 // all types of parameters are handled.
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
463 codeBuffer.setPosition(0);
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
464 codeBuffer.close(false);
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
465 return;
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
466 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
467 // 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
468 try {
13318
da0851712519 moved emitting code for LIR and queries about whether an edge goes to its lexical successor "inside" CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents: 13234
diff changeset
469 crb.emit(lirGen.lir);
12456
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
470 } catch (GraalInternalError e) {
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
471 e.printStackTrace();
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
472 // 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
473 // all types of parameters are handled.
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
474 codeBuffer.setPosition(0);
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
475 codeBuffer.close(false);
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
476 return;
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
477 }
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
478
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
479 // Emit the epilogue
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
480 codeBuffer.emitString0("}");
f87c68d79f07 improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents: 12431
diff changeset
481 codeBuffer.emitString("");
11596
91e5f927af63 Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff changeset
482 }
13819
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
483
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
484 /**
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
485 * 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
486 */
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
487 public int getAvailableProcessors() {
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
488 if (!deviceInitialized) {
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
489 return 0;
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
490 }
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
491 return getAvailableProcessors0();
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
492 }
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
493
49db2c1e3bee added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents: 13778
diff changeset
494 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
495
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
496 }