[OpenMP][WIP] Initial support for `begin/end declare variant`

[jdoerfert]

Summary

NOTE: This is a WIP patch to foster a discussion. Please do consider that when browsing the code. Details will be discussed in individual commits once we agreed on the overall model. This is also the reason why test coverage, documentation, TODOs, etc. is lacking.

This patch provides initial support for omp begin/end declare variant,
as defined in OpenMP technical report 8 (TR8).

A major purpose of this patch is to provide proper math.h/cmath support
for OpenMP target offloading. See PR42061, PR42798, PR42799.
The three tests included in this patch show that these bugs (should be)
fixed in this scheme.

In contrast to the declare variant handling we already have, this patch
makes use of the multi-version handling in clang. This is especially
useful as the variants have the same name as the base functions. We
should try to port all OpenMP variant handling to this scheme, see the
TODO in CodeGenModule for a proposed way towards this goal. Other than
that, we tried to reuse the existing multi-version and OpenMP variant
handling as much as possible.

NOTE: There are various TODOs that need fixing and switches that need additional cases.

[merge_guards_bot]

Build result: FAILURE -
Log files: console-log.txt, CMakeCache.txt

[jdoerfert]

clang/lib/AST/StmtOpenMP.cpp

2243

This code was basically only moved, not written for this patch. It needs to life somewhere accessible from Parser to CodeGen, see the TODOs below.

clang/lib/Headers/__clang_cuda_cmath.h

38

NOTE: It might be cleaner to revert the patches that put the OpenMP handling code here first.

70

As far as I can tell, fpclassify is not available in CUDA so it is unclear if we want to have it here or not. I removed it due to the TODO above. Consequently I also had to remove other fpclassify occurrences. If it turns out the host version is not usable on the device and we need the builtins, we add them back but under the opposite guard, that is #ifdef _OPENMP.

[jdoerfert]

clang/lib/Parse/ParseOpenMP.cpp

1048

The diff is confusing here. I actually extracted some code into a helper function (ParseOMPDeclareVariantMatchClause on the right) which I can reuse in the begin/end handling. The code "deleted" here is below ParseOMPDeclareVariantMatchClause on the right.

[jdoerfert]

Add (missing) include. (Worked locally just fine).

[merge_guards_bot]

Build result: FAILURE -
Log files: console-log.txt, CMakeCache.txt

[ABataev]

I read the spec and don't think that we need all this complex stuff for the implementation. Yiu need judt to check at the codegen phase if the function must be emitted or not. We don't even need to move context checksnfrom codegen, because with the current semantics all the checkscan be safely performed at the codegen phase.

[jdoerfert]

In D71179#1774444, @ABataev wrote:

I read the spec and don't think that we need all this complex stuff for the implementation. Yiu need judt to check at the codegen phase if the function must be emitted or not. We don't even need to move context checksnfrom codegen, because with the current semantics all the checkscan be safely performed at the codegen phase.

For better or worse we need this and it is actually a natural reuse of the multi-versioning code. We need this because:

1. For the begin/end version we cannot even parse anything in a context that does not match at encounter time, e.g. the kind(fpga) context in clang/test/AST/ast-dump-openmp-begin-declare-variant.c.
2. For the 5.0 version we cannot use the replaceAllUses approach currently implemented in tryEmitDeclareVariant as soon as we have the construct context selector trait. That means we will have to resolve the call target earlier anyway.

(FWIW, I wrote this part of the SPEC.)

[ABataev]

In D71179#1774448, @jdoerfert wrote:

In D71179#1774444, @ABataev wrote:

I read the spec and don't think that we need all this complex stuff for the implementation. Yiu need judt to check at the codegen phase if the function must be emitted or not. We don't even need to move context checksnfrom codegen, because with the current semantics all the checkscan be safely performed at the codegen phase.

For better or worse we need this and it is actually a natural reuse of the multi-versioning code. We need this because:

1. For the begin/end version we cannot even parse anything in a context that does not match at encounter time, e.g. the kind(fpga) context in clang/test/AST/ast-dump-openmp-begin-declare-variant.c.

Ok, here we can check the context and just skip everything between begin/end pragmas just like if something like #ifdef...#endif is seen.

2. For the 5.0 version we cannot use the replaceAllUses approach currently implemented in tryEmitDeclareVariant as soon as we have the construct context selector trait. That means we will have to resolve the call target earlier anyway.

I thought about this. Here we need to use a little bit different method, but again everything can be reolved at the codegen phase, no need to resolve it at parsing/sema. Plus, this is completely different problem and must be solved in the different patch.

(FWIW, I wrote this part of the SPEC.)

[jdoerfert]

In D71179#1774469, @ABataev wrote:

In D71179#1774448, @jdoerfert wrote:

In D71179#1774444, @ABataev wrote:

I read the spec and don't think that we need all this complex stuff for the implementation. Yiu need judt to check at the codegen phase if the function must be emitted or not. We don't even need to move context checksnfrom codegen, because with the current semantics all the checkscan be safely performed at the codegen phase.

For better or worse we need this and it is actually a natural reuse of the multi-versioning code. We need this because:

1. For the begin/end version we cannot even parse anything in a context that does not match at encounter time, e.g. the kind(fpga) context in clang/test/AST/ast-dump-openmp-begin-declare-variant.c.

Ok, here we can check the context and just skip everything between begin/end pragmas just like if something like #ifdef...#endif is seen.

Agreed.

2. For the 5.0 version we cannot use the replaceAllUses approach currently implemented in tryEmitDeclareVariant as soon as we have the construct context selector trait. That means we will have to resolve the call target earlier anyway.

I thought about this. Here we need to use a little bit different method, but again everything can be reolved at the codegen phase, no need to resolve it at parsing/sema.

I doubt we can, yet alone want to do (basically) overload resolution during codegen.
Depending on what function we resolve to, we get different instantiations which require everything from semantic analysis to run on that code again, right?
Could you elaborate why we should not do all the overload resolution at the same time and with the same mechanism that is already present? I mean, SemaOverload deals with multi-versioning already.

Plus, this is completely different problem and must be solved in the different patch.

As I noted in the very beginning of the commit message, this is not supposed to be a commited like this but split into multiple patches. Let's not mix discussions here.

[ABataev]

In D71179#1774470, @jdoerfert wrote:

In D71179#1774469, @ABataev wrote:

In D71179#1774448, @jdoerfert wrote:

In D71179#1774444, @ABataev wrote:

I read the spec and don't think that we need all this complex stuff for the implementation. Yiu need judt to check at the codegen phase if the function must be emitted or not. We don't even need to move context checksnfrom codegen, because with the current semantics all the checkscan be safely performed at the codegen phase.

For better or worse we need this and it is actually a natural reuse of the multi-versioning code. We need this because:

1. For the begin/end version we cannot even parse anything in a context that does not match at encounter time, e.g. the kind(fpga) context in clang/test/AST/ast-dump-openmp-begin-declare-variant.c.

Ok, here we can check the context and just skip everything between begin/end pragmas just like if something like #ifdef...#endif is seen.

Agreed.

2. For the 5.0 version we cannot use the replaceAllUses approach currently implemented in tryEmitDeclareVariant as soon as we have the construct context selector trait. That means we will have to resolve the call target earlier anyway.

I thought about this. Here we need to use a little bit different method, but again everything can be reolved at the codegen phase, no need to resolve it at parsing/sema.

I doubt we can, yet alone want to do (basically) overload resolution during codegen.
Depending on what function we resolve to, we get different instantiations which require everything from semantic analysis to run on that code again, right?
Could you elaborate why we should not do all the overload resolution at the same time and with the same mechanism that is already present? I mean, SemaOverload deals with multi-versioning already.

They do this because they have several function definitions with the same name. In our case, we have several different functions with different names and for us no need to worry about overloading resolution, the compiler will do everything for us.

Plus, this is completely different problem and must be solved in the different patch.

As I noted in the very beginning of the commit message, this is not supposed to be a commited like this but split into multiple patches. Let's not mix discussions here.

[jdoerfert]

In D71179#1774471, @ABataev wrote:

They do this because they have several function definitions with the same name. In our case, we have several different functions with different names and for us no need to worry about overloading resolution, the compiler will do everything for us.

I think we talk past each other again. This is the implementation of omp begin/end declare variant as described in TR8. Bt definition, the new variant mechanism will result in several different function definitions with the same name. See the two tests for examples.

[hfinkel]

clang/lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h

17

Should we use a more-specific selector and then get rid of this __NVPTX__ check?

clang/lib/Parse/ParseOpenMP.cpp

1489

Will this just inf-loop if the file ends?

[hfinkel]

In D71179#1774487, @jdoerfert wrote:

In D71179#1774471, @ABataev wrote:

They do this because they have several function definitions with the same name. In our case, we have several different functions with different names and for us no need to worry about overloading resolution, the compiler will do everything for us.

I think we talk past each other again. This is the implementation of omp begin/end declare variant as described in TR8. Bt definition, the new variant mechanism will result in several different function definitions with the same name. See the two tests for examples.

The intent of this feature is to allow us to include the device-function headers and the system headers simultaneously, giving preference to the device functions when compiling for the device, thus fixing a number of outstanding math.h OpenMP offloading problems. This definitely means that we'll have multiple functions with the same name and we need to pick the right ones during overload resolution.

@jdoerfert , how does the ".ompvariant" work with external functions? I see the part of the spec which says, "The symbol name of a function definition that appears between a begin declare variant...", but, if we append this name to, for example, the names of functions present in the device math library, won't we have a problem with linking?

[JonChesterfield]

Great to see the fragile math.h stuff disappear.

I'm not sure about the CPU/GPU/other granularity. An openmp program with x86 as the host and target offload regions for amdgcn and for nvptx seems like a reasonable aspiration. Or for a couple of different generations from the same vendor.

More ambitiously, one might want a GPU to be the host, and offload kernels for I/O to an aarch64 "target".

We don't need to wire such combinations in up front, and I don't think they're excluded by this design. A future 'x86-64' variant would presumably be chosen over a 'cpu' variant when compiling for x86-64.

[ABataev]

In D71179#1774487, @jdoerfert wrote:

In D71179#1774471, @ABataev wrote:

They do this because they have several function definitions with the same name. In our case, we have several different functions with different names and for us no need to worry about overloading resolution, the compiler will do everything for us.

I think we talk past each other again. This is the implementation of omp begin/end declare variant as described in TR8. Bt definition, the new variant mechanism will result in several different function definitions with the same name. See the two tests for examples.

I just don't get it. If begin/end is just a something like #ifdef...endif, why you just can't skip everything between begin/end if the context does not match?

[JonChesterfield]

clang/test/OpenMP/begin_declare_variant_codegen.cpp

72

The name mangling should probably append the device kind, .e.g. _Z3foov.ompvariant.gpu

[JonChesterfield]

clang/lib/Headers/__clang_cuda_cmath.h

70

We could call __builtin_fpclassify for nvptx, e.g. from https://github.com/ROCm-Developer-Tools/aomp-extras/blob/0.7-6/aomp-device-libs/libm/src/libm-nvptx.cpp

int fpclassify(float __x) {
  return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL, FP_ZERO, __x);
}
int fpclassify(double __x) {
  return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL, FP_ZERO, __x);
}

[hfinkel]

In D71179#1774678, @ABataev wrote:

In D71179#1774487, @jdoerfert wrote:

In D71179#1774471, @ABataev wrote:

They do this because they have several function definitions with the same name. In our case, we have several different functions with different names and for us no need to worry about overloading resolution, the compiler will do everything for us.

I think we talk past each other again. This is the implementation of omp begin/end declare variant as described in TR8. Bt definition, the new variant mechanism will result in several different function definitions with the same name. See the two tests for examples.

I just don't get it. If begin/end is just a something like #ifdef...endif, why you just can't skip everything between begin/end if the context does not match?

The patch does this (see in ParseOpenMP.cpp where I asked about the potential inf-loop). But when the definitions are not skipped, then we have to worry about having multiple decls/defs of the same name and the overload priorities.

[ABataev]

In D71179#1775066, @hfinkel wrote:

In D71179#1774678, @ABataev wrote:

In D71179#1774487, @jdoerfert wrote:

In D71179#1774471, @ABataev wrote:

They do this because they have several function definitions with the same name. In our case, we have several different functions with different names and for us no need to worry about overloading resolution, the compiler will do everything for us.

I think we talk past each other again. This is the implementation of omp begin/end declare variant as described in TR8. Bt definition, the new variant mechanism will result in several different function definitions with the same name. See the two tests for examples.

I just don't get it. If begin/end is just a something like #ifdef...endif, why you just can't skip everything between begin/end if the context does not match?

The patch does this (see in ParseOpenMP.cpp where I asked about the potential inf-loop). But when the definitions are not skipped, then we have to worry about having multiple decls/defs of the same name and the overload priorities.

I would recommend to drop all this extra stuff from the patch and focus on the initial patch. We'll need something similar to multiversion in case of the construct context selectors, but at first we need to solve all the problems with the simple versions of the construct rather that try to solve all the problems in the world in one patch. It is almost impossible to review.

clang/lib/AST/StmtOpenMP.cpp

2243

I don't think this is the right place for this code. Will try to move it to Basic directory in my patch.

clang/lib/Parse/ParseOpenMP.cpp

1489

It will.

[jdoerfert]

clang/lib/AST/StmtOpenMP.cpp

2243

Sure. As noted in the TODOs, finding a place for this is needed.

clang/lib/Headers/__clang_cuda_cmath.h

70

Agreed. Assuming it works, I'll put the fpclassify code back in but only remove the todo and OPENMP macro.

clang/lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h

17

For now, this is CUDA after all. I was going to revisit this once we know how the AMD solution looks (I guess via HIP).
That said, I'd put a pin on it for now.

(The kind(gpu) selector below is only because we don't have anything more specific and it matches all our one GPU targets for now.)

clang/lib/Parse/ParseOpenMP.cpp

1489

We'll add a check and test.

clang/test/OpenMP/begin_declare_variant_codegen.cpp

72

There is already a TODO for that (I think CodeGenModule). Mangling right now is hardcoded and needs to be revisited :)

[jdoerfert]

@jdoerfert , how does the ".ompvariant" work with external functions? I see the part of the spec which says, "The symbol name of a function definition that appears between a begin declare variant...", but, if we append this name to, for example, the names of functions present in the device math library, won't we have a problem with linking?

We restricted it for now to function definitions so we don't need to define the mangling as you cannot expect linking. (I did this to get it in TR8 while I figured it will solve all our math.h problems already).
However, we need to avoid collisions with user code, e.g., through the use of symbols in the name that are not allowed to be used by the user (I thought "." is one of them).

In D71179#1774639, @JonChesterfield wrote:

Great to see the fragile math.h stuff disappear.

I'm not sure about the CPU/GPU/other granularity. An openmp program with x86 as the host and target offload regions for amdgcn and for nvptx seems like a reasonable aspiration. Or for a couple of different generations from the same vendor.

More ambitiously, one might want a GPU to be the host, and offload kernels for I/O to an aarch64 "target".

We don't need to wire such combinations in up front, and I don't think they're excluded by this design. A future 'x86-64' variant would presumably be chosen over a 'cpu' variant when compiling for x86-64.

As I wrote in the inline comment somewhere, kind(gpu) is an artifact due to missing fine-grained context selectors. If that wasn't the core of your issue, please elaborate.

In D71179#1775157, @ABataev wrote:

In D71179#1775066, @hfinkel wrote:

In D71179#1774678, @ABataev wrote:

In D71179#1774487, @jdoerfert wrote:

In D71179#1774471, @ABataev wrote:

They do this because they have several function definitions with the same name. In our case, we have several different functions with different names and for us no need to worry about overloading resolution, the compiler will do everything for us.

I think we talk past each other again. This is the implementation of omp begin/end declare variant as described in TR8. Bt definition, the new variant mechanism will result in several different function definitions with the same name. See the two tests for examples.

I just don't get it. If begin/end is just a something like #ifdef...endif, why you just can't skip everything between begin/end if the context does not match?

The patch does this (see in ParseOpenMP.cpp where I asked about the potential inf-loop). But when the definitions are not skipped, then we have to worry about having multiple decls/defs of the same name and the overload priorities.

I would recommend to drop all this extra stuff from the patch and focus on the initial patch. We'll need something similar to multiversion in case of the construct context selectors, but at first we need to solve all the problems with the simple versions of the construct rather that try to solve all the problems in the world in one patch. It is almost impossible to review.

I agree with you to the point that this is not supposed to be reviewed. That's why I wrote that in the commit message. I did this so we can make sure the general path is clear and people (myself included) can see how/that it works.
I also agree that construct context selectors are very close to multi-versioned functions. That is why I said earlier we should move all variant handling into this scheme.

My plan:

