Enumerate all known GPU variants during libdevice detection instead of
hardcoding them.

In D77451#1964971, @sammccall wrote:

NVCC uses different options that should be properly translated

Interested to see how this will work. Is clang itself going to support these args (act compatibly with nvcc, or is the idea that just tools will be?

Please add some more details to the bug description. This change is to make clangd work when compilation database sees CUDA sources compiled with nvcc. NVCC uses different options that should be properly translated. This patch only deals with recognizing the sources as CUDA, but does not handle the compiler options. While not perfect, it's still a useful improvement.

LGTM for NVPTX tests.

In D77240#1957386, @jdoerfert wrote:

I just noticed those as well. I forgot to put the new definitions into the forward declare header. Will do it in a second. The OpenMP math overlay doesn't have one so I forgot :(

We do have a problem. With your patch I see a lot of errors about function redefinitions conflicting with the ones in CUDA's math_functions.hpp:

In D77240#1955724, @jdoerfert wrote:

At least that one is defined in what is "now" __clang_cuda_math.h:

The other macro uses a unqualified lookup already and the qualified one
will cause problems in the OpenMP overlay.

We'll need to make sure that all of these new functions are vailable in all supported CUDA versions.
E.g. acoshf does not seem to be present in CUDA-9.

LGTM for NVPTX changes.

+ @echristo who OK'ed the idea conditional on the actual patch. :-)

In D76365#1947479, @hliao wrote:

Looks OK to me, but I'm not that familiar with the legalizer.

In D76365#1946938, @hliao wrote:

The new revision is accepted, right? Just want to confirm as it seems you accept it before I posted the new change.

In D76948#1946878, @hliao wrote:

I tried that before submitting this one. But, as it's in the closed state, I cannot submit that anymore. I will attach the difference against the previous change somewhere.

Reopened for further work

Would it be possible to update the old review with the new diff? It would make it easier to see the incremental changes you've made. If the old review can be reopened that would be great as it would keep all relevant info in one place, but I'm fine doing the review here, too, if phabricator does not let you do it.

In D76365#1946415, @hliao wrote:

That's a partial template specialization needs handling. I am revising that patch. Please revert it first. Thanks.

In D76365#1946345, @tra wrote:

Looks like the change breaks compilation for us:

In file included from <built-in>:1:
In file included from llvm_unstable/toolchain/lib/clang/google3-trunk/include/__clang_cuda_runtime_wrapper.h:104:
In file included from cuda/include/cuda_runtime.h:116: cuda/include/cuda_surface_types.h:91:42: error: illegal device builtin surface reference type 'surface<void, dim>' declared here
struct  __device_builtin_surface_type__  surface<void, dim> : public surfaceReference
                                         ^
cuda/include/cuda_surface_types.h:91:42: note: 'surface<void, dim>' needs to be instantiated from a class template with the 2nd template argument as an integral value
1 error generated when compiling for sm_60.

I'm investigating, but we may need to roll back this patch. Stay tuned.

Looks like the change breaks compilation for us:

LGTM. Next step is to figure out what various __nv_tex_surf_handler(<string>...) maps to for various strings (there are ~110 of them in CUDA-10.2) and implement its replacement. I think we should be able to do it in the wrapper file.

In D76520#1937294, @yaxunl wrote:

-Xarch_ does not work for passing -cc1 options in the beginning. This patch does not change that.

This requires some further changes about how the options after -Xarch_ are handled. I would suggest to do that in another patch.

In D76520#1936837, @yaxunl wrote:

-Xarch_ works with driver options having value, e.g. -fcf-protection=branch. I added a test for that.

-mframe-pointer=none is a cc1 option. That's why it cannot be passed by -Xarch_. If it is made a driver option it can be passed.

Does it handle options with values? E.g. if I want to pass -mframe-pointer=none ? I vaguely recall the current -Xarch_* implementation had some limitations.
It may be worth adding a test for that.

In D76365#1932439, @hliao wrote:

I believe LLVM does have nvvm.texsurf.handle implemented: https://github.com/llvm/llvm-project/blob/d9972f848294b06807c8764615852ba2bc1e8a74/llvm/include/llvm/IR/IntrinsicsNVVM.td#L1150

Note that, clang-based one needs defining texture fetch functions as they could not be reused from CUDA SDK. That part is enclosed with #if defined(clang).

+1 for refactoring, but what's the long term plan.
Long time ago echristo@ and I had a vague idea to change clang's option parsing to allow something like -Xarch_host <host-only args> -Xarch_device <args for all GPU compilations...> -Xarch=<target> <options for <target> only...>

@hans -- this should be cherry-picked into 10 if it's not too late yet.

In D75811#1923281, @tambre wrote:

Your help here and over on CMake's side has been very helpful. Thank you!
I'll @ you on CMake's side if I need any help while working on CUDA support. Hopefully you won't mind. :)

In D75811#1919368, @tambre wrote:

After some work on my CMake changes, Clang detection as a CUDA compiler works and I can compile CUDA code.

LGTM, but I'm not familiar with the details of clang-cl. I've added Reid who'd have a better idea.

clang-formatted the changes.

Added a test.

This patch should fix it: https://reviews.llvm.org/D76030

In D75817#1913351, @thomasfaingnaert wrote:

I did some more testing, and it appears that that the InstCombine pass (opt -instcombine) is responsible for this.
Maybe it would be better to ensure InstCombine doesn't reorder instead of fixing this in instruction selection as I do now?
I'm not sure if that would influence other backends that rely on this reordering, though.

In D75811#1913148, @tambre wrote:

Magically changing compiler target based on something external to compiler is a bad idea IMO. I would expect a compilation with exactly the same compiler options to do exactly the same thing. If we magically change default target, that will not be the case.

It'd be the same behaviour as NVCC, which compiles for the lowest architecture it supports.

While such reordering may be beneficial, in general, GEP(ASC(x)) is not always equivalent of ASC(GEP(x)), so we can't just blindly do it. I believe this has been discussed few times in the past, though I can't find relevant emails now.

This looks like an elaborate way to achieve the effect of -Wno-unknown_cuda_version. :-)

