Overall the patch look OK. I'll take a closer look on Monday.

Added '>=' and '>>=' to the list of tokens that may indicate the end of the
empty template argument list.

The patch does not seem to match the description and appears to have nothing to do with rearranging include paths. Could you check if these are the changes you intended to send for review.

Renamed last instance of 'Matches' -> 'PreventedBy'.

Updated assertion message.

• Check that D is non-null

In D52179#1237194, @JonasToth wrote:

Is the condition for this assertion checked beforehand or could this create runtime failures?

Addressed Richard's comments.
Moved clang-tidy changes into separate review D52179.

@rsmith ping.

@jlebar Justin, can you take a look?

In D51663#1224057, @MatzeB wrote:

LGTM, sorry for the breakage.

Updated the changes according to Matthias' comments.

How about D51663?

Your patch above *almost* works, except that each test variant wants to create the same symlink $BUILD/ExternalCUDA/$TEST.teferece_output -> $SRC/External/CUDA/$TEST.reference_output. If symlinks are created at different points in time ninja manages to avoid conflicts, but typically I get one or two attempts to create the symlink launches simultaneously and one of them fails with "symlink already exists". Perhaps the symlink target name should be uniquified in some way.

In D51048#1221432, @MatzeB wrote:

Though thinking about it, I don't really understand your problem:

How can there be multiple outputs with the same name? They have different names in the repository (algorithm.reference_output, assert.reference_output, ...)

Not needed anymore after the reverts in rC341115 and rC341118, right?

I've ran into an unexpected problem after this patches have landed.

Nice. So, in effect, for optimized builds we'll generate pre-DWARF line info only, unless --cuda-noopt-device-debug is specified.
Will this deal with the warnings about back-end being unable to handle particular debug options?

Tests reverted in rL341118.

Reverted in rL341115 & rL341118.

Tests reverted in rL341118

Reverted in rL341115.

Reverted in rL341115

That, and r340967 D51441. I'm running check-clang now and will land reverted changes shortly.

In D50845#1219819, @Hahnfeld wrote:

Ok, the top preprocessor condition for that function is #ifndef __SSE2_MATH__ - the exact same macro that was part of the motivation. Can you please test compiling a simple C file (including math.h) with -mno-sse? My guess would be that this is broken as well.
If yes I'm fine with reverting because I need to teach Clang to allow anonymous unions in type specifiers to make that weird system header work with this patch.

I've sent out D51501. It unbreaks CUDA compilation and keeps OpenMP unchanged.

In D50845#1219733, @Hahnfeld wrote:

In D50845#1219726, @gtbercea wrote:

In general, it looks like this patch leads to some host macros having to be defined again for the auxiliary triple case. It is not clear to me how to exhaustively identify the missing macros, so far it's been just trial and error.

Well, that's the point of this patch, isn't it? Again, the current approach is to just define all macros which is definitely broken.

In our case the headers from a relatively old glibc and compiler errors out on this:

/* This function is used in the `isfinite' macro.  */
__MATH_INLINE int
__NTH (__finite (double __x))
{
  return (__extension__
	  (((((union { double __d; int __i[2]; }) {__d: __x}).__i[1]
	     | 0x800fffffu) + 1) >> 31));
}

FYI. This breaks our CUDA compilation. I haven't figured out what exactly is wrong yet. I may need to unroll the patch if the fix is not obvious.

Looks great. Thank you for updating the docs.

While we're here, perhaps Builder.defineMacro("__linux__") should be changed to DefineStd("linux") which defines linux/__linux/__linux__?

I could not find anything about PLTs in AMDGPU-ABI, nor could I find anything relevant on google.
I still have no idea why PLTs are required in this case. Without that info, the problem may as well be due to unintended requirement for PLT that this patch would hide.

Could you elaborate on what exactly is the problem this patch fixes?
I don't see how internalizing the symbols connects to PLTs. My understanding is that PLTs are used to provide stubs for symbols to be resolved by dynamic linker at runtime. AFAICT AMD does not use shared libs on device side. What do I miss?

Nice. LGTM.

Please keep an eye on CUDA buildbot http://lab.llvm.org:8011/builders/clang-cuda-build.
It runs fair amount of tests with libc++ and handful of libstdc++ versions and may a canary if these changes break something.

I've confirmed that the patch does not break anything in our CUDA code, so it's good to go as far as CUDA is concerned.

__clang_cuda_device_functions.h is not intended to be a device-side math.h, despite having a lot of overlap/similarities. It may change at any time we get new CUDA version.
I would suggest writing an OpenMP-specific replacement for math.h which would map to whatever device-specific function OpenMP needs. For NVPTX that may be libdevice, for which you have declarations in __clang_cuda_libdevice_declares.h. Using part of __clang_cuda_device_functions.h may be a decent starting point for NVPTX, but OpenMP will likely need to provide an equivalent for other back-ends, too.

In D51042#1207920, @bkramer wrote:

In D51042#1207769, @tra wrote:

