<html>
<head>
<meta http-equiv="Content-Type" content="text/html; charset=UTF-8">
</head>
<body>
<p>Hi Artem, here is what I found about this.</p>
<p>These labels are emitted only if DWARF 4 or 5 is used. 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>
<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. 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>
<pre class="moz-signature" cols="72">-------------
Best regards,
Alexey Bataev</pre>
<div class="moz-cite-prefix">12/2/2020 6:00 PM, Artem Belevich
пишет:<br>
</div>
<blockquote type="cite"
cite="mid:CA+wKYkPETHF8o30MdnqFdGYhFoBuwG+z6HWgDiVeijiyn1K4uw@mail.gmail.com">
<div dir="ltr">
<div class="gmail_default">Hi, Alexey!</div>
<div class="gmail_default"><br>
</div>
<div class="gmail_default">I've ran into an odd case with debug
info generation in NVPTX.</div>
<div class="gmail_default"><br>
</div>
<div class="gmail_default">Reproduction:</div>
<div class="gmail_default">------------------------</div>
<div class="gmail_default">__device__ __attribute__((noinline))
void bar() { printf("Hi!"); }<br>
__global__ void foo() { bar(); }<br>
int main(){}<br>
</div>
<div class="gmail_default">------------------------</div>
<div class="gmail_default"><br>
</div>
<div class="gmail_default">$ clang++ -v --cuda-gpu-arch=sm_70
--cuda-device-only -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 class="gmail_default">Compilation fails due to a syntax
error reported by ptxas.</div>
<div class="gmail_default">The reason for the error is that
clang generates a label in the middle of a `call.uni`
instuction. E.g:</div>
<div class="gmail_default"><br>
</div>
<div class="gmail_default"> { // callseq 1, 0<br>
.reg .b32 temp_param_reg;<br>
call.uni<br>
Ltmp14:<br>
_Z3barv,<br>
(<br>
);<br>
} // callseq 1<br>
</div>
<div class="gmail_default"><br>
</div>
<div class="gmail_default">The odd part is that we're only
generating line info and there is no DWARF in the generated
PTX.</div>
<div class="gmail_default">It appears that this behavior is
triggered by `-dwarf-version=5` passed to cc1. </div>
<div class="gmail_default">Looks like another case where PTX
syntax breaks DWARF generator assumptions.</div>
<div class="gmail_default"><br>
</div>
</div>
<div class="gmail_default">It's possible to work around it with
an additional `-Xarch_device -fdebug-default-version=2`, </div>
<div class="gmail_default">but I'd appreciate it if you could
take a look and see if that could be fixed.</div>
<div class="gmail_default"><br>
</div>
-- <br>
<div dir="ltr" data-smartmail="gmail_signature">
<div dir="ltr">--Artem Belevich</div>
</div>
</div>
</blockquote>
</body>
</html>