The error reported is "Use of undeclared identifier cudaConfigureCall": axpy_kernel <<<1, 2>>> (a, y, x); ^ There is no such error when compiling by clang against CUDA 11.0 (before Update 1) and against previous versions. The error occurs when compiling CUDA code containing kernel launch in CUDAKernelCallExpr syntax by all versions of clang, which initially supported CUDA 11.0: 10.x - 11.x, 12.0.0git (trunk).
Kernel launch is the core functionality for CUDA, thus setting priority to "release blocker". The issue has to be fixed not only in LLVM trunk but also in LLVM 10.x and 11.x releases.
Add Artem for awareness. I remember clang has added a new kernel launching API for CUDA. Can this issue be fixed by using the new kernel launching API by default? Thanks.
The issue is that CUDA-11.1 no longer ships with version.txt. As the result, clang does not know which CUDA version it is and defaults to 7.0 which was the last revision w/o the vestion.txt. Quick workaround would be to add a version'txt file in your CUDA installation directory with "CUDA Version 11.1.1" in it. With no explicit version file in the install dir, we'll need a different method for figuring out CUDA version. We may also need to add a --cuda-version=X.Y option to allow users to override/specify the version if it's not detected.
Changing the assumed default w/o version.txt file present will effectively make it impossible to use CUDA-7. I doubt there are many (or any) CUDA-7 users left, but I'm somewhat wary of decommissioning it on the emergency basis, though, admittedly, it's probably more useful to be able to compile with new CUDA version than with an ancient one. The workaround with manually placed version.txt file should work for CUDA-7 if someone happens to still use it. Also, technically clang's not working with CUDA-11.1 is also "works as intended". Clang even issues a warning about any CUDA version newer than 10.1. The fact that it *mostly* worked it not a guarantee that it will magically always work with any newer version. We've just been lucky that it (mostly) worked with 10.2 and 11.0. I'll send the patch for changing the default and we can discuss the merits/downsides there.
It's actually not that bad. We can allow detection of CUDA-11.1 as a new unknow version without affecting the old ones. https://reviews.llvm.org/D89752 +cc: hans@ for release cherry-picks.
> +cc: hans@ for release cherry-picks. The release shipped last week :) File a blocker on PR47800 for 11.0.1.
> The release shipped last week :) > File a blocker on PR47800 for 11.0.1. The patch is needed for LLVM 10 as well. Can we expect the patch in 10.0.2?
> The issue is that CUDA-11.1 no longer ships with version.txt. Actually, CUDA-11.0 Update 1 no longer ships it.
Should be fixed by these commits: D89752 65d206484c54177641d4b11d42cab1f1acc8c0c7 D89832 e7fe125b776bf08d95e60ff3354a5c836218a0e6
https://reviews.llvm.org/D89752 https://reviews.llvm.org/D89832