149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski//===-- NVPTXAsmPrinter.cpp - NVPTX LLVM assembly writer ------------------===// 249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// 349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// The LLVM Compiler Infrastructure 449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// 549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// This file is distributed under the University of Illinois Open Source 649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// License. See LICENSE.TXT for details. 749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// 849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski//===----------------------------------------------------------------------===// 949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// 1049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// This file contains a printer that converts from our internal representation 1149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// of machine-dependent LLVM code to NVPTX assembly language. 1249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// 1349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski//===----------------------------------------------------------------------===// 1449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 150bcbd1df7a204e1e512f1a27066d725309de1b13Bill Wendling#include "NVPTXAsmPrinter.h" 16d04a8d4b33ff316ca4cf961e06c9e312eff8e64fChandler Carruth#include "MCTargetDesc/NVPTXMCAsmInfo.h" 1749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski#include "NVPTX.h" 1849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski#include "NVPTXInstrInfo.h" 19d04a8d4b33ff316ca4cf961e06c9e312eff8e64fChandler Carruth#include "NVPTXNumRegisters.h" 2049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski#include "NVPTXRegisterInfo.h" 21d04a8d4b33ff316ca4cf961e06c9e312eff8e64fChandler Carruth#include "NVPTXTargetMachine.h" 220bcbd1df7a204e1e512f1a27066d725309de1b13Bill Wendling#include "NVPTXUtilities.h" 23d04a8d4b33ff316ca4cf961e06c9e312eff8e64fChandler Carruth#include "cl_common_defines.h" 2449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski#include "llvm/ADT/StringExtras.h" 25d04a8d4b33ff316ca4cf961e06c9e312eff8e64fChandler Carruth#include "llvm/Analysis/ConstantFolding.h" 26d04a8d4b33ff316ca4cf961e06c9e312eff8e64fChandler Carruth#include "llvm/Assembly/Writer.h" 2749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski#include "llvm/CodeGen/Analysis.h" 2849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski#include "llvm/CodeGen/MachineFrameInfo.h" 2949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski#include "llvm/CodeGen/MachineModuleInfo.h" 30d04a8d4b33ff316ca4cf961e06c9e312eff8e64fChandler Carruth#include "llvm/CodeGen/MachineRegisterInfo.h" 31d04a8d4b33ff316ca4cf961e06c9e312eff8e64fChandler Carruth#include "llvm/DebugInfo.h" 320b8c9a80f20772c3793201ab5b251d3520b9cea3Chandler Carruth#include "llvm/IR/DerivedTypes.h" 330b8c9a80f20772c3793201ab5b251d3520b9cea3Chandler Carruth#include "llvm/IR/Function.h" 340b8c9a80f20772c3793201ab5b251d3520b9cea3Chandler Carruth#include "llvm/IR/GlobalVariable.h" 350b8c9a80f20772c3793201ab5b251d3520b9cea3Chandler Carruth#include "llvm/IR/Module.h" 360b8c9a80f20772c3793201ab5b251d3520b9cea3Chandler Carruth#include "llvm/IR/Operator.h" 3749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski#include "llvm/MC/MCStreamer.h" 3849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski#include "llvm/MC/MCSymbol.h" 39d04a8d4b33ff316ca4cf961e06c9e312eff8e64fChandler Carruth#include "llvm/Support/CommandLine.h" 4049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski#include "llvm/Support/ErrorHandling.h" 4149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski#include "llvm/Support/FormattedStream.h" 4249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski#include "llvm/Support/Path.h" 43d04a8d4b33ff316ca4cf961e06c9e312eff8e64fChandler Carruth#include "llvm/Support/TargetRegistry.h" 44d04a8d4b33ff316ca4cf961e06c9e312eff8e64fChandler Carruth#include "llvm/Support/TimeValue.h" 45d04a8d4b33ff316ca4cf961e06c9e312eff8e64fChandler Carruth#include "llvm/Target/Mangler.h" 46d04a8d4b33ff316ca4cf961e06c9e312eff8e64fChandler Carruth#include "llvm/Target/TargetLoweringObjectFile.h" 470bcbd1df7a204e1e512f1a27066d725309de1b13Bill Wendling#include <sstream> 4849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskiusing namespace llvm; 4949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 5049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 5149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski#include "NVPTXGenAsmWriter.inc" 5249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 5349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskibool RegAllocNilUsed = true; 5449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 5549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski#define DEPOTNAME "__local_depot" 5649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 5749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskistatic cl::opt<bool> 5849683f3c961379fbc088871a5d6304950f1f1cbcJustin HolewinskiEmitLineNumbers("nvptx-emit-line-numbers", 5949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski cl::desc("NVPTX Specific: Emit Line numbers even without -G"), 6049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski cl::init(true)); 6149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 6249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskinamespace llvm { 6349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskibool InterleaveSrcInPtx = false; 6449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 6549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 6649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskistatic cl::opt<bool, true>InterleaveSrc("nvptx-emit-src", 6749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski cl::ZeroOrMore, 6849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski cl::desc("NVPTX Specific: Emit source line in ptx file"), 6949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski cl::location(llvm::InterleaveSrcInPtx)); 7049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 7149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 722085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinskinamespace { 732085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski/// DiscoverDependentGlobals - Return a set of GlobalVariables on which \p V 742085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski/// depends. 752085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinskivoid DiscoverDependentGlobals(Value *V, 762085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski DenseSet<GlobalVariable*> &Globals) { 772085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski if (GlobalVariable *GV = dyn_cast<GlobalVariable>(V)) 782085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski Globals.insert(GV); 792085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski else { 802085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski if (User *U = dyn_cast<User>(V)) { 812085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski for (unsigned i = 0, e = U->getNumOperands(); i != e; ++i) { 822085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski DiscoverDependentGlobals(U->getOperand(i), Globals); 832085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski } 842085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski } 852085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski } 862085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski} 8749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 882085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski/// VisitGlobalVariableForEmission - Add \p GV to the list of GlobalVariable 892085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski/// instances to be emitted, but only after any dependents have been added 902085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski/// first. 912085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinskivoid VisitGlobalVariableForEmission(GlobalVariable *GV, 922085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski SmallVectorImpl<GlobalVariable*> &Order, 932085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski DenseSet<GlobalVariable*> &Visited, 942085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski DenseSet<GlobalVariable*> &Visiting) { 952085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski // Have we already visited this one? 962085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski if (Visited.count(GV)) return; 972085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski 982085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski // Do we have a circular dependency? 992085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski if (Visiting.count(GV)) 1002085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski report_fatal_error("Circular dependency found in global variable set"); 1012085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski 1022085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski // Start visiting this global 1032085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski Visiting.insert(GV); 1042085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski 1052085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski // Make sure we visit all dependents first 1062085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski DenseSet<GlobalVariable*> Others; 1072085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski for (unsigned i = 0, e = GV->getNumOperands(); i != e; ++i) 1082085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski DiscoverDependentGlobals(GV->getOperand(i), Others); 1092085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski 1102085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski for (DenseSet<GlobalVariable*>::iterator I = Others.begin(), 1112085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski E = Others.end(); I != E; ++I) 1122085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski VisitGlobalVariableForEmission(*I, Order, Visited, Visiting); 1132085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski 1142085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski // Now we can visit ourself 1152085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski Order.push_back(GV); 1162085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski Visited.insert(GV); 1172085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski Visiting.erase(GV); 1182085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski} 1192085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski} 12049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 12149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// @TODO: This is a copy from AsmPrinter.cpp. The function is static, so we 12249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// cannot just link to the existing version. 12349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski/// LowerConstant - Lower the specified LLVM Constant to an MCExpr. 12449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski/// 12549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskiusing namespace nvptx; 12649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskiconst MCExpr *nvptx::LowerConstant(const Constant *CV, AsmPrinter &AP) { 12749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski MCContext &Ctx = AP.OutContext; 12849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 12949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (CV->isNullValue() || isa<UndefValue>(CV)) 13049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return MCConstantExpr::Create(0, Ctx); 13149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 13249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (const ConstantInt *CI = dyn_cast<ConstantInt>(CV)) 13349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return MCConstantExpr::Create(CI->getZExtValue(), Ctx); 13449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 13549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (const GlobalValue *GV = dyn_cast<GlobalValue>(CV)) 13649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return MCSymbolRefExpr::Create(AP.Mang->getSymbol(GV), Ctx); 13749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 13849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (const BlockAddress *BA = dyn_cast<BlockAddress>(CV)) 13949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return MCSymbolRefExpr::Create(AP.GetBlockAddressSymbol(BA), Ctx); 14049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 14149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const ConstantExpr *CE = dyn_cast<ConstantExpr>(CV); 14249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (CE == 0) 14349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski llvm_unreachable("Unknown constant value to lower!"); 14449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 14549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 14649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski switch (CE->getOpcode()) { 14749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski default: 14849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // If the code isn't optimized, there may be outstanding folding 1493574eca1b02600bac4e625297f4ecf745f4c4f32Micah Villmow // opportunities. Attempt to fold the expression using DataLayout as a 15049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // last resort before giving up. 15149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (Constant *C = 1523574eca1b02600bac4e625297f4ecf745f4c4f32Micah Villmow ConstantFoldConstantExpression(CE, AP.TM.getDataLayout())) 15349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (C != CE) 15449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return LowerConstant(C, AP); 15549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 15649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Otherwise report the problem to the user. 15749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski { 15849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski std::string S; 15949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_string_ostream OS(S); 16049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski OS << "Unsupported expression in static initializer: "; 16149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski WriteAsOperand(OS, CE, /*PrintType=*/false, 16249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski !AP.MF ? 0 : AP.MF->getFunction()->getParent()); 16349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski report_fatal_error(OS.str()); 16449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 16549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::GetElementPtr: { 1663574eca1b02600bac4e625297f4ecf745f4c4f32Micah Villmow const DataLayout &TD = *AP.TM.getDataLayout(); 16749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Generate a symbolic expression for the byte address 16898281a20503896349bd152e2dfe87435d3a6aadaNuno Lopes APInt OffsetAI(TD.getPointerSizeInBits(), 0); 16998281a20503896349bd152e2dfe87435d3a6aadaNuno Lopes cast<GEPOperator>(CE)->accumulateConstantOffset(TD, OffsetAI); 17049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 17149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const MCExpr *Base = LowerConstant(CE->getOperand(0), AP); 17298281a20503896349bd152e2dfe87435d3a6aadaNuno Lopes if (!OffsetAI) 17349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return Base; 17449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 17598281a20503896349bd152e2dfe87435d3a6aadaNuno Lopes int64_t Offset = OffsetAI.getSExtValue(); 17649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return MCBinaryExpr::CreateAdd(Base, MCConstantExpr::Create(Offset, Ctx), 17749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Ctx); 17849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 17949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 18049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::Trunc: 18149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // We emit the value and depend on the assembler to truncate the generated 18249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // expression properly. This is important for differences between 18349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // blockaddress labels. Since the two labels are in the same function, it 18449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // is reasonable to treat their delta as a 32-bit value. 18549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // FALL THROUGH. 18649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::BitCast: 18749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return LowerConstant(CE->getOperand(0), AP); 18849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 18949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::IntToPtr: { 1903574eca1b02600bac4e625297f4ecf745f4c4f32Micah Villmow const DataLayout &TD = *AP.TM.getDataLayout(); 19149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Handle casts to pointers by changing them into casts to the appropriate 19249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // integer type. This promotes constant folding and simplifies this code. 19349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Constant *Op = CE->getOperand(0); 194ece6c6bb6329748b92403c06ac87f45c43485911Chandler Carruth Op = ConstantExpr::getIntegerCast(Op, TD.getIntPtrType(CV->getContext()), 19549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski false/*ZExt*/); 19649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return LowerConstant(Op, AP); 19749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 19849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 19949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::PtrToInt: { 2003574eca1b02600bac4e625297f4ecf745f4c4f32Micah Villmow const DataLayout &TD = *AP.TM.getDataLayout(); 20149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Support only foldable casts to/from pointers that can be eliminated by 20249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // changing the pointer to the appropriately sized integer type. 20349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Constant *Op = CE->getOperand(0); 20449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Type *Ty = CE->getType(); 20549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 20649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const MCExpr *OpExpr = LowerConstant(Op, AP); 20749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 20849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // We can emit the pointer value into this slot if the slot is an 20949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // integer slot equal to the size of the pointer. 21049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (TD.getTypeAllocSize(Ty) == TD.getTypeAllocSize(Op->getType())) 21149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return OpExpr; 21249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 21349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Otherwise the pointer is smaller than the resultant integer, mask off 21449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // the high bits so we are sure to get a proper truncation if the input is 21549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // a constant expr. 21649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned InBits = TD.getTypeAllocSizeInBits(Op->getType()); 21749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const MCExpr *MaskExpr = MCConstantExpr::Create(~0ULL >> (64-InBits), Ctx); 21849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return MCBinaryExpr::CreateAnd(OpExpr, MaskExpr, Ctx); 21949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 22049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 22149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // The MC library also has a right-shift operator, but it isn't consistently 22249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // signed or unsigned between different targets. 22349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::Add: 22449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::Sub: 22549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::Mul: 22649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::SDiv: 22749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::SRem: 22849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::Shl: 22949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::And: 23049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::Or: 23149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::Xor: { 23249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const MCExpr *LHS = LowerConstant(CE->getOperand(0), AP); 23349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const MCExpr *RHS = LowerConstant(CE->getOperand(1), AP); 23449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski switch (CE->getOpcode()) { 23549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski default: llvm_unreachable("Unknown binary operator constant cast expr"); 23649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::Add: return MCBinaryExpr::CreateAdd(LHS, RHS, Ctx); 23749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::Sub: return MCBinaryExpr::CreateSub(LHS, RHS, Ctx); 23849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::Mul: return MCBinaryExpr::CreateMul(LHS, RHS, Ctx); 23949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::SDiv: return MCBinaryExpr::CreateDiv(LHS, RHS, Ctx); 24049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::SRem: return MCBinaryExpr::CreateMod(LHS, RHS, Ctx); 24149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::Shl: return MCBinaryExpr::CreateShl(LHS, RHS, Ctx); 24249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::And: return MCBinaryExpr::CreateAnd(LHS, RHS, Ctx); 24349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::Or: return MCBinaryExpr::CreateOr (LHS, RHS, Ctx); 24449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Instruction::Xor: return MCBinaryExpr::CreateXor(LHS, RHS, Ctx); 24549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 24649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 24749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 24849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 24949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 25049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 25149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::emitLineNumberAsDotLoc(const MachineInstr &MI) 25249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski{ 25349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!EmitLineNumbers) 25449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 25549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (ignoreLoc(MI)) 25649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 25749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 25849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski DebugLoc curLoc = MI.getDebugLoc(); 25949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 26049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (prevDebugLoc.isUnknown() && curLoc.isUnknown()) 26149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 26249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 26349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (prevDebugLoc == curLoc) 26449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 26549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 26649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski prevDebugLoc = curLoc; 26749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 26849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (curLoc.isUnknown()) 26949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 27049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 27149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 27249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const MachineFunction *MF = MI.getParent()->getParent(); 27349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski //const TargetMachine &TM = MF->getTarget(); 27449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 27549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const LLVMContext &ctx = MF->getFunction()->getContext(); 27649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski DIScope Scope(curLoc.getScope(ctx)); 27749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 27849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!Scope.Verify()) 27949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 28049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 28149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski StringRef fileName(Scope.getFilename()); 28249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski StringRef dirName(Scope.getDirectory()); 28349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski SmallString<128> FullPathName = dirName; 28449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!dirName.empty() && !sys::path::is_absolute(fileName)) { 28549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski sys::path::append(FullPathName, fileName); 28649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski fileName = FullPathName.str(); 28749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 28849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 28949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (filenameMap.find(fileName.str()) == filenameMap.end()) 29049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 29149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 29249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 29349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Emit the line from the source file. 29449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (llvm::InterleaveSrcInPtx) 29549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski this->emitSrcInText(fileName.str(), curLoc.getLine()); 29649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 29749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski std::stringstream temp; 29849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski temp << "\t.loc " << filenameMap[fileName.str()] 29949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski << " " << curLoc.getLine() << " " << curLoc.getCol(); 30049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski OutStreamer.EmitRawText(Twine(temp.str().c_str())); 30149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 30249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 30349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::EmitInstruction(const MachineInstr *MI) { 30449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski SmallString<128> Str; 30549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_svector_ostream OS(Str); 30649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA) 30749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitLineNumberAsDotLoc(*MI); 30849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printInstruction(MI, OS); 30949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski OutStreamer.EmitRawText(OS.str()); 31049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 31149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 31249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::printReturnValStr(const Function *F, 31349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_ostream &O) 31449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski{ 3153574eca1b02600bac4e625297f4ecf745f4c4f32Micah Villmow const DataLayout *TD = TM.getDataLayout(); 31649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const TargetLowering *TLI = TM.getTargetLowering(); 31749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 31849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Type *Ty = F->getReturnType(); 31949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 32049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski bool isABI = (nvptxSubtarget.getSmVersion() >= 20); 32149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 32249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (Ty->getTypeID() == Type::VoidTyID) 32349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 32449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 32549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " ("; 32649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 32749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (isABI) { 32849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (Ty->isPrimitiveType() || Ty->isIntegerTy()) { 32949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned size = 0; 33049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (const IntegerType *ITy = dyn_cast<IntegerType>(Ty)) { 33149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski size = ITy->getBitWidth(); 33249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (size < 32) size = 32; 33349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else { 33449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski assert(Ty->isFloatingPointTy() && 33549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski "Floating point type expected here"); 33649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski size = Ty->getPrimitiveSizeInBits(); 33749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 33849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 33949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".param .b" << size << " func_retval0"; 34049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 34149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else if (isa<PointerType>(Ty)) { 34249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".param .b" << TLI->getPointerTy().getSizeInBits() 34349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski << " func_retval0"; 34449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else { 34549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if ((Ty->getTypeID() == Type::StructTyID) || 34649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski isa<VectorType>(Ty)) { 34749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski SmallVector<EVT, 16> vtparts; 34849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ComputeValueVTs(*TLI, Ty, vtparts); 34949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned totalsz = 0; 35049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (unsigned i=0,e=vtparts.size(); i!=e; ++i) { 35149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned elems = 1; 35249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski EVT elemtype = vtparts[i]; 35349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (vtparts[i].isVector()) { 35449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski elems = vtparts[i].getVectorNumElements(); 35549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski elemtype = vtparts[i].getVectorElementType(); 35649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 35749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (unsigned j=0, je=elems; j!=je; ++j) { 35849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned sz = elemtype.getSizeInBits(); 35949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (elemtype.isInteger() && (sz < 8)) sz = 8; 36049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski totalsz += sz/8; 36149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 36249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 36349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned retAlignment = 0; 36449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!llvm::getAlign(*F, 0, retAlignment)) 36549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski retAlignment = TD->getABITypeAlignment(Ty); 36649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".param .align " 36749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski << retAlignment 36849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski << " .b8 func_retval0[" 36949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski << totalsz << "]"; 37049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else 37149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski assert(false && 37249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski "Unknown return type"); 37349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 37449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else { 37549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski SmallVector<EVT, 16> vtparts; 37649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ComputeValueVTs(*TLI, Ty, vtparts); 37749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned idx = 0; 37849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (unsigned i=0,e=vtparts.size(); i!=e; ++i) { 37949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned elems = 1; 38049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski EVT elemtype = vtparts[i]; 38149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (vtparts[i].isVector()) { 38249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski elems = vtparts[i].getVectorNumElements(); 38349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski elemtype = vtparts[i].getVectorElementType(); 38449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 38549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 38649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (unsigned j=0, je=elems; j!=je; ++j) { 38749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned sz = elemtype.getSizeInBits(); 38849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (elemtype.isInteger() && (sz < 32)) sz = 32; 38949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".reg .b" << sz << " func_retval" << idx; 39049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (j<je-1) O << ", "; 39149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ++idx; 39249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 39349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (i < e-1) 39449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ", "; 39549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 39649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 39749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ") "; 39849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 39949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 40049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 40149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::printReturnValStr(const MachineFunction &MF, 40249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_ostream &O) { 40349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const Function *F = MF.getFunction(); 40449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printReturnValStr(F, O); 40549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 40649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 40749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::EmitFunctionEntryLabel() { 40849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski SmallString<128> Str; 40949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_svector_ostream O(Str); 41049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 41149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Set up 41249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski MRI = &MF->getRegInfo(); 41349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski F = MF->getFunction(); 41449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitLinkageDirective(F,O); 41549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (llvm::isKernelFunction(*F)) 41649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".entry "; 41749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else { 41849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".func "; 41949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printReturnValStr(*MF, O); 42049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 42149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 42249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << *CurrentFnSym; 42349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 42449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitFunctionParamList(*MF, O); 42549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 42649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (llvm::isKernelFunction(*F)) 42749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitKernelFunctionDirectives(*F, O); 42849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 42949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski OutStreamer.EmitRawText(O.str()); 43049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 43149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski prevDebugLoc = DebugLoc(); 43249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 43349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 43449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::EmitFunctionBodyStart() { 43549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const TargetRegisterInfo &TRI = *TM.getRegisterInfo(); 43649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned numRegClasses = TRI.getNumRegClasses(); 43749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski VRidGlobal2LocalMap = new std::map<unsigned, unsigned>[numRegClasses+1]; 43849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski OutStreamer.EmitRawText(StringRef("{\n")); 43949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski setAndEmitFunctionVirtualRegisters(*MF); 44049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 44149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski SmallString<128> Str; 44249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_svector_ostream O(Str); 44349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitDemotedVars(MF->getFunction(), O); 44449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski OutStreamer.EmitRawText(O.str()); 44549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 44649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 44749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::EmitFunctionBodyEnd() { 44849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski OutStreamer.EmitRawText(StringRef("}\n")); 44949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski delete []VRidGlobal2LocalMap; 45049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 45149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 45249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 45349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid 45449683f3c961379fbc088871a5d6304950f1f1cbcJustin HolewinskiNVPTXAsmPrinter::emitKernelFunctionDirectives(const Function& F, 45549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_ostream &O) const { 45649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // If the NVVM IR has some of reqntid* specified, then output 45749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // the reqntid directive, and set the unspecified ones to 1. 45849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // If none of reqntid* is specified, don't output reqntid directive. 45949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned reqntidx, reqntidy, reqntidz; 46049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski bool specified = false; 46149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (llvm::getReqNTIDx(F, reqntidx) == false) reqntidx = 1; 46249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else specified = true; 46349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (llvm::getReqNTIDy(F, reqntidy) == false) reqntidy = 1; 46449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else specified = true; 46549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (llvm::getReqNTIDz(F, reqntidz) == false) reqntidz = 1; 46649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else specified = true; 46749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 46849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (specified) 46949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".reqntid " << reqntidx << ", " 47049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski << reqntidy << ", " << reqntidz << "\n"; 47149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 47249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // If the NVVM IR has some of maxntid* specified, then output 47349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // the maxntid directive, and set the unspecified ones to 1. 47449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // If none of maxntid* is specified, don't output maxntid directive. 47549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned maxntidx, maxntidy, maxntidz; 47649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski specified = false; 47749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (llvm::getMaxNTIDx(F, maxntidx) == false) maxntidx = 1; 47849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else specified = true; 47949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (llvm::getMaxNTIDy(F, maxntidy) == false) maxntidy = 1; 48049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else specified = true; 48149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (llvm::getMaxNTIDz(F, maxntidz) == false) maxntidz = 1; 48249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else specified = true; 48349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 48449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (specified) 48549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".maxntid " << maxntidx << ", " 48649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski << maxntidy << ", " << maxntidz << "\n"; 48749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 48849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned mincta; 48949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (llvm::getMinCTASm(F, mincta)) 49049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".minnctapersm " << mincta << "\n"; 49149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 49249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 49349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid 49449683f3c961379fbc088871a5d6304950f1f1cbcJustin HolewinskiNVPTXAsmPrinter::getVirtualRegisterName(unsigned vr, bool isVec, 49549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_ostream &O) { 49649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const TargetRegisterClass * RC = MRI->getRegClass(vr); 49749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned id = RC->getID(); 49849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 49949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski std::map<unsigned, unsigned> ®map = VRidGlobal2LocalMap[id]; 50049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned mapped_vr = regmap[vr]; 50149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 50249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!isVec) { 50349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << getNVPTXRegClassStr(RC) << mapped_vr; 50449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 50549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 5067eacad03efda36e09ebd96e95d7891cadaaa9087Justin Holewinski report_fatal_error("Bad register!"); 50749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 50849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 50949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid 51049683f3c961379fbc088871a5d6304950f1f1cbcJustin HolewinskiNVPTXAsmPrinter::emitVirtualRegister(unsigned int vr, bool isVec, 51149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_ostream &O) { 51249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski getVirtualRegisterName(vr, isVec, O); 51349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 51449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 51549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::printVecModifiedImmediate(const MachineOperand &MO, 51649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const char *Modifier, 51749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_ostream &O) { 5186fcf129cf5fcf4cd0ceb2afa8e2537a4333f896dCraig Topper static const char vecelem[] = {'0', '1', '2', '3', '0', '1', '2', '3'}; 51949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski int Imm = (int)MO.getImm(); 52049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if(0 == strcmp(Modifier, "vecelem")) 52149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "_" << vecelem[Imm]; 52249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else if(0 == strcmp(Modifier, "vecv4comm1")) { 52349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if((Imm < 0) || (Imm > 3)) 52449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "//"; 52549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 52649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else if(0 == strcmp(Modifier, "vecv4comm2")) { 52749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if((Imm < 4) || (Imm > 7)) 52849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "//"; 52949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 53049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else if(0 == strcmp(Modifier, "vecv4pos")) { 53149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if(Imm < 0) Imm = 0; 53249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "_" << vecelem[Imm%4]; 53349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 53449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else if(0 == strcmp(Modifier, "vecv2comm1")) { 53549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if((Imm < 0) || (Imm > 1)) 53649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "//"; 53749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 53849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else if(0 == strcmp(Modifier, "vecv2comm2")) { 53949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if((Imm < 2) || (Imm > 3)) 54049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "//"; 54149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 54249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else if(0 == strcmp(Modifier, "vecv2pos")) { 54349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if(Imm < 0) Imm = 0; 54449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "_" << vecelem[Imm%2]; 54549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 54649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 5476366361998599a63a74ee7cc5f4ba900711c7e7aCraig Topper llvm_unreachable("Unknown Modifier on immediate operand"); 54849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 54949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 55049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::printOperand(const MachineInstr *MI, int opNum, 55149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_ostream &O, const char *Modifier) { 55249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const MachineOperand &MO = MI->getOperand(opNum); 55349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski switch (MO.getType()) { 55449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case MachineOperand::MO_Register: 55549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (TargetRegisterInfo::isPhysicalRegister(MO.getReg())) { 55649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (MO.getReg() == NVPTX::VRDepot) 55749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << DEPOTNAME << getFunctionNumber(); 55849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 55949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << getRegisterName(MO.getReg()); 56049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else { 56149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!Modifier) 56249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitVirtualRegister(MO.getReg(), false, O); 56349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else { 56449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (strcmp(Modifier, "vecfull") == 0) 56549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitVirtualRegister(MO.getReg(), true, O); 56649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 5676366361998599a63a74ee7cc5f4ba900711c7e7aCraig Topper llvm_unreachable( 56849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski "Don't know how to handle the modifier on virtual register."); 56949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 57049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 57149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 57249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 57349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case MachineOperand::MO_Immediate: 57449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!Modifier) 57549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << MO.getImm(); 57649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else if (strstr(Modifier, "vec") == Modifier) 57749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printVecModifiedImmediate(MO, Modifier, O); 57849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 5796366361998599a63a74ee7cc5f4ba900711c7e7aCraig Topper llvm_unreachable("Don't know how to handle modifier on immediate operand"); 58049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 58149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 58249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case MachineOperand::MO_FPImmediate: 58349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printFPConstant(MO.getFPImm(), O); 58449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 58549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 58649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case MachineOperand::MO_GlobalAddress: 58749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << *Mang->getSymbol(MO.getGlobal()); 58849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 58949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 59049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case MachineOperand::MO_ExternalSymbol: { 59149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const char * symbname = MO.getSymbolName(); 59249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (strstr(symbname, ".PARAM") == symbname) { 59349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned index; 59449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski sscanf(symbname+6, "%u[];", &index); 59549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printParamName(index, O); 59649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 59749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else if (strstr(symbname, ".HLPPARAM") == symbname) { 59849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned index; 59949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski sscanf(symbname+9, "%u[];", &index); 60049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << *CurrentFnSym << "_param_" << index << "_offset"; 60149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 60249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 60349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << symbname; 60449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 60549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 60649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 60749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case MachineOperand::MO_MachineBasicBlock: 60849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << *MO.getMBB()->getSymbol(); 60949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 61049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 61149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski default: 6126366361998599a63a74ee7cc5f4ba900711c7e7aCraig Topper llvm_unreachable("Operand type not supported."); 61349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 61449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 61549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 61649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter:: 61749683f3c961379fbc088871a5d6304950f1f1cbcJustin HolewinskiprintImplicitDef(const MachineInstr *MI, raw_ostream &O) const { 61849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski#ifndef __OPTIMIZE__ 61949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t// Implicit def :"; 62049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski //printOperand(MI, 0); 62149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\n"; 62249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski#endif 62349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 62449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 62549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::printMemOperand(const MachineInstr *MI, int opNum, 62649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_ostream &O, const char *Modifier) { 62749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printOperand(MI, opNum, O); 62849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 62949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (Modifier && !strcmp(Modifier, "add")) { 63049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ", "; 63149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printOperand(MI, opNum+1, O); 63249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else { 63349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (MI->getOperand(opNum+1).isImm() && 63449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski MI->getOperand(opNum+1).getImm() == 0) 63549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; // don't print ',0' or '+0' 63649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "+"; 63749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printOperand(MI, opNum+1, O); 63849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 63949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 64049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 64149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::printLdStCode(const MachineInstr *MI, int opNum, 64249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_ostream &O, const char *Modifier) 64349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski{ 64449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (Modifier) { 64549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const MachineOperand &MO = MI->getOperand(opNum); 64649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski int Imm = (int)MO.getImm(); 64749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!strcmp(Modifier, "volatile")) { 64849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (Imm) 64949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".volatile"; 65049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else if (!strcmp(Modifier, "addsp")) { 65149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski switch (Imm) { 65249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::PTXLdStInstCode::GLOBAL: O << ".global"; break; 65349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::PTXLdStInstCode::SHARED: O << ".shared"; break; 65449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::PTXLdStInstCode::LOCAL: O << ".local"; break; 65549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::PTXLdStInstCode::PARAM: O << ".param"; break; 65649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::PTXLdStInstCode::CONSTANT: O << ".const"; break; 65749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::PTXLdStInstCode::GENERIC: 65849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!nvptxSubtarget.hasGenericLdSt()) 65949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".global"; 66049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 66149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski default: 6627454fc2e87ef5638f3644b86a4350a44513e5185Jakub Staszak llvm_unreachable("Wrong Address Space"); 66349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 66449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 66549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else if (!strcmp(Modifier, "sign")) { 66649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (Imm==NVPTX::PTXLdStInstCode::Signed) 66749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "s"; 66849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else if (Imm==NVPTX::PTXLdStInstCode::Unsigned) 66949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "u"; 67049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 67149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "f"; 67249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 67349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else if (!strcmp(Modifier, "vec")) { 67449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (Imm==NVPTX::PTXLdStInstCode::V2) 67549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".v2"; 67649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else if (Imm==NVPTX::PTXLdStInstCode::V4) 67749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".v4"; 67849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 67949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 6807454fc2e87ef5638f3644b86a4350a44513e5185Jakub Staszak llvm_unreachable("Unknown Modifier"); 68149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 68249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 6837454fc2e87ef5638f3644b86a4350a44513e5185Jakub Staszak llvm_unreachable("Empty Modifier"); 68449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 68549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 68649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::emitDeclaration (const Function *F, raw_ostream &O) { 68749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 68849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitLinkageDirective(F,O); 68949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (llvm::isKernelFunction(*F)) 69049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".entry "; 69149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 69249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".func "; 69349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printReturnValStr(F, O); 69449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << *CurrentFnSym << "\n"; 69549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitFunctionParamList(F, O); 69649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ";\n"; 69749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 69849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 69949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskistatic bool usedInGlobalVarDef(const Constant *C) 70049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski{ 70149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!C) 70249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return false; 70349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 70449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (const GlobalVariable *GV = dyn_cast<GlobalVariable>(C)) { 70549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (GV->getName().str() == "llvm.used") 70649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return false; 70749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return true; 70849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 70949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 71049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (Value::const_use_iterator ui=C->use_begin(), ue=C->use_end(); 71149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ui!=ue; ++ui) { 71249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const Constant *C = dyn_cast<Constant>(*ui); 71349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (usedInGlobalVarDef(C)) 71449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return true; 71549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 71649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return false; 71749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 71849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 71949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskistatic bool usedInOneFunc(const User *U, Function const *&oneFunc) 72049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski{ 72149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (const GlobalVariable *othergv = dyn_cast<GlobalVariable>(U)) { 72249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (othergv->getName().str() == "llvm.used") 72349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return true; 72449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 72549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 72649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (const Instruction *instr = dyn_cast<Instruction>(U)) { 72749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (instr->getParent() && instr->getParent()->getParent()) { 72849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const Function *curFunc = instr->getParent()->getParent(); 72949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (oneFunc && (curFunc != oneFunc)) 73049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return false; 73149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski oneFunc = curFunc; 73249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return true; 73349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 73449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 73549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return false; 73649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 73749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 73849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (const MDNode *md = dyn_cast<MDNode>(U)) 73949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (md->hasName() && ((md->getName().str() == "llvm.dbg.gv") || 74049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski (md->getName().str() == "llvm.dbg.sp"))) 74149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return true; 74249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 74349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 74449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (User::const_use_iterator ui=U->use_begin(), ue=U->use_end(); 74549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ui!=ue; ++ui) { 74649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (usedInOneFunc(*ui, oneFunc) == false) 74749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return false; 74849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 74949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return true; 75049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 75149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 75249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski/* Find out if a global variable can be demoted to local scope. 75349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski * Currently, this is valid for CUDA shared variables, which have local 75449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski * scope and global lifetime. So the conditions to check are : 75549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski * 1. Is the global variable in shared address space? 75649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski * 2. Does it have internal linkage? 75749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski * 3. Is the global variable referenced only in one function? 75849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski */ 75949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskistatic bool canDemoteGlobalVar(const GlobalVariable *gv, Function const *&f) { 76049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (gv->hasInternalLinkage() == false) 76149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return false; 76249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const PointerType *Pty = gv->getType(); 76349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (Pty->getAddressSpace() != llvm::ADDRESS_SPACE_SHARED) 76449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return false; 76549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 76649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const Function *oneFunc = 0; 76749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 76849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski bool flag = usedInOneFunc(gv, oneFunc); 76949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (flag == false) 77049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return false; 77149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!oneFunc) 77249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return false; 77349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski f = oneFunc; 77449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return true; 77549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 77649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 77749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskistatic bool useFuncSeen(const Constant *C, 77849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski llvm::DenseMap<const Function *, bool> &seenMap) { 77949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (Value::const_use_iterator ui=C->use_begin(), ue=C->use_end(); 78049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ui!=ue; ++ui) { 78149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (const Constant *cu = dyn_cast<Constant>(*ui)) { 78249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (useFuncSeen(cu, seenMap)) 78349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return true; 78449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else if (const Instruction *I = dyn_cast<Instruction>(*ui)) { 78549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const BasicBlock *bb = I->getParent(); 78649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!bb) continue; 78749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const Function *caller = bb->getParent(); 78849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!caller) continue; 78949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (seenMap.find(caller) != seenMap.end()) 79049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return true; 79149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 79249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 79349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return false; 79449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 79549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 79649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::emitDeclarations (Module &M, raw_ostream &O) { 79749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski llvm::DenseMap<const Function *, bool> seenMap; 79849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (Module::const_iterator FI=M.begin(), FE=M.end(); 79949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski FI!=FE; ++FI) { 80049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const Function *F = FI; 80149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 80249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (F->isDeclaration()) { 80349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (F->use_empty()) 80449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski continue; 80549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (F->getIntrinsicID()) 80649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski continue; 80749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski CurrentFnSym = Mang->getSymbol(F); 80849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitDeclaration(F, O); 80949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski continue; 81049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 81149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (Value::const_use_iterator iter=F->use_begin(), 81249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski iterEnd=F->use_end(); iter!=iterEnd; ++iter) { 81349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (const Constant *C = dyn_cast<Constant>(*iter)) { 81449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (usedInGlobalVarDef(C)) { 81549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // The use is in the initialization of a global variable 81649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // that is a function pointer, so print a declaration 81749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // for the original function 81849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski CurrentFnSym = Mang->getSymbol(F); 81949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitDeclaration(F, O); 82049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 82149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 82249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Emit a declaration of this function if the function that 82349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // uses this constant expr has already been seen. 82449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (useFuncSeen(C, seenMap)) { 82549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski CurrentFnSym = Mang->getSymbol(F); 82649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitDeclaration(F, O); 82749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 82849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 82949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 83049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 83149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!isa<Instruction>(*iter)) continue; 83249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const Instruction *instr = cast<Instruction>(*iter); 83349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const BasicBlock *bb = instr->getParent(); 83449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!bb) continue; 83549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const Function *caller = bb->getParent(); 83649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!caller) continue; 83749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 83849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // If a caller has already been seen, then the caller is 83949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // appearing in the module before the callee. so print out 84049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // a declaration for the callee. 84149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (seenMap.find(caller) != seenMap.end()) { 84249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski CurrentFnSym = Mang->getSymbol(F); 84349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitDeclaration(F, O); 84449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 84549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 84649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 84749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski seenMap[F] = true; 84849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 84949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 85049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 85149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::recordAndEmitFilenames(Module &M) { 85249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski DebugInfoFinder DbgFinder; 85349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski DbgFinder.processModule(M); 85449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 85549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned i=1; 85649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (DebugInfoFinder::iterator I = DbgFinder.compile_unit_begin(), 85749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski E = DbgFinder.compile_unit_end(); I != E; ++I) { 85849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski DICompileUnit DIUnit(*I); 85949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski StringRef Filename(DIUnit.getFilename()); 86049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski StringRef Dirname(DIUnit.getDirectory()); 86149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski SmallString<128> FullPathName = Dirname; 86249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!Dirname.empty() && !sys::path::is_absolute(Filename)) { 86349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski sys::path::append(FullPathName, Filename); 86449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Filename = FullPathName.str(); 86549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 86649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (filenameMap.find(Filename.str()) != filenameMap.end()) 86749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski continue; 86849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski filenameMap[Filename.str()] = i; 86949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski OutStreamer.EmitDwarfFileDirective(i, "", Filename.str()); 87049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ++i; 87149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 87249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 87349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (DebugInfoFinder::iterator I = DbgFinder.subprogram_begin(), 87449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski E = DbgFinder.subprogram_end(); I != E; ++I) { 87549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski DISubprogram SP(*I); 87649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski StringRef Filename(SP.getFilename()); 87749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski StringRef Dirname(SP.getDirectory()); 87849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski SmallString<128> FullPathName = Dirname; 87949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!Dirname.empty() && !sys::path::is_absolute(Filename)) { 88049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski sys::path::append(FullPathName, Filename); 88149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Filename = FullPathName.str(); 88249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 88349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (filenameMap.find(Filename.str()) != filenameMap.end()) 88449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski continue; 88549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski filenameMap[Filename.str()] = i; 88649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ++i; 88749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 88849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 88949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 89049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskibool NVPTXAsmPrinter::doInitialization (Module &M) { 89149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 89249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski SmallString<128> Str1; 89349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_svector_ostream OS1(Str1); 89449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 89549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski MMI = getAnalysisIfAvailable<MachineModuleInfo>(); 89649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski MMI->AnalyzeModule(M); 89749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 89849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // We need to call the parent's one explicitly. 89949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski //bool Result = AsmPrinter::doInitialization(M); 90049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 90149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Initialize TargetLoweringObjectFile. 90249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const_cast<TargetLoweringObjectFile&>(getObjFileLowering()) 90349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski .Initialize(OutContext, TM); 90449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 9053574eca1b02600bac4e625297f4ecf745f4c4f32Micah Villmow Mang = new Mangler(OutContext, *TM.getDataLayout()); 90649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 90749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Emit header before any dwarf directives are emitted below. 90849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitHeader(M, OS1); 90949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski OutStreamer.EmitRawText(OS1.str()); 91049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 91149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 91249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Already commented out 91349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski //bool Result = AsmPrinter::doInitialization(M); 91449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 91549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 91649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA) 91749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski recordAndEmitFilenames(M); 91849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 91949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski SmallString<128> Str2; 92049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_svector_ostream OS2(Str2); 92149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 92249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitDeclarations(M, OS2); 92349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 9242085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski // As ptxas does not support forward references of globals, we need to first 9252085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski // sort the list of module-level globals in def-use order. We visit each 9262085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski // global variable in order, and ensure that we emit it *after* its dependent 9272085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski // globals. We use a little extra memory maintaining both a set and a list to 9282085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski // have fast searches while maintaining a strict ordering. 9292085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski SmallVector<GlobalVariable*,8> Globals; 9302085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski DenseSet<GlobalVariable*> GVVisited; 9312085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski DenseSet<GlobalVariable*> GVVisiting; 9322085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski 9332085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski // Visit each global variable, in order 93449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (Module::global_iterator I = M.global_begin(), E = M.global_end(); 9352085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski I != E; ++I) 9362085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski VisitGlobalVariableForEmission(I, Globals, GVVisited, GVVisiting); 9372085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski 9382085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski assert(GVVisited.size() == M.getGlobalList().size() && 9392085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski "Missed a global variable"); 9402085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski assert(GVVisiting.size() == 0 && "Did not fully process a global variable"); 9412085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski 9422085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski // Print out module-level global variables in proper order 9432085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski for (unsigned i = 0, e = Globals.size(); i != e; ++i) 9442085d00d09f4f3678a6c67da46df3a04e31b499cJustin Holewinski printModuleLevelGV(Globals[i], OS2); 94549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 94649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski OS2 << '\n'; 94749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 94849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski OutStreamer.EmitRawText(OS2.str()); 94949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return false; // success 95049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 95149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 95249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::emitHeader (Module &M, raw_ostream &O) { 95349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "//\n"; 95449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "// Generated by LLVM NVPTX Back-End\n"; 95549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "//\n"; 95649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\n"; 95749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 95808e9cb46feb0c8e08e3d309a0f9fd75a04ca54fbJustin Holewinski unsigned PTXVersion = nvptxSubtarget.getPTXVersion(); 95908e9cb46feb0c8e08e3d309a0f9fd75a04ca54fbJustin Holewinski O << ".version " << (PTXVersion / 10) << "." << (PTXVersion % 10) << "\n"; 96049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 96149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".target "; 96249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << nvptxSubtarget.getTargetName(); 96349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 96449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (nvptxSubtarget.getDrvInterface() == NVPTX::NVCL) 96549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ", texmode_independent"; 96649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA) { 96749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!nvptxSubtarget.hasDouble()) 96849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ", map_f64_to_f32"; 96949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 97049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 97149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (MAI->doesSupportDebugInformation()) 97249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ", debug"; 97349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 97449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\n"; 97549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 97649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".address_size "; 97749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (nvptxSubtarget.is64Bit()) 97849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "64"; 97949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 98049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "32"; 98149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\n"; 98249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 98349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\n"; 98449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 98549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 98649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskibool NVPTXAsmPrinter::doFinalization(Module &M) { 98749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // XXX Temproarily remove global variables so that doFinalization() will not 98849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // emit them again (global variables are emitted at beginning). 98949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 99049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Module::GlobalListType &global_list = M.getGlobalList(); 99149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski int i, n = global_list.size(); 99249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski GlobalVariable **gv_array = new GlobalVariable* [n]; 99349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 99449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // first, back-up GlobalVariable in gv_array 99549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski i = 0; 99649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (Module::global_iterator I = global_list.begin(), E = global_list.end(); 99749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski I != E; ++I) 99849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski gv_array[i++] = &*I; 99949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 100049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // second, empty global_list 100149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski while (!global_list.empty()) 100249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski global_list.remove(global_list.begin()); 100349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 100449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // call doFinalization 100549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski bool ret = AsmPrinter::doFinalization(M); 100649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 100749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // now we restore global variables 100849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (i = 0; i < n; i ++) 100949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski global_list.insert(global_list.end(), gv_array[i]); 101049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 101149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski delete[] gv_array; 101249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return ret; 101349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 101449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 101549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski //bool Result = AsmPrinter::doFinalization(M); 101649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Instead of calling the parents doFinalization, we may 101749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // clone parents doFinalization and customize here. 101849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Currently, we if NVISA out the EmitGlobals() in 101949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // parent's doFinalization, which is too intrusive. 102049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // 102149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Same for the doInitialization. 102249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski //return Result; 102349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 102449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 102549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// This function emits appropriate linkage directives for 102649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// functions and global variables. 102749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// 102849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// extern function declaration -> .extern 102949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// extern function definition -> .visible 103049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// external global variable with init -> .visible 103149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// external without init -> .extern 103249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// appending -> not allowed, assert. 103349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 103449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::emitLinkageDirective(const GlobalValue* V, raw_ostream &O) 103549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski{ 103649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA) { 103749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (V->hasExternalLinkage()) { 103849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (isa<GlobalVariable>(V)) { 103949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const GlobalVariable *GVar = cast<GlobalVariable>(V); 104049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (GVar) { 104149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (GVar->hasInitializer()) 104249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".visible "; 104349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 104449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".extern "; 104549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 104649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else if (V->isDeclaration()) 104749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".extern "; 104849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 104949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".visible "; 105049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else if (V->hasAppendingLinkage()) { 105149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski std::string msg; 105249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski msg.append("Error: "); 105349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski msg.append("Symbol "); 105449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (V->hasName()) 105549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski msg.append(V->getName().str()); 105649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski msg.append("has unsupported appending linkage type"); 105749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski llvm_unreachable(msg.c_str()); 105849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 105949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 106049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 106149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 106249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 106349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::printModuleLevelGV(GlobalVariable* GVar, raw_ostream &O, 106449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski bool processDemoted) { 106549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 106649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Skip meta data 106749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (GVar->hasSection()) { 106849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (GVar->getSection() == "llvm.metadata") 106949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 107049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 107149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 10723574eca1b02600bac4e625297f4ecf745f4c4f32Micah Villmow const DataLayout *TD = TM.getDataLayout(); 107349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 107449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // GlobalVariables are always constant pointers themselves. 107549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const PointerType *PTy = GVar->getType(); 107649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Type *ETy = PTy->getElementType(); 107749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 107849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (GVar->hasExternalLinkage()) { 107949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (GVar->hasInitializer()) 108049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".visible "; 108149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 108249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".extern "; 108349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 108449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 108549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (llvm::isTexture(*GVar)) { 108649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".global .texref " << llvm::getTextureName(*GVar) << ";\n"; 108749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 108849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 108949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 109049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (llvm::isSurface(*GVar)) { 109149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".global .surfref " << llvm::getSurfaceName(*GVar) << ";\n"; 109249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 109349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 109449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 109549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (GVar->isDeclaration()) { 109649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // (extern) declarations, no definition or initializer 109749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Currently the only known declaration is for an automatic __local 109849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // (.shared) promoted to global. 109949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitPTXGlobalVariable(GVar, O); 110049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ";\n"; 110149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 110249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 110349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 110449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (llvm::isSampler(*GVar)) { 110549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".global .samplerref " << llvm::getSamplerName(*GVar); 110649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 110749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Constant *Initializer = NULL; 110849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (GVar->hasInitializer()) 110949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Initializer = GVar->getInitializer(); 111049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ConstantInt *CI = NULL; 111149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (Initializer) 111249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski CI = dyn_cast<ConstantInt>(Initializer); 111349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (CI) { 111449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned sample=CI->getZExtValue(); 111549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 111649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " = { "; 111749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 111849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (int i =0, addr=((sample & __CLK_ADDRESS_MASK ) >> 111949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski __CLK_ADDRESS_BASE) ; i < 3 ; i++) { 112049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "addr_mode_" << i << " = "; 112149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski switch (addr) { 112249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case 0: O << "wrap"; break; 112349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case 1: O << "clamp_to_border"; break; 112449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case 2: O << "clamp_to_edge"; break; 112549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case 3: O << "wrap"; break; 112649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case 4: O << "mirror"; break; 112749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 112849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O <<", "; 112949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 113049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "filter_mode = "; 113149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski switch (( sample & __CLK_FILTER_MASK ) >> __CLK_FILTER_BASE ) { 113249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case 0: O << "nearest"; break; 113349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case 1: O << "linear"; break; 113449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case 2: assert ( 0 && "Anisotropic filtering is not supported"); 113549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski default: O << "nearest"; break; 113649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 113749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!(( sample &__CLK_NORMALIZED_MASK ) >> __CLK_NORMALIZED_BASE)) { 113849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ", force_unnormalized_coords = 1"; 113949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 114049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " }"; 114149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 114249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 114349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ";\n"; 114449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 114549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 114649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 114749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (GVar->hasPrivateLinkage()) { 114849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 114949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!strncmp(GVar->getName().data(), "unrollpragma", 12)) 115049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 115149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 115249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // FIXME - need better way (e.g. Metadata) to avoid generating this global 115349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!strncmp(GVar->getName().data(), "filename", 8)) 115449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 115549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (GVar->use_empty()) 115649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 115749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 115849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 115949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const Function *demotedFunc = 0; 116049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!processDemoted && canDemoteGlobalVar(GVar, demotedFunc)) { 116149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "// " << GVar->getName().str() << " has been demoted\n"; 116249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (localDecls.find(demotedFunc) != localDecls.end()) 116349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski localDecls[demotedFunc].push_back(GVar); 116449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else { 116549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski std::vector<GlobalVariable *> temp; 116649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski temp.push_back(GVar); 116749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski localDecls[demotedFunc] = temp; 116849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 116949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 117049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 117149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 117249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "."; 117349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitPTXAddressSpace(PTy->getAddressSpace(), O); 117449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (GVar->getAlignment() == 0) 117549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " .align " << (int) TD->getPrefTypeAlignment(ETy); 117649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 117749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " .align " << GVar->getAlignment(); 117849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 117949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 118049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (ETy->isPrimitiveType() || ETy->isIntegerTy() || isa<PointerType>(ETy)) { 118149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " ."; 118249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << getPTXFundamentalTypeStr(ETy, false); 118349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " "; 118449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << *Mang->getSymbol(GVar); 118549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 118649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Ptx allows variable initilization only for constant and global state 118749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // spaces. 118849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (((PTy->getAddressSpace() == llvm::ADDRESS_SPACE_GLOBAL) || 118949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski (PTy->getAddressSpace() == llvm::ADDRESS_SPACE_CONST_NOT_GEN) || 119049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski (PTy->getAddressSpace() == llvm::ADDRESS_SPACE_CONST)) 119149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski && GVar->hasInitializer()) { 119249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Constant *Initializer = GVar->getInitializer(); 119349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!Initializer->isNullValue()) { 119449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " = " ; 119549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printScalarConstant(Initializer, O); 119649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 119749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 119849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else { 119949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned int ElementSize =0; 120049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 120149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Although PTX has direct support for struct type and array type and 120249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // LLVM IR is very similar to PTX, the LLVM CodeGen does not support for 120349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // targets that support these high level field accesses. Structs, arrays 120449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // and vectors are lowered into arrays of bytes. 120549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski switch (ETy->getTypeID()) { 120649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Type::StructTyID: 120749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Type::ArrayTyID: 120849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Type::VectorTyID: 120949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ElementSize = TD->getTypeStoreSize(ETy); 121049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Ptx allows variable initilization only for constant and 121149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // global state spaces. 121249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (((PTy->getAddressSpace() == llvm::ADDRESS_SPACE_GLOBAL) || 121349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski (PTy->getAddressSpace() == llvm::ADDRESS_SPACE_CONST_NOT_GEN) || 121449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski (PTy->getAddressSpace() == llvm::ADDRESS_SPACE_CONST)) 121549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski && GVar->hasInitializer()) { 121649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Constant *Initializer = GVar->getInitializer(); 121749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!isa<UndefValue>(Initializer) && 121849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski !Initializer->isNullValue()) { 121949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski AggBuffer aggBuffer(ElementSize, O, *this); 122049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski bufferAggregateConstant(Initializer, &aggBuffer); 122149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (aggBuffer.numSymbols) { 122249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (nvptxSubtarget.is64Bit()) { 122349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " .u64 " << *Mang->getSymbol(GVar) <<"[" ; 122449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ElementSize/8; 122549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 122649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else { 122749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " .u32 " << *Mang->getSymbol(GVar) <<"[" ; 122849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ElementSize/4; 122949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 123049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "]"; 123149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 123249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else { 123349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " .b8 " << *Mang->getSymbol(GVar) <<"[" ; 123449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ElementSize; 123549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "]"; 123649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 123749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " = {" ; 123849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer.print(); 123949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "}"; 124049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 124149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else { 124249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " .b8 " << *Mang->getSymbol(GVar) ; 124349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (ElementSize) { 124449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O <<"[" ; 124549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ElementSize; 124649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "]"; 124749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 124849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 124949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 125049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else { 125149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " .b8 " << *Mang->getSymbol(GVar); 125249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (ElementSize) { 125349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O <<"[" ; 125449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ElementSize; 125549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "]"; 125649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 125749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 125849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 125949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski default: 126049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski assert( 0 && "type not supported yet"); 126149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 126249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 126349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 126449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ";\n"; 126549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 126649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 126749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::emitDemotedVars(const Function *f, raw_ostream &O) { 126849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (localDecls.find(f) == localDecls.end()) 126949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 127049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 127149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski std::vector<GlobalVariable *> &gvars = localDecls[f]; 127249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 127349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (unsigned i=0, e=gvars.size(); i!=e; ++i) { 127449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t// demoted variable\n\t"; 127549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printModuleLevelGV(gvars[i], O, true); 127649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 127749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 127849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 127949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::emitPTXAddressSpace(unsigned int AddressSpace, 128049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_ostream &O) const { 128149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski switch (AddressSpace) { 128249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case llvm::ADDRESS_SPACE_LOCAL: 128349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "local" ; 128449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 128549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case llvm::ADDRESS_SPACE_GLOBAL: 128649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "global" ; 128749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 128849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case llvm::ADDRESS_SPACE_CONST: 128949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // This logic should be consistent with that in 129049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // getCodeAddrSpace() (NVPTXISelDATToDAT.cpp) 129149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (nvptxSubtarget.hasGenericLdSt()) 129249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "global" ; 129349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 129449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "const" ; 129549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 129649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case llvm::ADDRESS_SPACE_CONST_NOT_GEN: 129749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "const" ; 129849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 129949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case llvm::ADDRESS_SPACE_SHARED: 130049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "shared" ; 130149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 130249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski default: 130300d9da1ac45429cb69c8b298b9f25e10a4b57813Justin Holewinski report_fatal_error("Bad address space found while emitting PTX"); 130400d9da1ac45429cb69c8b298b9f25e10a4b57813Justin Holewinski break; 130549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 130649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 130749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 130849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskistd::string NVPTXAsmPrinter::getPTXFundamentalTypeStr(const Type *Ty, 130949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski bool useB4PTR) const { 131049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski switch (Ty->getTypeID()) { 131149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski default: 131249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski llvm_unreachable("unexpected type"); 131349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 131449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Type::IntegerTyID: { 131549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned NumBits = cast<IntegerType>(Ty)->getBitWidth(); 131649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (NumBits == 1) 131749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return "pred"; 131849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else if (NumBits <= 64) { 131949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski std::string name = "u"; 132049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return name + utostr(NumBits); 132149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else { 132249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski llvm_unreachable("Integer too large"); 132349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 132449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 132549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 132649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 132749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Type::FloatTyID: 132849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return "f32"; 132949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Type::DoubleTyID: 133049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return "f64"; 133149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Type::PointerTyID: 133249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (nvptxSubtarget.is64Bit()) 133349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (useB4PTR) return "b64"; 133449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else return "u64"; 133549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 133649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (useB4PTR) return "b32"; 133749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else return "u32"; 133849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 133949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski llvm_unreachable("unexpected type"); 134049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return NULL; 134149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 134249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 134349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable* GVar, 134449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_ostream &O) { 134549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 13463574eca1b02600bac4e625297f4ecf745f4c4f32Micah Villmow const DataLayout *TD = TM.getDataLayout(); 134749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 134849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // GlobalVariables are always constant pointers themselves. 134949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const PointerType *PTy = GVar->getType(); 135049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Type *ETy = PTy->getElementType(); 135149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 135249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "."; 135349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitPTXAddressSpace(PTy->getAddressSpace(), O); 135449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (GVar->getAlignment() == 0) 135549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " .align " << (int) TD->getPrefTypeAlignment(ETy); 135649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 135749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " .align " << GVar->getAlignment(); 135849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 135949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (ETy->isPrimitiveType() || ETy->isIntegerTy() || isa<PointerType>(ETy)) { 136049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " ."; 136149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << getPTXFundamentalTypeStr(ETy); 136249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " "; 136349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << *Mang->getSymbol(GVar); 136449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 136549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 136649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 136749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski int64_t ElementSize =0; 136849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 136949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Although PTX has direct support for struct type and array type and LLVM IR 137049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // is very similar to PTX, the LLVM CodeGen does not support for targets that 137149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // support these high level field accesses. Structs and arrays are lowered 137249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // into arrays of bytes. 137349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski switch (ETy->getTypeID()) { 137449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Type::StructTyID: 137549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Type::ArrayTyID: 137649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Type::VectorTyID: 137749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ElementSize = TD->getTypeStoreSize(ETy); 137849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << " .b8 " << *Mang->getSymbol(GVar) <<"[" ; 137949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (ElementSize) { 138049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << itostr(ElementSize) ; 138149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 138249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "]"; 138349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 138449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski default: 138549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski assert( 0 && "type not supported yet"); 138649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 138749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return ; 138849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 138949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 139049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 139149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskistatic unsigned int 13923574eca1b02600bac4e625297f4ecf745f4c4f32Micah VillmowgetOpenCLAlignment(const DataLayout *TD, 139349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Type *Ty) { 139449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (Ty->isPrimitiveType() || Ty->isIntegerTy() || isa<PointerType>(Ty)) 139549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return TD->getPrefTypeAlignment(Ty); 139649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 139749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const ArrayType *ATy = dyn_cast<ArrayType>(Ty); 139849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (ATy) 139949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return getOpenCLAlignment(TD, ATy->getElementType()); 140049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 140149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const VectorType *VTy = dyn_cast<VectorType>(Ty); 140249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (VTy) { 140349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Type *ETy = VTy->getElementType(); 140449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned int numE = VTy->getNumElements(); 140549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned int alignE = TD->getPrefTypeAlignment(ETy); 140649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (numE == 3) 140749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return 4*alignE; 140849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 140949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return numE*alignE; 141049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 141149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 141249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const StructType *STy = dyn_cast<StructType>(Ty); 141349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (STy) { 141449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned int alignStruct = 1; 141549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Go through each element of the struct and find the 141649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // largest alignment. 141749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (unsigned i=0, e=STy->getNumElements(); i != e; i++) { 141849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Type *ETy = STy->getElementType(i); 141949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned int align = getOpenCLAlignment(TD, ETy); 142049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (align > alignStruct) 142149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski alignStruct = align; 142249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 142349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return alignStruct; 142449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 142549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 142649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const FunctionType *FTy = dyn_cast<FunctionType>(Ty); 142749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (FTy) 1428426c2bf5cdd2173e4a33aea8cb92cf684a724f4bChandler Carruth return TD->getPointerPrefAlignment(); 142949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return TD->getPrefTypeAlignment(Ty); 143049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 143149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 143249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::printParamName(Function::const_arg_iterator I, 143349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski int paramIndex, raw_ostream &O) { 143449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if ((nvptxSubtarget.getDrvInterface() == NVPTX::NVCL) || 143549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA)) 143649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << *CurrentFnSym << "_param_" << paramIndex; 143749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else { 143849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski std::string argName = I->getName(); 143949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const char *p = argName.c_str(); 144049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski while (*p) { 144149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (*p == '.') 144249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "_"; 144349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 144449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << *p; 144549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski p++; 144649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 144749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 144849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 144949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 145049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::printParamName(int paramIndex, raw_ostream &O) { 145149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Function::const_arg_iterator I, E; 145249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski int i = 0; 145349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 145449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if ((nvptxSubtarget.getDrvInterface() == NVPTX::NVCL) || 145549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski (nvptxSubtarget.getDrvInterface() == NVPTX::CUDA)) { 145649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << *CurrentFnSym << "_param_" << paramIndex; 145749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 145849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 145949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 146049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (I = F->arg_begin(), E = F->arg_end(); I != E; ++I, i++) { 146149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (i==paramIndex) { 146249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printParamName(I, paramIndex, O); 146349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 146449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 146549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 146649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski llvm_unreachable("paramIndex out of bound"); 146749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 146849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 146949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::emitFunctionParamList(const Function *F, 147049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_ostream &O) { 14713574eca1b02600bac4e625297f4ecf745f4c4f32Micah Villmow const DataLayout *TD = TM.getDataLayout(); 147299faa3b4ec6d03ac7808fe4ff3fbf3d04e375502Bill Wendling const AttributeSet &PAL = F->getAttributes(); 147349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const TargetLowering *TLI = TM.getTargetLowering(); 147449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Function::const_arg_iterator I, E; 147549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned paramIndex = 0; 147649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski bool first = true; 147749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski bool isKernelFunc = llvm::isKernelFunction(*F); 147849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski bool isABI = (nvptxSubtarget.getSmVersion() >= 20); 147949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski MVT thePointerTy = TLI->getPointerTy(); 148049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 148149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "(\n"; 148249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 148349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (I = F->arg_begin(), E = F->arg_end(); I != E; ++I, paramIndex++) { 148449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const Type *Ty = I->getType(); 148549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 148649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (!first) 148749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ",\n"; 148849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 148949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski first = false; 149049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 149149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Handle image/sampler parameters 149249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (llvm::isSampler(*I) || llvm::isImage(*I)) { 149349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (llvm::isImage(*I)) { 149449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski std::string sname = I->getName(); 149549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (llvm::isImageWriteOnly(*I)) 149649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.param .surfref " << *CurrentFnSym << "_param_" << paramIndex; 149749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else // Default image is read_only 149849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.param .texref " << *CurrentFnSym << "_param_" << paramIndex; 149949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 150049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else // Should be llvm::isSampler(*I) 150149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.param .samplerref " << *CurrentFnSym << "_param_" 150249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski << paramIndex; 150349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski continue; 150449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 150549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 150694e94b350652d3a71993bbc7d44afbe3b304605eBill Wendling if (PAL.hasAttribute(paramIndex+1, Attribute::ByVal) == false) { 150749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Just a scalar 150849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const PointerType *PTy = dyn_cast<PointerType>(Ty); 150949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (isKernelFunc) { 151049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (PTy) { 151149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Special handling for pointer arguments to kernel 151249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.param .u" << thePointerTy.getSizeInBits() << " "; 151349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 151449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (nvptxSubtarget.getDrvInterface() != NVPTX::CUDA) { 151549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Type *ETy = PTy->getElementType(); 151649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski int addrSpace = PTy->getAddressSpace(); 151749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski switch(addrSpace) { 151849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski default: 151949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".ptr "; 152049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 152149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case llvm::ADDRESS_SPACE_CONST_NOT_GEN: 152249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".ptr .const "; 152349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 152449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case llvm::ADDRESS_SPACE_SHARED: 152549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".ptr .shared "; 152649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 152749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case llvm::ADDRESS_SPACE_GLOBAL: 152849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case llvm::ADDRESS_SPACE_CONST: 152949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".ptr .global "; 153049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 153149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 153249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ".align " << (int)getOpenCLAlignment(TD, ETy) << " "; 153349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 153449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printParamName(I, paramIndex, O); 153549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski continue; 153649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 153749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 153849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // non-pointer scalar to kernel func 153949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.param ." 154049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski << getPTXFundamentalTypeStr(Ty) << " "; 154149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printParamName(I, paramIndex, O); 154249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski continue; 154349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 154449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Non-kernel function, just print .param .b<size> for ABI 154549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // and .reg .b<size> for non ABY 154649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned sz = 0; 154749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (isa<IntegerType>(Ty)) { 154849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski sz = cast<IntegerType>(Ty)->getBitWidth(); 154949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (sz < 32) sz = 32; 155049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 155149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else if (isa<PointerType>(Ty)) 155249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski sz = thePointerTy.getSizeInBits(); 155349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 155449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski sz = Ty->getPrimitiveSizeInBits(); 155549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (isABI) 155649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.param .b" << sz << " "; 155749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 155849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.reg .b" << sz << " "; 155949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printParamName(I, paramIndex, O); 156049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski continue; 156149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 156249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 156349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // param has byVal attribute. So should be a pointer 156449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const PointerType *PTy = dyn_cast<PointerType>(Ty); 156549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski assert(PTy && 156649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski "Param with byval attribute should be a pointer type"); 156749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Type *ETy = PTy->getElementType(); 156849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 156949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (isABI || isKernelFunc) { 157049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Just print .param .b8 .align <a> .param[size]; 157149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // <a> = PAL.getparamalignment 157249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // size = typeallocsize of element type 157349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned align = PAL.getParamAlignment(paramIndex+1); 157489443ff7ae81e6a439b5b67e5dd587a450090599Justin Holewinski if (align == 0) 157589443ff7ae81e6a439b5b67e5dd587a450090599Justin Holewinski align = TD->getABITypeAlignment(ETy); 157689443ff7ae81e6a439b5b67e5dd587a450090599Justin Holewinski 157749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned sz = TD->getTypeAllocSize(ETy); 157849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.param .align " << align 157949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski << " .b8 "; 158049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printParamName(I, paramIndex, O); 158149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "[" << sz << "]"; 158249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski continue; 158349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else { 158449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Split the ETy into constituent parts and 158549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // print .param .b<size> <name> for each part. 158649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Further, if a part is vector, print the above for 158749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // each vector element. 158849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski SmallVector<EVT, 16> vtparts; 158949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ComputeValueVTs(*TLI, ETy, vtparts); 159049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (unsigned i=0,e=vtparts.size(); i!=e; ++i) { 159149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned elems = 1; 159249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski EVT elemtype = vtparts[i]; 159349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (vtparts[i].isVector()) { 159449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski elems = vtparts[i].getVectorNumElements(); 159549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski elemtype = vtparts[i].getVectorElementType(); 159649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 159749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 159849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (unsigned j=0,je=elems; j!=je; ++j) { 159949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned sz = elemtype.getSizeInBits(); 160049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (elemtype.isInteger() && (sz < 32)) sz = 32; 160149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.reg .b" << sz << " "; 160249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printParamName(I, paramIndex, O); 160349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (j<je-1) O << ",\n"; 160449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ++paramIndex; 160549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 160649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (i<e-1) 160749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ",\n"; 160849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 160949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski --paramIndex; 161049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski continue; 161149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 161249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 161349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 161449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\n)\n"; 161549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 161649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 161749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::emitFunctionParamList(const MachineFunction &MF, 161849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_ostream &O) { 161949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const Function *F = MF.getFunction(); 162049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski emitFunctionParamList(F, O); 162149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 162249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 162349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 162449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter:: 162549683f3c961379fbc088871a5d6304950f1f1cbcJustin HolewinskisetAndEmitFunctionVirtualRegisters(const MachineFunction &MF) { 162649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski SmallString<128> Str; 162749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_svector_ostream O(Str); 162849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 162949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Map the global virtual register number to a register class specific 163049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // virtual register number starting from 1 with that class. 163149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const TargetRegisterInfo *TRI = MF.getTarget().getRegisterInfo(); 163249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski //unsigned numRegClasses = TRI->getNumRegClasses(); 163349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 163449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Emit the Fake Stack Object 163549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const MachineFrameInfo *MFI = MF.getFrameInfo(); 163649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski int NumBytes = (int) MFI->getStackSize(); 163749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (NumBytes) { 163849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.local .align " << MFI->getMaxAlignment() << " .b8 \t" 163949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski << DEPOTNAME 164049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski << getFunctionNumber() << "[" << NumBytes << "];\n"; 164149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (nvptxSubtarget.is64Bit()) { 164249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.reg .b64 \t%SP;\n"; 164349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.reg .b64 \t%SPL;\n"; 164449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 164549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else { 164649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.reg .b32 \t%SP;\n"; 164749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.reg .b32 \t%SPL;\n"; 164849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 164949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 165049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 165149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Go through all virtual registers to establish the mapping between the 165249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // global virtual 165349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // register number and the per class virtual register number. 165449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // We use the per class virtual register number in the ptx output. 165549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned int numVRs = MRI->getNumVirtRegs(); 165649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (unsigned i=0; i< numVRs; i++) { 165749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned int vr = TRI->index2VirtReg(i); 165849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const TargetRegisterClass *RC = MRI->getRegClass(vr); 165949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski std::map<unsigned, unsigned> ®map = VRidGlobal2LocalMap[RC->getID()]; 166049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski int n = regmap.size(); 166149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski regmap.insert(std::make_pair(vr, n+1)); 166249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 166349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 166449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Emit register declarations 166549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // @TODO: Extract out the real register usage 166649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.reg .pred %p<" << NVPTXNumRegisters << ">;\n"; 166749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.reg .s16 %rc<" << NVPTXNumRegisters << ">;\n"; 166849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.reg .s16 %rs<" << NVPTXNumRegisters << ">;\n"; 166949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.reg .s32 %r<" << NVPTXNumRegisters << ">;\n"; 167049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.reg .s64 %rl<" << NVPTXNumRegisters << ">;\n"; 167149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.reg .f32 %f<" << NVPTXNumRegisters << ">;\n"; 167249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "\t.reg .f64 %fl<" << NVPTXNumRegisters << ">;\n"; 167349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 167449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Emit declaration of the virtual registers or 'physical' registers for 167549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // each register class 167649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski //for (unsigned i=0; i< numRegClasses; i++) { 167749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // std::map<unsigned, unsigned> ®map = VRidGlobal2LocalMap[i]; 167849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // const TargetRegisterClass *RC = TRI->getRegClass(i); 167949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // std::string rcname = getNVPTXRegClassName(RC); 168049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // std::string rcStr = getNVPTXRegClassStr(RC); 168149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // //int n = regmap.size(); 168249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // if (!isNVPTXVectorRegClass(RC)) { 168349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // O << "\t.reg " << rcname << " \t" << rcStr << "<" 168449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // << NVPTXNumRegisters << ">;\n"; 168549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // } 168649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 168749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Only declare those registers that may be used. And do not emit vector 168849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // registers as 168949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // they are all elementized to scalar registers. 169049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski //if (n && !isNVPTXVectorRegClass(RC)) { 169149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // if (RegAllocNilUsed) { 169249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // O << "\t.reg " << rcname << " \t" << rcStr << "<" << (n+1) 169349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // << ">;\n"; 169449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // } 169549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // else { 169649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // O << "\t.reg " << rcname << " \t" << StrToUpper(rcStr) 169749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // << "<" << 32 << ">;\n"; 169849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // } 169949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski //} 170049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski //} 170149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 170249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski OutStreamer.EmitRawText(O.str()); 170349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 170449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 170549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 170649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::printFPConstant(const ConstantFP *Fp, raw_ostream &O) { 170749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski APFloat APF = APFloat(Fp->getValueAPF()); // make a copy 170849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski bool ignored; 170949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned int numHex; 171049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const char *lead; 171149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 171249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (Fp->getType()->getTypeID()==Type::FloatTyID) { 171349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski numHex = 8; 171449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski lead = "0f"; 171549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski APF.convert(APFloat::IEEEsingle, APFloat::rmNearestTiesToEven, 171649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski &ignored); 171749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else if (Fp->getType()->getTypeID() == Type::DoubleTyID) { 171849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski numHex = 16; 171949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski lead = "0d"; 172049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski APF.convert(APFloat::IEEEdouble, APFloat::rmNearestTiesToEven, 172149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski &ignored); 172249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else 172349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski llvm_unreachable("unsupported fp type"); 172449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 172549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski APInt API = APF.bitcastToAPInt(); 172649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski std::string hexstr(utohexstr(API.getZExtValue())); 172749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << lead; 172849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (hexstr.length() < numHex) 172949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << std::string(numHex - hexstr.length(), '0'); 173049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << utohexstr(API.getZExtValue()); 173149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 173249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 173349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::printScalarConstant(Constant *CPV, raw_ostream &O) { 173449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) { 173549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << CI->getValue(); 173649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 173749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 173849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (ConstantFP *CFP = dyn_cast<ConstantFP>(CPV)) { 173949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printFPConstant(CFP, O); 174049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 174149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 174249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (isa<ConstantPointerNull>(CPV)) { 174349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << "0"; 174449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 174549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 174649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) { 174749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << *Mang->getSymbol(GVar); 174849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 174949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 175049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) { 175149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Value *v = Cexpr->stripPointerCasts(); 175249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (GlobalValue *GVar = dyn_cast<GlobalValue>(v)) { 175349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << *Mang->getSymbol(GVar); 175449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 175549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else { 175649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << *LowerConstant(CPV, *this); 175749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 175849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 175949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 176049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski llvm_unreachable("Not scalar type found in printScalarConstant()"); 176149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 176249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 176349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 176449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::bufferLEByte(Constant *CPV, int Bytes, 176549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski AggBuffer *aggBuffer) { 176649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 17673574eca1b02600bac4e625297f4ecf745f4c4f32Micah Villmow const DataLayout *TD = TM.getDataLayout(); 176849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 176949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (isa<UndefValue>(CPV) || CPV->isNullValue()) { 177049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski int s = TD->getTypeAllocSize(CPV->getType()); 177149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (s<Bytes) 177249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski s = Bytes; 177349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer->addZeros(s); 177449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 177549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 177649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 177749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned char *ptr; 177849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski switch (CPV->getType()->getTypeID()) { 177949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 178049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Type::IntegerTyID: { 178149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const Type *ETy = CPV->getType(); 178249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if ( ETy == Type::getInt8Ty(CPV->getContext()) ){ 178349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned char c = 178449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski (unsigned char)(dyn_cast<ConstantInt>(CPV))->getZExtValue(); 178549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ptr = &c; 178649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer->addBytes(ptr, 1, Bytes); 178749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else if ( ETy == Type::getInt16Ty(CPV->getContext()) ) { 178849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski short int16 = 178949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski (short)(dyn_cast<ConstantInt>(CPV))->getZExtValue(); 179049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ptr = (unsigned char*)&int16; 179149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer->addBytes(ptr, 2, Bytes); 179249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else if ( ETy == Type::getInt32Ty(CPV->getContext()) ) { 179349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (ConstantInt *constInt = dyn_cast<ConstantInt>(CPV)) { 179449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski int int32 =(int)(constInt->getZExtValue()); 179549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ptr = (unsigned char*)&int32; 179649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer->addBytes(ptr, 4, Bytes); 179749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 17986366361998599a63a74ee7cc5f4ba900711c7e7aCraig Topper } else if (ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) { 179949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (ConstantInt *constInt = 180049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski dyn_cast<ConstantInt>(ConstantFoldConstantExpression( 180149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Cexpr, TD))) { 180249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski int int32 =(int)(constInt->getZExtValue()); 180349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ptr = (unsigned char*)&int32; 180449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer->addBytes(ptr, 4, Bytes); 180549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 180649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 180749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (Cexpr->getOpcode() == Instruction::PtrToInt) { 180849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Value *v = Cexpr->getOperand(0)->stripPointerCasts(); 180949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer->addSymbol(v); 181049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer->addZeros(4); 181149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 181249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 181349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 18146366361998599a63a74ee7cc5f4ba900711c7e7aCraig Topper llvm_unreachable("unsupported integer const type"); 181549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else if (ETy == Type::getInt64Ty(CPV->getContext()) ) { 181649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (ConstantInt *constInt = dyn_cast<ConstantInt>(CPV)) { 181749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski long long int64 =(long long)(constInt->getZExtValue()); 181849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ptr = (unsigned char*)&int64; 181949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer->addBytes(ptr, 8, Bytes); 182049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 18216366361998599a63a74ee7cc5f4ba900711c7e7aCraig Topper } else if (ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) { 182249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (ConstantInt *constInt = dyn_cast<ConstantInt>( 182349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ConstantFoldConstantExpression(Cexpr, TD))) { 182449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski long long int64 =(long long)(constInt->getZExtValue()); 182549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ptr = (unsigned char*)&int64; 182649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer->addBytes(ptr, 8, Bytes); 182749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 182849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 182949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (Cexpr->getOpcode() == Instruction::PtrToInt) { 183049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Value *v = Cexpr->getOperand(0)->stripPointerCasts(); 183149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer->addSymbol(v); 183249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer->addZeros(8); 183349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 183449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 183549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 183649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski llvm_unreachable("unsupported integer const type"); 18376366361998599a63a74ee7cc5f4ba900711c7e7aCraig Topper } else 183849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski llvm_unreachable("unsupported integer const type"); 183949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 184049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 184149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Type::FloatTyID: 184249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Type::DoubleTyID: { 184349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ConstantFP *CFP = dyn_cast<ConstantFP>(CPV); 184449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const Type* Ty = CFP->getType(); 184549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (Ty == Type::getFloatTy(CPV->getContext())) { 184649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski float float32 = (float)CFP->getValueAPF().convertToFloat(); 184749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ptr = (unsigned char*)&float32; 184849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer->addBytes(ptr, 4, Bytes); 184949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } else if (Ty == Type::getDoubleTy(CPV->getContext())) { 185049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski double float64 = CFP->getValueAPF().convertToDouble(); 185149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski ptr = (unsigned char*)&float64; 185249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer->addBytes(ptr, 8, Bytes); 185349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 185449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else { 185549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski llvm_unreachable("unsupported fp const type"); 185649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 185749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 185849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 185949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Type::PointerTyID: { 186049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) { 186149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer->addSymbol(GVar); 186249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 186349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else if (ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) { 186449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Value *v = Cexpr->stripPointerCasts(); 186549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer->addSymbol(v); 186649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 186749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned int s = TD->getTypeAllocSize(CPV->getType()); 186849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer->addZeros(s); 186949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 187049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 187149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 187249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Type::ArrayTyID: 187349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Type::VectorTyID: 187449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case Type::StructTyID: { 187549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (isa<ConstantArray>(CPV) || isa<ConstantVector>(CPV) || 187649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski isa<ConstantStruct>(CPV)) { 187749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski int ElementSize = TD->getTypeAllocSize(CPV->getType()); 187849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski bufferAggregateConstant(CPV, aggBuffer); 187949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if ( Bytes > ElementSize ) 188049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer->addZeros(Bytes-ElementSize); 188149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 188249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else if (isa<ConstantAggregateZero>(CPV)) 188349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer->addZeros(Bytes); 188449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 188549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski llvm_unreachable("Unexpected Constant type"); 188649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 188749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 188849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 188949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski default: 189049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski llvm_unreachable("unsupported type"); 189149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 189249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 189349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 189449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::bufferAggregateConstant(Constant *CPV, 189549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski AggBuffer *aggBuffer) { 18963574eca1b02600bac4e625297f4ecf745f4c4f32Micah Villmow const DataLayout *TD = TM.getDataLayout(); 189749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski int Bytes; 189849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 189949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski // Old constants 190049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (isa<ConstantArray>(CPV) || isa<ConstantVector>(CPV)) { 190149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (CPV->getNumOperands()) 190249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (unsigned i = 0, e = CPV->getNumOperands(); i != e; ++i) 190349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski bufferLEByte(cast<Constant>(CPV->getOperand(i)), 0, aggBuffer); 190449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 190549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 190649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 190749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (const ConstantDataSequential *CDS = 190849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski dyn_cast<ConstantDataSequential>(CPV)) { 190949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (CDS->getNumElements()) 191049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (unsigned i = 0; i < CDS->getNumElements(); ++i) 191149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski bufferLEByte(cast<Constant>(CDS->getElementAsConstant(i)), 0, 191249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer); 191349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 191449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 191549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 191649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 191749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (isa<ConstantStruct>(CPV)) { 191849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (CPV->getNumOperands()) { 191949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski StructType *ST = cast<StructType>(CPV->getType()); 192049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski for (unsigned i = 0, e = CPV->getNumOperands(); i != e; ++i) { 192149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if ( i == (e - 1)) 192249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Bytes = TD->getStructLayout(ST)->getElementOffset(0) + 192349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski TD->getTypeAllocSize(ST) 192449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski - TD->getStructLayout(ST)->getElementOffset(i); 192549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski else 192649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski Bytes = TD->getStructLayout(ST)->getElementOffset(i+1) - 192749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski TD->getStructLayout(ST)->getElementOffset(i); 192849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski bufferLEByte(cast<Constant>(CPV->getOperand(i)), Bytes, 192949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski aggBuffer); 193049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 193149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 193249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return; 193349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 19346366361998599a63a74ee7cc5f4ba900711c7e7aCraig Topper llvm_unreachable("unsupported constant type in printAggregateConstant()"); 193549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 193649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 193749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// buildTypeNameMap - Run through symbol table looking for type names. 193849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// 193949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 194049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 194149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskibool NVPTXAsmPrinter::isImageType(const Type *Ty) { 194249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 194349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski std::map<const Type *, std::string>::iterator PI = TypeNameMap.find(Ty); 194449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 194549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (PI != TypeNameMap.end() && 194649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski (!PI->second.compare("struct._image1d_t") || 194749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski !PI->second.compare("struct._image2d_t") || 194849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski !PI->second.compare("struct._image3d_t"))) 194949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return true; 195049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 195149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return false; 195249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 195349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 195449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski/// PrintAsmOperand - Print out an operand for an inline asm expression. 195549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski/// 195649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskibool NVPTXAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo, 195749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned AsmVariant, 195849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const char *ExtraCode, 195949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_ostream &O) { 196049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (ExtraCode && ExtraCode[0]) { 196149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (ExtraCode[1] != 0) return true; // Unknown modifier. 196249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 196349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski switch (ExtraCode[0]) { 19640518fca843ff87d069ecb07fc00d306c1f587d58Jack Carter default: 19650518fca843ff87d069ecb07fc00d306c1f587d58Jack Carter // See if this is a generic print operand 19660518fca843ff87d069ecb07fc00d306c1f587d58Jack Carter return AsmPrinter::PrintAsmOperand(MI, OpNo, AsmVariant, ExtraCode, O); 196749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case 'r': 196849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski break; 196949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 197049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 197149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 197249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printOperand(MI, OpNo, O); 197349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 197449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return false; 197549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 197649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 197749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskibool NVPTXAsmPrinter::PrintAsmMemoryOperand(const MachineInstr *MI, 197849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned OpNo, 197949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski unsigned AsmVariant, 198049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski const char *ExtraCode, 198149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski raw_ostream &O) { 198249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (ExtraCode && ExtraCode[0]) 198349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return true; // Unknown modifier 198449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 198549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << '['; 198649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski printMemOperand(MI, OpNo, O); 198749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski O << ']'; 198849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 198949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return false; 199049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 199149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 199249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskibool NVPTXAsmPrinter::ignoreLoc(const MachineInstr &MI) 199349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski{ 199449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski switch(MI.getOpcode()) { 199549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski default: 199649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return false; 199749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::CallArgBeginInst: case NVPTX::CallArgEndInst0: 199849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::CallArgEndInst1: case NVPTX::CallArgF32: 199949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::CallArgF64: case NVPTX::CallArgI16: 200049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::CallArgI32: case NVPTX::CallArgI32imm: 200149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::CallArgI64: case NVPTX::CallArgI8: 200249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::CallArgParam: case NVPTX::CallVoidInst: 200349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::CallVoidInstReg: case NVPTX::Callseq_End: 200449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::CallVoidInstReg64: 200549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::DeclareParamInst: case NVPTX::DeclareRetMemInst: 200649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::DeclareRetRegInst: case NVPTX::DeclareRetScalarInst: 200749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::DeclareScalarParamInst: case NVPTX::DeclareScalarRegInst: 200849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::StoreParamF32: case NVPTX::StoreParamF64: 200949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::StoreParamI16: case NVPTX::StoreParamI32: 201049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::StoreParamI64: case NVPTX::StoreParamI8: 201149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::StoreParamS32I8: case NVPTX::StoreParamU32I8: 201249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::StoreParamS32I16: case NVPTX::StoreParamU32I16: 201349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::StoreRetvalF32: case NVPTX::StoreRetvalF64: 201449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::StoreRetvalI16: case NVPTX::StoreRetvalI32: 201549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::StoreRetvalI64: case NVPTX::StoreRetvalI8: 201649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::LastCallArgF32: case NVPTX::LastCallArgF64: 201749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::LastCallArgI16: case NVPTX::LastCallArgI32: 201849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::LastCallArgI32imm: case NVPTX::LastCallArgI64: 201949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::LastCallArgI8: case NVPTX::LastCallArgParam: 202049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::LoadParamMemF32: case NVPTX::LoadParamMemF64: 202149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::LoadParamMemI16: case NVPTX::LoadParamMemI32: 202249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::LoadParamMemI64: case NVPTX::LoadParamMemI8: 202349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::LoadParamRegF32: case NVPTX::LoadParamRegF64: 202449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::LoadParamRegI16: case NVPTX::LoadParamRegI32: 202549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::LoadParamRegI64: case NVPTX::LoadParamRegI8: 202649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski case NVPTX::PrototypeInst: case NVPTX::DBG_VALUE: 202749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return true; 202849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 202949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return false; 203049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 203149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 203249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// Force static initialization. 203349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskiextern "C" void LLVMInitializeNVPTXBackendAsmPrinter() { 203449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski RegisterAsmPrinter<NVPTXAsmPrinter> X(TheNVPTXTarget32); 203549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski RegisterAsmPrinter<NVPTXAsmPrinter> Y(TheNVPTXTarget64); 203649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 203749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 203849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 203949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskivoid NVPTXAsmPrinter::emitSrcInText(StringRef filename, unsigned line) { 204049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski std::stringstream temp; 204149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski LineReader * reader = this->getReader(filename.str()); 204249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski temp << "\n//"; 204349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski temp << filename.str(); 204449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski temp << ":"; 204549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski temp << line; 204649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski temp << " "; 204749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski temp << reader->readLine(line); 204849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski temp << "\n"; 204949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski this->OutStreamer.EmitRawText(Twine(temp.str())); 205049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 205149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 205249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 205349683f3c961379fbc088871a5d6304950f1f1cbcJustin HolewinskiLineReader *NVPTXAsmPrinter::getReader(std::string filename) { 205449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (reader == NULL) { 205549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski reader = new LineReader(filename); 205649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 205749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 205849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (reader->fileName() != filename) { 205949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski delete reader; 206049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski reader = new LineReader(filename); 206149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 206249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 206349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return reader; 206449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 206549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 206649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 206749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskistd::string 206849683f3c961379fbc088871a5d6304950f1f1cbcJustin HolewinskiLineReader::readLine(unsigned lineNum) { 206949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski if (lineNum < theCurLine) { 207049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski theCurLine = 0; 207149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski fstr.seekg(0,std::ios::beg); 207249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 207349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski while (theCurLine < lineNum) { 207449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski fstr.getline(buff,500); 207549683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski theCurLine++; 207649683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski } 207749683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski return buff; 207849683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 207949683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski 208049683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski// Force static initialization. 208149683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinskiextern "C" void LLVMInitializeNVPTXAsmPrinter() { 208249683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski RegisterAsmPrinter<NVPTXAsmPrinter> X(TheNVPTXTarget32); 208349683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski RegisterAsmPrinter<NVPTXAsmPrinter> Y(TheNVPTXTarget64); 208449683f3c961379fbc088871a5d6304950f1f1cbcJustin Holewinski} 2085