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

[Alexey Bataev via cfe-dev]

Hi Artem, will check this tomorrow.

Best regards,
Alexey Bataev

2 дек. 2020 г., в 18:00, Artem Belevich <tra at google.com> написал(а):


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