2012-05-05 04:18:50 +08:00
|
|
|
//===-- NVPTXAsmPrinter.cpp - NVPTX LLVM assembly writer ------------------===//
|
|
|
|
//
|
|
|
|
// The LLVM Compiler Infrastructure
|
|
|
|
//
|
|
|
|
// This file is distributed under the University of Illinois Open Source
|
|
|
|
// License. See LICENSE.TXT for details.
|
|
|
|
//
|
|
|
|
//===----------------------------------------------------------------------===//
|
|
|
|
//
|
|
|
|
// This file contains a printer that converts from our internal representation
|
|
|
|
// of machine-dependent LLVM code to NVPTX assembly language.
|
|
|
|
//
|
|
|
|
//===----------------------------------------------------------------------===//
|
|
|
|
|
2014-01-07 19:48:04 +08:00
|
|
|
#include "InstPrinter/NVPTXInstPrinter.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include "MCTargetDesc/NVPTXBaseInfo.h"
|
2012-12-04 00:50:05 +08:00
|
|
|
#include "MCTargetDesc/NVPTXMCAsmInfo.h"
|
2012-05-05 04:18:50 +08:00
|
|
|
#include "NVPTX.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include "NVPTXAsmPrinter.h"
|
2013-08-06 22:13:27 +08:00
|
|
|
#include "NVPTXMCExpr.h"
|
2015-01-14 19:23:27 +08:00
|
|
|
#include "NVPTXMachineFunctionInfo.h"
|
2012-05-05 04:18:50 +08:00
|
|
|
#include "NVPTXRegisterInfo.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include "NVPTXSubtarget.h"
|
2012-12-04 00:50:05 +08:00
|
|
|
#include "NVPTXTargetMachine.h"
|
2012-06-28 08:05:13 +08:00
|
|
|
#include "NVPTXUtilities.h"
|
2012-12-04 00:50:05 +08:00
|
|
|
#include "cl_common_defines.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include "llvm/ADT/APFloat.h"
|
|
|
|
#include "llvm/ADT/APInt.h"
|
|
|
|
#include "llvm/ADT/DenseMap.h"
|
|
|
|
#include "llvm/ADT/DenseSet.h"
|
|
|
|
#include "llvm/ADT/SmallString.h"
|
|
|
|
#include "llvm/ADT/SmallVector.h"
|
2012-05-05 04:18:50 +08:00
|
|
|
#include "llvm/ADT/StringExtras.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include "llvm/ADT/StringRef.h"
|
|
|
|
#include "llvm/ADT/Triple.h"
|
|
|
|
#include "llvm/ADT/Twine.h"
|
2012-12-04 00:50:05 +08:00
|
|
|
#include "llvm/Analysis/ConstantFolding.h"
|
2012-05-05 04:18:50 +08:00
|
|
|
#include "llvm/CodeGen/Analysis.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include "llvm/CodeGen/MachineBasicBlock.h"
|
2012-05-05 04:18:50 +08:00
|
|
|
#include "llvm/CodeGen/MachineFrameInfo.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include "llvm/CodeGen/MachineFunction.h"
|
|
|
|
#include "llvm/CodeGen/MachineInstr.h"
|
2015-02-01 10:27:45 +08:00
|
|
|
#include "llvm/CodeGen/MachineLoopInfo.h"
|
2012-05-05 04:18:50 +08:00
|
|
|
#include "llvm/CodeGen/MachineModuleInfo.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include "llvm/CodeGen/MachineOperand.h"
|
2012-12-04 00:50:05 +08:00
|
|
|
#include "llvm/CodeGen/MachineRegisterInfo.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include "llvm/CodeGen/MachineValueType.h"
|
|
|
|
#include "llvm/CodeGen/ValueTypes.h"
|
|
|
|
#include "llvm/IR/Attributes.h"
|
|
|
|
#include "llvm/IR/BasicBlock.h"
|
|
|
|
#include "llvm/IR/Constant.h"
|
|
|
|
#include "llvm/IR/Constants.h"
|
|
|
|
#include "llvm/IR/DataLayout.h"
|
2014-03-06 08:46:21 +08:00
|
|
|
#include "llvm/IR/DebugInfo.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include "llvm/IR/DebugInfoMetadata.h"
|
|
|
|
#include "llvm/IR/DebugLoc.h"
|
2013-01-02 19:36:10 +08:00
|
|
|
#include "llvm/IR/DerivedTypes.h"
|
|
|
|
#include "llvm/IR/Function.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include "llvm/IR/GlobalValue.h"
|
2013-01-02 19:36:10 +08:00
|
|
|
#include "llvm/IR/GlobalVariable.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include "llvm/IR/Instruction.h"
|
|
|
|
#include "llvm/IR/LLVMContext.h"
|
2013-01-02 19:36:10 +08:00
|
|
|
#include "llvm/IR/Module.h"
|
|
|
|
#include "llvm/IR/Operator.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include "llvm/IR/Type.h"
|
|
|
|
#include "llvm/IR/User.h"
|
|
|
|
#include "llvm/MC/MCExpr.h"
|
2015-05-16 06:19:42 +08:00
|
|
|
#include "llvm/MC/MCInst.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include "llvm/MC/MCInstrDesc.h"
|
2012-05-05 04:18:50 +08:00
|
|
|
#include "llvm/MC/MCStreamer.h"
|
|
|
|
#include "llvm/MC/MCSymbol.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include "llvm/Support/Casting.h"
|
2012-12-04 00:50:05 +08:00
|
|
|
#include "llvm/Support/CommandLine.h"
|
2012-05-05 04:18:50 +08:00
|
|
|
#include "llvm/Support/ErrorHandling.h"
|
|
|
|
#include "llvm/Support/Path.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include "llvm/Support/raw_ostream.h"
|
2012-12-04 00:50:05 +08:00
|
|
|
#include "llvm/Support/TargetRegistry.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include "llvm/Target/TargetLowering.h"
|
2012-12-04 00:50:05 +08:00
|
|
|
#include "llvm/Target/TargetLoweringObjectFile.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include "llvm/Target/TargetMachine.h"
|
|
|
|
#include "llvm/Target/TargetRegisterInfo.h"
|
2015-02-01 10:27:45 +08:00
|
|
|
#include "llvm/Transforms/Utils/UnrollLoop.h"
|
2017-01-10 06:16:51 +08:00
|
|
|
#include <cassert>
|
|
|
|
#include <cstdint>
|
|
|
|
#include <cstring>
|
|
|
|
#include <new>
|
2012-06-28 08:05:13 +08:00
|
|
|
#include <sstream>
|
2017-01-10 06:16:51 +08:00
|
|
|
#include <string>
|
|
|
|
#include <utility>
|
|
|
|
#include <vector>
|
|
|
|
|
2012-05-05 04:18:50 +08:00
|
|
|
using namespace llvm;
|
|
|
|
|
|
|
|
#define DEPOTNAME "__local_depot"
|
|
|
|
|
|
|
|
static cl::opt<bool>
|
2013-10-19 07:38:13 +08:00
|
|
|
EmitLineNumbers("nvptx-emit-line-numbers", cl::Hidden,
|
2012-05-05 04:18:50 +08:00
|
|
|
cl::desc("NVPTX Specific: Emit Line numbers even without -G"),
|
|
|
|
cl::init(true));
|
|
|
|
|
2013-10-27 19:31:46 +08:00
|
|
|
static cl::opt<bool>
|
2013-10-19 07:38:13 +08:00
|
|
|
InterleaveSrc("nvptx-emit-src", cl::ZeroOrMore, cl::Hidden,
|
2013-03-30 22:29:21 +08:00
|
|
|
cl::desc("NVPTX Specific: Emit source line in ptx file"),
|
2013-10-27 19:31:46 +08:00
|
|
|
cl::init(false));
|
2012-05-05 04:18:50 +08:00
|
|
|
|
2012-11-17 05:03:51 +08:00
|
|
|
/// DiscoverDependentGlobals - Return a set of GlobalVariables on which \p V
|
|
|
|
/// depends.
|
2017-01-10 06:16:51 +08:00
|
|
|
static void
|
|
|
|
DiscoverDependentGlobals(const Value *V,
|
|
|
|
DenseSet<const GlobalVariable *> &Globals) {
|
2013-05-20 20:13:32 +08:00
|
|
|
if (const GlobalVariable *GV = dyn_cast<GlobalVariable>(V))
|
2012-11-17 05:03:51 +08:00
|
|
|
Globals.insert(GV);
|
|
|
|
else {
|
2013-05-20 20:13:32 +08:00
|
|
|
if (const User *U = dyn_cast<User>(V)) {
|
2012-11-17 05:03:51 +08:00
|
|
|
for (unsigned i = 0, e = U->getNumOperands(); i != e; ++i) {
|
|
|
|
DiscoverDependentGlobals(U->getOperand(i), Globals);
|
|
|
|
}
|
|
|
|
}
|
|
|
|
}
|
|
|
|
}
|
2012-05-05 04:18:50 +08:00
|
|
|
|
2012-11-17 05:03:51 +08:00
|
|
|
/// VisitGlobalVariableForEmission - Add \p GV to the list of GlobalVariable
|
|
|
|
/// instances to be emitted, but only after any dependents have been added
|
2017-01-10 06:16:51 +08:00
|
|
|
/// first.s
|
|
|
|
static void
|
|
|
|
VisitGlobalVariableForEmission(const GlobalVariable *GV,
|
|
|
|
SmallVectorImpl<const GlobalVariable *> &Order,
|
|
|
|
DenseSet<const GlobalVariable *> &Visited,
|
|
|
|
DenseSet<const GlobalVariable *> &Visiting) {
|
2012-11-17 05:03:51 +08:00
|
|
|
// Have we already visited this one?
|
2013-03-30 22:29:21 +08:00
|
|
|
if (Visited.count(GV))
|
|
|
|
return;
|
2012-11-17 05:03:51 +08:00
|
|
|
|
|
|
|
// Do we have a circular dependency?
|
2014-10-10 23:32:50 +08:00
|
|
|
if (!Visiting.insert(GV).second)
|
2012-11-17 05:03:51 +08:00
|
|
|
report_fatal_error("Circular dependency found in global variable set");
|
|
|
|
|
|
|
|
// Make sure we visit all dependents first
|
2013-05-20 20:13:32 +08:00
|
|
|
DenseSet<const GlobalVariable *> Others;
|
2012-11-17 05:03:51 +08:00
|
|
|
for (unsigned i = 0, e = GV->getNumOperands(); i != e; ++i)
|
|
|
|
DiscoverDependentGlobals(GV->getOperand(i), Others);
|
2013-03-30 22:29:21 +08:00
|
|
|
|
2013-05-20 20:13:32 +08:00
|
|
|
for (DenseSet<const GlobalVariable *>::iterator I = Others.begin(),
|
|
|
|
E = Others.end();
|
2013-03-30 22:29:21 +08:00
|
|
|
I != E; ++I)
|
2012-11-17 05:03:51 +08:00
|
|
|
VisitGlobalVariableForEmission(*I, Order, Visited, Visiting);
|
|
|
|
|
|
|
|
// Now we can visit ourself
|
|
|
|
Order.push_back(GV);
|
|
|
|
Visited.insert(GV);
|
|
|
|
Visiting.erase(GV);
|
|
|
|
}
|
2012-05-05 04:18:50 +08:00
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
void NVPTXAsmPrinter::emitLineNumberAsDotLoc(const MachineInstr &MI) {
|
2012-05-05 04:18:50 +08:00
|
|
|
if (!EmitLineNumbers)
|
|
|
|
return;
|
|
|
|
if (ignoreLoc(MI))
|
|
|
|
return;
|
|
|
|
|
2016-05-27 20:30:51 +08:00
|
|
|
const DebugLoc &curLoc = MI.getDebugLoc();
|
2012-05-05 04:18:50 +08:00
|
|
|
|
2015-03-31 03:14:47 +08:00
|
|
|
if (!prevDebugLoc && !curLoc)
|
2012-05-05 04:18:50 +08:00
|
|
|
return;
|
|
|
|
|
|
|
|
if (prevDebugLoc == curLoc)
|
|
|
|
return;
|
|
|
|
|
|
|
|
prevDebugLoc = curLoc;
|
|
|
|
|
2015-03-31 03:14:47 +08:00
|
|
|
if (!curLoc)
|
2012-05-05 04:18:50 +08:00
|
|
|
return;
|
|
|
|
|
2015-04-30 00:38:44 +08:00
|
|
|
auto *Scope = cast_or_null<DIScope>(curLoc.getScope());
|
2013-06-28 13:43:10 +08:00
|
|
|
if (!Scope)
|
|
|
|
return;
|
2012-05-05 04:18:50 +08:00
|
|
|
|
2015-04-16 09:37:00 +08:00
|
|
|
StringRef fileName(Scope->getFilename());
|
|
|
|
StringRef dirName(Scope->getDirectory());
|
2012-05-05 04:18:50 +08:00
|
|
|
SmallString<128> FullPathName = dirName;
|
|
|
|
if (!dirName.empty() && !sys::path::is_absolute(fileName)) {
|
|
|
|
sys::path::append(FullPathName, fileName);
|
2015-03-30 23:42:36 +08:00
|
|
|
fileName = FullPathName;
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
2015-03-30 23:42:36 +08:00
|
|
|
if (filenameMap.find(fileName) == filenameMap.end())
|
2012-05-05 04:18:50 +08:00
|
|
|
return;
|
|
|
|
|
|
|
|
// Emit the line from the source file.
|
2013-10-27 19:31:46 +08:00
|
|
|
if (InterleaveSrc)
|
2015-03-30 23:42:36 +08:00
|
|
|
this->emitSrcInText(fileName, curLoc.getLine());
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
std::stringstream temp;
|
2015-03-30 23:42:36 +08:00
|
|
|
temp << "\t.loc " << filenameMap[fileName] << " " << curLoc.getLine()
|
2013-03-30 22:29:21 +08:00
|
|
|
<< " " << curLoc.getCol();
|
2015-04-25 03:11:51 +08:00
|
|
|
OutStreamer->EmitRawText(temp.str());
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
void NVPTXAsmPrinter::EmitInstruction(const MachineInstr *MI) {
|
|
|
|
SmallString<128> Str;
|
|
|
|
raw_svector_ostream OS(Str);
|
2015-02-19 08:08:23 +08:00
|
|
|
if (static_cast<NVPTXTargetMachine &>(TM).getDrvInterface() == NVPTX::CUDA)
|
2012-05-05 04:18:50 +08:00
|
|
|
emitLineNumberAsDotLoc(*MI);
|
2013-08-06 22:13:27 +08:00
|
|
|
|
|
|
|
MCInst Inst;
|
|
|
|
lowerToMCInst(MI, Inst);
|
2015-04-25 03:11:51 +08:00
|
|
|
EmitToStreamer(*OutStreamer, Inst);
|
2013-08-06 22:13:27 +08:00
|
|
|
}
|
|
|
|
|
2014-04-09 23:39:15 +08:00
|
|
|
// Handle symbol backtracking for targets that do not support image handles
|
|
|
|
bool NVPTXAsmPrinter::lowerImageHandleOperand(const MachineInstr *MI,
|
|
|
|
unsigned OpNo, MCOperand &MCOp) {
|
|
|
|
const MachineOperand &MO = MI->getOperand(OpNo);
|
2014-07-17 19:59:04 +08:00
|
|
|
const MCInstrDesc &MCID = MI->getDesc();
|
2014-04-09 23:39:15 +08:00
|
|
|
|
2014-07-17 19:59:04 +08:00
|
|
|
if (MCID.TSFlags & NVPTXII::IsTexFlag) {
|
2014-04-09 23:39:15 +08:00
|
|
|
// This is a texture fetch, so operand 4 is a texref and operand 5 is
|
|
|
|
// a samplerref
|
2014-07-17 19:59:04 +08:00
|
|
|
if (OpNo == 4 && MO.isImm()) {
|
2014-04-09 23:39:15 +08:00
|
|
|
lowerImageHandleSymbol(MO.getImm(), MCOp);
|
|
|
|
return true;
|
|
|
|
}
|
2014-07-17 19:59:04 +08:00
|
|
|
if (OpNo == 5 && MO.isImm() && !(MCID.TSFlags & NVPTXII::IsTexModeUnifiedFlag)) {
|
2014-04-09 23:39:15 +08:00
|
|
|
lowerImageHandleSymbol(MO.getImm(), MCOp);
|
|
|
|
return true;
|
|
|
|
}
|
|
|
|
|
|
|
|
return false;
|
2014-07-17 19:59:04 +08:00
|
|
|
} else if (MCID.TSFlags & NVPTXII::IsSuldMask) {
|
|
|
|
unsigned VecSize =
|
|
|
|
1 << (((MCID.TSFlags & NVPTXII::IsSuldMask) >> NVPTXII::IsSuldShift) - 1);
|
2014-04-09 23:39:15 +08:00
|
|
|
|
2014-07-17 19:59:04 +08:00
|
|
|
// For a surface load of vector size N, the Nth operand will be the surfref
|
|
|
|
if (OpNo == VecSize && MO.isImm()) {
|
2014-04-09 23:39:15 +08:00
|
|
|
lowerImageHandleSymbol(MO.getImm(), MCOp);
|
|
|
|
return true;
|
|
|
|
}
|
|
|
|
|
|
|
|
return false;
|
2014-07-17 19:59:04 +08:00
|
|
|
} else if (MCID.TSFlags & NVPTXII::IsSustFlag) {
|
2014-04-09 23:39:15 +08:00
|
|
|
// This is a surface store, so operand 0 is a surfref
|
2014-07-17 19:59:04 +08:00
|
|
|
if (OpNo == 0 && MO.isImm()) {
|
2014-04-09 23:39:15 +08:00
|
|
|
lowerImageHandleSymbol(MO.getImm(), MCOp);
|
|
|
|
return true;
|
|
|
|
}
|
|
|
|
|
|
|
|
return false;
|
2014-07-17 19:59:04 +08:00
|
|
|
} else if (MCID.TSFlags & NVPTXII::IsSurfTexQueryFlag) {
|
2014-04-09 23:39:15 +08:00
|
|
|
// This is a query, so operand 1 is a surfref/texref
|
2014-07-17 19:59:04 +08:00
|
|
|
if (OpNo == 1 && MO.isImm()) {
|
2014-04-09 23:39:15 +08:00
|
|
|
lowerImageHandleSymbol(MO.getImm(), MCOp);
|
|
|
|
return true;
|
|
|
|
}
|
|
|
|
|
|
|
|
return false;
|
|
|
|
}
|
2014-07-17 19:59:04 +08:00
|
|
|
|
|
|
|
return false;
|
2014-04-09 23:39:15 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
void NVPTXAsmPrinter::lowerImageHandleSymbol(unsigned Index, MCOperand &MCOp) {
|
|
|
|
// Ewwww
|
|
|
|
TargetMachine &TM = const_cast<TargetMachine&>(MF->getTarget());
|
|
|
|
NVPTXTargetMachine &nvTM = static_cast<NVPTXTargetMachine&>(TM);
|
|
|
|
const NVPTXMachineFunctionInfo *MFI = MF->getInfo<NVPTXMachineFunctionInfo>();
|
|
|
|
const char *Sym = MFI->getImageHandleSymbol(Index);
|
|
|
|
std::string *SymNamePtr =
|
|
|
|
nvTM.getManagedStrPool()->getManagedString(Sym);
|
2016-11-03 00:43:50 +08:00
|
|
|
MCOp = GetSymbolRef(OutContext.getOrCreateSymbol(StringRef(*SymNamePtr)));
|
2014-04-09 23:39:15 +08:00
|
|
|
}
|
|
|
|
|
2013-08-06 22:13:27 +08:00
|
|
|
void NVPTXAsmPrinter::lowerToMCInst(const MachineInstr *MI, MCInst &OutMI) {
|
|
|
|
OutMI.setOpcode(MI->getOpcode());
|
2013-11-15 20:30:04 +08:00
|
|
|
// Special: Do not mangle symbol operand of CALL_PROTOTYPE
|
|
|
|
if (MI->getOpcode() == NVPTX::CALL_PROTOTYPE) {
|
|
|
|
const MachineOperand &MO = MI->getOperand(0);
|
2014-04-09 23:39:15 +08:00
|
|
|
OutMI.addOperand(GetSymbolRef(
|
2015-05-19 02:43:14 +08:00
|
|
|
OutContext.getOrCreateSymbol(Twine(MO.getSymbolName()))));
|
2013-11-15 20:30:04 +08:00
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2013-08-06 22:13:27 +08:00
|
|
|
for (unsigned i = 0, e = MI->getNumOperands(); i != e; ++i) {
|
|
|
|
const MachineOperand &MO = MI->getOperand(i);
|
|
|
|
|
|
|
|
MCOperand MCOp;
|
2015-02-19 08:08:14 +08:00
|
|
|
if (!nvptxSubtarget->hasImageHandles()) {
|
2014-04-09 23:39:15 +08:00
|
|
|
if (lowerImageHandleOperand(MI, i, MCOp)) {
|
|
|
|
OutMI.addOperand(MCOp);
|
|
|
|
continue;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2013-08-06 22:13:27 +08:00
|
|
|
if (lowerOperand(MO, MCOp))
|
|
|
|
OutMI.addOperand(MCOp);
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
bool NVPTXAsmPrinter::lowerOperand(const MachineOperand &MO,
|
|
|
|
MCOperand &MCOp) {
|
|
|
|
switch (MO.getType()) {
|
|
|
|
default: llvm_unreachable("unknown operand type");
|
|
|
|
case MachineOperand::MO_Register:
|
2015-05-14 02:37:00 +08:00
|
|
|
MCOp = MCOperand::createReg(encodeVirtualRegister(MO.getReg()));
|
2013-08-06 22:13:27 +08:00
|
|
|
break;
|
|
|
|
case MachineOperand::MO_Immediate:
|
2015-05-14 02:37:00 +08:00
|
|
|
MCOp = MCOperand::createImm(MO.getImm());
|
2013-08-06 22:13:27 +08:00
|
|
|
break;
|
|
|
|
case MachineOperand::MO_MachineBasicBlock:
|
2015-05-30 09:25:56 +08:00
|
|
|
MCOp = MCOperand::createExpr(MCSymbolRefExpr::create(
|
2013-08-06 22:13:27 +08:00
|
|
|
MO.getMBB()->getSymbol(), OutContext));
|
|
|
|
break;
|
|
|
|
case MachineOperand::MO_ExternalSymbol:
|
2014-04-09 23:39:15 +08:00
|
|
|
MCOp = GetSymbolRef(GetExternalSymbolSymbol(MO.getSymbolName()));
|
2013-08-06 22:13:27 +08:00
|
|
|
break;
|
|
|
|
case MachineOperand::MO_GlobalAddress:
|
2014-04-09 23:39:15 +08:00
|
|
|
MCOp = GetSymbolRef(getSymbol(MO.getGlobal()));
|
2013-08-06 22:13:27 +08:00
|
|
|
break;
|
|
|
|
case MachineOperand::MO_FPImmediate: {
|
|
|
|
const ConstantFP *Cnt = MO.getFPImm();
|
2016-06-08 18:01:20 +08:00
|
|
|
const APFloat &Val = Cnt->getValueAPF();
|
2013-08-06 22:13:27 +08:00
|
|
|
|
|
|
|
switch (Cnt->getType()->getTypeID()) {
|
|
|
|
default: report_fatal_error("Unsupported FP type"); break;
|
2017-01-14 04:56:17 +08:00
|
|
|
case Type::HalfTyID:
|
|
|
|
MCOp = MCOperand::createExpr(
|
|
|
|
NVPTXFloatMCExpr::createConstantFPHalf(Val, OutContext));
|
|
|
|
break;
|
2013-08-06 22:13:27 +08:00
|
|
|
case Type::FloatTyID:
|
2015-05-14 02:37:00 +08:00
|
|
|
MCOp = MCOperand::createExpr(
|
2015-05-30 09:25:56 +08:00
|
|
|
NVPTXFloatMCExpr::createConstantFPSingle(Val, OutContext));
|
2013-08-06 22:13:27 +08:00
|
|
|
break;
|
|
|
|
case Type::DoubleTyID:
|
2015-05-14 02:37:00 +08:00
|
|
|
MCOp = MCOperand::createExpr(
|
2015-05-30 09:25:56 +08:00
|
|
|
NVPTXFloatMCExpr::createConstantFPDouble(Val, OutContext));
|
2013-08-06 22:13:27 +08:00
|
|
|
break;
|
|
|
|
}
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
return true;
|
|
|
|
}
|
|
|
|
|
|
|
|
unsigned NVPTXAsmPrinter::encodeVirtualRegister(unsigned Reg) {
|
2013-08-06 22:13:31 +08:00
|
|
|
if (TargetRegisterInfo::isVirtualRegister(Reg)) {
|
|
|
|
const TargetRegisterClass *RC = MRI->getRegClass(Reg);
|
|
|
|
|
|
|
|
DenseMap<unsigned, unsigned> &RegMap = VRegMapping[RC];
|
|
|
|
unsigned RegNum = RegMap[Reg];
|
|
|
|
|
|
|
|
// Encode the register class in the upper 4 bits
|
|
|
|
// Must be kept in sync with NVPTXInstPrinter::printRegName
|
|
|
|
unsigned Ret = 0;
|
|
|
|
if (RC == &NVPTX::Int1RegsRegClass) {
|
|
|
|
Ret = (1 << 28);
|
|
|
|
} else if (RC == &NVPTX::Int16RegsRegClass) {
|
|
|
|
Ret = (2 << 28);
|
|
|
|
} else if (RC == &NVPTX::Int32RegsRegClass) {
|
|
|
|
Ret = (3 << 28);
|
|
|
|
} else if (RC == &NVPTX::Int64RegsRegClass) {
|
|
|
|
Ret = (4 << 28);
|
|
|
|
} else if (RC == &NVPTX::Float32RegsRegClass) {
|
|
|
|
Ret = (5 << 28);
|
|
|
|
} else if (RC == &NVPTX::Float64RegsRegClass) {
|
|
|
|
Ret = (6 << 28);
|
2017-01-14 04:56:17 +08:00
|
|
|
} else if (RC == &NVPTX::Float16RegsRegClass) {
|
|
|
|
Ret = (7 << 28);
|
2017-02-24 06:38:24 +08:00
|
|
|
} else if (RC == &NVPTX::Float16x2RegsRegClass) {
|
|
|
|
Ret = (8 << 28);
|
2013-08-06 22:13:31 +08:00
|
|
|
} else {
|
|
|
|
report_fatal_error("Bad register class");
|
|
|
|
}
|
|
|
|
|
|
|
|
// Insert the vreg number
|
|
|
|
Ret |= (RegNum & 0x0FFFFFFF);
|
|
|
|
return Ret;
|
2013-08-06 22:13:27 +08:00
|
|
|
} else {
|
2013-08-06 22:13:31 +08:00
|
|
|
// Some special-use registers are actually physical registers.
|
|
|
|
// Encode this as the register class ID of 0 and the real register ID.
|
|
|
|
return Reg & 0x0FFFFFFF;
|
2013-08-06 22:13:27 +08:00
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2014-04-09 23:39:15 +08:00
|
|
|
MCOperand NVPTXAsmPrinter::GetSymbolRef(const MCSymbol *Symbol) {
|
2013-08-06 22:13:27 +08:00
|
|
|
const MCExpr *Expr;
|
2015-05-30 09:25:56 +08:00
|
|
|
Expr = MCSymbolRefExpr::create(Symbol, MCSymbolRefExpr::VK_None,
|
2013-08-07 07:06:42 +08:00
|
|
|
OutContext);
|
2015-05-14 02:37:00 +08:00
|
|
|
return MCOperand::createExpr(Expr);
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
void NVPTXAsmPrinter::printReturnValStr(const Function *F, raw_ostream &O) {
|
2015-07-09 09:57:34 +08:00
|
|
|
const DataLayout &DL = getDataLayout();
|
2015-02-19 08:08:14 +08:00
|
|
|
const TargetLowering *TLI = nvptxSubtarget->getTargetLowering();
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
Type *Ty = F->getReturnType();
|
|
|
|
|
2015-02-19 08:08:14 +08:00
|
|
|
bool isABI = (nvptxSubtarget->getSmVersion() >= 20);
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
if (Ty->getTypeID() == Type::VoidTyID)
|
|
|
|
return;
|
|
|
|
|
|
|
|
O << " (";
|
|
|
|
|
|
|
|
if (isABI) {
|
2013-12-08 03:34:20 +08:00
|
|
|
if (Ty->isFloatingPointTy() || Ty->isIntegerTy()) {
|
2012-05-05 04:18:50 +08:00
|
|
|
unsigned size = 0;
|
2015-08-02 06:20:21 +08:00
|
|
|
if (auto *ITy = dyn_cast<IntegerType>(Ty)) {
|
2012-05-05 04:18:50 +08:00
|
|
|
size = ITy->getBitWidth();
|
|
|
|
} else {
|
2013-03-30 22:29:21 +08:00
|
|
|
assert(Ty->isFloatingPointTy() && "Floating point type expected here");
|
2012-05-05 04:18:50 +08:00
|
|
|
size = Ty->getPrimitiveSizeInBits();
|
|
|
|
}
|
2017-01-14 04:56:17 +08:00
|
|
|
// PTX ABI requires all scalar return values to be at least 32
|
|
|
|
// bits in size. fp16 normally uses .b16 as its storage type in
|
|
|
|
// PTX, so its size must be adjusted here, too.
|
|
|
|
if (size < 32)
|
|
|
|
size = 32;
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
O << ".param .b" << size << " func_retval0";
|
2013-03-30 22:29:21 +08:00
|
|
|
} else if (isa<PointerType>(Ty)) {
|
2015-07-09 10:09:04 +08:00
|
|
|
O << ".param .b" << TLI->getPointerTy(DL).getSizeInBits()
|
2013-03-30 22:29:21 +08:00
|
|
|
<< " func_retval0";
|
2016-07-21 02:39:52 +08:00
|
|
|
} else if (Ty->isAggregateType() || Ty->isVectorTy()) {
|
2015-07-09 09:57:34 +08:00
|
|
|
unsigned totalsz = DL.getTypeAllocSize(Ty);
|
2016-07-21 02:39:52 +08:00
|
|
|
unsigned retAlignment = 0;
|
2017-01-10 06:16:51 +08:00
|
|
|
if (!getAlign(*F, 0, retAlignment))
|
2016-07-21 02:39:52 +08:00
|
|
|
retAlignment = DL.getABITypeAlignment(Ty);
|
|
|
|
O << ".param .align " << retAlignment << " .b8 func_retval0[" << totalsz
|
|
|
|
<< "]";
|
2015-01-05 18:15:49 +08:00
|
|
|
} else
|
|
|
|
llvm_unreachable("Unknown return type");
|
2012-05-05 04:18:50 +08:00
|
|
|
} else {
|
|
|
|
SmallVector<EVT, 16> vtparts;
|
2015-07-09 09:57:34 +08:00
|
|
|
ComputeValueVTs(*TLI, DL, Ty, vtparts);
|
2012-05-05 04:18:50 +08:00
|
|
|
unsigned idx = 0;
|
2013-03-30 22:29:21 +08:00
|
|
|
for (unsigned i = 0, e = vtparts.size(); i != e; ++i) {
|
2012-05-05 04:18:50 +08:00
|
|
|
unsigned elems = 1;
|
|
|
|
EVT elemtype = vtparts[i];
|
|
|
|
if (vtparts[i].isVector()) {
|
|
|
|
elems = vtparts[i].getVectorNumElements();
|
|
|
|
elemtype = vtparts[i].getVectorElementType();
|
|
|
|
}
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
for (unsigned j = 0, je = elems; j != je; ++j) {
|
2012-05-05 04:18:50 +08:00
|
|
|
unsigned sz = elemtype.getSizeInBits();
|
2013-03-30 22:29:21 +08:00
|
|
|
if (elemtype.isInteger() && (sz < 32))
|
|
|
|
sz = 32;
|
2012-05-05 04:18:50 +08:00
|
|
|
O << ".reg .b" << sz << " func_retval" << idx;
|
2013-03-30 22:29:21 +08:00
|
|
|
if (j < je - 1)
|
|
|
|
O << ", ";
|
2012-05-05 04:18:50 +08:00
|
|
|
++idx;
|
|
|
|
}
|
2013-03-30 22:29:21 +08:00
|
|
|
if (i < e - 1)
|
2012-05-05 04:18:50 +08:00
|
|
|
O << ", ";
|
|
|
|
}
|
|
|
|
}
|
|
|
|
O << ") ";
|
|
|
|
}
|
|
|
|
|
|
|
|
void NVPTXAsmPrinter::printReturnValStr(const MachineFunction &MF,
|
|
|
|
raw_ostream &O) {
|
|
|
|
const Function *F = MF.getFunction();
|
|
|
|
printReturnValStr(F, O);
|
|
|
|
}
|
|
|
|
|
2015-02-01 10:27:45 +08:00
|
|
|
// Return true if MBB is the header of a loop marked with
|
|
|
|
// llvm.loop.unroll.disable.
|
2015-02-04 01:57:38 +08:00
|
|
|
// TODO: consider "#pragma unroll 1" which is equivalent to "#pragma nounroll".
|
2015-02-01 10:27:45 +08:00
|
|
|
bool NVPTXAsmPrinter::isLoopHeaderOfNoUnroll(
|
|
|
|
const MachineBasicBlock &MBB) const {
|
|
|
|
MachineLoopInfo &LI = getAnalysis<MachineLoopInfo>();
|
|
|
|
// We insert .pragma "nounroll" only to the loop header.
|
2015-06-02 23:28:27 +08:00
|
|
|
if (!LI.isLoopHeader(&MBB))
|
2015-02-01 10:27:45 +08:00
|
|
|
return false;
|
|
|
|
|
|
|
|
// llvm.loop.unroll.disable is marked on the back edges of a loop. Therefore,
|
|
|
|
// we iterate through each back edge of the loop with header MBB, and check
|
|
|
|
// whether its metadata contains llvm.loop.unroll.disable.
|
|
|
|
for (auto I = MBB.pred_begin(); I != MBB.pred_end(); ++I) {
|
|
|
|
const MachineBasicBlock *PMBB = *I;
|
|
|
|
if (LI.getLoopFor(PMBB) != LI.getLoopFor(&MBB)) {
|
|
|
|
// Edges from other loops to MBB are not back edges.
|
|
|
|
continue;
|
|
|
|
}
|
|
|
|
if (const BasicBlock *PBB = PMBB->getBasicBlock()) {
|
2016-03-25 08:35:38 +08:00
|
|
|
if (MDNode *LoopID =
|
|
|
|
PBB->getTerminator()->getMetadata(LLVMContext::MD_loop)) {
|
2015-02-01 10:27:45 +08:00
|
|
|
if (GetUnrollMetadata(LoopID, "llvm.loop.unroll.disable"))
|
|
|
|
return true;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
}
|
|
|
|
return false;
|
|
|
|
}
|
|
|
|
|
|
|
|
void NVPTXAsmPrinter::EmitBasicBlockStart(const MachineBasicBlock &MBB) const {
|
|
|
|
AsmPrinter::EmitBasicBlockStart(MBB);
|
|
|
|
if (isLoopHeaderOfNoUnroll(MBB))
|
2015-04-25 03:11:51 +08:00
|
|
|
OutStreamer->EmitRawText(StringRef("\t.pragma \"nounroll\";\n"));
|
2015-02-01 10:27:45 +08:00
|
|
|
}
|
|
|
|
|
2012-05-05 04:18:50 +08:00
|
|
|
void NVPTXAsmPrinter::EmitFunctionEntryLabel() {
|
|
|
|
SmallString<128> Str;
|
|
|
|
raw_svector_ostream O(Str);
|
|
|
|
|
2013-05-20 20:13:32 +08:00
|
|
|
if (!GlobalsEmitted) {
|
|
|
|
emitGlobals(*MF->getFunction()->getParent());
|
|
|
|
GlobalsEmitted = true;
|
|
|
|
}
|
|
|
|
|
2012-05-05 04:18:50 +08:00
|
|
|
// Set up
|
|
|
|
MRI = &MF->getRegInfo();
|
|
|
|
F = MF->getFunction();
|
2013-03-30 22:29:21 +08:00
|
|
|
emitLinkageDirective(F, O);
|
2017-01-10 06:16:51 +08:00
|
|
|
if (isKernelFunction(*F))
|
2012-05-05 04:18:50 +08:00
|
|
|
O << ".entry ";
|
|
|
|
else {
|
|
|
|
O << ".func ";
|
|
|
|
printReturnValStr(*MF, O);
|
|
|
|
}
|
|
|
|
|
2015-06-09 08:31:39 +08:00
|
|
|
CurrentFnSym->print(O, MAI);
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
emitFunctionParamList(*MF, O);
|
|
|
|
|
2017-01-10 06:16:51 +08:00
|
|
|
if (isKernelFunction(*F))
|
2012-05-05 04:18:50 +08:00
|
|
|
emitKernelFunctionDirectives(*F, O);
|
|
|
|
|
2015-04-25 03:11:51 +08:00
|
|
|
OutStreamer->EmitRawText(O.str());
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
prevDebugLoc = DebugLoc();
|
|
|
|
}
|
|
|
|
|
|
|
|
void NVPTXAsmPrinter::EmitFunctionBodyStart() {
|
2013-05-31 20:14:49 +08:00
|
|
|
VRegMapping.clear();
|
2015-04-25 03:11:51 +08:00
|
|
|
OutStreamer->EmitRawText(StringRef("{\n"));
|
2012-05-05 04:18:50 +08:00
|
|
|
setAndEmitFunctionVirtualRegisters(*MF);
|
|
|
|
|
|
|
|
SmallString<128> Str;
|
|
|
|
raw_svector_ostream O(Str);
|
|
|
|
emitDemotedVars(MF->getFunction(), O);
|
2015-04-25 03:11:51 +08:00
|
|
|
OutStreamer->EmitRawText(O.str());
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
void NVPTXAsmPrinter::EmitFunctionBodyEnd() {
|
2015-04-25 03:11:51 +08:00
|
|
|
OutStreamer->EmitRawText(StringRef("}\n"));
|
2013-05-31 20:14:49 +08:00
|
|
|
VRegMapping.clear();
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
2013-10-11 20:39:36 +08:00
|
|
|
void NVPTXAsmPrinter::emitImplicitDef(const MachineInstr *MI) const {
|
|
|
|
unsigned RegNo = MI->getOperand(0).getReg();
|
2015-03-25 07:37:10 +08:00
|
|
|
if (TargetRegisterInfo::isVirtualRegister(RegNo)) {
|
2015-04-25 03:11:51 +08:00
|
|
|
OutStreamer->AddComment(Twine("implicit-def: ") +
|
|
|
|
getVirtualRegisterName(RegNo));
|
2013-10-11 20:39:36 +08:00
|
|
|
} else {
|
2015-04-25 03:11:51 +08:00
|
|
|
OutStreamer->AddComment(Twine("implicit-def: ") +
|
|
|
|
nvptxSubtarget->getRegisterInfo()->getName(RegNo));
|
2013-10-11 20:39:36 +08:00
|
|
|
}
|
2015-04-25 03:11:51 +08:00
|
|
|
OutStreamer->AddBlankLine();
|
2013-10-11 20:39:36 +08:00
|
|
|
}
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
void NVPTXAsmPrinter::emitKernelFunctionDirectives(const Function &F,
|
|
|
|
raw_ostream &O) const {
|
2012-05-05 04:18:50 +08:00
|
|
|
// If the NVVM IR has some of reqntid* specified, then output
|
|
|
|
// the reqntid directive, and set the unspecified ones to 1.
|
|
|
|
// If none of reqntid* is specified, don't output reqntid directive.
|
|
|
|
unsigned reqntidx, reqntidy, reqntidz;
|
|
|
|
bool specified = false;
|
2017-01-10 06:16:51 +08:00
|
|
|
if (!getReqNTIDx(F, reqntidx))
|
2013-03-30 22:29:21 +08:00
|
|
|
reqntidx = 1;
|
|
|
|
else
|
|
|
|
specified = true;
|
2017-01-10 06:16:51 +08:00
|
|
|
if (!getReqNTIDy(F, reqntidy))
|
2013-03-30 22:29:21 +08:00
|
|
|
reqntidy = 1;
|
|
|
|
else
|
|
|
|
specified = true;
|
2017-01-10 06:16:51 +08:00
|
|
|
if (!getReqNTIDz(F, reqntidz))
|
2013-03-30 22:29:21 +08:00
|
|
|
reqntidz = 1;
|
|
|
|
else
|
|
|
|
specified = true;
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
if (specified)
|
2013-03-30 22:29:21 +08:00
|
|
|
O << ".reqntid " << reqntidx << ", " << reqntidy << ", " << reqntidz
|
|
|
|
<< "\n";
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
// If the NVVM IR has some of maxntid* specified, then output
|
|
|
|
// the maxntid directive, and set the unspecified ones to 1.
|
|
|
|
// If none of maxntid* is specified, don't output maxntid directive.
|
|
|
|
unsigned maxntidx, maxntidy, maxntidz;
|
|
|
|
specified = false;
|
2017-01-10 06:16:51 +08:00
|
|
|
if (!getMaxNTIDx(F, maxntidx))
|
2013-03-30 22:29:21 +08:00
|
|
|
maxntidx = 1;
|
|
|
|
else
|
|
|
|
specified = true;
|
2017-01-10 06:16:51 +08:00
|
|
|
if (!getMaxNTIDy(F, maxntidy))
|
2013-03-30 22:29:21 +08:00
|
|
|
maxntidy = 1;
|
|
|
|
else
|
|
|
|
specified = true;
|
2017-01-10 06:16:51 +08:00
|
|
|
if (!getMaxNTIDz(F, maxntidz))
|
2013-03-30 22:29:21 +08:00
|
|
|
maxntidz = 1;
|
|
|
|
else
|
|
|
|
specified = true;
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
if (specified)
|
2013-03-30 22:29:21 +08:00
|
|
|
O << ".maxntid " << maxntidx << ", " << maxntidy << ", " << maxntidz
|
|
|
|
<< "\n";
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
unsigned mincta;
|
2017-01-10 06:16:51 +08:00
|
|
|
if (getMinCTASm(F, mincta))
|
2012-05-05 04:18:50 +08:00
|
|
|
O << ".minnctapersm " << mincta << "\n";
|
2016-12-15 06:32:50 +08:00
|
|
|
|
|
|
|
unsigned maxnreg;
|
2017-01-10 06:16:51 +08:00
|
|
|
if (getMaxNReg(F, maxnreg))
|
2016-12-15 06:32:50 +08:00
|
|
|
O << ".maxnreg " << maxnreg << "\n";
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
2013-10-11 20:39:36 +08:00
|
|
|
std::string
|
|
|
|
NVPTXAsmPrinter::getVirtualRegisterName(unsigned Reg) const {
|
|
|
|
const TargetRegisterClass *RC = MRI->getRegClass(Reg);
|
|
|
|
|
|
|
|
std::string Name;
|
|
|
|
raw_string_ostream NameStr(Name);
|
2012-05-05 04:18:50 +08:00
|
|
|
|
2013-10-11 20:39:36 +08:00
|
|
|
VRegRCMap::const_iterator I = VRegMapping.find(RC);
|
|
|
|
assert(I != VRegMapping.end() && "Bad register class");
|
|
|
|
const DenseMap<unsigned, unsigned> &RegMap = I->second;
|
2012-05-05 04:18:50 +08:00
|
|
|
|
2013-10-11 20:39:36 +08:00
|
|
|
VRegMap::const_iterator VI = RegMap.find(Reg);
|
|
|
|
assert(VI != RegMap.end() && "Bad virtual register");
|
|
|
|
unsigned MappedVR = VI->second;
|
|
|
|
|
|
|
|
NameStr << getNVPTXRegClassStr(RC) << MappedVR;
|
|
|
|
|
|
|
|
NameStr.flush();
|
|
|
|
return Name;
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
2013-10-11 20:39:36 +08:00
|
|
|
void NVPTXAsmPrinter::emitVirtualRegister(unsigned int vr,
|
2013-03-30 22:29:21 +08:00
|
|
|
raw_ostream &O) {
|
2013-10-11 20:39:36 +08:00
|
|
|
O << getVirtualRegisterName(vr);
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
void NVPTXAsmPrinter::printVecModifiedImmediate(
|
|
|
|
const MachineOperand &MO, const char *Modifier, raw_ostream &O) {
|
|
|
|
static const char vecelem[] = { '0', '1', '2', '3', '0', '1', '2', '3' };
|
|
|
|
int Imm = (int) MO.getImm();
|
|
|
|
if (0 == strcmp(Modifier, "vecelem"))
|
2012-05-05 04:18:50 +08:00
|
|
|
O << "_" << vecelem[Imm];
|
2013-03-30 22:29:21 +08:00
|
|
|
else if (0 == strcmp(Modifier, "vecv4comm1")) {
|
|
|
|
if ((Imm < 0) || (Imm > 3))
|
2012-05-05 04:18:50 +08:00
|
|
|
O << "//";
|
2013-03-30 22:29:21 +08:00
|
|
|
} else if (0 == strcmp(Modifier, "vecv4comm2")) {
|
|
|
|
if ((Imm < 4) || (Imm > 7))
|
2012-05-05 04:18:50 +08:00
|
|
|
O << "//";
|
2013-03-30 22:29:21 +08:00
|
|
|
} else if (0 == strcmp(Modifier, "vecv4pos")) {
|
|
|
|
if (Imm < 0)
|
|
|
|
Imm = 0;
|
|
|
|
O << "_" << vecelem[Imm % 4];
|
|
|
|
} else if (0 == strcmp(Modifier, "vecv2comm1")) {
|
|
|
|
if ((Imm < 0) || (Imm > 1))
|
2012-05-05 04:18:50 +08:00
|
|
|
O << "//";
|
2013-03-30 22:29:21 +08:00
|
|
|
} else if (0 == strcmp(Modifier, "vecv2comm2")) {
|
|
|
|
if ((Imm < 2) || (Imm > 3))
|
2012-05-05 04:18:50 +08:00
|
|
|
O << "//";
|
2013-03-30 22:29:21 +08:00
|
|
|
} else if (0 == strcmp(Modifier, "vecv2pos")) {
|
|
|
|
if (Imm < 0)
|
|
|
|
Imm = 0;
|
|
|
|
O << "_" << vecelem[Imm % 2];
|
|
|
|
} else
|
2012-05-24 15:02:50 +08:00
|
|
|
llvm_unreachable("Unknown Modifier on immediate operand");
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
void NVPTXAsmPrinter::emitDeclaration(const Function *F, raw_ostream &O) {
|
|
|
|
emitLinkageDirective(F, O);
|
2017-01-10 06:16:51 +08:00
|
|
|
if (isKernelFunction(*F))
|
2012-05-05 04:18:50 +08:00
|
|
|
O << ".entry ";
|
|
|
|
else
|
|
|
|
O << ".func ";
|
|
|
|
printReturnValStr(F, O);
|
2015-06-09 08:31:39 +08:00
|
|
|
getSymbol(F)->print(O, MAI);
|
|
|
|
O << "\n";
|
2012-05-05 04:18:50 +08:00
|
|
|
emitFunctionParamList(F, O);
|
|
|
|
O << ";\n";
|
|
|
|
}
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
static bool usedInGlobalVarDef(const Constant *C) {
|
2012-05-05 04:18:50 +08:00
|
|
|
if (!C)
|
|
|
|
return false;
|
|
|
|
|
|
|
|
if (const GlobalVariable *GV = dyn_cast<GlobalVariable>(C)) {
|
Refactor: Simplify boolean conditional return statements in lib/Target/NVPTX
Summary: Use clang-tidy to simplify boolean conditional return statements
Reviewers: rafael, echristo, chandlerc, bkramer, craig.topper, dexonsmith, chapuni, eliben, jingyue, jholewinski
Subscribers: llvm-commits, jholewinski
Differential Revision: http://reviews.llvm.org/D9983
llvm-svn: 243734
2015-07-31 13:09:47 +08:00
|
|
|
return GV->getName() != "llvm.used";
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
2014-03-09 11:16:01 +08:00
|
|
|
for (const User *U : C->users())
|
|
|
|
if (const Constant *C = dyn_cast<Constant>(U))
|
|
|
|
if (usedInGlobalVarDef(C))
|
|
|
|
return true;
|
|
|
|
|
2012-05-05 04:18:50 +08:00
|
|
|
return false;
|
|
|
|
}
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
static bool usedInOneFunc(const User *U, Function const *&oneFunc) {
|
2012-05-05 04:18:50 +08:00
|
|
|
if (const GlobalVariable *othergv = dyn_cast<GlobalVariable>(U)) {
|
2015-03-30 23:42:36 +08:00
|
|
|
if (othergv->getName() == "llvm.used")
|
2012-05-05 04:18:50 +08:00
|
|
|
return true;
|
|
|
|
}
|
|
|
|
|
|
|
|
if (const Instruction *instr = dyn_cast<Instruction>(U)) {
|
|
|
|
if (instr->getParent() && instr->getParent()->getParent()) {
|
|
|
|
const Function *curFunc = instr->getParent()->getParent();
|
|
|
|
if (oneFunc && (curFunc != oneFunc))
|
|
|
|
return false;
|
|
|
|
oneFunc = curFunc;
|
|
|
|
return true;
|
2013-03-30 22:29:21 +08:00
|
|
|
} else
|
2012-05-05 04:18:50 +08:00
|
|
|
return false;
|
|
|
|
}
|
|
|
|
|
2014-03-09 11:16:01 +08:00
|
|
|
for (const User *UU : U->users())
|
2015-03-24 00:26:23 +08:00
|
|
|
if (!usedInOneFunc(UU, oneFunc))
|
2012-05-05 04:18:50 +08:00
|
|
|
return false;
|
2014-03-09 11:16:01 +08:00
|
|
|
|
2012-05-05 04:18:50 +08:00
|
|
|
return true;
|
|
|
|
}
|
|
|
|
|
|
|
|
/* Find out if a global variable can be demoted to local scope.
|
|
|
|
* Currently, this is valid for CUDA shared variables, which have local
|
|
|
|
* scope and global lifetime. So the conditions to check are :
|
|
|
|
* 1. Is the global variable in shared address space?
|
|
|
|
* 2. Does it have internal linkage?
|
|
|
|
* 3. Is the global variable referenced only in one function?
|
|
|
|
*/
|
|
|
|
static bool canDemoteGlobalVar(const GlobalVariable *gv, Function const *&f) {
|
2015-03-24 00:26:23 +08:00
|
|
|
if (!gv->hasInternalLinkage())
|
2012-05-05 04:18:50 +08:00
|
|
|
return false;
|
2015-08-02 06:20:21 +08:00
|
|
|
PointerType *Pty = gv->getType();
|
2017-01-10 06:16:51 +08:00
|
|
|
if (Pty->getAddressSpace() != ADDRESS_SPACE_SHARED)
|
2012-05-05 04:18:50 +08:00
|
|
|
return false;
|
|
|
|
|
2014-04-25 13:30:21 +08:00
|
|
|
const Function *oneFunc = nullptr;
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
bool flag = usedInOneFunc(gv, oneFunc);
|
2015-03-24 00:26:23 +08:00
|
|
|
if (!flag)
|
2012-05-05 04:18:50 +08:00
|
|
|
return false;
|
|
|
|
if (!oneFunc)
|
|
|
|
return false;
|
|
|
|
f = oneFunc;
|
|
|
|
return true;
|
|
|
|
}
|
|
|
|
|
|
|
|
static bool useFuncSeen(const Constant *C,
|
2017-01-10 06:16:51 +08:00
|
|
|
DenseMap<const Function *, bool> &seenMap) {
|
2014-03-09 11:16:01 +08:00
|
|
|
for (const User *U : C->users()) {
|
|
|
|
if (const Constant *cu = dyn_cast<Constant>(U)) {
|
2012-05-05 04:18:50 +08:00
|
|
|
if (useFuncSeen(cu, seenMap))
|
|
|
|
return true;
|
2014-03-09 11:16:01 +08:00
|
|
|
} else if (const Instruction *I = dyn_cast<Instruction>(U)) {
|
2012-05-05 04:18:50 +08:00
|
|
|
const BasicBlock *bb = I->getParent();
|
2013-03-30 22:29:21 +08:00
|
|
|
if (!bb)
|
|
|
|
continue;
|
2012-05-05 04:18:50 +08:00
|
|
|
const Function *caller = bb->getParent();
|
2013-03-30 22:29:21 +08:00
|
|
|
if (!caller)
|
|
|
|
continue;
|
2012-05-05 04:18:50 +08:00
|
|
|
if (seenMap.find(caller) != seenMap.end())
|
|
|
|
return true;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
return false;
|
|
|
|
}
|
|
|
|
|
2013-05-20 20:13:32 +08:00
|
|
|
void NVPTXAsmPrinter::emitDeclarations(const Module &M, raw_ostream &O) {
|
2017-01-10 06:16:51 +08:00
|
|
|
DenseMap<const Function *, bool> seenMap;
|
2013-03-30 22:29:21 +08:00
|
|
|
for (Module::const_iterator FI = M.begin(), FE = M.end(); FI != FE; ++FI) {
|
2015-10-20 08:54:09 +08:00
|
|
|
const Function *F = &*FI;
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
if (F->isDeclaration()) {
|
|
|
|
if (F->use_empty())
|
|
|
|
continue;
|
|
|
|
if (F->getIntrinsicID())
|
|
|
|
continue;
|
|
|
|
emitDeclaration(F, O);
|
|
|
|
continue;
|
|
|
|
}
|
2014-03-09 11:16:01 +08:00
|
|
|
for (const User *U : F->users()) {
|
|
|
|
if (const Constant *C = dyn_cast<Constant>(U)) {
|
2012-05-05 04:18:50 +08:00
|
|
|
if (usedInGlobalVarDef(C)) {
|
|
|
|
// The use is in the initialization of a global variable
|
|
|
|
// that is a function pointer, so print a declaration
|
|
|
|
// for the original function
|
|
|
|
emitDeclaration(F, O);
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
// Emit a declaration of this function if the function that
|
|
|
|
// uses this constant expr has already been seen.
|
|
|
|
if (useFuncSeen(C, seenMap)) {
|
|
|
|
emitDeclaration(F, O);
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2014-03-09 11:16:01 +08:00
|
|
|
if (!isa<Instruction>(U))
|
2013-03-30 22:29:21 +08:00
|
|
|
continue;
|
2014-03-09 11:16:01 +08:00
|
|
|
const Instruction *instr = cast<Instruction>(U);
|
2012-05-05 04:18:50 +08:00
|
|
|
const BasicBlock *bb = instr->getParent();
|
2013-03-30 22:29:21 +08:00
|
|
|
if (!bb)
|
|
|
|
continue;
|
2012-05-05 04:18:50 +08:00
|
|
|
const Function *caller = bb->getParent();
|
2013-03-30 22:29:21 +08:00
|
|
|
if (!caller)
|
|
|
|
continue;
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
// If a caller has already been seen, then the caller is
|
|
|
|
// appearing in the module before the callee. so print out
|
|
|
|
// a declaration for the callee.
|
|
|
|
if (seenMap.find(caller) != seenMap.end()) {
|
|
|
|
emitDeclaration(F, O);
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
seenMap[F] = true;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
void NVPTXAsmPrinter::recordAndEmitFilenames(Module &M) {
|
|
|
|
DebugInfoFinder DbgFinder;
|
|
|
|
DbgFinder.processModule(M);
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
unsigned i = 1;
|
2015-04-30 00:38:44 +08:00
|
|
|
for (const DICompileUnit *DIUnit : DbgFinder.compile_units()) {
|
2015-04-16 07:19:27 +08:00
|
|
|
StringRef Filename = DIUnit->getFilename();
|
|
|
|
StringRef Dirname = DIUnit->getDirectory();
|
2012-05-05 04:18:50 +08:00
|
|
|
SmallString<128> FullPathName = Dirname;
|
|
|
|
if (!Dirname.empty() && !sys::path::is_absolute(Filename)) {
|
|
|
|
sys::path::append(FullPathName, Filename);
|
2015-03-30 23:42:36 +08:00
|
|
|
Filename = FullPathName;
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
2015-03-30 23:42:36 +08:00
|
|
|
if (filenameMap.find(Filename) != filenameMap.end())
|
2012-05-05 04:18:50 +08:00
|
|
|
continue;
|
2015-03-30 23:42:36 +08:00
|
|
|
filenameMap[Filename] = i;
|
2015-04-25 03:11:51 +08:00
|
|
|
OutStreamer->EmitDwarfFileDirective(i, "", Filename);
|
2012-05-05 04:18:50 +08:00
|
|
|
++i;
|
|
|
|
}
|
|
|
|
|
2015-04-30 00:38:44 +08:00
|
|
|
for (DISubprogram *SP : DbgFinder.subprograms()) {
|
2015-04-14 11:40:37 +08:00
|
|
|
StringRef Filename = SP->getFilename();
|
|
|
|
StringRef Dirname = SP->getDirectory();
|
2012-05-05 04:18:50 +08:00
|
|
|
SmallString<128> FullPathName = Dirname;
|
|
|
|
if (!Dirname.empty() && !sys::path::is_absolute(Filename)) {
|
|
|
|
sys::path::append(FullPathName, Filename);
|
2015-03-30 23:42:36 +08:00
|
|
|
Filename = FullPathName;
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
2015-03-30 23:42:36 +08:00
|
|
|
if (filenameMap.find(Filename) != filenameMap.end())
|
2012-05-05 04:18:50 +08:00
|
|
|
continue;
|
2015-03-30 23:42:36 +08:00
|
|
|
filenameMap[Filename] = i;
|
2016-02-12 02:21:47 +08:00
|
|
|
OutStreamer->EmitDwarfFileDirective(i, "", Filename);
|
2012-05-05 04:18:50 +08:00
|
|
|
++i;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2016-01-30 09:07:38 +08:00
|
|
|
static bool isEmptyXXStructor(GlobalVariable *GV) {
|
|
|
|
if (!GV) return true;
|
|
|
|
const ConstantArray *InitList = dyn_cast<ConstantArray>(GV->getInitializer());
|
|
|
|
if (!InitList) return true; // Not an array; we don't know how to parse.
|
|
|
|
return InitList->getNumOperands() == 0;
|
|
|
|
}
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
bool NVPTXAsmPrinter::doInitialization(Module &M) {
|
2015-02-19 08:08:14 +08:00
|
|
|
// Construct a default subtarget off of the TargetMachine defaults. The
|
|
|
|
// rest of NVPTX isn't friendly to change subtargets per function and
|
|
|
|
// so the default TargetMachine will have all of the options.
|
2015-06-16 23:44:21 +08:00
|
|
|
const Triple &TT = TM.getTargetTriple();
|
2015-02-19 08:08:14 +08:00
|
|
|
StringRef CPU = TM.getTargetCPU();
|
|
|
|
StringRef FS = TM.getTargetFeatureString();
|
|
|
|
const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
|
2015-02-19 08:08:27 +08:00
|
|
|
const NVPTXSubtarget STI(TT, CPU, FS, NTM);
|
2012-05-05 04:18:50 +08:00
|
|
|
|
2016-01-24 05:12:20 +08:00
|
|
|
if (M.alias_size()) {
|
|
|
|
report_fatal_error("Module has aliases, which NVPTX does not support.");
|
|
|
|
return true; // error
|
|
|
|
}
|
2016-01-30 09:07:38 +08:00
|
|
|
if (!isEmptyXXStructor(M.getNamedGlobal("llvm.global_ctors"))) {
|
|
|
|
report_fatal_error(
|
|
|
|
"Module has a nontrivial global ctor, which NVPTX does not support.");
|
|
|
|
return true; // error
|
|
|
|
}
|
|
|
|
if (!isEmptyXXStructor(M.getNamedGlobal("llvm.global_dtors"))) {
|
|
|
|
report_fatal_error(
|
|
|
|
"Module has a nontrivial global dtor, which NVPTX does not support.");
|
|
|
|
return true; // error
|
|
|
|
}
|
2016-01-24 05:12:20 +08:00
|
|
|
|
2012-05-05 04:18:50 +08:00
|
|
|
SmallString<128> Str1;
|
|
|
|
raw_svector_ostream OS1(Str1);
|
|
|
|
|
|
|
|
MMI = getAnalysisIfAvailable<MachineModuleInfo>();
|
|
|
|
|
|
|
|
// We need to call the parent's one explicitly.
|
|
|
|
//bool Result = AsmPrinter::doInitialization(M);
|
|
|
|
|
2016-09-29 10:03:47 +08:00
|
|
|
// Initialize TargetLoweringObjectFile since we didn't do in
|
|
|
|
// AsmPrinter::doInitialization either right above or where it's commented out
|
|
|
|
// below.
|
2013-03-30 22:29:21 +08:00
|
|
|
const_cast<TargetLoweringObjectFile &>(getObjFileLowering())
|
|
|
|
.Initialize(OutContext, TM);
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
// Emit header before any dwarf directives are emitted below.
|
2015-02-19 08:08:14 +08:00
|
|
|
emitHeader(M, OS1, STI);
|
2015-04-25 03:11:51 +08:00
|
|
|
OutStreamer->EmitRawText(OS1.str());
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
// Already commented out
|
|
|
|
//bool Result = AsmPrinter::doInitialization(M);
|
|
|
|
|
2013-07-01 21:00:14 +08:00
|
|
|
// Emit module-level inline asm if it exists.
|
|
|
|
if (!M.getModuleInlineAsm().empty()) {
|
2015-04-25 03:11:51 +08:00
|
|
|
OutStreamer->AddComment("Start of file scope inline assembly");
|
|
|
|
OutStreamer->AddBlankLine();
|
|
|
|
OutStreamer->EmitRawText(StringRef(M.getModuleInlineAsm()));
|
|
|
|
OutStreamer->AddBlankLine();
|
|
|
|
OutStreamer->AddComment("End of file scope inline assembly");
|
|
|
|
OutStreamer->AddBlankLine();
|
2013-07-01 21:00:14 +08:00
|
|
|
}
|
|
|
|
|
2015-02-19 08:08:14 +08:00
|
|
|
// If we're not NVCL we're CUDA, go ahead and emit filenames.
|
2015-06-16 23:44:21 +08:00
|
|
|
if (TM.getTargetTriple().getOS() != Triple::NVCL)
|
2012-05-05 04:18:50 +08:00
|
|
|
recordAndEmitFilenames(M);
|
|
|
|
|
2013-05-20 20:13:32 +08:00
|
|
|
GlobalsEmitted = false;
|
|
|
|
|
|
|
|
return false; // success
|
|
|
|
}
|
|
|
|
|
|
|
|
void NVPTXAsmPrinter::emitGlobals(const Module &M) {
|
2012-05-05 04:18:50 +08:00
|
|
|
SmallString<128> Str2;
|
|
|
|
raw_svector_ostream OS2(Str2);
|
|
|
|
|
|
|
|
emitDeclarations(M, OS2);
|
|
|
|
|
2012-11-17 05:03:51 +08:00
|
|
|
// As ptxas does not support forward references of globals, we need to first
|
|
|
|
// sort the list of module-level globals in def-use order. We visit each
|
|
|
|
// global variable in order, and ensure that we emit it *after* its dependent
|
|
|
|
// globals. We use a little extra memory maintaining both a set and a list to
|
|
|
|
// have fast searches while maintaining a strict ordering.
|
2013-05-20 20:13:32 +08:00
|
|
|
SmallVector<const GlobalVariable *, 8> Globals;
|
|
|
|
DenseSet<const GlobalVariable *> GVVisited;
|
|
|
|
DenseSet<const GlobalVariable *> GVVisiting;
|
2012-11-17 05:03:51 +08:00
|
|
|
|
|
|
|
// Visit each global variable, in order
|
2015-10-20 08:54:09 +08:00
|
|
|
for (const GlobalVariable &I : M.globals())
|
|
|
|
VisitGlobalVariableForEmission(&I, Globals, GVVisited, GVVisiting);
|
2012-11-17 05:03:51 +08:00
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
assert(GVVisited.size() == M.getGlobalList().size() &&
|
2012-11-17 05:03:51 +08:00
|
|
|
"Missed a global variable");
|
|
|
|
assert(GVVisiting.size() == 0 && "Did not fully process a global variable");
|
|
|
|
|
|
|
|
// Print out module-level global variables in proper order
|
|
|
|
for (unsigned i = 0, e = Globals.size(); i != e; ++i)
|
|
|
|
printModuleLevelGV(Globals[i], OS2);
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
OS2 << '\n';
|
|
|
|
|
2015-04-25 03:11:51 +08:00
|
|
|
OutStreamer->EmitRawText(OS2.str());
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
2015-02-19 08:08:14 +08:00
|
|
|
void NVPTXAsmPrinter::emitHeader(Module &M, raw_ostream &O,
|
|
|
|
const NVPTXSubtarget &STI) {
|
2012-05-05 04:18:50 +08:00
|
|
|
O << "//\n";
|
|
|
|
O << "// Generated by LLVM NVPTX Back-End\n";
|
|
|
|
O << "//\n";
|
|
|
|
O << "\n";
|
|
|
|
|
2015-02-19 08:08:14 +08:00
|
|
|
unsigned PTXVersion = STI.getPTXVersion();
|
2012-11-12 11:16:43 +08:00
|
|
|
O << ".version " << (PTXVersion / 10) << "." << (PTXVersion % 10) << "\n";
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
O << ".target ";
|
2015-02-19 08:08:14 +08:00
|
|
|
O << STI.getTargetName();
|
2012-05-05 04:18:50 +08:00
|
|
|
|
2015-02-19 08:22:47 +08:00
|
|
|
const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
|
|
|
|
if (NTM.getDrvInterface() == NVPTX::NVCL)
|
2012-05-05 04:18:50 +08:00
|
|
|
O << ", texmode_independent";
|
2015-02-19 08:08:23 +08:00
|
|
|
else {
|
2015-02-19 08:08:14 +08:00
|
|
|
if (!STI.hasDouble())
|
2012-05-05 04:18:50 +08:00
|
|
|
O << ", map_f64_to_f32";
|
|
|
|
}
|
|
|
|
|
|
|
|
if (MAI->doesSupportDebugInformation())
|
|
|
|
O << ", debug";
|
|
|
|
|
|
|
|
O << "\n";
|
|
|
|
|
|
|
|
O << ".address_size ";
|
2015-02-19 08:22:47 +08:00
|
|
|
if (NTM.is64Bit())
|
2012-05-05 04:18:50 +08:00
|
|
|
O << "64";
|
|
|
|
else
|
|
|
|
O << "32";
|
|
|
|
O << "\n";
|
|
|
|
|
|
|
|
O << "\n";
|
|
|
|
}
|
|
|
|
|
|
|
|
bool NVPTXAsmPrinter::doFinalization(Module &M) {
|
2013-05-20 20:13:32 +08:00
|
|
|
// If we did not emit any functions, then the global declarations have not
|
|
|
|
// yet been emitted.
|
|
|
|
if (!GlobalsEmitted) {
|
|
|
|
emitGlobals(M);
|
|
|
|
GlobalsEmitted = true;
|
|
|
|
}
|
|
|
|
|
2012-05-05 04:18:50 +08:00
|
|
|
// XXX Temproarily remove global variables so that doFinalization() will not
|
|
|
|
// emit them again (global variables are emitted at beginning).
|
|
|
|
|
|
|
|
Module::GlobalListType &global_list = M.getGlobalList();
|
|
|
|
int i, n = global_list.size();
|
2014-08-26 10:03:35 +08:00
|
|
|
GlobalVariable **gv_array = new GlobalVariable *[n];
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
// first, back-up GlobalVariable in gv_array
|
|
|
|
i = 0;
|
|
|
|
for (Module::global_iterator I = global_list.begin(), E = global_list.end();
|
2013-03-30 22:29:21 +08:00
|
|
|
I != E; ++I)
|
2012-05-05 04:18:50 +08:00
|
|
|
gv_array[i++] = &*I;
|
|
|
|
|
|
|
|
// second, empty global_list
|
|
|
|
while (!global_list.empty())
|
|
|
|
global_list.remove(global_list.begin());
|
|
|
|
|
|
|
|
// call doFinalization
|
|
|
|
bool ret = AsmPrinter::doFinalization(M);
|
|
|
|
|
|
|
|
// now we restore global variables
|
2013-03-30 22:29:21 +08:00
|
|
|
for (i = 0; i < n; i++)
|
2012-05-05 04:18:50 +08:00
|
|
|
global_list.insert(global_list.end(), gv_array[i]);
|
|
|
|
|
2014-04-09 23:38:52 +08:00
|
|
|
clearAnnotationCache(&M);
|
2014-08-26 10:03:35 +08:00
|
|
|
|
|
|
|
delete[] gv_array;
|
2012-05-05 04:18:50 +08:00
|
|
|
return ret;
|
|
|
|
|
|
|
|
//bool Result = AsmPrinter::doFinalization(M);
|
|
|
|
// Instead of calling the parents doFinalization, we may
|
|
|
|
// clone parents doFinalization and customize here.
|
|
|
|
// Currently, we if NVISA out the EmitGlobals() in
|
|
|
|
// parent's doFinalization, which is too intrusive.
|
|
|
|
//
|
|
|
|
// Same for the doInitialization.
|
|
|
|
//return Result;
|
|
|
|
}
|
|
|
|
|
|
|
|
// This function emits appropriate linkage directives for
|
|
|
|
// functions and global variables.
|
|
|
|
//
|
|
|
|
// extern function declaration -> .extern
|
|
|
|
// extern function definition -> .visible
|
|
|
|
// external global variable with init -> .visible
|
|
|
|
// external without init -> .extern
|
|
|
|
// appending -> not allowed, assert.
|
2014-06-28 02:35:10 +08:00
|
|
|
// for any linkage other than
|
|
|
|
// internal, private, linker_private,
|
|
|
|
// linker_private_weak, linker_private_weak_def_auto,
|
|
|
|
// we emit -> .weak.
|
2012-05-05 04:18:50 +08:00
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
void NVPTXAsmPrinter::emitLinkageDirective(const GlobalValue *V,
|
|
|
|
raw_ostream &O) {
|
2015-02-19 08:08:23 +08:00
|
|
|
if (static_cast<NVPTXTargetMachine &>(TM).getDrvInterface() == NVPTX::CUDA) {
|
2012-05-05 04:18:50 +08:00
|
|
|
if (V->hasExternalLinkage()) {
|
|
|
|
if (isa<GlobalVariable>(V)) {
|
|
|
|
const GlobalVariable *GVar = cast<GlobalVariable>(V);
|
|
|
|
if (GVar) {
|
|
|
|
if (GVar->hasInitializer())
|
|
|
|
O << ".visible ";
|
|
|
|
else
|
|
|
|
O << ".extern ";
|
|
|
|
}
|
|
|
|
} else if (V->isDeclaration())
|
|
|
|
O << ".extern ";
|
|
|
|
else
|
|
|
|
O << ".visible ";
|
|
|
|
} else if (V->hasAppendingLinkage()) {
|
|
|
|
std::string msg;
|
|
|
|
msg.append("Error: ");
|
|
|
|
msg.append("Symbol ");
|
|
|
|
if (V->hasName())
|
2015-03-30 23:42:36 +08:00
|
|
|
msg.append(V->getName());
|
2012-05-05 04:18:50 +08:00
|
|
|
msg.append("has unsupported appending linkage type");
|
|
|
|
llvm_unreachable(msg.c_str());
|
2014-06-28 02:35:10 +08:00
|
|
|
} else if (!V->hasInternalLinkage() &&
|
|
|
|
!V->hasPrivateLinkage()) {
|
|
|
|
O << ".weak ";
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2013-05-20 20:13:32 +08:00
|
|
|
void NVPTXAsmPrinter::printModuleLevelGV(const GlobalVariable *GVar,
|
|
|
|
raw_ostream &O,
|
2012-05-05 04:18:50 +08:00
|
|
|
bool processDemoted) {
|
|
|
|
// Skip meta data
|
|
|
|
if (GVar->hasSection()) {
|
2016-05-12 02:21:59 +08:00
|
|
|
if (GVar->getSection() == "llvm.metadata")
|
2012-05-05 04:18:50 +08:00
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2014-06-28 02:35:53 +08:00
|
|
|
// Skip LLVM intrinsic global variables
|
|
|
|
if (GVar->getName().startswith("llvm.") ||
|
|
|
|
GVar->getName().startswith("nvvm."))
|
|
|
|
return;
|
|
|
|
|
2015-07-16 14:11:10 +08:00
|
|
|
const DataLayout &DL = getDataLayout();
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
// GlobalVariables are always constant pointers themselves.
|
2015-08-02 06:20:21 +08:00
|
|
|
PointerType *PTy = GVar->getType();
|
2016-01-17 04:30:46 +08:00
|
|
|
Type *ETy = GVar->getValueType();
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
if (GVar->hasExternalLinkage()) {
|
|
|
|
if (GVar->hasInitializer())
|
|
|
|
O << ".visible ";
|
|
|
|
else
|
|
|
|
O << ".extern ";
|
2014-06-28 02:35:56 +08:00
|
|
|
} else if (GVar->hasLinkOnceLinkage() || GVar->hasWeakLinkage() ||
|
|
|
|
GVar->hasAvailableExternallyLinkage() ||
|
|
|
|
GVar->hasCommonLinkage()) {
|
|
|
|
O << ".weak ";
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
2017-01-10 06:16:51 +08:00
|
|
|
if (isTexture(*GVar)) {
|
|
|
|
O << ".global .texref " << getTextureName(*GVar) << ";\n";
|
2012-05-05 04:18:50 +08:00
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2017-01-10 06:16:51 +08:00
|
|
|
if (isSurface(*GVar)) {
|
|
|
|
O << ".global .surfref " << getSurfaceName(*GVar) << ";\n";
|
2012-05-05 04:18:50 +08:00
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
|
|
|
if (GVar->isDeclaration()) {
|
|
|
|
// (extern) declarations, no definition or initializer
|
|
|
|
// Currently the only known declaration is for an automatic __local
|
|
|
|
// (.shared) promoted to global.
|
|
|
|
emitPTXGlobalVariable(GVar, O);
|
|
|
|
O << ";\n";
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2017-01-10 06:16:51 +08:00
|
|
|
if (isSampler(*GVar)) {
|
|
|
|
O << ".global .samplerref " << getSamplerName(*GVar);
|
2012-05-05 04:18:50 +08:00
|
|
|
|
2014-04-25 13:30:21 +08:00
|
|
|
const Constant *Initializer = nullptr;
|
2012-05-05 04:18:50 +08:00
|
|
|
if (GVar->hasInitializer())
|
|
|
|
Initializer = GVar->getInitializer();
|
2014-04-25 13:30:21 +08:00
|
|
|
const ConstantInt *CI = nullptr;
|
2012-05-05 04:18:50 +08:00
|
|
|
if (Initializer)
|
|
|
|
CI = dyn_cast<ConstantInt>(Initializer);
|
|
|
|
if (CI) {
|
2013-03-30 22:29:21 +08:00
|
|
|
unsigned sample = CI->getZExtValue();
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
O << " = { ";
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
for (int i = 0,
|
|
|
|
addr = ((sample & __CLK_ADDRESS_MASK) >> __CLK_ADDRESS_BASE);
|
|
|
|
i < 3; i++) {
|
2012-05-05 04:18:50 +08:00
|
|
|
O << "addr_mode_" << i << " = ";
|
|
|
|
switch (addr) {
|
2013-03-30 22:29:21 +08:00
|
|
|
case 0:
|
|
|
|
O << "wrap";
|
|
|
|
break;
|
|
|
|
case 1:
|
|
|
|
O << "clamp_to_border";
|
|
|
|
break;
|
|
|
|
case 2:
|
|
|
|
O << "clamp_to_edge";
|
|
|
|
break;
|
|
|
|
case 3:
|
|
|
|
O << "wrap";
|
|
|
|
break;
|
|
|
|
case 4:
|
|
|
|
O << "mirror";
|
|
|
|
break;
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
2013-03-30 22:29:21 +08:00
|
|
|
O << ", ";
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
O << "filter_mode = ";
|
2013-03-30 22:29:21 +08:00
|
|
|
switch ((sample & __CLK_FILTER_MASK) >> __CLK_FILTER_BASE) {
|
|
|
|
case 0:
|
|
|
|
O << "nearest";
|
|
|
|
break;
|
|
|
|
case 1:
|
|
|
|
O << "linear";
|
|
|
|
break;
|
|
|
|
case 2:
|
2014-06-18 13:05:13 +08:00
|
|
|
llvm_unreachable("Anisotropic filtering is not supported");
|
2013-03-30 22:29:21 +08:00
|
|
|
default:
|
|
|
|
O << "nearest";
|
|
|
|
break;
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
2013-03-30 22:29:21 +08:00
|
|
|
if (!((sample & __CLK_NORMALIZED_MASK) >> __CLK_NORMALIZED_BASE)) {
|
2012-05-05 04:18:50 +08:00
|
|
|
O << ", force_unnormalized_coords = 1";
|
|
|
|
}
|
|
|
|
O << " }";
|
|
|
|
}
|
|
|
|
|
|
|
|
O << ";\n";
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
|
|
|
if (GVar->hasPrivateLinkage()) {
|
2017-01-10 06:16:51 +08:00
|
|
|
if (strncmp(GVar->getName().data(), "unrollpragma", 12) == 0)
|
2012-05-05 04:18:50 +08:00
|
|
|
return;
|
|
|
|
|
|
|
|
// FIXME - need better way (e.g. Metadata) to avoid generating this global
|
2017-01-10 06:16:51 +08:00
|
|
|
if (strncmp(GVar->getName().data(), "filename", 8) == 0)
|
2012-05-05 04:18:50 +08:00
|
|
|
return;
|
|
|
|
if (GVar->use_empty())
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2014-04-25 13:30:21 +08:00
|
|
|
const Function *demotedFunc = nullptr;
|
2012-05-05 04:18:50 +08:00
|
|
|
if (!processDemoted && canDemoteGlobalVar(GVar, demotedFunc)) {
|
2015-03-30 23:42:36 +08:00
|
|
|
O << "// " << GVar->getName() << " has been demoted\n";
|
2012-05-05 04:18:50 +08:00
|
|
|
if (localDecls.find(demotedFunc) != localDecls.end())
|
|
|
|
localDecls[demotedFunc].push_back(GVar);
|
|
|
|
else {
|
2013-05-20 20:13:32 +08:00
|
|
|
std::vector<const GlobalVariable *> temp;
|
2012-05-05 04:18:50 +08:00
|
|
|
temp.push_back(GVar);
|
|
|
|
localDecls[demotedFunc] = temp;
|
|
|
|
}
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
|
|
|
O << ".";
|
|
|
|
emitPTXAddressSpace(PTy->getAddressSpace(), O);
|
2014-06-28 02:35:58 +08:00
|
|
|
|
|
|
|
if (isManaged(*GVar)) {
|
|
|
|
O << " .attribute(.managed)";
|
|
|
|
}
|
|
|
|
|
2012-05-05 04:18:50 +08:00
|
|
|
if (GVar->getAlignment() == 0)
|
2015-07-16 14:11:10 +08:00
|
|
|
O << " .align " << (int)DL.getPrefTypeAlignment(ETy);
|
2012-05-05 04:18:50 +08:00
|
|
|
else
|
|
|
|
O << " .align " << GVar->getAlignment();
|
|
|
|
|
2017-01-18 08:29:53 +08:00
|
|
|
if (ETy->isFloatingPointTy() || ETy->isPointerTy() ||
|
|
|
|
(ETy->isIntegerTy() && ETy->getScalarSizeInBits() <= 64)) {
|
2012-05-05 04:18:50 +08:00
|
|
|
O << " .";
|
2013-05-20 20:13:28 +08:00
|
|
|
// Special case: ABI requires that we use .u8 for predicates
|
|
|
|
if (ETy->isIntegerTy(1))
|
|
|
|
O << "u8";
|
|
|
|
else
|
|
|
|
O << getPTXFundamentalTypeStr(ETy, false);
|
2012-05-05 04:18:50 +08:00
|
|
|
O << " ";
|
2015-06-09 08:31:39 +08:00
|
|
|
getSymbol(GVar)->print(O, MAI);
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
// Ptx allows variable initilization only for constant and global state
|
|
|
|
// spaces.
|
2014-06-28 02:36:01 +08:00
|
|
|
if (GVar->hasInitializer()) {
|
2017-01-10 06:16:51 +08:00
|
|
|
if ((PTy->getAddressSpace() == ADDRESS_SPACE_GLOBAL) ||
|
|
|
|
(PTy->getAddressSpace() == ADDRESS_SPACE_CONST)) {
|
2014-06-28 02:36:01 +08:00
|
|
|
const Constant *Initializer = GVar->getInitializer();
|
2015-04-24 10:57:30 +08:00
|
|
|
// 'undef' is treated as there is no value specified.
|
2014-06-28 02:36:01 +08:00
|
|
|
if (!Initializer->isNullValue() && !isa<UndefValue>(Initializer)) {
|
|
|
|
O << " = ";
|
|
|
|
printScalarConstant(Initializer, O);
|
|
|
|
}
|
|
|
|
} else {
|
2015-08-22 13:40:26 +08:00
|
|
|
// The frontend adds zero-initializer to device and constant variables
|
|
|
|
// that don't have an initial value, and UndefValue to shared
|
|
|
|
// variables, so skip warning for this case.
|
|
|
|
if (!GVar->getInitializer()->isNullValue() &&
|
|
|
|
!isa<UndefValue>(GVar->getInitializer())) {
|
2015-05-28 19:24:24 +08:00
|
|
|
report_fatal_error("initial value of '" + GVar->getName() +
|
|
|
|
"' is not allowed in addrspace(" +
|
|
|
|
Twine(PTy->getAddressSpace()) + ")");
|
2014-06-28 02:36:01 +08:00
|
|
|
}
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
}
|
|
|
|
} else {
|
2013-03-30 22:29:21 +08:00
|
|
|
unsigned int ElementSize = 0;
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
// Although PTX has direct support for struct type and array type and
|
|
|
|
// LLVM IR is very similar to PTX, the LLVM CodeGen does not support for
|
|
|
|
// targets that support these high level field accesses. Structs, arrays
|
|
|
|
// and vectors are lowered into arrays of bytes.
|
|
|
|
switch (ETy->getTypeID()) {
|
2017-01-18 08:29:53 +08:00
|
|
|
case Type::IntegerTyID: // Integers larger than 64 bits
|
2012-05-05 04:18:50 +08:00
|
|
|
case Type::StructTyID:
|
|
|
|
case Type::ArrayTyID:
|
|
|
|
case Type::VectorTyID:
|
2015-07-16 14:11:10 +08:00
|
|
|
ElementSize = DL.getTypeStoreSize(ETy);
|
2012-05-05 04:18:50 +08:00
|
|
|
// Ptx allows variable initilization only for constant and
|
|
|
|
// global state spaces.
|
2017-01-10 06:16:51 +08:00
|
|
|
if (((PTy->getAddressSpace() == ADDRESS_SPACE_GLOBAL) ||
|
|
|
|
(PTy->getAddressSpace() == ADDRESS_SPACE_CONST)) &&
|
2013-03-30 22:29:21 +08:00
|
|
|
GVar->hasInitializer()) {
|
2013-05-20 20:13:32 +08:00
|
|
|
const Constant *Initializer = GVar->getInitializer();
|
2013-03-30 22:29:21 +08:00
|
|
|
if (!isa<UndefValue>(Initializer) && !Initializer->isNullValue()) {
|
2012-05-05 04:18:50 +08:00
|
|
|
AggBuffer aggBuffer(ElementSize, O, *this);
|
|
|
|
bufferAggregateConstant(Initializer, &aggBuffer);
|
|
|
|
if (aggBuffer.numSymbols) {
|
2015-02-19 08:08:14 +08:00
|
|
|
if (static_cast<const NVPTXTargetMachine &>(TM).is64Bit()) {
|
2015-06-09 08:31:39 +08:00
|
|
|
O << " .u64 ";
|
|
|
|
getSymbol(GVar)->print(O, MAI);
|
|
|
|
O << "[";
|
2013-03-30 22:29:21 +08:00
|
|
|
O << ElementSize / 8;
|
|
|
|
} else {
|
2015-06-09 08:31:39 +08:00
|
|
|
O << " .u32 ";
|
|
|
|
getSymbol(GVar)->print(O, MAI);
|
|
|
|
O << "[";
|
2013-03-30 22:29:21 +08:00
|
|
|
O << ElementSize / 4;
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
O << "]";
|
2013-03-30 22:29:21 +08:00
|
|
|
} else {
|
2015-06-09 08:31:39 +08:00
|
|
|
O << " .b8 ";
|
|
|
|
getSymbol(GVar)->print(O, MAI);
|
|
|
|
O << "[";
|
2012-05-05 04:18:50 +08:00
|
|
|
O << ElementSize;
|
|
|
|
O << "]";
|
|
|
|
}
|
2013-03-30 22:29:21 +08:00
|
|
|
O << " = {";
|
2012-05-05 04:18:50 +08:00
|
|
|
aggBuffer.print();
|
|
|
|
O << "}";
|
2013-03-30 22:29:21 +08:00
|
|
|
} else {
|
2015-06-09 08:31:39 +08:00
|
|
|
O << " .b8 ";
|
|
|
|
getSymbol(GVar)->print(O, MAI);
|
2012-05-05 04:18:50 +08:00
|
|
|
if (ElementSize) {
|
2013-03-30 22:29:21 +08:00
|
|
|
O << "[";
|
2012-05-05 04:18:50 +08:00
|
|
|
O << ElementSize;
|
|
|
|
O << "]";
|
|
|
|
}
|
|
|
|
}
|
2013-03-30 22:29:21 +08:00
|
|
|
} else {
|
2015-06-09 08:31:39 +08:00
|
|
|
O << " .b8 ";
|
|
|
|
getSymbol(GVar)->print(O, MAI);
|
2012-05-05 04:18:50 +08:00
|
|
|
if (ElementSize) {
|
2013-03-30 22:29:21 +08:00
|
|
|
O << "[";
|
2012-05-05 04:18:50 +08:00
|
|
|
O << ElementSize;
|
|
|
|
O << "]";
|
|
|
|
}
|
|
|
|
}
|
|
|
|
break;
|
|
|
|
default:
|
2014-06-18 13:05:13 +08:00
|
|
|
llvm_unreachable("type not supported yet");
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
}
|
|
|
|
O << ";\n";
|
|
|
|
}
|
|
|
|
|
|
|
|
void NVPTXAsmPrinter::emitDemotedVars(const Function *f, raw_ostream &O) {
|
|
|
|
if (localDecls.find(f) == localDecls.end())
|
|
|
|
return;
|
|
|
|
|
2013-05-20 20:13:32 +08:00
|
|
|
std::vector<const GlobalVariable *> &gvars = localDecls[f];
|
2012-05-05 04:18:50 +08:00
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
for (unsigned i = 0, e = gvars.size(); i != e; ++i) {
|
2012-05-05 04:18:50 +08:00
|
|
|
O << "\t// demoted variable\n\t";
|
|
|
|
printModuleLevelGV(gvars[i], O, true);
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
void NVPTXAsmPrinter::emitPTXAddressSpace(unsigned int AddressSpace,
|
|
|
|
raw_ostream &O) const {
|
|
|
|
switch (AddressSpace) {
|
2017-01-10 06:16:51 +08:00
|
|
|
case ADDRESS_SPACE_LOCAL:
|
2013-03-30 22:29:21 +08:00
|
|
|
O << "local";
|
2012-05-05 04:18:50 +08:00
|
|
|
break;
|
2017-01-10 06:16:51 +08:00
|
|
|
case ADDRESS_SPACE_GLOBAL:
|
2013-03-30 22:29:21 +08:00
|
|
|
O << "global";
|
2012-05-05 04:18:50 +08:00
|
|
|
break;
|
2017-01-10 06:16:51 +08:00
|
|
|
case ADDRESS_SPACE_CONST:
|
2013-03-30 22:29:21 +08:00
|
|
|
O << "const";
|
2012-05-05 04:18:50 +08:00
|
|
|
break;
|
2017-01-10 06:16:51 +08:00
|
|
|
case ADDRESS_SPACE_SHARED:
|
2013-03-30 22:29:21 +08:00
|
|
|
O << "shared";
|
2012-05-05 04:18:50 +08:00
|
|
|
break;
|
|
|
|
default:
|
2013-02-09 21:34:15 +08:00
|
|
|
report_fatal_error("Bad address space found while emitting PTX");
|
|
|
|
break;
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
std::string
|
2015-08-02 06:20:21 +08:00
|
|
|
NVPTXAsmPrinter::getPTXFundamentalTypeStr(Type *Ty, bool useB4PTR) const {
|
2012-05-05 04:18:50 +08:00
|
|
|
switch (Ty->getTypeID()) {
|
|
|
|
default:
|
|
|
|
llvm_unreachable("unexpected type");
|
|
|
|
break;
|
|
|
|
case Type::IntegerTyID: {
|
|
|
|
unsigned NumBits = cast<IntegerType>(Ty)->getBitWidth();
|
|
|
|
if (NumBits == 1)
|
|
|
|
return "pred";
|
|
|
|
else if (NumBits <= 64) {
|
|
|
|
std::string name = "u";
|
|
|
|
return name + utostr(NumBits);
|
|
|
|
} else {
|
|
|
|
llvm_unreachable("Integer too large");
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
break;
|
|
|
|
}
|
2017-01-14 04:56:17 +08:00
|
|
|
case Type::HalfTyID:
|
|
|
|
// fp16 is stored as .b16 for compatibility with pre-sm_53 PTX assembly.
|
|
|
|
return "b16";
|
2012-05-05 04:18:50 +08:00
|
|
|
case Type::FloatTyID:
|
|
|
|
return "f32";
|
|
|
|
case Type::DoubleTyID:
|
|
|
|
return "f64";
|
|
|
|
case Type::PointerTyID:
|
2015-02-19 08:08:14 +08:00
|
|
|
if (static_cast<const NVPTXTargetMachine &>(TM).is64Bit())
|
2013-03-30 22:29:21 +08:00
|
|
|
if (useB4PTR)
|
|
|
|
return "b64";
|
|
|
|
else
|
|
|
|
return "u64";
|
|
|
|
else if (useB4PTR)
|
|
|
|
return "b32";
|
2012-05-05 04:18:50 +08:00
|
|
|
else
|
2013-03-30 22:29:21 +08:00
|
|
|
return "u32";
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
llvm_unreachable("unexpected type");
|
2014-04-25 13:30:21 +08:00
|
|
|
return nullptr;
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
void NVPTXAsmPrinter::emitPTXGlobalVariable(const GlobalVariable *GVar,
|
2012-05-05 04:18:50 +08:00
|
|
|
raw_ostream &O) {
|
2015-07-16 14:11:10 +08:00
|
|
|
const DataLayout &DL = getDataLayout();
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
// GlobalVariables are always constant pointers themselves.
|
2016-01-17 04:30:46 +08:00
|
|
|
Type *ETy = GVar->getValueType();
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
O << ".";
|
2016-01-17 04:30:46 +08:00
|
|
|
emitPTXAddressSpace(GVar->getType()->getAddressSpace(), O);
|
2012-05-05 04:18:50 +08:00
|
|
|
if (GVar->getAlignment() == 0)
|
2015-07-16 14:11:10 +08:00
|
|
|
O << " .align " << (int)DL.getPrefTypeAlignment(ETy);
|
2012-05-05 04:18:50 +08:00
|
|
|
else
|
|
|
|
O << " .align " << GVar->getAlignment();
|
|
|
|
|
2014-12-18 01:59:04 +08:00
|
|
|
if (ETy->isFloatingPointTy() || ETy->isIntegerTy() || ETy->isPointerTy()) {
|
2012-05-05 04:18:50 +08:00
|
|
|
O << " .";
|
|
|
|
O << getPTXFundamentalTypeStr(ETy);
|
|
|
|
O << " ";
|
2015-06-09 08:31:39 +08:00
|
|
|
getSymbol(GVar)->print(O, MAI);
|
2012-05-05 04:18:50 +08:00
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
int64_t ElementSize = 0;
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
// Although PTX has direct support for struct type and array type and LLVM IR
|
|
|
|
// is very similar to PTX, the LLVM CodeGen does not support for targets that
|
|
|
|
// support these high level field accesses. Structs and arrays are lowered
|
|
|
|
// into arrays of bytes.
|
|
|
|
switch (ETy->getTypeID()) {
|
|
|
|
case Type::StructTyID:
|
|
|
|
case Type::ArrayTyID:
|
|
|
|
case Type::VectorTyID:
|
2015-07-16 14:11:10 +08:00
|
|
|
ElementSize = DL.getTypeStoreSize(ETy);
|
2015-06-09 08:31:39 +08:00
|
|
|
O << " .b8 ";
|
|
|
|
getSymbol(GVar)->print(O, MAI);
|
|
|
|
O << "[";
|
2012-05-05 04:18:50 +08:00
|
|
|
if (ElementSize) {
|
2015-05-28 19:24:24 +08:00
|
|
|
O << ElementSize;
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
O << "]";
|
|
|
|
break;
|
|
|
|
default:
|
2014-06-18 13:05:13 +08:00
|
|
|
llvm_unreachable("type not supported yet");
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2015-07-16 14:11:10 +08:00
|
|
|
static unsigned int getOpenCLAlignment(const DataLayout &DL, Type *Ty) {
|
2013-12-08 03:34:20 +08:00
|
|
|
if (Ty->isSingleValueType())
|
2015-07-16 14:11:10 +08:00
|
|
|
return DL.getPrefTypeAlignment(Ty);
|
2012-05-05 04:18:50 +08:00
|
|
|
|
2015-08-02 06:20:21 +08:00
|
|
|
auto *ATy = dyn_cast<ArrayType>(Ty);
|
2012-05-05 04:18:50 +08:00
|
|
|
if (ATy)
|
2015-07-16 14:11:10 +08:00
|
|
|
return getOpenCLAlignment(DL, ATy->getElementType());
|
2012-05-05 04:18:50 +08:00
|
|
|
|
2015-08-02 06:20:21 +08:00
|
|
|
auto *STy = dyn_cast<StructType>(Ty);
|
2012-05-05 04:18:50 +08:00
|
|
|
if (STy) {
|
|
|
|
unsigned int alignStruct = 1;
|
|
|
|
// Go through each element of the struct and find the
|
|
|
|
// largest alignment.
|
2013-03-30 22:29:21 +08:00
|
|
|
for (unsigned i = 0, e = STy->getNumElements(); i != e; i++) {
|
2012-05-05 04:18:50 +08:00
|
|
|
Type *ETy = STy->getElementType(i);
|
2015-07-16 14:11:10 +08:00
|
|
|
unsigned int align = getOpenCLAlignment(DL, ETy);
|
2012-05-05 04:18:50 +08:00
|
|
|
if (align > alignStruct)
|
|
|
|
alignStruct = align;
|
|
|
|
}
|
|
|
|
return alignStruct;
|
|
|
|
}
|
|
|
|
|
2015-08-02 06:20:21 +08:00
|
|
|
auto *FTy = dyn_cast<FunctionType>(Ty);
|
2012-05-05 04:18:50 +08:00
|
|
|
if (FTy)
|
2015-07-16 14:11:10 +08:00
|
|
|
return DL.getPointerPrefAlignment();
|
|
|
|
return DL.getPrefTypeAlignment(Ty);
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
void NVPTXAsmPrinter::printParamName(Function::const_arg_iterator I,
|
|
|
|
int paramIndex, raw_ostream &O) {
|
2015-06-09 08:31:39 +08:00
|
|
|
getSymbol(I->getParent())->print(O, MAI);
|
|
|
|
O << "_param_" << paramIndex;
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
void NVPTXAsmPrinter::emitFunctionParamList(const Function *F, raw_ostream &O) {
|
2015-07-16 14:11:10 +08:00
|
|
|
const DataLayout &DL = getDataLayout();
|
Rename AttributeSet to AttributeList
Summary:
This class is a list of AttributeSetNodes corresponding the function
prototype of a call or function declaration. This class used to be
called ParamAttrListPtr, then AttrListPtr, then AttributeSet. It is
typically accessed by parameter and return value index, so
"AttributeList" seems like a more intuitive name.
Rename AttributeSetImpl to AttributeListImpl to follow suit.
It's useful to rename this class so that we can rename AttributeSetNode
to AttributeSet later. AttributeSet is the set of attributes that apply
to a single function, argument, or return value.
Reviewers: sanjoy, javed.absar, chandlerc, pete
Reviewed By: pete
Subscribers: pete, jholewinski, arsenm, dschuff, mehdi_amini, jfb, nhaehnle, sbc100, void, llvm-commits
Differential Revision: https://reviews.llvm.org/D31102
llvm-svn: 298393
2017-03-22 00:57:19 +08:00
|
|
|
const AttributeList &PAL = F->getAttributes();
|
2015-02-19 08:08:14 +08:00
|
|
|
const TargetLowering *TLI = nvptxSubtarget->getTargetLowering();
|
2012-05-05 04:18:50 +08:00
|
|
|
Function::const_arg_iterator I, E;
|
|
|
|
unsigned paramIndex = 0;
|
|
|
|
bool first = true;
|
2017-01-10 06:16:51 +08:00
|
|
|
bool isKernelFunc = isKernelFunction(*F);
|
2015-02-19 08:08:14 +08:00
|
|
|
bool isABI = (nvptxSubtarget->getSmVersion() >= 20);
|
2015-07-16 14:11:10 +08:00
|
|
|
MVT thePointerTy = TLI->getPointerTy(DL);
|
2012-05-05 04:18:50 +08:00
|
|
|
|
2016-01-24 05:12:17 +08:00
|
|
|
if (F->arg_empty()) {
|
|
|
|
O << "()\n";
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2012-05-05 04:18:50 +08:00
|
|
|
O << "(\n";
|
|
|
|
|
|
|
|
for (I = F->arg_begin(), E = F->arg_end(); I != E; ++I, paramIndex++) {
|
2013-03-25 05:17:47 +08:00
|
|
|
Type *Ty = I->getType();
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
if (!first)
|
|
|
|
O << ",\n";
|
|
|
|
|
|
|
|
first = false;
|
|
|
|
|
|
|
|
// Handle image/sampler parameters
|
2014-04-09 23:39:15 +08:00
|
|
|
if (isKernelFunction(*F)) {
|
|
|
|
if (isSampler(*I) || isImage(*I)) {
|
|
|
|
if (isImage(*I)) {
|
|
|
|
std::string sname = I->getName();
|
|
|
|
if (isImageWriteOnly(*I) || isImageReadWrite(*I)) {
|
2015-02-19 08:08:14 +08:00
|
|
|
if (nvptxSubtarget->hasImageHandles())
|
2014-04-09 23:39:15 +08:00
|
|
|
O << "\t.param .u64 .ptr .surfref ";
|
|
|
|
else
|
|
|
|
O << "\t.param .surfref ";
|
2015-06-09 08:31:39 +08:00
|
|
|
CurrentFnSym->print(O, MAI);
|
|
|
|
O << "_param_" << paramIndex;
|
2014-04-09 23:39:15 +08:00
|
|
|
}
|
|
|
|
else { // Default image is read_only
|
2015-02-19 08:08:14 +08:00
|
|
|
if (nvptxSubtarget->hasImageHandles())
|
2014-04-09 23:39:15 +08:00
|
|
|
O << "\t.param .u64 .ptr .texref ";
|
|
|
|
else
|
|
|
|
O << "\t.param .texref ";
|
2015-06-09 08:31:39 +08:00
|
|
|
CurrentFnSym->print(O, MAI);
|
|
|
|
O << "_param_" << paramIndex;
|
2014-04-09 23:39:15 +08:00
|
|
|
}
|
|
|
|
} else {
|
2015-02-19 08:08:14 +08:00
|
|
|
if (nvptxSubtarget->hasImageHandles())
|
2014-04-09 23:39:15 +08:00
|
|
|
O << "\t.param .u64 .ptr .samplerref ";
|
|
|
|
else
|
|
|
|
O << "\t.param .samplerref ";
|
2015-06-09 08:31:39 +08:00
|
|
|
CurrentFnSym->print(O, MAI);
|
|
|
|
O << "_param_" << paramIndex;
|
2014-04-09 23:39:15 +08:00
|
|
|
}
|
|
|
|
continue;
|
|
|
|
}
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
2017-04-14 07:12:13 +08:00
|
|
|
if (!PAL.hasParamAttribute(paramIndex, Attribute::ByVal)) {
|
2014-01-29 02:35:29 +08:00
|
|
|
if (Ty->isAggregateType() || Ty->isVectorTy()) {
|
|
|
|
// Just print .param .align <a> .b8 .param[size];
|
2013-03-25 05:17:47 +08:00
|
|
|
// <a> = PAL.getparamalignment
|
|
|
|
// size = typeallocsize of element type
|
2017-04-29 04:34:27 +08:00
|
|
|
unsigned align = PAL.getParamAlignment(paramIndex);
|
2013-03-25 05:17:47 +08:00
|
|
|
if (align == 0)
|
2015-07-16 14:11:10 +08:00
|
|
|
align = DL.getABITypeAlignment(Ty);
|
2013-03-25 05:17:47 +08:00
|
|
|
|
2015-07-16 14:11:10 +08:00
|
|
|
unsigned sz = DL.getTypeAllocSize(Ty);
|
2013-03-30 22:29:21 +08:00
|
|
|
O << "\t.param .align " << align << " .b8 ";
|
2013-03-25 05:17:47 +08:00
|
|
|
printParamName(I, paramIndex, O);
|
|
|
|
O << "[" << sz << "]";
|
|
|
|
|
|
|
|
continue;
|
|
|
|
}
|
2012-05-05 04:18:50 +08:00
|
|
|
// Just a scalar
|
2015-08-02 06:20:21 +08:00
|
|
|
auto *PTy = dyn_cast<PointerType>(Ty);
|
2012-05-05 04:18:50 +08:00
|
|
|
if (isKernelFunc) {
|
|
|
|
if (PTy) {
|
|
|
|
// Special handling for pointer arguments to kernel
|
|
|
|
O << "\t.param .u" << thePointerTy.getSizeInBits() << " ";
|
|
|
|
|
2015-02-19 08:08:23 +08:00
|
|
|
if (static_cast<NVPTXTargetMachine &>(TM).getDrvInterface() !=
|
|
|
|
NVPTX::CUDA) {
|
2012-05-05 04:18:50 +08:00
|
|
|
Type *ETy = PTy->getElementType();
|
|
|
|
int addrSpace = PTy->getAddressSpace();
|
2013-03-30 22:29:21 +08:00
|
|
|
switch (addrSpace) {
|
2012-05-05 04:18:50 +08:00
|
|
|
default:
|
|
|
|
O << ".ptr ";
|
|
|
|
break;
|
2017-01-10 06:16:51 +08:00
|
|
|
case ADDRESS_SPACE_CONST:
|
2012-05-05 04:18:50 +08:00
|
|
|
O << ".ptr .const ";
|
|
|
|
break;
|
2017-01-10 06:16:51 +08:00
|
|
|
case ADDRESS_SPACE_SHARED:
|
2012-05-05 04:18:50 +08:00
|
|
|
O << ".ptr .shared ";
|
|
|
|
break;
|
2017-01-10 06:16:51 +08:00
|
|
|
case ADDRESS_SPACE_GLOBAL:
|
2012-05-05 04:18:50 +08:00
|
|
|
O << ".ptr .global ";
|
|
|
|
break;
|
|
|
|
}
|
2015-07-16 14:11:10 +08:00
|
|
|
O << ".align " << (int)getOpenCLAlignment(DL, ETy) << " ";
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
printParamName(I, paramIndex, O);
|
|
|
|
continue;
|
|
|
|
}
|
|
|
|
|
|
|
|
// non-pointer scalar to kernel func
|
2013-05-20 20:13:28 +08:00
|
|
|
O << "\t.param .";
|
|
|
|
// Special case: predicate operands become .u8 types
|
|
|
|
if (Ty->isIntegerTy(1))
|
|
|
|
O << "u8";
|
|
|
|
else
|
|
|
|
O << getPTXFundamentalTypeStr(Ty);
|
|
|
|
O << " ";
|
2012-05-05 04:18:50 +08:00
|
|
|
printParamName(I, paramIndex, O);
|
|
|
|
continue;
|
|
|
|
}
|
|
|
|
// Non-kernel function, just print .param .b<size> for ABI
|
2013-12-05 13:44:44 +08:00
|
|
|
// and .reg .b<size> for non-ABI
|
2012-05-05 04:18:50 +08:00
|
|
|
unsigned sz = 0;
|
|
|
|
if (isa<IntegerType>(Ty)) {
|
|
|
|
sz = cast<IntegerType>(Ty)->getBitWidth();
|
2013-03-30 22:29:21 +08:00
|
|
|
if (sz < 32)
|
|
|
|
sz = 32;
|
|
|
|
} else if (isa<PointerType>(Ty))
|
2012-05-05 04:18:50 +08:00
|
|
|
sz = thePointerTy.getSizeInBits();
|
2017-01-14 04:56:17 +08:00
|
|
|
else if (Ty->isHalfTy())
|
|
|
|
// PTX ABI requires all scalar parameters to be at least 32
|
|
|
|
// bits in size. fp16 normally uses .b16 as its storage type
|
|
|
|
// in PTX, so its size must be adjusted here, too.
|
|
|
|
sz = 32;
|
2012-05-05 04:18:50 +08:00
|
|
|
else
|
|
|
|
sz = Ty->getPrimitiveSizeInBits();
|
|
|
|
if (isABI)
|
|
|
|
O << "\t.param .b" << sz << " ";
|
|
|
|
else
|
|
|
|
O << "\t.reg .b" << sz << " ";
|
|
|
|
printParamName(I, paramIndex, O);
|
|
|
|
continue;
|
|
|
|
}
|
|
|
|
|
|
|
|
// param has byVal attribute. So should be a pointer
|
2015-08-02 06:20:21 +08:00
|
|
|
auto *PTy = dyn_cast<PointerType>(Ty);
|
2013-03-30 22:29:21 +08:00
|
|
|
assert(PTy && "Param with byval attribute should be a pointer type");
|
2012-05-05 04:18:50 +08:00
|
|
|
Type *ETy = PTy->getElementType();
|
|
|
|
|
|
|
|
if (isABI || isKernelFunc) {
|
2014-01-29 02:35:29 +08:00
|
|
|
// Just print .param .align <a> .b8 .param[size];
|
2012-05-05 04:18:50 +08:00
|
|
|
// <a> = PAL.getparamalignment
|
|
|
|
// size = typeallocsize of element type
|
2017-04-29 04:34:27 +08:00
|
|
|
unsigned align = PAL.getParamAlignment(paramIndex);
|
2012-11-10 07:50:24 +08:00
|
|
|
if (align == 0)
|
2015-07-16 14:11:10 +08:00
|
|
|
align = DL.getABITypeAlignment(ETy);
|
2016-07-19 03:54:56 +08:00
|
|
|
// Work around a bug in ptxas. When PTX code takes address of
|
|
|
|
// byval parameter with alignment < 4, ptxas generates code to
|
|
|
|
// spill argument into memory. Alas on sm_50+ ptxas generates
|
|
|
|
// SASS code that fails with misaligned access. To work around
|
|
|
|
// the problem, make sure that we align byval parameters by at
|
|
|
|
// least 4. Matching change must be made in LowerCall() where we
|
|
|
|
// prepare parameters for the call.
|
|
|
|
//
|
|
|
|
// TODO: this will need to be undone when we get to support multi-TU
|
|
|
|
// device-side compilation as it breaks ABI compatibility with nvcc.
|
|
|
|
// Hopefully ptxas bug is fixed by then.
|
|
|
|
if (!isKernelFunc && align < 4)
|
|
|
|
align = 4;
|
2015-07-16 14:11:10 +08:00
|
|
|
unsigned sz = DL.getTypeAllocSize(ETy);
|
2013-03-30 22:29:21 +08:00
|
|
|
O << "\t.param .align " << align << " .b8 ";
|
2012-05-05 04:18:50 +08:00
|
|
|
printParamName(I, paramIndex, O);
|
|
|
|
O << "[" << sz << "]";
|
|
|
|
continue;
|
|
|
|
} else {
|
|
|
|
// Split the ETy into constituent parts and
|
|
|
|
// print .param .b<size> <name> for each part.
|
|
|
|
// Further, if a part is vector, print the above for
|
|
|
|
// each vector element.
|
|
|
|
SmallVector<EVT, 16> vtparts;
|
2015-07-16 14:11:10 +08:00
|
|
|
ComputeValueVTs(*TLI, DL, ETy, vtparts);
|
2013-03-30 22:29:21 +08:00
|
|
|
for (unsigned i = 0, e = vtparts.size(); i != e; ++i) {
|
2012-05-05 04:18:50 +08:00
|
|
|
unsigned elems = 1;
|
|
|
|
EVT elemtype = vtparts[i];
|
|
|
|
if (vtparts[i].isVector()) {
|
|
|
|
elems = vtparts[i].getVectorNumElements();
|
|
|
|
elemtype = vtparts[i].getVectorElementType();
|
|
|
|
}
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
for (unsigned j = 0, je = elems; j != je; ++j) {
|
2012-05-05 04:18:50 +08:00
|
|
|
unsigned sz = elemtype.getSizeInBits();
|
2013-03-30 22:29:21 +08:00
|
|
|
if (elemtype.isInteger() && (sz < 32))
|
|
|
|
sz = 32;
|
2012-05-05 04:18:50 +08:00
|
|
|
O << "\t.reg .b" << sz << " ";
|
|
|
|
printParamName(I, paramIndex, O);
|
2013-03-30 22:29:21 +08:00
|
|
|
if (j < je - 1)
|
|
|
|
O << ",\n";
|
2012-05-05 04:18:50 +08:00
|
|
|
++paramIndex;
|
|
|
|
}
|
2013-03-30 22:29:21 +08:00
|
|
|
if (i < e - 1)
|
2012-05-05 04:18:50 +08:00
|
|
|
O << ",\n";
|
|
|
|
}
|
|
|
|
--paramIndex;
|
|
|
|
continue;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
O << "\n)\n";
|
|
|
|
}
|
|
|
|
|
|
|
|
void NVPTXAsmPrinter::emitFunctionParamList(const MachineFunction &MF,
|
|
|
|
raw_ostream &O) {
|
|
|
|
const Function *F = MF.getFunction();
|
|
|
|
emitFunctionParamList(F, O);
|
|
|
|
}
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
void NVPTXAsmPrinter::setAndEmitFunctionVirtualRegisters(
|
|
|
|
const MachineFunction &MF) {
|
2012-05-05 04:18:50 +08:00
|
|
|
SmallString<128> Str;
|
|
|
|
raw_svector_ostream O(Str);
|
|
|
|
|
|
|
|
// Map the global virtual register number to a register class specific
|
|
|
|
// virtual register number starting from 1 with that class.
|
2014-08-05 10:39:49 +08:00
|
|
|
const TargetRegisterInfo *TRI = MF.getSubtarget().getRegisterInfo();
|
2012-05-05 04:18:50 +08:00
|
|
|
//unsigned numRegClasses = TRI->getNumRegClasses();
|
|
|
|
|
|
|
|
// Emit the Fake Stack Object
|
2016-07-29 02:40:00 +08:00
|
|
|
const MachineFrameInfo &MFI = MF.getFrameInfo();
|
|
|
|
int NumBytes = (int) MFI.getStackSize();
|
2012-05-05 04:18:50 +08:00
|
|
|
if (NumBytes) {
|
2016-07-29 02:40:00 +08:00
|
|
|
O << "\t.local .align " << MFI.getMaxAlignment() << " .b8 \t" << DEPOTNAME
|
2013-03-30 22:29:21 +08:00
|
|
|
<< getFunctionNumber() << "[" << NumBytes << "];\n";
|
2015-02-19 08:08:27 +08:00
|
|
|
if (static_cast<const NVPTXTargetMachine &>(MF.getTarget()).is64Bit()) {
|
2012-05-05 04:18:50 +08:00
|
|
|
O << "\t.reg .b64 \t%SP;\n";
|
|
|
|
O << "\t.reg .b64 \t%SPL;\n";
|
2013-03-30 22:29:21 +08:00
|
|
|
} else {
|
2012-05-05 04:18:50 +08:00
|
|
|
O << "\t.reg .b32 \t%SP;\n";
|
|
|
|
O << "\t.reg .b32 \t%SPL;\n";
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
// Go through all virtual registers to establish the mapping between the
|
|
|
|
// global virtual
|
|
|
|
// register number and the per class virtual register number.
|
|
|
|
// We use the per class virtual register number in the ptx output.
|
|
|
|
unsigned int numVRs = MRI->getNumVirtRegs();
|
2013-03-30 22:29:21 +08:00
|
|
|
for (unsigned i = 0; i < numVRs; i++) {
|
2012-05-05 04:18:50 +08:00
|
|
|
unsigned int vr = TRI->index2VirtReg(i);
|
|
|
|
const TargetRegisterClass *RC = MRI->getRegClass(vr);
|
2013-05-31 20:14:49 +08:00
|
|
|
DenseMap<unsigned, unsigned> ®map = VRegMapping[RC];
|
2012-05-05 04:18:50 +08:00
|
|
|
int n = regmap.size();
|
2013-03-30 22:29:21 +08:00
|
|
|
regmap.insert(std::make_pair(vr, n + 1));
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
// Emit register declarations
|
|
|
|
// @TODO: Extract out the real register usage
|
2013-05-31 20:14:49 +08:00
|
|
|
// O << "\t.reg .pred %p<" << NVPTXNumRegisters << ">;\n";
|
|
|
|
// O << "\t.reg .s16 %rc<" << NVPTXNumRegisters << ">;\n";
|
|
|
|
// O << "\t.reg .s16 %rs<" << NVPTXNumRegisters << ">;\n";
|
|
|
|
// O << "\t.reg .s32 %r<" << NVPTXNumRegisters << ">;\n";
|
2014-07-17 00:26:58 +08:00
|
|
|
// O << "\t.reg .s64 %rd<" << NVPTXNumRegisters << ">;\n";
|
2013-05-31 20:14:49 +08:00
|
|
|
// O << "\t.reg .f32 %f<" << NVPTXNumRegisters << ">;\n";
|
2014-07-17 00:26:58 +08:00
|
|
|
// O << "\t.reg .f64 %fd<" << NVPTXNumRegisters << ">;\n";
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
// Emit declaration of the virtual registers or 'physical' registers for
|
|
|
|
// each register class
|
2013-05-31 20:14:49 +08:00
|
|
|
for (unsigned i=0; i< TRI->getNumRegClasses(); i++) {
|
|
|
|
const TargetRegisterClass *RC = TRI->getRegClass(i);
|
|
|
|
DenseMap<unsigned, unsigned> ®map = VRegMapping[RC];
|
|
|
|
std::string rcname = getNVPTXRegClassName(RC);
|
|
|
|
std::string rcStr = getNVPTXRegClassStr(RC);
|
|
|
|
int n = regmap.size();
|
|
|
|
|
|
|
|
// Only declare those registers that may be used.
|
|
|
|
if (n) {
|
|
|
|
O << "\t.reg " << rcname << " \t" << rcStr << "<" << (n+1)
|
|
|
|
<< ">;\n";
|
|
|
|
}
|
|
|
|
}
|
2012-05-05 04:18:50 +08:00
|
|
|
|
2015-04-25 03:11:51 +08:00
|
|
|
OutStreamer->EmitRawText(O.str());
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
void NVPTXAsmPrinter::printFPConstant(const ConstantFP *Fp, raw_ostream &O) {
|
2013-03-30 22:29:21 +08:00
|
|
|
APFloat APF = APFloat(Fp->getValueAPF()); // make a copy
|
2012-05-05 04:18:50 +08:00
|
|
|
bool ignored;
|
|
|
|
unsigned int numHex;
|
|
|
|
const char *lead;
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
if (Fp->getType()->getTypeID() == Type::FloatTyID) {
|
2012-05-05 04:18:50 +08:00
|
|
|
numHex = 8;
|
|
|
|
lead = "0f";
|
2016-12-14 19:57:17 +08:00
|
|
|
APF.convert(APFloat::IEEEsingle(), APFloat::rmNearestTiesToEven, &ignored);
|
2012-05-05 04:18:50 +08:00
|
|
|
} else if (Fp->getType()->getTypeID() == Type::DoubleTyID) {
|
|
|
|
numHex = 16;
|
|
|
|
lead = "0d";
|
2016-12-14 19:57:17 +08:00
|
|
|
APF.convert(APFloat::IEEEdouble(), APFloat::rmNearestTiesToEven, &ignored);
|
2012-05-05 04:18:50 +08:00
|
|
|
} else
|
|
|
|
llvm_unreachable("unsupported fp type");
|
|
|
|
|
|
|
|
APInt API = APF.bitcastToAPInt();
|
|
|
|
std::string hexstr(utohexstr(API.getZExtValue()));
|
|
|
|
O << lead;
|
|
|
|
if (hexstr.length() < numHex)
|
|
|
|
O << std::string(numHex - hexstr.length(), '0');
|
|
|
|
O << utohexstr(API.getZExtValue());
|
|
|
|
}
|
|
|
|
|
2013-05-20 20:13:32 +08:00
|
|
|
void NVPTXAsmPrinter::printScalarConstant(const Constant *CPV, raw_ostream &O) {
|
|
|
|
if (const ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) {
|
2012-05-05 04:18:50 +08:00
|
|
|
O << CI->getValue();
|
|
|
|
return;
|
|
|
|
}
|
2013-05-20 20:13:32 +08:00
|
|
|
if (const ConstantFP *CFP = dyn_cast<ConstantFP>(CPV)) {
|
2012-05-05 04:18:50 +08:00
|
|
|
printFPConstant(CFP, O);
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
if (isa<ConstantPointerNull>(CPV)) {
|
|
|
|
O << "0";
|
|
|
|
return;
|
|
|
|
}
|
2013-05-20 20:13:32 +08:00
|
|
|
if (const GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) {
|
2014-04-09 23:39:11 +08:00
|
|
|
bool IsNonGenericPointer = false;
|
2016-01-17 04:30:46 +08:00
|
|
|
if (GVar->getType()->getAddressSpace() != 0) {
|
2014-04-09 23:39:11 +08:00
|
|
|
IsNonGenericPointer = true;
|
|
|
|
}
|
|
|
|
if (EmitGeneric && !isa<Function>(CPV) && !IsNonGenericPointer) {
|
|
|
|
O << "generic(";
|
2015-06-09 08:31:39 +08:00
|
|
|
getSymbol(GVar)->print(O, MAI);
|
2014-04-09 23:39:11 +08:00
|
|
|
O << ")";
|
|
|
|
} else {
|
2015-06-09 08:31:39 +08:00
|
|
|
getSymbol(GVar)->print(O, MAI);
|
2014-04-09 23:39:11 +08:00
|
|
|
}
|
2012-05-05 04:18:50 +08:00
|
|
|
return;
|
|
|
|
}
|
2013-05-20 20:13:32 +08:00
|
|
|
if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
|
|
|
|
const Value *v = Cexpr->stripPointerCasts();
|
2014-04-09 23:39:11 +08:00
|
|
|
PointerType *PTy = dyn_cast<PointerType>(Cexpr->getType());
|
|
|
|
bool IsNonGenericPointer = false;
|
|
|
|
if (PTy && PTy->getAddressSpace() != 0) {
|
|
|
|
IsNonGenericPointer = true;
|
|
|
|
}
|
2013-05-20 20:13:32 +08:00
|
|
|
if (const GlobalValue *GVar = dyn_cast<GlobalValue>(v)) {
|
2014-04-09 23:39:11 +08:00
|
|
|
if (EmitGeneric && !isa<Function>(v) && !IsNonGenericPointer) {
|
|
|
|
O << "generic(";
|
2015-06-09 08:31:39 +08:00
|
|
|
getSymbol(GVar)->print(O, MAI);
|
2014-04-09 23:39:11 +08:00
|
|
|
O << ")";
|
|
|
|
} else {
|
2015-06-09 08:31:39 +08:00
|
|
|
getSymbol(GVar)->print(O, MAI);
|
2014-04-09 23:39:11 +08:00
|
|
|
}
|
2012-05-05 04:18:50 +08:00
|
|
|
return;
|
|
|
|
} else {
|
2015-06-09 08:31:39 +08:00
|
|
|
lowerConstant(CPV)->print(O, MAI);
|
2012-05-05 04:18:50 +08:00
|
|
|
return;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
llvm_unreachable("Not scalar type found in printScalarConstant()");
|
|
|
|
}
|
|
|
|
|
2015-06-10 00:29:34 +08:00
|
|
|
// These utility functions assure we get the right sequence of bytes for a given
|
|
|
|
// type even for big-endian machines
|
|
|
|
template <typename T> static void ConvertIntToBytes(unsigned char *p, T val) {
|
|
|
|
int64_t vp = (int64_t)val;
|
|
|
|
for (unsigned i = 0; i < sizeof(T); ++i) {
|
|
|
|
p[i] = (unsigned char)vp;
|
|
|
|
vp >>= 8;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
static void ConvertFloatToBytes(unsigned char *p, float val) {
|
|
|
|
int32_t *vp = (int32_t *)&val;
|
|
|
|
for (unsigned i = 0; i < sizeof(int32_t); ++i) {
|
|
|
|
p[i] = (unsigned char)*vp;
|
|
|
|
*vp >>= 8;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
static void ConvertDoubleToBytes(unsigned char *p, double val) {
|
|
|
|
int64_t *vp = (int64_t *)&val;
|
|
|
|
for (unsigned i = 0; i < sizeof(int64_t); ++i) {
|
|
|
|
p[i] = (unsigned char)*vp;
|
|
|
|
*vp >>= 8;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2013-05-20 20:13:32 +08:00
|
|
|
void NVPTXAsmPrinter::bufferLEByte(const Constant *CPV, int Bytes,
|
2012-05-05 04:18:50 +08:00
|
|
|
AggBuffer *aggBuffer) {
|
2015-07-16 14:11:10 +08:00
|
|
|
const DataLayout &DL = getDataLayout();
|
2012-05-05 04:18:50 +08:00
|
|
|
|
|
|
|
if (isa<UndefValue>(CPV) || CPV->isNullValue()) {
|
2015-07-16 14:11:10 +08:00
|
|
|
int s = DL.getTypeAllocSize(CPV->getType());
|
2013-03-30 22:29:21 +08:00
|
|
|
if (s < Bytes)
|
2012-05-05 04:18:50 +08:00
|
|
|
s = Bytes;
|
|
|
|
aggBuffer->addZeros(s);
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2015-06-10 00:29:34 +08:00
|
|
|
unsigned char ptr[8];
|
2012-05-05 04:18:50 +08:00
|
|
|
switch (CPV->getType()->getTypeID()) {
|
|
|
|
|
|
|
|
case Type::IntegerTyID: {
|
2015-08-02 06:20:21 +08:00
|
|
|
Type *ETy = CPV->getType();
|
2013-03-30 22:29:21 +08:00
|
|
|
if (ETy == Type::getInt8Ty(CPV->getContext())) {
|
2015-04-10 19:24:51 +08:00
|
|
|
unsigned char c = (unsigned char)cast<ConstantInt>(CPV)->getZExtValue();
|
2015-06-10 00:29:34 +08:00
|
|
|
ConvertIntToBytes<>(ptr, c);
|
2012-05-05 04:18:50 +08:00
|
|
|
aggBuffer->addBytes(ptr, 1, Bytes);
|
2013-03-30 22:29:21 +08:00
|
|
|
} else if (ETy == Type::getInt16Ty(CPV->getContext())) {
|
2015-04-10 19:24:51 +08:00
|
|
|
short int16 = (short)cast<ConstantInt>(CPV)->getZExtValue();
|
2015-06-10 00:29:34 +08:00
|
|
|
ConvertIntToBytes<>(ptr, int16);
|
2012-05-05 04:18:50 +08:00
|
|
|
aggBuffer->addBytes(ptr, 2, Bytes);
|
2013-03-30 22:29:21 +08:00
|
|
|
} else if (ETy == Type::getInt32Ty(CPV->getContext())) {
|
2013-05-20 20:13:32 +08:00
|
|
|
if (const ConstantInt *constInt = dyn_cast<ConstantInt>(CPV)) {
|
2013-03-30 22:29:21 +08:00
|
|
|
int int32 = (int)(constInt->getZExtValue());
|
2015-06-10 00:29:34 +08:00
|
|
|
ConvertIntToBytes<>(ptr, int32);
|
2012-05-05 04:18:50 +08:00
|
|
|
aggBuffer->addBytes(ptr, 4, Bytes);
|
|
|
|
break;
|
2016-07-29 11:27:26 +08:00
|
|
|
} else if (const auto *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
|
|
|
|
if (const auto *constInt = dyn_cast_or_null<ConstantInt>(
|
|
|
|
ConstantFoldConstant(Cexpr, DL))) {
|
2013-03-30 22:29:21 +08:00
|
|
|
int int32 = (int)(constInt->getZExtValue());
|
2015-06-10 00:29:34 +08:00
|
|
|
ConvertIntToBytes<>(ptr, int32);
|
2012-05-05 04:18:50 +08:00
|
|
|
aggBuffer->addBytes(ptr, 4, Bytes);
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
if (Cexpr->getOpcode() == Instruction::PtrToInt) {
|
|
|
|
Value *v = Cexpr->getOperand(0)->stripPointerCasts();
|
2015-04-24 10:57:30 +08:00
|
|
|
aggBuffer->addSymbol(v, Cexpr->getOperand(0));
|
2012-05-05 04:18:50 +08:00
|
|
|
aggBuffer->addZeros(4);
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
}
|
2012-05-24 15:02:50 +08:00
|
|
|
llvm_unreachable("unsupported integer const type");
|
2013-03-30 22:29:21 +08:00
|
|
|
} else if (ETy == Type::getInt64Ty(CPV->getContext())) {
|
2013-05-20 20:13:32 +08:00
|
|
|
if (const ConstantInt *constInt = dyn_cast<ConstantInt>(CPV)) {
|
2013-03-30 22:29:21 +08:00
|
|
|
long long int64 = (long long)(constInt->getZExtValue());
|
2015-06-10 00:29:34 +08:00
|
|
|
ConvertIntToBytes<>(ptr, int64);
|
2012-05-05 04:18:50 +08:00
|
|
|
aggBuffer->addBytes(ptr, 8, Bytes);
|
|
|
|
break;
|
2013-05-20 20:13:32 +08:00
|
|
|
} else if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
|
2016-07-29 11:27:26 +08:00
|
|
|
if (const auto *constInt = dyn_cast_or_null<ConstantInt>(
|
|
|
|
ConstantFoldConstant(Cexpr, DL))) {
|
2013-03-30 22:29:21 +08:00
|
|
|
long long int64 = (long long)(constInt->getZExtValue());
|
2015-06-10 00:29:34 +08:00
|
|
|
ConvertIntToBytes<>(ptr, int64);
|
2012-05-05 04:18:50 +08:00
|
|
|
aggBuffer->addBytes(ptr, 8, Bytes);
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
if (Cexpr->getOpcode() == Instruction::PtrToInt) {
|
|
|
|
Value *v = Cexpr->getOperand(0)->stripPointerCasts();
|
2015-04-24 10:57:30 +08:00
|
|
|
aggBuffer->addSymbol(v, Cexpr->getOperand(0));
|
2012-05-05 04:18:50 +08:00
|
|
|
aggBuffer->addZeros(8);
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
llvm_unreachable("unsupported integer const type");
|
2012-05-24 15:02:50 +08:00
|
|
|
} else
|
2012-05-05 04:18:50 +08:00
|
|
|
llvm_unreachable("unsupported integer const type");
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
case Type::FloatTyID:
|
|
|
|
case Type::DoubleTyID: {
|
2013-05-20 20:13:32 +08:00
|
|
|
const ConstantFP *CFP = dyn_cast<ConstantFP>(CPV);
|
2015-08-02 06:20:21 +08:00
|
|
|
Type *Ty = CFP->getType();
|
2012-05-05 04:18:50 +08:00
|
|
|
if (Ty == Type::getFloatTy(CPV->getContext())) {
|
2013-03-30 22:29:21 +08:00
|
|
|
float float32 = (float) CFP->getValueAPF().convertToFloat();
|
2015-06-10 00:29:34 +08:00
|
|
|
ConvertFloatToBytes(ptr, float32);
|
2012-05-05 04:18:50 +08:00
|
|
|
aggBuffer->addBytes(ptr, 4, Bytes);
|
|
|
|
} else if (Ty == Type::getDoubleTy(CPV->getContext())) {
|
|
|
|
double float64 = CFP->getValueAPF().convertToDouble();
|
2015-06-10 00:29:34 +08:00
|
|
|
ConvertDoubleToBytes(ptr, float64);
|
2012-05-05 04:18:50 +08:00
|
|
|
aggBuffer->addBytes(ptr, 8, Bytes);
|
2013-03-30 22:29:21 +08:00
|
|
|
} else {
|
2012-05-05 04:18:50 +08:00
|
|
|
llvm_unreachable("unsupported fp const type");
|
|
|
|
}
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
case Type::PointerTyID: {
|
2013-05-20 20:13:32 +08:00
|
|
|
if (const GlobalValue *GVar = dyn_cast<GlobalValue>(CPV)) {
|
2015-04-24 10:57:30 +08:00
|
|
|
aggBuffer->addSymbol(GVar, GVar);
|
2013-05-20 20:13:32 +08:00
|
|
|
} else if (const ConstantExpr *Cexpr = dyn_cast<ConstantExpr>(CPV)) {
|
|
|
|
const Value *v = Cexpr->stripPointerCasts();
|
2015-04-24 10:57:30 +08:00
|
|
|
aggBuffer->addSymbol(v, Cexpr);
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
2015-07-16 14:11:10 +08:00
|
|
|
unsigned int s = DL.getTypeAllocSize(CPV->getType());
|
2012-05-05 04:18:50 +08:00
|
|
|
aggBuffer->addZeros(s);
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
|
|
|
|
case Type::ArrayTyID:
|
|
|
|
case Type::VectorTyID:
|
|
|
|
case Type::StructTyID: {
|
2016-04-06 05:10:45 +08:00
|
|
|
if (isa<ConstantAggregate>(CPV) || isa<ConstantDataSequential>(CPV)) {
|
2015-07-16 14:11:10 +08:00
|
|
|
int ElementSize = DL.getTypeAllocSize(CPV->getType());
|
2012-05-05 04:18:50 +08:00
|
|
|
bufferAggregateConstant(CPV, aggBuffer);
|
2013-03-30 22:29:21 +08:00
|
|
|
if (Bytes > ElementSize)
|
|
|
|
aggBuffer->addZeros(Bytes - ElementSize);
|
|
|
|
} else if (isa<ConstantAggregateZero>(CPV))
|
2012-05-05 04:18:50 +08:00
|
|
|
aggBuffer->addZeros(Bytes);
|
|
|
|
else
|
|
|
|
llvm_unreachable("Unexpected Constant type");
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
|
|
|
|
default:
|
|
|
|
llvm_unreachable("unsupported type");
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2013-05-20 20:13:32 +08:00
|
|
|
void NVPTXAsmPrinter::bufferAggregateConstant(const Constant *CPV,
|
2012-05-05 04:18:50 +08:00
|
|
|
AggBuffer *aggBuffer) {
|
2015-07-16 14:11:10 +08:00
|
|
|
const DataLayout &DL = getDataLayout();
|
2012-05-05 04:18:50 +08:00
|
|
|
int Bytes;
|
|
|
|
|
2017-01-18 08:29:53 +08:00
|
|
|
// Integers of arbitrary width
|
|
|
|
if (const ConstantInt *CI = dyn_cast<ConstantInt>(CPV)) {
|
|
|
|
APInt Val = CI->getValue();
|
|
|
|
for (unsigned I = 0, E = DL.getTypeAllocSize(CPV->getType()); I < E; ++I) {
|
|
|
|
uint8_t Byte = Val.getLoBits(8).getZExtValue();
|
|
|
|
aggBuffer->addBytes(&Byte, 1, 1);
|
2017-04-19 01:14:21 +08:00
|
|
|
Val.lshrInPlace(8);
|
2017-01-18 08:29:53 +08:00
|
|
|
}
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
2012-05-05 04:18:50 +08:00
|
|
|
// Old constants
|
|
|
|
if (isa<ConstantArray>(CPV) || isa<ConstantVector>(CPV)) {
|
|
|
|
if (CPV->getNumOperands())
|
|
|
|
for (unsigned i = 0, e = CPV->getNumOperands(); i != e; ++i)
|
|
|
|
bufferLEByte(cast<Constant>(CPV->getOperand(i)), 0, aggBuffer);
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
|
|
|
if (const ConstantDataSequential *CDS =
|
2013-03-30 22:29:21 +08:00
|
|
|
dyn_cast<ConstantDataSequential>(CPV)) {
|
2012-05-05 04:18:50 +08:00
|
|
|
if (CDS->getNumElements())
|
|
|
|
for (unsigned i = 0; i < CDS->getNumElements(); ++i)
|
|
|
|
bufferLEByte(cast<Constant>(CDS->getElementAsConstant(i)), 0,
|
|
|
|
aggBuffer);
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
|
|
|
if (isa<ConstantStruct>(CPV)) {
|
|
|
|
if (CPV->getNumOperands()) {
|
|
|
|
StructType *ST = cast<StructType>(CPV->getType());
|
|
|
|
for (unsigned i = 0, e = CPV->getNumOperands(); i != e; ++i) {
|
2013-03-30 22:29:21 +08:00
|
|
|
if (i == (e - 1))
|
2015-07-16 14:11:10 +08:00
|
|
|
Bytes = DL.getStructLayout(ST)->getElementOffset(0) +
|
|
|
|
DL.getTypeAllocSize(ST) -
|
|
|
|
DL.getStructLayout(ST)->getElementOffset(i);
|
2012-05-05 04:18:50 +08:00
|
|
|
else
|
2015-07-16 14:11:10 +08:00
|
|
|
Bytes = DL.getStructLayout(ST)->getElementOffset(i + 1) -
|
|
|
|
DL.getStructLayout(ST)->getElementOffset(i);
|
2013-03-30 22:29:21 +08:00
|
|
|
bufferLEByte(cast<Constant>(CPV->getOperand(i)), Bytes, aggBuffer);
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
}
|
|
|
|
return;
|
|
|
|
}
|
2012-05-24 15:02:50 +08:00
|
|
|
llvm_unreachable("unsupported constant type in printAggregateConstant()");
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
// buildTypeNameMap - Run through symbol table looking for type names.
|
|
|
|
//
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
bool NVPTXAsmPrinter::ignoreLoc(const MachineInstr &MI) {
|
|
|
|
switch (MI.getOpcode()) {
|
2012-05-05 04:18:50 +08:00
|
|
|
default:
|
|
|
|
return false;
|
2013-03-30 22:29:21 +08:00
|
|
|
case NVPTX::CallArgBeginInst:
|
|
|
|
case NVPTX::CallArgEndInst0:
|
|
|
|
case NVPTX::CallArgEndInst1:
|
|
|
|
case NVPTX::CallArgF32:
|
|
|
|
case NVPTX::CallArgF64:
|
|
|
|
case NVPTX::CallArgI16:
|
|
|
|
case NVPTX::CallArgI32:
|
|
|
|
case NVPTX::CallArgI32imm:
|
|
|
|
case NVPTX::CallArgI64:
|
|
|
|
case NVPTX::CallArgParam:
|
|
|
|
case NVPTX::CallVoidInst:
|
|
|
|
case NVPTX::CallVoidInstReg:
|
|
|
|
case NVPTX::Callseq_End:
|
2012-05-05 04:18:50 +08:00
|
|
|
case NVPTX::CallVoidInstReg64:
|
2013-03-30 22:29:21 +08:00
|
|
|
case NVPTX::DeclareParamInst:
|
|
|
|
case NVPTX::DeclareRetMemInst:
|
|
|
|
case NVPTX::DeclareRetRegInst:
|
|
|
|
case NVPTX::DeclareRetScalarInst:
|
|
|
|
case NVPTX::DeclareScalarParamInst:
|
|
|
|
case NVPTX::DeclareScalarRegInst:
|
|
|
|
case NVPTX::StoreParamF32:
|
|
|
|
case NVPTX::StoreParamF64:
|
|
|
|
case NVPTX::StoreParamI16:
|
|
|
|
case NVPTX::StoreParamI32:
|
|
|
|
case NVPTX::StoreParamI64:
|
|
|
|
case NVPTX::StoreParamI8:
|
|
|
|
case NVPTX::StoreRetvalF32:
|
|
|
|
case NVPTX::StoreRetvalF64:
|
|
|
|
case NVPTX::StoreRetvalI16:
|
|
|
|
case NVPTX::StoreRetvalI32:
|
|
|
|
case NVPTX::StoreRetvalI64:
|
|
|
|
case NVPTX::StoreRetvalI8:
|
|
|
|
case NVPTX::LastCallArgF32:
|
|
|
|
case NVPTX::LastCallArgF64:
|
|
|
|
case NVPTX::LastCallArgI16:
|
|
|
|
case NVPTX::LastCallArgI32:
|
|
|
|
case NVPTX::LastCallArgI32imm:
|
|
|
|
case NVPTX::LastCallArgI64:
|
|
|
|
case NVPTX::LastCallArgParam:
|
|
|
|
case NVPTX::LoadParamMemF32:
|
|
|
|
case NVPTX::LoadParamMemF64:
|
|
|
|
case NVPTX::LoadParamMemI16:
|
|
|
|
case NVPTX::LoadParamMemI32:
|
|
|
|
case NVPTX::LoadParamMemI64:
|
|
|
|
case NVPTX::LoadParamMemI8:
|
|
|
|
case NVPTX::PrototypeInst:
|
|
|
|
case NVPTX::DBG_VALUE:
|
2012-05-05 04:18:50 +08:00
|
|
|
return true;
|
|
|
|
}
|
|
|
|
return false;
|
|
|
|
}
|
|
|
|
|
2015-04-29 01:18:30 +08:00
|
|
|
/// lowerConstantForGV - Return an MCExpr for the given Constant. This is mostly
|
|
|
|
/// a copy from AsmPrinter::lowerConstant, except customized to only handle
|
|
|
|
/// expressions that are representable in PTX and create
|
|
|
|
/// NVPTXGenericMCSymbolRefExpr nodes for addrspacecast instructions.
|
|
|
|
const MCExpr *
|
|
|
|
NVPTXAsmPrinter::lowerConstantForGV(const Constant *CV, bool ProcessingGeneric) {
|
|
|
|
MCContext &Ctx = OutContext;
|
|
|
|
|
|
|
|
if (CV->isNullValue() || isa<UndefValue>(CV))
|
2015-05-30 09:25:56 +08:00
|
|
|
return MCConstantExpr::create(0, Ctx);
|
2015-04-29 01:18:30 +08:00
|
|
|
|
|
|
|
if (const ConstantInt *CI = dyn_cast<ConstantInt>(CV))
|
2015-05-30 09:25:56 +08:00
|
|
|
return MCConstantExpr::create(CI->getZExtValue(), Ctx);
|
2015-04-29 01:18:30 +08:00
|
|
|
|
|
|
|
if (const GlobalValue *GV = dyn_cast<GlobalValue>(CV)) {
|
|
|
|
const MCSymbolRefExpr *Expr =
|
2015-05-30 09:25:56 +08:00
|
|
|
MCSymbolRefExpr::create(getSymbol(GV), Ctx);
|
2015-04-29 01:18:30 +08:00
|
|
|
if (ProcessingGeneric) {
|
2015-05-30 09:25:56 +08:00
|
|
|
return NVPTXGenericMCSymbolRefExpr::create(Expr, Ctx);
|
2015-04-29 01:18:30 +08:00
|
|
|
} else {
|
|
|
|
return Expr;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
const ConstantExpr *CE = dyn_cast<ConstantExpr>(CV);
|
|
|
|
if (!CE) {
|
|
|
|
llvm_unreachable("Unknown constant value to lower!");
|
|
|
|
}
|
|
|
|
|
|
|
|
switch (CE->getOpcode()) {
|
|
|
|
default:
|
|
|
|
// If the code isn't optimized, there may be outstanding folding
|
|
|
|
// opportunities. Attempt to fold the expression using DataLayout as a
|
|
|
|
// last resort before giving up.
|
2016-07-29 11:27:26 +08:00
|
|
|
if (Constant *C = ConstantFoldConstant(CE, getDataLayout()))
|
|
|
|
if (C && C != CE)
|
2015-04-29 01:18:30 +08:00
|
|
|
return lowerConstantForGV(C, ProcessingGeneric);
|
|
|
|
|
|
|
|
// Otherwise report the problem to the user.
|
|
|
|
{
|
|
|
|
std::string S;
|
|
|
|
raw_string_ostream OS(S);
|
|
|
|
OS << "Unsupported expression in static initializer: ";
|
|
|
|
CE->printAsOperand(OS, /*PrintType=*/false,
|
|
|
|
!MF ? nullptr : MF->getFunction()->getParent());
|
|
|
|
report_fatal_error(OS.str());
|
|
|
|
}
|
|
|
|
|
|
|
|
case Instruction::AddrSpaceCast: {
|
|
|
|
// Strip the addrspacecast and pass along the operand
|
|
|
|
PointerType *DstTy = cast<PointerType>(CE->getType());
|
|
|
|
if (DstTy->getAddressSpace() == 0) {
|
|
|
|
return lowerConstantForGV(cast<const Constant>(CE->getOperand(0)), true);
|
|
|
|
}
|
|
|
|
std::string S;
|
|
|
|
raw_string_ostream OS(S);
|
|
|
|
OS << "Unsupported expression in static initializer: ";
|
|
|
|
CE->printAsOperand(OS, /*PrintType=*/ false,
|
2017-01-10 06:16:51 +08:00
|
|
|
!MF ? nullptr : MF->getFunction()->getParent());
|
2015-04-29 01:18:30 +08:00
|
|
|
report_fatal_error(OS.str());
|
|
|
|
}
|
|
|
|
|
|
|
|
case Instruction::GetElementPtr: {
|
2015-07-16 14:11:10 +08:00
|
|
|
const DataLayout &DL = getDataLayout();
|
2015-04-29 01:18:30 +08:00
|
|
|
|
|
|
|
// Generate a symbolic expression for the byte address
|
|
|
|
APInt OffsetAI(DL.getPointerTypeSizeInBits(CE->getType()), 0);
|
|
|
|
cast<GEPOperator>(CE)->accumulateConstantOffset(DL, OffsetAI);
|
|
|
|
|
|
|
|
const MCExpr *Base = lowerConstantForGV(CE->getOperand(0),
|
|
|
|
ProcessingGeneric);
|
|
|
|
if (!OffsetAI)
|
|
|
|
return Base;
|
|
|
|
|
|
|
|
int64_t Offset = OffsetAI.getSExtValue();
|
2015-05-30 09:25:56 +08:00
|
|
|
return MCBinaryExpr::createAdd(Base, MCConstantExpr::create(Offset, Ctx),
|
2015-04-29 01:18:30 +08:00
|
|
|
Ctx);
|
|
|
|
}
|
|
|
|
|
|
|
|
case Instruction::Trunc:
|
|
|
|
// We emit the value and depend on the assembler to truncate the generated
|
|
|
|
// expression properly. This is important for differences between
|
|
|
|
// blockaddress labels. Since the two labels are in the same function, it
|
|
|
|
// is reasonable to treat their delta as a 32-bit value.
|
2016-08-18 04:30:52 +08:00
|
|
|
LLVM_FALLTHROUGH;
|
2015-04-29 01:18:30 +08:00
|
|
|
case Instruction::BitCast:
|
|
|
|
return lowerConstantForGV(CE->getOperand(0), ProcessingGeneric);
|
|
|
|
|
|
|
|
case Instruction::IntToPtr: {
|
2015-07-16 14:11:10 +08:00
|
|
|
const DataLayout &DL = getDataLayout();
|
2015-04-29 01:18:30 +08:00
|
|
|
|
|
|
|
// Handle casts to pointers by changing them into casts to the appropriate
|
|
|
|
// integer type. This promotes constant folding and simplifies this code.
|
|
|
|
Constant *Op = CE->getOperand(0);
|
|
|
|
Op = ConstantExpr::getIntegerCast(Op, DL.getIntPtrType(CV->getType()),
|
|
|
|
false/*ZExt*/);
|
|
|
|
return lowerConstantForGV(Op, ProcessingGeneric);
|
|
|
|
}
|
|
|
|
|
|
|
|
case Instruction::PtrToInt: {
|
2015-07-16 14:11:10 +08:00
|
|
|
const DataLayout &DL = getDataLayout();
|
2015-04-29 01:18:30 +08:00
|
|
|
|
|
|
|
// Support only foldable casts to/from pointers that can be eliminated by
|
|
|
|
// changing the pointer to the appropriately sized integer type.
|
|
|
|
Constant *Op = CE->getOperand(0);
|
|
|
|
Type *Ty = CE->getType();
|
|
|
|
|
|
|
|
const MCExpr *OpExpr = lowerConstantForGV(Op, ProcessingGeneric);
|
|
|
|
|
|
|
|
// We can emit the pointer value into this slot if the slot is an
|
|
|
|
// integer slot equal to the size of the pointer.
|
|
|
|
if (DL.getTypeAllocSize(Ty) == DL.getTypeAllocSize(Op->getType()))
|
|
|
|
return OpExpr;
|
|
|
|
|
|
|
|
// Otherwise the pointer is smaller than the resultant integer, mask off
|
|
|
|
// the high bits so we are sure to get a proper truncation if the input is
|
|
|
|
// a constant expr.
|
|
|
|
unsigned InBits = DL.getTypeAllocSizeInBits(Op->getType());
|
2015-05-30 09:25:56 +08:00
|
|
|
const MCExpr *MaskExpr = MCConstantExpr::create(~0ULL >> (64-InBits), Ctx);
|
|
|
|
return MCBinaryExpr::createAnd(OpExpr, MaskExpr, Ctx);
|
2015-04-29 01:18:30 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
// The MC library also has a right-shift operator, but it isn't consistently
|
|
|
|
// signed or unsigned between different targets.
|
|
|
|
case Instruction::Add: {
|
|
|
|
const MCExpr *LHS = lowerConstantForGV(CE->getOperand(0), ProcessingGeneric);
|
|
|
|
const MCExpr *RHS = lowerConstantForGV(CE->getOperand(1), ProcessingGeneric);
|
|
|
|
switch (CE->getOpcode()) {
|
|
|
|
default: llvm_unreachable("Unknown binary operator constant cast expr");
|
2015-05-30 09:25:56 +08:00
|
|
|
case Instruction::Add: return MCBinaryExpr::createAdd(LHS, RHS, Ctx);
|
2015-04-29 01:18:30 +08:00
|
|
|
}
|
|
|
|
}
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
// Copy of MCExpr::print customized for NVPTX
|
|
|
|
void NVPTXAsmPrinter::printMCExpr(const MCExpr &Expr, raw_ostream &OS) {
|
|
|
|
switch (Expr.getKind()) {
|
|
|
|
case MCExpr::Target:
|
2015-06-09 08:31:39 +08:00
|
|
|
return cast<MCTargetExpr>(&Expr)->printImpl(OS, MAI);
|
2015-04-29 01:18:30 +08:00
|
|
|
case MCExpr::Constant:
|
|
|
|
OS << cast<MCConstantExpr>(Expr).getValue();
|
|
|
|
return;
|
|
|
|
|
|
|
|
case MCExpr::SymbolRef: {
|
|
|
|
const MCSymbolRefExpr &SRE = cast<MCSymbolRefExpr>(Expr);
|
|
|
|
const MCSymbol &Sym = SRE.getSymbol();
|
2015-06-09 08:31:39 +08:00
|
|
|
Sym.print(OS, MAI);
|
2015-04-29 01:18:30 +08:00
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
|
|
|
case MCExpr::Unary: {
|
|
|
|
const MCUnaryExpr &UE = cast<MCUnaryExpr>(Expr);
|
|
|
|
switch (UE.getOpcode()) {
|
|
|
|
case MCUnaryExpr::LNot: OS << '!'; break;
|
|
|
|
case MCUnaryExpr::Minus: OS << '-'; break;
|
|
|
|
case MCUnaryExpr::Not: OS << '~'; break;
|
|
|
|
case MCUnaryExpr::Plus: OS << '+'; break;
|
|
|
|
}
|
|
|
|
printMCExpr(*UE.getSubExpr(), OS);
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
|
|
|
|
case MCExpr::Binary: {
|
|
|
|
const MCBinaryExpr &BE = cast<MCBinaryExpr>(Expr);
|
|
|
|
|
|
|
|
// Only print parens around the LHS if it is non-trivial.
|
|
|
|
if (isa<MCConstantExpr>(BE.getLHS()) || isa<MCSymbolRefExpr>(BE.getLHS()) ||
|
|
|
|
isa<NVPTXGenericMCSymbolRefExpr>(BE.getLHS())) {
|
|
|
|
printMCExpr(*BE.getLHS(), OS);
|
|
|
|
} else {
|
|
|
|
OS << '(';
|
|
|
|
printMCExpr(*BE.getLHS(), OS);
|
|
|
|
OS<< ')';
|
|
|
|
}
|
|
|
|
|
|
|
|
switch (BE.getOpcode()) {
|
|
|
|
case MCBinaryExpr::Add:
|
|
|
|
// Print "X-42" instead of "X+-42".
|
|
|
|
if (const MCConstantExpr *RHSC = dyn_cast<MCConstantExpr>(BE.getRHS())) {
|
|
|
|
if (RHSC->getValue() < 0) {
|
|
|
|
OS << RHSC->getValue();
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
OS << '+';
|
|
|
|
break;
|
|
|
|
default: llvm_unreachable("Unhandled binary operator");
|
|
|
|
}
|
|
|
|
|
|
|
|
// Only print parens around the LHS if it is non-trivial.
|
|
|
|
if (isa<MCConstantExpr>(BE.getRHS()) || isa<MCSymbolRefExpr>(BE.getRHS())) {
|
|
|
|
printMCExpr(*BE.getRHS(), OS);
|
|
|
|
} else {
|
|
|
|
OS << '(';
|
|
|
|
printMCExpr(*BE.getRHS(), OS);
|
|
|
|
OS << ')';
|
|
|
|
}
|
|
|
|
return;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
llvm_unreachable("Invalid expression kind!");
|
|
|
|
}
|
|
|
|
|
2013-08-24 09:17:23 +08:00
|
|
|
/// PrintAsmOperand - Print out an operand for an inline asm expression.
|
|
|
|
///
|
|
|
|
bool NVPTXAsmPrinter::PrintAsmOperand(const MachineInstr *MI, unsigned OpNo,
|
|
|
|
unsigned AsmVariant,
|
|
|
|
const char *ExtraCode, raw_ostream &O) {
|
|
|
|
if (ExtraCode && ExtraCode[0]) {
|
|
|
|
if (ExtraCode[1] != 0)
|
|
|
|
return true; // Unknown modifier.
|
|
|
|
|
|
|
|
switch (ExtraCode[0]) {
|
|
|
|
default:
|
|
|
|
// See if this is a generic print operand
|
|
|
|
return AsmPrinter::PrintAsmOperand(MI, OpNo, AsmVariant, ExtraCode, O);
|
|
|
|
case 'r':
|
|
|
|
break;
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
printOperand(MI, OpNo, O);
|
|
|
|
|
|
|
|
return false;
|
|
|
|
}
|
|
|
|
|
|
|
|
bool NVPTXAsmPrinter::PrintAsmMemoryOperand(
|
|
|
|
const MachineInstr *MI, unsigned OpNo, unsigned AsmVariant,
|
|
|
|
const char *ExtraCode, raw_ostream &O) {
|
|
|
|
if (ExtraCode && ExtraCode[0])
|
|
|
|
return true; // Unknown modifier
|
|
|
|
|
|
|
|
O << '[';
|
|
|
|
printMemOperand(MI, OpNo, O);
|
|
|
|
O << ']';
|
|
|
|
|
|
|
|
return false;
|
|
|
|
}
|
|
|
|
|
|
|
|
void NVPTXAsmPrinter::printOperand(const MachineInstr *MI, int opNum,
|
|
|
|
raw_ostream &O, const char *Modifier) {
|
|
|
|
const MachineOperand &MO = MI->getOperand(opNum);
|
|
|
|
switch (MO.getType()) {
|
|
|
|
case MachineOperand::MO_Register:
|
|
|
|
if (TargetRegisterInfo::isPhysicalRegister(MO.getReg())) {
|
|
|
|
if (MO.getReg() == NVPTX::VRDepot)
|
|
|
|
O << DEPOTNAME << getFunctionNumber();
|
|
|
|
else
|
|
|
|
O << NVPTXInstPrinter::getRegisterName(MO.getReg());
|
|
|
|
} else {
|
2013-10-11 20:39:36 +08:00
|
|
|
emitVirtualRegister(MO.getReg(), O);
|
2013-08-24 09:17:23 +08:00
|
|
|
}
|
|
|
|
return;
|
|
|
|
|
|
|
|
case MachineOperand::MO_Immediate:
|
|
|
|
if (!Modifier)
|
|
|
|
O << MO.getImm();
|
|
|
|
else if (strstr(Modifier, "vec") == Modifier)
|
|
|
|
printVecModifiedImmediate(MO, Modifier, O);
|
|
|
|
else
|
|
|
|
llvm_unreachable(
|
|
|
|
"Don't know how to handle modifier on immediate operand");
|
|
|
|
return;
|
|
|
|
|
|
|
|
case MachineOperand::MO_FPImmediate:
|
|
|
|
printFPConstant(MO.getFPImm(), O);
|
|
|
|
break;
|
|
|
|
|
|
|
|
case MachineOperand::MO_GlobalAddress:
|
2015-06-09 08:31:39 +08:00
|
|
|
getSymbol(MO.getGlobal())->print(O, MAI);
|
2013-08-24 09:17:23 +08:00
|
|
|
break;
|
|
|
|
|
|
|
|
case MachineOperand::MO_MachineBasicBlock:
|
2015-06-09 08:31:39 +08:00
|
|
|
MO.getMBB()->getSymbol()->print(O, MAI);
|
2013-08-24 09:17:23 +08:00
|
|
|
return;
|
|
|
|
|
|
|
|
default:
|
|
|
|
llvm_unreachable("Operand type not supported.");
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
|
|
|
void NVPTXAsmPrinter::printMemOperand(const MachineInstr *MI, int opNum,
|
|
|
|
raw_ostream &O, const char *Modifier) {
|
|
|
|
printOperand(MI, opNum, O);
|
|
|
|
|
2017-01-10 06:16:51 +08:00
|
|
|
if (Modifier && strcmp(Modifier, "add") == 0) {
|
2013-08-24 09:17:23 +08:00
|
|
|
O << ", ";
|
|
|
|
printOperand(MI, opNum + 1, O);
|
|
|
|
} else {
|
|
|
|
if (MI->getOperand(opNum + 1).isImm() &&
|
|
|
|
MI->getOperand(opNum + 1).getImm() == 0)
|
|
|
|
return; // don't print ',0' or '+0'
|
|
|
|
O << "+";
|
|
|
|
printOperand(MI, opNum + 1, O);
|
|
|
|
}
|
|
|
|
}
|
|
|
|
|
2012-05-05 04:18:50 +08:00
|
|
|
void NVPTXAsmPrinter::emitSrcInText(StringRef filename, unsigned line) {
|
|
|
|
std::stringstream temp;
|
2015-03-30 23:42:36 +08:00
|
|
|
LineReader *reader = this->getReader(filename);
|
2012-05-05 04:18:50 +08:00
|
|
|
temp << "\n//";
|
|
|
|
temp << filename.str();
|
|
|
|
temp << ":";
|
|
|
|
temp << line;
|
|
|
|
temp << " ";
|
|
|
|
temp << reader->readLine(line);
|
|
|
|
temp << "\n";
|
2015-04-25 03:11:51 +08:00
|
|
|
this->OutStreamer->EmitRawText(temp.str());
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
2016-06-09 03:09:22 +08:00
|
|
|
LineReader *NVPTXAsmPrinter::getReader(const std::string &filename) {
|
2014-04-25 13:30:21 +08:00
|
|
|
if (!reader) {
|
2013-03-30 22:29:21 +08:00
|
|
|
reader = new LineReader(filename);
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
if (reader->fileName() != filename) {
|
|
|
|
delete reader;
|
2013-03-30 22:29:21 +08:00
|
|
|
reader = new LineReader(filename);
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
|
|
|
|
return reader;
|
|
|
|
}
|
|
|
|
|
2013-03-30 22:29:21 +08:00
|
|
|
std::string LineReader::readLine(unsigned lineNum) {
|
2012-05-05 04:18:50 +08:00
|
|
|
if (lineNum < theCurLine) {
|
|
|
|
theCurLine = 0;
|
2013-03-30 22:29:21 +08:00
|
|
|
fstr.seekg(0, std::ios::beg);
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|
|
|
|
while (theCurLine < lineNum) {
|
2013-03-30 22:29:21 +08:00
|
|
|
fstr.getline(buff, 500);
|
2012-05-05 04:18:50 +08:00
|
|
|
theCurLine++;
|
|
|
|
}
|
|
|
|
return buff;
|
|
|
|
}
|
|
|
|
|
|
|
|
// Force static initialization.
|
|
|
|
extern "C" void LLVMInitializeNVPTXAsmPrinter() {
|
2016-10-10 07:00:34 +08:00
|
|
|
RegisterAsmPrinter<NVPTXAsmPrinter> X(getTheNVPTXTarget32());
|
|
|
|
RegisterAsmPrinter<NVPTXAsmPrinter> Y(getTheNVPTXTarget64());
|
2012-05-05 04:18:50 +08:00
|
|
|
}
|