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

[Artem Belevich via cfe-dev]

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 -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
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-dev/attachments/20201202/23cc950d/attachment.html>