[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/NVPTXAsmPrinter.cpp

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 in 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]+}}//, debug
+; 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;