Clang fails to compile cuda code in C++20 mode, works in C++17

[bmanga]

Bugzilla Link	48228
Resolution	FIXED
Resolved on	Dec 09, 2020 09:48
Version	11.0
OS	Linux
Blocks	#47144
CC	@bmanga,@jdoerfert,@Artem-B,@tstellar
Fixed by commit(s)	9a46505 4326792 aa29049 59012b6

Extended Description

Trying to compile the following cuda code:

__device__ void foo() {}
int main(){}

Works in C++17 mode, fails in C++20 mode with a series of errors about unknown type name '__device__' from the cuda_wrappers headers.

[bmanga]

assigned to @Artem-B

[Artem-B]

Ouch. In c++20 mode one of the the standard C++ headers includes which ends up including cuda_wrappers/new before we've got the standard CUDA macros ready.

Should be fixed by https://reviews.llvm.org/D91807

[Artem-B]

Fixed in https://reviews.llvm.org/rG9a465057a64dba8a8614424d26136f5c0452bcc3

[Artem-B]

This would be a good low-risk fix to cherry-pick into 11.0.1

[bmanga]

Hi, I applied the patch you posted, but now I am getting the following errors:

/usr/lib/llvm-11/lib/clang/11.0.1/include/cuda_wrappers/new:50:12: error: reference to host function 'malloc' in device function
return ::malloc(size);
^
/usr/include/stdlib.h:539:14: note: 'malloc' declared here
extern void *malloc (size_t __size) __THROW attribute_malloc __wur;

/usr/lib/llvm-11/lib/clang/11.0.1/include/cuda_wrappers/new:67:7: error: reference to host function 'free' in device function
::free(ptr);
^
/usr/include/stdlib.h:563:13: note: 'free' declared here
extern void free (void *__ptr) __THROW;

[tstellar]

I'm going to hold off on backporting this until the failure in comment #​4 is addressed.

[Artem-B]

Interesting.

Can you tell me which OS and libstdc++ version you're using?

Did the patched cuda_wrappers/new from the source tree make it into the /usr/lib/llvm-11/lib/clang/11.0.1/include/cuda_wrappers/new where your clang is looking for them?

[bmanga]

Ubuntu 18.04
libstdc++ should be from gcc 10.1:
_GLIBCXX_RELEASE 10
GLIBCXX 20200515

I applied your phabricator patch to that file, yes.

[bmanga]

I have just tested it with libc++ and it works fine (but of course it would be nice to have it working with libstdc++).

[Artem-B]

Interesting. My version of libstdc++ is 20200918 and I do not see this issue.

Would you be able to capture the output of the following command and attach it to the bug? Adjust --cuda-path to point to the CUDA version you're using.

$ bin/clang++ -v --cuda-path=$HOME/local/cuda-11.0 --cuda-gpu-arch=sm_70 --std=c++20 -x cuda --stdlib=libstdc++ --cuda-device-only /dev/null -o - -dD -E

[bmanga]

clang command log
Here it is (I used arch sm_61, not sure if it makes a difference)

[Artem-B]

Thank you for the log. Indeed the device-side declarations for malloc/free are missing when cuda_wrappers/new has been includes.

Clang in HEAD does not complain unless we must emit operator new which hides the issue. I can reproduce it by using new in device-side code.

Looks like D91807 does not really fix the problem, just hides it sometimes.

I'll need to find a better way to deal with this.

[tstellar]

Do I understand correctly that this patch does not cause a regression, it just fixes a bug in some cases but not others?

Do we still want to try to backport this?

[Artem-B]

Alas, the patch only hides the issue, and only in a rather useless toy cases.
Unfortunately I do not think it would do any good to backport it to 11.0.1 and I do not have a better fix for it yet.

If I were to attempt fixing it for 11.0.1, how much time do I have? I should have a better idea what can be done by the end of the day tomorrow.

[Artem-B]

https://reviews.llvm.org/D91807 has been updated with a better fix.

[Artem-B]

Landed in 4326792

@​tstellar if/when you cherry-pick it, it will need to apply on top of 9a46505. I didn't think to revert the first patch first. :-(

[Artem-B]

I've enabled testing w/ c++17 and c++20 on CUDA test bots and the tests compile and pass.
http://lab.llvm.org:8011/#/builders/55/builds/2569

I think the patch is ready for cherry-picking into 11.0.1, if the train is not gone yet.

[tstellar]

Merged: 59012b6

[wrvsrx]

This problem reappear in clangd. Now clang can compile cuda programs with c++20, but clangd raises error.

clangd/clangd#1815
clangd/clangd#1452