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

[Alexey Bataev via Phabricator via llvm-commits]

ABataev created this revision.
ABataev added reviewers: tra, jlebar, echristo.
Herald added subscribers: JDevlieghere, aprantl, jholewinski.

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.


Repository:
  rL LLVM

https://reviews.llvm.org/D45782

Files:
  lib/Target/NVPTX/NVPTXAsmPrinter.cpp
  test/DebugInfo/NVPTX/debug-info.ll


Index: test/DebugInfo/NVPTX/debug-info.ll
===================================================================
--- test/DebugInfo/NVPTX/debug-info.ll
+++ test/DebugInfo/NVPTX/debug-info.ll
@@ -1,4 +1,5 @@
-; 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; }
@@ -9,7 +10,8 @@
 ;    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 {{.+}},
Index: lib/Target/NVPTX/NVPTXAsmPrinter.cpp
===================================================================
--- lib/Target/NVPTX/NVPTXAsmPrinter.cpp
+++ lib/Target/NVPTX/NVPTXAsmPrinter.cpp
@@ -93,6 +93,11 @@
 
 #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
@@ -876,7 +881,7 @@
     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";


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D45782.142975.patch
Type: text/x-patch
Size: 1636 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20180418/ab53b9eb/attachment.bin>