<html>
<head>
<meta http-equiv="Content-Type" content="text/html; charset=utf-8">
</head>
<body dir="auto">
Hi Artem, will check this tomorrow.<br>
<br>
<div dir="ltr">Best regards,
<div>Alexey Bataev</div>
</div>
<div dir="ltr"><br>
<blockquote type="cite">2 дек. 2020 г., в 18:00, Artem Belevich <tra@google.com> написал(а):<br>
<br>
</blockquote>
</div>
<blockquote type="cite">
<div dir="ltr">
<div dir="ltr">
<div class="gmail_default" style="font-family:verdana,sans-serif">Hi, Alexey!</div>
<div class="gmail_default" style="font-family:verdana,sans-serif"><br>
</div>
<div class="gmail_default" style="font-family:verdana,sans-serif">I've ran into an odd case with debug info generation in NVPTX.</div>
<div class="gmail_default" style="font-family:verdana,sans-serif"><br>
</div>
<div class="gmail_default" style="font-family:verdana,sans-serif">Reproduction:</div>
<div class="gmail_default" style="font-family:verdana,sans-serif">------------------------</div>
<div class="gmail_default" style="font-family:verdana,sans-serif">__device__ __attribute__((noinline)) void bar() { printf("Hi!"); }<br>
__global__ void foo() { bar(); }<br>
int main(){}<br>
</div>
<div class="gmail_default" style="font-family:verdana,sans-serif">------------------------</div>
<div class="gmail_default" style="font-family:verdana,sans-serif"><br>
</div>
<div class="gmail_default" style="font-family:verdana,sans-serif">$ clang++ -v --cuda-gpu-arch=sm_70 --cuda-device-only -fdebug-default-version=5
<a href="http://a.cu" target="_blank">a.cu</a> -gmlt -O1 -c</div>
<div><br>
</div>
<div>
<div class="gmail_default" style="font-family:verdana,sans-serif">Compilation fails due to a syntax error reported by ptxas.</div>
<div class="gmail_default" style="font-family:verdana,sans-serif">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" style="font-family:verdana,sans-serif"><br>
</div>
<div class="gmail_default" style="font-family:verdana,sans-serif"> { // 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" style="font-family:verdana,sans-serif"><br>
</div>
<div class="gmail_default" style="font-family:verdana,sans-serif">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" style="font-family:verdana,sans-serif">It appears that this behavior is triggered by `-dwarf-version=5` passed to cc1. </div>
<div class="gmail_default" style="font-family:verdana,sans-serif">Looks like another case where PTX syntax breaks DWARF generator assumptions.</div>
<div class="gmail_default" style="font-family:verdana,sans-serif"><br>
</div>
</div>
<div class="gmail_default" style="font-family:verdana,sans-serif">It's possible to work around it with an additional `-Xarch_device -fdebug-default-version=2`, </div>
<div class="gmail_default" style="font-family:verdana,sans-serif">but I'd appreciate it if you could take a look and see if that could be fixed.</div>
<div class="gmail_default" style="font-family:verdana,sans-serif"><br>
</div>
-- <br>
<div dir="ltr" data-smartmail="gmail_signature">
<div dir="ltr">--Artem Belevich</div>
</div>
</div>
</div>
</blockquote>
</body>
</html>