maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

Revert "Temporarily revert "[DEBUG] Initial adaptation of NVPTX target for debug info emission."" 

This reapplies commits: r330271, r330592, r330779.

    [DEBUG] Initial adaptation of NVPTX target for debug info emission.

    Summary:
    Patch adds initial emission of the debug info for NVPTX target.
    Currently, only .file and .loc directives are emitted, everything else is
    commented out to not break the compilation of Cuda.

llvm-svn: 332689
This commit is contained in:
Eric Christopher 2018-05-18 03:13:08 +00:00
parent 3aede9283c
commit 68f2218e1e
16 changed files with 9151 additions and 309 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

22
llvm/lib/CodeGen/AsmPrinter/DwarfDebug.cpp
View File

@ -312,7 +312,11 @@ DwarfDebug::DwarfDebug(AsmPrinter *A, Module *M)
} else
TheAccelTableKind = AccelTables;
UseInlineStrings = DwarfInlinedStrings == Enable;
if (DwarfInlinedStrings == Default)
UseInlineStrings = TT.isNVPTX();
else
UseInlineStrings = DwarfInlinedStrings == Enable;
HasAppleExtensionAttributes = tuneForLLDB();
// Handle split DWARF.
@ -327,14 +331,18 @@ DwarfDebug::DwarfDebug(AsmPrinter *A, Module *M)
unsigned DwarfVersionNumber = Asm->TM.Options.MCOptions.DwarfVersion;
unsigned DwarfVersion = DwarfVersionNumber ? DwarfVersionNumber
: MMI->getModule()->getDwarfVersion();
// Use dwarf 4 by default if nothing is requested.
DwarfVersion = DwarfVersion ? DwarfVersion : dwarf::DWARF_VERSION;
// Use dwarf 4 by default if nothing is requested. For NVPTX, use dwarf 2.
DwarfVersion =
TT.isNVPTX() ? 2 : (DwarfVersion ? DwarfVersion : dwarf::DWARF_VERSION);
UsePubSections = !NoDwarfPubSections;
UseRangesSection = !NoDwarfRangesSection;
UsePubSections = !NoDwarfPubSections && !TT.isNVPTX();
UseRangesSection = !NoDwarfRangesSection && !TT.isNVPTX();
// Use sections as references.
UseSectionsAsReferences = DwarfSectionsAsReferences == Enable;
// Use sections as references. Force for NVPTX.
if (DwarfSectionsAsReferences == Default)
UseSectionsAsReferences = TT.isNVPTX();
else
UseSectionsAsReferences = DwarfSectionsAsReferences == Enable;
// Work around a GDB bug. GDB doesn't support the standard opcode;
// SCE doesn't support GNU's; LLDB prefers the standard opcode, which

1
llvm/lib/Target/NVPTX/MCTargetDesc/CMakeLists.txt
View File

@ -1,4 +1,5 @@
add_llvm_library(LLVMNVPTXDesc
NVPTXMCAsmInfo.cpp
NVPTXMCTargetDesc.cpp
NVPTXTargetStreamer.cpp
)

26
llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp
View File

