[llvm] r330271 - [DEBUG] Initial adaptation of NVPTX target for debug info emission.
Eric Christopher via llvm-commits
llvm-commits at lists.llvm.org
Thu May 17 20:17:41 PDT 2018
I've now reapplied this in r332689. Thank you Alexey for your patience
while we investigated!
-eric
On Mon, Apr 30, 2018 at 5:14 PM Eric Christopher <echristo at gmail.com> wrote:
> 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/20180517/e38e1891/attachment.html>
More information about the llvm-commits
mailing list