<div dir="ltr"><div dir="ltr"><br></div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Thu, Dec 3, 2020 at 12:29 PM Artem Belevich via cfe-dev <<a href="mailto:cfe-dev@lists.llvm.org">cfe-dev@lists.llvm.org</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 dir="ltr"><div dir="ltr"><div 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" target="_blank">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 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 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></div></div></div></blockquote><div><br></div><div>I believe it's immediately after the call.<br><br>Testing an x86 example:<br><br>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">$ cat test.c</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">void f1();</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">void f2();</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">void f3() {</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures"><span class="gmail-Apple-converted-space"> </span>f1();</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures"><span class="gmail-Apple-converted-space"> </span>f2();</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">}</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">$ clang-tot -gdwarf-5 test.c -c -O3 && llvm-dwarfdump-tot test.o | grep "DW_TAG\|_call_" && llvm-objdump -d test.o</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">0x0000000c: </span><span class="gmail-s4" style="font-variant-ligatures:no-common-ligatures;color:rgb(202,51,35)"><b>DW_TAG</b></span><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">_compile_unit</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">0x00000023: <span class="gmail-Apple-converted-space"> </span></span><span class="gmail-s4" style="font-variant-ligatures:no-common-ligatures;color:rgb(202,51,35)"><b>DW_TAG</b></span><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">_subprogram</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures"><span class="gmail-Apple-converted-space"> </span>DW_AT</span><span class="gmail-s4" style="font-variant-ligatures:no-common-ligatures;color:rgb(202,51,35)"><b>_call_</b></span><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">all_calls<span class="gmail-Apple-converted-space"> </span>(true)</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">0x0000002e: <span class="gmail-Apple-converted-space"> </span></span><span class="gmail-s4" style="font-variant-ligatures:no-common-ligatures;color:rgb(202,51,35)"><b>DW_TAG_call_</b></span><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">site</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures"><span class="gmail-Apple-converted-space"> </span>DW_AT</span><span class="gmail-s4" style="font-variant-ligatures:no-common-ligatures;color:rgb(202,51,35)"><b>_call_</b></span><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">origin <span class="gmail-Apple-converted-space"> </span>(0x0000003b)</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures"><span class="gmail-Apple-converted-space"> </span>DW_AT</span><span class="gmail-s4" style="font-variant-ligatures:no-common-ligatures;color:rgb(202,51,35)"><b>_call_</b></span><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">return_pc<span class="gmail-Apple-converted-space"> </span>(0x0000000000000008)</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">0x00000034: <span class="gmail-Apple-converted-space"> </span></span><span class="gmail-s4" style="font-variant-ligatures:no-common-ligatures;color:rgb(202,51,35)"><b>DW_TAG_call_</b></span><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">site</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures"><span class="gmail-Apple-converted-space"> </span>DW_AT</span><span class="gmail-s4" style="font-variant-ligatures:no-common-ligatures;color:rgb(202,51,35)"><b>_call_</b></span><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">origin <span class="gmail-Apple-converted-space"> </span>(0x00000041)</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures"><span class="gmail-Apple-converted-space"> </span>DW_AT</span><span class="gmail-s4" style="font-variant-ligatures:no-common-ligatures;color:rgb(202,51,35)"><b>_call_</b></span><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">tail_call<span class="gmail-Apple-converted-space"> </span>(true)</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures"><span class="gmail-Apple-converted-space"> </span>DW_AT</span><span class="gmail-s4" style="font-variant-ligatures:no-common-ligatures;color:rgb(202,51,35)"><b>_call_</b></span><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">pc (0x000000000000000b)</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">0x0000003b: <span class="gmail-Apple-converted-space"> </span></span><span class="gmail-s4" style="font-variant-ligatures:no-common-ligatures;color:rgb(202,51,35)"><b>DW_TAG</b></span><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">_subprogram</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">0x0000003f: <span class="gmail-Apple-converted-space"> </span></span><span class="gmail-s4" style="font-variant-ligatures:no-common-ligatures;color:rgb(202,51,35)"><b>DW_TAG</b></span><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">_unspecified_parameters</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">0x00000041: <span class="gmail-Apple-converted-space"> </span></span><span class="gmail-s4" style="font-variant-ligatures:no-common-ligatures;color:rgb(202,51,35)"><b>DW_TAG</b></span><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">_subprogram</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">0x00000045: <span class="gmail-Apple-converted-space"> </span></span><span class="gmail-s4" style="font-variant-ligatures:no-common-ligatures;color:rgb(202,51,35)"><b>DW_TAG</b></span><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">_unspecified_parameters</span></p>
<p class="gmail-p2" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0);min-height:13px"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures"></span><br></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">test.o: file format elf64-x86-64</span></p>
<p class="gmail-p2" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0);min-height:13px"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures"></span><br></p>
<p class="gmail-p2" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0);min-height:13px"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures"></span><br></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">Disassembly of section .text:</span></p>
<p class="gmail-p2" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0);min-height:13px"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures"></span><br></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures">0000000000000000 <f3>:</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures"><span class="gmail-Apple-converted-space"> </span>0: 50<span class="gmail-Apple-converted-space"> </span>pushq <span class="gmail-Apple-converted-space"> </span>%rax</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures"><span class="gmail-Apple-converted-space"> </span>1: 31 c0 <span class="gmail-Apple-converted-space"> </span>xorl<span class="gmail-Apple-converted-space"> </span>%eax, %eax</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures"><span class="gmail-Apple-converted-space"> </span>3: e8 00 00 00 00<span class="gmail-Apple-converted-space"> </span>callq <span class="gmail-Apple-converted-space"> </span>0x8 <f3+0x8></span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures"><span class="gmail-Apple-converted-space"> </span>8: 31 c0 <span class="gmail-Apple-converted-space"> </span>xorl<span class="gmail-Apple-converted-space"> </span>%eax, %eax</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures"><span class="gmail-Apple-converted-space"> </span>a: 59<span class="gmail-Apple-converted-space"> </span>popq<span class="gmail-Apple-converted-space"> </span>%rcx</span></p>
<p class="gmail-p1" style="margin:0px;font-variant-numeric:normal;font-variant-east-asian:normal;font-stretch:normal;font-size:11px;line-height:normal;font-family:Menlo;color:rgb(0,0,0)"><span class="gmail-s1" style="font-variant-ligatures:no-common-ligatures"><span class="gmail-Apple-converted-space"> </span>b: e9 00 00 00 00<span class="gmail-Apple-converted-space"> </span>jmp <span class="gmail-Apple-converted-space"> </span>0x10 <f3+0x10></span></p><br><br>But the tail call uses the jump location, because it can't do anything else.<br><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 dir="ltr"><div class="gmail_quote"><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 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 style="font-family:verdana,sans-serif">Let me see if I can fix the lineinfo generation first. </div></div></div></div></blockquote><div><br></div><div>Yeah, -gdwarf-5 and -fdebug-default-version=5 -g should behave the same. So whatever it is that's disabling/downgrading to DWARFv2 for NVPTX when the user uses -gdwarf-5 shoudl do the same for -fdebug-default-version=5</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 dir="ltr"><div class="gmail_quote"><div><br></div><div 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"><div dir="ltr">--Artem Belevich</div></div></div>
_______________________________________________<br>
cfe-dev mailing list<br>
<a href="mailto:cfe-dev@lists.llvm.org" target="_blank">cfe-dev@lists.llvm.org</a><br>
<a href="https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-dev" rel="noreferrer" target="_blank">https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-dev</a><br>
</blockquote></div></div>