rgreenblatt added a comment.
In D94337#2491515 <https://reviews.llvm.org/D94337#2491515>, @tra wrote:
> Does it make sense?
Yep, most of what your saying makes sense to me. Thanks for taking the time to
review this.
A few notes:
> That would still require properly defined CUDA macros. CUDA in clang relies
> on various CUDA attributes, currently wrapped in
> `__{host|device|global|etc.}__` macros in order to compile the code. Without
> them, `-fsyntax-only` will not be give you correct results on most of the
> CUDA code.
> Most likely you'll see tons of errors when compiler sees `__device__` and has
> no idea what to do with it. Hence my suggestion that clang needs at least a
> minimum subset of CUDA headers to provide the critical subset of macros
> sufficient to convey critical semantics of CUDA code.
Actually this already works (roughly) with the changes made so far.
For example consider the following header:
#pragma once
__global__ void a() {
unsigned block_idx = blockIdx.x;
unsigned thread_idx = threadIdx.x;
__shfl_down_sync(1, 2, 1);
}
When saved as a .cuh and compiled as `clang++ file.cuh -fsyntax-only
--cuda-gpu-arch=sm_75` (using clang++ built from this commit), this works fine
other than an invalid diagnostic for the #pragma once.
Obviously there are several errors when building this with -x c++-header.
The reason why this "works" is because of the change to Driver.cpp. This change
makes it so that -x cuda-header is similar (identical?) to that of -x cuda.
Further changes to Driver.cpp will need to happen to avoid warning about
#pragma once and (if desired) to actually output a precompiled header.
>> A secondary goal was to make it so that header tab completion recognizes
>> .cuh files.
>> This change doesn't depend on the other changes - it only requires a minor
>> edit to clang/lib/Sema/SemaCodeComplete.cpp.
> Maybe. It depends on how well clang can recover from the errors induced by
> the unexpanded CUDA macros. This could range from OK on simple code to rather
> badly if we fail to instantiate templated code.
I am pretty sure that header tab completion is totally unrelated to the
syntactic validity of headers; it's just finding the list of files in the
include path which match the text entered so far and then filtering out files
without a accepted extension. This change just adds ".cuh" to the list of
acceptable extensions.
> We may not need the special CUDA header type after that. The tooling
> heuristic should be able to deal with detecting whether it deals with CUDA
> sources better than clang -- compilation database + user input should help.
> If/when tooling knows it deals with CUDA, it can tell so to clang with -x
> cuda.
This seems like a decent approach to me, but this will result in incorrectly
issuing a diagnostic for #pragma once.
This can of course be fixed by directly disabling the warning, but this does
seem a bit hacky. There may also be other header specific behavior, but I can't
think of any.
On the whole, it does seem a bit gross for tooling to have to compile headers
as though they are main files.
================
Comment at: clang/lib/Driver/Driver.cpp:2466
if (!(IA->getType() == types::TY_CUDA ||
+ IA->getType() == types::TY_CUDAHeader ||
IA->getType() == types::TY_HIP ||
----------------
From my understanding this basically makes the treatment of TY_CUDAHeader
(-xcuda-header) identical to TY_CUDA (-xcuda).
Some changes will need to happen below this for correct handling of header
files (for example not warning about using #pragma once).
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D94337/new/
https://reviews.llvm.org/D94337
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
