[DEBUGINFO, NVPTX] Allow to disable debug info from command line.

[ABataev]

Summary

For better support of the CUDA profiling tools with the optimized code
we should be able to emit .target directive without debug option.
This is required as .line and .file directives must be emitted, but
the DWARF sections should be ignored. ptxas does not support compilation
of the modules with the enabled debug info when optimization level is

0, but allows to emit linenumbers for the profilers.

-no-cuda-debug option disables emission of the debug option, but not
the emission of the debug info itself.

[aprantl]

Why can't you use the kind=NoDebug flag on DICompileUnit like the other profiling modes do?

[ABataev]

In D45782#1071334, @aprantl wrote:

Why can't you use the kind=NoDebug flag on DICompileUnit like the other profiling modes do?

We need to emit .loc and .file directives, NoDebug disables debug info completely.

[aprantl]

What's the difference between that and -gline-tables-only?

[ABataev]

In D45782#1071339, @aprantl wrote:

What's the difference between that and -gline-tables-only?

-gline-table-only still emits .debug_info DWARF sections and because of that debug option is emitted. But it is not allowed to have debug options on .target directive when you need to optimize the code. Without debug option DWARF sections are just ignored.

[tra]

For better support of the CUDA profiling tools with the optimized code
we should be able to emit .target directive without debug option.
This is required as .line and .file directives must be emitted, but
the DWARF sections should be ignored. ptxas does not support compilation
of the modules with the enabled debug info when optimization level is

I'm not sure I understand how removing ", debug" helps profiling tools. It does allow compiling PTX with -O1+ even if dwarf bits are present and that is useful, but I'm not sure how it's connected to profiling. Are you saying that nvprof (or something else) may look at the dwarf bits in the PTX?

If dwarf info is not going to make it to SASS, I think we'd be better off not generating dwarf at all.
What do I miss?

[ABataev]

In D45782#1072425, @tra wrote:

For better support of the CUDA profiling tools with the optimized code
we should be able to emit .target directive without debug option.
This is required as .line and .file directives must be emitted, but
the DWARF sections should be ignored. ptxas does not support compilation
of the modules with the enabled debug info when optimization level is

I'm not sure I understand how removing ", debug" helps profiling tools. It does allow compiling PTX with -O1+ even if dwarf bits are present and that is useful, but I'm not sure how it's connected to profiling. Are you saying that nvprof (or something else) may look at the dwarf bits in the PTX?

If dwarf info is not going to make it to SASS, I think we'd be better off not generating dwarf at all.
What do I miss?

, debug should be removed to enable opt level > 0. I checked the ptxas, without , debug it just ignores DWARF sections, but with -lineinfo option it may generate correct data for the nvprof from .loc and .file directives.

[tra]

, debug should be removed to enable opt level > 0. I checked the ptxas, without , debug it just ignores DWARF sections, but with -lineinfo option it may generate correct data for the nvprof from .loc and .file directives.

Are you saying that enabling lineinfo debug also produces DWARF sections now and that's why we need the flag to ignore them if optimizations are enabled? Previously there was no need because we were emitting only .loc and .file.

[ABataev]

In D45782#1072437, @tra wrote:

, debug should be removed to enable opt level > 0. I checked the ptxas, without , debug it just ignores DWARF sections, but with -lineinfo option it may generate correct data for the nvprof from .loc and .file directives.

Are you saying that enabling lineinfo debug also produces DWARF sections now and that's why we need the flag to ignore them if optimizations are enabled? Previously there was no need because we were emitting only .loc and .file.

Yes, currently it is impossible to emit .loc and .file directives without the emission of the DWARF sections as we're using general DebugInfo codebase. Later we could try to disable emission of the DWARF sections if it is not required.

[tra]

In D45782#1072442, @ABataev wrote:

Yes, currently it is impossible to emit .loc and .file directives without the emission of the DWARF sections as we're using general DebugInfo codebase. Later we could try to disable emission of the DWARF sections if it is not required.

Got it. This should be included in the patch description. LGTM otherwise.

[dblaikie]

Sounds like sort of "-gno-really-line-tables-only" (I kid, somewhat)

What's the end goal for this functionality? Does it make sense to plumb it
through as another level alongside GMLT/limited/full?

[ABataev]

Need to rework the solution after discussion with David Blaike and Paul Robinson.

lib/Target/NVPTX/NVPTXAsmPrinter.cpp

 #include <string>
 #include <utility>
 #include <vector>
 
 using namespace llvm;
 
 #define DEPOTNAME "__local_depot"
 
+static cl::opt<bool>
+    NoCudaDebug("no-cuda-debug",
+                cl::desc("Do not mark ptx file as having debug info"),
+                cl::init(false));
+
 /// DiscoverDependentGlobals - Return a set of GlobalVariables on which \p V
 /// depends.
 static void
 DiscoverDependentGlobals(const Value *V,
                          DenseSet<const GlobalVariable *> &Globals) {
   if (const GlobalVariable *GV = dyn_cast<GlobalVariable>(V))
     Globals.insert(GV);
   else {
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";
 
   // FIXME: remove comment once debug info is properly supported.
-  if (MMI && MMI->hasDebugInfo())
+  if (MMI && MMI->hasDebugInfo() && !NoCudaDebug)
     O << "//, debug";
 
   O << "\n";
 
   O << ".address_size ";
   if (NTM.is64Bit())
     O << "64";
   else

test/DebugInfo/NVPTX/debug-info.ll

-; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s
+; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s --check-prefixes=CHECK,DEBUG
+; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -no-cuda-debug | FileCheck %s --check-prefixes=CHECK,NODEBUG
 
 ; // Bitcode int this test case is reduced version of compiled code below:
 ;__device__ inline void res(float x, float y, float *res) { *res = x + y; }
 ;
 ;__global__ void saxpy(int n, float a, float *x, float *y) {
 ;  int i = blockIdx.x * blockDim.x + threadIdx.x;
 ;  if (i < n)
 ;    res(a * x[i], y[i], &y[i]);
 ;}
 
-; CHECK: .target sm_{{[0-9]+}}//, debug
+; DEBUG: .target sm_{{[0-9]+}}//, debug
+; NODEBUG: .target sm_{{[0-9]+$}}
 
 ; CHECK: .visible .entry _Z5saxpyifPfS_(
 ; CHECK: .param .u32 {{.+}},
 ; CHECK: .param .f32 {{.+}},
 ; CHECK: .param .u64 {{.+}},
 ; CHECK: .param .u64 {{.+}}
 ; CHECK: )
 ; CHECK: {