[llvm] r330271 - [DEBUG] Initial adaptation of NVPTX target for debug info emission.
Eric Christopher via llvm-commits
llvm-commits at lists.llvm.org
Mon Apr 30 17:14:40 PDT 2018
FWIW This has been temporarily reverted in r331237 after some discussion on
the patch review thread. Still working up a public testcase that can show
the issue.
-eric
On Wed, Apr 18, 2018 at 9:16 AM Alexey Bataev via llvm-commits <
llvm-commits at lists.llvm.org> wrote:
> Author: abataev
> Date: Wed Apr 18 09:13:41 2018
> New Revision: 330271
>
> URL: http://llvm.org/viewvc/llvm-project?rev=330271&view=rev
> Log:
> [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.
>
> Reviewers: echristo, jlebar, tra, jholewinski
>
> Subscribers: mgorny, aprantl, JDevlieghere, llvm-commits
>
> Differential Revision: https://reviews.llvm.org/D41827
>
> Added:
> llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp
> llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h
> llvm/trunk/test/DebugInfo/NVPTX/cu-range-hole.ll
> Removed:
> llvm/trunk/lib/Target/NVPTX/NVPTXSection.h
> Modified:
> llvm/trunk/lib/CodeGen/AsmPrinter/DwarfDebug.cpp
> llvm/trunk/lib/Target/NVPTX/MCTargetDesc/CMakeLists.txt
> llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp
> llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.h
> llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp
> llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
> llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.h
> llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp
> llvm/trunk/lib/Target/NVPTX/NVPTXTargetObjectFile.h
> llvm/trunk/test/DebugInfo/NVPTX/debug-file-loc.ll
> llvm/trunk/test/DebugInfo/NVPTX/debug-info.ll
>
> Modified: llvm/trunk/lib/CodeGen/AsmPrinter/DwarfDebug.cpp
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/CodeGen/AsmPrinter/DwarfDebug.cpp?rev=330271&r1=330270&r2=330271&view=diff
>
> ==============================================================================
> --- llvm/trunk/lib/CodeGen/AsmPrinter/DwarfDebug.cpp (original)
> +++ llvm/trunk/lib/CodeGen/AsmPrinter/DwarfDebug.cpp Wed Apr 18 09:13:41
> 2018
> @@ -312,7 +312,11 @@ DwarfDebug::DwarfDebug(AsmPrinter *A, Mo
> } 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, Mo
> 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;
> -
> - UsePubSections = !NoDwarfPubSections;
> - UseRangesSection = !NoDwarfRangesSection;
> -
> - // Use sections as references.
> - UseSectionsAsReferences = DwarfSectionsAsReferences == Enable;
> + // 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 && !TT.isNVPTX();
> + UseRangesSection = !NoDwarfRangesSection && !TT.isNVPTX();
> +
> + // 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
>
> Modified: llvm/trunk/lib/Target/NVPTX/MCTargetDesc/CMakeLists.txt
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/CMakeLists.txt?rev=330271&r1=330270&r2=330271&view=diff
>
> ==============================================================================
> --- llvm/trunk/lib/Target/NVPTX/MCTargetDesc/CMakeLists.txt (original)
> +++ llvm/trunk/lib/Target/NVPTX/MCTargetDesc/CMakeLists.txt Wed Apr 18
> 09:13:41 2018
> @@ -1,4 +1,5 @@
> add_llvm_library(LLVMNVPTXDesc
> NVPTXMCAsmInfo.cpp
> NVPTXMCTargetDesc.cpp
> + NVPTXTargetStreamer.cpp
> )
>
> Modified: llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp?rev=330271&r1=330270&r2=330271&view=diff
>
> ==============================================================================
> --- llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp (original)
> +++ llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp Wed Apr 18
> 09:13:41 2018
> @@ -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 Tri
> 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 Tri
> 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";
>
> Modified: llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.h
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.h?rev=330271&r1=330270&r2=330271&view=diff
>
> ==============================================================================
> --- llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.h (original)
> +++ llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.h Wed Apr 18
> 09:13:41 2018
> @@ -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
>
>
> Modified: llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp?rev=330271&r1=330270&r2=330271&view=diff
>
> ==============================================================================
> --- llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp
> (original)
> +++ llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp Wed Apr
> 18 09:13:41 2018
> @@ -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 *createNVPTXMCInstP
> 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 LLVMInitializeNVPTXTarge
>
> // Register the MCInstPrinter.
> TargetRegistry::RegisterMCInstPrinter(*T, createNVPTXMCInstPrinter);
> +
> + // Register the MCTargetStreamer.
> + TargetRegistry::RegisterAsmTargetStreamer(*T,
> createTargetAsmStreamer);
> }
> }
>
> Added: llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp?rev=330271&view=auto
>
> ==============================================================================
> --- llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp
> (added)
> +++ llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp Wed
> Apr 18 09:13:41 2018
> @@ -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";
> + }
> +}
>
> Added: llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h?rev=330271&view=auto
>
> ==============================================================================
> --- llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h (added)
> +++ llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h Wed Apr
> 18 09:13:41 2018
> @@ -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
>
> Modified: llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp?rev=330271&r1=330270&r2=330271&view=diff
>
> ==============================================================================
> --- llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp (original)
> +++ llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp Wed Apr 18 09:13:41
> 2018
> @@ -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 Glo
> 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::EmitFunctionEntryL
> emitGlobals(*MF->getFunction().getParent());
> GlobalsEmitted = true;
> }
> -
> +
> // Set up
> MRI = &MF->getRegInfo();
> F = &MF->getFunction();
> @@ -526,14 +467,25 @@ void NVPTXAsmPrinter::EmitFunctionEntryL
>
> OutStreamer->EmitRawText(O.str());
>
> - prevDebugLoc = DebugLoc();
> -}
> -
> -void NVPTXAsmPrinter::EmitFunctionBodyStart() {
> 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() {
> SmallString<128> Str;
> raw_svector_ostream O(Str);
> emitDemotedVars(&MF->getFunction(), O);
> @@ -541,7 +493,6 @@ void NVPTXAsmPrinter::EmitFunctionBodySt
> }
>
> void NVPTXAsmPrinter::EmitFunctionBodyEnd() {
> - OutStreamer->EmitRawText(StringRef("}\n"));
> VRegMapping.clear();
> }
>
> @@ -818,42 +769,6 @@ void NVPTXAsmPrinter::emitDeclarations(c
> }
> }
>
> -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(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(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
> 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
> }
>
> 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(Mod
> 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);
>
> Modified: llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.h
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.h?rev=330271&r1=330270&r2=330271&view=diff
>
> ==============================================================================
> --- llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.h (original)
> +++ llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.h Wed Apr 18 09:13:41 2018
> @@ -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
>
> Modified: llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp?rev=330271&r1=330270&r2=330271&view=diff
>
> ==============================================================================
> --- llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp (original)
> +++ llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp Wed Apr 18 09:13:41
> 2018
> @@ -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"
> @@ -4701,31 +4700,8 @@ void NVPTXTargetLowering::ReplaceNodeRes
> }
> }
>
> -// 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 {
>
> Removed: llvm/trunk/lib/Target/NVPTX/NVPTXSection.h
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXSection.h?rev=330270&view=auto
>
> ==============================================================================
> --- llvm/trunk/lib/Target/NVPTX/NVPTXSection.h (original)
> +++ llvm/trunk/lib/Target/NVPTX/NVPTXSection.h (removed)
> @@ -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
>
> Modified: llvm/trunk/lib/Target/NVPTX/NVPTXTargetObjectFile.h
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXTargetObjectFile.h?rev=330271&r1=330270&r2=330271&view=diff
>
> ==============================================================================
> --- llvm/trunk/lib/Target/NVPTX/NVPTXTargetObjectFile.h (original)
> +++ llvm/trunk/lib/Target/NVPTX/NVPTXTargetObjectFile.h Wed Apr 18
> 09:13:41 2018
> @@ -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,
>
> Added: llvm/trunk/test/DebugInfo/NVPTX/cu-range-hole.ll
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/test/DebugInfo/NVPTX/cu-range-hole.ll?rev=330271&view=auto
>
> ==============================================================================
> --- llvm/trunk/test/DebugInfo/NVPTX/cu-range-hole.ll (added)
> +++ llvm/trunk/test/DebugInfo/NVPTX/cu-range-hole.ll Wed Apr 18 09:13:41
> 2018
> @@ -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, variables: !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, variables: !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)
>
> Modified: llvm/trunk/test/DebugInfo/NVPTX/debug-file-loc.ll
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/test/DebugInfo/NVPTX/debug-file-loc.ll?rev=330271&r1=330270&r2=330271&view=diff
>
> ==============================================================================
> --- llvm/trunk/test/DebugInfo/NVPTX/debug-file-loc.ll (original)
> +++ llvm/trunk/test/DebugInfo/NVPTX/debug-file-loc.ll Wed Apr 18 09:13:41
> 2018
> @@ -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, variables: !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)
>
> Modified: llvm/trunk/test/DebugInfo/NVPTX/debug-info.ll
> URL:
> http://llvm.org/viewvc/llvm-project/llvm/trunk/test/DebugInfo/NVPTX/debug-info.ll?rev=330271&r1=330270&r2=330271&view=diff
>
> ==============================================================================
> --- llvm/trunk/test/DebugInfo/NVPTX/debug-info.ll (original)
> +++ llvm/trunk/test/DebugInfo/NVPTX/debug-info.ll Wed Apr 18 09:13:41 2018
> @@ -9,17 +9,8 @@
> ; res(a * x[i], y[i], &y[i]);
> ;}
>
> -; CHECK: .file 1 "/some/directory{{.+}}debug-info.cu"
> -; CHECK: .file 2
> "/some/directory{{.+}}clang/include{{.+}}__clang_cuda_math_forward_declares.h"
> -; CHECK: .file 3 "{{.*}}/usr/include{{.+}}mathcalls.h"
> -; CHECK: .file 4 "{{.*}}/usr/include{{.+}}stdlib.h"
> -; CHECK: .file 5 "{{.*}}/usr/include{{.+}}stdlib-float.h"
> -; CHECK: .file 6 "{{.*}}/usr/include{{.+}}stdlib-bsearch.h"
> -; CHECK: .file 7
> "{{.*}}/usr/lib/gcc/4.8/../../../../include/c++/4.8{{.+}}cstdlib"
> -; CHECK: .file 8 "{{.*}}/usr/local/cuda/include{{.+}}math_functions.hpp"
> -; CHECK: .file 9
> "{{.*}}/usr/local/cuda/include{{.+}}device_functions.hpp"
> -; CHECK: .file 10
> "/some/directory{{.+}}clang/include{{.+}}__clang_cuda_builtin_vars.h"
> -; CHECK: .file 11 "{{.*}}/usr/local/cuda/include{{.+}}vector_types.h"
> +; CHECK: .target sm_{{[0-9]+}}//, debug
> +
> ; CHECK: .visible .entry _Z5saxpyifPfS_(
> ; CHECK: .param .u32 {{.+}},
> ; CHECK: .param .f32 {{.+}},
> @@ -31,18 +22,19 @@
> ; CHECK: .reg .f32 %f<5>;
> ; CHECK: .reg .b32 %r<6>;
> ; CHECK: .reg .b64 %rd<8>;
> +; CHECK: .loc [[DEBUG_INFO_CU:[0-9]+]] 5 0
> ; CHECK: ld.param.u32 %r{{.+}}, [{{.+}}];
> -; CHECK: .loc 10 78 180
> +; CHECK: .loc [[BUILTUIN_VARS_H:[0-9]+]] 78 180
> ; CHECK: mov.u32 %r{{.+}}, %ctaid.x;
> -; CHECK: .loc 10 89 180
> +; CHECK: .loc [[BUILTUIN_VARS_H]] 89 180
> ; CHECK: mov.u32 %r{{.+}}, %ntid.x;
> -; CHECK: .loc 10 67 180
> +; CHECK: .loc [[BUILTUIN_VARS_H]] 67 180
> ; CHECK: mov.u32 %r{{.+}}, %tid.x;
> -; CHECK: .loc 1 6 35
> +; CHECK: .loc [[DEBUG_INFO_CU]] 6 35
> ; CHECK: mad.lo.s32 %r{{.+}}, %r{{.+}}, %r{{.+}}, %r{{.+}};
> -; CHECK: .loc 1 7 9
> +; CHECK: .loc [[DEBUG_INFO_CU]] 7 9
> ; CHECK: setp.ge.s32 %p{{.+}}, %r{{.+}}, %r{{.+}};
> -; CHECK: .loc 1 7 7
> +; CHECK: .loc [[DEBUG_INFO_CU]] 7 7
> ; CHECK: @%p{{.+}} bra [[BB:.+]];
> ; CHECK: ld.param.f32 %f{{.+}}, [{{.+}}];
> ; CHECK: ld.param.u64 %rd{{.+}}, [{{.+}}];
> @@ -51,17 +43,17 @@
> ; CHECK: cvta.to.global.u64 %rd{{.+}}, %rd{{.+}};
> ; CHECK: mul.wide.u32 %rd{{.+}}, %r{{.+}}, 4;
> ; CHECK: add.s64 %rd{{.+}}, %rd{{.+}}, %rd{{.+}};
> -; CHECK: .loc 1 8 13
> +; CHECK: .loc [[DEBUG_INFO_CU]] 8 13
> ; CHECK: ld.global.f32 %f{{.+}}, [%rd{{.+}}];
> ; CHECK: add.s64 %rd{{.+}}, %rd{{.+}}, %rd{{.+}};
> -; CHECK: .loc 1 8 19
> +; CHECK: .loc [[DEBUG_INFO_CU]] 8 19
> ; CHECK: ld.global.f32 %f{{.+}}, [%rd{{.+}}];
> -; CHECK: .loc 1 3 82
> +; CHECK: .loc [[DEBUG_INFO_CU]] 3 82
> ; CHECK: fma.rn.f32 %f{{.+}}, %f{{.+}}, %f{{.+}}, %f{{.+}};
> -; CHECK: .loc 1 3 78
> +; CHECK: .loc [[DEBUG_INFO_CU]] 3 78
> ; CHECK: st.global.f32 [%rd{{.+}}], %f{{.+}};
> ; CHECK: [[BB]]:
> -; CHECK: .loc 1 9 1
> +; CHECK: .loc [[DEBUG_INFO_CU]] 9 1
> ; CHECK: ret;
> ; CHECK: }
>
> @@ -99,6 +91,8307 @@ if.end:
> ret void, !dbg !718
> }
>
> +; CHECK-DAG: .file {{[0-9]+}}
> "{{.*}}clang/include{{/|\\\\}}__clang_cuda_math_forward_declares.h"
> +; CHECK-DAG: .file {{[0-9]+}} "{{.*}}/usr/include{{/|\\\\}}mathcalls.h"
> +; CHECK-DAG: .file {{[0-9]+}}
> "{{.*}}/usr/lib/gcc/4.8/../../../../include/c++/4.8{{/|\\\\}}cmath"
> +; CHECK-DAG: .file {{[0-9]+}} "{{.*}}/usr/include{{/|\\\\}}stdlib.h"
> +; CHECK-DAG: .file {{[0-9]+}}
> "{{.*}}/usr/lib/gcc/4.8/../../../../include/c++/4.8{{/|\\\\}}cstdlib"
> +; CHECK-DAG: .file {{[0-9]+}} "{{.*}}/usr/include{{/|\\\\}}stdlib-float.h"
> +; CHECK-DAG: .file {{[0-9]+}}
> "{{.*}}/usr/include{{/|\\\\}}stdlib-bsearch.h"
> +; CHECK-DAG: .file {{[0-9]+}} "{{.*}}clang/include{{/|\\\\}}stddef.h"
> +; CHECK-DAG: .file {{[0-9]+}}
> "{{.*}}/usr/local/cuda/include{{/|\\\\}}math_functions.hpp"
> +; CHECK_DAG: .file {{[0-9]+}}
> "{{.*}}clang/include{{/|\\\\}}__clang_cuda_cmath.h"
> +; CHECK-DAG: .file {{[0-9]+}}
> "{{.*}}/usr/local/cuda/include{{/|\\\\}}device_functions.hpp"
> +; CHECK-DAG: .file [[DEBUG_INFO_CU]] "{{.*}}debug-info.cu"
> +; CHECK-DAG: .file [[BUILTUIN_VARS_H]]
> "{{.*}}clang/include{{/|\\\\}}__clang_cuda_builtin_vars.h"
> +
> +; 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 57 // DW_TAG_namespace
> +; CHECK: // .b8 1 // DW_CHILDREN_yes
> +; CHECK: // .b8 3 // DW_AT_name
> +; CHECK: // .b8 8 // DW_FORM_string
> +; CHECK: // .b8 0 // EOM(1)
> +; CHECK: // .b8 0 // EOM(2)
> +; CHECK: // .b8 3 // Abbreviation Code
> +; CHECK: // .b8 8 //
> DW_TAG_imported_declaration
> +; CHECK: // .b8 0 // DW_CHILDREN_no
> +; 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 24 // DW_AT_import
> +; CHECK: // .b8 19 // DW_FORM_ref4
> +; CHECK: // .b8 0 // EOM(1)
> +; CHECK: // .b8 0 // EOM(2)
> +; CHECK: // .b8 4 // Abbreviation Code
> +; CHECK: // .b8 8 //
> DW_TAG_imported_declaration
> +; CHECK: // .b8 0 // DW_CHILDREN_no
> +; CHECK: // .b8 58 // DW_AT_decl_file
> +; CHECK: // .b8 11 // DW_FORM_data1
> +; CHECK: // .b8 59 // DW_AT_decl_line
> +; CHECK: // .b8 5 // DW_FORM_data2
> +; CHECK: // .b8 24 // DW_AT_import
> +; CHECK: // .b8 19 // DW_FORM_ref4
> +; CHECK: // .b8 0 // EOM(1)
> +; CHECK: // .b8 0 // EOM(2)
> +; CHECK: // .b8 5 // Abbreviation Code
> +; CHECK: // .b8 46 // DW_TAG_subprogram
> +; CHECK: // .b8 1 // DW_CHILDREN_yes
> +; CHECK: // .b8 135 //
> DW_AT_MIPS_linkage_name
> +; CHECK: // .b8 64
> +; CHECK: // .b8 8 // DW_FORM_string
> +; 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 60 // DW_AT_declaration
> +; CHECK: // .b8 12 // DW_FORM_flag
> +; CHECK: // .b8 0 // EOM(1)
> +; CHECK: // .b8 0 // EOM(2)
> +; CHECK: // .b8 6 // Abbreviation Code
> +; CHECK: // .b8 5 //
> DW_TAG_formal_parameter
> +; CHECK: // .b8 0 // DW_CHILDREN_no
> +; CHECK: // .b8 73 // DW_AT_type
> +; CHECK: // .b8 19 // DW_FORM_ref4
> +; CHECK: // .b8 0 // EOM(1)
> +; CHECK: // .b8 0 // EOM(2)
> +; CHECK: // .b8 7 // 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 8 // Abbreviation Code
> +; CHECK: // .b8 15 // DW_TAG_pointer_type
> +; CHECK: // .b8 0 // DW_CHILDREN_no
> +; CHECK: // .b8 73 // DW_AT_type
> +; CHECK: // .b8 19 // DW_FORM_ref4
> +; CHECK: // .b8 0 // EOM(1)
> +; CHECK: // .b8 0 //
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20180501/a90dd0f4/attachment.html>
More information about the llvm-commits
mailing list