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

[Ryan Greenblatt via Phabricator via cfe-commits]

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