search

llvm / llvm-project

The LLVM Project is a collection of modular and reusable compiler and toolchain technologies.
http://llvm.org
Other
28.55k stars 11.8k forks source link

CUDA pure virtual inheritance issue #49183

Open llvmbot opened 3 years ago

llvmbot commented 3 years ago
Bugzilla Link 49839
Version trunk
OS Linux
Attachments Program showing issue
Reporter LLVM Bugzilla Contributor
CC @Artem-B

Extended Description

CUDA pure virtual inheritance issue:

Summary:

The attached program initializes a class with a pure virtual method and overrides this pure virtual method in a child class. This program fails to compile in Clang CUDA with no optimization applied. However, this program does compile under certain optimization settings: -O1, and -O2.

Description:

The attached program initializes a class with a pure virtual method, and a child class which overrides this pure virtual method.

We use Clang Trunk (SHA: 1cc9d949a1233e8b17b3b345ccb67ca7296c1a6c) to compile the attached program.

Compiling the attached program in Clang CUDA fails, while also outputting the following errors: ptxas fatal : Cannot take address of function '__cxa_pure_virtual' clang-13: error: ptxas command failed with exit code 255 (use -v to see invocation) clang version 13.0.0 (https://github.com/llvm/llvm-project.git 1cc9d949a1233e8b17b3b345ccb67ca7296c1a6c) Target: x86_64-unknown-linux-gnu Thread model: posix InstalledDir: /home/kjain/projects/llvm-project/build/bin clang-13: note: diagnostic msg: 

PLEASE ATTACH THE FOLLOWING FILES TO THE BUG REPORT:
Preprocessed source(s) and associated run script(s) are located at:
clang-13: note: diagnostic msg: /tmp/22-3e3111.cu
clang-13: note: diagnostic msg: /tmp/22-aee9af.cu
clang-13: note: diagnostic msg: /tmp/22-3e3111.sh
clang-13: note: diagnostic msg: 

********************

This same program compiles successfully under optimization settings -O1 and -O2, while failing under the optimization setting -O0.

This program also compiles successfully in nvcc.

The files outputted by clang-13 are also attached.

Artem-B commented 3 years ago

ptxas fatal : Cannot take address of function '__cxa_pure_virtual'

AFAICT, the failure is due to the fact that we do not have the standard runtime support library on the GPU side of the compilation, so there's nothing that would provide __cxa_pure_virtual.

If compiler failed to eliminate such a library call, there's nothing we can do other than complain during the final linking stage, which is effectively what ptxas does in this case.

As a short-term workaround you can add something like this somewhere in the CUDA code you're compiling:

extern "C" __device__ void __cxa_pure_virtual() { __trap(); }

This may be sufficient to work around your problem. Longer term we should probably consider adding some C++ runtime support functions to the CUDA header wrappers that we ship with clang.

llvmbot commented 3 years ago

Wasn't able to attach the files shown in the diagnostic message, as they are too large (>1000kb).