• We play around with this prototype now, make sure there are no major problems with it (so far it didn't seem so).
• We split it up (This doesn't necessarily need to be only done by me, as that often slows down these processes).
• We review the parts with proper test coverage, etc. and get it in.

[hfinkel]

In D71179#1775157, @ABataev wrote:

In D71179#1775066, @hfinkel wrote:

In D71179#1774678, @ABataev wrote:

In D71179#1774487, @jdoerfert wrote:

In D71179#1774471, @ABataev wrote:

They do this because they have several function definitions with the same name. In our case, we have several different functions with different names and for us no need to worry about overloading resolution, the compiler will do everything for us.

I think we talk past each other again. This is the implementation of omp begin/end declare variant as described in TR8. Bt definition, the new variant mechanism will result in several different function definitions with the same name. See the two tests for examples.

I just don't get it. If begin/end is just a something like #ifdef...endif, why you just can't skip everything between begin/end if the context does not match?

The patch does this (see in ParseOpenMP.cpp where I asked about the potential inf-loop). But when the definitions are not skipped, then we have to worry about having multiple decls/defs of the same name and the overload priorities.

I would recommend to drop all this extra stuff from the patch and focus on the initial patch. We'll need something similar to multiversion in case of the construct context selectors, but at first we need to solve all the problems with the simple versions of the construct rather that try to solve all the problems in the world in one patch. It is almost impossible to review.

I agree. We should split this into several patches (e.g., basic handling, skipping parsing for incompatible selectors, overload things). I think that @jdoerfert posted this so that people can see the high-level direction and provide feedback (including feedback on how to stage the functionality for review).

@jdoerfert , also, do we have tests that can go into the test suite / libomptarget regression tests demonstrating the collection of problems people have currently opened bugs on regarding math.h? I recall we still had problems with host code needing the long-double overloads, with constants from the system headers, etc.

[ABataev]

In D71179#1775442, @jdoerfert wrote:

@jdoerfert , how does the ".ompvariant" work with external functions? I see the part of the spec which says, "The symbol name of a function definition that appears between a begin declare variant...", but, if we append this name to, for example, the names of functions present in the device math library, won't we have a problem with linking?

We restricted it for now to function definitions so we don't need to define the mangling as you cannot expect linking. (I did this to get it in TR8 while I figured it will solve all our math.h problems already).
However, we need to avoid collisions with user code, e.g., through the use of symbols in the name that are not allowed to be used by the user (I thought "." is one of them).

In D71179#1774639, @JonChesterfield wrote:

Great to see the fragile math.h stuff disappear.

I'm not sure about the CPU/GPU/other granularity. An openmp program with x86 as the host and target offload regions for amdgcn and for nvptx seems like a reasonable aspiration. Or for a couple of different generations from the same vendor.

More ambitiously, one might want a GPU to be the host, and offload kernels for I/O to an aarch64 "target".

We don't need to wire such combinations in up front, and I don't think they're excluded by this design. A future 'x86-64' variant would presumably be chosen over a 'cpu' variant when compiling for x86-64.

As I wrote in the inline comment somewhere, kind(gpu) is an artifact due to missing fine-grained context selectors. If that wasn't the core of your issue, please elaborate.

In D71179#1775157, @ABataev wrote:

In D71179#1775066, @hfinkel wrote:

In D71179#1774678, @ABataev wrote:

In D71179#1774487, @jdoerfert wrote:

In D71179#1774471, @ABataev wrote:

They do this because they have several function definitions with the same name. In our case, we have several different functions with different names and for us no need to worry about overloading resolution, the compiler will do everything for us.

I think we talk past each other again. This is the implementation of omp begin/end declare variant as described in TR8. Bt definition, the new variant mechanism will result in several different function definitions with the same name. See the two tests for examples.

I just don't get it. If begin/end is just a something like #ifdef...endif, why you just can't skip everything between begin/end if the context does not match?

The patch does this (see in ParseOpenMP.cpp where I asked about the potential inf-loop). But when the definitions are not skipped, then we have to worry about having multiple decls/defs of the same name and the overload priorities.

I would recommend to drop all this extra stuff from the patch and focus on the initial patch. We'll need something similar to multiversion in case of the construct context selectors, but at first we need to solve all the problems with the simple versions of the construct rather that try to solve all the problems in the world in one patch. It is almost impossible to review.

I agree with you to the point that this is not supposed to be reviewed. That's why I wrote that in the commit message. I did this so we can make sure the general path is clear and people (myself included) can see how/that it works.
I also agree that construct context selectors are very close to multi-versioned functions. That is why I said earlier we should move all variant handling into this scheme.

I don't think we should do this. Something similar to multiversioning is required only for a small subset. Everything else can be implemented in a more straightforward and simple way. Plus, I'm not sure that we'll need full reuse of the multiversioning. Seems to me, we can implement codegen in a different way. Multiversioning is supported only by x86 in clang/LLVM. I think we can try to implement a more portable and universal scheme.

My plan:

• We play around with this prototype now, make sure there are no major problems with it (so far it didn't seem so).
• We split it up (This doesn't necessarily need to be only done by me, as that often slows down these processes).
• We review the parts with proper test coverage, etc. and get it in.

[hfinkel]

In D71179#1775442, @jdoerfert wrote:

@jdoerfert , how does the ".ompvariant" work with external functions? I see the part of the spec which says, "The symbol name of a function definition that appears between a begin declare variant...", but, if we append this name to, for example, the names of functions present in the device math library, won't we have a problem with linking?

We restricted it for now to function definitions so we don't need to define the mangling as you cannot expect linking. (I did this to get it in TR8 while I figured it will solve all our math.h problems already).
However, we need to avoid collisions with user code, e.g., through the use of symbols in the name that are not allowed to be used by the user (I thought "." is one of them).

Okay, but how to we distinguish functions for which there is a declaration and we need the mangling because the user has provided a definition elsewhere, from those for which there is a declaration, and we don't want mangling because we need to link to some system library?

[jdoerfert]

Add one more test sin(long double), and fix some rebase issues

[jdoerfert]

@jdoerfert , also, do we have tests that can go into the test suite / libomptarget regression tests demonstrating the collection of problems people have currently opened bugs on regarding math.h? I recall we still had problems with host code needing the long-double overloads, with constants from the system headers, etc.

The three tests I have in here show already that almost all of the known problems are solved by this (e.g. constants from the system headers). The rest can be easily added as lit test. The test suite situation is evolving but far from me being resolved. I would prefer not to mix these discussions and focus on lit tests with this patch (once split).

We restricted it for now to function definitions so we don't need to define the mangling as you cannot expect linking. (I did this to get it in TR8 while I figured it will solve all our math.h problems already).
However, we need to avoid collisions with user code, e.g., through the use of symbols in the name that are not allowed to be used by the user (I thought "." is one of them).

Okay, but how to we distinguish functions for which there is a declaration and we need the mangling because the user has provided a definition elsewhere, from those for which there is a declaration, and we don't want mangling because we need to link to some system library?

The idea is, declarations inside begin/end declare variant are supposed to be not affected by the begin/end declare variant. That is, if you have declarations you cannot expect variant multi-versioning to happen. Having declarations inside or outside the begin/end declare variant is still fine if they all denote the same function.

I don't think we should do this. Something similar to multiversioning is required only for a small subset.

This is neither true, nor relevant. It is not true because OpenMP 5.0 declare variant is so broken it cannot be used for what it was intended for. That means people (as for example we for math) will inevitably use begin/end declare variant.

Everything else can be implemented in a more straightforward and simple way.

Having a single scheme is arguably simpler than maintaining multiple schemes. There is no additional overhead in using the more powerful and available multi-version scheme for everything.

Plus, I'm not sure that we'll need full reuse of the multiversioning. Seems to me, we can implement codegen in a different way.

Please provide actual details with statements like this. It is impossible to tell what you mean.

Multiversioning is supported only by x86 in clang/LLVM. I think we can try to implement a more portable and universal scheme.

This is not true, at least not from a conceptual standpoint. While cpu_supports and cpu_is multi-versioning is restricted to X86, see supportsMultiVersioning in TargetInfo.h, the new kind of OpenMP multi-versioning is a portable and universal scheme (see the uses of supportsMultiVersioning)

[hfinkel]

In D71179#1775687, @jdoerfert wrote:

...

We restricted it for now to function definitions so we don't need to define the mangling as you cannot expect linking. (I did this to get it in TR8 while I figured it will solve all our math.h problems already).
However, we need to avoid collisions with user code, e.g., through the use of symbols in the name that are not allowed to be used by the user (I thought "." is one of them).

Okay, but how to we distinguish functions for which there is a declaration and we need the mangling because the user has provided a definition elsewhere, from those for which there is a declaration, and we don't want mangling because we need to link to some system library?

The idea is, declarations inside begin/end declare variant are supposed to be not affected by the begin/end declare variant. That is, if you have declarations you cannot expect variant multi-versioning to happen. Having declarations inside or outside the begin/end declare variant is still fine if they all denote the same function.

Thanks, now I understand. This seems like it will work.

[merge_guards_bot]

Build result: fail - 60637 tests passed, 24 failed and 726 were skipped.

failed: Clang.CXX/dcl_dcl/basic_namespace/namespace_udecl/p11.cpp
failed: Clang.CXX/drs/dr5xx.cpp
failed: Clang.CXX/modules-ts/basic/basic_def_odr/p6/global-vs-module.cpp
failed: Clang.CXX/special/class_inhctor/p3.cpp
failed: Clang.Headers/nvptx_device_cmath_functions.c
failed: Clang.Headers/nvptx_device_cmath_functions.cpp
failed: Clang.Headers/nvptx_device_cmath_functions_cxx17.cpp
failed: Clang.Headers/nvptx_device_math_functions.c
failed: Clang.Headers/nvptx_device_math_functions.cpp
failed: Clang.Headers/nvptx_device_math_functions_cxx17.cpp
failed: Clang.OpenMP/declare_variant_ast_print.cpp
failed: Clang.OpenMP/declare_variant_device_kind_codegen.cpp
failed: Clang.OpenMP/declare_variant_implementation_vendor_codegen.cpp
failed: Clang.OpenMP/declare_variant_messages.c
failed: Clang.OpenMP/declare_variant_messages.cpp
failed: Clang.OpenMP/declare_variant_mixed_codegen.cpp
failed: Clang.OpenMP/math_codegen.cpp
failed: Clang.OpenMP/math_fp_macro.cpp
failed: Clang.OpenMP/nvptx_declare_variant_device_kind_codegen.cpp
failed: Clang.OpenMP/nvptx_declare_variant_implementation_vendor_codegen.cpp
failed: Clang.SemaCXX/attr-cpuspecific.cpp
failed: Clang.SemaCXX/attr-target-mv.cpp
failed: Clang.SemaCXX/friend.cpp
failed: Clang.SemaCXX/using-decl-1.cpp

Log files: console-log.txt, CMakeCache.txt

[tra]

clang/lib/Headers/__clang_cuda_cmath.h

72

Please keep fpclassify in place. It's been available in this header for a long time and it *is* needed.

462

I think only #ifdef should be removed here. scalblnf itself should remain.

clang/lib/Headers/__clang_cuda_device_functions.h

1724

Ditto here. Only preprocessor statements should be removed.

[jdoerfert]

minor fix

[merge_guards_bot]

Build result: fail - 60641 tests passed, 20 failed and 726 were skipped.

failed: Clang.CXX/dcl_dcl/basic_namespace/namespace_udecl/p11.cpp
failed: Clang.CXX/drs/dr5xx.cpp
failed: Clang.CXX/modules-ts/basic/basic_def_odr/p6/global-vs-module.cpp
failed: Clang.CXX/special/class_inhctor/p3.cpp
failed: Clang.Headers/nvptx_device_cmath_functions.c
failed: Clang.Headers/nvptx_device_math_functions.c
failed: Clang.OpenMP/declare_variant_ast_print.cpp
failed: Clang.OpenMP/declare_variant_device_kind_codegen.cpp
failed: Clang.OpenMP/declare_variant_implementation_vendor_codegen.cpp
failed: Clang.OpenMP/declare_variant_messages.c
failed: Clang.OpenMP/declare_variant_messages.cpp
failed: Clang.OpenMP/declare_variant_mixed_codegen.cpp
failed: Clang.OpenMP/math_codegen.cpp
failed: Clang.OpenMP/math_fp_macro.cpp
failed: Clang.OpenMP/nvptx_declare_variant_device_kind_codegen.cpp
failed: Clang.OpenMP/nvptx_declare_variant_implementation_vendor_codegen.cpp
failed: Clang.SemaCXX/attr-cpuspecific.cpp
failed: Clang.SemaCXX/attr-target-mv.cpp
failed: Clang.SemaCXX/friend.cpp
failed: Clang.SemaCXX/using-decl-1.cpp

Log files: console-log.txt, CMakeCache.txt

[ABataev]

In D71179#1775687, @jdoerfert wrote:

@jdoerfert , also, do we have tests that can go into the test suite / libomptarget regression tests demonstrating the collection of problems people have currently opened bugs on regarding math.h? I recall we still had problems with host code needing the long-double overloads, with constants from the system headers, etc.

The three tests I have in here show already that almost all of the known problems are solved by this (e.g. constants from the system headers). The rest can be easily added as lit test. The test suite situation is evolving but far from me being resolved. I would prefer not to mix these discussions and focus on lit tests with this patch (once split).

We restricted it for now to function definitions so we don't need to define the mangling as you cannot expect linking. (I did this to get it in TR8 while I figured it will solve all our math.h problems already).
However, we need to avoid collisions with user code, e.g., through the use of symbols in the name that are not allowed to be used by the user (I thought "." is one of them).

Okay, but how to we distinguish functions for which there is a declaration and we need the mangling because the user has provided a definition elsewhere, from those for which there is a declaration, and we don't want mangling because we need to link to some system library?

The idea is, declarations inside begin/end declare variant are supposed to be not affected by the begin/end declare variant. That is, if you have declarations you cannot expect variant multi-versioning to happen. Having declarations inside or outside the begin/end declare variant is still fine if they all denote the same function.

I don't think we should do this. Something similar to multiversioning is required only for a small subset.

This is neither true, nor relevant. It is not true because OpenMP 5.0 declare variant is so broken it cannot be used for what it was intended for. That means people (as for example we for math) will inevitably use begin/end declare variant.

I rather doubt that it is so much broken. The fact, that you need some new construct to express some functionality does not mean that the previous one is incorrect. It is incomplete, maybe. But not broken. And even for begin/end stuff, multiversioning is only required for construct traits, for all other traits we can reuse the existing implementation.

Everything else can be implemented in a more straightforward and simple way.

Having a single scheme is arguably simpler than maintaining multiple schemes. There is no additional overhead in using the more powerful and available multi-version scheme for everything.

Plus, I'm not sure that we'll need full reuse of the multiversioning. Seems to me, we can implement codegen in a different way.

Please provide actual details with statements like this. It is impossible to tell what you mean.

Multiversioning is supported only by x86 in clang/LLVM. I think we can try to implement a more portable and universal scheme.

This is not true, at least not from a conceptual standpoint. While cpu_supports and cpu_is multi-versioning is restricted to X86, see supportsMultiVersioning in TargetInfo.h, the new kind of OpenMP multi-versioning is a portable and universal scheme (see the uses of supportsMultiVersioning)

[jdoerfert]

Undo math function removal (fpclassify & scalblnf), reorder includes (host
first) The latter is the "natural way" but also necessary because fpclassify
uses macros and we did not copy the complex cuda_runtime_wrapper include magic.
However, the sin(long double) is back if it is called in a function that has
a target region. This is an artifact unrelated to any of this (I would argue).
The problem is that we parse + type check *host* code surrounding the target
region when we compile for the target. This has various down sites and can
easily break without math involvement. Long story short, we need to fix this
later separately.

clang/lib/Headers/__clang_cuda_cmath.h

72

Done.

462

I misinterpreted the TODOs, here and above. That is why I removed code. Sorry for the noise.

clang/lib/Headers/__clang_cuda_device_functions.h

1724

Yeah, my bad.

[jdoerfert]

This is neither true, nor relevant. It is not true because OpenMP 5.0 declare variant is so broken it cannot be used for what it was intended for. That means people (as for example we for math) will inevitably use begin/end declare variant.

I rather doubt that it is so much broken. The fact, that you need some new construct to express some functionality does not mean that the previous one is incorrect. It is incomplete, maybe. But not broken.

Broken in the sense that we (in the OpenMP accelerator subcommittee) don't think it can be used for what we envisioned it initially. It can be used for certain things though.

And even for begin/end stuff, multiversioning is only required for construct traits, for all other traits we can reuse the existing implementation.

Again, this is not the case. begin/end *always* caused multiple definitions with the same name. Even if we ignore that for a second, why should we not use the powerful infrastructure we have (=multi-versioning) that supports construct traits and not use it for the other traits? Or asked differently, why should we have a second codegen rewriting scheme?

[merge_guards_bot]

Build result: fail - 60639 tests passed, 24 failed and 726 were skipped.

failed: Clang.CXX/dcl_dcl/basic_namespace/namespace_udecl/p11.cpp
failed: Clang.CXX/drs/dr5xx.cpp
failed: Clang.CXX/modules-ts/basic/basic_def_odr/p6/global-vs-module.cpp
failed: Clang.CXX/special/class_inhctor/p3.cpp
failed: Clang.Headers/nvptx_device_cmath_functions.c
failed: Clang.Headers/nvptx_device_cmath_functions.cpp
failed: Clang.Headers/nvptx_device_cmath_functions_cxx17.cpp
failed: Clang.Headers/nvptx_device_math_functions.c
failed: Clang.Headers/nvptx_device_math_functions.cpp
failed: Clang.Headers/nvptx_device_math_functions_cxx17.cpp
failed: Clang.OpenMP/declare_variant_ast_print.cpp
failed: Clang.OpenMP/declare_variant_device_kind_codegen.cpp
failed: Clang.OpenMP/declare_variant_implementation_vendor_codegen.cpp
failed: Clang.OpenMP/declare_variant_messages.c
failed: Clang.OpenMP/declare_variant_messages.cpp
failed: Clang.OpenMP/declare_variant_mixed_codegen.cpp
failed: Clang.OpenMP/math_codegen.cpp
failed: Clang.OpenMP/math_fp_macro.cpp
failed: Clang.OpenMP/nvptx_declare_variant_device_kind_codegen.cpp
failed: Clang.OpenMP/nvptx_declare_variant_implementation_vendor_codegen.cpp
failed: Clang.SemaCXX/attr-cpuspecific.cpp
failed: Clang.SemaCXX/attr-target-mv.cpp
failed: Clang.SemaCXX/friend.cpp
failed: Clang.SemaCXX/using-decl-1.cpp

Log files: console-log.txt, CMakeCache.txt

[ABataev]

In D71179#1775834, @jdoerfert wrote:

This is neither true, nor relevant. It is not true because OpenMP 5.0 declare variant is so broken it cannot be used for what it was intended for. That means people (as for example we for math) will inevitably use begin/end declare variant.

I rather doubt that it is so much broken. The fact, that you need some new construct to express some functionality does not mean that the previous one is incorrect. It is incomplete, maybe. But not broken.

Broken in the sense that we (in the OpenMP accelerator subcommittee) don't think it can be used for what we envisioned it initially. It can be used for certain things though.

And even for begin/end stuff, multiversioning is only required for construct traits, for all other traits we can reuse the existing implementation.

Again, this is not the case. begin/end *always* caused multiple definitions with the same name. Even if we ignore that for a second, why should we not use the powerful infrastructure we have (=multi-versioning) that supports construct traits and not use it for the other traits? Or asked differently, why should we have a second codegen rewriting scheme?

Not always. If we see that the context selector does not match, we can skip everything between begin/end. It means exactly what I said - multiversioning is needed only for construct because all other traits can be easily resolved at the compile time. Generally speaking, there are 2 kinds of traits - global traits (like vendor, kind, isa, etc.), which can be resolved completely statically and do not need multiversioning, and local traits, like construct, which depend on the OpenMP directives and require something similar to the multiversioning.

[jdoerfert]

Not always. If we see that the context selector does not match, we can skip everything between begin/end. It means exactly what I said - multiversioning is needed only for construct because all other traits can be easily resolved at the compile time. Generally speaking, there are 2 kinds of traits - global traits (like vendor, kind, isa, etc.), which can be resolved completely statically and do not need multiversioning, and local traits, like construct, which depend on the OpenMP directives and require something similar to the multiversioning.

The case where the code is skipped is easy, sure. However, if we "could easily resolve" the other case, we could have implemented an #ifdef solution for math.h/cmath. This was not the case and still is not. We basically populate the namespace with multiple versions of the same function (with the same name) and then select the appropriate one for each call site.

Instead of trying to argue why this is not needed for some cases, could you argue why we should have multiple schemes to resolve all types of variants? It seems you inherently assume the codegen patching scheme implemented right now is useful even if we need something else to complement it. I don't think so, thus there is little reason for me to distinguish between the types of variants that need multi-version support ant the types that can be implemented with multi-versions but don't need it.

[ABataev]

In D71179#1776034, @jdoerfert wrote:

Not always. If we see that the context selector does not match, we can skip everything between begin/end. It means exactly what I said - multiversioning is needed only for construct because all other traits can be easily resolved at the compile time. Generally speaking, there are 2 kinds of traits - global traits (like vendor, kind, isa, etc.), which can be resolved completely statically and do not need multiversioning, and local traits, like construct, which depend on the OpenMP directives and require something similar to the multiversioning.

The case where the code is skipped is easy, sure. However, if we "could easily resolve" the other case, we could have implemented an #ifdef solution for math.h/cmath. This was not the case and still is not. We basically populate the namespace with multiple versions of the same function (with the same name) and then select the appropriate one for each call site.

Instead of trying to argue why this is not needed for some cases, could you argue why we should have multiple schemes to resolve all types of variants? It seems you inherently assume the codegen patching scheme implemented right now is useful even if we need something else to complement it. I don't think so, thus there is little reason for me to distinguish between the types of variants that need multi-version support ant the types that can be implemented with multi-versions but don't need it.

Because each particular problem requires its own solution and it is always a bad idea to use the microscope to hammer the nails.

[jdoerfert]

In D71179#1776046, @ABataev wrote:

In D71179#1776034, @jdoerfert wrote:

Not always. If we see that the context selector does not match, we can skip everything between begin/end. It means exactly what I said - multiversioning is needed only for construct because all other traits can be easily resolved at the compile time. Generally speaking, there are 2 kinds of traits - global traits (like vendor, kind, isa, etc.), which can be resolved completely statically and do not need multiversioning, and local traits, like construct, which depend on the OpenMP directives and require something similar to the multiversioning.

The case where the code is skipped is easy, sure. However, if we "could easily resolve" the other case, we could have implemented an #ifdef solution for math.h/cmath. This was not the case and still is not. We basically populate the namespace with multiple versions of the same function (with the same name) and then select the appropriate one for each call site.

Instead of trying to argue why this is not needed for some cases, could you argue why we should have multiple schemes to resolve all types of variants? It seems you inherently assume the codegen patching scheme implemented right now is useful even if we need something else to complement it. I don't think so, thus there is little reason for me to distinguish between the types of variants that need multi-version support ant the types that can be implemented with multi-versions but don't need it.

Because each particular problem requires its own solution and it is always a bad idea to use the microscope to hammer the nails.

While I see where you are coming from, I disagree. We have a generic framework available that we already need to use in some cases, there is no harm in using it for all cases. It would be different if we wouldn't need the generic framework at all, but that is not the case. All I ask is to literally share existing code, no additional complexity needed. Your suggestion will complicate the setup, duplicate logic, and make it overall harder to maintain and compose in the future. If you still disagree, please provide some arguments (and details) why we would benefit from your setup.

[ABataev]

In D71179#1776108, @jdoerfert wrote:

In D71179#1776046, @ABataev wrote:

In D71179#1776034, @jdoerfert wrote:

Not always. If we see that the context selector does not match, we can skip everything between begin/end. It means exactly what I said - multiversioning is needed only for construct because all other traits can be easily resolved at the compile time. Generally speaking, there are 2 kinds of traits - global traits (like vendor, kind, isa, etc.), which can be resolved completely statically and do not need multiversioning, and local traits, like construct, which depend on the OpenMP directives and require something similar to the multiversioning.

The case where the code is skipped is easy, sure. However, if we "could easily resolve" the other case, we could have implemented an #ifdef solution for math.h/cmath. This was not the case and still is not. We basically populate the namespace with multiple versions of the same function (with the same name) and then select the appropriate one for each call site.

Instead of trying to argue why this is not needed for some cases, could you argue why we should have multiple schemes to resolve all types of variants? It seems you inherently assume the codegen patching scheme implemented right now is useful even if we need something else to complement it. I don't think so, thus there is little reason for me to distinguish between the types of variants that need multi-version support ant the types that can be implemented with multi-versions but don't need it.

Because each particular problem requires its own solution and it is always a bad idea to use the microscope to hammer the nails.

While I see where you are coming from, I disagree. We have a generic framework available that we already need to use in some cases, there is no harm in using it for all cases. It would be different if we wouldn't need the generic framework at all, but that is not the case. All I ask is to literally share existing code, no additional complexity needed. Your suggestion will complicate the setup, duplicate logic, and make it overall harder to maintain and compose in the future. If you still disagree, please provide some arguments (and details) why we would benefit from your setup.

I have different opinion. You can reuse existing codegen for declare variant functions with global context selectors only. You just need to iterate through all the variants and choose the best one.
That's why you don't need the dispatching in your scheme. You're doing absolutely the same thing as the original declare variant implementation.

We cannot use multiversioning for the original declare variant construct since there is no multiversioning at all. We have a single function with many different aliasing functions, having different names. They are completely different functions. And I don't think it would be correct to add them as multiversiin variants to the original function.

[jdoerfert]

You're doing absolutely the same thing as the original declare variant implementation.

I don't think so but if you do why do you oppose this approach?

And I don't think it would be correct to add them as multiversiin variants to the original function.

Why wouldn't it be correct to pick the version through the overload resolution instead of the code generation?
How this could work is already described in the TODO (CodeGenModule.cpp):

// TODO: We should introduce function aliases for `omp declare variant`
//       directives such that we can treat them through the same overload
//       resolution scheme (via multi versioning) as `omp begin declare
//       variant` functions. For an `omp declare variant(VARIANT) ...`
//       that is attached to a BASE function we would create a global alias
//       VARIANT = BASE which will participate in the multi version overload
//       resolution. If picked, here is no need to emit them explicitly.

---

I still haven't understood why we cannot/should not reuse the existing multi-version support and instead duplicate the logic in some custom scheme.
We have this patch that shows how we can reuse the logic in Clang. It works on a per-call basis, so it will work for all context selector (incl. construct).
If you think there is something conceptually not working, I'd like to hear about it. However, just saying "it wouldn't be correct" is not sufficient. You need to provide details about the situation, what you think would not work, and why.

[jdoerfert]

clang/lib/Headers/__clang_cuda_math_forward_declares.h

41

I have to double check what abs declarations where here and which were not.

[ABataev]

In D71179#1776457, @jdoerfert wrote:

You're doing absolutely the same thing as the original declare variant implementation.

I don't think so but if you do why do you oppose this approach?

And I don't think it would be correct to add them as multiversiin variants to the original function.

Why wouldn't it be correct to pick the version through the overload resolution instead of the code generation?
How this could work is already described in the TODO (CodeGenModule.cpp):

// TODO: We should introduce function aliases for `omp declare variant`
//       directives such that we can treat them through the same overload
//       resolution scheme (via multi versioning) as `omp begin declare
//       variant` functions. For an `omp declare variant(VARIANT) ...`
//       that is attached to a BASE function we would create a global alias
//       VARIANT = BASE which will participate in the multi version overload
//       resolution. If picked, here is no need to emit them explicitly.

---

I still haven't understood why we cannot/should not reuse the existing multi-version support and instead duplicate the logic in some custom scheme.
We have this patch that shows how we can reuse the logic in Clang. It works on a per-call basis, so it will work for all context selector (incl. construct).
If you think there is something conceptually not working, I'd like to hear about it. However, just saying "it wouldn't be correct" is not sufficient. You need to provide details about the situation, what you think would not work, and why.

I explayned already: declare variant cannot be represented as mutiversion functiin, for example.

[hfinkel]

In D71179#1776467, @ABataev wrote:

In D71179#1776457, @jdoerfert wrote:

You're doing absolutely the same thing as the original declare variant implementation.

I don't think so but if you do why do you oppose this approach?

And I don't think it would be correct to add them as multiversiin variants to the original function.

Why wouldn't it be correct to pick the version through the overload resolution instead of the code generation?
How this could work is already described in the TODO (CodeGenModule.cpp):

// TODO: We should introduce function aliases for `omp declare variant`
//       directives such that we can treat them through the same overload
//       resolution scheme (via multi versioning) as `omp begin declare
//       variant` functions. For an `omp declare variant(VARIANT) ...`
//       that is attached to a BASE function we would create a global alias
//       VARIANT = BASE which will participate in the multi version overload
//       resolution. If picked, here is no need to emit them explicitly.

---

I still haven't understood why we cannot/should not reuse the existing multi-version support and instead duplicate the logic in some custom scheme.
We have this patch that shows how we can reuse the logic in Clang. It works on a per-call basis, so it will work for all context selector (incl. construct).
If you think there is something conceptually not working, I'd like to hear about it. However, just saying "it wouldn't be correct" is not sufficient. You need to provide details about the situation, what you think would not work, and why.

I explayned already: declare variant cannot be represented as mutiversion functiin, for example.

@ABataev, can you please elaborate? It's not obvious to me that we cannot handle the existing declare variant with the same scheme (as @jdoerfert highlighted above). In general, I believe it's preferable to have one generic scheme and use it to handle all cases as opposed to continuing to use a more-limited scheme in addition to the generic scheme.

[ABataev]

In D71179#1776487, @hfinkel wrote:

In D71179#1776467, @ABataev wrote:

In D71179#1776457, @jdoerfert wrote:

You're doing absolutely the same thing as the original declare variant implementation.

I don't think so but if you do why do you oppose this approach?

And I don't think it would be correct to add them as multiversiin variants to the original function.

Why wouldn't it be correct to pick the version through the overload resolution instead of the code generation?
How this could work is already described in the TODO (CodeGenModule.cpp):

// TODO: We should introduce function aliases for `omp declare variant`
//       directives such that we can treat them through the same overload
//       resolution scheme (via multi versioning) as `omp begin declare
//       variant` functions. For an `omp declare variant(VARIANT) ...`
//       that is attached to a BASE function we would create a global alias
//       VARIANT = BASE which will participate in the multi version overload
//       resolution. If picked, here is no need to emit them explicitly.

---

I still haven't understood why we cannot/should not reuse the existing multi-version support and instead duplicate the logic in some custom scheme.
We have this patch that shows how we can reuse the logic in Clang. It works on a per-call basis, so it will work for all context selector (incl. construct).
If you think there is something conceptually not working, I'd like to hear about it. However, just saying "it wouldn't be correct" is not sufficient. You need to provide details about the situation, what you think would not work, and why.

I explayned already: declare variant cannot be represented as mutiversion functiin, for example.

@ABataev, can you please elaborate? It's not obvious to me that we cannot handle the existing declare variant with the same scheme (as @jdoerfert highlighted above). In general, I believe it's preferable to have one generic scheme and use it to handle all cases as opposed to continuing to use a more-limited scheme in addition to the generic scheme.

Eaine already. Current version of declare variant cannot be represented as multiversiin functions, because it is not. We have a function that is the alias to other functions with different names. They just are not multiversion functions by definition.

[hfinkel]

In D71179#1776491, @ABataev wrote:

In D71179#1776487, @hfinkel wrote:

In D71179#1776467, @ABataev wrote:

In D71179#1776457, @jdoerfert wrote:

You're doing absolutely the same thing as the original declare variant implementation.

I don't think so but if you do why do you oppose this approach?

And I don't think it would be correct to add them as multiversiin variants to the original function.

Why wouldn't it be correct to pick the version through the overload resolution instead of the code generation?
How this could work is already described in the TODO (CodeGenModule.cpp):

// TODO: We should introduce function aliases for `omp declare variant`
//       directives such that we can treat them through the same overload
//       resolution scheme (via multi versioning) as `omp begin declare
//       variant` functions. For an `omp declare variant(VARIANT) ...`
//       that is attached to a BASE function we would create a global alias
//       VARIANT = BASE which will participate in the multi version overload
//       resolution. If picked, here is no need to emit them explicitly.

---

I still haven't understood why we cannot/should not reuse the existing multi-version support and instead duplicate the logic in some custom scheme.
We have this patch that shows how we can reuse the logic in Clang. It works on a per-call basis, so it will work for all context selector (incl. construct).
If you think there is something conceptually not working, I'd like to hear about it. However, just saying "it wouldn't be correct" is not sufficient. You need to provide details about the situation, what you think would not work, and why.

I explayned already: declare variant cannot be represented as mutiversion functiin, for example.

@ABataev, can you please elaborate? It's not obvious to me that we cannot handle the existing declare variant with the same scheme (as @jdoerfert highlighted above). In general, I believe it's preferable to have one generic scheme and use it to handle all cases as opposed to continuing to use a more-limited scheme in addition to the generic scheme.

Eaine already. Current version of declare variant cannot be represented as multiversiin functions, because it is not. We have a function that is the alias to other functions with different names. They just are not multiversion functions by definition.

I understand that they have different names. I don't see why we that means that they can't be added to the overload set as multi-version candidates if we add logic which does exactly that.

[jdoerfert]

So, D71241 shows how declare variant (5.0) would look like if we implement it through SemaLookup. I will actually revisit this patch tomorrow as I might be able to make it even simpler. (D71241 is saving ~250 lines and from what I've seen in the tests actually fixing things.)

[ABataev]

In D71179#1776528, @hfinkel wrote:

In D71179#1776491, @ABataev wrote:

In D71179#1776487, @hfinkel wrote:

In D71179#1776467, @ABataev wrote:

In D71179#1776457, @jdoerfert wrote:

You're doing absolutely the same thing as the original declare variant implementation.

I don't think so but if you do why do you oppose this approach?

And I don't think it would be correct to add them as multiversiin variants to the original function.

Why wouldn't it be correct to pick the version through the overload resolution instead of the code generation?
How this could work is already described in the TODO (CodeGenModule.cpp):

// TODO: We should introduce function aliases for `omp declare variant`
//       directives such that we can treat them through the same overload
//       resolution scheme (via multi versioning) as `omp begin declare
//       variant` functions. For an `omp declare variant(VARIANT) ...`
//       that is attached to a BASE function we would create a global alias
//       VARIANT = BASE which will participate in the multi version overload
//       resolution. If picked, here is no need to emit them explicitly.

---

I still haven't understood why we cannot/should not reuse the existing multi-version support and instead duplicate the logic in some custom scheme.
We have this patch that shows how we can reuse the logic in Clang. It works on a per-call basis, so it will work for all context selector (incl. construct).
If you think there is something conceptually not working, I'd like to hear about it. However, just saying "it wouldn't be correct" is not sufficient. You need to provide details about the situation, what you think would not work, and why.

I explayned already: declare variant cannot be represented as mutiversion functiin, for example.

@ABataev, can you please elaborate? It's not obvious to me that we cannot handle the existing declare variant with the same scheme (as @jdoerfert highlighted above). In general, I believe it's preferable to have one generic scheme and use it to handle all cases as opposed to continuing to use a more-limited scheme in addition to the generic scheme.

Eaine already. Current version of declare variant cannot be represented as multiversiin functions, because it is not. We have a function that is the alias to other functions with different names. They just are not multiversion functions by definition.

I understand that they have different names. I don't see why we that means that they can't be added to the overload set as multi-version candidates if we add logic which does exactly that.

Because this is exactly what I said- you want to reuse the exiwting solution for completely different purpose just because you want to you reuse though even semantically it has nothing to do with multiversioning. And I think it is bad idead to break the semantics of the existing solution. It requires some addition changes like merging of different functiins with different names. And here I want to ask - why do you think it is better than my proposal to reuse the codegen for the already implemented declare variant stuff for the OpenMP multiversioned functions? It really requires less work, bdcause you just need to add a loop over all varinants and call tryEmit... function.

[hfinkel]

In D71179#1776761, @ABataev wrote:

In D71179#1776528, @hfinkel wrote:

In D71179#1776491, @ABataev wrote:

In D71179#1776487, @hfinkel wrote:

In D71179#1776467, @ABataev wrote:

In D71179#1776457, @jdoerfert wrote:

You're doing absolutely the same thing as the original declare variant implementation.

I don't think so but if you do why do you oppose this approach?

And I don't think it would be correct to add them as multiversiin variants to the original function.

Why wouldn't it be correct to pick the version through the overload resolution instead of the code generation?
How this could work is already described in the TODO (CodeGenModule.cpp):

// TODO: We should introduce function aliases for `omp declare variant`
//       directives such that we can treat them through the same overload
//       resolution scheme (via multi versioning) as `omp begin declare
//       variant` functions. For an `omp declare variant(VARIANT) ...`
//       that is attached to a BASE function we would create a global alias
//       VARIANT = BASE which will participate in the multi version overload
//       resolution. If picked, here is no need to emit them explicitly.

---

I still haven't understood why we cannot/should not reuse the existing multi-version support and instead duplicate the logic in some custom scheme.
We have this patch that shows how we can reuse the logic in Clang. It works on a per-call basis, so it will work for all context selector (incl. construct).
If you think there is something conceptually not working, I'd like to hear about it. However, just saying "it wouldn't be correct" is not sufficient. You need to provide details about the situation, what you think would not work, and why.

I explayned already: declare variant cannot be represented as mutiversion functiin, for example.

@ABataev, can you please elaborate? It's not obvious to me that we cannot handle the existing declare variant with the same scheme (as @jdoerfert highlighted above). In general, I believe it's preferable to have one generic scheme and use it to handle all cases as opposed to continuing to use a more-limited scheme in addition to the generic scheme.

Eaine already. Current version of declare variant cannot be represented as multiversiin functions, because it is not. We have a function that is the alias to other functions with different names. They just are not multiversion functions by definition.

I understand that they have different names. I don't see why we that means that they can't be added to the overload set as multi-version candidates if we add logic which does exactly that.

@jdoerfert posted a prototype implementation in D71241, so we don't need to just have a theoretical discussion, but I'd like to address a high-level issue here:

Because this is exactly what I said- you want to reuse the exiwting solution for completely different purpose just because you want to you reuse though even semantically it has nothing to do with multiversioning.

This kind of comment really isn't appropriate. We're all experienced developers here, and no one is proposing to reuse code in an inappropriate manner "just because" or for any other reason. I ask you to reconsider your reasoning here for two reasons:

1. "Reus[ing] the existing solution for a completely different purpose", which I'll classify as structural code reuse, is not necessarily bad. Structural code reuse, where you reuse code with a similar structure, but different purpose, from what you need, is often a useful impetus for the creation of new abstractions. The trade off relevant here, in my experience, is against future structural divergence. In the future, is it likely that the abstraction will break down because the different purposes will tend to require the code structure to change in the future in divergent ways? If so, that can be a good argument against code reuse.
2. Your statement of differing purpose, that declare variant has "nothing to do with multiversioning", it not obviously true. Declare variant, as the spec says, "declares a specialized variant of a base function and specifies the context in which that specialized variant is used." multiversioning, according to the GCC docs, makes it so that "you may specify multiple versions of a function, where each function is specialized for a specific target feature. At runtime, the appropriate version of the function is automatically executed depending on the characteristics of the execution platform." These two concepts do share some conceptual relationship.

It certainly seems fair to say that the AST representation desired for a call with runtime dispatch might be sufficiently different from a call resolved at compile time to make the code reuse inadvisable. However, the requirements of which I'm aware are: the representation should be unambiguous and faithful to the source and language structure, and also, the statically-resolvable callee should be referenced by the call site in the AST. As can be seen in this patch and the associated tests, both of these requirements are satisfied.

And I think it is bad idead to break the semantics of the existing solution. It requires some addition changes like merging of different functiins with different names. And here I want to ask - why do you think it is better than my proposal to reuse the codegen for the already implemented declare variant stuff for the OpenMP multiversioned functions? It really requires less work, bdcause you just need to add a loop over all varinants and call tryEmit... function.

When you say "semantics of the existing solution", do you mean the extent to which it satisfies the standard, non-standard user-visible behaviors of the existing implementation, or the internal structure of the implementation? It's certainly not a bad idea to change, or even replace, the current implementation if the result is better in some way (e.g., more general, supports more features, conceptually cleaner). The proposed solution seems to have a number of advantages over the current solution in codegen, and in addition, naturally handles the new features that we would like to support. Generally, resolution of static calls is something that should happen in Sema, not in CodeGen, in part so that static analysis tools can easily understand the call semantics. This new approach naturally provides this implementation property.

[ABataev]

In D71179#1778512, @hfinkel wrote:

In D71179#1776761, @ABataev wrote:

In D71179#1776528, @hfinkel wrote:

In D71179#1776491, @ABataev wrote:

In D71179#1776487, @hfinkel wrote:

In D71179#1776467, @ABataev wrote:

In D71179#1776457, @jdoerfert wrote:

You're doing absolutely the same thing as the original declare variant implementation.

I don't think so but if you do why do you oppose this approach?

And I don't think it would be correct to add them as multiversiin variants to the original function.

Why wouldn't it be correct to pick the version through the overload resolution instead of the code generation?
How this could work is already described in the TODO (CodeGenModule.cpp):

// TODO: We should introduce function aliases for `omp declare variant`
//       directives such that we can treat them through the same overload
//       resolution scheme (via multi versioning) as `omp begin declare
//       variant` functions. For an `omp declare variant(VARIANT) ...`
//       that is attached to a BASE function we would create a global alias
//       VARIANT = BASE which will participate in the multi version overload
//       resolution. If picked, here is no need to emit them explicitly.

---

I still haven't understood why we cannot/should not reuse the existing multi-version support and instead duplicate the logic in some custom scheme.
We have this patch that shows how we can reuse the logic in Clang. It works on a per-call basis, so it will work for all context selector (incl. construct).
If you think there is something conceptually not working, I'd like to hear about it. However, just saying "it wouldn't be correct" is not sufficient. You need to provide details about the situation, what you think would not work, and why.

I explayned already: declare variant cannot be represented as mutiversion functiin, for example.

@ABataev, can you please elaborate? It's not obvious to me that we cannot handle the existing declare variant with the same scheme (as @jdoerfert highlighted above). In general, I believe it's preferable to have one generic scheme and use it to handle all cases as opposed to continuing to use a more-limited scheme in addition to the generic scheme.

Eaine already. Current version of declare variant cannot be represented as multiversiin functions, because it is not. We have a function that is the alias to other functions with different names. They just are not multiversion functions by definition.

I understand that they have different names. I don't see why we that means that they can't be added to the overload set as multi-version candidates if we add logic which does exactly that.

@jdoerfert posted a prototype implementation in D71241, so we don't need to just have a theoretical discussion, but I'd like to address a high-level issue here:

Because this is exactly what I said- you want to reuse the exiwting solution for completely different purpose just because you want to you reuse though even semantically it has nothing to do with multiversioning.

This kind of comment really isn't appropriate. We're all experienced developers here, and no one is proposing to reuse code in an inappropriate manner "just because" or for any other reason. I ask you to reconsider your reasoning here for two reasons:

1. "Reus[ing] the existing solution for a completely different purpose", which I'll classify as structural code reuse, is not necessarily bad. Structural code reuse, where you reuse code with a similar structure, but different purpose, from what you need, is often a useful impetus for the creation of new abstractions. The trade off relevant here, in my experience, is against future structural divergence. In the future, is it likely that the abstraction will break down because the different purposes will tend to require the code structure to change in the future in divergent ways? If so, that can be a good argument against code reuse.

I agree that reusing is a good idea but not in this case. I already wrote that Johanmes reuses just a single feature of the multiversioning - handling of the multiple definitons for the same function. Nothing else. Everything else is a new functionality to support declare variant stuff.

2. Your statement of differing purpose, that declare variant has "nothing to do with multiversioning", it not obviously true. Declare variant, as the spec says, "declares a specialized variant of a base function and specifies the context in which that specialized variant is used." multiversioning, according to the GCC docs, makes it so that "you may specify multiple versions of a function, where each function is specialized for a specific target feature. At runtime, the appropriate version of the function is automatically executed depending on the characteristics of the execution platform." These two concepts do share some conceptual relationship.

It certainly seems fair to say that the AST representation desired for a call with runtime dispatch might be sufficiently different from a call resolved at compile time to make the code reuse inadvisable. However, the requirements of which I'm aware are: the representation should be unambiguous and faithful to the source and language structure, and also, the statically-resolvable callee should be referenced by the call site in the AST. As can be seen in this patch and the associated tests, both of these requirements are satisfied.

Current solution allows to do everything correctly and requires just a small rework. The actual selection of the fumction can (and must be done) at the codegen. There is no benefits inchoosing correct variant functiin in sema. Moreover, it leads to breaking of the AST in the way that the original function call is replaced by a new function call in AST. And dump/printing works differently than the user expected. L

And I think it is bad idead to break the semantics of the existing solution. It requires some addition changes like merging of different functiins with different names. And here I want to ask - why do you think it is better than my proposal to reuse the codegen for the already implemented declare variant stuff for the OpenMP multiversioned functions? It really requires less work, bdcause you just need to add a loop over all varinants and call tryEmit... function.

When you say "semantics of the existing solution", do you mean the extent to which it satisfies the standard, non-standard user-visible behaviors of the existing implementation, or the internal structure of the implementation? It's certainly not a bad idea to change, or even replace, the current implementation if the result is better in some way (e.g., more general, supports more features, conceptually cleaner). The proposed solution seems to have a number of advantages over the current solution in codegen, and in addition, naturally handles the new features that we would like to support. Generally, resolution of static calls is something that should happen in Sema, not in CodeGen, in part so that static analysis tools can easily understand the call semantics. This new approach naturally provides this implementation property.

There are no advantages at all. This solution has absolutely the same power as the existing one. And I see not a single reason why a function resolution should happen in sema. Moreover, there are real problems with the proposed solution with AST.

[jdoerfert]

Consistent overload based solution.

[jdoerfert]

Diff against TOT

[merge_guards_bot]

Build result: FAILURE - Could not check out parent git hash "9a3d576b08c13533597182498ba5e739924f892f". It was not found in the repository. Did you configure the "Parent Revision" in Phabricator properly? Trying to apply the patch to the master branch instead...

ERROR: arc patch failed with error code 1. Check build log for details.
Log files: console-log.txt, CMakeCache.txt

[merge_guards_bot]

Build result: FAILURE - Could not check out parent git hash "9a3d576b08c13533597182498ba5e739924f892f". It was not found in the repository. Did you configure the "Parent Revision" in Phabricator properly? Trying to apply the patch to the master branch instead...

ERROR: arc patch failed with error code 1. Check build log for details.
Log files: console-log.txt, CMakeCache.txt

[jdoerfert]

Fix math problem

[merge_guards_bot]

Build result: FAILURE - Could not check out parent git hash "9a3d576b08c13533597182498ba5e739924f892f". It was not found in the repository. Did you configure the "Parent Revision" in Phabricator properly? Trying to apply the patch to the master branch instead...

ERROR: arc patch failed with error code 1. Check build log for details.
Log files: console-log.txt, CMakeCache.txt

[jdoerfert]

D75779 is the proper implementation of the OpenMP standard.

clang/include/clang/AST/Decl.h

private:
   void setParameterIndexLarge(unsigned parameterIndex);
   unsigned getParameterIndexLarge() const;
 };
 
 enum class MultiVersionKind {
   None,
   Target,
   CPUSpecific,
-  CPUDispatch
+  CPUDispatch,
+  OMPVariant,
 };
 
 /// Represents a function declaration or definition.
 ///
 /// Since a given function can be declared several times in a program,
 /// there may be several FunctionDecls that correspond to that
 /// function. Only one of those FunctionDecls will be found when
 /// traversing the list of declarations in the context of the
public:
   /// True if this function is a multiversioned processor specific function as a
   /// part of the cpu_specific/cpu_dispatch functionality.
   bool isCPUSpecificMultiVersion() const;
 
   /// True if this function is a multiversioned dispatch function as a part of
   /// the target functionality.
   bool isTargetMultiVersion() const;
 
+  /// True if this function is a multiversioned function as a part of
+  /// the OpenMP begin/end declare variant functionality.
+  bool isOpenMPMultiVersion() const;
+
   void setPreviousDeclaration(FunctionDecl * PrevDecl);
 
   FunctionDecl *getCanonicalDecl() override;
   const FunctionDecl *getCanonicalDecl() const {
     return const_cast<FunctionDecl*>(this)->getCanonicalDecl();
   }
 
   unsigned getBuiltinID(bool ConsiderWrapperFunctions = false) const;

clang/include/clang/AST/StmtOpenMP.h

public:
   CreateEmpty(const ASTContext &C, unsigned NumClauses, unsigned CollapsedNum,
               EmptyShell);
 
   static bool classof(const Stmt *T) {
     return T->getStmtClass() == OMPTargetTeamsDistributeSimdDirectiveClass;
   }
 };
 
+class OMPDeclareVariantAttr;
+
+/// Helper to determine the best of two potential context matches. Note that
+/// nullptr are valid inputs but also valid outputs, e.g., if neither attribute
+/// describes a matching context.
+const OMPDeclareVariantAttr *
+getBetterOpenMPContextMatch(ASTContext &C, const OMPDeclareVariantAttr *LHSAttr,
+                            const OMPDeclareVariantAttr *RHSAttr);
+
+/// Return true if the context described by \p A matches.
+bool isOpenMPContextMatch(ASTContext &C, const OMPDeclareVariantAttr *A);
+
 } // end namespace clang
 
 #endif

clang/include/clang/Basic/OpenMPKinds.def

 OPENMP_DIRECTIVE_EXT(target_teams_distribute_parallel_for_simd, "target teams distribute parallel for simd")
 OPENMP_DIRECTIVE_EXT(target_teams_distribute_simd, "target teams distribute simd")
 OPENMP_DIRECTIVE(allocate)
 OPENMP_DIRECTIVE_EXT(declare_variant, "declare variant")
 OPENMP_DIRECTIVE_EXT(master_taskloop, "master taskloop")
 OPENMP_DIRECTIVE_EXT(parallel_master_taskloop, "parallel master taskloop")
 OPENMP_DIRECTIVE_EXT(master_taskloop_simd, "master taskloop simd")
 OPENMP_DIRECTIVE_EXT(parallel_master_taskloop_simd, "parallel master taskloop simd")
+OPENMP_DIRECTIVE_EXT(begin_declare_variant, "begin declare variant")
+OPENMP_DIRECTIVE_EXT(end_declare_variant, "end declare variant")
 
 // OpenMP clauses.
 OPENMP_CLAUSE(allocator, OMPAllocatorClause)
 OPENMP_CLAUSE(if, OMPIfClause)
 OPENMP_CLAUSE(final, OMPFinalClause)
 OPENMP_CLAUSE(num_threads, OMPNumThreadsClause)
 OPENMP_CLAUSE(safelen, OMPSafelenClause)
 OPENMP_CLAUSE(simdlen, OMPSimdlenClause)

clang/include/clang/Parse/Parser.h

DeclGroupPtrTy ParseOMPDeclareSimdClauses(DeclGroupPtrTy Ptr,
                                             CachedTokens &Toks,
                                             SourceLocation Loc);
   /// Parses OpenMP context selectors and calls \p Callback for each
   /// successfully parsed context selector.
   bool
   parseOpenMPContextSelectors(SourceLocation Loc,
                               SmallVectorImpl<Sema::OMPCtxSelectorData> &Data);
 
+  /// Parse match clause of '#pragma omp [begin] declare variant'.
+  void ParseOMPDeclareVariantMatchClause(
+      SourceLocation Loc, SmallVectorImpl<Sema::OMPCtxSelectorData> &Data);
+
   /// Parse clauses for '#pragma omp declare variant'.
   void ParseOMPDeclareVariantClauses(DeclGroupPtrTy Ptr, CachedTokens &Toks,
                                      SourceLocation Loc);
   /// Parse clauses for '#pragma omp declare target'.
   DeclGroupPtrTy ParseOMPDeclareTargetClauses();
   /// Parse '#pragma omp end declare target'.
   void ParseOMPEndDeclareTargetDirective(OpenMPDirectiveKind DKind,
                                          SourceLocation Loc);

clang/include/clang/Sema/Sema.h

public:
   /// emit diagnostics.
   /// \return true if type is disabled.
   bool checkOpenCLDisabledDecl(const NamedDecl &D, const Expr &E);
 
   //===--------------------------------------------------------------------===//
   // OpenMP directives and clauses.
   //
 private:
+  /// Copies declare variant attributes from the template TD to the function FD.
+  void inheritOpenMPVariantAttrs(FunctionDecl *FD,
+                                 const FunctionTemplateDecl &TD);
   void *VarDataSharingAttributesStack;
   /// Number of nested '#pragma omp declare target' directives.
   unsigned DeclareTargetNestingLevel = 0;
   /// Initialization of data-sharing attributes stack.
   void InitDataSharingAttributesStack();
   void DestroyDataSharingAttributesStack();
   ExprResult
   VerifyPositiveIntegerConstantInClause(Expr *Op, OpenMPClauseKind CKind,
void markOpenMPDeclareVariantFuncsReferenced(SourceLocation Loc,
                                                bool MightBeOdrUse);
 
 public:
   /// Struct to store the context selectors info for declare variant directive.
   using OMPCtxStringType = SmallString<8>;
   using OMPCtxSelectorData =
       OpenMPCtxSelectorData<SmallVector<OMPCtxStringType, 4>, ExprResult>;
 
+  /// A declare variant attribute if we are inside a begin/end declare variant
+  OMPDeclareVariantAttr *DeclareVariantScopeAttr = nullptr;
+
   /// Checks if the variant/multiversion functions are compatible.
   bool areMultiversionVariantFunctionsCompatible(
       const FunctionDecl *OldFD, const FunctionDecl *NewFD,
       const PartialDiagnostic &NoProtoDiagID,
       const PartialDiagnosticAt &NoteCausedDiagIDAt,
       const PartialDiagnosticAt &NoSupportDiagIDAt,
       const PartialDiagnosticAt &DiffDiagIDAt, bool TemplatesSupported,
-      bool ConstexprSupported, bool CLinkageMayDiffer);
+      bool ConstexprSupported, bool CLinkageMayDiffer,
+      bool StorageClassMayDiffer, bool ConstexprSpecMayDiffer,
+      bool InlineSpecificationMayDiffer);
 
   /// Function tries to capture lambda's captured variables in the OpenMP region
   /// before the original lambda is captured.
   void tryCaptureOpenMPLambdas(ValueDecl *V);
 
   /// Return true if the provided declaration \a VD should be captured by
   /// reference.
   /// \param Level Relative level of nested OpenMP construct for that the check
public:
   /// Called on well-formed '\#pragma omp declare variant' after parsing of
   /// the associated method/function.
   /// \param FD Function declaration to which declare variant directive is
   /// applied to.
   /// \param VariantRef Expression that references the variant function, which
   /// must be used instead of the original one, specified in \p DG.
   /// \param Data Set of context-specific data for the specified context
   /// selector.
-  void ActOnOpenMPDeclareVariantDirective(FunctionDecl *FD, Expr *VariantRef,
+  bool ActOnOpenMPDeclareVariantDirective(FunctionDecl *FD, Expr *VariantRef,
                                           SourceRange SR,
                                           ArrayRef<OMPCtxSelectorData> Data);
 
   OMPClause *ActOnOpenMPSingleExprClause(OpenMPClauseKind Kind,
                                          Expr *Expr,
                                          SourceLocation StartLoc,
                                          SourceLocation LParenLoc,
                                          SourceLocation EndLoc);

clang/lib/AST/Decl.cpp

 
 MultiVersionKind FunctionDecl::getMultiVersionKind() const {
   if (hasAttr<TargetAttr>())
     return MultiVersionKind::Target;
   if (hasAttr<CPUDispatchAttr>())
     return MultiVersionKind::CPUDispatch;
   if (hasAttr<CPUSpecificAttr>())
     return MultiVersionKind::CPUSpecific;
+  if (hasAttr<OMPDeclareVariantAttr>() &&
+      !getAttr<OMPDeclareVariantAttr>()->getVariantFuncRef())
+    return MultiVersionKind::OMPVariant;
   return MultiVersionKind::None;
 }
 
 bool FunctionDecl::isCPUDispatchMultiVersion() const {
   return isMultiVersion() && hasAttr<CPUDispatchAttr>();
 }
 
 bool FunctionDecl::isCPUSpecificMultiVersion() const {
   return isMultiVersion() && hasAttr<CPUSpecificAttr>();
 }
 
 bool FunctionDecl::isTargetMultiVersion() const {
   return isMultiVersion() && hasAttr<TargetAttr>();
 }
 
+bool FunctionDecl::isOpenMPMultiVersion() const {
+  return isMultiVersion() && hasAttr<OMPDeclareVariantAttr>() &&
+         !getAttr<OMPDeclareVariantAttr>()->getVariantFuncRef();
+}
+
 void
 FunctionDecl::setPreviousDeclaration(FunctionDecl *PrevDecl) {
   redeclarable_base::setPreviousDecl(PrevDecl);
 
   if (FunctionTemplateDecl *FunTmpl = getDescribedFunctionTemplate()) {
     FunctionTemplateDecl *PrevFunTmpl
       = PrevDecl? PrevDecl->getDescribedFunctionTemplate() : nullptr;
     assert((!PrevDecl || PrevFunTmpl) && "Function/function template mismatch");

clang/lib/AST/StmtOpenMP.cpp

 //===--- StmtOpenMP.cpp - Classes for OpenMP directives -------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file implements the subclesses of Stmt class declared in StmtOpenMP.h
 //
 //===----------------------------------------------------------------------===//
 
 #include "clang/AST/StmtOpenMP.h"
 
 #include "clang/AST/ASTContext.h"
+#include "llvm/ADT/SetOperations.h"
 
 using namespace clang;
 
 void OMPExecutableDirective::setClauses(ArrayRef<OMPClause *> Clauses) {
   assert(Clauses.size() == getNumClauses() &&
          "Number of clauses is not the same as the preallocated buffer");
   std::copy(Clauses.begin(), Clauses.end(), getClauses().begin());
 }
auto Size = llvm::alignTo(sizeof(OMPTargetTeamsDistributeSimdDirective),
                             alignof(OMPClause *));
   void *Mem = C.Allocate(
       Size + sizeof(OMPClause *) * NumClauses +
       sizeof(Stmt *) *
           numLoopChildren(CollapsedNum, OMPD_target_teams_distribute_simd));
   return new (Mem)
       OMPTargetTeamsDistributeSimdDirective(CollapsedNum, NumClauses);
 }
+
+// TODO: We have various representations for the same data, it might help to
+//       reuse some instead of converting them.
+// TODO: It is unclear where this checking code should live. It is used all over
+//       the place and would probably fit bet in OMPDeclareVariantAttr.
+using OMPContextSelectorData =
+    OpenMPCtxSelectorData<ArrayRef<StringRef>, llvm::APSInt>;
+using CompleteOMPContextSelectorData = SmallVector<OMPContextSelectorData, 4>;
+
+/// Checks current context and returns true if it matches the context selector.
+template <OpenMPContextSelectorSetKind CtxSet, OpenMPContextSelectorKind Ctx,
+          typename... Arguments>
+static bool checkContext(const OMPContextSelectorData &Data,
+                         Arguments... Params) {
+  assert(Data.CtxSet != OMP_CTX_SET_unknown && Data.Ctx != OMP_CTX_unknown &&
+         "Unknown context selector or context selector set.");
+  return false;
+}
+
+/// Checks for implementation={vendor(<vendor>)} context selector.
+/// \returns true iff <vendor>="llvm", false otherwise.
+template <>
+bool checkContext<OMP_CTX_SET_implementation, OMP_CTX_vendor>(
+    const OMPContextSelectorData &Data) {
+  return llvm::all_of(Data.Names,
+                      [](StringRef S) { return !S.compare_lower("llvm"); });
+}
+
+/// Checks for device={kind(<kind>)} context selector.
+/// \returns true if <kind>="host" and compilation is for host.
+/// true if <kind>="nohost" and compilation is for device.
+/// true if <kind>="cpu" and compilation is for Arm, X86 or PPC CPU.
+/// true if <kind>="gpu" and compilation is for NVPTX or AMDGCN.
+/// false otherwise.
+template <>
+bool checkContext<OMP_CTX_SET_device, OMP_CTX_kind, const LangOptions &,
+                  const TargetInfo &>(const OMPContextSelectorData &Data,
+                                      const LangOptions &LO,
+                                      const TargetInfo &TI) {
+  for (StringRef Name : Data.Names) {
+    if (!Name.compare_lower("host")) {
+      if (LO.OpenMPIsDevice)
+        return false;
+      continue;
+    }
+    if (!Name.compare_lower("nohost")) {
+      if (!LO.OpenMPIsDevice)
+        return false;
+      continue;
+    }
+    switch (TI.getTriple().getArch()) {
+    case llvm::Triple::arm:
+    case llvm::Triple::armeb:
+    case llvm::Triple::aarch64:
+    case llvm::Triple::aarch64_be:
+    case llvm::Triple::aarch64_32:
+    case llvm::Triple::ppc:
+    case llvm::Triple::ppc64:
+    case llvm::Triple::ppc64le:
+    case llvm::Triple::x86:
+    case llvm::Triple::x86_64:
+      if (Name.compare_lower("cpu"))
+        return false;
+      break;
+    case llvm::Triple::amdgcn:
+    case llvm::Triple::nvptx:
+    case llvm::Triple::nvptx64:
+      if (Name.compare_lower("gpu"))
+        return false;
+      break;
+    case llvm::Triple::UnknownArch:
+    case llvm::Triple::arc:
+    case llvm::Triple::avr:
+    case llvm::Triple::bpfel:
+    case llvm::Triple::bpfeb:
+    case llvm::Triple::hexagon:
+    case llvm::Triple::mips:
+    case llvm::Triple::mipsel:
+    case llvm::Triple::mips64:
+    case llvm::Triple::mips64el:
+    case llvm::Triple::msp430:
+    case llvm::Triple::r600:
+    case llvm::Triple::riscv32:
+    case llvm::Triple::riscv64:
+    case llvm::Triple::sparc:
+    case llvm::Triple::sparcv9:
+    case llvm::Triple::sparcel:
+    case llvm::Triple::systemz:
+    case llvm::Triple::tce:
+    case llvm::Triple::tcele:
+    case llvm::Triple::thumb:
+    case llvm::Triple::thumbeb:
+    case llvm::Triple::xcore:
+    case llvm::Triple::le32:
+    case llvm::Triple::le64:
+    case llvm::Triple::amdil:
+    case llvm::Triple::amdil64:
+    case llvm::Triple::hsail:
+    case llvm::Triple::hsail64:
+    case llvm::Triple::spir:
+    case llvm::Triple::spir64:
+    case llvm::Triple::kalimba:
+    case llvm::Triple::shave:
+    case llvm::Triple::lanai:
+    case llvm::Triple::wasm32:
+    case llvm::Triple::wasm64:
+    case llvm::Triple::renderscript32:
+    case llvm::Triple::renderscript64:
+      return false;
+    }
+  }
+  return true;
+}
+
+static CompleteOMPContextSelectorData
+translateAttrToContextSelectorData(ASTContext &C,
+                                   const OMPDeclareVariantAttr *A) {
+  CompleteOMPContextSelectorData Data;
+  if (!A)
+    return Data;
+  for (unsigned I = 0, E = A->scores_size(); I < E; ++I) {
+    Data.emplace_back();
+    auto CtxSet = static_cast<OpenMPContextSelectorSetKind>(
+        *std::next(A->ctxSelectorSets_begin(), I));
+    auto Ctx = static_cast<OpenMPContextSelectorKind>(
+        *std::next(A->ctxSelectors_begin(), I));
+    Data.back().CtxSet = CtxSet;
+    Data.back().Ctx = Ctx;
+    const Expr *Score = *std::next(A->scores_begin(), I);
+    Data.back().Score = Score->EvaluateKnownConstInt(C);
+    switch (Ctx) {
+    case OMP_CTX_vendor:
+      assert(CtxSet == OMP_CTX_SET_implementation &&
+             "Expected implementation context selector set.");
+      Data.back().Names =
+          llvm::makeArrayRef(A->implVendors_begin(), A->implVendors_end());
+      break;
+    case OMP_CTX_kind:
+      assert(CtxSet == OMP_CTX_SET_device &&
+             "Expected device context selector set.");
+      Data.back().Names =
+          llvm::makeArrayRef(A->deviceKinds_begin(), A->deviceKinds_end());
+      break;
+    case OMP_CTX_unknown:
+      llvm_unreachable("Unknown context selector kind.");
+    }
+  }
+  return Data;
+}
+
+static bool
+matchesOpenMPContextImpl(const CompleteOMPContextSelectorData &ContextData,
+                         const LangOptions &LO, const TargetInfo &TI) {
+  for (const OMPContextSelectorData &Data : ContextData) {
+    switch (Data.Ctx) {
+    case OMP_CTX_vendor:
+      assert(Data.CtxSet == OMP_CTX_SET_implementation &&
+             "Expected implementation context selector set.");
+      if (!checkContext<OMP_CTX_SET_implementation, OMP_CTX_vendor>(Data))
+        return false;
+      break;
+    case OMP_CTX_kind:
+      assert(Data.CtxSet == OMP_CTX_SET_device &&
+             "Expected device context selector set.");
+      if (!checkContext<OMP_CTX_SET_device, OMP_CTX_kind, const LangOptions &,
+                        const TargetInfo &>(Data, LO, TI))
+        return false;
+      break;
+    case OMP_CTX_unknown:
+      llvm_unreachable("Unknown context selector kind.");
+    }
+  }
+  return true;
+}
+
+static bool isStrictSubset(const CompleteOMPContextSelectorData &LHS,
+                           const CompleteOMPContextSelectorData &RHS) {
+  llvm::SmallDenseMap<std::pair<int, int>, llvm::StringSet<>, 4> RHSData;
+  for (const OMPContextSelectorData &D : RHS) {
+    auto &Pair = RHSData.FindAndConstruct(std::make_pair(D.CtxSet, D.Ctx));
+    Pair.getSecond().insert(D.Names.begin(), D.Names.end());
+  }
+  bool AllSetsAreEqual = true;
+  for (const OMPContextSelectorData &D : LHS) {
+    auto It = RHSData.find(std::make_pair(D.CtxSet, D.Ctx));
+    if (It == RHSData.end())
+      return false;
+    if (D.Names.size() > It->getSecond().size())
+      return false;
+    if (llvm::set_union(It->getSecond(), D.Names))
+      return false;
+    AllSetsAreEqual =
+        AllSetsAreEqual && (D.Names.size() == It->getSecond().size());
+  }
+
+  return LHS.size() != RHS.size() || !AllSetsAreEqual;
+}
+
+const OMPDeclareVariantAttr *
+clang::getBetterOpenMPContextMatch(ASTContext &C,
+                                   const OMPDeclareVariantAttr *LHSAttr,
+                                   const OMPDeclareVariantAttr *RHSAttr) {
+  if (LHSAttr->isInherited() && RHSAttr->isInherited())
+    return nullptr;
+  if (LHSAttr->isInherited())
+    return RHSAttr;
+  if (RHSAttr->isInherited())
+    return LHSAttr;
+  const CompleteOMPContextSelectorData LHS =
+      translateAttrToContextSelectorData(C, LHSAttr);
+  const CompleteOMPContextSelectorData RHS =
+      translateAttrToContextSelectorData(C, RHSAttr);
+  bool LHSMatch =
+      matchesOpenMPContextImpl(LHS, C.getLangOpts(), C.getTargetInfo());
+  bool RHSMatch =
+      matchesOpenMPContextImpl(RHS, C.getLangOpts(), C.getTargetInfo());
+  if (!LHSMatch && !RHSMatch)
+    return nullptr;
+  if (LHSMatch && !RHSMatch)
+    return LHSAttr;
+  if (!LHSMatch && RHSMatch)
+    return RHSAttr;
+  assert(LHSMatch && RHSMatch && "broken invariant");
+
+  // Score is calculated as sum of all scores + 1.
+  llvm::APSInt LHSScore(llvm::APInt(64, 1), /*isUnsigned=*/false);
+  bool RHSIsSubsetOfLHS = isStrictSubset(RHS, LHS);
+  if (RHSIsSubsetOfLHS) {
+    LHSScore = llvm::APSInt::get(0);
+  } else {
+    for (const OMPContextSelectorData &Data : LHS) {
+      if (Data.Score.getBitWidth() > LHSScore.getBitWidth()) {
+        LHSScore = LHSScore.extend(Data.Score.getBitWidth()) + Data.Score;
+      } else if (Data.Score.getBitWidth() < LHSScore.getBitWidth()) {
+        LHSScore += Data.Score.extend(LHSScore.getBitWidth());
+      } else {
+        LHSScore += Data.Score;
+      }
+    }
+  }
+  llvm::APSInt RHSScore(llvm::APInt(64, 1), /*isUnsigned=*/false);
+  if (!RHSIsSubsetOfLHS && isStrictSubset(LHS, RHS)) {
+    RHSScore = llvm::APSInt::get(0);
+  } else {
+    for (const OMPContextSelectorData &Data : RHS) {
+      if (Data.Score.getBitWidth() > RHSScore.getBitWidth()) {
+        RHSScore = RHSScore.extend(Data.Score.getBitWidth()) + Data.Score;
+      } else if (Data.Score.getBitWidth() < RHSScore.getBitWidth()) {
+        RHSScore += Data.Score.extend(RHSScore.getBitWidth());
+      } else {
+        RHSScore += Data.Score;
+      }
+    }
+  }
+  return llvm::APSInt::compareValues(LHSScore, RHSScore) >= 0 ? LHSAttr
+                                                              : RHSAttr;
+}
+
+bool clang::isOpenMPContextMatch(ASTContext &C,
+                                 const OMPDeclareVariantAttr *A) {
+  const CompleteOMPContextSelectorData Data =
+      translateAttrToContextSelectorData(C, A);
+  return matchesOpenMPContextImpl(Data, C.getLangOpts(), C.getTargetInfo());
+}

clang/lib/CodeGen/CGOpenMPRuntime.cpp

CGF.EHStack.pushCleanup<OMPAllocateCleanupTy>(NormalAndEHCleanup, FiniRTLFn,
                                                 llvm::makeArrayRef(FiniArgs));
   Addr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
       Addr,
       CGF.ConvertTypeForMem(CGM.getContext().getPointerType(CVD->getType())),
       CVD->getName() + ".addr");
   return Address(Addr, Align);
 }
 
-namespace {
-using OMPContextSelectorData =
-    OpenMPCtxSelectorData<ArrayRef<StringRef>, llvm::APSInt>;
-using CompleteOMPContextSelectorData = SmallVector<OMPContextSelectorData, 4>;
-} // anonymous namespace
-
-/// Checks current context and returns true if it matches the context selector.
-template <OpenMPContextSelectorSetKind CtxSet, OpenMPContextSelectorKind Ctx,
-          typename... Arguments>
-static bool checkContext(const OMPContextSelectorData &Data,
-                         Arguments... Params) {
-  assert(Data.CtxSet != OMP_CTX_SET_unknown && Data.Ctx != OMP_CTX_unknown &&
-         "Unknown context selector or context selector set.");
-  return false;
-}
-
-/// Checks for implementation={vendor(<vendor>)} context selector.
-/// \returns true iff <vendor>="llvm", false otherwise.
-template <>
-bool checkContext<OMP_CTX_SET_implementation, OMP_CTX_vendor>(
-    const OMPContextSelectorData &Data) {
-  return llvm::all_of(Data.Names,
-                      [](StringRef S) { return !S.compare_lower("llvm"); });
-}
-
-/// Checks for device={kind(<kind>)} context selector.
-/// \returns true if <kind>="host" and compilation is for host.
-/// true if <kind>="nohost" and compilation is for device.
-/// true if <kind>="cpu" and compilation is for Arm, X86 or PPC CPU.
-/// true if <kind>="gpu" and compilation is for NVPTX or AMDGCN.
-/// false otherwise.
-template <>
-bool checkContext<OMP_CTX_SET_device, OMP_CTX_kind, CodeGenModule &>(
-    const OMPContextSelectorData &Data, CodeGenModule &CGM) {
-  for (StringRef Name : Data.Names) {
-    if (!Name.compare_lower("host")) {
-      if (CGM.getLangOpts().OpenMPIsDevice)
-        return false;
-      continue;
-    }
-    if (!Name.compare_lower("nohost")) {
-      if (!CGM.getLangOpts().OpenMPIsDevice)
-        return false;
-      continue;
-    }
-    switch (CGM.getTriple().getArch()) {
-    case llvm::Triple::arm:
-    case llvm::Triple::armeb:
-    case llvm::Triple::aarch64:
-    case llvm::Triple::aarch64_be:
-    case llvm::Triple::aarch64_32:
-    case llvm::Triple::ppc:
-    case llvm::Triple::ppc64:
-    case llvm::Triple::ppc64le:
-    case llvm::Triple::x86:
-    case llvm::Triple::x86_64:
-      if (Name.compare_lower("cpu"))
-        return false;
-      break;
-    case llvm::Triple::amdgcn:
-    case llvm::Triple::nvptx:
-    case llvm::Triple::nvptx64:
-      if (Name.compare_lower("gpu"))
-        return false;
-      break;
-    case llvm::Triple::UnknownArch:
-    case llvm::Triple::arc:
-    case llvm::Triple::avr:
-    case llvm::Triple::bpfel:
-    case llvm::Triple::bpfeb:
-    case llvm::Triple::hexagon:
-    case llvm::Triple::mips:
-    case llvm::Triple::mipsel:
-    case llvm::Triple::mips64:
-    case llvm::Triple::mips64el:
-    case llvm::Triple::msp430:
-    case llvm::Triple::r600:
-    case llvm::Triple::riscv32:
-    case llvm::Triple::riscv64:
-    case llvm::Triple::sparc:
-    case llvm::Triple::sparcv9:
-    case llvm::Triple::sparcel:
-    case llvm::Triple::systemz:
-    case llvm::Triple::tce:
-    case llvm::Triple::tcele:
-    case llvm::Triple::thumb:
-    case llvm::Triple::thumbeb:
-    case llvm::Triple::xcore:
-    case llvm::Triple::le32:
-    case llvm::Triple::le64:
-    case llvm::Triple::amdil:
-    case llvm::Triple::amdil64:
-    case llvm::Triple::hsail:
-    case llvm::Triple::hsail64:
-    case llvm::Triple::spir:
-    case llvm::Triple::spir64:
-    case llvm::Triple::kalimba:
-    case llvm::Triple::shave:
-    case llvm::Triple::lanai:
-    case llvm::Triple::wasm32:
-    case llvm::Triple::wasm64:
-    case llvm::Triple::renderscript32:
-    case llvm::Triple::renderscript64:
-      return false;
-    }
-  }
-  return true;
-}
-
-bool matchesContext(CodeGenModule &CGM,
-                    const CompleteOMPContextSelectorData &ContextData) {
-  for (const OMPContextSelectorData &Data : ContextData) {
-    switch (Data.Ctx) {
-    case OMP_CTX_vendor:
-      assert(Data.CtxSet == OMP_CTX_SET_implementation &&
-             "Expected implementation context selector set.");
-      if (!checkContext<OMP_CTX_SET_implementation, OMP_CTX_vendor>(Data))
-        return false;
-      break;
-    case OMP_CTX_kind:
-      assert(Data.CtxSet == OMP_CTX_SET_device &&
-             "Expected device context selector set.");
-      if (!checkContext<OMP_CTX_SET_device, OMP_CTX_kind, CodeGenModule &>(Data,
-                                                                           CGM))
-        return false;
-      break;
-    case OMP_CTX_unknown:
-      llvm_unreachable("Unknown context selector kind.");
-    }
-  }
-  return true;
-}
-
-static CompleteOMPContextSelectorData
-translateAttrToContextSelectorData(ASTContext &C,
-                                   const OMPDeclareVariantAttr *A) {
-  CompleteOMPContextSelectorData Data;
-  for (unsigned I = 0, E = A->scores_size(); I < E; ++I) {
-    Data.emplace_back();
-    auto CtxSet = static_cast<OpenMPContextSelectorSetKind>(
-        *std::next(A->ctxSelectorSets_begin(), I));
-    auto Ctx = static_cast<OpenMPContextSelectorKind>(
-        *std::next(A->ctxSelectors_begin(), I));
-    Data.back().CtxSet = CtxSet;
-    Data.back().Ctx = Ctx;
-    const Expr *Score = *std::next(A->scores_begin(), I);
-    Data.back().Score = Score->EvaluateKnownConstInt(C);
-    switch (Ctx) {
-    case OMP_CTX_vendor:
-      assert(CtxSet == OMP_CTX_SET_implementation &&
-             "Expected implementation context selector set.");
-      Data.back().Names =
-          llvm::makeArrayRef(A->implVendors_begin(), A->implVendors_end());
-      break;
-    case OMP_CTX_kind:
-      assert(CtxSet == OMP_CTX_SET_device &&
-             "Expected device context selector set.");
-      Data.back().Names =
-          llvm::makeArrayRef(A->deviceKinds_begin(), A->deviceKinds_end());
-      break;
-    case OMP_CTX_unknown:
-      llvm_unreachable("Unknown context selector kind.");
-    }
-  }
-  return Data;
-}
-
-static bool isStrictSubset(const CompleteOMPContextSelectorData &LHS,
-                           const CompleteOMPContextSelectorData &RHS) {
-  llvm::SmallDenseMap<std::pair<int, int>, llvm::StringSet<>, 4> RHSData;
-  for (const OMPContextSelectorData &D : RHS) {
-    auto &Pair = RHSData.FindAndConstruct(std::make_pair(D.CtxSet, D.Ctx));
-    Pair.getSecond().insert(D.Names.begin(), D.Names.end());
-  }
-  bool AllSetsAreEqual = true;
-  for (const OMPContextSelectorData &D : LHS) {
-    auto It = RHSData.find(std::make_pair(D.CtxSet, D.Ctx));
-    if (It == RHSData.end())
-      return false;
-    if (D.Names.size() > It->getSecond().size())
-      return false;
-    if (llvm::set_union(It->getSecond(), D.Names))
-      return false;
-    AllSetsAreEqual =
-        AllSetsAreEqual && (D.Names.size() == It->getSecond().size());
-  }
-
-  return LHS.size() != RHS.size() || !AllSetsAreEqual;
-}
-
-static bool greaterCtxScore(const CompleteOMPContextSelectorData &LHS,
-                            const CompleteOMPContextSelectorData &RHS) {
-  // Score is calculated as sum of all scores + 1.
-  llvm::APSInt LHSScore(llvm::APInt(64, 1), /*isUnsigned=*/false);
-  bool RHSIsSubsetOfLHS = isStrictSubset(RHS, LHS);
-  if (RHSIsSubsetOfLHS) {
-    LHSScore = llvm::APSInt::get(0);
-  } else {
-    for (const OMPContextSelectorData &Data : LHS) {
-      if (Data.Score.getBitWidth() > LHSScore.getBitWidth()) {
-        LHSScore = LHSScore.extend(Data.Score.getBitWidth()) + Data.Score;
-      } else if (Data.Score.getBitWidth() < LHSScore.getBitWidth()) {
-        LHSScore += Data.Score.extend(LHSScore.getBitWidth());
-      } else {
-        LHSScore += Data.Score;
-      }
-    }
-  }
-  llvm::APSInt RHSScore(llvm::APInt(64, 1), /*isUnsigned=*/false);
-  if (!RHSIsSubsetOfLHS && isStrictSubset(LHS, RHS)) {
-    RHSScore = llvm::APSInt::get(0);
-  } else {
-    for (const OMPContextSelectorData &Data : RHS) {
-      if (Data.Score.getBitWidth() > RHSScore.getBitWidth()) {
-        RHSScore = RHSScore.extend(Data.Score.getBitWidth()) + Data.Score;
-      } else if (Data.Score.getBitWidth() < RHSScore.getBitWidth()) {
-        RHSScore += Data.Score.extend(RHSScore.getBitWidth());
-      } else {
-        RHSScore += Data.Score;
-      }
-    }
-  }
-  return llvm::APSInt::compareValues(LHSScore, RHSScore) >= 0;
-}
-
 /// Finds the variant function that matches current context with its context
 /// selector.
 static const FunctionDecl *getDeclareVariantFunction(CodeGenModule &CGM,
                                                      const FunctionDecl *FD) {
   if (!FD->hasAttrs() || !FD->hasAttr<OMPDeclareVariantAttr>())
     return FD;
   // Iterate through all DeclareVariant attributes and check context selectors.
   const OMPDeclareVariantAttr *TopMostAttr = nullptr;
-  CompleteOMPContextSelectorData TopMostData;
-  for (const auto *A : FD->specific_attrs<OMPDeclareVariantAttr>()) {
-    CompleteOMPContextSelectorData Data =
-        translateAttrToContextSelectorData(CGM.getContext(), A);
-    if (!matchesContext(CGM, Data))
-      continue;
-    // If the attribute matches the context, find the attribute with the highest
-    // score.
-    if (!TopMostAttr || !greaterCtxScore(TopMostData, Data)) {
-      TopMostAttr = A;
-      TopMostData.swap(Data);
-    }
-  }
+  for (const auto *A : FD->specific_attrs<OMPDeclareVariantAttr>())
+    TopMostAttr = getBetterOpenMPContextMatch(CGM.getContext(), TopMostAttr, A);
   if (!TopMostAttr)
     return FD;
   return cast<FunctionDecl>(
       cast<DeclRefExpr>(TopMostAttr->getVariantFuncRef()->IgnoreParenImpCasts())
           ->getDecl());
 }
 
 bool CGOpenMPRuntime::emitDeclareVariant(GlobalDecl GD, bool IsForDefinition) {

clang/lib/CodeGen/CodeGenModule.cpp

static void AppendTargetMangling(const CodeGenModule &CGM,
   for (StringRef Feat : Info.Features) {
     if (!IsFirst)
       Out << '_';
     IsFirst = false;
     Out << Feat.substr(1);
   }
 }
 
+static void AppendOpenMPVariantMangling(const CodeGenModule &CGM,
+                                        const FunctionDecl *FD,
+                                        raw_ostream &Out) {
+  for (const OMPDeclareVariantAttr *Attr :
+       FD->specific_attrs<OMPDeclareVariantAttr>()) {
+    if (Attr->isInherited())
+      continue;
+    // TODO: Mangle the name based on the context
+    Out << ".ompvariant";
+  }
+}
+
 static std::string getMangledNameImpl(const CodeGenModule &CGM, GlobalDecl GD,
                                       const NamedDecl *ND,
                                       bool OmitMultiVersionMangling = false) {
   SmallString<256> Buffer;
   llvm::raw_svector_ostream Out(Buffer);
   MangleContext &MC = CGM.getCXXABI().getMangleContext();
   if (MC.shouldMangleDeclName(ND)) {
     llvm::raw_svector_ostream Out(Buffer);
if (FD->isMultiVersion() && !OmitMultiVersionMangling) {
       case MultiVersionKind::CPUSpecific:
         AppendCPUSpecificCPUDispatchMangling(CGM,
                                              FD->getAttr<CPUSpecificAttr>(),
                                              GD.getMultiVersionIndex(), Out);
         break;
       case MultiVersionKind::Target:
         AppendTargetMangling(CGM, FD->getAttr<TargetAttr>(), Out);
         break;
+      case MultiVersionKind::OMPVariant:
+        AppendOpenMPVariantMangling(CGM, FD, Out);
+        break;
       case MultiVersionKind::None:
         llvm_unreachable("None multiversion type isn't valid here");
       }
     }
 
   return Out.str();
 }
 
void CodeGenModule::EmitGlobal(GlobalDecl GD) {
   // function. If the global must always be emitted, do it eagerly if possible
   // to benefit from cache locality.
   if (MustBeEmitted(Global) && MayBeEmittedEagerly(Global)) {
     // Emit the definition if it can't be deferred.
     EmitGlobalDefinition(GD);
     return;
   }
 
-    // Check if this must be emitted as declare variant.
+  // TODO: See the TODO at the other emitDeclareVariant call.
   if (LangOpts.OpenMP && isa<FunctionDecl>(Global) && OpenMPRuntime &&
       OpenMPRuntime->emitDeclareVariant(GD, /*IsForDefinition=*/false))
     return;
 
   // If we're deferring emission of a C++ variable with an
   // initializer, remember the order in which it appeared in the file.
   if (getLangOpts().CPlusPlus && isa<VarDecl>(Global) &&
       cast<VarDecl>(Global)->hasInit()) {
Priority = std::max(
         Priority, TI.multiVersionSortPriority(RO.Conditions.Architecture));
   return Priority;
 }
 
 void CodeGenModule::emitMultiVersionFunctions() {
   for (GlobalDecl GD : MultiVersionFuncs) {
     SmallVector<CodeGenFunction::MultiVersionResolverOption, 10> Options;
     const FunctionDecl *FD = cast<FunctionDecl>(GD.getDecl());
+    // OpenMP multi versioning is (for now) resolved at compile time, no
+    // resolver function necessary (yet).
+    if (FD->isOpenMPMultiVersion())
+      continue;
     getContext().forEachMultiversionedFunctionVersion(
         FD, [this, &GD, &Options](const FunctionDecl *CurFD) {
           GlobalDecl CurGD{
               (CurFD->isDefined() ? CurFD->getDefinition() : CurFD)};
           StringRef MangledName = getMangledName(CurGD);
           llvm::Constant *Func = GetGlobalValue(MangledName);
           if (!Func) {
             if (CurFD->isDefined()) {
if (getLangOpts().OpenMPIsDevice && OpenMPRuntime &&
           GDDef = GlobalDecl(DD, GD.getDtorType());
         else
           GDDef = GlobalDecl(FDDef);
         EmitGlobal(GDDef);
       }
     }
     // Check if this must be emitted as declare variant and emit reference to
     // the the declare variant function.
+    // TODO: We should introduce function aliases for `omp declare variant`
+    //       directives such that we can treat them through the same overload
+    //       resolution scheme (via multi versioning) as `omp begin declare
+    //       variant` functions. For an `omp declare variant(VARIANT) ...`
+    //       that is attached to a BASE function we would create a global alias
+    //       VARIANT = BASE which will participate in the multi version overload
+    //       resolution. If picked, here is no need to emit them explicitly.
     if (LangOpts.OpenMP && OpenMPRuntime)
       (void)OpenMPRuntime->emitDeclareVariant(GD, /*IsForDefinition=*/true);
 
-    if (FD->isMultiVersion()) {
+    if (FD->isMultiVersion() && !FD->isOpenMPMultiVersion()) {
       const auto *TA = FD->getAttr<TargetAttr>();
       if (TA && TA->isDefaultVersion())
         UpdateMultiVersionNames(GD, FD);
       if (!IsForDefinition)
         return GetOrCreateMultiVersionResolver(GD, Ty, FD);
     }
   }
 
if (VD->getDefinition() && TSK == TSK_ExplicitInstantiationDefinition)
     GetAddrOfGlobalVar(VD);
 
   EmitTopLevelDecl(VD);
 }
 
 void CodeGenModule::EmitGlobalFunctionDefinition(GlobalDecl GD,
                                                  llvm::GlobalValue *GV) {
   // Check if this must be emitted as declare variant.
+  // TODO: See the TODO at the other emitDeclareVariant call.
   if (LangOpts.OpenMP && OpenMPRuntime &&
       OpenMPRuntime->emitDeclareVariant(GD, /*IsForDefinition=*/true))
     return;
 
   const auto *D = cast<FunctionDecl>(GD.getDecl());
 
   // Compute the function info and LLVM type.
   const CGFunctionInfo &FI = getTypes().arrangeGlobalDeclaration(GD);

clang/lib/Headers/__clang_cuda_cmath.h

 // these functions in namespace std, because our device functions need to
 // overload the standard library functions, which may be declared in the global
 // namespace or in std, depending on the degree of conformance of the stdlib
 // implementation.  Declaring in the global namespace and pulling into namespace
 // std covers all of the known knowns.
 
 #ifdef _OPENMP
 #define __DEVICE__ static __attribute__((always_inline))
+#pragma omp begin declare variant match(device = {kind(gpu)})
 #else
 #define __DEVICE__ static __device__ __inline__ __attribute__((always_inline))
 #endif
 
-// For C++ 17 we need to include noexcept attribute to be compatible
-// with the header-defined version. This may be removed once
-// variant is supported.
-#if defined(_OPENMP) && defined(__cplusplus) && __cplusplus >= 201703L
-#define __NOEXCEPT noexcept
-#else
-#define __NOEXCEPT
-#endif
-
-#if !(defined(_OPENMP) && defined(__cplusplus))
 __DEVICE__ long long abs(long long __n) { return ::llabs(__n); }
 __DEVICE__ long abs(long __n) { return ::labs(__n); }
 __DEVICE__ float abs(float __x) { return ::fabsf(__x); }
 __DEVICE__ double abs(double __x) { return ::fabs(__x); }
-#endif
-// TODO: remove once variat is supported.
-#if defined(_OPENMP) && defined(__cplusplus)
-__DEVICE__ const float abs(const float __x) { return ::fabsf((float)__x); }
-__DEVICE__ const double abs(const double __x) { return ::fabs((double)__x); }
-#endif
 __DEVICE__ float acos(float __x) { return ::acosf(__x); }
 __DEVICE__ float asin(float __x) { return ::asinf(__x); }
 __DEVICE__ float atan(float __x) { return ::atanf(__x); }
 __DEVICE__ float atan2(float __x, float __y) { return ::atan2f(__x, __y); }
 __DEVICE__ float ceil(float __x) { return ::ceilf(__x); }
 __DEVICE__ float cos(float __x) { return ::cosf(__x); }
 __DEVICE__ float cosh(float __x) { return ::coshf(__x); }
 __DEVICE__ float exp(float __x) { return ::expf(__x); }
-__DEVICE__ float fabs(float __x) __NOEXCEPT { return ::fabsf(__x); }
+__DEVICE__ float fabs(float __x) { return ::fabsf(__x); }
 __DEVICE__ float floor(float __x) { return ::floorf(__x); }
 __DEVICE__ float fmod(float __x, float __y) { return ::fmodf(__x, __y); }
-// TODO: remove when variant is supported
-#ifndef _OPENMP
-__DEVICE__ int fpclassify(float __x) {
-  return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
-                              FP_ZERO, __x);
-}
-__DEVICE__ int fpclassify(double __x) {
-  return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
-                              FP_ZERO, __x);
-}
-#endif
 __DEVICE__ float frexp(float __arg, int *__exp) {
   return ::frexpf(__arg, __exp);
 }
 
 // For inscrutable reasons, the CUDA headers define these functions for us on
 // Windows.
 #ifndef _MSC_VER
 __DEVICE__ bool isinf(float __x) { return ::__isinff(__x); }

 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, exp2)
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, expm1)
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, fabs)
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(double, fdim);
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(double, floor)
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(double, fmax);
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(double, fmin);
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(double, fmod);
-__CUDA_CLANG_FN_INTEGER_OVERLOAD_1(int, fpclassify)
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(double, hypot);
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(int, ilogb)
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(bool, isfinite)
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(bool, isgreater);
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(bool, isgreaterequal);
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_1(bool, isinf);
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(bool, isless);
 __CUDA_CLANG_FN_INTEGER_OVERLOAD_2(bool, islessequal);

 using ::expm1;
 using ::fabs;
 using ::fdim;
 using ::floor;
 using ::fma;
 using ::fmax;
 using ::fmin;
 using ::fmod;
-using ::fpclassify;
 using ::frexp;
 using ::hypot;
 using ::ilogb;
 using ::isfinite;
 using ::isgreater;
 using ::isgreaterequal;
 using ::isless;
 using ::islessequal;

 using ::modff;
 using ::nearbyintf;
 using ::nextafterf;
 using ::powf;
 using ::remainderf;
 using ::remquof;
 using ::rintf;
 using ::roundf;
-// TODO: remove once variant is supported
-#ifndef _OPENMP
-using ::scalblnf;
-#endif
 using ::scalbnf;
 using ::sinf;
 using ::sinhf;
 using ::sqrtf;
 using ::tanf;
 using ::tanhf;
 using ::tgammaf;
 using ::truncf;
 
 #ifdef _LIBCPP_END_NAMESPACE_STD
 _LIBCPP_END_NAMESPACE_STD
 #else
 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
 _GLIBCXX_END_NAMESPACE_VERSION
 #endif
 } // namespace std
 #endif
 
+#ifdef _OPENMP
+#pragma omp end declare variant
+#endif
+
 #undef __NOEXCEPT
 #undef __DEVICE__
 
 #endif

clang/lib/Headers/__clang_cuda_device_functions.h

 // -ffast-math or -fcuda-approx-transcendentals are in effect.
 #pragma push_macro("__FAST_OR_SLOW")
 #if defined(__CLANG_CUDA_APPROX_TRANSCENDENTALS__)
 #define __FAST_OR_SLOW(fast, slow) fast
 #else
 #define __FAST_OR_SLOW(fast, slow) slow
 #endif
 
-// For C++ 17 we need to include noexcept attribute to be compatible
-// with the header-defined version. This may be removed once
-// variant is supported.
-#if defined(_OPENMP) && defined(__cplusplus) && __cplusplus >= 201703L
-#define __NOEXCEPT noexcept
-#else
-#define __NOEXCEPT
-#endif
-
 __DEVICE__ int __all(int __a) { return __nvvm_vote_all(__a); }
 __DEVICE__ int __any(int __a) { return __nvvm_vote_any(__a); }
 __DEVICE__ unsigned int __ballot(int __a) { return __nvvm_vote_ballot(__a); }
 __DEVICE__ unsigned int __brev(unsigned int __a) { return __nv_brev(__a); }
 __DEVICE__ unsigned long long __brevll(unsigned long long __a) {
   return __nv_brevll(__a);
 }
-#if defined(__cplusplus)
 __DEVICE__ void __brkpt() { asm volatile("brkpt;"); }
 __DEVICE__ void __brkpt(int __a) { __brkpt(); }
-#else
-__DEVICE__ void __attribute__((overloadable)) __brkpt(void) { asm volatile("brkpt;"); }
-__DEVICE__ void __attribute__((overloadable)) __brkpt(int __a) { __brkpt(); }
-#endif
 __DEVICE__ unsigned int __byte_perm(unsigned int __a, unsigned int __b,
                                     unsigned int __c) {
   return __nv_byte_perm(__a, __b, __c);
 }
 __DEVICE__ int __clz(int __a) { return __nv_clz(__a); }
 __DEVICE__ int __clzll(long long __a) { return __nv_clzll(__a); }
 __DEVICE__ float __cosf(float __a) { return __nv_fast_cosf(__a); }
 __DEVICE__ double __dAtomicAdd(double *__p, double __v) {

 __DEVICE__ unsigned int __vsubus4(unsigned int __a, unsigned int __b) {
   unsigned int r;
   asm("vsub4.u32.u32.u32.sat %0,%1,%2,%3;"
       : "=r"(r)
       : "r"(__a), "r"(__b), "r"(0));
   return r;
 }
 #endif // CUDA_VERSION >= 9020
-__DEVICE__ int abs(int __a) __NOEXCEPT { return __nv_abs(__a); }
-__DEVICE__ double fabs(double __a) __NOEXCEPT { return __nv_fabs(__a); }
+__DEVICE__ int abs(int __a) { return __nv_abs(__a); }
+__DEVICE__ double fabs(double __a) { return __nv_fabs(__a); }
 __DEVICE__ double acos(double __a) { return __nv_acos(__a); }
 __DEVICE__ float acosf(float __a) { return __nv_acosf(__a); }
 __DEVICE__ double acosh(double __a) { return __nv_acosh(__a); }
 __DEVICE__ float acoshf(float __a) { return __nv_acoshf(__a); }
 __DEVICE__ double asin(double __a) { return __nv_asin(__a); }
 __DEVICE__ float asinf(float __a) { return __nv_asinf(__a); }
 __DEVICE__ double asinh(double __a) { return __nv_asinh(__a); }
 __DEVICE__ float asinhf(float __a) { return __nv_asinhf(__a); }

 __DEVICE__ int ilogbf(float __a) { return __nv_ilogbf(__a); }
 __DEVICE__ double j0(double __a) { return __nv_j0(__a); }
 __DEVICE__ float j0f(float __a) { return __nv_j0f(__a); }
 __DEVICE__ double j1(double __a) { return __nv_j1(__a); }
 __DEVICE__ float j1f(float __a) { return __nv_j1f(__a); }
 __DEVICE__ double jn(int __n, double __a) { return __nv_jn(__n, __a); }
 __DEVICE__ float jnf(int __n, float __a) { return __nv_jnf(__n, __a); }
 #if defined(__LP64__) || defined(_WIN64)
-__DEVICE__ long labs(long __a) __NOEXCEPT { return __nv_llabs(__a); };
+__DEVICE__ long labs(long __a) { return __nv_llabs(__a); };
 #else
-__DEVICE__ long labs(long __a) __NOEXCEPT { return __nv_abs(__a); };
+__DEVICE__ long labs(long __a) { return __nv_abs(__a); };
 #endif
 __DEVICE__ double ldexp(double __a, int __b) { return __nv_ldexp(__a, __b); }
 __DEVICE__ float ldexpf(float __a, int __b) { return __nv_ldexpf(__a, __b); }
 __DEVICE__ double lgamma(double __a) { return __nv_lgamma(__a); }
 __DEVICE__ float lgammaf(float __a) { return __nv_lgammaf(__a); }
-__DEVICE__ long long llabs(long long __a) __NOEXCEPT { return __nv_llabs(__a); }
+__DEVICE__ long long llabs(long long __a) { return __nv_llabs(__a); }
 __DEVICE__ long long llmax(long long __a, long long __b) {
   return __nv_llmax(__a, __b);
 }
 __DEVICE__ long long llmin(long long __a, long long __b) {
   return __nv_llmin(__a, __b);
 }
 __DEVICE__ long long llrint(double __a) { return __nv_llrint(__a); }
 __DEVICE__ long long llrintf(float __a) { return __nv_llrintf(__a); }
__DEVICE__ float rnormf(int __dim, const float *__t) {
   return __nv_rnormf(__dim, __t);
 }
 __DEVICE__ double round(double __a) { return __nv_round(__a); }
 __DEVICE__ float roundf(float __a) { return __nv_roundf(__a); }
 __DEVICE__ double rsqrt(double __a) { return __nv_rsqrt(__a); }
 __DEVICE__ float rsqrtf(float __a) { return __nv_rsqrtf(__a); }
 __DEVICE__ double scalbn(double __a, int __b) { return __nv_scalbn(__a, __b); }
 __DEVICE__ float scalbnf(float __a, int __b) { return __nv_scalbnf(__a, __b); }
-// TODO: remove once variant is supported
-#ifndef _OPENMP
-__DEVICE__ double scalbln(double __a, long __b) {
-  if (__b > INT_MAX)
-    return __a > 0 ? HUGE_VAL : -HUGE_VAL;
-  if (__b < INT_MIN)
-    return __a > 0 ? 0.0 : -0.0;
-  return scalbn(__a, (int)__b);
-}
-__DEVICE__ float scalblnf(float __a, long __b) {
-  if (__b > INT_MAX)
-    return __a > 0 ? HUGE_VALF : -HUGE_VALF;
-  if (__b < INT_MIN)
-    return __a > 0 ? 0.f : -0.f;
-  return scalbnf(__a, (int)__b);
-}
-#endif
 __DEVICE__ double sin(double __a) { return __nv_sin(__a); }
 __DEVICE__ void sincos(double __a, double *__s, double *__c) {
   return __nv_sincos(__a, __s, __c);
 }
 __DEVICE__ void sincosf(float __a, float *__s, float *__c) {
   return __FAST_OR_SLOW(__nv_fast_sincosf, __nv_sincosf)(__a, __s, __c);
 }
 __DEVICE__ void sincospi(double __a, double *__s, double *__c) {

 }
 __DEVICE__ double y0(double __a) { return __nv_y0(__a); }
 __DEVICE__ float y0f(float __a) { return __nv_y0f(__a); }
 __DEVICE__ double y1(double __a) { return __nv_y1(__a); }
 __DEVICE__ float y1f(float __a) { return __nv_y1f(__a); }
 __DEVICE__ double yn(int __a, double __b) { return __nv_yn(__a, __b); }
 __DEVICE__ float ynf(int __a, float __b) { return __nv_ynf(__a, __b); }
 
-#undef __NOEXCEPT
 #pragma pop_macro("__DEVICE__")
 #pragma pop_macro("__FAST_OR_SLOW")
+
 #endif // __CLANG_CUDA_DEVICE_FUNCTIONS_H__

clang/lib/Headers/__clang_cuda_math_forward_declares.h

 #pragma push_macro("__DEVICE__")
 #ifdef _OPENMP
 #define __DEVICE__ static __inline__ __attribute__((always_inline))
 #else
 #define __DEVICE__                                                             \
   static __inline__ __attribute__((always_inline)) __attribute__((device))
 #endif
 
-// For C++ 17 we need to include noexcept attribute to be compatible
-// with the header-defined version. This may be removed once
-// variant is supported.
-#if defined(_OPENMP) && defined(__cplusplus) && __cplusplus >= 201703L
-#define __NOEXCEPT noexcept
-#else
-#define __NOEXCEPT
-#endif
-
-#if !(defined(_OPENMP) && defined(__cplusplus))
 __DEVICE__ long abs(long);
 __DEVICE__ long long abs(long long);
-__DEVICE__ double abs(double);
-__DEVICE__ float abs(float);
-#endif
-// While providing the CUDA declarations and definitions for math functions,
-// we may manually define additional functions.
-// TODO: Once variant is supported the additional functions will have
-// to be removed.
-#if defined(_OPENMP) && defined(__cplusplus)
-__DEVICE__ const double abs(const double);
-__DEVICE__ const float abs(const float);
-#endif
-__DEVICE__ int abs(int) __NOEXCEPT;
 __DEVICE__ double acos(double);
 __DEVICE__ float acos(float);
 __DEVICE__ double acosh(double);
 __DEVICE__ float acosh(float);
 __DEVICE__ double asin(double);
 __DEVICE__ float asin(float);
 __DEVICE__ double asinh(double);
 __DEVICE__ float asinh(float);

 __DEVICE__ double erf(double);
 __DEVICE__ float erf(float);
 __DEVICE__ double exp2(double);
 __DEVICE__ float exp2(float);
 __DEVICE__ double exp(double);
 __DEVICE__ float exp(float);
 __DEVICE__ double expm1(double);
 __DEVICE__ float expm1(float);
-__DEVICE__ double fabs(double) __NOEXCEPT;
-__DEVICE__ float fabs(float) __NOEXCEPT;
+__DEVICE__ double fabs(double);
+__DEVICE__ float fabs(float);
 __DEVICE__ double fdim(double, double);
 __DEVICE__ float fdim(float, float);
 __DEVICE__ double floor(double);
 __DEVICE__ float floor(float);
 __DEVICE__ double fma(double, double, double);
 __DEVICE__ float fma(float, float, float);
 __DEVICE__ double fmax(double, double);
 __DEVICE__ float fmax(float, float);
 __DEVICE__ double fmin(double, double);
 __DEVICE__ float fmin(float, float);
 __DEVICE__ double fmod(double, double);
 __DEVICE__ float fmod(float, float);
-__DEVICE__ int fpclassify(double);
-__DEVICE__ int fpclassify(float);
 __DEVICE__ double frexp(double, int *);
 __DEVICE__ float frexp(float, int *);
 __DEVICE__ double hypot(double, double);
 __DEVICE__ float hypot(float, float);
 __DEVICE__ int ilogb(double);
 __DEVICE__ int ilogb(float);
 #ifdef _MSC_VER
 __DEVICE__ bool isfinite(long double);

 __DEVICE__ bool isnan(long double);
 #endif
 __DEVICE__ bool isnan(double);
 __DEVICE__ bool isnan(float);
 __DEVICE__ bool isnormal(double);
 __DEVICE__ bool isnormal(float);
 __DEVICE__ bool isunordered(double, double);
 __DEVICE__ bool isunordered(float, float);
-__DEVICE__ long labs(long) __NOEXCEPT;
+__DEVICE__ long labs(long);
 __DEVICE__ double ldexp(double, int);
 __DEVICE__ float ldexp(float, int);
 __DEVICE__ double lgamma(double);
 __DEVICE__ float lgamma(float);
-__DEVICE__ long long llabs(long long) __NOEXCEPT;
+__DEVICE__ long long llabs(long long);
 __DEVICE__ long long llrint(double);
 __DEVICE__ long long llrint(float);
 __DEVICE__ double log10(double);
 __DEVICE__ float log10(float);
 __DEVICE__ double log1p(double);
 __DEVICE__ float log1p(float);
 __DEVICE__ double log2(double);
 __DEVICE__ float log2(float);
 __DEVICE__ double logb(double);
 __DEVICE__ float logb(float);
-#if defined(_OPENMP) && defined(__cplusplus)
 __DEVICE__ long double log(long double);
-#endif
 __DEVICE__ double log(double);
 __DEVICE__ float log(float);
 __DEVICE__ long lrint(double);
 __DEVICE__ long lrint(float);
 __DEVICE__ long lround(double);
 __DEVICE__ long lround(float);
 __DEVICE__ long long llround(float); // No llround(double).
 __DEVICE__ double modf(double, double *);

 using ::expm1;
 using ::fabs;
 using ::fdim;
 using ::floor;
 using ::fma;
 using ::fmax;
 using ::fmin;
 using ::fmod;
-using ::fpclassify;
 using ::frexp;
 using ::hypot;
 using ::ilogb;
 using ::isfinite;
 using ::isgreater;
 using ::isgreaterequal;
 using ::isinf;
 using ::isless;

 _LIBCPP_END_NAMESPACE_STD
 #else
 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
 _GLIBCXX_END_NAMESPACE_VERSION
 #endif
 } // namespace std
 #endif
 
-#undef __NOEXCEPT
 #pragma pop_macro("__DEVICE__")
 
 #endif

clang/lib/Headers/openmp_wrappers/__clang_openmp_math.h

 /*===---- __clang_openmp_math.h - OpenMP target math support ---------------===
  *
  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
  * See https://llvm.org/LICENSE.txt for license information.
  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
  *
  *===-----------------------------------------------------------------------===
  */
 
 #if defined(__NVPTX__) && defined(_OPENMP)
-/// TODO:
-/// We are currently reusing the functionality of the Clang-CUDA code path
-/// as an alternative to the host declarations provided by math.h and cmath.
-/// This is suboptimal.
-///
-/// We should instead declare the device functions in a similar way, e.g.,
-/// through OpenMP 5.0 variants, and afterwards populate the module with the
-/// host declarations by unconditionally including the host math.h or cmath,
-/// respectively. This is actually what the Clang-CUDA code path does, using
-/// __device__ instead of variants to avoid redeclarations and get the desired
-/// overload resolution.
 
 #define __CUDA__
 
 #if defined(__cplusplus)
   #include <__clang_cuda_cmath.h>
 #endif
 
 #undef __CUDA__
 
-/// Magic macro for stopping the math.h/cmath host header from being included.
-#define __CLANG_NO_HOST_MATH__
-
 #endif
 

clang/lib/Headers/openmp_wrappers/__clang_openmp_math_declares.h

 /*===---- __clang_openmp_math_declares.h - OpenMP math declares ------------===
  *
  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
  * See https://llvm.org/LICENSE.txt for license information.
  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
  *
  *===-----------------------------------------------------------------------===
  */
 
 #ifndef __CLANG_OPENMP_MATH_DECLARES_H__
 #define __CLANG_OPENMP_MATH_DECLARES_H__
 
 #ifndef _OPENMP
 #error "This file is for OpenMP compilation only."
 #endif
 
 #if defined(__NVPTX__) && defined(_OPENMP)
 
 #define __CUDA__
 
+#pragma omp begin declare variant match(device = {kind(gpu)})
+
 #if defined(__cplusplus)
   #include <__clang_cuda_math_forward_declares.h>
 #endif
 
 /// Include declarations for libdevice functions.
 #include <__clang_cuda_libdevice_declares.h>
 /// Provide definitions for these functions.
 #include <__clang_cuda_device_functions.h>
 
+#pragma omp end declare variant
+
 #undef __CUDA__
 
 #endif
 #endif

clang/lib/Headers/openmp_wrappers/cmath

 /*===-------------- cmath - Alternative cmath header -----------------------===
  *
  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
  * See https://llvm.org/LICENSE.txt for license information.
  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
  *
  *===-----------------------------------------------------------------------===
  */
 
 #include <__clang_openmp_math.h>
 
-#ifndef __CLANG_NO_HOST_MATH__
 #include_next <cmath>
-#else
-#undef __CLANG_NO_HOST_MATH__
-#endif

clang/lib/Headers/openmp_wrappers/math.h

 /*===------------- math.h - Alternative math.h header ----------------------===
  *
  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
  * See https://llvm.org/LICENSE.txt for license information.
  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
  *
  *===-----------------------------------------------------------------------===
  */
 
 #include <__clang_openmp_math.h>
 
-#ifndef __CLANG_NO_HOST_MATH__
 #include_next <math.h>
-#else
-#undef __CLANG_NO_HOST_MATH__
-#endif
-

clang/lib/Parse/ParseOpenMP.cpp

enum OpenMPDirectiveKindEx {
   OMPD_target_enter,
   OMPD_target_exit,
   OMPD_update,
   OMPD_distribute_parallel,
   OMPD_teams_distribute_parallel,
   OMPD_target_teams_distribute_parallel,
   OMPD_mapper,
   OMPD_variant,
+  OMPD_begin,
+  OMPD_begin_declare,
 };
 
 class DeclDirectiveListParserHelper final {
   SmallVector<Expr *, 4> Identifiers;
   Parser *P;
   OpenMPDirectiveKind Kind;
 
 public:
return llvm::StringSwitch<unsigned>(S)
       .Case("end", OMPD_end)
       .Case("enter", OMPD_enter)
       .Case("exit", OMPD_exit)
       .Case("point", OMPD_point)
       .Case("reduction", OMPD_reduction)
       .Case("update", OMPD_update)
       .Case("mapper", OMPD_mapper)
       .Case("variant", OMPD_variant)
+      .Case("begin", OMPD_begin)
       .Default(OMPD_unknown);
 }
 
 static OpenMPDirectiveKind parseOpenMPDirectiveKind(Parser &P) {
   // Array of foldings: F[i][0] F[i][1] ===> F[i][2].
   // E.g.: OMPD_for OMPD_simd ===> OMPD_for_simd
   // TODO: add other combined directives in topological order.
   static const unsigned F[][3] = {
+      {OMPD_begin, OMPD_declare, OMPD_begin_declare},
+      {OMPD_end, OMPD_declare, OMPD_end_declare},
       {OMPD_cancellation, OMPD_point, OMPD_cancellation_point},
       {OMPD_declare, OMPD_reduction, OMPD_declare_reduction},
       {OMPD_declare, OMPD_mapper, OMPD_declare_mapper},
       {OMPD_declare, OMPD_simd, OMPD_declare_simd},
       {OMPD_declare, OMPD_target, OMPD_declare_target},
       {OMPD_declare, OMPD_variant, OMPD_declare_variant},
+      {OMPD_begin_declare, OMPD_variant, OMPD_begin_declare_variant},
+      {OMPD_end_declare, OMPD_variant, OMPD_end_declare_variant},
       {OMPD_distribute, OMPD_parallel, OMPD_distribute_parallel},
       {OMPD_distribute_parallel, OMPD_for, OMPD_distribute_parallel_for},
       {OMPD_distribute_parallel_for, OMPD_simd,
        OMPD_distribute_parallel_for_simd},
       {OMPD_distribute, OMPD_simd, OMPD_distribute_simd},
-      {OMPD_end, OMPD_declare, OMPD_end_declare},
       {OMPD_end_declare, OMPD_target, OMPD_end_declare_target},
       {OMPD_target, OMPD_data, OMPD_target_data},
       {OMPD_target, OMPD_enter, OMPD_target_enter},
       {OMPD_target, OMPD_exit, OMPD_target_exit},
       {OMPD_target, OMPD_update, OMPD_target_update},
       {OMPD_target_enter, OMPD_data, OMPD_target_enter_data},
       {OMPD_target_exit, OMPD_data, OMPD_target_exit_data},
       {OMPD_for, OMPD_simd, OMPD_for_simd},
// Unknown selector - just ignore it completely.
       (void)TBr.consumeClose();
     }
     // Consume ','
     if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::annot_pragma_openmp_end))
       (void)ExpectAndConsume(tok::comma);
   } while (Tok.isAnyIdentifier());
   return false;
 }
 
