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

[Alexey Bataev via llvm-commits]

Author: abataev
Date: Thu Dec  6 08:25:35 2018
New Revision: 348497

URL: http://llvm.org/viewvc/llvm-project?rev=348497&view=rev
Log:
[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

Modified:
    llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
    llvm/trunk/test/DebugInfo/NVPTX/debug-file-loc-only.ll

Modified: llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp?rev=348497&r1=348496&r2=348497&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp (original)
+++ llvm/trunk/lib/Target/NVPTX/NVPTXAsmPrinter.cpp Thu Dec  6 08:25:35 2018
@@ -881,8 +881,22 @@ void NVPTXAsmPrinter::emitHeader(Module
   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";

Modified: llvm/trunk/test/DebugInfo/NVPTX/debug-file-loc-only.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/DebugInfo/NVPTX/debug-file-loc-only.ll?rev=348497&r1=348496&r2=348497&view=diff
==============================================================================
--- llvm/trunk/test/DebugInfo/NVPTX/debug-file-loc-only.ll (original)
+++ llvm/trunk/test/DebugInfo/NVPTX/debug-file-loc-only.ll Thu Dec  6 08:25:35 2018
@@ -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