This is a surprise. PTX ISA does not mention that .ftz is not applicable to cvt.*.f16.* instructions.
Is it only cvt that does not support .ftz or does it impact other instructions? PTX spec has add/sub/mul/fma/set/setp instructions that support f16 and have .ftz variant.

It's only cvt with an explicit rounding mode. I actually ran the output of f16-instructions.ll with FTZ through ptxas and removed instructions until it compiled it. This might even be a bug in ptxas.

This is a surprise. PTX ISA does not mention that .ftz is not applicable to cvt.*.f16.* instructions.
Is it only cvt that does not support .ftz or does it impact other instructions? PTX spec has add/sub/mul/fma/set/setp instructions that support f16 and have .ftz variant.

Talked to @ahatanak over IRC. It appears that this patch may have exposed a preexisting bug.
Apparently delete t; in test/SemaCUDA/call-host-fn-from-device.cu does actually end up calling __host__ operator delete. It should've picked __device__ operator delete, but it does not, so reporting an error here appears to be correct.

In D47757#1204621, @ahatanak wrote:

In D47757#1204561, @tra wrote:

It's a regression. There's a decent chance it breaks someone and this patch, if committed by itself, will end up being rolled back.

Is the regression you are referring to about the static function case? I don't see a difference between ToT clang and my patch in the diagnostics they produce when I compile the following code:

__host__ void f();
static __host__ __device__ void g() { f(); }
__host__ __device__ void g2() { g(); }

Both error out when -fcuda-is-device is provided. If I comment out the definition of g2, it compiles fine.

In D47757#1204545, @ahatanak wrote:

@tra and @rsmith: Can we move forward and fix the incorrect cuda diagnostics in a separate patch?

It appears that libcxx/include/CMakeLists.txt needs to be updated to include bit file into the file set.

In D50845#1203031, @gtbercea wrote:

In D50845#1202991, @hfinkel wrote:

In D50845#1202965, @Hahnfeld wrote:

In D50845#1202963, @hfinkel wrote:

As a result, we should really have a separate header that has those actually-available functions. When targeting NVPTX, why don't we have the included math.h be CUDA's math.h? In the end, those are the functions we need to call when we generate code. Right?

That's what D47849 deals with.

Yes, but it doesn't get CUDA's math.h. Maybe I misunderstand how this works (and I very well might, because it's not clear that CUDA has a math.h by that name), but that patch tries to avoid problems with the host's math.h and then also injects __clang_cuda_device_functions.h into the device compilation. How does this compare to when you include math.h in Clang's CUDA mode? It seems to be that we want to somehow map standard includes, where applicable, to include files in CUDA's include/crt directory (e.g., crt/math_functions.h and crt/common_functions.h for stdio.h for printf), and nothing else ends up being available (because it is, in fact, not available).

There's no CUDA specific math.h unless you want to regard clang_cuda_device_functions.h as a math header.

In D50845#1202551, @ABataev wrote:

In D50845#1202550, @Hahnfeld wrote:

In D50845#1202540, @ABataev wrote:

Maybe for device compilation we also should define __NO_MATH_INLINES and __NO_STRING_INLINES macros to disable inline assembly in glibc?

The problem is that __NO_MATH_INLINES doesn't even avoid all inline assembly from bits/mathinline.h :-( incidentally Clang already defines __NO_MATH_INLINES for x86 (due to an old bug which has been fixed long ago) - and on CentOS we still have problems as described in PR38464.

As a second thought: This might be valid for NVPTX, but I don't think it's a good idea for x86-like offloading targets - they might well profit from inline assembly code.

I'm not saying that we should define those macros for all targets, only for NVPTX. But still, it may disable some inline assembly for other architectures.

Is there a way to control this from clang command line?

In general .relaxed.sys semantics does appear to match guarantees provided by llvm's monotonic ordering, so the patch overall looks like the right thing to do.

Thank you for the analysis. The issue is not a showstopper for us, so I'm OK with waiting for your patches to land.

Added 'explicit' keyword to contructor.

In D50207#1186491, @jlebar wrote:

Just to check, the notion is that it's OK if I report a sm version less than what I end up running on?

It appears that this patch crashes llvm under some circumstances.

We normally do not need to deviate from the host options all that often. I would argue that keeping options identical is a reasonable default for most options.
For some options the driver may be able to derive a sensible value based on the host options. E.g. some options can be ignored. Some can be downgraded. Some can be replaced with a target-specific equivalent.
For others we must require the user to provide the value.

I wonder, what's the right thing to do to silence the warnings. For instance, we compile everything with -Werror and the warnings result in build breaks.

It appears that the issue that originally prompted this change is due to suspected bug in glibc triggered by specific details of our internal build.

Looks OK overall except for the huge if below.

This patch also allows function-scope static const variable without device memory qualifier and emits it as a global variable in constant address space.

Use [] to escape special characters in the regex.
Otherwise I have to use "\\\\" which is way too many characters for my taste.