[DEBUGINFO, NVPTX] Disable emission of ',debug' option if only debug directives are allowed.

[ABataev]

Summary

If the output of debug directives only is requested, we should drop
emission of ',debug' option from the target directive. Required for
supporting of nvprof profiler.

[ABataev]

Updated and reworked to latest changes

[ABataev]

Ping

[echristo]

Do you mean basically just line tables or something else? (Minus the need in normal dwarf for a minimal cu to associate the line table with). Also, how does this handle inlining etc?

[ABataev]

In D46061#1224055, @echristo wrote:

Do you mean basically just line tables or something else? (Minus the need in normal dwarf for a minimal cu to associate the line table with). Also, how does this handle inlining etc?

When we have full debug info, we need to add debug option to the target directive. When we do not emit debug info or emit just debug directives (i.e. only lineinfo), this options should not be added to the target directive. Plus we need to render all the debug directives (.line|.file) to be able to get the correct data from the profiler.
The inlining is handled only in terms of files/location debug directives, nothing else.

[ABataev]

Ping!

[ABataev]

Removed an option, used emission kind for the check of the full debug info.

[probinson]

LGTM

test/DebugInfo/NVPTX/debug-file-loc-only.ll

3 ↗

(On Diff #172164)

Typo: "Bitcode in this test..." (not "int").

[dblaikie]

Looks like Eric still had some questions here? Or at least hadn't commented
about whether the answers/changes addressed his concerns?

[probinson]

In D46061#1285431, @dblaikie wrote:

Looks like Eric still had some questions here? Or at least hadn't commented
about whether the answers/changes addressed his concerns?

My bad. It looked answered to me, but yes should have waited on his confirmation.

[echristo]

Some inline comments - I'm still not happy with some of how this patch is and would like to see some changes and elaborations of how we split things out. Mostly it's bikeshed naming things, but the current state is a bit more confusing than without.

Thanks!

-eric

llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp

888–890

This seems to conflict with what you said above:

"When we have full debug info, we need to add debug option to the target directive. When we do not emit debug info or emit just debug directives (i.e. only lineinfo),"

as LineTablesOnly is line tables. If we need to split line tables into "things that handle inlining and things that don't" we should handle that discrepancy, but DebugDirectivesOnly isn't very helpful as a descriptor of what the compiler should be doing and paying attention to.

956–958

What's going on here?

[ABataev]

Update after revert.

[ABataev]

Removed extra functionality from the patch.

[ABataev]

Fixed test comment.

[ABataev]

llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp

888–890

Yes, because I reworked the patch for the frontend. LineTablesOnly is considered as normal debug info option now. Plus, even before, when we saw -line-table-only option for NVPTX target, we still emitted DICompileUnit::DebugDirectivesOnly debug info kind. So, this is ok even for the previous description.

956–958

I split this patch into 2 different patches.

llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.h

 class NVPTXTargetStreamer : public MCTargetStreamer {
 private:
   SmallVector<std::string, 4> DwarfFiles;
 
 public:
   NVPTXTargetStreamer(MCStreamer &S);
   ~NVPTXTargetStreamer() override;
 
+  /// Outputs the list of the DWARF '.file' directives to the streamer.
+  void outputDwarfFileDirectives();
+
   /// 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/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXTargetStreamer.cpp

 
 //
 // NVPTXTargetStreamer Implemenation
 //
 NVPTXTargetStreamer::NVPTXTargetStreamer(MCStreamer &S) : MCTargetStreamer(S) {}
 
 NVPTXTargetStreamer::~NVPTXTargetStreamer() = default;
 
+void NVPTXTargetStreamer::outputDwarfFileDirectives() {
+  for (const std::string &S : DwarfFiles)
+    getStreamer().EmitRawText(S.data());
+  DwarfFiles.clear();
+}
+
 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.
void NVPTXTargetStreamer::changeSection(const MCSection *CurSection,
   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();
+    outputDwarfFileDirectives();
     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";
   }
 }
 

llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp

 // of machine-dependent LLVM code to NVPTX assembly language.
 //
 //===----------------------------------------------------------------------===//
 
 #include "NVPTXAsmPrinter.h"
 #include "InstPrinter/NVPTXInstPrinter.h"
 #include "MCTargetDesc/NVPTXBaseInfo.h"
 #include "MCTargetDesc/NVPTXMCAsmInfo.h"
+#include "MCTargetDesc/NVPTXTargetStreamer.h"
 #include "NVPTX.h"
 #include "NVPTXMCExpr.h"
 #include "NVPTXMachineFunctionInfo.h"
 #include "NVPTXRegisterInfo.h"
 #include "NVPTXSubtarget.h"
 #include "NVPTXTargetMachine.h"
 #include "NVPTXUtilities.h"
 #include "cl_common_defines.h"
void NVPTXAsmPrinter::emitHeader(Module &M, raw_ostream &O,
 
   O << ".target ";
   O << STI.getTargetName();
 
   const NVPTXTargetMachine &NTM = static_cast<const NVPTXTargetMachine &>(TM);
   if (NTM.getDrvInterface() == NVPTX::NVCL)
     O << ", texmode_independent";
 
+  bool HasFullDebugInfo = false;
+  for (DICompileUnit *CU : M.debug_compile_units()) {
+    switch(CU->getEmissionKind()) {
+    case DICompileUnit::NoDebug:
+    case DICompileUnit::DebugDirectivesOnly:
+      break;
+    case DICompileUnit::LineTablesOnly:
+    case DICompileUnit::FullDebug:
+      HasFullDebugInfo = true;
+      break;
+    }
+    if (HasFullDebugInfo)
+      break;
+  }
   // FIXME: remove comment once debug info is properly supported.
-  if (MMI && MMI->hasDebugInfo())
+  if (MMI && MMI->hasDebugInfo() && HasFullDebugInfo)
     O << "//, debug";
 
   O << "\n";
 
   O << ".address_size ";
   if (NTM.is64Bit())
     O << "64";
   else
bool NVPTXAsmPrinter::doFinalization(Module &M) {
   clearAnnotationCache(&M);
 
   delete[] gv_array;
   // FIXME: remove comment once debug info is properly supported.
   // Close the last emitted section
   if (HasDebugInfo)
     OutStreamer->EmitRawText("//\t}");
 
+  // Output last DWARF .file directives, if any.
+  static_cast<NVPTXTargetStreamer *>(OutStreamer->getTargetStreamer())
+      ->outputDwarfFileDirectives();
+
   return ret;
 
   //bool Result = AsmPrinter::doFinalization(M);
   // Instead of calling the parents doFinalization, we may
   // clone parents doFinalization and customize here.
   // Currently, we if NVISA out the EmitGlobals() in
   // parent's doFinalization, which is too intrusive.
   //

llvm/trunk/test/DebugInfo/NVPTX/debug-file-loc-only.ll

+; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s
+
+; // Bitcode int this test case is reduced version of compiled code below:
+;extern "C" {
+;#line 1 "/source/dir/foo.h"
+;__device__ void foo() {}
+;#line 2 "/source/dir/bar.cu"
+;__device__ void bar() {}
+;}
+
+; CHECK: .target sm_{{[0-9]+$}}
+
+; CHECK: .visible .func foo()
+; CHECK: .loc [[FOO:[0-9]+]] 1 31
+; CHECK:  ret;
+; CHECK: .visible .func bar()
+; CHECK: .loc [[BAR:[0-9]+]] 2 31
+; CHECK:  ret;
+
+define void @foo() !dbg !4 {
+bb:
+  ret void, !dbg !10
+}
+
+define void @bar() !dbg !7 {
+bb:
+  ret void, !dbg !11
+}
+
+; CHECK-DAG: .file [[FOO]] "{{.*}}foo.h"
+; CHECK-DAG: .file [[BAR]] "{{.*}}bar.cu"
+
+; CHECK-NOT: .section .debug{{.*}}
+
+!llvm.dbg.cu = !{!0}
+!llvm.module.flags = !{!8, !9}
+
+!0 = distinct !DICompileUnit(language: DW_LANG_C_plus_plus, file: !1, producer: "", isOptimized: false, runtimeVersion: 0, emissionKind: DebugDirectivesOnly, enums: !2)
+!1 = !DIFile(filename: "bar.cu", directory: "/source/dir")
+!2 = !{}
+!4 = distinct !DISubprogram(name: "foo", scope: !5, file: !5, line: 1, type: !6, isLocal: false, isDefinition: true, scopeLine: 1, flags: DIFlagPrototyped, isOptimized: false, unit: !0, retainedNodes: !2)
+!5 = !DIFile(filename: "foo.h", directory: "/source/dir")
+!6 = !DISubroutineType(types: !2)
+!7 = distinct !DISubprogram(name: "bar", scope: !1, file: !1, line: 2, type: !6, isLocal: false, isDefinition: true, scopeLine: 2, flags: DIFlagPrototyped, isOptimized: false, unit: !0, retainedNodes: !2)
+!8 = !{i32 2, !"Dwarf Version", i32 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)