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

Alexey Bataev via cfe-dev cfe-dev at lists.llvm.org
Wed Dec 2 15:20:21 PST 2020

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.

__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;
        } // 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>

More information about the cfe-dev mailing list