@ -13,16 +13,9 @@
#include "NVPTXMCAsmInfo.h"
#include "llvm/ADT/Triple.h"
#include "llvm/Support/CommandLine.h"
using namespace llvm;
// -debug-compile - Command line option to inform opt and llc passes to
// compile for debugging
static cl::opt<bool> CompileForDebugging("debug-compile",
cl::desc("Compile for debugging"),
cl::Hidden, cl::init(false));
void NVPTXMCAsmInfo::anchor() {}
NVPTXMCAsmInfo::NVPTXMCAsmInfo(const Triple &TheTriple) {
@ -37,7 +30,7 @@ NVPTXMCAsmInfo::NVPTXMCAsmInfo(const Triple &TheTriple) {
InlineAsmStart = " begin inline asm";
InlineAsmEnd = " end inline asm";
SupportsDebugInformation = CompileForDebugging;
SupportsDebugInformation = true;
// PTX does not allow .align on functions.
HasFunctionAlignment = false;
HasDotTypeDotSizeDirective = false;
@ -45,13 +38,16 @@ NVPTXMCAsmInfo::NVPTXMCAsmInfo(const Triple &TheTriple) {
HiddenDeclarationVisibilityAttr = HiddenVisibilityAttr = MCSA_Invalid;
ProtectedVisibilityAttr = MCSA_Invalid;
Data8bitsDirective = " .b8 ";
Data16bitsDirective = " .b16 ";
Data32bitsDirective = " .b32 ";
Data64bitsDirective = " .b64 ";
ZeroDirective = " .b8";
AsciiDirective = " .b8";
AscizDirective = " .b8";
// FIXME: remove comment once debug info is properly supported.
Data8bitsDirective = "// .b8 ";
Data16bitsDirective = nullptr; // not supported
Data32bitsDirective = "// .b32 ";
Data64bitsDirective = "// .b64 ";
ZeroDirective = "// .b8";
AsciiDirective = nullptr; // not supported
AscizDirective = nullptr; // not supported
SupportsQuotedNames = false;
SupportsExtendedDwarfLocDirective = false;
// @TODO: Can we just disable this?
WeakDirective = "\t// .weak\t";

11
llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.h
View File

@ -25,6 +25,17 @@ class NVPTXMCAsmInfo : public MCAsmInfo {
public:
explicit NVPTXMCAsmInfo(const Triple &TheTriple);
/// Return true if the .section directive should be omitted when
/// emitting \p SectionName. For example:
///
/// shouldOmitSectionDirective(".text")
///
/// returns false => .section .text,#alloc,#execinstr
/// returns true => .text
bool shouldOmitSectionDirective(StringRef SectionName) const override {
return true;
}
};
} // namespace llvm

12
llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp
View File

@ -11,9 +11,10 @@
//
//===----------------------------------------------------------------------===//
#include "NVPTXMCTargetDesc.h"
#include "InstPrinter/NVPTXInstPrinter.h"
#include "NVPTXMCAsmInfo.h"
#include "NVPTXMCTargetDesc.h"
#include "NVPTXTargetStreamer.h"
#include "llvm/MC/MCInstrInfo.h"
#include "llvm/MC/MCRegisterInfo.h"
#include "llvm/MC/MCSubtargetInfo.h"
@ -58,6 +59,12 @@ static MCInstPrinter *createNVPTXMCInstPrinter(const Triple &T,
return nullptr;
}
static MCTargetStreamer *createTargetAsmStreamer(MCStreamer &S,
formatted_raw_ostream &,
MCInstPrinter *, bool) {
return new NVPTXTargetStreamer(S);
}
// Force static initialization.
extern "C" void LLVMInitializeNVPTXTargetMC() {
for (Target *T : {&getTheNVPTXTarget32(), &getTheNVPTXTarget64()}) {
@ -75,5 +82,8 @@ extern "C" void LLVMInitializeNVPTXTargetMC() {
// Register the MCInstPrinter.
TargetRegistry::RegisterMCInstPrinter(*T, createNVPTXMCInstPrinter);
// Register the MCTargetStreamer.
TargetRegistry::RegisterAsmTargetStreamer(*T, createTargetAsmStreamer);
}
}

94
llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp Normal file
View File

@ -0,0 +1,94 @@
//=====- NVPTXTargetStreamer.cpp - NVPTXTargetStreamer class ------------=====//
//
// The LLVM Compiler Infrastructure
//
// This file is distributed under the University of Illinois Open Source
// License. See LICENSE.TXT for details.
//
//===----------------------------------------------------------------------===//
//
// This file implements the NVPTXTargetStreamer class.
//
//===----------------------------------------------------------------------===//
#include "NVPTXTargetStreamer.h"
#include "llvm/MC/MCAsmInfo.h"
#include "llvm/MC/MCContext.h"
#include "llvm/MC/MCObjectFileInfo.h"
using namespace llvm;
//
// NVPTXTargetStreamer Implemenation
//
NVPTXTargetStreamer::NVPTXTargetStreamer(MCStreamer &S) : MCTargetStreamer(S) {}
NVPTXTargetStreamer::~NVPTXTargetStreamer() = default;
void NVPTXTargetStreamer::emitDwarfFileDirective(StringRef Directive) {
DwarfFiles.emplace_back(Directive);
}
static bool isDwarfSection(const MCObjectFileInfo *FI,
const MCSection *Section) {
// FIXME: the checks for the DWARF sections are very fragile and should be
// fixed up in a followup patch.
if (!Section || Section->getKind().isText() ||
Section->getKind().isWriteable())
return false;
return Section == FI->getDwarfAbbrevSection() ||
Section == FI->getDwarfInfoSection() ||
Section == FI->getDwarfMacinfoSection() ||
Section == FI->getDwarfFrameSection() ||
Section == FI->getDwarfAddrSection() ||
Section == FI->getDwarfRangesSection() ||
Section == FI->getDwarfARangesSection() ||
Section == FI->getDwarfLocSection() ||
Section == FI->getDwarfStrSection() ||
Section == FI->getDwarfLineSection() ||
Section == FI->getDwarfStrOffSection() ||
Section == FI->getDwarfLineStrSection() ||
Section == FI->getDwarfPubNamesSection() ||
Section == FI->getDwarfPubTypesSection() ||
Section == FI->getDwarfSwiftASTSection() ||
Section == FI->getDwarfTypesDWOSection() ||
Section == FI->getDwarfAbbrevDWOSection() ||
Section == FI->getDwarfAccelObjCSection() ||
Section == FI->getDwarfAccelNamesSection() ||
Section == FI->getDwarfAccelTypesSection() ||
Section == FI->getDwarfAccelNamespaceSection() ||
Section == FI->getDwarfLocDWOSection() ||
Section == FI->getDwarfStrDWOSection() ||
Section == FI->getDwarfCUIndexSection() ||
Section == FI->getDwarfInfoDWOSection() ||
Section == FI->getDwarfLineDWOSection() ||
Section == FI->getDwarfTUIndexSection() ||
Section == FI->getDwarfStrOffDWOSection() ||
Section == FI->getDwarfDebugNamesSection() ||
Section == FI->getDwarfDebugInlineSection() ||
Section == FI->getDwarfGnuPubNamesSection() ||
Section == FI->getDwarfGnuPubTypesSection();
}
void NVPTXTargetStreamer::changeSection(const MCSection *CurSection,
MCSection *Section,
const MCExpr *SubSection,
raw_ostream &OS) {
assert(!SubSection && "SubSection is not null!");
const MCObjectFileInfo *FI = getStreamer().getContext().getObjectFileInfo();
// FIXME: remove comment once debug info is properly supported.
// Emit closing brace for DWARF sections only.
if (isDwarfSection(FI, CurSection))
OS << "//\t}\n";
if (isDwarfSection(FI, Section)) {
// Emit DWARF .file directives in the outermost scope.
for (const std::string &S : DwarfFiles)
getStreamer().EmitRawText(S.data());
DwarfFiles.clear();
OS << "//\t.section";
Section->PrintSwitchToSection(*getStreamer().getContext().getAsmInfo(),
FI->getTargetTriple(), OS, SubSection);
// DWARF sections are enclosed into braces - emit the open one.
OS << "//\t{\n";
}
}

46
llvm/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h Normal file
View File

@ -0,0 +1,46 @@
//=====-- NVPTXTargetStreamer.h - NVPTX Target Streamer ------*- C++ -*--=====//
//
// The LLVM Compiler Infrastructure
//
// This file is distributed under the University of Illinois Open Source
// License. See LICENSE.TXT for details.
//
//===----------------------------------------------------------------------===//
#ifndef LLVM_LIB_TARGET_NVPTX_MCTARGETDESC_NVPTXTARGETSTREAMER_H
#define LLVM_LIB_TARGET_NVPTX_MCTARGETDESC_NVPTXTARGETSTREAMER_H
#include "llvm/MC/MCStreamer.h"
namespace llvm {
class MCSection;
/// Implments NVPTX-specific streamer.
class NVPTXTargetStreamer : public MCTargetStreamer {
private:
SmallVector<std::string, 4> DwarfFiles;
public:
NVPTXTargetStreamer(MCStreamer &S);
~NVPTXTargetStreamer() override;
/// Record DWARF file directives for later output.
/// According to PTX ISA, CUDA Toolkit documentation, 11.5.3. Debugging
/// Directives: .file
/// (http://docs.nvidia.com/cuda/parallel-thread-execution/index.html#debugging-directives-file),
/// The .file directive is allowed only in the outermost scope, i.e., at the
/// same level as kernel and device function declarations. Also, the order of
/// the .loc and .file directive does not matter, .file directives may follow
/// the .loc directives where the file is referenced.
/// LLVM emits .file directives immediately the location debug info is
/// emitted, i.e. they may be emitted inside functions. We gather all these
/// directives and emit them outside of the sections and, thus, outside of the
/// functions.
void emitDwarfFileDirective(StringRef Directive) override;
void changeSection(const MCSection *CurSection, MCSection *Section,
const MCExpr *SubSection, raw_ostream &OS) override;
};
} // end namespace llvm
#endif

152
llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
View File

@ -93,16 +93,6 @@ using namespace llvm;
#define DEPOTNAME "__local_depot"
static cl::opt<bool>
EmitLineNumbers("nvptx-emit-line-numbers", cl::Hidden,
cl::desc("NVPTX Specific: Emit Line numbers even without -G"),
cl::init(true));
static cl::opt<bool>
InterleaveSrc("nvptx-emit-src", cl::ZeroOrMore, cl::Hidden,
cl::desc("NVPTX Specific: Emit source line in ptx file"),
cl::init(false));
/// DiscoverDependentGlobals - Return a set of GlobalVariables on which \p V
/// depends.
static void
@ -151,56 +141,7 @@ VisitGlobalVariableForEmission(const GlobalVariable *GV,
Visiting.erase(GV);
}
void NVPTXAsmPrinter::emitLineNumberAsDotLoc(const MachineInstr &MI) {
if (!EmitLineNumbers)
return;
if (ignoreLoc(MI))
return;
const DebugLoc &curLoc = MI.getDebugLoc();
if (!prevDebugLoc && !curLoc)
return;
if (prevDebugLoc == curLoc)
return;
prevDebugLoc = curLoc;
if (!curLoc)
return;
auto *Scope = cast_or_null<DIScope>(curLoc.getScope());
if (!Scope)
return;
StringRef fileName(Scope->getFilename());
StringRef dirName(Scope->getDirectory());
SmallString<128> FullPathName = dirName;
if (!dirName.empty() && !sys::path::is_absolute(fileName)) {
sys::path::append(FullPathName, fileName);
fileName = FullPathName;
}
if (filenameMap.find(fileName) == filenameMap.end())
return;
// Emit the line from the source file.
if (InterleaveSrc)
this->emitSrcInText(fileName, curLoc.getLine());
std::stringstream temp;
temp << "\t.loc " << filenameMap[fileName] << " " << curLoc.getLine()
<< " " << curLoc.getCol();
OutStreamer->EmitRawText(temp.str());
}
void NVPTXAsmPrinter::EmitInstruction(const MachineInstr *MI) {
SmallString<128> Str;
raw_svector_ostream OS(Str);
if (static_cast<NVPTXTargetMachine &>(TM).getDrvInterface() == NVPTX::CUDA)
emitLineNumberAsDotLoc(*MI);
MCInst Inst;
lowerToMCInst(MI, Inst);
EmitToStreamer(*OutStreamer, Inst);
@ -505,7 +446,7 @@ void NVPTXAsmPrinter::EmitFunctionEntryLabel() {
emitGlobals(*MF->getFunction().getParent());
GlobalsEmitted = true;
}
// Set up
MRI = &MF->getRegInfo();
F = &MF->getFunction();
@ -526,14 +467,25 @@ void NVPTXAsmPrinter::EmitFunctionEntryLabel() {
OutStreamer->EmitRawText(O.str());
prevDebugLoc = DebugLoc();
VRegMapping.clear();
// Emit open brace for function body.
OutStreamer->EmitRawText(StringRef("{\n"));
setAndEmitFunctionVirtualRegisters(*MF);
}
bool NVPTXAsmPrinter::runOnMachineFunction(MachineFunction &F) {
nvptxSubtarget = &F.getSubtarget<NVPTXSubtarget>();
bool Result = AsmPrinter::runOnMachineFunction(F);
// Emit closing brace for the body of function F.
// The closing brace must be emitted here because we need to emit additional
// debug labels/data after the last basic block.
// We need to emit the closing brace here because we don't have function that
// finished emission of the function body.
OutStreamer->EmitRawText(StringRef("}\n"));
return Result;
}
void NVPTXAsmPrinter::EmitFunctionBodyStart() {
VRegMapping.clear();
OutStreamer->EmitRawText(StringRef("{\n"));
setAndEmitFunctionVirtualRegisters(*MF);
SmallString<128> Str;
raw_svector_ostream O(Str);
emitDemotedVars(&MF->getFunction(), O);
@ -541,7 +493,6 @@ void NVPTXAsmPrinter::EmitFunctionBodyStart() {
}
void NVPTXAsmPrinter::EmitFunctionBodyEnd() {
OutStreamer->EmitRawText(StringRef("}\n"));
VRegMapping.clear();
}
@ -818,42 +769,6 @@ void NVPTXAsmPrinter::emitDeclarations(const Module &M, raw_ostream &O) {
}
}
void NVPTXAsmPrinter::recordAndEmitFilenames(Module &M) {
DebugInfoFinder DbgFinder;
DbgFinder.processModule(M);
unsigned i = 1;
for (const DICompileUnit *DIUnit : DbgFinder.compile_units()) {
StringRef Filename = DIUnit->getFilename();
StringRef Dirname = DIUnit->getDirectory();
SmallString<128> FullPathName = Dirname;
if (!Dirname.empty() && !sys::path::is_absolute(Filename)) {
sys::path::append(FullPathName, Filename);
Filename = FullPathName;
}
if (filenameMap.find(Filename) != filenameMap.end())
continue;
filenameMap[Filename] = i;
OutStreamer->EmitDwarfFileDirective(i, "", Filename);
++i;
}
for (DISubprogram *SP : DbgFinder.subprograms()) {
StringRef Filename = SP->getFilename();
StringRef Dirname = SP->getDirectory();
SmallString<128> FullPathName = Dirname;
if (!Dirname.empty() && !sys::path::is_absolute(Filename)) {
sys::path::append(FullPathName, Filename);
Filename = FullPathName;
}
if (filenameMap.find(Filename) != filenameMap.end())
continue;
filenameMap[Filename] = i;
OutStreamer->EmitDwarfFileDirective(i, "", Filename);
++i;
}
}
static bool isEmptyXXStructor(GlobalVariable *GV) {
if (!GV) return true;
const ConstantArray *InitList = dyn_cast<ConstantArray>(GV->getInitializer());
@ -889,24 +804,13 @@ bool NVPTXAsmPrinter::doInitialization(Module &M) {
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);
// Initialize TargetLoweringObjectFile since we didn't do in
// AsmPrinter::doInitialization either right above or where it's commented out
// below.
const_cast<TargetLoweringObjectFile &>(getObjFileLowering())
.Initialize(OutContext, TM);
bool Result = AsmPrinter::doInitialization(M);
// Emit header before any dwarf directives are emitted below.
emitHeader(M, OS1, STI);
OutStreamer->EmitRawText(OS1.str());
// Already commented out
//bool Result = AsmPrinter::doInitialization(M);
// Emit module-level inline asm if it exists.
if (!M.getModuleInlineAsm().empty()) {
OutStreamer->AddComment("Start of file scope inline assembly");
@ -917,13 +821,9 @@ bool NVPTXAsmPrinter::doInitialization(Module &M) {
OutStreamer->AddBlankLine();
}
// If we're not NVCL we're CUDA, go ahead and emit filenames.
if (TM.getTargetTriple().getOS() != Triple::NVCL)
recordAndEmitFilenames(M);
GlobalsEmitted = false;
return false; // success
return Result;
}
void NVPTXAsmPrinter::emitGlobals(const Module &M) {
@ -975,8 +875,9 @@ void NVPTXAsmPrinter::emitHeader(Module &M, raw_ostream &O,
if (NTM.getDrvInterface() == NVPTX::NVCL)
O << ", texmode_independent";
if (MAI->doesSupportDebugInformation())
O << ", debug";
// FIXME: remove comment once debug info is properly supported.
if (MMI && MMI->hasDebugInfo())
O << "//, debug";
O << "\n";
@ -991,6 +892,8 @@ void NVPTXAsmPrinter::emitHeader(Module &M, raw_ostream &O,
}
bool NVPTXAsmPrinter::doFinalization(Module &M) {
bool HasDebugInfo = MMI && MMI->hasDebugInfo();
// If we did not emit any functions, then the global declarations have not
// yet been emitted.
if (!GlobalsEmitted) {
@ -1025,6 +928,11 @@ bool NVPTXAsmPrinter::doFinalization(Module &M) {
clearAnnotationCache(&M);
delete[] gv_array;
// FIXME: remove comment once debug info is properly supported.
// Close the last emitted section
if (HasDebugInfo)
OutStreamer->EmitRawText("//\t}");
return ret;
//bool Result = AsmPrinter::doFinalization(M);

8
llvm/lib/Target/NVPTX/NVPTXAsmPrinter.h
View File

@ -344,10 +344,7 @@ public:
delete reader;
}
bool runOnMachineFunction(MachineFunction &F) override {
nvptxSubtarget = &F.getSubtarget<NVPTXSubtarget>();
return AsmPrinter::runOnMachineFunction(F);
}
bool runOnMachineFunction(MachineFunction &F) override;
void getAnalysisUsage(AnalysisUsage &AU) const override {
AU.addRequired<MachineLoopInfo>();
@ -357,9 +354,6 @@ public:
bool ignoreLoc(const MachineInstr &);
std::string getVirtualRegisterName(unsigned) const;
DebugLoc prevDebugLoc;
void emitLineNumberAsDotLoc(const MachineInstr &);
};
} // end namespace llvm

28
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
View File

@ -15,7 +15,6 @@
#include "NVPTXISelLowering.h"
#include "MCTargetDesc/NVPTXBaseInfo.h"
#include "NVPTX.h"
#include "NVPTXSection.h"
#include "NVPTXSubtarget.h"
#include "NVPTXTargetMachine.h"
#include "NVPTXTargetObjectFile.h"
@ -4738,31 +4737,8 @@ void NVPTXTargetLowering::ReplaceNodeResults(
}
}
// Pin NVPTXSection's and NVPTXTargetObjectFile's vtables to this file.
void NVPTXSection::anchor() {}
NVPTXTargetObjectFile::~NVPTXTargetObjectFile() {
delete static_cast<NVPTXSection *>(TextSection);
delete static_cast<NVPTXSection *>(DataSection);
delete static_cast<NVPTXSection *>(BSSSection);
delete static_cast<NVPTXSection *>(ReadOnlySection);
delete static_cast<NVPTXSection *>(StaticCtorSection);
delete static_cast<NVPTXSection *>(StaticDtorSection);
delete static_cast<NVPTXSection *>(LSDASection);
delete static_cast<NVPTXSection *>(EHFrameSection);
delete static_cast<NVPTXSection *>(DwarfAbbrevSection);
delete static_cast<NVPTXSection *>(DwarfInfoSection);
delete static_cast<NVPTXSection *>(DwarfLineSection);
delete static_cast<NVPTXSection *>(DwarfFrameSection);
delete static_cast<NVPTXSection *>(DwarfPubTypesSection);
delete static_cast<const NVPTXSection *>(DwarfDebugInlineSection);
delete static_cast<NVPTXSection *>(DwarfStrSection);
delete static_cast<NVPTXSection *>(DwarfLocSection);
delete static_cast<NVPTXSection *>(DwarfARangesSection);
delete static_cast<NVPTXSection *>(DwarfRangesSection);
delete static_cast<NVPTXSection *>(DwarfMacinfoSection);
}
// Pin NVPTXTargetObjectFile's vtables to this file.
NVPTXTargetObjectFile::~NVPTXTargetObjectFile() {}
MCSection *NVPTXTargetObjectFile::SelectSectionForGlobal(
const GlobalObject *GO, SectionKind Kind, const TargetMachine &TM) const {

45
llvm/lib/Target/NVPTX/NVPTXSection.h
View File

@ -1,45 +0,0 @@
//===- NVPTXSection.h - NVPTX-specific section representation ---*- C++ -*-===//
//
// The LLVM Compiler Infrastructure
//
// This file is distributed under the University of Illinois Open Source
// License. See LICENSE.TXT for details.
//
//===----------------------------------------------------------------------===//
//
// This file declares the NVPTXSection class.
//
//===----------------------------------------------------------------------===//
#ifndef LLVM_LIB_TARGET_NVPTX_NVPTXSECTION_H
#define LLVM_LIB_TARGET_NVPTX_NVPTXSECTION_H
#include "llvm/MC/MCSection.h"
#include "llvm/MC/SectionKind.h"
namespace llvm {
/// Represents a section in PTX PTX does not have sections. We create this class
/// in order to use the ASMPrint interface.
///
class NVPTXSection final : public MCSection {
virtual void anchor();
public:
NVPTXSection(SectionVariant V, SectionKind K) : MCSection(V, K, nullptr) {}
~NVPTXSection() = default;
/// Override this as NVPTX has its own way of printing switching
/// to a section.
void PrintSwitchToSection(const MCAsmInfo &MAI, const Triple &T,
raw_ostream &OS,
const MCExpr *Subsection) const override {}
/// Base address of PTX sections is zero.
bool UseCodeAlign() const override { return false; }
bool isVirtualSection() const override { return false; }
};
} // end namespace llvm
#endif // LLVM_LIB_TARGET_NVPTX_NVPTXSECTION_H

59
llvm/lib/Target/NVPTX/NVPTXTargetObjectFile.h
View File

@ -10,7 +10,6 @@
#ifndef LLVM_LIB_TARGET_NVPTX_NVPTXTARGETOBJECTFILE_H
#define LLVM_LIB_TARGET_NVPTX_NVPTXTARGETOBJECTFILE_H
#include "NVPTXSection.h"
#include "llvm/MC/MCSection.h"
#include "llvm/MC/SectionKind.h"
#include "llvm/Target/TargetLoweringObjectFile.h"
@ -19,68 +18,12 @@ namespace llvm {
class NVPTXTargetObjectFile : public TargetLoweringObjectFile {
public:
NVPTXTargetObjectFile() {
TextSection = nullptr;
DataSection = nullptr;
BSSSection = nullptr;
ReadOnlySection = nullptr;
StaticCtorSection = nullptr;
StaticDtorSection = nullptr;
LSDASection = nullptr;
EHFrameSection = nullptr;
DwarfAbbrevSection = nullptr;
DwarfInfoSection = nullptr;
DwarfLineSection = nullptr;
DwarfFrameSection = nullptr;
DwarfPubTypesSection = nullptr;
DwarfDebugInlineSection = nullptr;
DwarfStrSection = nullptr;
DwarfLocSection = nullptr;
DwarfARangesSection = nullptr;
DwarfRangesSection = nullptr;
DwarfMacinfoSection = nullptr;
}
NVPTXTargetObjectFile() : TargetLoweringObjectFile() {}
~NVPTXTargetObjectFile() override;
void Initialize(MCContext &ctx, const TargetMachine &TM) override {
TargetLoweringObjectFile::Initialize(ctx, TM);
TextSection = new NVPTXSection(MCSection::SV_ELF, SectionKind::getText());
DataSection = new NVPTXSection(MCSection::SV_ELF, SectionKind::getData());
BSSSection = new NVPTXSection(MCSection::SV_ELF, SectionKind::getBSS());
ReadOnlySection =
new NVPTXSection(MCSection::SV_ELF, SectionKind::getReadOnly());
StaticCtorSection =
new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());
StaticDtorSection =
new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());
LSDASection =
new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());
EHFrameSection =
new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());
DwarfAbbrevSection =
new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());
DwarfInfoSection =
new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());
DwarfLineSection =
new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());
DwarfFrameSection =
new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());
DwarfPubTypesSection =
new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());
DwarfDebugInlineSection =
new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());
DwarfStrSection =
new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());
DwarfLocSection =
new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());
DwarfARangesSection =
new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());
DwarfRangesSection =
new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());
DwarfMacinfoSection =
new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());
}
MCSection *getSectionForConstant(const DataLayout &DL, SectionKind Kind,

291
llvm/test/DebugInfo/NVPTX/cu-range-hole.ll Normal file
View File

@ -0,0 +1,291 @@
; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s
; CHECK: .target sm_{{[0-9]+}}//, debug
; CHECK: .visible .func (.param .b32 func_retval0) b(
; CHECK: .param .b32 b_param_0
; CHECK: )
; CHECK: {
; CHECK: Lfunc_begin0:
; CHECK: .loc 1 1 0
; CHECK: .loc 1 1 0
; CHECK: ret;
; CHECK: Lfunc_end0:
; CHECK: }
; CHECK: .visible .func (.param .b32 func_retval0) a(
; CHECK: .param .b32 a_param_0
; CHECK: )
; CHECK: {
; CHECK: Lfunc_begin1:
; CHECK-NOT: .loc
; CHECK: ret;
; CHECK: Lfunc_end1:
; CHECK: }
; CHECK: .visible .func (.param .b32 func_retval0) d(
; CHECK: .param .b32 d_param_0
; CHECK: )
; CHECK: {
; CHECK: Lfunc_begin2:
; CHECK: .loc 1 3 0
; CHECK: ret;
; CHECK: Lfunc_end2:
; CHECK: }
; CHECK: .file 1 "{{.*}}b.c"
; Function Attrs: nounwind uwtable
define i32 @b(i32 %c) #0 !dbg !5 {
entry:
%c.addr = alloca i32, align 4
store i32 %c, i32* %c.addr, align 4
call void @llvm.dbg.declare(metadata i32* %c.addr, metadata !13, metadata !DIExpression()), !dbg !14
%0 = load i32, i32* %c.addr, align 4, !dbg !14
%add = add nsw i32 %0, 1, !dbg !14
ret i32 %add, !dbg !14
}
; Function Attrs: nounwind uwtable
define i32 @a(i32 %b) #0 {
entry:
%b.addr = alloca i32, align 4
store i32 %b, i32* %b.addr, align 4
%0 = load i32, i32* %b.addr, align 4
%add = add nsw i32 %0, 1
ret i32 %add
}
; Function Attrs: nounwind readnone
declare void @llvm.dbg.declare(metadata, metadata, metadata) #1
; Function Attrs: nounwind uwtable
define i32 @d(i32 %e) #0 !dbg !10 {
entry:
%e.addr = alloca i32, align 4
store i32 %e, i32* %e.addr, align 4
call void @llvm.dbg.declare(metadata i32* %e.addr, metadata !15, metadata !DIExpression()), !dbg !16
%0 = load i32, i32* %e.addr, align 4, !dbg !16
%add = add nsw i32 %0, 1, !dbg !16
ret i32 %add, !dbg !16
}
; CHECK: // .section .debug_abbrev
; CHECK: // {
; CHECK: // .b8 1 // Abbreviation Code
; CHECK: // .b8 17 // DW_TAG_compile_unit
; CHECK: // .b8 1 // DW_CHILDREN_yes
; CHECK: // .b8 37 // DW_AT_producer
; CHECK: // .b8 8 // DW_FORM_string
; CHECK: // .b8 19 // DW_AT_language
; CHECK: // .b8 5 // DW_FORM_data2
; CHECK: // .b8 3 // DW_AT_name
; CHECK: // .b8 8 // DW_FORM_string
; CHECK: // .b8 16 // DW_AT_stmt_list
; CHECK: // .b8 6 // DW_FORM_data4
; CHECK: // .b8 27 // DW_AT_comp_dir
; CHECK: // .b8 8 // DW_FORM_string
; CHECK: // .b8 17 // DW_AT_low_pc
; CHECK: // .b8 1 // DW_FORM_addr
; CHECK: // .b8 18 // DW_AT_high_pc
; CHECK: // .b8 1 // DW_FORM_addr
; CHECK: // .b8 0 // EOM(1)
; CHECK: // .b8 0 // EOM(2)
; CHECK: // .b8 2 // Abbreviation Code
; CHECK: // .b8 46 // DW_TAG_subprogram
; CHECK: // .b8 1 // DW_CHILDREN_yes
; CHECK: // .b8 17 // DW_AT_low_pc
; CHECK: // .b8 1 // DW_FORM_addr
; CHECK: // .b8 18 // DW_AT_high_pc
; CHECK: // .b8 1 // DW_FORM_addr
; CHECK: // .b8 3 // DW_AT_name
; CHECK: // .b8 8 // DW_FORM_string
; CHECK: // .b8 58 // DW_AT_decl_file
; CHECK: // .b8 11 // DW_FORM_data1
; CHECK: // .b8 59 // DW_AT_decl_line
; CHECK: // .b8 11 // DW_FORM_data1
; CHECK: // .b8 39 // DW_AT_prototyped
; CHECK: // .b8 12 // DW_FORM_flag
; CHECK: // .b8 73 // DW_AT_type
; CHECK: // .b8 19 // DW_FORM_ref4
; CHECK: // .b8 63 // DW_AT_external
; CHECK: // .b8 12 // DW_FORM_flag
; CHECK: // .b8 0 // EOM(1)
; CHECK: // .b8 0 // EOM(2)
; CHECK: // .b8 3 // Abbreviation Code
; CHECK: // .b8 5 // DW_TAG_formal_parameter
; CHECK: // .b8 0 // DW_CHILDREN_no
; CHECK: // .b8 3 // DW_AT_name
; CHECK: // .b8 8 // DW_FORM_string
; CHECK: // .b8 58 // DW_AT_decl_file
; CHECK: // .b8 11 // DW_FORM_data1
; CHECK: // .b8 59 // DW_AT_decl_line
; CHECK: // .b8 11 // DW_FORM_data1
; CHECK: // .b8 73 // DW_AT_type
; CHECK: // .b8 19 // DW_FORM_ref4
; CHECK: // .b8 0 // EOM(1)
; CHECK: // .b8 0 // EOM(2)
; CHECK: // .b8 4 // Abbreviation Code
; CHECK: // .b8 36 // DW_TAG_base_type
; CHECK: // .b8 0 // DW_CHILDREN_no
; CHECK: // .b8 3 // DW_AT_name
; CHECK: // .b8 8 // DW_FORM_string
; CHECK: // .b8 62 // DW_AT_encoding
; CHECK: // .b8 11 // DW_FORM_data1
; CHECK: // .b8 11 // DW_AT_byte_size
; CHECK: // .b8 11 // DW_FORM_data1
; CHECK: // .b8 0 // EOM(1)
; CHECK: // .b8 0 // EOM(2)
; CHECK: // .b8 0 // EOM(3)
; CHECK: // }
; CHECK: // .section .debug_info
; CHECK: // {
; CHECK: // .b32 179 // Length of Unit
; CHECK: // .b8 2 // DWARF version number
; CHECK: // .b8 0
; CHECK: // .b32 .debug_abbrev // Offset Into Abbrev. Section
; CHECK: // .b8 8 // Address Size (in bytes)
; CHECK: // .b8 1 // Abbrev [1] 0xb:0xac DW_TAG_compile_unit
; CHECK: // .b8 99 // DW_AT_producer
; CHECK: // .b8 108
; CHECK: // .b8 97
; CHECK: // .b8 110
; CHECK: // .b8 103
; CHECK: // .b8 32
; CHECK: // .b8 118
; CHECK: // .b8 101
; CHECK: // .b8 114
; CHECK: // .b8 115
; CHECK: // .b8 105
; CHECK: // .b8 111
; CHECK: // .b8 110
; CHECK: // .b8 32
; CHECK: // .b8 51
; CHECK: // .b8 46
; CHECK: // .b8 53
; CHECK: // .b8 46
; CHECK: // .b8 48
; CHECK: // .b8 32
; CHECK: // .b8 40
; CHECK: // .b8 116
; CHECK: // .b8 114
; CHECK: // .b8 117
; CHECK: // .b8 110
; CHECK: // .b8 107
; CHECK: // .b8 32
; CHECK: // .b8 50
; CHECK: // .b8 48
; CHECK: // .b8 52
; CHECK: // .b8 49
; CHECK: // .b8 54
; CHECK: // .b8 52
; CHECK: // .b8 41
; CHECK: // .b8 32
; CHECK: // .b8 40
; CHECK: // .b8 108
; CHECK: // .b8 108
; CHECK: // .b8 118
; CHECK: // .b8 109
; CHECK: // .b8 47
; CHECK: // .b8 116
; CHECK: // .b8 114
; CHECK: // .b8 117
; CHECK: // .b8 110
; CHECK: // .b8 107
; CHECK: // .b8 32
; CHECK: // .b8 50
; CHECK: // .b8 48
; CHECK: // .b8 52
; CHECK: // .b8 49
; CHECK: // .b8 56
; CHECK: // .b8 51
; CHECK: // .b8 41
; CHECK: // .b8 0
; CHECK: // .b8 12 // DW_AT_language
; CHECK: // .b8 0
; CHECK: // .b8 98 // DW_AT_name
; CHECK: // .b8 46
; CHECK: // .b8 99
; CHECK: // .b8 0
; CHECK: // .b32 .debug_line // DW_AT_stmt_list
; CHECK: // .b8 47 // DW_AT_comp_dir
; CHECK: // .b8 115
; CHECK: // .b8 111
; CHECK: // .b8 117
; CHECK: // .b8 114
; CHECK: // .b8 99
; CHECK: // .b8 101
; CHECK: // .b8 0
; CHECK: // .b64 Lfunc_begin0 // DW_AT_low_pc
; CHECK: // .b64 Lfunc_end2 // DW_AT_high_pc
; CHECK: // .b8 2 // Abbrev [2] 0x65:0x25 DW_TAG_subprogram
; CHECK: // .b64 Lfunc_begin0 // DW_AT_low_pc
; CHECK: // .b64 Lfunc_end0 // DW_AT_high_pc
; CHECK: // .b8 98 // DW_AT_name
; CHECK: // .b8 0
; CHECK: // .b8 1 // DW_AT_decl_file
; CHECK: // .b8 1 // DW_AT_decl_line
; CHECK: // .b8 1 // DW_AT_prototyped
; CHECK: // .b32 175 // DW_AT_type
; CHECK: // .b8 1 // DW_AT_external
; CHECK: // .b8 3 // Abbrev [3] 0x80:0x9 DW_TAG_formal_parameter
; CHECK: // .b8 99 // DW_AT_name
; CHECK: // .b8 0
; CHECK: // .b8 1 // DW_AT_decl_file
; CHECK: // .b8 1 // DW_AT_decl_line
; CHECK: // .b32 175 // DW_AT_type
; CHECK: // .b8 0 // End Of Children Mark
; CHECK: // .b8 2 // Abbrev [2] 0x8a:0x25 DW_TAG_subprogram
; CHECK: // .b64 Lfunc_begin2 // DW_AT_low_pc
; CHECK: // .b64 Lfunc_end2 // DW_AT_high_pc
; CHECK: // .b8 100 // DW_AT_name
; CHECK: // .b8 0
; CHECK: // .b8 1 // DW_AT_decl_file
; CHECK: // .b8 3 // DW_AT_decl_line
; CHECK: // .b8 1 // DW_AT_prototyped
; CHECK: // .b32 175 // DW_AT_type
; CHECK: // .b8 1 // DW_AT_external
; CHECK: // .b8 3 // Abbrev [3] 0xa5:0x9 DW_TAG_formal_parameter
; CHECK: // .b8 101 // DW_AT_name
; CHECK: // .b8 0
; CHECK: // .b8 1 // DW_AT_decl_file
; CHECK: // .b8 3 // DW_AT_decl_line
; CHECK: // .b32 175 // DW_AT_type
; CHECK: // .b8 0 // End Of Children Mark
; CHECK: // .b8 4 // Abbrev [4] 0xaf:0x7 DW_TAG_base_type
; CHECK: // .b8 105 // DW_AT_name
; CHECK: // .b8 110
; CHECK: // .b8 116
; CHECK: // .b8 0
; CHECK: // .b8 5 // DW_AT_encoding
; CHECK: // .b8 4 // DW_AT_byte_size
; CHECK: // .b8 0 // End Of Children Mark
; CHECK: // }
; CHECK: // .section .debug_macinfo
; CHECK: // {
; CHECK: // .b8 0 // End Of Macro List Mark
; CHECK: // }
attributes #0 = { nounwind uwtable "less-precise-fpmad"="false" "no-frame-pointer-elim"="true" "no-frame-pointer-elim-non-leaf" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" }
attributes #1 = { nounwind readnone }
!llvm.ident = !{!0, !0}
!llvm.dbg.cu = !{!1}
!llvm.module.flags = !{!11, !12}
!0 = !{!"clang version 3.5.0 (trunk 204164) (llvm/trunk 204183)"}
!1 = distinct !DICompileUnit(language: DW_LANG_C99, producer: "clang version 3.5.0 (trunk 204164) (llvm/trunk 204183)", isOptimized: false, emissionKind: FullDebug, file: !2, enums: !3, retainedTypes: !3, globals: !3, imports: !3)
!2 = !DIFile(filename: "b.c", directory: "/source")
!3 = !{}
!5 = distinct !DISubprogram(name: "b", line: 1, isLocal: false, isDefinition: true, virtualIndex: 6, flags: DIFlagPrototyped, isOptimized: false, unit: !1, scopeLine: 1, file: !2, scope: !6, type: !7, retainedNodes: !3)
!6 = !DIFile(filename: "b.c", directory: "/source")
!7 = !DISubroutineType(types: !8)
!8 = !{!9, !9}
!9 = !DIBasicType(tag: DW_TAG_base_type, name: "int", size: 32, align: 32, encoding: DW_ATE_signed)
!10 = distinct !DISubprogram(name: "d", line: 3, isLocal: false, isDefinition: true, virtualIndex: 6, flags: DIFlagPrototyped, isOptimized: false, unit: !1, scopeLine: 3, file: !2, scope: !6, type: !7, retainedNodes: !3)
!11 = !{i32 2, !"Dwarf Version", i32 2}
!12 = !{i32 1, !"Debug Info Version", i32 3}
!13 = !DILocalVariable(name: "c", line: 1, arg: 1, scope: !5, file: !6, type: !9)
!14 = !DILocation(line: 1, scope: !5)
!15 = !DILocalVariable(name: "e", line: 3, arg: 1, scope: !10, file: !6, type: !9)
!16 = !DILocation(line: 3, scope: !10)

