Mercurial > hg > truffle
annotate graal/com.oracle.graal.hotspot.ptx/src/com/oracle/graal/hotspot/ptx/PTXHotSpotBackend.java @ 14143:8836f566b0bc
Remove withScope parameter on GraalCompiler.compileGraph.
author | Thomas Wuerthinger <thomas.wuerthinger@oracle.com> |
---|---|
date | Tue, 11 Mar 2014 16:28:43 +0100 |
parents | 4eac66a9b87d |
children | 4ff08c0366ae |
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.*; |
13649
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
40 import com.oracle.graal.debug.*; |
13766
1a81855dae10
removed unused PTX code
Doug Simon <doug.simon@oracle.com>
parents:
13758
diff
changeset
|
41 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
|
42 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
|
43 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
|
44 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
|
45 import com.oracle.graal.hotspot.meta.*; |
91e5f927af63
Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff
changeset
|
46 import com.oracle.graal.lir.*; |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
47 import com.oracle.graal.lir.LIRInstruction.OperandFlag; |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
48 import com.oracle.graal.lir.LIRInstruction.OperandMode; |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
49 import com.oracle.graal.lir.LIRInstruction.ValueProcedure; |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
50 import com.oracle.graal.lir.StandardOp.LabelOp; |
11596
91e5f927af63
Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff
changeset
|
51 import com.oracle.graal.lir.asm.*; |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
52 import com.oracle.graal.lir.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
|
53 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
|
54 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
|
55 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
|
56 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
|
57 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
|
58 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
|
59 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
|
60 |
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 * 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
|
63 */ |
91e5f927af63
Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff
changeset
|
64 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
|
65 |
13819
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
66 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
|
67 |
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
|
68 /** |
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
|
69 * 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
|
70 * |
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 * <pre> |
13649
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
72 * jlong (JavaThread* thread, |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
73 * jlong kernel, |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
74 * jint dimX, |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
75 * jint dimY, |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
76 * jint dimZ, |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
77 * jlong parametersAndReturnValueBuffer, |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
78 * 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
|
79 * 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
|
80 * 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
|
81 * jlong pinnedObjects, |
13649
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
82 * 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
|
83 * </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
|
84 */ |
13649
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
85 // @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
|
86 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
|
87 Word.class, // thread |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
88 long.class, // kernel |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
89 int.class, // dimX |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
90 int.class, // dimY |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
91 int.class, // dimZ |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
92 long.class, // parametersAndReturnValueBuffer |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
93 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
|
94 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
|
95 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
|
96 long.class, // pinnedObjects |
13649
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
97 int.class); // encodedReturnTypeSize |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
98 // @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
|
99 |
12431
7080a96be216
rename: graalRuntime -> runtime, getGraalRuntime -> getRuntime
Doug Simon <doug.simon@oracle.com>
parents:
12429
diff
changeset
|
100 public PTXHotSpotBackend(HotSpotGraalRuntime runtime, HotSpotProviders providers) { |
7080a96be216
rename: graalRuntime -> runtime, getGraalRuntime -> getRuntime
Doug Simon <doug.simon@oracle.com>
parents:
12429
diff
changeset
|
101 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
|
102 if (OmitDeviceInit) { |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
103 deviceInitialized = true; |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
104 } else { |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
105 boolean init = false; |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
106 try { |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
107 init = initialize(); |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
108 } catch (UnsatisfiedLinkError e) { |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
109 } |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
110 deviceInitialized = init; |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
111 } |
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
|
112 } |
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 |
13819
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
114 /** |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
115 * 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
|
116 * |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
117 * @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
|
118 */ |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
119 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
|
120 |
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
|
121 @Override |
11775
b04b94b71649
Finished PTX assembler and Register -> Variable conversion
Morris Meyer <morris.meyer@oracle.com>
parents:
11596
diff
changeset
|
122 public boolean shouldAllocateRegisters() { |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
123 return false; |
11775
b04b94b71649
Finished PTX assembler and Register -> Variable conversion
Morris Meyer <morris.meyer@oracle.com>
parents:
11596
diff
changeset
|
124 } |
b04b94b71649
Finished PTX assembler and Register -> Variable conversion
Morris Meyer <morris.meyer@oracle.com>
parents:
11596
diff
changeset
|
125 |
13649
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
126 /** |
13819
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
127 * 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
|
128 */ |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
129 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
|
130 |
11775
b04b94b71649
Finished PTX assembler and Register -> Variable conversion
Morris Meyer <morris.meyer@oracle.com>
parents:
11596
diff
changeset
|
131 @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
|
132 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
|
133 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
|
134 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
|
135 if (deviceInitialized) { |
13819
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
136 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
|
137 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
|
138 } |
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
|
139 /* 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
|
140 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
|
141 @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
|
142 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
|
143 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
|
144 } |
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 }); |
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
|
146 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
|
147 } |
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 |
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
|
149 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
|
150 |
13819
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
151 /** |
13888
51584f76462d
pulled Ptx and Hsail classes out of gpu class namespace
Doug Simon <doug.simon@oracle.com>
parents:
13839
diff
changeset
|
152 * 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
|
153 */ |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
154 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
|
155 |
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
|
156 @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
|
157 public FrameMap newFrameMap() { |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
158 return new PTXFrameMap(getCodeCache()); |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
159 } |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
160 |
13819
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
161 /** |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
162 * 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
|
163 */ |
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
|
164 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
|
165 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
|
166 } |
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 @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
|
169 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
|
170 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
|
171 // 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
|
172 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
|
173 } |
13651
1274feb0efe6
renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents:
13649
diff
changeset
|
174 return new GraphProducer() { |
1274feb0efe6
renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents:
13649
diff
changeset
|
175 |
1274feb0efe6
renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents:
13649
diff
changeset
|
176 public StructuredGraph getGraphFor(ResolvedJavaMethod method) { |
1274feb0efe6
renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents:
13649
diff
changeset
|
177 if (canOffloadToGPU(method)) { |
1274feb0efe6
renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents:
13649
diff
changeset
|
178 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
|
179 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
|
180 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
|
181 } |
1274feb0efe6
renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents:
13649
diff
changeset
|
182 return null; |
1274feb0efe6
renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents:
13649
diff
changeset
|
183 } |
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 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
|
186 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
|
187 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
|
188 } |
1274feb0efe6
renamed PTXLaunchKernelGraphKit to PTXWrapperBuilder
Doug Simon <doug.simon@oracle.com>
parents:
13649
diff
changeset
|
189 }; |
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
|
190 } |
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 |
13649
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
192 /** |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
193 * 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
|
194 * |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
195 * @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
|
196 * 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
|
197 * {@linkplain ExternalCompilationResult#getEntryPoint() entry point}. |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
198 * @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
|
199 */ |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
200 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
|
201 StructuredGraph graph = new StructuredGraph(method); |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
202 HotSpotProviders providers = getProviders(); |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
203 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
|
204 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
|
205 graphBuilderSuite.appendPhase(new NonNullParametersPhase()); |
13649
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
206 Suites suites = providers.getSuites().getDefaultSuites(); |
14142
4eac66a9b87d
Remove reference to graph in LIRGenerator.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
14065
diff
changeset
|
207 ExternalCompilationResult ptxCode = compileGraph(graph, null, cc, method, providers, this, this.getTarget(), null, graphBuilderSuite, OptimisticOptimizations.NONE, getProfilingInfo(graph), |
14143
8836f566b0bc
Remove withScope parameter on GraalCompiler.compileGraph.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
14142
diff
changeset
|
208 null, suites, new ExternalCompilationResult(), CompilationResultBuilderFactory.Default); |
13649
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
209 if (makeBinary) { |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
210 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
|
211 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
|
212 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
|
213 ptxCode.setEntryPoint(kernel); |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
214 } catch (Throwable e) { |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
215 throw Debug.handle(e); |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
216 } |
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 return ptxCode; |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
219 } |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
220 |
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
|
221 /** |
13819
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
222 * 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
|
223 */ |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
224 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
|
225 |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
226 /** |
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
|
227 * 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
|
228 * 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
|
229 * 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
|
230 * 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
|
231 */ |
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 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
|
233 |
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 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
|
235 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
|
236 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
|
237 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
|
238 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
|
239 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
|
240 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
|
241 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
|
242 } |
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 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
|
245 } |
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 return kernel; |
13649
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
247 } |
36d4faef2c56
moved all PTX compilation harness logic in PTXHotSpotBackend
Doug Simon <doug.simon@oracle.com>
parents:
13638
diff
changeset
|
248 |
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
|
249 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
|
250 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
|
251 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
|
252 |
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 // 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
|
254 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
|
255 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
|
256 |
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 // 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
|
258 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
|
259 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
|
260 |
1a7e7011a341
* PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents:
12456
diff
changeset
|
261 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
|
262 |
14032
d1c1f103d42c
renamed com.oracle.graal.asm.AbstractAssembler to com.oracle.graal.asm.Assembler
twisti
parents:
14031
diff
changeset
|
263 void emitDeclarations(Assembler asm) { |
12762
884bee435276
Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents:
12653
diff
changeset
|
264 for (Integer i : unsigned8) { |
14031
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
265 asm.emitString(".reg .u8 %r" + i.intValue() + ";"); |
12762
884bee435276
Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents:
12653
diff
changeset
|
266 } |
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
|
267 for (Integer i : signed32) { |
14031
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
268 asm.emitString(".reg .s32 %r" + i.intValue() + ";"); |
12653
1a7e7011a341
* PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents:
12456
diff
changeset
|
269 } |
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 for (Integer i : signed64) { |
14031
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
271 asm.emitString(".reg .s64 %r" + i.intValue() + ";"); |
12653
1a7e7011a341
* PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents:
12456
diff
changeset
|
272 } |
1a7e7011a341
* PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents:
12456
diff
changeset
|
273 for (Integer i : unsigned64) { |
14031
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
274 asm.emitString(".reg .u64 %r" + i.intValue() + ";"); |
12653
1a7e7011a341
* PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents:
12456
diff
changeset
|
275 } |
1a7e7011a341
* PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents:
12456
diff
changeset
|
276 for (Integer i : float32) { |
14031
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
277 asm.emitString(".reg .f32 %r" + i.intValue() + ";"); |
12653
1a7e7011a341
* PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents:
12456
diff
changeset
|
278 } |
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 for (Integer i : float64) { |
14031
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
280 asm.emitString(".reg .f64 %r" + i.intValue() + ";"); |
12653
1a7e7011a341
* PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents:
12456
diff
changeset
|
281 } |
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 @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
|
285 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
|
286 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
|
287 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
|
288 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
|
289 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
|
290 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
|
291 } 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
|
292 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
|
293 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
|
294 // 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
|
295 // 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
|
296 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
|
297 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
|
298 } |
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 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
|
300 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
|
301 // 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
|
302 // 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
|
303 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
|
304 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
|
305 } |
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 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
|
307 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
|
308 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
|
309 // 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
|
310 // 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
|
311 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
|
312 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
|
313 } |
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 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
|
315 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
|
316 // 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
|
317 // 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
|
318 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
|
319 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
|
320 } |
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 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
|
322 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
|
323 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
|
324 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
|
325 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
|
326 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
|
327 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
|
328 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
|
329 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
|
330 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
|
331 } |
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 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
|
335 } |
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 |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
338 class PTXFrameContext implements FrameContext { |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
339 |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
340 @Override |
13227
1a66453f73db
renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents:
13225
diff
changeset
|
341 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
|
342 // 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
|
343 } |
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 @Override |
13227
1a66453f73db
renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents:
13225
diff
changeset
|
346 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
|
347 } |
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
|
348 |
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 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
|
350 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
|
351 } |
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
|
352 } |
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 @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
|
355 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
|
356 // 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
|
357 // - 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
|
358 // - 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
|
359 // - 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
|
360 // - 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
|
361 FrameMap frameMap = lirGen.frameMap; |
14032
d1c1f103d42c
renamed com.oracle.graal.asm.AbstractAssembler to com.oracle.graal.asm.Assembler
twisti
parents:
14031
diff
changeset
|
362 Assembler masm = createAssembler(frameMap); |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
363 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
|
364 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
|
365 crb.setFrameSize(0); |
1a66453f73db
renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents:
13225
diff
changeset
|
366 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
|
367 } |
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 @Override |
14032
d1c1f103d42c
renamed com.oracle.graal.asm.AbstractAssembler to com.oracle.graal.asm.Assembler
twisti
parents:
14031
diff
changeset
|
370 protected Assembler createAssembler(FrameMap frameMap) { |
12762
884bee435276
Implement support for passing byte arguments in Java methods to be compiled to PTX.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents:
12653
diff
changeset
|
371 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
|
372 } |
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 @Override |
14142
4eac66a9b87d
Remove reference to graph in LIRGenerator.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
14065
diff
changeset
|
375 public LIRGenerator newLIRGenerator(StructuredGraph graph, Object stub, 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
|
376 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
|
377 } |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
378 |
14065
5dec26f3d4a4
Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
14032
diff
changeset
|
379 private static void emitKernelEntry(CompilationResultBuilder crb, LIR lir, ResolvedJavaMethod codeCacheOwner) { |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
380 // 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
|
381 // 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
|
382 // 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
|
383 // facilitate seemless PTX code generation subsequently. |
14065
5dec26f3d4a4
Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
14032
diff
changeset
|
384 assert codeCacheOwner != null : lir + " is not associated with a method"; |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
385 final String name = codeCacheOwner.getName(); |
14032
d1c1f103d42c
renamed com.oracle.graal.asm.AbstractAssembler to com.oracle.graal.asm.Assembler
twisti
parents:
14031
diff
changeset
|
386 Assembler asm = crb.asm; |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
387 |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
388 // Emit initial boiler-plate directives. |
14031
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
389 asm.emitString(".version 3.0"); |
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
390 asm.emitString(".target sm_30"); |
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
391 asm.emitString0(".entry " + name + " ("); |
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
392 asm.emitString(""); |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
393 |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
394 // Get the start block |
14065
5dec26f3d4a4
Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
14032
diff
changeset
|
395 Block startBlock = lir.cfg.getStartBlock(); |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
396 // 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
|
397 // 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
|
398 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
|
399 |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
400 // 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
|
401 // instruction. |
14065
5dec26f3d4a4
Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
14032
diff
changeset
|
402 for (LIRInstruction op : lir.lir(startBlock)) { |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
403 if (op instanceof PTXParameterOp) { |
13227
1a66453f73db
renamed TargetMethodAssembler to CompilationResultBuilder
Doug Simon <doug.simon@oracle.com>
parents:
13225
diff
changeset
|
404 op.emitCode(crb); |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
405 deleteOps.add(op); |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
406 } |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
407 } |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
408 |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
409 // Delete ParameterOp instructions. |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
410 for (LIRInstruction op : deleteOps) { |
14065
5dec26f3d4a4
Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
14032
diff
changeset
|
411 lir.lir(startBlock).remove(op); |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
412 } |
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 // Start emiting body of the PTX kernel. |
14031
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
415 asm.emitString0(") {"); |
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
416 asm.emitString(""); |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
417 } |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
418 |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
419 // Emit .reg space declarations |
14065
5dec26f3d4a4
Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
14032
diff
changeset
|
420 private static void emitRegisterDecl(CompilationResultBuilder crb, LIR lir, ResolvedJavaMethod codeCacheOwner) { |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
421 |
14065
5dec26f3d4a4
Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
14032
diff
changeset
|
422 assert codeCacheOwner != null : lir + " is not associated with a method"; |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
423 |
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
|
424 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
|
425 |
14065
5dec26f3d4a4
Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
14032
diff
changeset
|
426 for (Block b : lir.codeEmittingOrder()) { |
5dec26f3d4a4
Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
14032
diff
changeset
|
427 for (LIRInstruction op : lir.lir(b)) { |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
428 if (op instanceof LabelOp) { |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
429 // 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
|
430 } 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
|
431 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
|
432 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
|
433 op.forEachOutput(registerAnalysis); |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
434 } |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
435 } |
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 |
14032
d1c1f103d42c
renamed com.oracle.graal.asm.AbstractAssembler to com.oracle.graal.asm.Assembler
twisti
parents:
14031
diff
changeset
|
438 Assembler asm = crb.asm; |
14031
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
439 registerAnalysis.emitDeclarations(asm); |
12653
1a7e7011a341
* PTX kernel argument buffer now has naturally aligned arguments as required by PTX JIT compiler.
S.Bharadwaj Yadavalli <bharadwaj.yadavalli@oracle.com>
parents:
12456
diff
changeset
|
440 |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
441 // emit predicate register declaration |
14065
5dec26f3d4a4
Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
14032
diff
changeset
|
442 int maxPredRegNum = lir.numVariables(); |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
443 if (maxPredRegNum > 0) { |
14031
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
444 asm.emitString(".reg .pred %p<" + maxPredRegNum + ">;"); |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
445 } |
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 @Override |
14065
5dec26f3d4a4
Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
14032
diff
changeset
|
449 public void emitCode(CompilationResultBuilder crb, LIR lir, ResolvedJavaMethod codeCacheOwner) { |
5dec26f3d4a4
Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
14032
diff
changeset
|
450 assert codeCacheOwner != null : lir + " is not associated with a method"; |
14032
d1c1f103d42c
renamed com.oracle.graal.asm.AbstractAssembler to com.oracle.graal.asm.Assembler
twisti
parents:
14031
diff
changeset
|
451 Assembler asm = crb.asm; |
14031
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
452 |
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 |
14065
5dec26f3d4a4
Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
14032
diff
changeset
|
454 emitKernelEntry(crb, lir, codeCacheOwner); |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
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 { |
14065
5dec26f3d4a4
Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
14032
diff
changeset
|
458 emitRegisterDecl(crb, lir, codeCacheOwner); |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
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. |
14031
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
463 asm.close(false); |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
464 return; |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
465 } |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
466 // 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
|
467 try { |
14065
5dec26f3d4a4
Use LIR instead of LIRGenerator as parameter in emitCode.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
14032
diff
changeset
|
468 crb.emit(lir); |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
469 } catch (GraalInternalError e) { |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
470 e.printStackTrace(); |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
471 // 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
|
472 // all types of parameters are handled. |
14031
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
473 asm.close(false); |
12456
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
474 return; |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
475 } |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
476 |
f87c68d79f07
improved support for co-existing, multiple backends (GRAAL-363)
Doug Simon <doug.simon@oracle.com>
parents:
12431
diff
changeset
|
477 // Emit the epilogue |
14031
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
478 asm.emitString0("}"); |
390c4b742890
made com.oracle.graal.asm.Buffer non-public and a private field in AbstractAssembler
twisti
parents:
13922
diff
changeset
|
479 asm.emitString(""); |
11596
91e5f927af63
Initial implementation of PTXRuntime (RegisterConfig, PTX description etc); guarded with new flag UseGPU. Specify -XX:+UseGPU to exercise this new implementation.
bharadwaj
parents:
diff
changeset
|
480 } |
13819
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
481 |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
482 /** |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
483 * 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
|
484 */ |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
485 public int getAvailableProcessors() { |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
486 if (!deviceInitialized) { |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
487 return 0; |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
488 } |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
489 return getAvailableProcessors0(); |
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 |
49db2c1e3bee
added support for co-existing GPU backends (JBS:GRAAL-1)
Doug Simon <doug.simon@oracle.com>
parents:
13778
diff
changeset
|
492 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
|
493 |
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
|
494 } |