Mercurial > hg > truffle

annotate graal/com.oracle.graal.compiler.ptx/src/com/oracle/graal/compiler/ptx/PTXBackend.java @ 10880:7bd19a37f764

Find changesets by keywords (author, files, the commit message), revision number or hash, or revset expression.
PTX support for Linux
author Morris Meyer <morris.meyer@oracle.com>
date Thu, 25 Jul 2013 22:17:37 -0400
parents d55f24eac4b1
children f9215ee02538
Ignore whitespace changes - Everywhere: Within whitespace: At end of lines:
rev   line source
7804
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
1 /*
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
2 * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved.
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
3 * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER.
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
4 *
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
5 * This code is free software; you can redistribute it and/or modify it
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
6 * under the terms of the GNU General Public License version 2 only, as
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
7 * published by the Free Software Foundation.
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
8 *
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
9 * This code is distributed in the hope that it will be useful, but WITHOUT
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
10 * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
11 * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
12 * version 2 for more details (a copy is included in the LICENSE file that
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
13 * accompanied this code).
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
14 *
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
15 * You should have received a copy of the GNU General Public License version
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
16 * 2 along with this work; if not, write to the Free Software Foundation,
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
17 * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA.
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
18 *
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
19 * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
20 * or visit www.oracle.com if you need additional information or have any
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
21 * questions.
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
22 */
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
23 package com.oracle.graal.compiler.ptx;
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
24
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
25 import com.oracle.graal.api.code.*;
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
26 import com.oracle.graal.api.meta.*;
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
27 import com.oracle.graal.asm.*;
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
28 import com.oracle.graal.asm.ptx.*;
7805
0e58445d54df Integration fixes.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 7804
diff changeset
29 import com.oracle.graal.compiler.gen.*;
0e58445d54df Integration fixes.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 7804
diff changeset
30 import com.oracle.graal.compiler.target.*;
7804
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
31 import com.oracle.graal.lir.*;
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
32 import com.oracle.graal.lir.asm.*;
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
33 import com.oracle.graal.nodes.*;
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
34
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
35 /**
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
36 * PTX specific backend.
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
37 */
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
38 public class PTXBackend extends Backend {
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
39
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
40 public PTXBackend(CodeCacheProvider runtime, TargetDescription target) {
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
41 super(runtime, target);
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
42 }
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
43
7805
0e58445d54df Integration fixes.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 7804
diff changeset
44 @Override
9613
0c17815817a4 removed LIRGenerator.method field
Doug Simon <doug.simon@oracle.com>
parents: 9612
diff changeset
45 public LIRGenerator newLIRGenerator(StructuredGraph graph, FrameMap frameMap, CallingConvention cc, LIR lir) {
0c17815817a4 removed LIRGenerator.method field
Doug Simon <doug.simon@oracle.com>
parents: 9612
diff changeset
46 return new PTXLIRGenerator(graph, runtime(), target, frameMap, cc, lir);
7804
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
47 }
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
48
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
49 class HotSpotFrameContext implements FrameContext {
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
50
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
51 @Override
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
52 public void enter(TargetMethodAssembler tasm) {
7805
0e58445d54df Integration fixes.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 7804
diff changeset
53 // codeBuffer.emitString(".address_size 32"); // PTX ISA version 2.3
7804
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
54 }
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
55
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
56 @Override
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
57 public void leave(TargetMethodAssembler tasm) {
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
58 }
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
59 }
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
60
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
61 @Override
9432
6680389bd36f Make assembler creation in backend more extensible.
Roland Schatz <roland.schatz@oracle.com>
parents: 9430
diff changeset
62 protected AbstractAssembler createAssembler(FrameMap frameMap) {
6680389bd36f Make assembler creation in backend more extensible.
Roland Schatz <roland.schatz@oracle.com>
parents: 9430
diff changeset
63 return new PTXAssembler(target, frameMap.registerConfig);
6680389bd36f Make assembler creation in backend more extensible.
Roland Schatz <roland.schatz@oracle.com>
parents: 9430
diff changeset
64 }
6680389bd36f Make assembler creation in backend more extensible.
Roland Schatz <roland.schatz@oracle.com>
parents: 9430
diff changeset
65
6680389bd36f Make assembler creation in backend more extensible.
Roland Schatz <roland.schatz@oracle.com>
parents: 9430
diff changeset
66 @Override
8281
8fde1be81b2d LIRGenerator is transmitted across backend passes instead of just the LIR so that backend-specific, per-compilation information can be attached to the backend-specific LIRGenerator object
Doug Simon <doug.simon@oracle.com>
parents: 8201
diff changeset
67 public TargetMethodAssembler newAssembler(LIRGenerator lirGen, CompilationResult compilationResult) {
7804
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
68 // Omit the frame if the method:
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
69 // - has no spill slots or other slots allocated during register allocation
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
70 // - has no callee-saved registers
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
71 // - has no incoming arguments passed on the stack
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
72 // - has no instructions with debug info
8281
8fde1be81b2d LIRGenerator is transmitted across backend passes instead of just the LIR so that backend-specific, per-compilation information can be attached to the backend-specific LIRGenerator object
Doug Simon <doug.simon@oracle.com>
parents: 8201
diff changeset
73 FrameMap frameMap = lirGen.frameMap;
9432
6680389bd36f Make assembler creation in backend more extensible.
Roland Schatz <roland.schatz@oracle.com>
parents: 9430
diff changeset
74 AbstractAssembler masm = createAssembler(frameMap);
8281
8fde1be81b2d LIRGenerator is transmitted across backend passes instead of just the LIR so that backend-specific, per-compilation information can be attached to the backend-specific LIRGenerator object
Doug Simon <doug.simon@oracle.com>
parents: 8201
diff changeset
75 HotSpotFrameContext frameContext = new HotSpotFrameContext();
9430
147162b27799 GRAAL-234 - PTX code loading
Morris Meyer <morris.meyer@oracle.com>
parents: 8995
diff changeset
76 TargetMethodAssembler tasm = new PTXTargetMethodAssembler(target, runtime(), frameMap, masm, frameContext, compilationResult);
7804
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
77 tasm.setFrameSize(frameMap.frameSize());
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
78 return tasm;
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
79 }
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
80
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
81 @Override
9621
e97dc9bbfedc introduced installedCodeOwner parameter of type ResolvedJavaMethod to GraalCompiler.compileGraph to properly distinguish the source method of a graph from the method under which the code compiled for the graph will be installed
Doug Simon <doug.simon@oracle.com>
parents: 9615
diff changeset
82 public void emitCode(TargetMethodAssembler tasm, LIRGenerator lirGen, ResolvedJavaMethod codeCacheOwner) {
7804
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
83 // Emit the prologue
9621
e97dc9bbfedc introduced installedCodeOwner parameter of type ResolvedJavaMethod to GraalCompiler.compileGraph to properly distinguish the source method of a graph from the method under which the code compiled for the graph will be installed
Doug Simon <doug.simon@oracle.com>
parents: 9615
diff changeset
84 assert codeCacheOwner != null : lirGen.getGraph() + " is not associated with a method";
e97dc9bbfedc introduced installedCodeOwner parameter of type ResolvedJavaMethod to GraalCompiler.compileGraph to properly distinguish the source method of a graph from the method under which the code compiled for the graph will be installed
Doug Simon <doug.simon@oracle.com>
parents: 9615
diff changeset
85 final String name = codeCacheOwner.getName();
7804
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
86 Buffer codeBuffer = tasm.asm.codeBuffer;
9430
147162b27799 GRAAL-234 - PTX code loading
Morris Meyer <morris.meyer@oracle.com>
parents: 8995
diff changeset
87 codeBuffer.emitString(".version 1.4");
147162b27799 GRAAL-234 - PTX code loading
Morris Meyer <morris.meyer@oracle.com>
parents: 8995
diff changeset
88 codeBuffer.emitString(".target sm_10");
7804
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
89 codeBuffer.emitString0(".entry " + name + " (");
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
90 codeBuffer.emitString("");
7805
0e58445d54df Integration fixes.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 7804
diff changeset
91
9621
e97dc9bbfedc introduced installedCodeOwner parameter of type ResolvedJavaMethod to GraalCompiler.compileGraph to properly distinguish the source method of a graph from the method under which the code compiled for the graph will be installed
Doug Simon <doug.simon@oracle.com>
parents: 9615
diff changeset
92 Signature signature = codeCacheOwner.getSignature();
10879
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
93 int paramCount = signature.getParameterCount(false);
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
94 // TODO - Revisit this.
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
95 // Bit-size of registers to be declared and used by the kernel.
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
96 int regSize = 32;
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
97 for (int i = 0; i < paramCount; i++) {
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
98 String param;
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
99 // No unsigned types in Java. So using .s specifier
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
100 switch (signature.getParameterKind(i)) {
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
101 case Boolean:
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
102 case Byte:
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
103 param = ".param .s8 param" + i;
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
104 regSize = 8;
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
105 break;
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
106 case Char:
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
107 case Short:
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
108 param = ".param .s16 param" + i;
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
109 regSize = 16;
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
110 break;
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
111 case Int:
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
112 param = ".param .s32 param" + i;
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
113 regSize = 32;
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
114 break;
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
115 case Long:
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
116 case Float:
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
117 case Double:
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
118 case Void:
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
119 param = ".param .s64 param" + i;
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
120 regSize = 32;
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
121 break;
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
122 default:
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
123 // Not sure but specify 64-bit specifier??
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
124 param = ".param .s64 param" + i;
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
125 break;
10880
7bd19a37f764 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 10879
diff changeset
126 }
7bd19a37f764 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 10879
diff changeset
127 if (i != (paramCount -1 )) {
10879
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
128 param += ",";
d55f24eac4b1 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 9621
diff changeset
129 }
7804
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
130 codeBuffer.emitString(param);
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
131 }
7805
0e58445d54df Integration fixes.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 7804
diff changeset
132
7804
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
133 codeBuffer.emitString0(") {");
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
134 codeBuffer.emitString("");
7805
0e58445d54df Integration fixes.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 7804
diff changeset
135
7804
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
136 // XXX For now declare one predicate and all registers
8995
585cc62fcdc5 PTX enhancements - arithmetic, control, float, integer math, control and basic switch
Morris Meyer <morris.meyer@oracle.com>
parents: 8463
diff changeset
137 codeBuffer.emitString(" .reg .pred %p,%q;");
10880
7bd19a37f764 PTX support for Linux
Morris Meyer <morris.meyer@oracle.com>
parents: 10879
diff changeset
138 codeBuffer.emitString(" .reg .s" + regSize + " %r<16>;");
7805
0e58445d54df Integration fixes.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents: 7804
diff changeset
139
7804
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
140 // Emit code for the LIR
8281
8fde1be81b2d LIRGenerator is transmitted across backend passes instead of just the LIR so that backend-specific, per-compilation information can be attached to the backend-specific LIRGenerator object
Doug Simon <doug.simon@oracle.com>
parents: 8201
diff changeset
141 lirGen.lir.emitCode(tasm);
7804
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
142
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
143 // Emit the epilogue
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
144 codeBuffer.emitString0("}");
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
145 codeBuffer.emitString("");
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
146 }
447f9ba1962b Experimental PTX backend. Contribution by Christian Thalinger.
Thomas Wuerthinger <thomas.wuerthinger@oracle.com>
parents:
diff changeset
147 }