<html>
<head>
<meta http-equiv="Content-Type" content="text/html; charset=utf-8">
</head>
<body dir="auto">
We translate clang's On to the corresponding llvm and ptxas On. Llvm emits ptx file, then clang invokes ptxas to compile PTX file into binary<br>
<br>
<div id="AppleMailSignature">Best regards,
<div>Alexey Bataev</div>
</div>
<div><br>
23 апр. 2018 г., в 19:40, David Blaikie <<a href="mailto:dblaikie@gmail.com">dblaikie@gmail.com</a>> написал(а):<br>
<br>
</div>
<blockquote type="cite">
<div>
<div dir="ltr">& is the ptxas tool given a specific optimization level? How does it know whether the code is optimized?</div>
<br>
<div class="gmail_quote">
<div dir="ltr">On Mon, Apr 23, 2018 at 4:39 PM Alexey Bataev <<a href="mailto:a.bataev@hotmail.com">a.bataev@hotmail.com</a>> wrote:<br>
</div>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<div dir="auto">It is the limitation of the ptxas tool, which is used for the compilation of the PTX code.<br>
<br>
<div id="m_-1764899498981951189AppleMailSignature">Best regards,
<div>Alexey Bataev</div>
</div>
<div><br>
23 апр. 2018 г., в 19:02, David Blaikie <<a href="mailto:dblaikie@gmail.com" target="_blank">dblaikie@gmail.com</a>> написал(а):<br>
<br>
</div>
</div>
<div dir="auto">
<blockquote type="cite">
<div>
<div dir="ltr">But optimizations aren't really a defined thing to the assembler, right? How does nvptx know if optimizations are enabled? - I guess you ask it to do additional optimizations?</div>
<br>
<div class="gmail_quote">
<div dir="ltr">On Mon, Apr 23, 2018 at 3:49 PM Alexey Bataev <<a href="mailto:a.bataev@hotmail.com" target="_blank">a.bataev@hotmail.com</a>> wrote:<br>
</div>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<div dir="auto">If optimisation is disabled(O0), you may have ',debug' option, if optimisation is enabled (O1, O2,...), you are not allowed to have this option. <br>
<br>
<div id="m_-1764899498981951189m_-1302320397317942555AppleMailSignature">Best regards,
<div>Alexey Bataev</div>
</div>
<div><br>
23 апр. 2018 г., в 18:24, David Blaikie <<a href="mailto:dblaikie@gmail.com" target="_blank">dblaikie@gmail.com</a>> написал(а):<br>
<br>
</div>
</div>
<div dir="auto">
<blockquote type="cite">
<div>
<div dir="ltr">Is there a particularly strong/specific/narrow definition of "optimization" that's relevant here/that causes nvptx or whatnot to have problems?</div>
<br>
<div class="gmail_quote">
<div dir="ltr">On Mon, Apr 23, 2018 at 2:57 PM Alexey Bataev <<a href="mailto:a.bataev@hotmail.com" target="_blank">a.bataev@hotmail.com</a>> wrote:<br>
</div>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<div dir="auto">The end goal is have a way to emit '.loc' and '.file' directives only, without the ',debug' option. Maybe, we could check for the optimization level here: if it is 0, we can emit ',debug', otherwise - disable the emission of the option. <br>
<br>
<div id="m_-1764899498981951189m_-1302320397317942555m_7309642758832753996AppleMailSignature">
Best regards,
<div>Alexey Bataev</div>
</div>
<div><br>
23 апр. 2018 г., в 17:47, David Blaikie <<a href="mailto:dblaikie@gmail.com" target="_blank">dblaikie@gmail.com</a>> написал(а):<br>
<br>
</div>
</div>
<div dir="auto">
<blockquote type="cite">
<div>
<div dir="ltr">yeah, controlling the ',debug' (what is that anyway? Guess it's an nvptx thing?) would be a bit awkward from a DICompileUnit property - since there may be multiple DICompileUnits in a given module.<br>
<br>
Dunno what other debug info owners think about that - could use a module flag, could use a "scan all the CUs" sort of thing (but I guess that sort of "merge this property" functionality is what we have with module flags anyway). Other ideas? <br>
<br>
My original question was: what's the end goal here? Generally we don't want to add more backend options (they don't work through LTO, don't work well for LLVM-as-a-library, etc) & would usually either put something in the IR or weave it through the MCOptions
or the like. If this doesn't really need to be an option & should just be hard coded for cuda, maybe that's how it should be done.<br>
<br>
- Dave</div>
<br>
<div class="gmail_quote">
<div dir="ltr">On Mon, Apr 23, 2018 at 2:43 PM Alexey Bataev <<a href="mailto:a.bataev@hotmail.com" target="_blank">a.bataev@hotmail.com</a>> wrote:<br>
</div>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<div dir="auto">Ok, I see. In this case we can add a new mode that emits only '.loc' and '.file' directives without the dwarf sections, but we still need to disable the emission of ',debug' option in the target directive.<br>
<br>
<div id="m_-1764899498981951189m_-1302320397317942555m_7309642758832753996m_-2554786336117665563AppleMailSignature">
Best regards,
<div>Alexey Bataev</div>
</div>
<div><br>
23 апр. 2018 г., в 17:38, David Blaikie <<a href="mailto:dblaikie@gmail.com" target="_blank">dblaikie@gmail.com</a>> написал(а):<br>
<br>
</div>
</div>
<div dir="auto">
<blockquote type="cite">
<div>
<div dir="ltr">It sets the emission kind of the DICompileUnit to NoDebug - but it's still a DICompileUnit that goes into the IR, which is distinct from the absence of a DICompileUnit entirely.<br>
<br>
No DICompileUnit means no debug info in the IR at all.<br>
A DICompileUnit with NoDebug emission kind is debug info in the IR, but it tells the backend of LLVM not to emit any debug info for this DICompileUnit and the things inside it.</div>
<br>
<div class="gmail_quote">
<div dir="ltr">On Mon, Apr 23, 2018 at 2:35 PM Alexey Bataev <<a href="mailto:a.bataev@hotmail.com" target="_blank">a.bataev@hotmail.com</a>> wrote:<br>
</div>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<div dir="auto">Hmm, what I see in clang codegen that currently it just disables emission of any debug info, including debug locations. <br>
<br>
<div id="m_-1764899498981951189m_-1302320397317942555m_7309642758832753996m_-2554786336117665563m_6836801640579460067AppleMailSignature">
Best regards,
<div>Alexey Bataev</div>
</div>
<div><br>
23 апр. 2018 г., в 17:32, David Blaikie <<a href="mailto:dblaikie@gmail.com" target="_blank">dblaikie@gmail.com</a>> написал(а):<br>
<br>
</div>
</div>
<div dir="auto">
<blockquote type="cite">
<div>
<div dir="ltr">It keeps the debug locations on instructions, but does not emit debug info.<br>
<br>
This is used for sample-based profile guided optimization. Optimizations can use the locations on instructions to associate the samples in a profile with those instructions.<br>
<br>
No debug info need be emitted in this case.</div>
<br>
<div class="gmail_quote">
<div dir="ltr">On Mon, Apr 23, 2018 at 2:31 PM Alexey Bataev <<a href="mailto:a.bataev@hotmail.com" target="_blank">a.bataev@hotmail.com</a>> wrote:<br>
</div>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<div dir="auto">But why do we need LocTrackingOnly at all? If it just disables the emission of the debug info completely?<br>
<br>
<div id="m_-1764899498981951189m_-1302320397317942555m_7309642758832753996m_-2554786336117665563m_6836801640579460067m_-5819398429401669935AppleMailSignature">
Best regards,
<div>Alexey Bataev</div>
</div>
<div><br>
23 апр. 2018 г., в 17:28, David Blaikie <<a href="mailto:dblaikie@gmail.com" target="_blank">dblaikie@gmail.com</a>> написал(а):<br>
<br>
</div>
</div>
<div dir="auto">
<blockquote type="cite">
<div>
<div dir="ltr">Yes, that's why Paul mentioned what you want is somewhere between LocTrackingOnly and LineTablesOnly.</div>
<br>
<div class="gmail_quote">
<div dir="ltr">On Mon, Apr 23, 2018 at 2:25 PM Alexey Bataev <<a href="mailto:a.bataev@hotmail.com" target="_blank">a.bataev@hotmail.com</a>> wrote:<br>
</div>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<div dir="auto">Yes, and I could use LocTrackingOnly, but just like I said it currently does not work at all, it does not allow to emit any debug info at all. <br>
<br>
<div id="m_-1764899498981951189m_-1302320397317942555m_7309642758832753996m_-2554786336117665563m_6836801640579460067m_-5819398429401669935m_-5598698118171778765AppleMailSignature">
Best regards,
<div>Alexey Bataev</div>
</div>
<div><br>
23 апр. 2018 г., в 17:20, David Blaikie <<a href="mailto:dblaikie@gmail.com" target="_blank">dblaikie@gmail.com</a>> написал(а):<br>
<br>
</div>
</div>
<div dir="auto">
<blockquote type="cite">
<div>
<div dir="ltr">Yes, that's the purpose of LocTrackingOnly, which is why it doesn't sound like it'd do what you want? Sounds like you want the line table to be emitted for sample profiling?<br>
<br>
<div class="gmail_quote">
<div dir="ltr">On Mon, Apr 23, 2018 at 2:01 PM Alexey Bataev <<a href="mailto:a.bataev@hotmail.com" target="_blank">a.bataev@hotmail.com</a>> wrote:<br>
</div>
<blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">
<div dir="auto">Actually, LocTrackingOnly might work, but currently it just disables debug info completely and does not emit anything. <br>
<br>
<div id="m_-1764899498981951189m_-1302320397317942555m_7309642758832753996m_-2554786336117665563m_6836801640579460067m_-5819398429401669935m_-5598698118171778765m_-1249439991652957429AppleMailSignature">
Best regards,
<div>Alexey Bataev</div>
</div>
<div><br>
23 апр. 2018 г., в 16:58, "<a href="mailto:paul.robinson@sony.com" target="_blank">paul.robinson@sony.com</a>" <<a href="mailto:paul.robinson@sony.com" target="_blank">paul.robinson@sony.com</a>> написал(а):<br>
<br>
</div>
</div>
<div dir="auto">
<blockquote type="cite">
<div>
<div class="m_-1764899498981951189m_-1302320397317942555m_7309642758832753996m_-2554786336117665563m_6836801640579460067m_-5819398429401669935m_-5598698118171778765m_-1249439991652957429WordSection1">
<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).<u></u><u></u></span></p>
<p class="MsoNormal"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d">--paulr<u></u><u></u></span></p>
<p class="MsoNormal"><a name="m_-1764899498981951189_m_-1302320397317942555_m_7309642758832753996_m_-2554786336117665563_m_6836801640579460067_m_-5819398429401669935_m_-5598698118171778765_m_-1249439991652957429__MailEndCompose"><span style="font-size:11.0pt;font-family:"Calibri","sans-serif";color:#1f497d"><u></u> <u></u></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 [<a href="mailto:llvm-commits-bounces@lists.llvm.org" target="_blank">mailto:llvm-commits-bounces@lists.llvm.org</a>]
<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> <a href="mailto:reviews+d45782+public+181908537bbf2c34@reviews.llvm.org" target="_blank">
reviews+d45782+public+181908537bbf2c34@reviews.llvm.org</a>; <a href="mailto:jholewinski@nvidia.com" target="_blank">
jholewinski@nvidia.com</a>; <a href="mailto:llvm-commits@lists.llvm.org" target="_blank">
llvm-commits@lists.llvm.org</a>; Alexey Bataev<br>
<b>Subject:</b> Re: [PATCH] D45782: [DEBUGINFO, NVPTX] Allow to disable debug info from command line.<u></u><u></u></span></p>
</div>
</div>
<p class="MsoNormal"><u></u> <u></u></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).<u></u><u></u></pre>
<pre>The lineinfo is required for the nvprof, the CUDA profiler. So, our purpose is to emit these directives but without DWARF sections.<u></u><u></u></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.<u></u><u></u></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. <u></u><u></u></pre>
<pre>-------------<u></u><u></u></pre>
<pre>Best regards,<u></u><u></u></pre>
<pre>Alexey Bataev<u></u><u></u></pre>
<div>
<p class="MsoNormal">23.04.2018 14:02, David Blaikie пишет:<u></u><u></u></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?<u></u><u></u></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" target="_blank">llvm-commits@lists.llvm.org</a>> wrote:<u></u><u></u></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><u></u><u></u></p>
</blockquote>
</div>
</div>
</blockquote>
<p class="MsoNormal"><u></u> <u></u></p>
</div>
</div>
</div>
</blockquote>
</div>
</blockquote>
</div>
</div>
</div>
</blockquote>
</div>
</blockquote>
</div>
</div>
</blockquote>
</div>
</blockquote>
</div>
</div>
</blockquote>
</div>
</blockquote>
</div>
</div>
</blockquote>
</div>
</blockquote>
</div>
</div>
</blockquote>
</div>
</blockquote>
</div>
</div>
</blockquote>
</div>
</blockquote>
</div>
</div>
</blockquote>
</div>
</blockquote>
</div>
</div>
</blockquote>
</body>
</html>