<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">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? It appears that swarf injects the label after the call instruction line gets printed.</span> </div><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?</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></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">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>