[cfe-dev] compiling CUDA w/ -fdebug-default-version=5 generates invalid PTX

[Alexey.Bataev via cfe-dev]

-------------
Best regards,
Alexey Bataev

12/3/2020 3:28 PM, Artem Belevich пишет:
>
>
> On Thu, Dec 3, 2020 at 11:57 AM Alexey.Bataev <a.bataev at outlook.com
> <mailto:a.bataev at outlook.com>> wrote:
>
>     Hi Artem, here is what I found about this.
>
>     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).
>
> Thank you for looking into this.
>  
>
>     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.
>
> Would it help if we were to fold the whole call.uni into a single line?
If it would be a single instruction - yes, it should fix it.
> It appears that swarf injects the label after the call instruction
> line gets printed.
Actually, it may insert the label before (for tail calls) and after a call.
> What is the label supposed to point at? At the call instruction
> itself? Or at the return point?
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.
>
>     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
>
> David suggested not allowing -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.
Yes, that's a good idea. If we need to change it we still can use LLVM
option to override the DWARF version.
>
> Let me see if I can fix the lineinfo generation first. 
>
> --Artem
>
>     -------------
>     Best regards,
>     Alexey Bataev
>
>     12/2/2020 6:00 PM, Artem Belevich пишет:
>>     Hi, Alexey!
>>
>>     I've ran into an odd case with debug info generation in NVPTX.
>>
>>     Reproduction:
>>     ------------------------
>>     __device__ __attribute__((noinline)) void bar() { printf("Hi!"); }
>>     __global__ void foo() { bar(); }
>>     int main(){}
>>     ------------------------
>>
>>     $ clang++ -v --cuda-gpu-arch=sm_70 --cuda-device-only
>>     -fdebug-default-version=5 a.cu <http://a.cu> -gmlt -O1 -c
>>
>>     Compilation fails due to a syntax error reported by ptxas.
>>     The reason for the error is that clang generates a label in the
>>     middle of a `call.uni` instuction. E.g:
>>
>>             { // callseq 1, 0
>>             .reg .b32 temp_param_reg;
>>             call.uni
>>     Ltmp14:
>>             _Z3barv,
>>             (
>>             );
>>             } // callseq 1
>>
>>     The odd part is that we're only generating line info and there is
>>     no DWARF in the generated PTX.
>>     It appears that this behavior is triggered by `-dwarf-version=5`
>>     passed to cc1. 
>>     Looks like another case where PTX syntax breaks DWARF generator
>>     assumptions.
>>
>>     It's possible to work around it with an additional `-Xarch_device
>>     -fdebug-default-version=2`, 
>>     but I'd appreciate it if you could take a look and see if that
>>     could be fixed.
>>
>>     -- 
>>     --Artem Belevich
>
>
>
> -- 
> --Artem Belevich
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-dev/attachments/20201203/1ff4090e/attachment-0001.html>
-------------- next part --------------
A non-text attachment was scrubbed...
Name: OpenPGP_signature
Type: application/pgp-signature
Size: 840 bytes
Desc: OpenPGP digital signature
URL: <http://lists.llvm.org/pipermail/cfe-dev/attachments/20201203/1ff4090e/attachment-0001.sig>