<html>
<head>
<meta http-equiv="Content-Type" content="text/html; charset=UTF-8">
</head>
<body>
<p><br>
</p>
<pre class="moz-signature" cols="72">-------------
Best regards,
Alexey Bataev</pre>
<div class="moz-cite-prefix">12/3/2020 3:28 PM, Artem Belevich
пишет:<br>
</div>
<blockquote type="cite"
cite="mid:CA+wKYkOgQ5ijejnS8xhFa_URNszwV=XCy4d1jTrxyNBZntK3YQ@mail.gmail.com">
<meta http-equiv="Content-Type" content="text/html; charset=UTF-8">
<div dir="ltr">
<div dir="ltr">
<div class="gmail_default"
style="font-family:verdana,sans-serif"><br>
</div>
</div>
<br>
<div class="gmail_quote">
<div dir="ltr" class="gmail_attr">On Thu, Dec 3, 2020 at 11:57
AM Alexey.Bataev <<a href="mailto:a.bataev@outlook.com"
moz-do-not-send="true">a.bataev@outlook.com</a>> wrote:<br>
</div>
<blockquote class="gmail_quote" style="margin:0px 0px 0px
0.8ex;border-left:1px solid
rgb(204,204,204);padding-left:1ex">
<div>
<p>Hi Artem, here is what I found about this.</p>
<p>These labels are emitted only if DWARF 4 or 5 is used.<span
class="gmail_default"
style="font-family:verdana,sans-serif"> </span>They
are required for emission of the DW_AT_call_site
attribute. The info about callsites also emitted for
lineinfo emission with debug info for profiling like in
your example (-O1 -gmlt).</p>
</div>
</blockquote>
<div>
<div class="gmail_default"
style="font-family:verdana,sans-serif">Thank you for
looking into this.</div>
</div>
<div> </div>
<blockquote class="gmail_quote" style="margin:0px 0px 0px
0.8ex;border-left:1px solid
rgb(204,204,204);padding-left:1ex">
<div>
<p>Call/CallUni instruction is treated as a separate
instruction by the debug info generator and it treats it
as a separate call and emits labels for it. </p>
</div>
</blockquote>
<div><span class="gmail_default"
style="font-family:verdana,sans-serif">Would it help if we
were to fold the whole call.uni into a single line? </span></div>
</div>
</div>
</blockquote>
If it would be a single instruction - yes, it should fix it.<br>
<blockquote type="cite"
cite="mid:CA+wKYkOgQ5ijejnS8xhFa_URNszwV=XCy4d1jTrxyNBZntK3YQ@mail.gmail.com">
<div dir="ltr">
<div class="gmail_quote">
<div><span class="gmail_default"
style="font-family:verdana,sans-serif">It appears that
swarf injects the label after the call instruction line
gets printed.</span><br>
</div>
</div>
</div>
</blockquote>
Actually, it may insert the label before (for tail calls) and after
a call. <br>
<blockquote type="cite"
cite="mid:CA+wKYkOgQ5ijejnS8xhFa_URNszwV=XCy4d1jTrxyNBZntK3YQ@mail.gmail.com">
<div dir="ltr">
<div class="gmail_quote">
<div>
<div class="gmail_default"
style="font-family:verdana,sans-serif">What is the label
supposed to point at? At the call instruction itself? Or
at the return point?<br>
</div>
</div>
</div>
</div>
</blockquote>
The labels just point on the start and the end of function call,
i.e. it must be inserted before call (for tail calls) and right
after call.<br>
<blockquote type="cite"
cite="mid:CA+wKYkOgQ5ijejnS8xhFa_URNszwV=XCy4d1jTrxyNBZntK3YQ@mail.gmail.com">
<div dir="ltr">
<div class="gmail_quote">
<div><br>
</div>
<blockquote class="gmail_quote" style="margin:0px 0px 0px
0.8ex;border-left:1px solid
rgb(204,204,204);padding-left:1ex">
<div>
<p>You can try to mark the CallUni instructions as
hasDelaySlot = 1 or stop treating it as a call
instruction in NVPTXInstrInfo.td. Can't say which one is
better/correct. Looks like the representation for
call/call.uni is not quite compatible with the debug
info <br>
</p>
</div>
</blockquote>
<div>
<div class="gmail_default"
style="font-family:verdana,sans-serif">David suggested not
allowing <span class="gmail_default"></span><span
style="font-family:Arial,Helvetica,sans-serif">-fdebug-default-version=5
to override the DWARF version provided by NVPTX
back-end. That would make sense, considering that we
can't handle the newer DWARF versions anyways.</span><br>
</div>
</div>
</div>
</div>
</blockquote>
Yes, that's a good idea. If we need to change it we still can use
LLVM option to override the DWARF version.<br>
<blockquote type="cite"
cite="mid:CA+wKYkOgQ5ijejnS8xhFa_URNszwV=XCy4d1jTrxyNBZntK3YQ@mail.gmail.com">
<div dir="ltr">
<div class="gmail_quote">
<div><br>
</div>
<div>
<div class="gmail_default"
style="font-family:verdana,sans-serif">Let me see if I can
fix the lineinfo generation first. </div>
</div>
<div><br>
</div>
<div class="gmail_default"
style="font-family:verdana,sans-serif">--Artem</div>
<blockquote class="gmail_quote" style="margin:0px 0px 0px
0.8ex;border-left:1px solid
rgb(204,204,204);padding-left:1ex">
<div>
<p> </p>
<pre cols="72">-------------
Best regards,
Alexey Bataev</pre>
<div>12/2/2020 6:00 PM, Artem Belevich пишет:<br>
</div>
<blockquote type="cite">
<div dir="ltr">
<div>Hi, Alexey!</div>
<div><br>
</div>
<div>I've ran into an odd case with debug info
generation in NVPTX.</div>
<div><br>
</div>
<div>Reproduction:</div>
<div>------------------------</div>
<div>__device__ __attribute__((noinline)) void bar() {
printf("Hi!"); }<br>
__global__ void foo() { bar(); }<br>
int main(){}<br>
</div>
<div>------------------------</div>
<div><br>
</div>
<div>$ clang++ -v --cuda-gpu-arch=sm_70
--cuda-device-only <span class="gmail_default"
style="font-family:verdana,sans-serif"></span>-fdebug-default-version=5
<a href="http://a.cu" target="_blank"
moz-do-not-send="true">a.cu</a> -gmlt -O1 -c</div>
<div><br>
</div>
<div>
<div>Compilation fails due to a syntax error
reported by ptxas.</div>
<div>The reason for the error is that clang
generates a label in the middle of a `call.uni`
instuction. E.g:</div>
<div><br>
</div>
<div> { // callseq 1, 0<br>
.reg .b32 temp_param_reg;<br>
call.uni<br>
Ltmp14:<br>
_Z3barv,<br>
(<br>
);<br>
} // callseq 1<br>
</div>
<div><br>
</div>
<div>The odd part is that we're only generating line
info and there is no DWARF in the generated PTX.</div>
<div>It appears that this behavior is triggered by
`-dwarf-version=5` passed to cc1. </div>
<div>Looks like another case where PTX syntax breaks
DWARF generator assumptions.</div>
<div><br>
</div>
</div>
<div>It's possible to work around it with an
additional `-Xarch_device
-fdebug-default-version=2`, </div>
<div>but I'd appreciate it if you could take a look
and see if that could be fixed.</div>
<div><br>
</div>
-- <br>
<div dir="ltr">
<div dir="ltr">--Artem Belevich</div>
</div>
</div>
</blockquote>
</div>
</blockquote>
</div>
<br clear="all">
<div><br>
</div>
-- <br>
<div dir="ltr" class="gmail_signature">
<div dir="ltr">--Artem Belevich</div>
</div>
</div>
</blockquote>
</body>
</html>