LGTM

LGTM.

Added a test for the corner case Richard has pointed out in the comments.

Removed unneeded changes.

Addressed Richard's comments.

In D94884#2508000, @JonChesterfield wrote:

I won't object too strongly, as ultimately I don't care about cuda, but I view intertwining the two implementations as technical debt.

LGTM for __clang_cuda_builtin_vars.h.

Presumably, __managed__ variables would have to be memory-mapped into the host address space.

I'd propose splitting the patch into two. One with the addition of CUID and the other that changes the way we havdle static vars.
CUID is useful on its own and is relatively uncontroversial.

Do you have any measurements showing actual impact on performance or memory consumption?

Go ahead.

In D94337#2496627, @rgreenblatt wrote:

... The goal of __clang_cuda_standalone_defs.h is to make it possible to parse CUDA sources at all w/o having to rely on CUDA SDK. ...

Should __clang_cuda_standalone_defs.h depend on the arch? For example, __match_all_sync doesn't exist in sm_35 but does exist in sm_75.
For tooling this might not matter because the host cuda includes don't depend on the arch.

For CUDA we have tests in the test-suite (https://github.com/llvm/llvm-test-suite/tree/main/External/CUDA) and a handful of buildbots running them (e.g. http://lab.llvm.org:8011/#/builders/55).
AMD should probably set up some public build/test bots for HIP, too.
In this case, manual testing and rollback if something breaks is about all we can do.

In D94337#2492108, @rgreenblatt wrote:

In D94337#2491825, @tra wrote:

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

No, this part is definitely working, the full set of SDK headers is included. I have confirmed this by compiling a .cu and a .cuh file with -E and checking that the output is identical. Further, I have confirmed that global<<<_, _>>>() builds, that __host__ only functions can't be used on the device, and that __device__ functions can't be used on the host.

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

In D94337#2491269, @rgreenblatt wrote:

My primary goal for this change was to allow for language servers and other tooling to properly handle cuda header files. From my understanding the way that language servers handle c++ header files is by compiling them with -xc++-header and -fsyntax-only. This is certainly true for ccls and it seems to be true for clangd.
So this can be accomplished without actually able to produce preprocessed output for cuda headers - it only requires handling the "-fsyntax-only" use case.

This adds a cuda header type with file extension "cuh". The output type file
extension is "cuhi" - not sure if this is a good choice. This allows
language servers to properly handle cuh files without additional arguments.

LGTM

In D93638#2483913, @hliao wrote:

Forget that C function could be overloaded on Clang with overloadable
extension. With that, we don't need to mark functions from <ymath.h> as HD.
Instead, we could provide their device-side implementation directly.

LGTM overall.

LGTM.

In D94123#2480848, @jhuber6 wrote:

Okay, so without that flag Clang will not create debug symbols in the PTX assembly output.

There's --cuda-noopt-device-debug option specifically to allow compiling GPU code with full debug info. Clang will generate optimized PTX, but ptxas optimizations will be disabled.

LGTM.

@jdoerfert -- do you happen to know how archives with bitcode are used by OpenMP? Does OpenMP ever link just the archive alone w/o any other bitcode files?

The change look OK. But we should still wait for @jsjodin to confirm that the first-file-is-different for archive files is unintentional.

In D92720#2448890, @yaxunl wrote:

Output of -E for HIP combined host/device compilation is a plain text. It has C++ comments inserted between preprocessor outputs for host and different GPU arch's. The C++ comments follow the format of clang-offload-bundler bundled text files therefore clang-offload-bundler is able to unbundle it.

The patch could use an OK with OMP folks, considering that we've changed the way offload bunder is invoked for OMP.

Use device in the test case.

My first variant of the patch only helped with some cases when the surface/texture attribute type was used.
Trying to reduce real-world failure resulted in an example that I've added as the test case which was still failing with this patch applied.

Found another corner case (reference within a template with the surface/texture
attibute.) and figured out a better fix.

In D92918#2444389, @sdmitriev wrote:

And, by the way, this function is called as a static member (i.e. Linker::linkModules()) everywhere in LLVM sources except llvm-link.cpp:

Reorganized tests for unsupported debug options & dwarf version clamping.

In D92535#2441741, @sdmitriev wrote:

Sure. I believe llvm-link works incorrectly when linking --only-needed symbols from archives with bitcode files. As it is implemented now, llvm-link, when dealing with archives, first links archive modules together into an intermediate module and then tries to link required symbols from that intermediate module into the result. The problem is that archive modules are linked together with --only-needed flag as well, so we always end up with getting empty intermediate module because archive linking starts from scratch (i.e. nothing gets imported into archive module because there are no dependencies).

llvm::Linker::linkModules() is a static member, so there is no need
to pass reference to llvm::Linker instance to loadArFile() function.

In D91281#2441147, @tra wrote:

I think isCUDADeviceBuiltinTextureType has problem handling texture refs within templates.

It appears that we need to add special handling for texture/surface references. Nominally they are host-side objects, but they are accessed/used from device functions as far as Sema is concerned.

The patch description describes what the patch does, but does not tell us much about the problem is is supposed to fix.
Could you give us more details on why the patch is needed?

Adjusted openmp test for the changed -gembed-source warning.

Addressed comments.

Updated to address the comments. PTAL.

-E by default prints preprocessed output to stdout. CUDA will print preprocessed output from all subcompilations. What does HIP do in this case? Printing out the bundle is probably not what the user will expect.
IMO preprocessed output is frequently used as a debugging tool, so it's important for users to be able to read it. Bundled output is rather cumbersome to deal with. It's possible to manually unbundle it, but the tool is not documented well and it's not particularly suitable for human use.

Removed forgotten pop_macro(DEVICE)

@jlebar PTAL. The previous fix did not quite fix the problem. __device__ operator new/delete was still not defined, but the issue was not reported due to deferred diags if the functions were not used.

Better fix for the early inclusion of <new>

Reopening the tracker for a better fix.

Simplified dwarf version clamping.

Updated according to Devid's feedback.

In D92363#2431812, @yaxunl wrote:

In D80450#2428631, @yaxunl wrote:

Also, naming. -ffix-overload-resolution is rather non-specific. I didn't mean to use it literally. The problem is that I can't think of a good descriptive name for what we do here. -fgpu-fix-wrong-side-overloads ? Something else?

How about -fgpu-exclude-wrong-side-overloads? Since what this patch does is always excluding wrong side overloads whereas previously only excluding wrong side overloads if there are same side overloads.

I'd suggest adding more details on the background of this change to the commit log (point to the comment in the isBetterOverloadCandidate ?) and outline the intention to enable the new way to do overloading after some soak time.

While I agree that the default GPU choice is not likely to be correct, or usable, for everyone, but the warning seems to be a half-measure.
If the default is not usable, then it should not be the default. If it's usable, then we don't need a warning.

LGTM in general.