I'm not sure that's the problem worth solving.

Few nits. LGTM otherwise.

Couple of nits below. LGTM for CUDA headers otherwise.

In D75001#1891861, @kkwli0 wrote:

@tra Will it also include -fopenmp-targets=nvptx64-nvidia-cuda?

In D75001#1891284, @Hahnfeld wrote:

In D75001#1887876, @jdoerfert wrote:

I like this way better. I was hoping we could do it in our cmake only :)

Give others a day or so to comment before your commit but I'm fine with this.

Well, that doesn't really work if openmp-commits is only subscribed on commit. That said, the solution is a bit ugly but I don't have an alternative right now.

Somewhat related, that means Clang issues a warning for every compilation should there be a "unsupported" CUDA version around, even if it's not used? @tra maybe we can only issue the warning if CUDA is going to be used?

Nice! Thank you for making these changes.

In D74571#1876647, @jdoerfert wrote:

That sounds like the right approach for OpenMP. We require a minimal CUDA version, based on what we use internally, but no maximal version if possible since we don't use new features anyway.

In D74571#1875296, @JonChesterfield wrote:

Interesting distinction.

Should compiling without warning indicate comprehensive support, or merely that we ran the tests and they passed?

In D74571#1875002, @kkwli0 wrote:

In D74571#1874947, @JonChesterfield wrote:

Do the in tree tests all pass with the 10.2 toolchain? That's not exactly the same as whether it works but is the closest approximation we have available.

Assuming yes, this patch seems uncontroversial.

Yes, in tree tests pass with 10.2.

It's a bit premature to call CUDA-10.2 supported. We can compile using it, but clang/llvm has no support for the new things introduced by CUDA-10.2.
E.g. CUDA-10.2 introduces new PTX version with new instructions (and matching clang builtins)
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#ptx-isa-version-6-5

Minor test nits. LGTM otherwise.
I'm curious, what does generated ptx for the function look before/after the patch.

Tensorflow folks report that the __try() here does not compile on windows:

ERROR: T:/tmp/nsz6drem/external/llvm-project/llvm/BUILD:3716:1: C++ compilation of rule '@llvm-project//llvm:support' failed (Exit 2)
external/llvm-project/llvm/lib/Support/CrashRecoveryContext.cpp(218): error C2712: Cannot use __try in functions that require object unwinding

LGTM to unblock our build/integration.

LGTM, but wait a bit before landing in case @nhaehnle has concerns.

In D74339#1867894, @MaskRay wrote:

+@tra because of D47396

What types of relocations are from/to .nv_fatbin?

In D73979#1857728, @yaxunl wrote:

All extern shared vars are sharing the same address, however, they may be used as different types in different functions.

For example,

__device__ int foo() {
  extern __shared__ int a;
  for (...) a+=...;
  return a;
}

__device__ double bar(int x) {
  extern __shared__ double b[10];
  for(...) b[x]+=...;
  return b[0];
}

__global__ void k() {
  foo();
  //...
  bar();
}

Thank you for adding the escape hatch option.

In D73979#1857536, @yaxunl wrote:

Based on CUDA usage of extern shared var (https://devblogs.nvidia.com/using-shared-memory-cuda-cc/), it seems CUDA also assumes all extern shared vars have the same address, therefore HIP and CUDA have similar behavior.

A better description for the change would be helpful.

On one hand the change makes sense to me and fits well with what we've done so far.
On the other hand, I worry that this is likely to break things.
We sort of have been implicitly relying on not having the macros related to advanced CPU features enabled on device side which typically results in device-side compilation seeing a more portable/simpler version of the code.
Defining the macros that enable more host-side CPU-specific code may trigger interesting compatibility features. Postponed diagnostics combined with ignore (some) errors in the wrong-side-only code should probably deal with most of them, but I suspect we'll see new interesting failure cases. We would not know until we try.

LGTM for CUDA.

@hans : that's another candidate for 10.x cherry-pick, if you're OK with it.

Use std::string instead of Twine which can't be stored.