252
llvm/test/DebugInfo/NVPTX/dbg-declare-alloca.ll Normal file
View File

@ -0,0 +1,252 @@
; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s
; CHECK: .target sm_20//, debug
; CHECK: .visible .func use_dbg_declare()
; CHECK: .local .align 8 .b8 __local_depot0[8];
; CHECK: mov.u64 %SPL, __local_depot0;
; CHECK: add.u64 %rd1, %SP, 0;
; CHECK: .loc 1 5 3 // t.c:5:3
; CHECK: { // callseq 0, 0
; CHECK: .reg .b32 temp_param_reg;
; CHECK: .param .b64 param0;
; CHECK: st.param.b64 [param0+0], %rd1;
; CHECK: call.uni
; CHECK: escape_foo,
; CHECK: (
; CHECK: param0
; CHECK: );
; CHECK: } // callseq 0
; CHECK: .loc 1 6 1 // t.c:6:1
; CHECK: ret;
; CHECK: }
; CHECK: .file 1 "test{{(/|\\\\)}}t.c"
; CHECK: // .section .debug_abbrev
; CHECK: // {
; CHECK: // .b8 1 // Abbreviation Code
; CHECK: // .b8 17 // DW_TAG_compile_unit
; CHECK: // .b8 1 // DW_CHILDREN_yes
; CHECK: // .b8 37 // DW_AT_producer
; CHECK: // .b8 8 // DW_FORM_string
; CHECK: // .b8 19 // DW_AT_language
; CHECK: // .b8 5 // DW_FORM_data2
; CHECK: // .b8 3 // DW_AT_name
; CHECK: // .b8 8 // DW_FORM_string
; CHECK: // .b8 16 // DW_AT_stmt_list
; CHECK: // .b8 6 // DW_FORM_data4
; CHECK: // .b8 27 // DW_AT_comp_dir
; CHECK: // .b8 8 // DW_FORM_string
; CHECK: // .b8 17 // DW_AT_low_pc
; CHECK: // .b8 1 // DW_FORM_addr
; CHECK: // .b8 18 // DW_AT_high_pc
; CHECK: // .b8 1 // DW_FORM_addr
; CHECK: // .b8 0 // EOM(1)
; CHECK: // .b8 0 // EOM(2)
; CHECK: // .b8 2 // Abbreviation Code
; CHECK: // .b8 46 // DW_TAG_subprogram
; CHECK: // .b8 1 // DW_CHILDREN_yes
; CHECK: // .b8 17 // DW_AT_low_pc
; CHECK: // .b8 1 // DW_FORM_addr
; CHECK: // .b8 18 // DW_AT_high_pc
; CHECK: // .b8 1 // DW_FORM_addr
; CHECK: // .b8 3 // DW_AT_name
; CHECK: // .b8 8 // DW_FORM_string
; CHECK: // .b8 58 // DW_AT_decl_file
; CHECK: // .b8 11 // DW_FORM_data1
; CHECK: // .b8 59 // DW_AT_decl_line
; CHECK: // .b8 11 // DW_FORM_data1
; CHECK: // .b8 39 // DW_AT_prototyped
; CHECK: // .b8 12 // DW_FORM_flag
; CHECK: // .b8 63 // DW_AT_external
; CHECK: // .b8 12 // DW_FORM_flag
; CHECK: // .b8 0 // EOM(1)
; CHECK: // .b8 0 // EOM(2)
; CHECK: // .b8 3 // Abbreviation Code
; CHECK: // .b8 52 // DW_TAG_variable
; CHECK: // .b8 0 // DW_CHILDREN_no
; CHECK: // .b8 2 // DW_AT_location
; CHECK: // .b8 10 // DW_FORM_block1
; CHECK: // .b8 3 // DW_AT_name
; CHECK: // .b8 8 // DW_FORM_string
; CHECK: // .b8 58 // DW_AT_decl_file
; CHECK: // .b8 11 // DW_FORM_data1
; CHECK: // .b8 59 // DW_AT_decl_line
; CHECK: // .b8 11 // DW_FORM_data1
; CHECK: // .b8 73 // DW_AT_type
; CHECK: // .b8 19 // DW_FORM_ref4
; CHECK: // .b8 0 // EOM(1)
; CHECK: // .b8 0 // EOM(2)
; CHECK: // .b8 4 // Abbreviation Code
; CHECK: // .b8 19 // DW_TAG_structure_type
; CHECK: // .b8 1 // DW_CHILDREN_yes
; CHECK: // .b8 3 // DW_AT_name
; CHECK: // .b8 8 // DW_FORM_string
; CHECK: // .b8 11 // DW_AT_byte_size
; CHECK: // .b8 11 // DW_FORM_data1
; CHECK: // .b8 58 // DW_AT_decl_file
; CHECK: // .b8 11 // DW_FORM_data1
; CHECK: // .b8 59 // DW_AT_decl_line
; CHECK: // .b8 11 // DW_FORM_data1
; CHECK: // .b8 0 // EOM(1)
; CHECK: // .b8 0 // EOM(2)
; CHECK: // .b8 5 // Abbreviation Code
; CHECK: // .b8 13 // DW_TAG_member
; CHECK: // .b8 0 // DW_CHILDREN_no
; CHECK: // .b8 3 // DW_AT_name
; CHECK: // .b8 8 // DW_FORM_string
; CHECK: // .b8 73 // DW_AT_type
; CHECK: // .b8 19 // DW_FORM_ref4
; CHECK: // .b8 58 // DW_AT_decl_file
; CHECK: // .b8 11 // DW_FORM_data1
; CHECK: // .b8 59 // DW_AT_decl_line
; CHECK: // .b8 11 // DW_FORM_data1
; CHECK: // .b8 56 // DW_AT_data_member_location
; CHECK: // .b8 10 // DW_FORM_block1
; CHECK: // .b8 0 // EOM(1)
; CHECK: // .b8 0 // EOM(2)
; CHECK: // .b8 6 // Abbreviation Code
; CHECK: // .b8 36 // DW_TAG_base_type
; CHECK: // .b8 0 // DW_CHILDREN_no
; CHECK: // .b8 3 // DW_AT_name
; CHECK: // .b8 8 // DW_FORM_string
; CHECK: // .b8 62 // DW_AT_encoding
; CHECK: // .b8 11 // DW_FORM_data1
; CHECK: // .b8 11 // DW_AT_byte_size
; CHECK: // .b8 11 // DW_FORM_data1
; CHECK: // .b8 0 // EOM(1)
; CHECK: // .b8 0 // EOM(2)
; CHECK: // .b8 0 // EOM(3)
; CHECK: // }
; CHECK: // .section .debug_info
; CHECK: // {
; CHECK: // .b32 124 // Length of Unit
; CHECK: // .b8 2 // DWARF version number
; CHECK: // .b8 0
; CHECK: // .b32 .debug_abbrev // Offset Into Abbrev. Section
; CHECK: // .b8 8 // Address Size (in bytes)
; CHECK: // .b8 1 // Abbrev [1] 0xb:0x75 DW_TAG_compile_unit
; CHECK: // .b8 99 // DW_AT_producer
; CHECK: // .b8 108
; CHECK: // .b8 97
; CHECK: // .b8 110
; CHECK: // .b8 103
; CHECK: // .b8 0
; CHECK: // .b8 12 // DW_AT_language
; CHECK: // .b8 0
; CHECK: // .b8 116 // DW_AT_name
; CHECK: // .b8 46
; CHECK: // .b8 99
; CHECK: // .b8 0
; CHECK: // .b32 .debug_line // DW_AT_stmt_list
; CHECK: // .b8 116 // DW_AT_comp_dir
; CHECK: // .b8 101
; CHECK: // .b8 115
; CHECK: // .b8 116
; CHECK: // .b8 0
; CHECK: // .b64 Lfunc_begin0 // DW_AT_low_pc
; CHECK: // .b64 Lfunc_end0 // DW_AT_high_pc
; CHECK: // .b8 2 // Abbrev [2] 0x31:0x32 DW_TAG_subprogram
; CHECK: // .b64 Lfunc_begin0 // DW_AT_low_pc
; CHECK: // .b64 Lfunc_end0 // DW_AT_high_pc
; CHECK: // .b8 117 // DW_AT_name
; CHECK: // .b8 115
; CHECK: // .b8 101
; CHECK: // .b8 95
; CHECK: // .b8 100
; CHECK: // .b8 98
; CHECK: // .b8 103
; CHECK: // .b8 95
; CHECK: // .b8 100
; CHECK: // .b8 101
; CHECK: // .b8 99
; CHECK: // .b8 108
; CHECK: // .b8 97
; CHECK: // .b8 114
; CHECK: // .b8 101
; CHECK: // .b8 0
; CHECK: // .b8 1 // DW_AT_decl_file
; CHECK: // .b8 3 // DW_AT_decl_line
; CHECK: // .b8 1 // DW_AT_prototyped
; CHECK: // .b8 1 // DW_AT_external
; CHECK: // .b8 3 // Abbrev [3] 0x56:0xc DW_TAG_variable
; CHECK: // .b8 2 // DW_AT_location
; CHECK: // .b8 35
; CHECK: // .b8 8
; CHECK: // .b8 111 // DW_AT_name
; CHECK: // .b8 0
; CHECK: // .b8 1 // DW_AT_decl_file
; CHECK: // .b8 4 // DW_AT_decl_line
; CHECK: // .b32 99 // DW_AT_type
; CHECK: // .b8 0 // End Of Children Mark
; CHECK: // .b8 4 // Abbrev [4] 0x63:0x15 DW_TAG_structure_type
; CHECK: // .b8 70 // DW_AT_name
; CHECK: // .b8 111
; CHECK: // .b8 111
; CHECK: // .b8 0
; CHECK: // .b8 4 // DW_AT_byte_size
; CHECK: // .b8 1 // DW_AT_decl_file
; CHECK: // .b8 1 // DW_AT_decl_line
; CHECK: // .b8 5 // Abbrev [5] 0x6b:0xc DW_TAG_member
; CHECK: // .b8 120 // DW_AT_name
; CHECK: // .b8 0
; CHECK: // .b32 120 // DW_AT_type
; CHECK: // .b8 1 // DW_AT_decl_file
; CHECK: // .b8 1 // DW_AT_decl_line
; CHECK: // .b8 2 // DW_AT_data_member_location
; CHECK: // .b8 35
; CHECK: // .b8 0
; CHECK: // .b8 0 // End Of Children Mark
; CHECK: // .b8 6 // Abbrev [6] 0x78:0x7 DW_TAG_base_type
; CHECK: // .b8 105 // DW_AT_name
; CHECK: // .b8 110
; CHECK: // .b8 116
; CHECK: // .b8 0
; CHECK: // .b8 5 // DW_AT_encoding
; CHECK: // .b8 4 // DW_AT_byte_size
; CHECK: // .b8 0 // End Of Children Mark
; CHECK: // }
%struct.Foo = type { i32 }
; Function Attrs: noinline nounwind uwtable
define void @use_dbg_declare() #0 !dbg !7 {
entry:
%o = alloca %struct.Foo, align 4
call void @llvm.dbg.declare(metadata %struct.Foo* %o, metadata !10, metadata !15), !dbg !16
call void @escape_foo(%struct.Foo* %o), !dbg !17
ret void, !dbg !18
}
; Function Attrs: nounwind readnone speculatable
declare void @llvm.dbg.declare(metadata, metadata, metadata) #1
declare void @escape_foo(%struct.Foo*)
attributes #0 = { noinline nounwind uwtable }
attributes #1 = { nounwind readnone speculatable }
!llvm.dbg.cu = !{!0}
!llvm.module.flags = !{!3, !4, !5}
!llvm.ident = !{!6}
!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang", isOptimized: false, runtimeVersion: 0, emissionKind: FullDebug, enums: !2)
!1 = !DIFile(filename: "t.c", directory: "test")
!2 = !{}
!3 = !{i32 2, !"Dwarf Version", i32 2}
!4 = !{i32 2, !"Debug Info Version", i32 3}
!5 = !{i32 1, !"wchar_size", i32 4}
!6 = !{!"clang"}
!7 = distinct !DISubprogram(name: "use_dbg_declare", scope: !1, file: !1, line: 3, type: !8, isLocal: false, isDefinition: true, scopeLine: 3, flags: DIFlagPrototyped, isOptimized: false, unit: !0, retainedNodes: !2)
!8 = !DISubroutineType(types: !9)
!9 = !{null}
!10 = !DILocalVariable(name: "o", scope: !7, file: !1, line: 4, type: !11)
!11 = distinct !DICompositeType(tag: DW_TAG_structure_type, name: "Foo", file: !1, line: 1, size: 32, elements: !12)
!12 = !{!13}
!13 = !DIDerivedType(tag: DW_TAG_member, name: "x", scope: !11, file: !1, line: 1, baseType: !14, size: 32)
!14 = !DIBasicType(name: "int", size: 32, encoding: DW_ATE_signed)
!15 = !DIExpression()
!16 = !DILocation(line: 4, column: 14, scope: !7)
!17 = !DILocation(line: 5, column: 3, scope: !7)
!18 = !DILocation(line: 6, column: 1, scope: !7)

