1 /* 2 * Copyright (c) 2013, Oracle and/or its affiliates. All rights reserved. 3 * DO NOT ALTER OR REMOVE COPYRIGHT NOTICES OR THIS FILE HEADER. 4 * 5 * This code is free software; you can redistribute it and/or modify it 6 * under the terms of the GNU General Public License version 2 only, as 7 * published by the Free Software Foundation. 8 * 9 * This code is distributed in the hope that it will be useful, but WITHOUT 10 * ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or 11 * FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License 12 * version 2 for more details (a copy is included in the LICENSE file that 13 * accompanied this code). 14 * 15 * You should have received a copy of the GNU General Public License version 16 * 2 along with this work; if not, write to the Free Software Foundation, 17 * Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301 USA. 18 * 19 * Please contact Oracle, 500 Oracle Parkway, Redwood Shores, CA 94065 USA 20 * or visit www.oracle.com if you need additional information or have any 21 * questions. 22 */ 23 package com.oracle.graal.compiler.ptx; 24 25 import com.oracle.graal.api.code.*; 26 import com.oracle.graal.api.meta.*; 27 import com.oracle.graal.asm.*; 28 import com.oracle.graal.asm.ptx.*; 29 import com.oracle.graal.compiler.gen.*; 30 import com.oracle.graal.compiler.target.*; 31 import com.oracle.graal.lir.*; 32 import com.oracle.graal.lir.asm.*; 33 import com.oracle.graal.nodes.*; 34 35 /** 36 * PTX specific backend. 37 */ 38 public class PTXBackend extends Backend { 39 40 public PTXBackend(CodeCacheProvider runtime, TargetDescription target) { 41 super(runtime, target); 42 } 43 44 @Override 45 public LIRGenerator newLIRGenerator(StructuredGraph graph, FrameMap frameMap, ResolvedJavaMethod method, LIR lir) { 46 return new PTXLIRGenerator(graph, runtime(), target, frameMap, method, lir); 47 } 48 49 class HotSpotFrameContext implements FrameContext { 50 51 @Override 52 public void enter(TargetMethodAssembler tasm) { 53 Buffer codeBuffer = tasm.asm.codeBuffer; 54 codeBuffer.emitString(".version 1.4"); 55 codeBuffer.emitString(".target sm_10"); 56 // codeBuffer.emitString(".address_size 32"); // PTX ISA version 2.3 57 } 58 59 @Override 60 public void leave(TargetMethodAssembler tasm) { 61 } 62 } 63 64 @Override 65 public TargetMethodAssembler newAssembler(LIRGenerator lirGen, CompilationResult compilationResult) { 66 // Omit the frame if the method: 67 // - has no spill slots or other slots allocated during register allocation 68 // - has no callee-saved registers 69 // - has no incoming arguments passed on the stack 70 // - has no instructions with debug info 71 FrameMap frameMap = lirGen.frameMap; 72 AbstractAssembler masm = new PTXAssembler(target, frameMap.registerConfig); 73 HotSpotFrameContext frameContext = new HotSpotFrameContext(); 74 TargetMethodAssembler tasm = new TargetMethodAssembler(target, runtime(), frameMap, masm, frameContext, compilationResult); 75 tasm.setFrameSize(frameMap.frameSize()); 76 return tasm; 77 } 78 79 @Override 80 public void emitCode(TargetMethodAssembler tasm, ResolvedJavaMethod method, LIRGenerator lirGen) { 81 // Emit the prologue 82 final String name = method.getName(); 83 Buffer codeBuffer = tasm.asm.codeBuffer; 84 codeBuffer.emitString0(".entry " + name + " ("); 85 codeBuffer.emitString(""); 86 87 Signature signature = method.getSignature(); 88 for (int i = 0; i < signature.getParameterCount(false); i++) { 89 String param = ".param .u32 param" + i; 90 codeBuffer.emitString(param); 91 } 92 93 codeBuffer.emitString0(") {"); 94 codeBuffer.emitString(""); 95 96 // XXX For now declare one predicate and all registers 97 codeBuffer.emitString(" .reg .pred %p,%q;"); 98 codeBuffer.emitString(" .reg .u32 %r<16>;"); 99 100 // Emit code for the LIR 101 lirGen.lir.emitCode(tasm); 102 103 // Emit the epilogue 104 codeBuffer.emitString0("}"); 105 codeBuffer.emitString(""); 106 } 107 }