-/// Parse clauses for '#pragma omp declare variant ( variant-func-id ) clause'.
-void Parser::ParseOMPDeclareVariantClauses(Parser::DeclGroupPtrTy Ptr,
-                                           CachedTokens &Toks,
-                                           SourceLocation Loc) {
-  PP.EnterToken(Tok, /*IsReinject*/ true);
-  PP.EnterTokenStream(Toks, /*DisableMacroExpansion=*/true,
-                      /*IsReinject*/ true);
-  // Consume the previously pushed token.
-  ConsumeAnyToken(/*ConsumeCodeCompletionTok=*/true);
-  ConsumeAnyToken(/*ConsumeCodeCompletionTok=*/true);
-
-  FNContextRAII FnContext(*this, Ptr);
-  // Parse function declaration id.
-  SourceLocation RLoc;
-  // Parse with IsAddressOfOperand set to true to parse methods as DeclRefExprs
-  // instead of MemberExprs.
-  ExprResult AssociatedFunction =
-      ParseOpenMPParensExpr(getOpenMPDirectiveName(OMPD_declare_variant), RLoc,
-                            /*IsAddressOfOperand=*/true);
-  if (!AssociatedFunction.isUsable()) {
-    if (!Tok.is(tok::annot_pragma_openmp_end))
-      while (!SkipUntil(tok::annot_pragma_openmp_end, StopBeforeMatch))
-        ;
-    // Skip the last annot_pragma_openmp_end.
-    (void)ConsumeAnnotationToken();
-    return;
-  }
-  Optional<std::pair<FunctionDecl *, Expr *>> DeclVarData =
-      Actions.checkOpenMPDeclareVariantFunction(
-          Ptr, AssociatedFunction.get(), SourceRange(Loc, Tok.getLocation()));
-
+void Parser::ParseOMPDeclareVariantMatchClause(
+    SourceLocation Loc, SmallVectorImpl<Sema::OMPCtxSelectorData> &Data) {
   // Parse 'match'.
   OpenMPClauseKind CKind = Tok.isAnnotation()
                                ? OMPC_unknown
                                : getOpenMPClauseKind(PP.getSpelling(Tok));
   if (CKind != OMPC_match) {
     Diag(Tok.getLocation(), diag::err_omp_declare_variant_wrong_clause)
         << getOpenMPClauseName(OMPC_match);
     while (!SkipUntil(tok::annot_pragma_openmp_end, Parser::StopBeforeMatch))
if (T.expectAndConsume(diag::err_expected_lparen_after,
     while (!SkipUntil(tok::annot_pragma_openmp_end, StopBeforeMatch))
       ;
     // Skip the last annot_pragma_openmp_end.
     (void)ConsumeAnnotationToken();
     return;
   }
 
   // Parse inner context selectors.
-  SmallVector<Sema::OMPCtxSelectorData, 4> Data;
   if (!parseOpenMPContextSelectors(Loc, Data)) {
     // Parse ')'.
     (void)T.consumeClose();
     // Need to check for extra tokens.
     if (Tok.isNot(tok::annot_pragma_openmp_end)) {
       Diag(Tok, diag::warn_omp_extra_tokens_at_eol)
           << getOpenMPDirectiveName(OMPD_declare_variant);
     }
   }
+}
+
+/// Parse clauses for '#pragma omp declare variant ( variant-func-id ) clause'.
+void Parser::ParseOMPDeclareVariantClauses(Parser::DeclGroupPtrTy Ptr,
+                                           CachedTokens &Toks,
+                                           SourceLocation Loc) {
+  PP.EnterToken(Tok, /*IsReinject*/ true);
+  PP.EnterTokenStream(Toks, /*DisableMacroExpansion=*/true,
+                      /*IsReinject*/ true);
+  // Consume the previously pushed token.
+  ConsumeAnyToken(/*ConsumeCodeCompletionTok=*/true);
+  ConsumeAnyToken(/*ConsumeCodeCompletionTok=*/true);
+
+  FNContextRAII FnContext(*this, Ptr);
+  // Parse function declaration id.
+  SourceLocation RLoc;
+  // Parse with IsAddressOfOperand set to true to parse methods as DeclRefExprs
+  // instead of MemberExprs.
+  ExprResult AssociatedFunction =
+      ParseOpenMPParensExpr(getOpenMPDirectiveName(OMPD_declare_variant), RLoc,
+                            /*IsAddressOfOperand=*/true);
+  if (!AssociatedFunction.isUsable()) {
+    if (!Tok.is(tok::annot_pragma_openmp_end))
+      while (!SkipUntil(tok::annot_pragma_openmp_end, StopBeforeMatch))
+        ;
+    // Skip the last annot_pragma_openmp_end.
+    (void)ConsumeAnnotationToken();
+    return;
+  }
+  Optional<std::pair<FunctionDecl *, Expr *>> DeclVarData =
+      Actions.checkOpenMPDeclareVariantFunction(
+          Ptr, AssociatedFunction.get(), SourceRange(Loc, Tok.getLocation()));
+
+  SmallVector<Sema::OMPCtxSelectorData, 4> Data;
+  ParseOMPDeclareVariantMatchClause(Loc, Data);
 
   // Skip last tokens.
   while (Tok.isNot(tok::annot_pragma_openmp_end))
     ConsumeAnyToken();
   if (DeclVarData.hasValue())
     Actions.ActOnOpenMPDeclareVariantDirective(
         DeclVarData.getValue().first, DeclVarData.getValue().second,
         SourceRange(Loc, Tok.getLocation()), Data);
case OMPD_declare_mapper: {
     ConsumeToken();
     if (DeclGroupPtrTy Res = ParseOpenMPDeclareMapperDirective(AS)) {
       // Skip the last annot_pragma_openmp_end.
       ConsumeAnnotationToken();
       return Res;
     }
     break;
   }
+  case OMPD_begin_declare_variant: {
+    // The syntax is:
+    // { #pragma omp begin declare variant clause }
+    // <function-declaration-or-definition-sequence>
+    // { #pragma omp end declare variant }
+    //
+    ConsumeToken();
+
+    SmallVector<Sema::OMPCtxSelectorData, 4> Data;
+    ParseOMPDeclareVariantMatchClause(Loc, Data);
+
+    // Skip last tokens.
+    while (Tok.isNot(tok::annot_pragma_openmp_end))
+      ConsumeAnyToken();
+
+    bool Elide = Actions.ActOnOpenMPDeclareVariantDirective(
+        nullptr, nullptr, SourceRange(Loc, Tok.getLocation()), Data);
+    if (!Elide)
+      break;
+
+    // Elide all the code till the matching end declare variant was found.
+    unsigned Nesting = 1;
+    do {
+      ConsumeAnyToken();
+      OpenMPDirectiveKind DK = parseOpenMPDirectiveKind(*this);
+      if (DK == OMPD_end_declare_variant)
+        --Nesting;
+      if (DK == OMPD_begin_declare_variant)
+        ++Nesting;
+    } while (Nesting);
+
+    LLVM_FALLTHROUGH;
+  }
+  case OMPD_end_declare_variant:
+    assert(getActions().DeclareVariantScopeAttr &&
+           "TODO error for unmatched end declare variant");
+    // TODO: verify DeclareVariantScopeAttr is null after parsing
+    // TODO: Make this a call in the SEMA
+    getActions().DeclareVariantScopeAttr = nullptr;
+    break;
   case OMPD_declare_variant:
   case OMPD_declare_simd: {
     // The syntax is:
     // { #pragma omp declare {simd|variant} }
     // <function-declaration-or-definition>
     //
     CachedTokens Toks;
     Toks.push_back(Tok);
case OMPD_target_teams_distribute_simd: {
     OMPDirectiveScope.Exit();
     break;
   }
   case OMPD_declare_simd:
   case OMPD_declare_target:
   case OMPD_end_declare_target:
   case OMPD_requires:
   case OMPD_declare_variant:
+  case OMPD_begin_declare_variant:
+  case OMPD_end_declare_variant:
     Diag(Tok, diag::err_omp_unexpected_directive)
         << 1 << getOpenMPDirectiveName(DKind);
     SkipUntil(tok::annot_pragma_openmp_end);
     break;
   case OMPD_unknown:
     Diag(Tok, diag::err_omp_unknown_directive);
     SkipUntil(tok::annot_pragma_openmp_end);
     break;

clang/lib/Sema/SemaDecl.cpp

if (getLangOpts().CPlusPlus) {
     //     typedef int I;
     //     typedef int I;
     //   };
     //
     // since that was the intent of DR56.
     if (!isa<TypedefNameDecl>(Old))
       return;
 
-    Diag(New->getLocation(), diag::err_redefinition)
-      << New->getDeclName();
+    Diag(New->getLocation(), diag::err_redefinition) << New->getDeclName();
     notePreviousDefinition(Old, New->getLocation());
     return New->setInvalidDecl();
   }
 
   // Modules always permit redefinition of typedefs, as does C11.
   if (getLangOpts().Modules || getLangOpts().C11)
     return;
 
Sema::ActOnFunctionDeclarator(Scope *S, Declarator &D, DeclContext *DC,
 
   DeclContext *OriginalDC = DC;
   bool IsLocalExternDecl = adjustContextForLocalExternDecl(DC);
 
   FunctionDecl *NewFD = CreateNewFunctionDecl(*this, D, DC, R, TInfo, SC,
                                               isVirtualOkay);
   if (!NewFD) return nullptr;
 
+  if (getLangOpts().OpenMP && DeclareVariantScopeAttr) {
+    OMPDeclareVariantAttr *DeclVarAttr =
+        DeclareVariantScopeAttr->clone(getASTContext());
+    DeclVarAttr->setInherited(false);
+    NewFD->addAttr(DeclVarAttr);
+    NewFD->setIsMultiVersion();
+  }
+
   if (OriginalLexicalContext && OriginalLexicalContext->isObjCContainer())
     NewFD->setTopLevelDeclInObjCContainer();
 
   // Set the lexical context. If this is a function-scope declaration, or has a
   // C++ scope specifier, or is the object of a friend declaration, the lexical
   // context will be different from the semantic context.
   NewFD->setLexicalDeclContext(CurContext);
 
case attr::CPUSpecific:
       if (MVType != MultiVersionKind::CPUDispatch &&
           MVType != MultiVersionKind::CPUSpecific)
         return true;
       break;
     case attr::Target:
       if (MVType != MultiVersionKind::Target)
         return true;
       break;
+    case attr::OMPDeclareVariant:
+      if (MVType != MultiVersionKind::OMPVariant)
+        return true;
+      break;
     default:
       return true;
     }
   }
   return false;
 }
 
 bool Sema::areMultiversionVariantFunctionsCompatible(
     const FunctionDecl *OldFD, const FunctionDecl *NewFD,
     const PartialDiagnostic &NoProtoDiagID,
     const PartialDiagnosticAt &NoteCausedDiagIDAt,
     const PartialDiagnosticAt &NoSupportDiagIDAt,
     const PartialDiagnosticAt &DiffDiagIDAt, bool TemplatesSupported,
-    bool ConstexprSupported, bool CLinkageMayDiffer) {
+    bool ConstexprSupported, bool CLinkageMayDiffer, bool StorageClassMayDiffer,
+    bool ConstexprSpecMayDiffer, bool InlineSpecificationMayDiffer) {
   enum DoesntSupport {
     FuncTemplates = 0,
     VirtFuncs = 1,
     DeducedReturn = 2,
     Constructors = 3,
     Destructors = 4,
     DeletedFuncs = 5,
     DefaultedFuncs = 6,
return Diag(NoSupportDiagIDAt.first, NoSupportDiagIDAt.second)
            << DefaultedFuncs;
 
   if (!ConstexprSupported && NewFD->isConstexpr())
     return Diag(NoSupportDiagIDAt.first, NoSupportDiagIDAt.second)
            << (NewFD->isConsteval() ? ConstevalFuncs : ConstexprFuncs);
 
   QualType NewQType = Context.getCanonicalType(NewFD->getType());
   const auto *NewType = cast<FunctionType>(NewQType);
-  QualType NewReturnType = NewType->getReturnType();
+  QualType NewReturnType = NewType->getReturnType().getUnqualifiedType();
 
   if (NewReturnType->isUndeducedType())
     return Diag(NoSupportDiagIDAt.first, NoSupportDiagIDAt.second)
            << DeducedReturn;
 
   // Ensure the return type is identical.
   if (OldFD) {
     QualType OldQType = Context.getCanonicalType(OldFD->getType());
     const auto *OldType = cast<FunctionType>(OldQType);
     FunctionType::ExtInfo OldTypeInfo = OldType->getExtInfo();
     FunctionType::ExtInfo NewTypeInfo = NewType->getExtInfo();
 
     if (OldTypeInfo.getCC() != NewTypeInfo.getCC())
       return Diag(DiffDiagIDAt.first, DiffDiagIDAt.second) << CallingConv;
 
-    QualType OldReturnType = OldType->getReturnType();
+    QualType OldReturnType = OldType->getReturnType().getUnqualifiedType();
 
     if (OldReturnType != NewReturnType)
       return Diag(DiffDiagIDAt.first, DiffDiagIDAt.second) << ReturnType;
 
-    if (OldFD->getConstexprKind() != NewFD->getConstexprKind())
+    if (!ConstexprSpecMayDiffer &&
+        OldFD->getConstexprKind() != NewFD->getConstexprKind())
       return Diag(DiffDiagIDAt.first, DiffDiagIDAt.second) << ConstexprSpec;
 
-    if (OldFD->isInlineSpecified() != NewFD->isInlineSpecified())
+    if (!InlineSpecificationMayDiffer &&
+        OldFD->isInlineSpecified() != NewFD->isInlineSpecified())
       return Diag(DiffDiagIDAt.first, DiffDiagIDAt.second) << InlineSpec;
 
-    if (OldFD->getStorageClass() != NewFD->getStorageClass())
+    if (!StorageClassMayDiffer &&
+        OldFD->getStorageClass() != NewFD->getStorageClass())
       return Diag(DiffDiagIDAt.first, DiffDiagIDAt.second) << StorageClass;
 
     if (!CLinkageMayDiffer && OldFD->isExternC() != NewFD->isExternC())
       return Diag(DiffDiagIDAt.first, DiffDiagIDAt.second) << Linkage;
 
     if (CheckEquivalentExceptionSpec(
             OldFD->getType()->getAs<FunctionProtoType>(), OldFD->getLocation(),
             NewFD->getType()->getAs<FunctionProtoType>(), NewFD->getLocation()))
       return true;
   }
   return false;
 }
 
 static bool CheckMultiVersionAdditionalRules(Sema &S, const FunctionDecl *OldFD,
                                              const FunctionDecl *NewFD,
                                              bool CausesMV,
                                              MultiVersionKind MVType) {
-  if (!S.getASTContext().getTargetInfo().supportsMultiVersioning()) {
+  bool IsOpenMPVariant = MVType == MultiVersionKind::OMPVariant;
+  if (!IsOpenMPVariant &&
+      !S.getASTContext().getTargetInfo().supportsMultiVersioning()) {
     S.Diag(NewFD->getLocation(), diag::err_multiversion_not_supported);
     if (OldFD)
       S.Diag(OldFD->getLocation(), diag::note_previous_declaration);
     return true;
   }
 
   bool IsCPUSpecificCPUDispatchMVType =
       MVType == MultiVersionKind::CPUDispatch ||
       MVType == MultiVersionKind::CPUSpecific;
 
   // For now, disallow all other attributes.  These should be opt-in, but
   // an analysis of all of them is a future FIXME.
-  if (CausesMV && OldFD && HasNonMultiVersionAttributes(OldFD, MVType)) {
+  if (CausesMV && OldFD && !IsOpenMPVariant &&
+      HasNonMultiVersionAttributes(OldFD, MVType)) {
     S.Diag(OldFD->getLocation(), diag::err_multiversion_no_other_attrs)
         << IsCPUSpecificCPUDispatchMVType;
     S.Diag(NewFD->getLocation(), diag::note_multiversioning_caused_here);
     return true;
   }
 
-  if (HasNonMultiVersionAttributes(NewFD, MVType))
+  if (!IsOpenMPVariant && HasNonMultiVersionAttributes(NewFD, MVType))
     return S.Diag(NewFD->getLocation(), diag::err_multiversion_no_other_attrs)
            << IsCPUSpecificCPUDispatchMVType;
 
   // Only allow transition to MultiVersion if it hasn't been used.
-  if (OldFD && CausesMV && OldFD->isUsed(false))
+  if (OldFD && CausesMV && !IsOpenMPVariant && OldFD->isUsed(false))
     return S.Diag(NewFD->getLocation(), diag::err_multiversion_after_used);
 
   return S.areMultiversionVariantFunctionsCompatible(
       OldFD, NewFD, S.PDiag(diag::err_multiversion_noproto),
       PartialDiagnosticAt(NewFD->getLocation(),
                           S.PDiag(diag::note_multiversioning_caused_here)),
       PartialDiagnosticAt(NewFD->getLocation(),
                           S.PDiag(diag::err_multiversion_doesnt_support)
                               << IsCPUSpecificCPUDispatchMVType),
       PartialDiagnosticAt(NewFD->getLocation(),
                           S.PDiag(diag::err_multiversion_diff)),
-      /*TemplatesSupported=*/false,
+      /*TemplatesSupported=*/IsOpenMPVariant,
       /*ConstexprSupported=*/!IsCPUSpecificCPUDispatchMVType,
-      /*CLinkageMayDiffer=*/false);
+      /*CLinkageMayDiffer=*/IsOpenMPVariant,
+      /*StorageClassMayDiffer=*/IsOpenMPVariant,
+      /*ConstexprSpecMayDiffer=*/IsOpenMPVariant,
+      /*InlineSpecificationMayDiffer=*/IsOpenMPVariant);
 }
 
 /// Check the validity of a multiversion function declaration that is the
 /// first of its kind. Also sets the multiversion'ness' of the function itself.
 ///
 /// This sets NewFD->isInvalidDecl() to true if there was an error.
 ///
 /// Returns true if there was an error, false otherwise.
 static bool CheckMultiVersionFirstFunction(Sema &S, FunctionDecl *FD,
                                            MultiVersionKind MVType,
-                                           const TargetAttr *TA) {
+                                           const TargetAttr *TA,
+                                           NamedDecl *OldDecl) {
   assert(MVType != MultiVersionKind::None &&
          "Function lacks multiversion attribute");
 
   // Target only causes MV if it is default, otherwise this is a normal
   // function.
   if (MVType == MultiVersionKind::Target && !TA->isDefaultVersion())
     return false;
 

 }
 
 /// Check the validity of a new function declaration being added to an existing
 /// multiversioned declaration collection.
 static bool CheckMultiVersionAdditionalDecl(
     Sema &S, FunctionDecl *OldFD, FunctionDecl *NewFD,
     MultiVersionKind NewMVType, const TargetAttr *NewTA,
     const CPUDispatchAttr *NewCPUDisp, const CPUSpecificAttr *NewCPUSpec,
-    bool &Redeclaration, NamedDecl *&OldDecl, bool &MergeTypeWithPrevious,
-    LookupResult &Previous) {
+    const OMPDeclareVariantAttr *NewOpenMPVariant, bool &Redeclaration,
+    NamedDecl *&OldDecl, bool &MergeTypeWithPrevious, LookupResult &Previous) {
 
   MultiVersionKind OldMVType = OldFD->getMultiVersionKind();
   // Disallow mixing of multiversioning types.
   if ((OldMVType == MultiVersionKind::Target &&
        NewMVType != MultiVersionKind::Target) ||
       (NewMVType == MultiVersionKind::Target &&
        OldMVType != MultiVersionKind::Target)) {
     S.Diag(NewFD->getLocation(), diag::err_multiversion_types_mixed);
     S.Diag(OldFD->getLocation(), diag::note_previous_declaration);
     NewFD->setInvalidDecl();
     return true;
   }
 
+  if (OldMVType == MultiVersionKind::OMPVariant &&
+      NewMVType == MultiVersionKind::None) {
+    assert(!NewOpenMPVariant && "Didn't expect variant attr!");
+    auto *OldOMPVariant = OldFD->getAttr<OMPDeclareVariantAttr>();
+    auto *NewOMPVariant = OldOMPVariant->clone(S.getASTContext());
+    NewOMPVariant->setInherited(true);
+    NewFD->addAttr(NewOMPVariant);
+    NewFD->setIsMultiVersion();
+    NewOpenMPVariant = NewOMPVariant;
+    NewMVType = MultiVersionKind::OMPVariant;
+  }
+
   TargetAttr::ParsedTargetAttr NewParsed;
   if (NewTA) {
     NewParsed = NewTA->parse();
     llvm::sort(NewParsed.Features);
   }
 
   bool UseMemberUsingDeclRules =
       S.CurContext->isRecord() && !NewFD->getFriendObjectKind();
if (NewMVType == MultiVersionKind::Target) {
       TargetAttr::ParsedTargetAttr CurParsed =
           CurTA->parse(std::less<std::string>());
       if (CurParsed == NewParsed) {
         S.Diag(NewFD->getLocation(), diag::err_multiversion_duplicate);
         S.Diag(CurFD->getLocation(), diag::note_previous_declaration);
         NewFD->setInvalidDecl();
         return true;
       }
+    } else if (NewMVType == MultiVersionKind::OMPVariant) {
+      auto *CurOMPVariant = CurFD->getAttr<OMPDeclareVariantAttr>();
+      if (!CurOMPVariant) {
+        CurOMPVariant = NewOpenMPVariant->clone(S.getASTContext());
+        CurOMPVariant->setInherited(true);
+        CurFD->addAttr(CurOMPVariant);
+        CurFD->setIsMultiVersion();
+      }
     } else {
       const auto *CurCPUSpec = CurFD->getAttr<CPUSpecificAttr>();
       const auto *CurCPUDisp = CurFD->getAttr<CPUDispatchAttr>();
       // Handle CPUDispatch/CPUSpecific versions.
       // Only 1 CPUDispatch function is allowed, this will make it go through
       // the redeclaration errors.
       if (NewMVType == MultiVersionKind::CPUDispatch &&
           CurFD->hasAttr<CPUDispatchAttr>()) {
static bool CheckMultiVersionAdditionalDecl(
   NewFD->setIsMultiVersion();
   Redeclaration = false;
   MergeTypeWithPrevious = false;
   OldDecl = nullptr;
   Previous.clear();
   return false;
 }
 
-
 /// Check the validity of a mulitversion function declaration.
 /// Also sets the multiversion'ness' of the function itself.
 ///
 /// This sets NewFD->isInvalidDecl() to true if there was an error.
 ///
 /// Returns true if there was an error, false otherwise.
 static bool CheckMultiVersionFunction(Sema &S, FunctionDecl *NewFD,
                                       bool &Redeclaration, NamedDecl *&OldDecl,
                                       bool &MergeTypeWithPrevious,
                                       LookupResult &Previous) {
   const auto *NewTA = NewFD->getAttr<TargetAttr>();
   const auto *NewCPUDisp = NewFD->getAttr<CPUDispatchAttr>();
   const auto *NewCPUSpec = NewFD->getAttr<CPUSpecificAttr>();
+  const auto *NewOpenMPVariant = NewFD->getAttr<OMPDeclareVariantAttr>();
+  unsigned NumMV = bool(NewTA) + bool(NewCPUDisp) + bool(NewCPUSpec) +
+                   bool(NewOpenMPVariant);
 
   // Mixing Multiversioning types is prohibited.
-  if ((NewTA && NewCPUDisp) || (NewTA && NewCPUSpec) ||
-      (NewCPUDisp && NewCPUSpec)) {
+  if (NumMV > 1) {
     S.Diag(NewFD->getLocation(), diag::err_multiversion_types_mixed);
     NewFD->setInvalidDecl();
     return true;
   }
 
   MultiVersionKind  MVType = NewFD->getMultiVersionKind();
 
   // Main isn't allowed to become a multiversion function, however it IS
   // permitted to have 'main' be marked with the 'target' optimization hint.
   if (NewFD->isMain()) {
     if ((MVType == MultiVersionKind::Target && NewTA->isDefaultVersion()) ||
         MVType == MultiVersionKind::CPUDispatch ||
         MVType == MultiVersionKind::CPUSpecific) {
       S.Diag(NewFD->getLocation(), diag::err_multiversion_not_allowed_on_main);
       NewFD->setInvalidDecl();
       return true;
     }
     return false;
   }
 
+  if (auto *USD = dyn_cast_or_null<UsingShadowDecl>(OldDecl))
+    OldDecl = USD->getTargetDecl();
+
   if (!OldDecl || !OldDecl->getAsFunction() ||
-      OldDecl->getDeclContext()->getRedeclContext() !=
-          NewFD->getDeclContext()->getRedeclContext()) {
+      (OldDecl->getDeclContext()->getRedeclContext() !=
+           NewFD->getDeclContext()->getRedeclContext() &&
+       !OldDecl->getAsFunction()->isOpenMPMultiVersion())) {
     // If there's no previous declaration, AND this isn't attempting to cause
     // multiversioning, this isn't an error condition.
     if (MVType == MultiVersionKind::None)
       return false;
-    return CheckMultiVersionFirstFunction(S, NewFD, MVType, NewTA);
+    return CheckMultiVersionFirstFunction(S, NewFD, MVType, NewTA, OldDecl);
   }
 
   FunctionDecl *OldFD = OldDecl->getAsFunction();
 
   if (!OldFD->isMultiVersion() && MVType == MultiVersionKind::None)
     return false;
 
-  if (OldFD->isMultiVersion() && MVType == MultiVersionKind::None) {
+  if (OldFD->isMultiVersion() && MVType == MultiVersionKind::None &&
+      !OldFD->isOpenMPMultiVersion()) {
     S.Diag(NewFD->getLocation(), diag::err_multiversion_required_in_redecl)
         << (OldFD->getMultiVersionKind() != MultiVersionKind::Target);
     NewFD->setInvalidDecl();
     return true;
   }
 
   // Handle the target potentially causes multiversioning case.
   if (!OldFD->isMultiVersion() && MVType == MultiVersionKind::Target)
     return CheckTargetCausesMultiVersioning(S, OldFD, NewFD, NewTA,
                                             Redeclaration, OldDecl,
                                             MergeTypeWithPrevious, Previous);
 
   // At this point, we have a multiversion function decl (in OldFD) AND an
   // appropriate attribute in the current function decl.  Resolve that these are
   // still compatible with previous declarations.
   return CheckMultiVersionAdditionalDecl(
-      S, OldFD, NewFD, MVType, NewTA, NewCPUDisp, NewCPUSpec, Redeclaration,
-      OldDecl, MergeTypeWithPrevious, Previous);
+      S, OldFD, NewFD, MVType, NewTA, NewCPUDisp, NewCPUSpec, NewOpenMPVariant,
+      Redeclaration, OldDecl, MergeTypeWithPrevious, Previous);
 }
 
 /// Perform semantic checking of a new function declaration.
 ///
 /// Performs semantic analysis of the new function declaration
 /// NewFD. This routine performs all semantic checking that does not
 /// require the actual declarator involved in the declaration, and is
 /// used both for the declaration of functions as they are parsed
bool Sema::CheckFunctionDeclaration(Scope *S, FunctionDecl *NewFD,
                                     LookupResult &Previous,
                                     bool IsMemberSpecialization) {
   assert(!NewFD->getReturnType()->isVariablyModifiedType() &&
          "Variably modified return types are not handled here");
 
   // Determine whether the type of this function should be merged with
   // a previous visible declaration. This never happens for functions in C++,
   // and always happens in C if the previous declaration was visible.
-  bool MergeTypeWithPrevious = !getLangOpts().CPlusPlus &&
-                               !Previous.isShadowed();
+  bool MergeTypeWithPrevious =
+      !getLangOpts().CPlusPlus && !Previous.isShadowed();
 
   bool Redeclaration = false;
   NamedDecl *OldDecl = nullptr;
   bool MayNeedOverloadableChecks = false;
 
   // Merge or overload the declaration with an existing declaration of
   // the same name, if appropriate.
   if (!Previous.empty()) {
Decl *Sema::ActOnStartOfFunctionDef(Scope *FnBodyScope, Decl *D,
 
   FunctionDecl *FD = nullptr;
 
   if (FunctionTemplateDecl *FunTmpl = dyn_cast<FunctionTemplateDecl>(D))
     FD = FunTmpl->getTemplatedDecl();
   else
     FD = cast<FunctionDecl>(D);
 
+  if (getLangOpts().OpenMP && DeclareVariantScopeAttr) {
+    OMPDeclareVariantAttr *DeclVarAttr = FD->getAttr<OMPDeclareVariantAttr>();
+    if (!DeclVarAttr) {
+      DeclVarAttr = DeclareVariantScopeAttr->clone(getASTContext());
+      FD->addAttr(DeclVarAttr);
+    }
+    DeclVarAttr->setInherited(false);
+    FD->setIsMultiVersion();
+  }
+
   // Do not push if it is a lambda because one is already pushed when building
   // the lambda in ActOnStartOfLambdaDefinition().
   if (!isLambdaCallOperator(FD))
     PushExpressionEvaluationContext(ExprEvalContexts.back().Context);
 
   // Check for defining attributes before the check for redefinition.
   if (const auto *Attr = FD->getAttr<AliasAttr>()) {
     Diag(Attr->getLocation(), diag::err_alias_is_definition) << FD << 0;

clang/lib/Sema/SemaOpenMP.cpp

Sema::checkOpenMPDeclareVariantFunction(Sema::DeclGroupPtrTy DG,
   const FunctionDecl *Definition;
   if (!FD->isThisDeclarationADefinition() && FD->isDefined(Definition) &&
       (LangOpts.EmitAllDecls || Context.DeclMustBeEmitted(Definition)))
     Diag(SR.getBegin(), diag::warn_omp_declare_variant_after_emitted)
         << FD->getLocation();
 
   // The VariantRef must point to function.
   if (!VariantRef) {
-    Diag(SR.getBegin(), diag::err_omp_function_expected) << VariantId;
-    return None;
+    // Diag(SR.getBegin(), diag::err_omp_function_expected) << VariantId;
+    return std::make_pair(FD, VariantRef);
   }
 
   // Do not check templates, wait until instantiation.
   if (VariantRef->isTypeDependent() || VariantRef->isValueDependent() ||
       VariantRef->containsUnexpandedParameterPack() ||
       VariantRef->isInstantiationDependent() || FD->isDependentContext())
     return std::make_pair(FD, VariantRef);
 
Diag(VariantRef->getExprLoc(),
          diag::warn_omp_declare_variant_marked_as_declare_variant)
         << VariantRef->getSourceRange();
     SourceRange SR =
         NewFD->specific_attr_begin<OMPDeclareVariantAttr>()->getRange();
     Diag(SR.getBegin(), diag::note_omp_marked_declare_variant_here) << SR;
     return None;
   }
 
+  // TODO check these for missing VariantRef as well
   enum DoesntSupport {
     VirtFuncs = 1,
     Constructors = 3,
     Destructors = 4,
     DeletedFuncs = 5,
     DefaultedFuncs = 6,
     ConstexprFuncs = 7,
     ConstevalFuncs = 8,
if (areMultiversionVariantFunctionsCompatible(
               PDiag(diag::note_omp_declare_variant_specified_here) << SR),
           PartialDiagnosticAt(
               VariantRef->getExprLoc(),
               PDiag(diag::err_omp_declare_variant_doesnt_support)),
           PartialDiagnosticAt(VariantRef->getExprLoc(),
                               PDiag(diag::err_omp_declare_variant_diff)
                                   << FD->getLocation()),
           /*TemplatesSupported=*/true, /*ConstexprSupported=*/false,
-          /*CLinkageMayDiffer=*/true))
+          /*CLinkageMayDiffer=*/true,
+          /*StorageClassMayDiffer=*/true,
+          /*ConstexprSpecMayDiffer=*/true,
+          /*InlineSpecificationMayDiffer=*/true))
     return None;
   return std::make_pair(FD, cast<Expr>(DRE));
 }
 
-void Sema::ActOnOpenMPDeclareVariantDirective(
+bool Sema::ActOnOpenMPDeclareVariantDirective(
     FunctionDecl *FD, Expr *VariantRef, SourceRange SR,
     ArrayRef<OMPCtxSelectorData> Data) {
   if (Data.empty())
-    return;
+    return false;
   SmallVector<Expr *, 4> CtxScores;
   SmallVector<unsigned, 4> CtxSets;
   SmallVector<unsigned, 4> Ctxs;
   SmallVector<StringRef, 4> ImplVendors, DeviceKinds;
   bool IsError = false;
   for (const OMPCtxSelectorData &D : Data) {
     OpenMPContextSelectorSetKind CtxSet = D.CtxSet;
     OpenMPContextSelectorKind Ctx = D.Ctx;
     if (CtxSet == OMP_CTX_SET_unknown || Ctx == OMP_CTX_unknown)
-      return;
+      return false;
     Expr *Score = nullptr;
     if (D.Score.isUsable()) {
       Score = D.Score.get();
       if (!Score->isTypeDependent() && !Score->isValueDependent() &&
           !Score->isInstantiationDependent() &&
           !Score->containsUnexpandedParameterPack()) {
         Score =
             PerformOpenMPImplicitIntegerConversion(Score->getExprLoc(), Score)
for (const OMPCtxSelectorData &D : Data) {
     CtxScores.push_back(Score);
   }
   if (!IsError) {
     auto *NewAttr = OMPDeclareVariantAttr::CreateImplicit(
         Context, VariantRef, CtxScores.begin(), CtxScores.size(),
         CtxSets.begin(), CtxSets.size(), Ctxs.begin(), Ctxs.size(),
         ImplVendors.begin(), ImplVendors.size(), DeviceKinds.begin(),
         DeviceKinds.size(), SR);
+    if (FD) {
       FD->addAttr(NewAttr);
+    } else {
+      assert(!DeclareVariantScopeAttr &&
+             "TODO nested begin/end declare varinat");
+      DeclareVariantScopeAttr = NewAttr;
+      return !isOpenMPContextMatch(getASTContext(), DeclareVariantScopeAttr);
     }
   }
+  return false;
+}
 
 void Sema::markOpenMPDeclareVariantFuncsReferenced(SourceLocation Loc,
                                                    FunctionDecl *Func,
                                                    bool MightBeOdrUse) {
   assert(LangOpts.OpenMP && "Expected OpenMP mode.");
 
   if (!Func->isDependentContext() && Func->hasAttrs()) {
     for (OMPDeclareVariantAttr *A :
          Func->specific_attrs<OMPDeclareVariantAttr>()) {
       // TODO: add checks for active OpenMP context where possible.
       Expr *VariantRef = A->getVariantFuncRef();
+      FunctionDecl *F = Func;
+      if (VariantRef) {
         auto *DRE = cast<DeclRefExpr>(VariantRef->IgnoreParenImpCasts());
-      auto *F = cast<FunctionDecl>(DRE->getDecl());
+        F = cast<FunctionDecl>(DRE->getDecl());
+      }
       if (!F->isDefined() && F->isTemplateInstantiation())
         InstantiateFunctionDefinition(Loc, F->getFirstDecl());
+      if (F != Func)
         MarkFunctionReferenced(Loc, F, MightBeOdrUse);
     }
   }
 }
 
 StmtResult Sema::ActOnOpenMPParallelDirective(ArrayRef<OMPClause *> Clauses,
                                               Stmt *AStmt,
                                               SourceLocation StartLoc,
                                               SourceLocation EndLoc) {
OMPClause *Sema::ActOnOpenMPAllocateClause(
   }
 
   if (Vars.empty())
     return nullptr;
 
   return OMPAllocateClause::Create(Context, StartLoc, LParenLoc, Allocator,
                                    ColonLoc, EndLoc, Vars);
 }
+
+template <typename AttrTy>
+static void copyAttrIfPresent(Sema &S, FunctionDecl *FD,
+                              const FunctionDecl &TemplateFD) {
+  if (!FD->hasAttr<AttrTy>())
+    if (AttrTy *Attribute = TemplateFD.getAttr<AttrTy>()) {
+      AttrTy *Clone = Attribute->clone(S.Context);
+      Clone->setInherited(true);
+      FD->addAttr(Clone);
+    }
+}
+
+void Sema::inheritOpenMPVariantAttrs(FunctionDecl *FD,
+                                     const FunctionTemplateDecl &TD) {
+  const FunctionDecl &TemplateFD = *TD.getTemplatedDecl();
+  copyAttrIfPresent<OMPDeclareVariantAttr>(*this, FD, TemplateFD);
+}

clang/lib/Sema/SemaOverload.cpp

 
 #include "clang/Sema/Overload.h"
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/CXXInheritance.h"
 #include "clang/AST/DeclObjC.h"
 #include "clang/AST/Expr.h"
 #include "clang/AST/ExprCXX.h"
 #include "clang/AST/ExprObjC.h"
+#include "clang/AST/StmtOpenMP.h"
 #include "clang/AST/TypeOrdering.h"
 #include "clang/Basic/Diagnostic.h"
 #include "clang/Basic/DiagnosticOptions.h"
 #include "clang/Basic/PartialDiagnostic.h"
 #include "clang/Basic/TargetInfo.h"
 #include "clang/Sema/Initialization.h"
 #include "clang/Sema/Lookup.h"
 #include "clang/Sema/SemaInternal.h"
for (auto Pair : zip_longest(Cand1Attrs, Cand2Attrs)) {
     (*Cand2A)->getCond()->Profile(Cand2ID, S.getASTContext(), true);
     if (Cand1ID != Cand2ID)
       return Comparison::Worse;
   }
 
   return Comparison::Equal;
 }
 
-static bool isBetterMultiversionCandidate(const OverloadCandidate &Cand1,
+static bool isBetterMultiversionCandidate(ASTContext &ASTCtx,
+                                          const OverloadCandidate &Cand1,
                                           const OverloadCandidate &Cand2) {
   if (!Cand1.Function || !Cand1.Function->isMultiVersion() || !Cand2.Function ||
       !Cand2.Function->isMultiVersion())
     return false;
 
   // If Cand1 is invalid, it cannot be a better match, if Cand2 is invalid, this
   // is obviously better.
   if (Cand1.Function->isInvalidDecl()) return false;
   if (Cand2.Function->isInvalidDecl()) return true;
 
+  // If we have an OpenMP declare variant attribute on either candidate we use
+  // it to order the candidates. The first is only better if it has a attribute
+  // that is considered better or if it has no attribute and the one on the
+  // second candidate is not a match.
+  auto *OMPVariantAttr1 = Cand1.Function->getAttr<OMPDeclareVariantAttr>();
+  auto *OMPVariantAttr2 = Cand2.Function->getAttr<OMPDeclareVariantAttr>();
+  if (OMPVariantAttr1 || OMPVariantAttr2) {
+    auto *OMPVariantAttrBest =
+        getBetterOpenMPContextMatch(ASTCtx, OMPVariantAttr1, OMPVariantAttr2);
+    return OMPVariantAttrBest == OMPVariantAttr1;
+  }
+
   // If this is a cpu_dispatch/cpu_specific multiversion situation, prefer
   // cpu_dispatch, else arbitrarily based on the identifiers.
   bool Cand1CPUDisp = Cand1.Function->hasAttr<CPUDispatchAttr>();
   bool Cand2CPUDisp = Cand2.Function->hasAttr<CPUDispatchAttr>();
   const auto *Cand1CPUSpec = Cand1.Function->getAttr<CPUSpecificAttr>();
   const auto *Cand2CPUSpec = Cand2.Function->getAttr<CPUSpecificAttr>();
 
   if (!Cand1CPUDisp && !Cand2CPUDisp && !Cand1CPUSpec && !Cand2CPUSpec)
bool clang::isBetterOverloadCandidate(
 
   bool HasPS1 = Cand1.Function != nullptr &&
                 functionHasPassObjectSizeParams(Cand1.Function);
   bool HasPS2 = Cand2.Function != nullptr &&
                 functionHasPassObjectSizeParams(Cand2.Function);
   if (HasPS1 != HasPS2 && HasPS1)
     return true;
 
-  return isBetterMultiversionCandidate(Cand1, Cand2);
+  return isBetterMultiversionCandidate(S.getASTContext(), Cand1, Cand2);
 }
 
 /// Determine whether two declarations are "equivalent" for the purposes of
 /// name lookup and overload resolution. This applies when the same internal/no
 /// linkage entity is defined by two modules (probably by textually including
 /// the same header). In such a case, we don't consider the declarations to
 /// declare the same entity, but we also don't want lookups with both
 /// declarations visible to be ambiguous in some cases (this happens when using
if (ContainsSameSideCandidate) {
         return Cand->Viable && Cand->Function &&
                S.IdentifyCUDAPreference(Caller, Cand->Function) ==
                    Sema::CFP_WrongSide;
       };
       llvm::erase_if(Candidates, IsWrongSideCandidate);
     }
   }
 
+  // [OpenMP] Similar to the CUDA code above, OpenMP declare variants might not
+  // be eligible at all so we need to filter them out early.
+  if (S.getLangOpts().OpenMP) {
+    // TODO use context information
+    auto IsNonMatchVariant = [&](OverloadCandidate *Cand) {
+      if (!Cand->Viable || !Cand->Function)
+        return false;
+      auto *OMPVariantAttr = Cand->Function->getAttr<OMPDeclareVariantAttr>();
+      return OMPVariantAttr &&
+             !isOpenMPContextMatch(S.getASTContext(), OMPVariantAttr);
+    };
+    llvm::erase_if(Candidates, IsNonMatchVariant);
+  }
+
   // Find the best viable function.
   Best = end();
   for (auto *Cand : Candidates) {
     Cand->Best = false;
     if (Cand->Viable)
       if (Best == end() ||
           isBetterOverloadCandidate(S, *Cand, *Best, Loc, Kind))
         Best = Cand;

clang/lib/Sema/SemaTemplate.cpp

bool Sema::CheckFunctionTemplateSpecialization(
   // A function template specialization inherits the target attributes
   // of its template.  (We require the attributes explicitly in the
   // code to match, but a template may have implicit attributes by
   // virtue e.g. of being constexpr, and it passes these implicit
   // attributes on to its specializations.)
   if (LangOpts.CUDA)
     inheritCUDATargetAttrs(FD, *Specialization->getPrimaryTemplate());
 
+  if (LangOpts.OpenMP)
+    inheritOpenMPVariantAttrs(FD, *Specialization->getPrimaryTemplate());
+
   // The "previous declaration" for this function template specialization is
   // the prior function template specialization.
   Previous.clear();
   Previous.addDecl(Specialization);
   return false;
 }
 
 /// Perform semantic analysis for the given non-template member

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

case OMP_CTX_SET_device:
       case OMP_CTX_unknown:
         llvm_unreachable("Unexpected context selector kind.");
       }
       break;
     case OMP_CTX_SET_unknown:
       llvm_unreachable("Unexpected context selector set kind.");
     }
   }
-  S.ActOnOpenMPDeclareVariantDirective(DeclVarData.getValue().first,
-                                       DeclVarData.getValue().second,
-                                       Attr.getRange(), Data);
+  FunctionDecl *VariantFD =
+      DeclVarData.hasValue() ? DeclVarData.getValue().first : FD;
+  Expr *VariantExpr =
+      DeclVarData.hasValue() ? DeclVarData.getValue().second : nullptr;
+  S.ActOnOpenMPDeclareVariantDirective(VariantFD, VariantExpr, Attr.getRange(),
+                                       Data);
+
+  // The new attribute on the instantiation is inherited if the template
+  // attribute was.
+  if (Attr.isInherited())
+    VariantFD->getAttr<OMPDeclareVariantAttr>()->setInherited(true);
 }
 
 static void instantiateDependentAMDGPUFlatWorkGroupSizeAttr(
     Sema &S, const MultiLevelTemplateArgumentList &TemplateArgs,
     const AMDGPUFlatWorkGroupSizeAttr &Attr, Decl *New) {
   // Both min and max expression are constant expressions.
   EnterExpressionEvaluationContext Unevaluated(
       S, Sema::ExpressionEvaluationContext::ConstantEvaluated);

clang/test/AST/ast-dump-openmp-begin-declare-variant.c

+// RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -ast-dump %s | FileCheck %s
+
+int also_before(void) {
+  return 0;
+}
+
+#pragma omp begin declare variant match(device={kind(cpu)})
+int also_after(void) {
+  return 1;
+}
+int also_before(void) {
+  return 1;
+}
+#pragma omp end declare variant
+
+#pragma omp begin declare variant match(device={kind(gpu)})
+int also_after(void) {
+  return 2;
+}
+int also_before(void) {
+  return 2;
+}
+#pragma omp end declare variant
+
+
+#pragma omp begin declare variant match(device={kind(fpga)})
+
+This text is never parsed!
+
+#pragma omp end declare variant
+
+int also_after(void) {
+  return 0;
+}
+
+int test() {
+  return also_after() + also_before();
+}
+
+// Make sure:
+// 1) we pick the right versions, that is test should reference the kind(cpu) versions.
+// 2) we do not see the ast nodes for the gpu kind
+// 3) we do not choke on the text in the kind(fpga) guarded scope.
+
+// CHECK:       -FunctionDecl {{.*}} <{{.*}}3:1, line:{{.*}}:1> line:{{.*}}:5 also_before 'int (void)'
+// CHECK-NEXT:  | |-CompoundStmt {{.*}} <col:23, line:{{.*}}:1>
+// CHECK-NEXT:  | | `-ReturnStmt {{.*}} <line:{{.*}}:3, col:10>
+// CHECK-NEXT:  | |   `-IntegerLiteral {{.*}} <col:10> 'int' 0
+// CHECK-NEXT:  | `-OMPDeclareVariantAttr {{.*}} <line:{{.*}}:1, col:60> Inherited Implicit 1 1 cpu
+// CHECK-NEXT:  |   |-<<<NULL>>>
+// CHECK-NEXT:  |   `-IntegerLiteral {{.*}} <<invalid sloc>> 'int' 1
+// CHECK-NEXT:  |-FunctionDecl [[GOOD_ALSO_AFTER:0x[a-z0-9]*]] <line:{{.*}}:1, line:{{.*}}:1> line:{{.*}}:5 used also_after 'int (void)'
+// CHECK-NEXT:  | |-CompoundStmt {{.*}} <col:22, line:{{.*}}:1>
+// CHECK-NEXT:  | | `-ReturnStmt {{.*}} <line:{{.*}}:3, col:10>
+// CHECK-NEXT:  | |   `-IntegerLiteral {{.*}} <col:10> 'int' 1
+// CHECK-NEXT:  | `-OMPDeclareVariantAttr {{.*}} <line:{{.*}}:1, col:60> Implicit 1 1 cpu
+// CHECK-NEXT:  |   |-<<<NULL>>>
+// CHECK-NEXT:  |   `-IntegerLiteral {{.*}} <<invalid sloc>> 'int' 1
+// CHECK-NEXT:  |-FunctionDecl [[GOOD_ALSO_BEFORE:0x[a-z0-9]*]] <line:{{.*}}:1, line:{{.*}}:1> line:{{.*}}:5 used also_before 'int (void)'
+// CHECK-NEXT:  | |-CompoundStmt {{.*}} <col:23, line:{{.*}}:1>
+// CHECK-NEXT:  | | `-ReturnStmt {{.*}} <line:{{.*}}:3, col:10>
+// CHECK-NEXT:  | |   `-IntegerLiteral {{.*}} <col:10> 'int' 1
+// CHECK-NEXT:  | `-OMPDeclareVariantAttr {{.*}} <line:{{.*}}:1, col:60> Implicit 1 1 cpu
+// CHECK-NEXT:  |   |-<<<NULL>>>
+// CHECK-NEXT:  |   `-IntegerLiteral {{.*}} <<invalid sloc>> 'int' 1
+// CHECK-NEXT:  |-FunctionDecl {{.*}} <line:{{.*}}:1, line:{{.*}}:1> line:{{.*}}:5 also_after 'int (void)'
+// CHECK-NEXT:  | |-CompoundStmt {{.*}} <col:22, line:{{.*}}:1>
+// CHECK-NEXT:  | | `-ReturnStmt {{.*}} <line:{{.*}}:3, col:10>
+// CHECK-NEXT:  | |   `-IntegerLiteral {{.*}} <col:10> 'int' 0
+// CHECK-NEXT:  | `-OMPDeclareVariantAttr {{.*}} <line:{{.*}}:1, col:60> Inherited Implicit 1 1 cpu
+// CHECK-NEXT:  |   |-<<<NULL>>>
+// CHECK-NEXT:  |   `-IntegerLiteral {{.*}} <<invalid sloc>> 'int' 1
+// CHECK-NEXT:  `-FunctionDecl {{.*}} <line:{{.*}}:1, line:{{.*}}:1> line:{{.*}}:5 test 'int ()'
+// CHECK-NEXT:    `-CompoundStmt {{.*}} <col:12, line:{{.*}}:1>
+// CHECK-NEXT:      `-ReturnStmt {{.*}} <line:{{.*}}:3, col:37>
+// CHECK-NEXT:        `-BinaryOperator {{.*}} <col:10, col:37> 'int' '+'
+// CHECK-NEXT:          |-CallExpr {{.*}} <col:10, col:21> 'int'
+// CHECK-NEXT:          | `-ImplicitCastExpr {{.*}} <col:10> 'int (*)(void)' <FunctionToPointerDecay>
+// CHECK-NEXT:          |   `-DeclRefExpr {{.*}} <col:10> 'int (void)' lvalue Function [[GOOD_ALSO_AFTER]] 'also_after' 'int (void)'
+// CHECK-NEXT:          `-CallExpr {{.*}} <col:25, col:37> 'int'
+// CHECK-NEXT:            `-ImplicitCastExpr {{.*}} <col:25> 'int (*)(void)' <FunctionToPointerDecay>
+// CHECK-NEXT:              `-DeclRefExpr {{.*}} <col:25> 'int (void)' lvalue Function [[GOOD_ALSO_BEFORE]] 'also_before' 'int (void)'
+

clang/test/OpenMP/begin_declare_variant_codegen.cpp

+// RUN: %clang_cc1 -verify -fopenmp -x c++ -emit-llvm %s -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -o - | FileCheck %s
+// expected-no-diagnostics
+
+int bar(void) {
+  return 0;
+}
+
+template <typename T>
+T baz(void) { return 0; }
+
+#pragma omp begin declare variant match(device={kind(cpu)})
+int foo(void) {
+  return 1;
+}
+int bar(void) {
+  return 1;
+}
+template <typename T>
+T baz(void) { return 1; }
+
+template <typename T>
+T biz(void) { return 1; }
+
+template <typename T>
+T buz(void) { return 3; }
+
+template <>
+char buz(void) { return 1; }
+
+template <typename T>
+T bez(void) { return 3; }
+#pragma omp end declare variant
+
+#pragma omp begin declare variant match(device={kind(gpu)})
+int foo(void) {
+  return 2;
+}
+int bar(void) {
+  return 2;
+}
+#pragma omp end declare variant
+
+
+#pragma omp begin declare variant match(device={kind(fpga)})
+
+This text is never parsed!
+
+#pragma omp end declare variant
+
+int foo(void) {
+  return 0;
+}
+
+template <typename T>
+T biz(void) { return 0; }
+
+template <>
+char buz(void) { return 0; }
+
+template <>
+long bez(void) { return 0; }
+
+#pragma omp begin declare variant match(device = {kind(cpu)})
+template <>
+long bez(void) { return 1; }
+#pragma omp end declare variant
+
+int test() {
+  return foo() + bar() + baz<int>() + biz<short>() + buz<char>() + bez<long>();
+}
+
+// Make sure all ompvariant functions return 1 and all others return 0.
+
+// CHECK:       ; Function Attrs:
+// CHECK-NEXT:  define i32 @_Z3barv()
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret i32 0
+// CHECK-NEXT:  }
+// CHECK:       ; Function Attrs:
+// CHECK-NEXT:  define i32 @_Z3foov.ompvariant()
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret i32 1
+// CHECK-NEXT:  }
+// CHECK:       ; Function Attrs:
+// CHECK-NEXT:  define i32 @_Z3barv.ompvariant()
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret i32 1
+// CHECK-NEXT:  }
+// CHECK:       ; Function Attrs:
+// CHECK-NEXT:  define signext i8 @_Z3buzIcET_v.ompvariant()
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret i8 1
+// CHECK-NEXT:  }
+// CHECK:       ; Function Attrs:
+// CHECK-NEXT:  define i32 @_Z3foov()
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret i32 0
+// CHECK-NEXT:  }
+// CHECK:       ; Function Attrs:
+// CHECK-NEXT:  define signext i8 @_Z3buzIcET_v()
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret i8 0
+// CHECK-NEXT:  }
+// CHECK:       ; Function Attrs:
+// CHECK-NEXT:  define i64 @_Z3bezIlET_v()
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret i64 0
+// CHECK-NEXT:  }
+// CHECK:       ; Function Attrs:
+// CHECK-NEXT:  define i64 @_Z3bezIlET_v.ompvariant()
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret i64 1
+// CHECK-NEXT:  }
+
+// Make sure we call only ompvariant functions
+
+// CHECK:  define i32 @_Z4testv()
+// CHECK:    %call = call i32 @_Z3foov.ompvariant()
+// CHECK:    %call1 = call i32 @_Z3barv.ompvariant()
+// CHECK:    %call2 = call i32 @_Z3bazIiET_v.ompvariant()
+// CHECK:    %call4 = call signext i16 @_Z3bizIsET_v.ompvariant()
+// CHECK:    %call6 = call signext i8 @_Z3buzIcET_v.ompvariant()
+// CHECK:    %call10 = call i64 @_Z3bezIlET_v.ompvariant()
+
+// CHECK:       ; Function Attrs:
+// CHECK-NEXT:  define linkonce_odr i32 @_Z3bazIiET_v.ompvariant()
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret i32 1
+// CHECK-NEXT:  }
+// CHECK:       ; Function Attrs:
+// CHECK-NEXT:  define linkonce_odr signext i16 @_Z3bizIsET_v.ompvariant()
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    ret i16 1
+// CHECK-NEXT:  }

clang/test/OpenMP/math_fp_macro.cpp

+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -x c++ -emit-llvm %s -triple %itanium_abi_triple -fexceptions -fcxx-exceptions -o - | FileCheck %s
+// expected-no-diagnostics
+
+#include <cmath>
+
+int main() {
+  double a(0);
+  return (std::fpclassify(a) != FP_ZERO);
+}