<div dir="ltr">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.<div><br></div><div>-eric</div></div><br><div class="gmail_quote"><div dir="ltr">On Wed, Apr 18, 2018 at 9:16 AM Alexey Bataev via llvm-commits <<a href="mailto:llvm-commits@lists.llvm.org">llvm-commits@lists.llvm.org</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Author: abataev<br>
Date: Wed Apr 18 09:13:41 2018<br>
New Revision: 330271<br>
<br>
URL: <a href="http://llvm.org/viewvc/llvm-project?rev=330271&view=rev" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project?rev=330271&view=rev</a><br>
Log:<br>
[DEBUG] Initial adaptation of NVPTX target for debug info emission.<br>
<br>
Summary:<br>
Patch adds initial emission of the debug info for NVPTX target.<br>
Currently, only .file and .loc directives are emitted, everything else is<br>
commented out to not break the compilation of Cuda.<br>
<br>
Reviewers: echristo, jlebar, tra, jholewinski<br>
<br>
Subscribers: mgorny, aprantl, JDevlieghere, llvm-commits<br>
<br>
Differential Revision: <a href="https://reviews.llvm.org/D41827" rel="noreferrer" target="_blank">https://reviews.llvm.org/D41827</a><br>
<br>
Added:<br>
llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp<br>
llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h<br>
llvm/trunk/test/DebugInfo/NVPTX/cu-range-hole.ll<br>
Removed:<br>
llvm/trunk/lib/Target/NVPTX/NVPTXSection.h<br>
Modified:<br>
llvm/trunk/lib/CodeGen/AsmPrinter/DwarfDebug.cpp<br>
llvm/trunk/lib/Target/NVPTX/MCTargetDesc/CMakeLists.txt<br>
llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp<br>
llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.h<br>
llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp<br>
llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp<br>
llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.h<br>
llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp<br>
llvm/trunk/lib/Target/NVPTX/NVPTXTargetObjectFile.h<br>
llvm/trunk/test/DebugInfo/NVPTX/debug-file-loc.ll<br>
llvm/trunk/test/DebugInfo/NVPTX/debug-info.ll<br>
<br>
Modified: llvm/trunk/lib/CodeGen/AsmPrinter/DwarfDebug.cpp<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/CodeGen/AsmPrinter/DwarfDebug.cpp?rev=330271&r1=330270&r2=330271&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/CodeGen/AsmPrinter/DwarfDebug.cpp?rev=330271&r1=330270&r2=330271&view=diff</a><br>
==============================================================================<br>
--- llvm/trunk/lib/CodeGen/AsmPrinter/DwarfDebug.cpp (original)<br>
+++ llvm/trunk/lib/CodeGen/AsmPrinter/DwarfDebug.cpp Wed Apr 18 09:13:41 2018<br>
@@ -312,7 +312,11 @@ DwarfDebug::DwarfDebug(AsmPrinter *A, Mo<br>
} else<br>
TheAccelTableKind = AccelTables;<br>
<br>
- UseInlineStrings = DwarfInlinedStrings == Enable;<br>
+ if (DwarfInlinedStrings == Default)<br>
+ UseInlineStrings = TT.isNVPTX();<br>
+ else<br>
+ UseInlineStrings = DwarfInlinedStrings == Enable;<br>
+<br>
HasAppleExtensionAttributes = tuneForLLDB();<br>
<br>
// Handle split DWARF.<br>
@@ -327,14 +331,18 @@ DwarfDebug::DwarfDebug(AsmPrinter *A, Mo<br>
unsigned DwarfVersionNumber = Asm->TM.Options.MCOptions.DwarfVersion;<br>
unsigned DwarfVersion = DwarfVersionNumber ? DwarfVersionNumber<br>
: MMI->getModule()->getDwarfVersion();<br>
- // Use dwarf 4 by default if nothing is requested.<br>
- DwarfVersion = DwarfVersion ? DwarfVersion : dwarf::DWARF_VERSION;<br>
-<br>
- UsePubSections = !NoDwarfPubSections;<br>
- UseRangesSection = !NoDwarfRangesSection;<br>
-<br>
- // Use sections as references.<br>
- UseSectionsAsReferences = DwarfSectionsAsReferences == Enable;<br>
+ // Use dwarf 4 by default if nothing is requested. For NVPTX, use dwarf 2.<br>
+ DwarfVersion =<br>
+ TT.isNVPTX() ? 2 : (DwarfVersion ? DwarfVersion : dwarf::DWARF_VERSION);<br>
+<br>
+ UsePubSections = !NoDwarfPubSections && !TT.isNVPTX();<br>
+ UseRangesSection = !NoDwarfRangesSection && !TT.isNVPTX();<br>
+<br>
+ // Use sections as references. Force for NVPTX.<br>
+ if (DwarfSectionsAsReferences == Default)<br>
+ UseSectionsAsReferences = TT.isNVPTX();<br>
+ else<br>
+ UseSectionsAsReferences = DwarfSectionsAsReferences == Enable;<br>
<br>
// Work around a GDB bug. GDB doesn't support the standard opcode;<br>
// SCE doesn't support GNU's; LLDB prefers the standard opcode, which<br>
<br>
Modified: llvm/trunk/lib/Target/NVPTX/MCTargetDesc/CMakeLists.txt<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/CMakeLists.txt?rev=330271&r1=330270&r2=330271&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/CMakeLists.txt?rev=330271&r1=330270&r2=330271&view=diff</a><br>
==============================================================================<br>
--- llvm/trunk/lib/Target/NVPTX/MCTargetDesc/CMakeLists.txt (original)<br>
+++ llvm/trunk/lib/Target/NVPTX/MCTargetDesc/CMakeLists.txt Wed Apr 18 09:13:41 2018<br>
@@ -1,4 +1,5 @@<br>
add_llvm_library(LLVMNVPTXDesc<br>
NVPTXMCAsmInfo.cpp<br>
NVPTXMCTargetDesc.cpp<br>
+ NVPTXTargetStreamer.cpp<br>
)<br>
<br>
Modified: llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp?rev=330271&r1=330270&r2=330271&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp?rev=330271&r1=330270&r2=330271&view=diff</a><br>
==============================================================================<br>
--- llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp (original)<br>
+++ llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.cpp Wed Apr 18 09:13:41 2018<br>
@@ -13,16 +13,9 @@<br>
<br>
#include "NVPTXMCAsmInfo.h"<br>
#include "llvm/ADT/Triple.h"<br>
-#include "llvm/Support/CommandLine.h"<br>
<br>
using namespace llvm;<br>
<br>
-// -debug-compile - Command line option to inform opt and llc passes to<br>
-// compile for debugging<br>
-static cl::opt<bool> CompileForDebugging("debug-compile",<br>
- cl::desc("Compile for debugging"),<br>
- cl::Hidden, cl::init(false));<br>
-<br>
void NVPTXMCAsmInfo::anchor() {}<br>
<br>
NVPTXMCAsmInfo::NVPTXMCAsmInfo(const Triple &TheTriple) {<br>
@@ -37,7 +30,7 @@ NVPTXMCAsmInfo::NVPTXMCAsmInfo(const Tri<br>
InlineAsmStart = " begin inline asm";<br>
InlineAsmEnd = " end inline asm";<br>
<br>
- SupportsDebugInformation = CompileForDebugging;<br>
+ SupportsDebugInformation = true;<br>
// PTX does not allow .align on functions.<br>
HasFunctionAlignment = false;<br>
HasDotTypeDotSizeDirective = false;<br>
@@ -45,13 +38,16 @@ NVPTXMCAsmInfo::NVPTXMCAsmInfo(const Tri<br>
HiddenDeclarationVisibilityAttr = HiddenVisibilityAttr = MCSA_Invalid;<br>
ProtectedVisibilityAttr = MCSA_Invalid;<br>
<br>
- Data8bitsDirective = " .b8 ";<br>
- Data16bitsDirective = " .b16 ";<br>
- Data32bitsDirective = " .b32 ";<br>
- Data64bitsDirective = " .b64 ";<br>
- ZeroDirective = " .b8";<br>
- AsciiDirective = " .b8";<br>
- AscizDirective = " .b8";<br>
+ // FIXME: remove comment once debug info is properly supported.<br>
+ Data8bitsDirective = "// .b8 ";<br>
+ Data16bitsDirective = nullptr; // not supported<br>
+ Data32bitsDirective = "// .b32 ";<br>
+ Data64bitsDirective = "// .b64 ";<br>
+ ZeroDirective = "// .b8";<br>
+ AsciiDirective = nullptr; // not supported<br>
+ AscizDirective = nullptr; // not supported<br>
+ SupportsQuotedNames = false;<br>
+ SupportsExtendedDwarfLocDirective = false;<br>
<br>
// @TODO: Can we just disable this?<br>
WeakDirective = "\t// .weak\t";<br>
<br>
Modified: llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.h<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.h?rev=330271&r1=330270&r2=330271&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.h?rev=330271&r1=330270&r2=330271&view=diff</a><br>
==============================================================================<br>
--- llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.h (original)<br>
+++ llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCAsmInfo.h Wed Apr 18 09:13:41 2018<br>
@@ -25,6 +25,17 @@ class NVPTXMCAsmInfo : public MCAsmInfo<br>
<br>
public:<br>
explicit NVPTXMCAsmInfo(const Triple &TheTriple);<br>
+<br>
+ /// Return true if the .section directive should be omitted when<br>
+ /// emitting \p SectionName. For example:<br>
+ ///<br>
+ /// shouldOmitSectionDirective(".text")<br>
+ ///<br>
+ /// returns false => .section .text,#alloc,#execinstr<br>
+ /// returns true => .text<br>
+ bool shouldOmitSectionDirective(StringRef SectionName) const override {<br>
+ return true;<br>
+ }<br>
};<br>
} // namespace llvm<br>
<br>
<br>
Modified: llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp?rev=330271&r1=330270&r2=330271&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp?rev=330271&r1=330270&r2=330271&view=diff</a><br>
==============================================================================<br>
--- llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp (original)<br>
+++ llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp Wed Apr 18 09:13:41 2018<br>
@@ -11,9 +11,10 @@<br>
//<br>
//===----------------------------------------------------------------------===//<br>
<br>
-#include "NVPTXMCTargetDesc.h"<br>
#include "InstPrinter/NVPTXInstPrinter.h"<br>
#include "NVPTXMCAsmInfo.h"<br>
+#include "NVPTXMCTargetDesc.h"<br>
+#include "NVPTXTargetStreamer.h"<br>
#include "llvm/MC/MCInstrInfo.h"<br>
#include "llvm/MC/MCRegisterInfo.h"<br>
#include "llvm/MC/MCSubtargetInfo.h"<br>
@@ -58,6 +59,12 @@ static MCInstPrinter *createNVPTXMCInstP<br>
return nullptr;<br>
}<br>
<br>
+static MCTargetStreamer *createTargetAsmStreamer(MCStreamer &S,<br>
+ formatted_raw_ostream &,<br>
+ MCInstPrinter *, bool) {<br>
+ return new NVPTXTargetStreamer(S);<br>
+}<br>
+<br>
// Force static initialization.<br>
extern "C" void LLVMInitializeNVPTXTargetMC() {<br>
for (Target *T : {&getTheNVPTXTarget32(), &getTheNVPTXTarget64()}) {<br>
@@ -75,5 +82,8 @@ extern "C" void LLVMInitializeNVPTXTarge<br>
<br>
// Register the MCInstPrinter.<br>
TargetRegistry::RegisterMCInstPrinter(*T, createNVPTXMCInstPrinter);<br>
+<br>
+ // Register the MCTargetStreamer.<br>
+ TargetRegistry::RegisterAsmTargetStreamer(*T, createTargetAsmStreamer);<br>
}<br>
}<br>
<br>
Added: llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp?rev=330271&view=auto" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp?rev=330271&view=auto</a><br>
==============================================================================<br>
--- llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp (added)<br>
+++ llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp Wed Apr 18 09:13:41 2018<br>
@@ -0,0 +1,94 @@<br>
+//=====- NVPTXTargetStreamer.cpp - NVPTXTargetStreamer class ------------=====//<br>
+//<br>
+// The LLVM Compiler Infrastructure<br>
+//<br>
+// This file is distributed under the University of Illinois Open Source<br>
+// License. See LICENSE.TXT for details.<br>
+//<br>
+//===----------------------------------------------------------------------===//<br>
+//<br>
+// This file implements the NVPTXTargetStreamer class.<br>
+//<br>
+//===----------------------------------------------------------------------===//<br>
+<br>
+#include "NVPTXTargetStreamer.h"<br>
+#include "llvm/MC/MCAsmInfo.h"<br>
+#include "llvm/MC/MCContext.h"<br>
+#include "llvm/MC/MCObjectFileInfo.h"<br>
+<br>
+using namespace llvm;<br>
+<br>
+//<br>
+// NVPTXTargetStreamer Implemenation<br>
+//<br>
+NVPTXTargetStreamer::NVPTXTargetStreamer(MCStreamer &S) : MCTargetStreamer(S) {}<br>
+<br>
+NVPTXTargetStreamer::~NVPTXTargetStreamer() = default;<br>
+<br>
+void NVPTXTargetStreamer::emitDwarfFileDirective(StringRef Directive) {<br>
+ DwarfFiles.emplace_back(Directive);<br>
+}<br>
+<br>
+static bool isDwarfSection(const MCObjectFileInfo *FI,<br>
+ const MCSection *Section) {<br>
+ // FIXME: the checks for the DWARF sections are very fragile and should be<br>
+ // fixed up in a followup patch.<br>
+ if (!Section || Section->getKind().isText() ||<br>
+ Section->getKind().isWriteable())<br>
+ return false;<br>
+ return Section == FI->getDwarfAbbrevSection() ||<br>
+ Section == FI->getDwarfInfoSection() ||<br>
+ Section == FI->getDwarfMacinfoSection() ||<br>
+ Section == FI->getDwarfFrameSection() ||<br>
+ Section == FI->getDwarfAddrSection() ||<br>
+ Section == FI->getDwarfRangesSection() ||<br>
+ Section == FI->getDwarfARangesSection() ||<br>
+ Section == FI->getDwarfLocSection() ||<br>
+ Section == FI->getDwarfStrSection() ||<br>
+ Section == FI->getDwarfLineSection() ||<br>
+ Section == FI->getDwarfStrOffSection() ||<br>
+ Section == FI->getDwarfLineStrSection() ||<br>
+ Section == FI->getDwarfPubNamesSection() ||<br>
+ Section == FI->getDwarfPubTypesSection() ||<br>
+ Section == FI->getDwarfSwiftASTSection() ||<br>
+ Section == FI->getDwarfTypesDWOSection() ||<br>
+ Section == FI->getDwarfAbbrevDWOSection() ||<br>
+ Section == FI->getDwarfAccelObjCSection() ||<br>
+ Section == FI->getDwarfAccelNamesSection() ||<br>
+ Section == FI->getDwarfAccelTypesSection() ||<br>
+ Section == FI->getDwarfAccelNamespaceSection() ||<br>
+ Section == FI->getDwarfLocDWOSection() ||<br>
+ Section == FI->getDwarfStrDWOSection() ||<br>
+ Section == FI->getDwarfCUIndexSection() ||<br>
+ Section == FI->getDwarfInfoDWOSection() ||<br>
+ Section == FI->getDwarfLineDWOSection() ||<br>
+ Section == FI->getDwarfTUIndexSection() ||<br>
+ Section == FI->getDwarfStrOffDWOSection() ||<br>
+ Section == FI->getDwarfDebugNamesSection() ||<br>
+ Section == FI->getDwarfDebugInlineSection() ||<br>
+ Section == FI->getDwarfGnuPubNamesSection() ||<br>
+ Section == FI->getDwarfGnuPubTypesSection();<br>
+}<br>
+<br>
+void NVPTXTargetStreamer::changeSection(const MCSection *CurSection,<br>
+ MCSection *Section,<br>
+ const MCExpr *SubSection,<br>
+ raw_ostream &OS) {<br>
+ assert(!SubSection && "SubSection is not null!");<br>
+ const MCObjectFileInfo *FI = getStreamer().getContext().getObjectFileInfo();<br>
+ // FIXME: remove comment once debug info is properly supported.<br>
+ // Emit closing brace for DWARF sections only.<br>
+ if (isDwarfSection(FI, CurSection))<br>
+ OS << "//\t}\n";<br>
+ if (isDwarfSection(FI, Section)) {<br>
+ // Emit DWARF .file directives in the outermost scope.<br>
+ for (const std::string &S : DwarfFiles)<br>
+ getStreamer().EmitRawText(S.data());<br>
+ DwarfFiles.clear();<br>
+ OS << "//\t.section";<br>
+ Section->PrintSwitchToSection(*getStreamer().getContext().getAsmInfo(),<br>
+ FI->getTargetTriple(), OS, SubSection);<br>
+ // DWARF sections are enclosed into braces - emit the open one.<br>
+ OS << "//\t{\n";<br>
+ }<br>
+}<br>
<br>
Added: llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h?rev=330271&view=auto" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h?rev=330271&view=auto</a><br>
==============================================================================<br>
--- llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h (added)<br>
+++ llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h Wed Apr 18 09:13:41 2018<br>
@@ -0,0 +1,46 @@<br>
+//=====-- NVPTXTargetStreamer.h - NVPTX Target Streamer ------*- C++ -*--=====//<br>
+//<br>
+// The LLVM Compiler Infrastructure<br>
+//<br>
+// This file is distributed under the University of Illinois Open Source<br>
+// License. See LICENSE.TXT for details.<br>
+//<br>
+//===----------------------------------------------------------------------===//<br>
+<br>
+#ifndef LLVM_LIB_TARGET_NVPTX_MCTARGETDESC_NVPTXTARGETSTREAMER_H<br>
+#define LLVM_LIB_TARGET_NVPTX_MCTARGETDESC_NVPTXTARGETSTREAMER_H<br>
+<br>
+#include "llvm/MC/MCStreamer.h"<br>
+<br>
+namespace llvm {<br>
+class MCSection;<br>
+<br>
+/// Implments NVPTX-specific streamer.<br>
+class NVPTXTargetStreamer : public MCTargetStreamer {<br>
+private:<br>
+ SmallVector<std::string, 4> DwarfFiles;<br>
+<br>
+public:<br>
+ NVPTXTargetStreamer(MCStreamer &S);<br>
+ ~NVPTXTargetStreamer() override;<br>
+<br>
+ /// Record DWARF file directives for later output.<br>
+ /// According to PTX ISA, CUDA Toolkit documentation, 11.5.3. Debugging<br>
+ /// Directives: .file<br>
+ /// (<a href="http://docs.nvidia.com/cuda/parallel-thread-execution/index.html#debugging-directives-file" rel="noreferrer" target="_blank">http://docs.nvidia.com/cuda/parallel-thread-execution/index.html#debugging-directives-file</a>),<br>
+ /// The .file directive is allowed only in the outermost scope, i.e., at the<br>
+ /// same level as kernel and device function declarations. Also, the order of<br>
+ /// the .loc and .file directive does not matter, .file directives may follow<br>
+ /// the .loc directives where the file is referenced.<br>
+ /// LLVM emits .file directives immediately the location debug info is<br>
+ /// emitted, i.e. they may be emitted inside functions. We gather all these<br>
+ /// directives and emit them outside of the sections and, thus, outside of the<br>
+ /// functions.<br>
+ void emitDwarfFileDirective(StringRef Directive) override;<br>
+ void changeSection(const MCSection *CurSection, MCSection *Section,<br>
+ const MCExpr *SubSection, raw_ostream &OS) override;<br>
+};<br>
+<br>
+} // end namespace llvm<br>
+<br>
+#endif<br>
<br>
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp?rev=330271&r1=330270&r2=330271&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp?rev=330271&r1=330270&r2=330271&view=diff</a><br>
==============================================================================<br>
--- llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp (original)<br>
+++ llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp Wed Apr 18 09:13:41 2018<br>
@@ -93,16 +93,6 @@ using namespace llvm;<br>
<br>
#define DEPOTNAME "__local_depot"<br>
<br>
-static cl::opt<bool><br>
-EmitLineNumbers("nvptx-emit-line-numbers", cl::Hidden,<br>
- cl::desc("NVPTX Specific: Emit Line numbers even without -G"),<br>
- cl::init(true));<br>
-<br>
-static cl::opt<bool><br>
-InterleaveSrc("nvptx-emit-src", cl::ZeroOrMore, cl::Hidden,<br>
- cl::desc("NVPTX Specific: Emit source line in ptx file"),<br>
- cl::init(false));<br>
-<br>
/// DiscoverDependentGlobals - Return a set of GlobalVariables on which \p V<br>
/// depends.<br>
static void<br>
@@ -151,56 +141,7 @@ VisitGlobalVariableForEmission(const Glo<br>
Visiting.erase(GV);<br>
}<br>
<br>
-void NVPTXAsmPrinter::emitLineNumberAsDotLoc(const MachineInstr &MI) {<br>
- if (!EmitLineNumbers)<br>
- return;<br>
- if (ignoreLoc(MI))<br>
- return;<br>
-<br>
- const DebugLoc &curLoc = MI.getDebugLoc();<br>
-<br>
- if (!prevDebugLoc && !curLoc)<br>
- return;<br>
-<br>
- if (prevDebugLoc == curLoc)<br>
- return;<br>
-<br>
- prevDebugLoc = curLoc;<br>
-<br>
- if (!curLoc)<br>
- return;<br>
-<br>
- auto *Scope = cast_or_null<DIScope>(curLoc.getScope());<br>
- if (!Scope)<br>
- return;<br>
-<br>
- StringRef fileName(Scope->getFilename());<br>
- StringRef dirName(Scope->getDirectory());<br>
- SmallString<128> FullPathName = dirName;<br>
- if (!dirName.empty() && !sys::path::is_absolute(fileName)) {<br>
- sys::path::append(FullPathName, fileName);<br>
- fileName = FullPathName;<br>
- }<br>
-<br>
- if (filenameMap.find(fileName) == filenameMap.end())<br>
- return;<br>
-<br>
- // Emit the line from the source file.<br>
- if (InterleaveSrc)<br>
- this->emitSrcInText(fileName, curLoc.getLine());<br>
-<br>
- std::stringstream temp;<br>
- temp << "\t.loc " << filenameMap[fileName] << " " << curLoc.getLine()<br>
- << " " << curLoc.getCol();<br>
- OutStreamer->EmitRawText(temp.str());<br>
-}<br>
-<br>
void NVPTXAsmPrinter::EmitInstruction(const MachineInstr *MI) {<br>
- SmallString<128> Str;<br>
- raw_svector_ostream OS(Str);<br>
- if (static_cast<NVPTXTargetMachine &>(TM).getDrvInterface() == NVPTX::CUDA)<br>
- emitLineNumberAsDotLoc(*MI);<br>
-<br>
MCInst Inst;<br>
lowerToMCInst(MI, Inst);<br>
EmitToStreamer(*OutStreamer, Inst);<br>
@@ -505,7 +446,7 @@ void NVPTXAsmPrinter::EmitFunctionEntryL<br>
emitGlobals(*MF->getFunction().getParent());<br>
GlobalsEmitted = true;<br>
}<br>
- <br>
+<br>
// Set up<br>
MRI = &MF->getRegInfo();<br>
F = &MF->getFunction();<br>
@@ -526,14 +467,25 @@ void NVPTXAsmPrinter::EmitFunctionEntryL<br>
<br>
OutStreamer->EmitRawText(O.str());<br>
<br>
- prevDebugLoc = DebugLoc();<br>
-}<br>
-<br>
-void NVPTXAsmPrinter::EmitFunctionBodyStart() {<br>
VRegMapping.clear();<br>
+ // Emit open brace for function body.<br>
OutStreamer->EmitRawText(StringRef("{\n"));<br>
setAndEmitFunctionVirtualRegisters(*MF);<br>
+}<br>
+<br>
+bool NVPTXAsmPrinter::runOnMachineFunction(MachineFunction &F) {<br>
+ nvptxSubtarget = &F.getSubtarget<NVPTXSubtarget>();<br>
+ bool Result = AsmPrinter::runOnMachineFunction(F);<br>
+ // Emit closing brace for the body of function F.<br>
+ // The closing brace must be emitted here because we need to emit additional<br>
+ // debug labels/data after the last basic block.<br>
+ // We need to emit the closing brace here because we don't have function that<br>
+ // finished emission of the function body.<br>
+ OutStreamer->EmitRawText(StringRef("}\n"));<br>
+ return Result;<br>
+}<br>
<br>
+void NVPTXAsmPrinter::EmitFunctionBodyStart() {<br>
SmallString<128> Str;<br>
raw_svector_ostream O(Str);<br>
emitDemotedVars(&MF->getFunction(), O);<br>
@@ -541,7 +493,6 @@ void NVPTXAsmPrinter::EmitFunctionBodySt<br>
}<br>
<br>
void NVPTXAsmPrinter::EmitFunctionBodyEnd() {<br>
- OutStreamer->EmitRawText(StringRef("}\n"));<br>
VRegMapping.clear();<br>
}<br>
<br>
@@ -818,42 +769,6 @@ void NVPTXAsmPrinter::emitDeclarations(c<br>
}<br>
}<br>
<br>
-void NVPTXAsmPrinter::recordAndEmitFilenames(Module &M) {<br>
- DebugInfoFinder DbgFinder;<br>
- DbgFinder.processModule(M);<br>
-<br>
- unsigned i = 1;<br>
- for (const DICompileUnit *DIUnit : DbgFinder.compile_units()) {<br>
- StringRef Filename = DIUnit->getFilename();<br>
- StringRef Dirname = DIUnit->getDirectory();<br>
- SmallString<128> FullPathName = Dirname;<br>
- if (!Dirname.empty() && !sys::path::is_absolute(Filename)) {<br>
- sys::path::append(FullPathName, Filename);<br>
- Filename = FullPathName;<br>
- }<br>
- if (filenameMap.find(Filename) != filenameMap.end())<br>
- continue;<br>
- filenameMap[Filename] = i;<br>
- OutStreamer->EmitDwarfFileDirective(i, "", Filename);<br>
- ++i;<br>
- }<br>
-<br>
- for (DISubprogram *SP : DbgFinder.subprograms()) {<br>
- StringRef Filename = SP->getFilename();<br>
- StringRef Dirname = SP->getDirectory();<br>
- SmallString<128> FullPathName = Dirname;<br>
- if (!Dirname.empty() && !sys::path::is_absolute(Filename)) {<br>
- sys::path::append(FullPathName, Filename);<br>
- Filename = FullPathName;<br>
- }<br>
- if (filenameMap.find(Filename) != filenameMap.end())<br>
- continue;<br>
- filenameMap[Filename] = i;<br>
- OutStreamer->EmitDwarfFileDirective(i, "", Filename);<br>
- ++i;<br>
- }<br>
-}<br>
-<br>
static bool isEmptyXXStructor(GlobalVariable *GV) {<br>
if (!GV) return true;<br>
const ConstantArray *InitList = dyn_cast<ConstantArray>(GV->getInitializer());<br>
@@ -889,24 +804,13 @@ bool NVPTXAsmPrinter::doInitialization(M<br>
SmallString<128> Str1;<br>
raw_svector_ostream OS1(Str1);<br>
<br>
- MMI = getAnalysisIfAvailable<MachineModuleInfo>();<br>
-<br>
// We need to call the parent's one explicitly.<br>
- //bool Result = AsmPrinter::doInitialization(M);<br>
-<br>
- // Initialize TargetLoweringObjectFile since we didn't do in<br>
- // AsmPrinter::doInitialization either right above or where it's commented out<br>
- // below.<br>
- const_cast<TargetLoweringObjectFile &>(getObjFileLowering())<br>
- .Initialize(OutContext, TM);<br>
+ bool Result = AsmPrinter::doInitialization(M);<br>
<br>
// Emit header before any dwarf directives are emitted below.<br>
emitHeader(M, OS1, STI);<br>
OutStreamer->EmitRawText(OS1.str());<br>
<br>
- // Already commented out<br>
- //bool Result = AsmPrinter::doInitialization(M);<br>
-<br>
// Emit module-level inline asm if it exists.<br>
if (!M.getModuleInlineAsm().empty()) {<br>
OutStreamer->AddComment("Start of file scope inline assembly");<br>
@@ -917,13 +821,9 @@ bool NVPTXAsmPrinter::doInitialization(M<br>
OutStreamer->AddBlankLine();<br>
}<br>
<br>
- // If we're not NVCL we're CUDA, go ahead and emit filenames.<br>
- if (TM.getTargetTriple().getOS() != Triple::NVCL)<br>
- recordAndEmitFilenames(M);<br>
-<br>
GlobalsEmitted = false;<br>
- <br>
- return false; // success<br>
+<br>
+ return Result;<br>
}<br>
<br>
void NVPTXAsmPrinter::emitGlobals(const Module &M) {<br>
@@ -975,8 +875,9 @@ void NVPTXAsmPrinter::emitHeader(Module<br>
if (NTM.getDrvInterface() == NVPTX::NVCL)<br>
O << ", texmode_independent";<br>
<br>
- if (MAI->doesSupportDebugInformation())<br>
- O << ", debug";<br>
+ // FIXME: remove comment once debug info is properly supported.<br>
+ if (MMI && MMI->hasDebugInfo())<br>
+ O << "//, debug";<br>
<br>
O << "\n";<br>
<br>
@@ -991,6 +892,8 @@ void NVPTXAsmPrinter::emitHeader(Module<br>
}<br>
<br>
bool NVPTXAsmPrinter::doFinalization(Module &M) {<br>
+ bool HasDebugInfo = MMI && MMI->hasDebugInfo();<br>
+<br>
// If we did not emit any functions, then the global declarations have not<br>
// yet been emitted.<br>
if (!GlobalsEmitted) {<br>
@@ -1025,6 +928,11 @@ bool NVPTXAsmPrinter::doFinalization(Mod<br>
clearAnnotationCache(&M);<br>
<br>
delete[] gv_array;<br>
+ // FIXME: remove comment once debug info is properly supported.<br>
+ // Close the last emitted section<br>
+ if (HasDebugInfo)<br>
+ OutStreamer->EmitRawText("//\t}");<br>
+<br>
return ret;<br>
<br>
//bool Result = AsmPrinter::doFinalization(M);<br>
<br>
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.h<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.h?rev=330271&r1=330270&r2=330271&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.h?rev=330271&r1=330270&r2=330271&view=diff</a><br>
==============================================================================<br>
--- llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.h (original)<br>
+++ llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.h Wed Apr 18 09:13:41 2018<br>
@@ -344,10 +344,7 @@ public:<br>
delete reader;<br>
}<br>
<br>
- bool runOnMachineFunction(MachineFunction &F) override {<br>
- nvptxSubtarget = &F.getSubtarget<NVPTXSubtarget>();<br>
- return AsmPrinter::runOnMachineFunction(F);<br>
- }<br>
+ bool runOnMachineFunction(MachineFunction &F) override;<br>
<br>
void getAnalysisUsage(AnalysisUsage &AU) const override {<br>
AU.addRequired<MachineLoopInfo>();<br>
@@ -357,9 +354,6 @@ public:<br>
bool ignoreLoc(const MachineInstr &);<br>
<br>
std::string getVirtualRegisterName(unsigned) const;<br>
-<br>
- DebugLoc prevDebugLoc;<br>
- void emitLineNumberAsDotLoc(const MachineInstr &);<br>
};<br>
<br>
} // end namespace llvm<br>
<br>
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp?rev=330271&r1=330270&r2=330271&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp?rev=330271&r1=330270&r2=330271&view=diff</a><br>
==============================================================================<br>
--- llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp (original)<br>
+++ llvm/trunk/lib/Target/NVPTX/NVPTXISelLowering.cpp Wed Apr 18 09:13:41 2018<br>
@@ -15,7 +15,6 @@<br>
#include "NVPTXISelLowering.h"<br>
#include "MCTargetDesc/NVPTXBaseInfo.h"<br>
#include "NVPTX.h"<br>
-#include "NVPTXSection.h"<br>
#include "NVPTXSubtarget.h"<br>
#include "NVPTXTargetMachine.h"<br>
#include "NVPTXTargetObjectFile.h"<br>
@@ -4701,31 +4700,8 @@ void NVPTXTargetLowering::ReplaceNodeRes<br>
}<br>
}<br>
<br>
-// Pin NVPTXSection's and NVPTXTargetObjectFile's vtables to this file.<br>
-void NVPTXSection::anchor() {}<br>
-<br>
-NVPTXTargetObjectFile::~NVPTXTargetObjectFile() {<br>
- delete static_cast<NVPTXSection *>(TextSection);<br>
- delete static_cast<NVPTXSection *>(DataSection);<br>
- delete static_cast<NVPTXSection *>(BSSSection);<br>
- delete static_cast<NVPTXSection *>(ReadOnlySection);<br>
-<br>
- delete static_cast<NVPTXSection *>(StaticCtorSection);<br>
- delete static_cast<NVPTXSection *>(StaticDtorSection);<br>
- delete static_cast<NVPTXSection *>(LSDASection);<br>
- delete static_cast<NVPTXSection *>(EHFrameSection);<br>
- delete static_cast<NVPTXSection *>(DwarfAbbrevSection);<br>
- delete static_cast<NVPTXSection *>(DwarfInfoSection);<br>
- delete static_cast<NVPTXSection *>(DwarfLineSection);<br>
- delete static_cast<NVPTXSection *>(DwarfFrameSection);<br>
- delete static_cast<NVPTXSection *>(DwarfPubTypesSection);<br>
- delete static_cast<const NVPTXSection *>(DwarfDebugInlineSection);<br>
- delete static_cast<NVPTXSection *>(DwarfStrSection);<br>
- delete static_cast<NVPTXSection *>(DwarfLocSection);<br>
- delete static_cast<NVPTXSection *>(DwarfARangesSection);<br>
- delete static_cast<NVPTXSection *>(DwarfRangesSection);<br>
- delete static_cast<NVPTXSection *>(DwarfMacinfoSection);<br>
-}<br>
+// Pin NVPTXTargetObjectFile's vtables to this file.<br>
+NVPTXTargetObjectFile::~NVPTXTargetObjectFile() {}<br>
<br>
MCSection *NVPTXTargetObjectFile::SelectSectionForGlobal(<br>
const GlobalObject *GO, SectionKind Kind, const TargetMachine &TM) const {<br>
<br>
Removed: llvm/trunk/lib/Target/NVPTX/NVPTXSection.h<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXSection.h?rev=330270&view=auto" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXSection.h?rev=330270&view=auto</a><br>
==============================================================================<br>
--- llvm/trunk/lib/Target/NVPTX/NVPTXSection.h (original)<br>
+++ llvm/trunk/lib/Target/NVPTX/NVPTXSection.h (removed)<br>
@@ -1,45 +0,0 @@<br>
-//===- NVPTXSection.h - NVPTX-specific section representation ---*- C++ -*-===//<br>
-//<br>
-// The LLVM Compiler Infrastructure<br>
-//<br>
-// This file is distributed under the University of Illinois Open Source<br>
-// License. See LICENSE.TXT for details.<br>
-//<br>
-//===----------------------------------------------------------------------===//<br>
-//<br>
-// This file declares the NVPTXSection class.<br>
-//<br>
-//===----------------------------------------------------------------------===//<br>
-<br>
-#ifndef LLVM_LIB_TARGET_NVPTX_NVPTXSECTION_H<br>
-#define LLVM_LIB_TARGET_NVPTX_NVPTXSECTION_H<br>
-<br>
-#include "llvm/MC/MCSection.h"<br>
-#include "llvm/MC/SectionKind.h"<br>
-<br>
-namespace llvm {<br>
-<br>
-/// Represents a section in PTX PTX does not have sections. We create this class<br>
-/// in order to use the ASMPrint interface.<br>
-///<br>
-class NVPTXSection final : public MCSection {<br>
- virtual void anchor();<br>
-<br>
-public:<br>
- NVPTXSection(SectionVariant V, SectionKind K) : MCSection(V, K, nullptr) {}<br>
- ~NVPTXSection() = default;<br>
-<br>
- /// Override this as NVPTX has its own way of printing switching<br>
- /// to a section.<br>
- void PrintSwitchToSection(const MCAsmInfo &MAI, const Triple &T,<br>
- raw_ostream &OS,<br>
- const MCExpr *Subsection) const override {}<br>
-<br>
- /// Base address of PTX sections is zero.<br>
- bool UseCodeAlign() const override { return false; }<br>
- bool isVirtualSection() const override { return false; }<br>
-};<br>
-<br>
-} // end namespace llvm<br>
-<br>
-#endif // LLVM_LIB_TARGET_NVPTX_NVPTXSECTION_H<br>
<br>
Modified: llvm/trunk/lib/Target/NVPTX/NVPTXTargetObjectFile.h<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXTargetObjectFile.h?rev=330271&r1=330270&r2=330271&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXTargetObjectFile.h?rev=330271&r1=330270&r2=330271&view=diff</a><br>
==============================================================================<br>
--- llvm/trunk/lib/Target/NVPTX/NVPTXTargetObjectFile.h (original)<br>
+++ llvm/trunk/lib/Target/NVPTX/NVPTXTargetObjectFile.h Wed Apr 18 09:13:41 2018<br>
@@ -10,7 +10,6 @@<br>
#ifndef LLVM_LIB_TARGET_NVPTX_NVPTXTARGETOBJECTFILE_H<br>
#define LLVM_LIB_TARGET_NVPTX_NVPTXTARGETOBJECTFILE_H<br>
<br>
-#include "NVPTXSection.h"<br>
#include "llvm/MC/MCSection.h"<br>
#include "llvm/MC/SectionKind.h"<br>
#include "llvm/Target/TargetLoweringObjectFile.h"<br>
@@ -19,68 +18,12 @@ namespace llvm {<br>
<br>
class NVPTXTargetObjectFile : public TargetLoweringObjectFile {<br>
public:<br>
- NVPTXTargetObjectFile() {<br>
- TextSection = nullptr;<br>
- DataSection = nullptr;<br>
- BSSSection = nullptr;<br>
- ReadOnlySection = nullptr;<br>
-<br>
- StaticCtorSection = nullptr;<br>
- StaticDtorSection = nullptr;<br>
- LSDASection = nullptr;<br>
- EHFrameSection = nullptr;<br>
- DwarfAbbrevSection = nullptr;<br>
- DwarfInfoSection = nullptr;<br>
- DwarfLineSection = nullptr;<br>
- DwarfFrameSection = nullptr;<br>
- DwarfPubTypesSection = nullptr;<br>
- DwarfDebugInlineSection = nullptr;<br>
- DwarfStrSection = nullptr;<br>
- DwarfLocSection = nullptr;<br>
- DwarfARangesSection = nullptr;<br>
- DwarfRangesSection = nullptr;<br>
- DwarfMacinfoSection = nullptr;<br>
- }<br>
+ NVPTXTargetObjectFile() : TargetLoweringObjectFile() {}<br>
<br>
~NVPTXTargetObjectFile() override;<br>
<br>
void Initialize(MCContext &ctx, const TargetMachine &TM) override {<br>
TargetLoweringObjectFile::Initialize(ctx, TM);<br>
- TextSection = new NVPTXSection(MCSection::SV_ELF, SectionKind::getText());<br>
- DataSection = new NVPTXSection(MCSection::SV_ELF, SectionKind::getData());<br>
- BSSSection = new NVPTXSection(MCSection::SV_ELF, SectionKind::getBSS());<br>
- ReadOnlySection =<br>
- new NVPTXSection(MCSection::SV_ELF, SectionKind::getReadOnly());<br>
- StaticCtorSection =<br>
- new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());<br>
- StaticDtorSection =<br>
- new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());<br>
- LSDASection =<br>
- new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());<br>
- EHFrameSection =<br>
- new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());<br>
- DwarfAbbrevSection =<br>
- new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());<br>
- DwarfInfoSection =<br>
- new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());<br>
- DwarfLineSection =<br>
- new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());<br>
- DwarfFrameSection =<br>
- new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());<br>
- DwarfPubTypesSection =<br>
- new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());<br>
- DwarfDebugInlineSection =<br>
- new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());<br>
- DwarfStrSection =<br>
- new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());<br>
- DwarfLocSection =<br>
- new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());<br>
- DwarfARangesSection =<br>
- new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());<br>
- DwarfRangesSection =<br>
- new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());<br>
- DwarfMacinfoSection =<br>
- new NVPTXSection(MCSection::SV_ELF, SectionKind::getMetadata());<br>
}<br>
<br>
MCSection *getSectionForConstant(const DataLayout &DL, SectionKind Kind,<br>
<br>
Added: llvm/trunk/test/DebugInfo/NVPTX/cu-range-hole.ll<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/test/DebugInfo/NVPTX/cu-range-hole.ll?rev=330271&view=auto" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/llvm/trunk/test/DebugInfo/NVPTX/cu-range-hole.ll?rev=330271&view=auto</a><br>
==============================================================================<br>
--- llvm/trunk/test/DebugInfo/NVPTX/cu-range-hole.ll (added)<br>
+++ llvm/trunk/test/DebugInfo/NVPTX/cu-range-hole.ll Wed Apr 18 09:13:41 2018<br>
@@ -0,0 +1,291 @@<br>
+; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s<br>
+<br>
+; CHECK: .target sm_{{[0-9]+}}//, debug<br>
+<br>
+; CHECK: .visible .func (.param .b32 func_retval0) b(<br>
+; CHECK: .param .b32 b_param_0<br>
+; CHECK: )<br>
+; CHECK: {<br>
+; CHECK: Lfunc_begin0:<br>
+; CHECK: .loc 1 1 0<br>
+; CHECK: .loc 1 1 0<br>
+; CHECK: ret;<br>
+; CHECK: Lfunc_end0:<br>
+; CHECK: }<br>
+<br>
+; CHECK: .visible .func (.param .b32 func_retval0) a(<br>
+; CHECK: .param .b32 a_param_0<br>
+; CHECK: )<br>
+; CHECK: {<br>
+; CHECK: Lfunc_begin1:<br>
+; CHECK-NOT: .loc<br>
+; CHECK: ret;<br>
+; CHECK: Lfunc_end1:<br>
+; CHECK: }<br>
+<br>
+; CHECK: .visible .func (.param .b32 func_retval0) d(<br>
+; CHECK: .param .b32 d_param_0<br>
+; CHECK: )<br>
+; CHECK: {<br>
+; CHECK: Lfunc_begin2:<br>
+; CHECK: .loc 1 3 0<br>
+; CHECK: ret;<br>
+; CHECK: Lfunc_end2:<br>
+; CHECK: }<br>
+<br>
+; CHECK: .file 1 "{{.*}}b.c"<br>
+<br>
+; Function Attrs: nounwind uwtable<br>
+define i32 @b(i32 %c) #0 !dbg !5 {<br>
+entry:<br>
+ %c.addr = alloca i32, align 4<br>
+ store i32 %c, i32* %c.addr, align 4<br>
+ call void @llvm.dbg.declare(metadata i32* %c.addr, metadata !13, metadata !DIExpression()), !dbg !14<br>
+ %0 = load i32, i32* %c.addr, align 4, !dbg !14<br>
+ %add = add nsw i32 %0, 1, !dbg !14<br>
+ ret i32 %add, !dbg !14<br>
+}<br>
+<br>
+; Function Attrs: nounwind uwtable<br>
+define i32 @a(i32 %b) #0 {<br>
+entry:<br>
+ %b.addr = alloca i32, align 4<br>
+ store i32 %b, i32* %b.addr, align 4<br>
+ %0 = load i32, i32* %b.addr, align 4<br>
+ %add = add nsw i32 %0, 1<br>
+ ret i32 %add<br>
+}<br>
+<br>
+; Function Attrs: nounwind readnone<br>
+declare void @llvm.dbg.declare(metadata, metadata, metadata) #1<br>
+<br>
+; Function Attrs: nounwind uwtable<br>
+define i32 @d(i32 %e) #0 !dbg !10 {<br>
+entry:<br>
+ %e.addr = alloca i32, align 4<br>
+ store i32 %e, i32* %e.addr, align 4<br>
+ call void @llvm.dbg.declare(metadata i32* %e.addr, metadata !15, metadata !DIExpression()), !dbg !16<br>
+ %0 = load i32, i32* %e.addr, align 4, !dbg !16<br>
+ %add = add nsw i32 %0, 1, !dbg !16<br>
+ ret i32 %add, !dbg !16<br>
+}<br>
+<br>
+; CHECK: // .section .debug_abbrev<br>
+; CHECK: // {<br>
+; CHECK: // .b8 1 // Abbreviation Code<br>
+; CHECK: // .b8 17 // DW_TAG_compile_unit<br>
+; CHECK: // .b8 1 // DW_CHILDREN_yes<br>
+; CHECK: // .b8 37 // DW_AT_producer<br>
+; CHECK: // .b8 8 // DW_FORM_string<br>
+; CHECK: // .b8 19 // DW_AT_language<br>
+; CHECK: // .b8 5 // DW_FORM_data2<br>
+; CHECK: // .b8 3 // DW_AT_name<br>
+; CHECK: // .b8 8 // DW_FORM_string<br>
+; CHECK: // .b8 16 // DW_AT_stmt_list<br>
+; CHECK: // .b8 6 // DW_FORM_data4<br>
+; CHECK: // .b8 27 // DW_AT_comp_dir<br>
+; CHECK: // .b8 8 // DW_FORM_string<br>
+; CHECK: // .b8 17 // DW_AT_low_pc<br>
+; CHECK: // .b8 1 // DW_FORM_addr<br>
+; CHECK: // .b8 18 // DW_AT_high_pc<br>
+; CHECK: // .b8 1 // DW_FORM_addr<br>
+; CHECK: // .b8 0 // EOM(1)<br>
+; CHECK: // .b8 0 // EOM(2)<br>
+; CHECK: // .b8 2 // Abbreviation Code<br>
+; CHECK: // .b8 46 // DW_TAG_subprogram<br>
+; CHECK: // .b8 1 // DW_CHILDREN_yes<br>
+; CHECK: // .b8 17 // DW_AT_low_pc<br>
+; CHECK: // .b8 1 // DW_FORM_addr<br>
+; CHECK: // .b8 18 // DW_AT_high_pc<br>
+; CHECK: // .b8 1 // DW_FORM_addr<br>
+; CHECK: // .b8 3 // DW_AT_name<br>
+; CHECK: // .b8 8 // DW_FORM_string<br>
+; CHECK: // .b8 58 // DW_AT_decl_file<br>
+; CHECK: // .b8 11 // DW_FORM_data1<br>
+; CHECK: // .b8 59 // DW_AT_decl_line<br>
+; CHECK: // .b8 11 // DW_FORM_data1<br>
+; CHECK: // .b8 39 // DW_AT_prototyped<br>
+; CHECK: // .b8 12 // DW_FORM_flag<br>
+; CHECK: // .b8 73 // DW_AT_type<br>
+; CHECK: // .b8 19 // DW_FORM_ref4<br>
+; CHECK: // .b8 63 // DW_AT_external<br>
+; CHECK: // .b8 12 // DW_FORM_flag<br>
+; CHECK: // .b8 0 // EOM(1)<br>
+; CHECK: // .b8 0 // EOM(2)<br>
+; CHECK: // .b8 3 // Abbreviation Code<br>
+; CHECK: // .b8 5 // DW_TAG_formal_parameter<br>
+; CHECK: // .b8 0 // DW_CHILDREN_no<br>
+; CHECK: // .b8 3 // DW_AT_name<br>
+; CHECK: // .b8 8 // DW_FORM_string<br>
+; CHECK: // .b8 58 // DW_AT_decl_file<br>
+; CHECK: // .b8 11 // DW_FORM_data1<br>
+; CHECK: // .b8 59 // DW_AT_decl_line<br>
+; CHECK: // .b8 11 // DW_FORM_data1<br>
+; CHECK: // .b8 73 // DW_AT_type<br>
+; CHECK: // .b8 19 // DW_FORM_ref4<br>
+; CHECK: // .b8 0 // EOM(1)<br>
+; CHECK: // .b8 0 // EOM(2)<br>
+; CHECK: // .b8 4 // Abbreviation Code<br>
+; CHECK: // .b8 36 // DW_TAG_base_type<br>
+; CHECK: // .b8 0 // DW_CHILDREN_no<br>
+; CHECK: // .b8 3 // DW_AT_name<br>
+; CHECK: // .b8 8 // DW_FORM_string<br>
+; CHECK: // .b8 62 // DW_AT_encoding<br>
+; CHECK: // .b8 11 // DW_FORM_data1<br>
+; CHECK: // .b8 11 // DW_AT_byte_size<br>
+; CHECK: // .b8 11 // DW_FORM_data1<br>
+; CHECK: // .b8 0 // EOM(1)<br>
+; CHECK: // .b8 0 // EOM(2)<br>
+; CHECK: // .b8 0 // EOM(3)<br>
+; CHECK: // }<br>
+; CHECK: // .section .debug_info<br>
+; CHECK: // {<br>
+; CHECK: // .b32 179 // Length of Unit<br>
+; CHECK: // .b8 2 // DWARF version number<br>
+; CHECK: // .b8 0<br>
+; CHECK: // .b32 .debug_abbrev // Offset Into Abbrev. Section<br>
+; CHECK: // .b8 8 // Address Size (in bytes)<br>
+; CHECK: // .b8 1 // Abbrev [1] 0xb:0xac DW_TAG_compile_unit<br>
+; CHECK: // .b8 99 // DW_AT_producer<br>
+; CHECK: // .b8 108<br>
+; CHECK: // .b8 97<br>
+; CHECK: // .b8 110<br>
+; CHECK: // .b8 103<br>
+; CHECK: // .b8 32<br>
+; CHECK: // .b8 118<br>
+; CHECK: // .b8 101<br>
+; CHECK: // .b8 114<br>
+; CHECK: // .b8 115<br>
+; CHECK: // .b8 105<br>
+; CHECK: // .b8 111<br>
+; CHECK: // .b8 110<br>
+; CHECK: // .b8 32<br>
+; CHECK: // .b8 51<br>
+; CHECK: // .b8 46<br>
+; CHECK: // .b8 53<br>
+; CHECK: // .b8 46<br>
+; CHECK: // .b8 48<br>
+; CHECK: // .b8 32<br>
+; CHECK: // .b8 40<br>
+; CHECK: // .b8 116<br>
+; CHECK: // .b8 114<br>
+; CHECK: // .b8 117<br>
+; CHECK: // .b8 110<br>
+; CHECK: // .b8 107<br>
+; CHECK: // .b8 32<br>
+; CHECK: // .b8 50<br>
+; CHECK: // .b8 48<br>
+; CHECK: // .b8 52<br>
+; CHECK: // .b8 49<br>
+; CHECK: // .b8 54<br>
+; CHECK: // .b8 52<br>
+; CHECK: // .b8 41<br>
+; CHECK: // .b8 32<br>
+; CHECK: // .b8 40<br>
+; CHECK: // .b8 108<br>
+; CHECK: // .b8 108<br>
+; CHECK: // .b8 118<br>
+; CHECK: // .b8 109<br>
+; CHECK: // .b8 47<br>
+; CHECK: // .b8 116<br>
+; CHECK: // .b8 114<br>
+; CHECK: // .b8 117<br>
+; CHECK: // .b8 110<br>
+; CHECK: // .b8 107<br>
+; CHECK: // .b8 32<br>
+; CHECK: // .b8 50<br>
+; CHECK: // .b8 48<br>
+; CHECK: // .b8 52<br>
+; CHECK: // .b8 49<br>
+; CHECK: // .b8 56<br>
+; CHECK: // .b8 51<br>
+; CHECK: // .b8 41<br>
+; CHECK: // .b8 0<br>
+; CHECK: // .b8 12 // DW_AT_language<br>
+; CHECK: // .b8 0<br>
+; CHECK: // .b8 98 // DW_AT_name<br>
+; CHECK: // .b8 46<br>
+; CHECK: // .b8 99<br>
+; CHECK: // .b8 0<br>
+; CHECK: // .b32 .debug_line // DW_AT_stmt_list<br>
+; CHECK: // .b8 47 // DW_AT_comp_dir<br>
+; CHECK: // .b8 115<br>
+; CHECK: // .b8 111<br>
+; CHECK: // .b8 117<br>
+; CHECK: // .b8 114<br>
+; CHECK: // .b8 99<br>
+; CHECK: // .b8 101<br>
+; CHECK: // .b8 0<br>
+; CHECK: // .b64 Lfunc_begin0 // DW_AT_low_pc<br>
+; CHECK: // .b64 Lfunc_end2 // DW_AT_high_pc<br>
+; CHECK: // .b8 2 // Abbrev [2] 0x65:0x25 DW_TAG_subprogram<br>
+; CHECK: // .b64 Lfunc_begin0 // DW_AT_low_pc<br>
+; CHECK: // .b64 Lfunc_end0 // DW_AT_high_pc<br>
+; CHECK: // .b8 98 // DW_AT_name<br>
+; CHECK: // .b8 0<br>
+; CHECK: // .b8 1 // DW_AT_decl_file<br>
+; CHECK: // .b8 1 // DW_AT_decl_line<br>
+; CHECK: // .b8 1 // DW_AT_prototyped<br>
+; CHECK: // .b32 175 // DW_AT_type<br>
+; CHECK: // .b8 1 // DW_AT_external<br>
+; CHECK: // .b8 3 // Abbrev [3] 0x80:0x9 DW_TAG_formal_parameter<br>
+; CHECK: // .b8 99 // DW_AT_name<br>
+; CHECK: // .b8 0<br>
+; CHECK: // .b8 1 // DW_AT_decl_file<br>
+; CHECK: // .b8 1 // DW_AT_decl_line<br>
+; CHECK: // .b32 175 // DW_AT_type<br>
+; CHECK: // .b8 0 // End Of Children Mark<br>
+; CHECK: // .b8 2 // Abbrev [2] 0x8a:0x25 DW_TAG_subprogram<br>
+; CHECK: // .b64 Lfunc_begin2 // DW_AT_low_pc<br>
+; CHECK: // .b64 Lfunc_end2 // DW_AT_high_pc<br>
+; CHECK: // .b8 100 // DW_AT_name<br>
+; CHECK: // .b8 0<br>
+; CHECK: // .b8 1 // DW_AT_decl_file<br>
+; CHECK: // .b8 3 // DW_AT_decl_line<br>
+; CHECK: // .b8 1 // DW_AT_prototyped<br>
+; CHECK: // .b32 175 // DW_AT_type<br>
+; CHECK: // .b8 1 // DW_AT_external<br>
+; CHECK: // .b8 3 // Abbrev [3] 0xa5:0x9 DW_TAG_formal_parameter<br>
+; CHECK: // .b8 101 // DW_AT_name<br>
+; CHECK: // .b8 0<br>
+; CHECK: // .b8 1 // DW_AT_decl_file<br>
+; CHECK: // .b8 3 // DW_AT_decl_line<br>
+; CHECK: // .b32 175 // DW_AT_type<br>
+; CHECK: // .b8 0 // End Of Children Mark<br>
+; CHECK: // .b8 4 // Abbrev [4] 0xaf:0x7 DW_TAG_base_type<br>
+; CHECK: // .b8 105 // DW_AT_name<br>
+; CHECK: // .b8 110<br>
+; CHECK: // .b8 116<br>
+; CHECK: // .b8 0<br>
+; CHECK: // .b8 5 // DW_AT_encoding<br>
+; CHECK: // .b8 4 // DW_AT_byte_size<br>
+; CHECK: // .b8 0 // End Of Children Mark<br>
+; CHECK: // }<br>
+; CHECK: // .section .debug_macinfo<br>
+; CHECK: // {<br>
+; CHECK: // .b8 0 // End Of Macro List Mark<br>
+; CHECK: // }<br>
+<br>
+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" }<br>
+attributes #1 = { nounwind readnone }<br>
+<br>
+!llvm.ident = !{!0, !0}<br>
+!<a href="http://llvm.dbg.cu" rel="noreferrer" target="_blank">llvm.dbg.cu</a> = !{!1}<br>
+!llvm.module.flags = !{!11, !12}<br>
+<br>
+!0 = !{!"clang version 3.5.0 (trunk 204164) (llvm/trunk 204183)"}<br>
+!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)<br>
+!2 = !DIFile(filename: "b.c", directory: "/source")<br>
+!3 = !{}<br>
+!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)<br>
+!6 = !DIFile(filename: "b.c", directory: "/source")<br>
+!7 = !DISubroutineType(types: !8)<br>
+!8 = !{!9, !9}<br>
+!9 = !DIBasicType(tag: DW_TAG_base_type, name: "int", size: 32, align: 32, encoding: DW_ATE_signed)<br>
+!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)<br>
+!11 = !{i32 2, !"Dwarf Version", i32 2}<br>
+!12 = !{i32 1, !"Debug Info Version", i32 3}<br>
+!13 = !DILocalVariable(name: "c", line: 1, arg: 1, scope: !5, file: !6, type: !9)<br>
+!14 = !DILocation(line: 1, scope: !5)<br>
+!15 = !DILocalVariable(name: "e", line: 3, arg: 1, scope: !10, file: !6, type: !9)<br>
+!16 = !DILocation(line: 3, scope: !10)<br>
<br>
Modified: llvm/trunk/test/DebugInfo/NVPTX/debug-file-loc.ll<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/test/DebugInfo/NVPTX/debug-file-loc.ll?rev=330271&r1=330270&r2=330271&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/llvm/trunk/test/DebugInfo/NVPTX/debug-file-loc.ll?rev=330271&r1=330270&r2=330271&view=diff</a><br>
==============================================================================<br>
--- llvm/trunk/test/DebugInfo/NVPTX/debug-file-loc.ll (original)<br>
+++ llvm/trunk/test/DebugInfo/NVPTX/debug-file-loc.ll Wed Apr 18 09:13:41 2018<br>
@@ -8,13 +8,13 @@<br>
;__device__ void bar() {}<br>
;}<br>
<br>
-; CHECK: .file 1 "/source/dir{{.+}}<a href="http://bar.cu" rel="noreferrer" target="_blank">bar.cu</a>"<br>
-; CHECK: .file 2 "/source/dir{{.+}}foo.h"<br>
+; CHECK: .target sm_{{[0-9]+}}//, debug<br>
+<br>
; CHECK: .visible .func foo()<br>
-; CHECK: .loc 2 1 31<br>
+; CHECK: .loc [[FOO:[0-9]+]] 1 31<br>
; CHECK: ret;<br>
; CHECK: .visible .func bar()<br>
-; CHECK: .loc 1 2 31<br>
+; CHECK: .loc [[BAR:[0-9]+]] 2 31<br>
; CHECK: ret;<br>
<br>
define void @foo() !dbg !4 {<br>
@@ -27,6 +27,70 @@ bb:<br>
ret void, !dbg !11<br>
}<br>
<br>
+; CHECK-DAG: .file [[FOO]] "{{.*}}foo.h"<br>
+; CHECK-DAG: .file [[BAR]] "{{.*}}<a href="http://bar.cu" rel="noreferrer" target="_blank">bar.cu</a>"<br>
+; CHECK: // .section .debug_abbrev<br>
+; CHECK: // {<br>
+; CHECK: // .b8 1 // Abbreviation Code<br>
+; CHECK: // .b8 17 // DW_TAG_compile_unit<br>
+; CHECK: // .b8 0 // DW_CHILDREN_no<br>
+; CHECK: // .b8 37 // DW_AT_producer<br>
+; CHECK: // .b8 8 // DW_FORM_string<br>
+; CHECK: // .b8 19 // DW_AT_language<br>
+; CHECK: // .b8 5 // DW_FORM_data2<br>
+; CHECK: // .b8 3 // DW_AT_name<br>
+; CHECK: // .b8 8 // DW_FORM_string<br>
+; CHECK: // .b8 16 // DW_AT_stmt_list<br>
+; CHECK: // .b8 6 // DW_FORM_data4<br>
+; CHECK: // .b8 27 // DW_AT_comp_dir<br>
+; CHECK: // .b8 8 // DW_FORM_string<br>
+; CHECK: // .b8 17 // DW_AT_low_pc<br>
+; CHECK: // .b8 1 // DW_FORM_addr<br>
+; CHECK: // .b8 18 // DW_AT_high_pc<br>
+; CHECK: // .b8 1 // DW_FORM_addr<br>
+; CHECK: // .b8 0 // EOM(1)<br>
+; CHECK: // .b8 0 // EOM(2)<br>
+; CHECK: // .b8 0 // EOM(3)<br>
+; CHECK: // }<br>
+; CHECK: // .section .debug_info<br>
+; CHECK: // {<br>
+; CHECK: // .b32 50 // Length of Unit<br>
+; CHECK: // .b8 2 // DWARF version number<br>
+; CHECK: // .b8 0<br>
+; CHECK: // .b32 .debug_abbrev // Offset Into Abbrev. Section<br>
+; CHECK: // .b8 8 // Address Size (in bytes)<br>
+; CHECK: // .b8 1 // Abbrev [1] 0xb:0x2b DW_TAG_compile_unit<br>
+; CHECK: // .b8 0 // DW_AT_producer<br>
+; CHECK: // .b8 4 // DW_AT_language<br>
+; CHECK: // .b8 0<br>
+; CHECK: // .b8 98 // DW_AT_name<br>
+; CHECK: // .b8 97<br>
+; CHECK: // .b8 114<br>
+; CHECK: // .b8 46<br>
+; CHECK: // .b8 99<br>
+; CHECK: // .b8 117<br>
+; CHECK: // .b8 0<br>
+; CHECK: // .b32 .debug_line // DW_AT_stmt_list<br>
+; CHECK: // .b8 47 // DW_AT_comp_dir<br>
+; CHECK: // .b8 115<br>
+; CHECK: // .b8 111<br>
+; CHECK: // .b8 117<br>
+; CHECK: // .b8 114<br>
+; CHECK: // .b8 99<br>
+; CHECK: // .b8 101<br>
+; CHECK: // .b8 47<br>
+; CHECK: // .b8 100<br>
+; CHECK: // .b8 105<br>
+; CHECK: // .b8 114<br>
+; CHECK: // .b8 0<br>
+; CHECK: // .b64 Lfunc_begin0 // DW_AT_low_pc<br>
+; CHECK: // .b64 Lfunc_end1 // DW_AT_high_pc<br>
+; CHECK: // }<br>
+; CHECK: // .section .debug_macinfo<br>
+; CHECK: // {<br>
+; CHECK: // .b8 0 // End Of Macro List Mark<br>
+; CHECK: // }<br>
+<br>
!<a href="http://llvm.dbg.cu" rel="noreferrer" target="_blank">llvm.dbg.cu</a> = !{!0}<br>
!llvm.module.flags = !{!8, !9}<br>
<br>
@@ -37,7 +101,7 @@ bb:<br>
!5 = !DIFile(filename: "foo.h", directory: "/source/dir")<br>
!6 = !DISubroutineType(types: !2)<br>
!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)<br>
-!8 = !{i32 2, !"Dwarf Version", i32 4}<br>
+!8 = !{i32 2, !"Dwarf Version", i32 2}<br>
!9 = !{i32 2, !"Debug Info Version", i32 3}<br>
!10 = !DILocation(line: 1, column: 31, scope: !4)<br>
!11 = !DILocation(line: 2, column: 31, scope: !7)<br>
<br>
Modified: llvm/trunk/test/DebugInfo/NVPTX/debug-info.ll<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/test/DebugInfo/NVPTX/debug-info.ll?rev=330271&r1=330270&r2=330271&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-project/llvm/trunk/test/DebugInfo/NVPTX/debug-info.ll?rev=330271&r1=330270&r2=330271&view=diff</a><br>
==============================================================================<br>
--- llvm/trunk/test/DebugInfo/NVPTX/debug-info.ll (original)<br>
+++ llvm/trunk/test/DebugInfo/NVPTX/debug-info.ll Wed Apr 18 09:13:41 2018<br>
@@ -9,17 +9,8 @@<br>
; res(a * x[i], y[i], &y[i]);<br>
;}<br>
<br>
-; CHECK: .file 1 "/some/directory{{.+}}<a href="http://debug-info.cu" rel="noreferrer" target="_blank">debug-info.cu</a>"<br>
-; CHECK: .file 2 "/some/directory{{.+}}clang/include{{.+}}__clang_cuda_math_forward_declares.h"<br>
-; CHECK: .file 3 "{{.*}}/usr/include{{.+}}mathcalls.h"<br>
-; CHECK: .file 4 "{{.*}}/usr/include{{.+}}stdlib.h"<br>
-; CHECK: .file 5 "{{.*}}/usr/include{{.+}}stdlib-float.h"<br>
-; CHECK: .file 6 "{{.*}}/usr/include{{.+}}stdlib-bsearch.h"<br>
-; CHECK: .file 7 "{{.*}}/usr/lib/gcc/4.8/../../../../include/c++/4.8{{.+}}cstdlib"<br>
-; CHECK: .file 8 "{{.*}}/usr/local/cuda/include{{.+}}math_functions.hpp"<br>
-; CHECK: .file 9 "{{.*}}/usr/local/cuda/include{{.+}}device_functions.hpp"<br>
-; CHECK: .file 10 "/some/directory{{.+}}clang/include{{.+}}__clang_cuda_builtin_vars.h"<br>
-; CHECK: .file 11 "{{.*}}/usr/local/cuda/include{{.+}}vector_types.h"<br>
+; CHECK: .target sm_{{[0-9]+}}//, debug<br>
+<br>
; CHECK: .visible .entry _Z5saxpyifPfS_(<br>
; CHECK: .param .u32 {{.+}},<br>
; CHECK: .param .f32 {{.+}},<br>
@@ -31,18 +22,19 @@<br>
; CHECK: .reg .f32 %f<5>;<br>
; CHECK: .reg .b32 %r<6>;<br>
; CHECK: .reg .b64 %rd<8>;<br>
+; CHECK: .loc [[DEBUG_INFO_CU:[0-9]+]] 5 0<br>
; CHECK: ld.param.u32 %r{{.+}}, [{{.+}}];<br>
-; CHECK: .loc 10 78 180<br>
+; CHECK: .loc [[BUILTUIN_VARS_H:[0-9]+]] 78 180<br>
; CHECK: mov.u32 %r{{.+}}, %ctaid.x;<br>
-; CHECK: .loc 10 89 180<br>
+; CHECK: .loc [[BUILTUIN_VARS_H]] 89 180<br>
; CHECK: mov.u32 %r{{.+}}, %ntid.x;<br>
-; CHECK: .loc 10 67 180<br>
+; CHECK: .loc [[BUILTUIN_VARS_H]] 67 180<br>
; CHECK: mov.u32 %r{{.+}}, %tid.x;<br>
-; CHECK: .loc 1 6 35<br>
+; CHECK: .loc [[DEBUG_INFO_CU]] 6 35<br>
; CHECK: mad.lo.s32 %r{{.+}}, %r{{.+}}, %r{{.+}}, %r{{.+}};<br>
-; CHECK: .loc 1 7 9<br>
+; CHECK: .loc [[DEBUG_INFO_CU]] 7 9<br>
; CHECK: setp.ge.s32 %p{{.+}}, %r{{.+}}, %r{{.+}};<br>
-; CHECK: .loc 1 7 7<br>
+; CHECK: .loc [[DEBUG_INFO_CU]] 7 7<br>
; CHECK: @%p{{.+}} bra [[BB:.+]];<br>
; CHECK: ld.param.f32 %f{{.+}}, [{{.+}}];<br>
; CHECK: ld.param.u64 %rd{{.+}}, [{{.+}}];<br>
@@ -51,17 +43,17 @@<br>
; CHECK: cvta.to.global.u64 %rd{{.+}}, %rd{{.+}};<br>
; CHECK: mul.wide.u32 %rd{{.+}}, %r{{.+}}, 4;<br>
; CHECK: add.s64 %rd{{.+}}, %rd{{.+}}, %rd{{.+}};<br>
-; CHECK: .loc 1 8 13<br>
+; CHECK: .loc [[DEBUG_INFO_CU]] 8 13<br>
; CHECK: ld.global.f32 %f{{.+}}, [%rd{{.+}}];<br>
; CHECK: add.s64 %rd{{.+}}, %rd{{.+}}, %rd{{.+}};<br>
-; CHECK: .loc 1 8 19<br>
+; CHECK: .loc [[DEBUG_INFO_CU]] 8 19<br>
; CHECK: ld.global.f32 %f{{.+}}, [%rd{{.+}}];<br>
-; CHECK: .loc 1 3 82<br>
+; CHECK: .loc [[DEBUG_INFO_CU]] 3 82<br>
; CHECK: fma.rn.f32 %f{{.+}}, %f{{.+}}, %f{{.+}}, %f{{.+}};<br>
-; CHECK: .loc 1 3 78<br>
+; CHECK: .loc [[DEBUG_INFO_CU]] 3 78<br>
; CHECK: st.global.f32 [%rd{{.+}}], %f{{.+}};<br>
; CHECK: [[BB]]:<br>
-; CHECK: .loc 1 9 1<br>
+; CHECK: .loc [[DEBUG_INFO_CU]] 9 1<br>
; CHECK: ret;<br>
; CHECK: }<br>
<br>
@@ -99,6 +91,8307 @@ if.end:<br>
ret void, !dbg !718<br>
}<br>
<br>
+; CHECK-DAG: .file {{[0-9]+}} "{{.*}}clang/include{{/|\\\\}}__clang_cuda_math_forward_declares.h"<br>
+; CHECK-DAG: .file {{[0-9]+}} "{{.*}}/usr/include{{/|\\\\}}mathcalls.h"<br>
+; CHECK-DAG: .file {{[0-9]+}} "{{.*}}/usr/lib/gcc/4.8/../../../../include/c++/4.8{{/|\\\\}}cmath"<br>
+; CHECK-DAG: .file {{[0-9]+}} "{{.*}}/usr/include{{/|\\\\}}stdlib.h"<br>
+; CHECK-DAG: .file {{[0-9]+}} "{{.*}}/usr/lib/gcc/4.8/../../../../include/c++/4.8{{/|\\\\}}cstdlib"<br>
+; CHECK-DAG: .file {{[0-9]+}} "{{.*}}/usr/include{{/|\\\\}}stdlib-float.h"<br>
+; CHECK-DAG: .file {{[0-9]+}} "{{.*}}/usr/include{{/|\\\\}}stdlib-bsearch.h"<br>
+; CHECK-DAG: .file {{[0-9]+}} "{{.*}}clang/include{{/|\\\\}}stddef.h"<br>
+; CHECK-DAG: .file {{[0-9]+}} "{{.*}}/usr/local/cuda/include{{/|\\\\}}math_functions.hpp"<br>
+; CHECK_DAG: .file {{[0-9]+}} "{{.*}}clang/include{{/|\\\\}}__clang_cuda_cmath.h"<br>
+; CHECK-DAG: .file {{[0-9]+}} "{{.*}}/usr/local/cuda/include{{/|\\\\}}device_functions.hpp"<br>
+; CHECK-DAG: .file [[DEBUG_INFO_CU]] "{{.*}}<a href="http://debug-info.cu" rel="noreferrer" target="_blank">debug-info.cu</a>"<br>
+; CHECK-DAG: .file [[BUILTUIN_VARS_H]] "{{.*}}clang/include{{/|\\\\}}__clang_cuda_builtin_vars.h"<br>
+<br>
+; CHECK: // .section .debug_abbrev<br>
+; CHECK: // {<br>
+; CHECK: // .b8 1 // Abbreviation Code<br>
+; CHECK: // .b8 17 // DW_TAG_compile_unit<br>
+; CHECK: // .b8 1 // DW_CHILDREN_yes<br>
+; CHECK: // .b8 37 // DW_AT_producer<br>
+; CHECK: // .b8 8 // DW_FORM_string<br>
+; CHECK: // .b8 19 // DW_AT_language<br>
+; CHECK: // .b8 5 // DW_FORM_data2<br>
+; CHECK: // .b8 3 // DW_AT_name<br>
+; CHECK: // .b8 8 // DW_FORM_string<br>
+; CHECK: // .b8 16 // DW_AT_stmt_list<br>
+; CHECK: // .b8 6 // DW_FORM_data4<br>
+; CHECK: // .b8 27 // DW_AT_comp_dir<br>
+; CHECK: // .b8 8 // DW_FORM_string<br>
+; CHECK: // .b8 17 // DW_AT_low_pc<br>
+; CHECK: // .b8 1 // DW_FORM_addr<br>
+; CHECK: // .b8 18 // DW_AT_high_pc<br>
+; CHECK: // .b8 1 // DW_FORM_addr<br>
+; CHECK: // .b8 0 // EOM(1)<br>
+; CHECK: // .b8 0 // EOM(2)<br>
+; CHECK: // .b8 2 // Abbreviation Code<br>
+; CHECK: // .b8 57 // DW_TAG_namespace<br>
+; CHECK: // .b8 1 // DW_CHILDREN_yes<br>
+; CHECK: // .b8 3 // DW_AT_name<br>
+; CHECK: // .b8 8 // DW_FORM_string<br>
+; CHECK: // .b8 0 // EOM(1)<br>
+; CHECK: // .b8 0 // EOM(2)<br>
+; CHECK: // .b8 3 // Abbreviation Code<br>
+; CHECK: // .b8 8 // DW_TAG_imported_declaration<br>
+; CHECK: // .b8 0 // DW_CHILDREN_no<br>
+; CHECK: // .b8 58 // DW_AT_decl_file<br>
+; CHECK: // .b8 11 // DW_FORM_data1<br>
+; CHECK: // .b8 59 // DW_AT_decl_line<br>
+; CHECK: // .b8 11 // DW_FORM_data1<br>
+; CHECK: // .b8 24 // DW_AT_import<br>
+; CHECK: // .b8 19 // DW_FORM_ref4<br>
+; CHECK: // .b8 0 // EOM(1)<br>
+; CHECK: // .b8 0 // EOM(2)<br>
+; CHECK: // .b8 4 // Abbreviation Code<br>
+; CHECK: // .b8 8 // DW_TAG_imported_declaration<br>
+; CHECK: // .b8 0 // DW_CHILDREN_no<br>
+; CHECK: // .b8 58 // DW_AT_decl_file<br>
+; CHECK: // .b8 11 // DW_FORM_data1<br>
+; CHECK: // .b8 59 // DW_AT_decl_line<br>
+; CHECK: // .b8 5 // DW_FORM_data2<br>
+; CHECK: // .b8 24 // DW_AT_import<br>
+; CHECK: // .b8 19 // DW_FORM_ref4<br>
+; CHECK: // .b8 0 // EOM(1)<br>
+; CHECK: // .b8 0 // EOM(2)<br>
+; CHECK: // .b8 5 // Abbreviation Code<br>
+; CHECK: // .b8 46 // DW_TAG_subprogram<br>
+; CHECK: // .b8 1 // DW_CHILDREN_yes<br>
+; CHECK: // .b8 135 // DW_AT_MIPS_linkage_name<br>
+; CHECK: // .b8 64<br>
+; CHECK: // .b8 8 // DW_FORM_string<br>
+; CHECK: // .b8 3 // DW_AT_name<br>
+; CHECK: // .b8 8 // DW_FORM_string<br>
+; CHECK: // .b8 58 // DW_AT_decl_file<br>
+; CHECK: // .b8 11 // DW_FORM_data1<br>
+; CHECK: // .b8 59 // DW_AT_decl_line<br>
+; CHECK: // .b8 11 // DW_FORM_data1<br>
+; CHECK: // .b8 73 // DW_AT_type<br>
+; CHECK: // .b8 19 // DW_FORM_ref4<br>
+; CHECK: // .b8 60 // DW_AT_declaration<br>
+; CHECK: // .b8 12 // DW_FORM_flag<br>
+; CHECK: // .b8 0 // EOM(1)<br>
+; CHECK: // .b8 0 // EOM(2)<br>
+; CHECK: // .b8 6 // Abbreviation Code<br>
+; CHECK: // .b8 5 // DW_TAG_formal_parameter<br>
+; CHECK: // .b8 0 // DW_CHILDREN_no<br>
+; CHECK: // .b8 73 // DW_AT_type<br>
+; CHECK: // .b8 19 // DW_FORM_ref4<br>
+; CHECK: // .b8 0 // EOM(1)<br>
+; CHECK: // .b8 0 // EOM(2)<br>
+; CHECK: // .b8 7 // Abbreviation Code<br>
+; CHECK: // .b8 36 // DW_TAG_base_type<br>
+; CHECK: // .b8 0 // DW_CHILDREN_no<br>
+; CHECK: // .b8 3 // DW_AT_name<br>
+; CHECK: // .b8 8 // DW_FORM_string<br>
+; CHECK: // .b8 62 // DW_AT_encoding<br>
+; CHECK: // .b8 11 // DW_FORM_data1<br>
+; CHECK: // .b8 11 // DW_AT_byte_size<br>
+; CHECK: // .b8 11 // DW_FORM_data1<br>
+; CHECK: // .b8 0 // EOM(1)<br>
+; CHECK: // .b8 0 // EOM(2)<br>
+; CHECK: // .b8 8 // Abbreviation Code<br>
+; CHECK: // .b8 15 // DW_TAG_pointer_type<br>
+; CHECK: // .b8 0 // DW_CHILDREN_no<br>
+; CHECK: // .b8 73 // DW_AT_type<br>
+; CHECK: // .b8 19 // DW_FORM_ref4<br>
+; CHECK: // .b8 0 // EOM(1)<br>
+; CHECK: // .b8 0 //</blockquote></div>