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

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.

Reviewers: echristo

Subscribers: llvm-commits

Differential Revision: https://reviews.llvm.org/D46061

llvm-svn: 348497

llvm/lib/Target/NVPTX/NVPTXAsmPrinter.cpp

@@ -881,8 +881,22 @@ void NVPTXAsmPrinter::emitHeader(Module &M, raw_ostream &O,
   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";

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

@@ -8,7 +8,7 @@
 ;__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