<html xmlns:v="urn:schemas-microsoft-com:vml" xmlns:o="urn:schemas-microsoft-com:office:office" xmlns:w="urn:schemas-microsoft-com:office:word" xmlns:m="http://schemas.microsoft.com/office/2004/12/omml" xmlns="http://www.w3.org/TR/REC-html40">
<head>
<meta http-equiv="Content-Type" content="text/html; charset=utf-8">
<meta name="Generator" content="Microsoft Word 14 (filtered medium)">
<style><!--
/* Font Definitions */
@font-face
{font-family:Calibri;
panose-1:2 15 5 2 2 2 4 3 2 4;}
@font-face
{font-family:Tahoma;
panose-1:2 11 6 4 3 5 4 4 2 4;}
@font-face
{font-family:Consolas;
panose-1:2 11 6 9 2 2 4 3 2 4;}
/* Style Definitions */
p.MsoNormal, li.MsoNormal, div.MsoNormal
{margin:0in;
margin-bottom:.0001pt;
font-size:12.0pt;
font-family:"Times New Roman","serif";}
a:link, span.MsoHyperlink
{mso-style-priority:99;
color:blue;
text-decoration:underline;}
a:visited, span.MsoHyperlinkFollowed
{mso-style-priority:99;
color:purple;
text-decoration:underline;}
pre
{mso-style-priority:99;
mso-style-link:"HTML Preformatted Char";
margin:0in;
margin-bottom:.0001pt;
font-size:10.0pt;
font-family:"Courier New";}
span.HTMLPreformattedChar
{mso-style-name:"HTML Preformatted Char";
mso-style-priority:99;
mso-style-link:"HTML Preformatted";
font-family:"Consolas","serif";}
span.EmailStyle19
{mso-style-type:personal-reply;
font-family:"Calibri","sans-serif";
color:#1F497D;}
.MsoChpDefault
{mso-style-type:export-only;
font-size:10.0pt;}
@page WordSection1
{size:8.5in 11.0in;
margin:1.0in 1.0in 1.0in 1.0in;}
div.WordSection1
{page:WordSection1;}
--></style><!--[if gte mso 9]><xml>
<o:shapedefaults v:ext="edit" spidmax="1026" />
</xml><![endif]--><!--[if gte mso 9]><xml>
<o:shapelayout v:ext="edit">
<o:idmap v:ext="edit" data="1" />
</o:shapelayout></xml><![endif]-->
</head>
<body lang="EN-US" link="blue" vlink="purple">
<div class="WordSection1">
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1F497D">I think David is suggesting that you add a new enumeration to DebugInfoKind, in clang/include/clang/Basic/DebugInfoOptions.h. This would go between LocTrackingOnly
(which remembers source locations but does not emit any debug sections) and DebugLineTablesOnly (which emits more than just the line tables, despite the name).<o:p></o:p></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1F497D">--paulr<o:p></o:p></span></p>
<p class="MsoNormal"><a name="_MailEndCompose"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1F497D"><o:p> </o:p></span></a></p>
<div style="border:none;border-left:solid blue 1.5pt;padding:0in 0in 0in 4.0pt">
<div>
<div style="border:none;border-top:solid #B5C4DF 1.0pt;padding:3.0pt 0in 0in 0in">
<p class="MsoNormal"><b><span style="font-size:10.0pt;font-family:"Tahoma","sans-serif"">From:</span></b><span style="font-size:10.0pt;font-family:"Tahoma","sans-serif""> llvm-commits [mailto:llvm-commits-bounces@lists.llvm.org]
<b>On Behalf Of </b>Alexey Bataev via llvm-commits<br>
<b>Sent:</b> Monday, April 23, 2018 2:19 PM<br>
<b>To:</b> David Blaikie<br>
<b>Cc:</b> reviews+d45782+public+181908537bbf2c34@reviews.llvm.org; jholewinski@nvidia.com; llvm-commits@lists.llvm.org; Alexey Bataev<br>
<b>Subject:</b> Re: [PATCH] D45782: [DEBUGINFO, NVPTX] Allow to disable debug info from command line.<o:p></o:p></span></p>
</div>
</div>
<p class="MsoNormal"><o:p> </o:p></p>
<pre>Cuda does not allow to compile code at optimization level >0 with debug info, but you can compile code at O1-3 with lineinfo only (i.e. when you have just `.loc` and `.file` directives in the PTX files).<o:p></o:p></pre>
<pre>The lineinfo is required for the nvprof, the CUDA profiler. So, our purpose is to emit these directives but without DWARF sections.<o:p></o:p></pre>
<pre>Actually, we can live with the DWARF sections in the code for now, the only thing we actually need is just to disable emission of the `,debug` option in the `.target` directive.<o:p></o:p></pre>
<pre>But it would be good to disable emission of the DWARF sections. It is required because CUDA tools include the PTX(assembler) file in the resulting binary, thus increasing its final size. <o:p></o:p></pre>
<pre>-------------<o:p></o:p></pre>
<pre>Best regards,<o:p></o:p></pre>
<pre>Alexey Bataev<o:p></o:p></pre>
<div>
<p class="MsoNormal">23.04.2018 14:02, David Blaikie пишет:<o:p></o:p></p>
</div>
<blockquote style="margin-top:5.0pt;margin-bottom:5.0pt">
<div>
<p class="MsoNormal" style="margin-bottom:12.0pt">Sounds like sort of "-gno-really-line-tables-only" (I kid, somewhat)<br>
<br>
What's the end goal for this functionality? Does it make sense to plumb it through as another level alongside GMLT/limited/full?<o:p></o:p></p>
<div>
<div>
<p class="MsoNormal">On Wed, Apr 18, 2018 at 11:43 AM Alexey Bataev via Phabricator via llvm-commits <<a href="mailto:llvm-commits@lists.llvm.org">llvm-commits@lists.llvm.org</a>> wrote:<o:p></o:p></p>
</div>
<blockquote style="border:none;border-left:solid #CCCCCC 1.0pt;padding:0in 0in 0in 6.0pt;margin-left:4.8pt;margin-right:0in">
<p class="MsoNormal">ABataev created this revision.<br>
ABataev added reviewers: tra, jlebar, echristo.<br>
Herald added subscribers: JDevlieghere, aprantl, jholewinski.<br>
<br>
For better support of the CUDA profiling tools with the optimized code<br>
we should be able to emit `.target` directive without `debug` option.<br>
This is required as `.line` and `.file` directives must be emitted, but<br>
the DWARF sections should be ignored. ptxas does not support compilation<br>
of the modules with the enabled debug info when optimization level is<br>
<br>
> 0, but allows to emit linenumbers for the profilers.<br>
<br>
`-no-cuda-debug` option disables emission of the `debug` option, but not<br>
the emission of the debug info itself.<br>
<br>
<br>
Repository:<br>
rL LLVM<br>
<br>
<a href="https://nam03.safelinks.protection.outlook.com/?url=https%3A%2F%2Freviews.llvm.org%2FD45782&data=02%7C01%7C%7C6d63caffde494d491d6608d5a9446866%7C84df9e7fe9f640afb435aaaaaaaaaaaa%7C1%7C0%7C636601033622234770&sdata=TtvUF%2Fc4kvsg1Fi%2B5SiDYofzUUzBo%2BkLIRYSwclIP4g%3D&reserved=0" target="_blank">https://reviews.llvm.org/D45782</a><br>
<br>
Files:<br>
lib/Target/NVPTX/NVPTXAsmPrinter.cpp<br>
test/DebugInfo/NVPTX/debug-info.ll<br>
<br>
<br>
Index: test/DebugInfo/NVPTX/debug-info.ll<br>
===================================================================<br>
--- test/DebugInfo/NVPTX/debug-info.ll<br>
+++ test/DebugInfo/NVPTX/debug-info.ll<br>
@@ -1,4 +1,5 @@<br>
-; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s<br>
+; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda | FileCheck %s --check-prefixes=CHECK,DEBUG<br>
+; RUN: llc < %s -mtriple=nvptx64-nvidia-cuda -no-cuda-debug | FileCheck %s --check-prefixes=CHECK,NODEBUG<br>
<br>
; // Bitcode int this test case is reduced version of compiled code below:<br>
;__device__ inline void res(float x, float y, float *res) { *res = x + y; }<br>
@@ -9,7 +10,8 @@<br>
; res(a * x[i], y[i], &y[i]);<br>
;}<br>
<br>
-; CHECK: .target sm_{{[0-9]+}}//, debug<br>
+; DEBUG: .target sm_{{[0-9]+}}//, debug<br>
+; NODEBUG: .target sm_{{[0-9]+$}}<br>
<br>
; CHECK: .visible .entry _Z5saxpyifPfS_(<br>
; CHECK: .param .u32 {{.+}},<br>
Index: lib/Target/NVPTX/NVPTXAsmPrinter.cpp<br>
===================================================================<br>
--- lib/Target/NVPTX/NVPTXAsmPrinter.cpp<br>
+++ lib/Target/NVPTX/NVPTXAsmPrinter.cpp<br>
@@ -93,6 +93,11 @@<br>
<br>
#define DEPOTNAME "__local_depot"<br>
<br>
+static cl::opt<bool><br>
+ NoCudaDebug("no-cuda-debug",<br>
+ cl::desc("Do not mark ptx file as having debug info"),<br>
+ cl::init(false));<br>
+<br>
/// DiscoverDependentGlobals - Return a set of GlobalVariables on which \p V<br>
/// depends.<br>
static void<br>
@@ -876,7 +881,7 @@<br>
O << ", texmode_independent";<br>
<br>
// FIXME: remove comment once debug info is properly supported.<br>
- if (MMI && MMI->hasDebugInfo())<br>
+ if (MMI && MMI->hasDebugInfo() && !NoCudaDebug)<br>
O << "//, debug";<br>
<br>
O << "\n";<br>
<br>
<br>
_______________________________________________<br>
llvm-commits mailing list<br>
<a href="mailto:llvm-commits@lists.llvm.org" target="_blank">llvm-commits@lists.llvm.org</a><br>
<a href="https://nam03.safelinks.protection.outlook.com/?url=http%3A%2F%2Flists.llvm.org%2Fcgi-bin%2Fmailman%2Flistinfo%2Fllvm-commits&data=02%7C01%7C%7C6d63caffde494d491d6608d5a9446866%7C84df9e7fe9f640afb435aaaaaaaaaaaa%7C1%7C0%7C636601033622234770&sdata=xFD8D4h0WRTTPYapBJ3jXszJOPDasL59Us0PwB4ZyyU%3D&reserved=0" target="_blank">http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-commits</a><o:p></o:p></p>
</blockquote>
</div>
</div>
</blockquote>
<p class="MsoNormal"><o:p> </o:p></p>
</div>
</div>
</body>
</html>