[PATCH] D94337: Add cuda header type for cuh files

[Artem Belevich via Phabricator via cfe-commits]

tra added a comment.

In D94337#2491761 <https://reviews.llvm.org/D94337#2491761>, @rgreenblatt wrote:

> 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.

'Works' is not exactly the same as 'works correctly'. This example makes `a()` look like a regular host function, instead of the kernel, and that affects how the rest of the TU get parsed.
I.e. you'll have further errors if somewhere down below the file has `a<<<1,1>>>()`. Similar story with ignoring `__host__` and `__device__` attributes -- that may lead to various overload resolution errors, or reporting conflicting redeclarations/redefinitions for the perfectly valid host/device function overloads. The list goes on.

> 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 handeled similarly (identically?) 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.

CUDA compilation in clang rather heavily depends on compiler understanding CUDA-related attributes and those come from the header files. You can make some things work without those attributes, but you will have more cases where you'll get wrong results ranging from slightly wrong to mostly wrong. :-/ If you're curious, search for CUDA in `lib/Sema/` and you'll see plenty of places where CUDA attributes affect how compiler interprets the sources.

>>> 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.

SGTM. I'm not against the change. It's just the patch does too much (automatically inferring thet .cuh is a CUDA file) and not enough (does not really handle stand-alone CUDA headers in consistent manner).
I'm fine adding CUDA header kind, but without automatic inference of the type via extension. E.g. assigning a `cuda-headers-are-not-really-supported-yet` as an extension would address my concern, while allowing the tools or interested users to use `-x cuda-header` if they need it.

>> 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.

`#pragma once` has been biting me with tooling used on C++ headers, too, so I suspect it's not a CUDA-specific issue.

> 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.

I'm not sure if it should matter much if we only care about `-fsyntax-only`. The warning for `#pragma once` is a rather unusual case where the distinction does matter. Introducing a whole new file kind just to work around the warning looks even more hacky than silencing the warning to me. :-/
In any case, just adding a file kind is simple enough. I'm OK with it as long as we have a way not to expose it to the end-users by default until we can make clang's behavior for that data type sensible.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D94337/new/

https://reviews.llvm.org/D94337