74
llvm/test/DebugInfo/NVPTX/debug-file-loc.ll
View File

@ -8,13 +8,13 @@
;__device__ void bar() {}
;}
; CHECK: .file 1 "/source/dir{{.+}}bar.cu"
; CHECK: .file 2 "/source/dir{{.+}}foo.h"
; CHECK: .target sm_{{[0-9]+}}//, debug
; CHECK: .visible .func foo()
; CHECK: .loc 2 1 31
; CHECK: .loc [[FOO:[0-9]+]] 1 31
; CHECK: ret;
; CHECK: .visible .func bar()
; CHECK: .loc 1 2 31
; CHECK: .loc [[BAR:[0-9]+]] 2 31
; CHECK: ret;
define void @foo() !dbg !4 {
@ -27,6 +27,70 @@ bb:
ret void, !dbg !11
}
; CHECK-DAG: .file [[FOO]] "{{.*}}foo.h"
; CHECK-DAG: .file [[BAR]] "{{.*}}bar.cu"
; CHECK: // .section .debug_abbrev
; CHECK: // {
; CHECK: // .b8 1 // Abbreviation Code
; CHECK: // .b8 17 // DW_TAG_compile_unit
; CHECK: // .b8 0 // DW_CHILDREN_no
; CHECK: // .b8 37 // DW_AT_producer
; CHECK: // .b8 8 // DW_FORM_string
; CHECK: // .b8 19 // DW_AT_language
; CHECK: // .b8 5 // DW_FORM_data2
; CHECK: // .b8 3 // DW_AT_name
; CHECK: // .b8 8 // DW_FORM_string
; CHECK: // .b8 16 // DW_AT_stmt_list
; CHECK: // .b8 6 // DW_FORM_data4
; CHECK: // .b8 27 // DW_AT_comp_dir
; CHECK: // .b8 8 // DW_FORM_string
; CHECK: // .b8 17 // DW_AT_low_pc
; CHECK: // .b8 1 // DW_FORM_addr
; CHECK: // .b8 18 // DW_AT_high_pc
; CHECK: // .b8 1 // DW_FORM_addr
; CHECK: // .b8 0 // EOM(1)
; CHECK: // .b8 0 // EOM(2)
; CHECK: // .b8 0 // EOM(3)
; CHECK: // }
; CHECK: // .section .debug_info
; CHECK: // {
; CHECK: // .b32 50 // Length of Unit
; CHECK: // .b8 2 // DWARF version number
; CHECK: // .b8 0
; CHECK: // .b32 .debug_abbrev // Offset Into Abbrev. Section
; CHECK: // .b8 8 // Address Size (in bytes)
; CHECK: // .b8 1 // Abbrev [1] 0xb:0x2b DW_TAG_compile_unit
; CHECK: // .b8 0 // DW_AT_producer
; CHECK: // .b8 4 // DW_AT_language
; CHECK: // .b8 0
; CHECK: // .b8 98 // DW_AT_name
; CHECK: // .b8 97
; CHECK: // .b8 114
; CHECK: // .b8 46
; CHECK: // .b8 99
; CHECK: // .b8 117
; CHECK: // .b8 0
; CHECK: // .b32 .debug_line // DW_AT_stmt_list
; CHECK: // .b8 47 // DW_AT_comp_dir
; CHECK: // .b8 115
; CHECK: // .b8 111
; CHECK: // .b8 117
; CHECK: // .b8 114
; CHECK: // .b8 99
; CHECK: // .b8 101
; CHECK: // .b8 47
; CHECK: // .b8 100
; CHECK: // .b8 105
; CHECK: // .b8 114
; CHECK: // .b8 0
; CHECK: // .b64 Lfunc_begin0 // DW_AT_low_pc
; CHECK: // .b64 Lfunc_end1 // DW_AT_high_pc
; CHECK: // }
; CHECK: // .section .debug_macinfo
; CHECK: // {
; CHECK: // .b8 0 // End Of Macro List Mark
; CHECK: // }
!llvm.dbg.cu = !{!0}
!llvm.module.flags = !{!8, !9}
@ -37,7 +101,7 @@ bb:
!5 = !DIFile(filename: "foo.h", directory: "/source/dir")
!6 = !DISubroutineType(types: !2)
!7 = distinct !DISubprogram(name: "bar", scope: !1, file: !1, line: 2, type: !6, isLocal: false, isDefinition: true, scopeLine: 2, flags: DIFlagPrototyped, isOptimized: false, unit: !0, retainedNodes: !2)
!8 = !{i32 2, !"Dwarf Version", i32 4}
!8 = !{i32 2, !"Dwarf Version", i32 2}
!9 = !{i32 2, !"Debug Info Version", i32 3}
!10 = !DILocation(line: 1, column: 31, scope: !4)
!11 = !DILocation(line: 2, column: 31, scope: !7)

8339
llvm/test/DebugInfo/NVPTX/debug-info.ll
View File

File diff suppressed because it is too large Load Diff