[OpenMP] Permit map with DSA on combined directive

[jdenny]

Summary

For map, the following restriction changed in OpenMP 5.0:

• OpenMP 4.5 [2.15.5.1, Restrictions]: "A list item cannot appear in both a map clause and a data-sharing attribute clause on the same construct.

• OpenMP 5.0 [2.19.7.1, Restrictions]: "A list item cannot appear in both a map clause and a data-sharing attribute clause on the same construct unless the construct is a combined construct."

This patch removes this restriction in the case of combined constructs
and OpenMP 5.0, and it updates Sema not to capture a scalar by copy in
the target region when firstprivate and map appear for that scalar
on a combined target construct.

This patch also adds a fixme to a test that now reveals that a
diagnostic about loop iteration variables is dropped in the case of
OpenMP 5.0. That bug exists regardless of this patch's changes.

[ABataev]

is_device_ptr can be considered as a kind of mapping clause (see 2.19.7 Data-Mapping Attribute Rules, Clauses, and Directives), so, I assume, clang correct here in terms of OpenMP 4.5.
Thus, I would not call this a "fix", this is just a new feature from OpenMP 5.0.
Plus, these changes are not enough to support this new feature from OpenMP 5.0. There definitely must be some changes in the codegen. If the variable is mapped in the target construct, we should not generate a code for the private clause of this variable on the target construct, since, in this case, private clauses are applied for the inner subdirectives, which are the part of the combined directive, but not the target part of the directive.

clang/lib/Sema/SemaOpenMP.cpp

10895 ↗

(On Diff #213749)

I would suggest to guard this change and limit this new functionality only for OpenMP 5.0.

[jdoerfert]

clang/test/OpenMP/target_parallel_for_is_device_ptr_messages.cpp

93 ↗

(On Diff #213749)

I think this should cause an error or at least a warning. Telling the compiler ps is a device pointer only to create a local uninitialized shadowing variable seems like an error to me.

[jdenny]

In D65835#1618865, @ABataev wrote:

is_device_ptr can be considered as a kind of mapping clause (see 2.19.7 Data-Mapping Attribute Rules, Clauses, and Directives), so, I assume, clang correct here in terms of OpenMP 4.5.

Maybe, but I haven't found any statement in either version that states that map restrictions apply to is_device_ptr.

Another question is whether the restriction would make sense. For example, does it ever make sense to specify both is_device_ptr and firstprivate for the same variable on a target construct? I think that would mean that modifications that are made to that device pointer within the target region should not be seen after the target region. I think that's reasonable, but that combination is not possible with the restriction. As Johannes points out, private plus is_device_ptr probably doesn't make sense.

Thus, I would not call this a "fix", this is just a new feature from OpenMP 5.0.

Understood.

I should have reported that the current implementation isn't complete for OpenMP 4.5. For example, on target teams, reduction(+:x) map(x) is an error but not map(x) reduction(+:x). So there are bugs to fix, and maybe this will evolve into multiple patches, but I want to be sure I'm on the right path first.

Plus, these changes are not enough to support this new feature from OpenMP 5.0. There definitely must be some changes in the codegen. If the variable is mapped in the target construct, we should not generate a code for the private clause of this variable on the target construct, since, in this case, private clauses are applied for the inner subdirectives, which are the part of the combined directive, but not the target part of the directive.

I'll look into it.

Thanks for the quick review.

clang/lib/Sema/SemaOpenMP.cpp

10895 ↗

(On Diff #213749)

Do you agree that this is strictly an extension to 4.5 that won't alter the behavior of 4.5-conforming applications?

Do we generally want to complain about the use of extensions, or is there another reason for the guard you suggest?

[ABataev]

clang/test/OpenMP/target_parallel_for_is_device_ptr_messages.cpp

93 ↗

(On Diff #213749)

It is allowed according to OpenMP 5.0. Private copy must be created in the context of the parallel region, not the target region. So, for OpenMP 5.0 we should not emit error here.

[jdoerfert]

clang/test/OpenMP/target_parallel_for_is_device_ptr_messages.cpp

93 ↗

(On Diff #213749)

What does that mean and how does that affect my reasoning?

[ABataev]

Maybe, but I haven't found any statement in either version that states that map restrictions apply to is_device_ptr.

is_device_ptr is a kind of mapping clause. I assume we can extend the restrictions for map clause to this clause too.

Another question is whether the restriction would make sense. For example, does it ever make sense to specify both is_device_ptr and firstprivate for the same variable on a target construct?

On a target construct - no. On a target parallel - yes. This may be important especially in unified shared memory mode.

clang/test/OpenMP/target_parallel_for_is_device_ptr_messages.cpp

93 ↗

(On Diff #213749)

It means, that for OpenMP 5.0 we should emit a warning/error here. It is allowed according to the standard, we should allow it too.

[ABataev]

I should have reported that the current implementation isn't complete for OpenMP 4.5. For example, on target teams, reduction(+:x) map(x) is an error but not map(x) reduction(+:x). So there are bugs to fix, and maybe this will evolve into multiple patches, but I want to be sure I'm on the right path first.

It is just a bug, not a missing feature. Just file a bug report for it.

clang/lib/Sema/SemaOpenMP.cpp

10895 ↗

(On Diff #213749)

No. It is incorrect according to OpenMP 4.5 and we shall emit diagnostics here.

[jdenny]

In D65835#1619526, @ABataev wrote:

Maybe, but I haven't found any statement in either version that states that map restrictions apply to is_device_ptr.

is_device_ptr is a kind of mapping clause. I assume we can extend the restrictions for map clause to this clause too.

I'd like to understand this better. Is there something from the spec we can quote in the code?

Another question is whether the restriction would make sense. For example, does it ever make sense to specify both is_device_ptr and firstprivate for the same variable on a target construct?

On a target construct - no.

Why not?

[ABataev]

In D65835#1619549, @jdenny wrote:

In D65835#1619526, @ABataev wrote:

Maybe, but I haven't found any statement in either version that states that map restrictions apply to is_device_ptr.

is_device_ptr is a kind of mapping clause. I assume we can extend the restrictions for map clause to this clause too.

I'd like to understand this better. Is there something from the spec we can quote in the code?

See 2.19.7 Data-Mapping Attribute Rules, Clauses, and Directives

Another question is whether the restriction would make sense. For example, does it ever make sense to specify both is_device_ptr and firstprivate for the same variable on a target construct?

On a target construct - no.

Why not?

It is meaningless. That's why it is prohibited in OpenMP 5.0 and allowed only for the combined constructs. These private clauses are applied to inner subconstructs.
For example, target parallel map(p) private(p). In the context of target region the variable is mapped while in the parallel context it is private.
For target map(p) private(p) it is absolutely meaningless.

[jdoerfert]

clang/test/OpenMP/target_parallel_for_is_device_ptr_messages.cpp

93 ↗

(On Diff #213749)

So, for OpenMP 5.0 we should not emit error here.
that for OpenMP 5.0 we should emit a warning/error here.

The last answer contradicts what you said earlier. I expect there is a *not* missing, correct?

Assuming you do not want an error, which is fine, I still think a warning is appropriate as it seems to me there is never a reason to have a is_device_ptr clause for a private value. I mean, it is either a bug by the programmer, e.g., 5 letters of firstprivate went missing, or simply nonsensical code for which we warn in other situations as well.

[ABataev]

clang/test/OpenMP/target_parallel_for_is_device_ptr_messages.cpp

93 ↗

(On Diff #213749)

Missed not.
These kind of construct are explicitly allowed in OpenMP. And we should obey the standard unconditionally.
Plus, there might be situations where user may require it explicitly. For example, the device pointer is dereferenced in one of the clauses for the subregions but in the deeper subregion it might be used as a private pointer. Why we should emit a warning here?

[jdoerfert]

clang/test/OpenMP/target_parallel_for_is_device_ptr_messages.cpp

93 ↗

(On Diff #213749)

If you have a different situation, e.g., the one you describe, you should not have a warning. Though, that is not the point. If you have the situation above (single directive), as per my reasoning, there doesn't seem to be a sensible use case. If you have one and we should create an explicit test for it.

These kind of construct are explicitly allowed in OpenMP.

explicitly allowed != not forbidded (yet)

And we should obey the standard unconditionally.

Nobody says we should not. We warn people all the time even if it is valid code.

[ABataev]

clang/test/OpenMP/target_parallel_for_is_device_ptr_messages.cpp

93 ↗

(On Diff #213749)

Warnings may prevent successful compilation in some cases, e.g. when warnings are treated as errors. Why we should emit a warning if the construct is allowed by the standard? Ask to change the standard if you did not agree with it.

[hfinkel]

clang/test/OpenMP/target_parallel_for_is_device_ptr_messages.cpp

93 ↗

(On Diff #213749)

Warnings are specifically for constructs which are legal, but likely wrong (i.e., will not do what the user likely intended). Treating warnings as errors is not a conforming compilation mode - by design (specifically because that will reject conforming programs). Thus...

Why we should emit a warning if the construct is allowed by the standard? Ask to change the standard if you did not agree with it.

This is the wrong way to approach this. Warnings are specifically for legal code. They help users prevent errors, however, in cases where that legal code is likely problematic or won't do what the user intends.

[ABataev]

clang/test/OpenMP/target_parallel_for_is_device_ptr_messages.cpp

93 ↗

(On Diff #213749)

Ok, we could emit wqrnings in some cases. But better to do it in the separate patches. Each particular case requires additional analysis.

This is the wrong way to approach this.

I don't think so. If some cases are really meaningless, better to ask to prohibit them in the standard. It is always a good idea to change the requirements first, if you think that some scenarios are not described correctly rather than do the changes in the code. It leads to different behavior of different compilers in the same situation and it is not good for the users.

[hfinkel]

clang/test/OpenMP/target_parallel_for_is_device_ptr_messages.cpp

93 ↗

(On Diff #213749)

I don't think so. If some cases are really meaningless, better to ask to prohibit them in the standard. It is always a good idea to change the requirements first, if you think that some scenarios are not described correctly rather than do the changes in the code. It leads to different behavior of different compilers in the same situation and it is not good for the users.

There are at least two relevant factors:

1. Language standards often express general concepts that can be combined in some regular set of ways. Some of these combinations are likely unintentional (e.g., user error), but standards don't explicitly prohibit them because: a) standards committees have limited bandwidth and need to concentrate on the highest-priority items and new features b) filling standards with a large number of special cases, even in the name of preventing user error, itself has a cost (in terms of maintenance of the standard, constraining conforming implementation techniques, and so on).
2. Even if a standards committee were to take up restricting some set of special cases, implementation experience with a warning is often very helpful. Saying, "we added a warning, and no one complained about it being a false positive" is good evidence in support of making that warning a mandated error.

In the end, standards committees depend on implementers to add value on top of the standard itself in creating an high-QoI products. This has always been a focus area of Clang, and Clang is well known for its high diagnostic quality - not just in error messages, but in warnings too.

I have plenty of users who specifically compile with multiple compilers specifically to get the warnings for each compiler. Is it sometimes true that some compilers generating some warnings ends up being problematic? yes. I think that we all have observed that. But warnings are very helpful in catching likely bugs, and implementations have more freedom with warnings than with errors, so many users depend on high-quality warnings to help quickly find bugs and, thus, increase their productivity.

[ABataev]

clang/test/OpenMP/target_parallel_for_is_device_ptr_messages.cpp

93 ↗

(On Diff #213749)

Just like I said, if you think there are some incorrect combinations we could generate a warning. But better to implement it in a different patch. There are many possible combinations and each one may have different preconditions.

[hfinkel]

clang/test/OpenMP/target_parallel_for_is_device_ptr_messages.cpp

93 ↗

(On Diff #213749)

I have no objection to adding warnings in separate patch. I simply wanted to provide some feedback on the general conditions under which we should consider adding warnings. Thanks, Alexey.

[jdenny]

In D65835#1619560, @ABataev wrote:

In D65835#1619549, @jdenny wrote:

In D65835#1619526, @ABataev wrote:

Maybe, but I haven't found any statement in either version that states that map restrictions apply to is_device_ptr.

is_device_ptr is a kind of mapping clause. I assume we can extend the restrictions for map clause to this clause too.

I'd like to understand this better. Is there something from the spec we can quote in the code?

See 2.19.7 Data-Mapping Attribute Rules, Clauses, and Directives

I looked again. I'm still not finding any text in that section that implies is_device_ptr follows the same restrictions as map. Can you please cite specific lines of text instead of an entire section? Thanks for your help.

clang/test/OpenMP/target_parallel_for_is_device_ptr_messages.cpp

93 ↗

(On Diff #213749)

Thanks for the discussion. It sounds like people are fine if (1) the diagnostic proposed here would be a warning not an error and (2) that warning would not be implemented by this patch.

[ABataev]

In D65835#1621168, @jdenny wrote:

In D65835#1619560, @ABataev wrote:

In D65835#1619549, @jdenny wrote:

In D65835#1619526, @ABataev wrote:

Maybe, but I haven't found any statement in either version that states that map restrictions apply to is_device_ptr.

is_device_ptr is a kind of mapping clause. I assume we can extend the restrictions for map clause to this clause too.

I'd like to understand this better. Is there something from the spec we can quote in the code?

See 2.19.7 Data-Mapping Attribute Rules, Clauses, and Directives

I looked again. I'm still not finding any text in that section that implies is_device_ptr follows the same restrictions as map. Can you please cite specific lines of text instead of an entire section? Thanks for your help.

Ah, it is only in OpenMP 5.0 anв restrictions for the map clause are for map clause only. Then we should allow is_device_ptr with the private clauses for OpenMP 4.5.
Better to do this in a separate patch only for is_device_ptr.

[jdoerfert]

clang/test/OpenMP/target_parallel_for_is_device_ptr_messages.cpp

93 ↗

(On Diff #213749)

yes

[jdenny]

In D65835#1621202, @ABataev wrote:

See 2.19.7 Data-Mapping Attribute Rules, Clauses, and Directives

I looked again. I'm still not finding any text in that section that implies is_device_ptr follows the same restrictions as map. Can you please cite specific lines of text instead of an entire section? Thanks for your help.

Ah, it is only in OpenMP 5.0 anв restrictions for the map clause are for map clause only. Then we should allow is_device_ptr with the private clauses for OpenMP 4.5.

I want to be sure we're on the same page: For OpenMP 5.0, should we allow is_device_ptr with the private clauses?

Better to do this in a separate patch only for is_device_ptr.

I'll remove it from this one. Thanks.

[ABataev]

I want to be sure we're on the same page: For OpenMP 5.0, should we allow is_device_ptr with the private clauses?

Yes, since it is allowed by the standard.

[kkwli0]

In D65835#1622042, @ABataev wrote:

I want to be sure we're on the same page: For OpenMP 5.0, should we allow is_device_ptr with the private clauses?

Yes, since it is allowed by the standard.

Umm ... I probably missed some earlier discussions! What would be the behavior of the following code?

p = omp_target_alloc(...);
#pragma omp target private(p) is_device_ptr(p)
  p[...] = ...;   // crash or not?

[ABataev]

In D65835#1623756, @kkwli0 wrote:

In D65835#1622042, @ABataev wrote:

I want to be sure we're on the same page: For OpenMP 5.0, should we allow is_device_ptr with the private clauses?

Yes, since it is allowed by the standard.

Umm ... I probably missed some earlier discussions! What would be the behavior of the following code?

p = omp_target_alloc(...);
#pragma omp target private(p) is_device_ptr(p)
  p[...] = ...;   // crash or not?

It must crush, I assume. The main problem is that this construct is allowed by the standard.

[hfinkel]

In D65835#1623772, @ABataev wrote:

In D65835#1623756, @kkwli0 wrote:

In D65835#1622042, @ABataev wrote:

I want to be sure we're on the same page: For OpenMP 5.0, should we allow is_device_ptr with the private clauses?

Yes, since it is allowed by the standard.

Umm ... I probably missed some earlier discussions! What would be the behavior of the following code?

p = omp_target_alloc(...);
#pragma omp target private(p) is_device_ptr(p)
  p[...] = ...;   // crash or not?

It must crush, I assume. The main problem is that this construct is allowed by the standard.

Yep. We should add a warning message for it.

[jdenny]

In D65835#1623814, @hfinkel wrote:

In D65835#1623772, @ABataev wrote:

In D65835#1623756, @kkwli0 wrote:

In D65835#1622042, @ABataev wrote:

I want to be sure we're on the same page: For OpenMP 5.0, should we allow is_device_ptr with the private clauses?

Yes, since it is allowed by the standard.

Umm ... I probably missed some earlier discussions! What would be the behavior of the following code?

p = omp_target_alloc(...);
#pragma omp target private(p) is_device_ptr(p)
  p[...] = ...;   // crash or not?

It must crush, I assume. The main problem is that this construct is allowed by the standard.

Yep. We should add a warning message for it.

Clang currently doesn't permit is_device_ptr with firstprivate either, but I'm not aware of any reason to diagnose that. Is there? And then there are privatization clauses like linear and reduction....

[ABataev]

In D65835#1623854, @jdenny wrote:

In D65835#1623814, @hfinkel wrote:

In D65835#1623772, @ABataev wrote:

In D65835#1623756, @kkwli0 wrote:

In D65835#1622042, @ABataev wrote:

I want to be sure we're on the same page: For OpenMP 5.0, should we allow is_device_ptr with the private clauses?

Yes, since it is allowed by the standard.

Umm ... I probably missed some earlier discussions! What would be the behavior of the following code?

p = omp_target_alloc(...);
#pragma omp target private(p) is_device_ptr(p)
  p[...] = ...;   // crash or not?

It must crush, I assume. The main problem is that this construct is allowed by the standard.

Yep. We should add a warning message for it.

Clang currently doesn't permit is_device_ptr with firstprivate either, but I'm not aware of any reason to diagnose that. Is there? And then there are privatization clauses like linear and reduction....

I assume, we should allow all allowed datasharing clauses along with is_device_ptr clause, including shared, linear, etc.

[jdenny]

In D65835#1623858, @ABataev wrote:

In D65835#1623854, @jdenny wrote:

In D65835#1623814, @hfinkel wrote:

In D65835#1623772, @ABataev wrote:

In D65835#1623756, @kkwli0 wrote:

In D65835#1622042, @ABataev wrote:

I want to be sure we're on the same page: For OpenMP 5.0, should we allow is_device_ptr with the private clauses?

Yes, since it is allowed by the standard.

Umm ... I probably missed some earlier discussions! What would be the behavior of the following code?

p = omp_target_alloc(...);
#pragma omp target private(p) is_device_ptr(p)
  p[...] = ...;   // crash or not?

It must crush, I assume. The main problem is that this construct is allowed by the standard.

Yep. We should add a warning message for it.

Clang currently doesn't permit is_device_ptr with firstprivate either, but I'm not aware of any reason to diagnose that. Is there? And then there are privatization clauses like linear and reduction....

I assume, we should allow all allowed datasharing clauses along with is_device_ptr clause, including shared, linear, etc.

Agreed. I was just wondering if anyone saw an obvious issue that might deserve a warning (as in the private plus is_device_ptr case). I'm not planning to implement warnings right now. I'm just curious for later.

[hfinkel]

In D65835#1623814, @hfinkel wrote:

In D65835#1623772, @ABataev wrote:

In D65835#1623756, @kkwli0 wrote:

In D65835#1622042, @ABataev wrote:

I want to be sure we're on the same page: For OpenMP 5.0, should we allow is_device_ptr with the private clauses?

Yes, since it is allowed by the standard.

Umm ... I probably missed some earlier discussions! What would be the behavior of the following code?

p = omp_target_alloc(...);
#pragma omp target private(p) is_device_ptr(p)
  p[...] = ...;   // crash or not?

It must crush, I assume. The main problem is that this construct is allowed by the standard.

Yep. We should add a warning message for it.

Upon further reflection, this is not clearly allowed by the standard. My experience is that, when reading standards, sometimes things are disallowed by contradiction (i.e., the standard does not define some behavior, and what the standard does say that's relevant is self contradictory). In this case, 2.19.3 says that list items which are privatized (and which are used) undergo replacement (with new items created as specified) while 2.12.5 says that "The is_device_ptr clause is used to indicate that a list item is a device pointer already in the device data environment and that it should be used directly." A given list item cannot simultaneously be "used directly" (2.12.5) and also undergo replacement: "Inside the construct, all references to the original list item are replaced by references to a new list item received by the task or SIMD lane" (2.19.3). Thus, it may be disallowed.

[kkwli0]

In D65835#1623883, @hfinkel wrote:

In D65835#1623814, @hfinkel wrote:

In D65835#1623772, @ABataev wrote:

In D65835#1623756, @kkwli0 wrote:

In D65835#1622042, @ABataev wrote:

I want to be sure we're on the same page: For OpenMP 5.0, should we allow is_device_ptr with the private clauses?

Yes, since it is allowed by the standard.

Umm ... I probably missed some earlier discussions! What would be the behavior of the following code?

p = omp_target_alloc(...);
#pragma omp target private(p) is_device_ptr(p)
  p[...] = ...;   // crash or not?

It must crush, I assume. The main problem is that this construct is allowed by the standard.

Yep. We should add a warning message for it.

Upon further reflection, this is not clearly allowed by the standard. My experience is that, when reading standards, sometimes things are disallowed by contradiction (i.e., the standard does not define some behavior, and what the standard does say that's relevant is self contradictory). In this case, 2.19.3 says that list items which are privatized (and which are used) undergo replacement (with new items created as specified) while 2.12.5 says that "The is_device_ptr clause is used to indicate that a list item is a device pointer already in the device data environment and that it should be used directly." A given list item cannot simultaneously be "used directly" (2.12.5) and also undergo replacement: "Inside the construct, all references to the original list item are replaced by references to a new list item received by the task or SIMD lane" (2.19.3). Thus, it may be disallowed.

That is what I thought. Specifying these two clauses on the target construct creates ambiguity on which p it referred to inside the construct. The private p or the pointer p that stores the device address? I will work with the committee to fix the spec.

[ABataev]

In D65835#1623883, @hfinkel wrote:

In D65835#1623814, @hfinkel wrote:

In D65835#1623772, @ABataev wrote:

In D65835#1623756, @kkwli0 wrote:

In D65835#1622042, @ABataev wrote:

I want to be sure we're on the same page: For OpenMP 5.0, should we allow is_device_ptr with the private clauses?

Yes, since it is allowed by the standard.

Umm ... I probably missed some earlier discussions! What would be the behavior of the following code?

p = omp_target_alloc(...);
#pragma omp target private(p) is_device_ptr(p)
  p[...] = ...;   // crash or not?

It must crush, I assume. The main problem is that this construct is allowed by the standard.

Yep. We should add a warning message for it.

Upon further reflection, this is not clearly allowed by the standard. My experience is that, when reading standards, sometimes things are disallowed by contradiction (i.e., the standard does not define some behavior, and what the standard does say that's relevant is self contradictory). In this case, 2.19.3 says that list items which are privatized (and which are used) undergo replacement (with new items created as specified) while 2.12.5 says that "The is_device_ptr clause is used to indicate that a list item is a device pointer already in the device data environment and that it should be used directly." A given list item cannot simultaneously be "used directly" (2.12.5) and also undergo replacement: "Inside the construct, all references to the original list item are replaced by references to a new list item received by the task or SIMD lane" (2.19.3). Thus, it may be disallowed.

I think, this combination is still allowed. I assume that

#Pragma omp target parallel is_device_ptr(a) <dsa_clause>(a)

is the same as

#pragma omp target is_device_ptr(a)
#pragma omp parallel <dsa_clause>(a)

i.e. datasharing clauses are applied to inner sub-regions, not the target region itself.

[hfinkel]

In D65835#1623903, @ABataev wrote:

In D65835#1623883, @hfinkel wrote:

In D65835#1623814, @hfinkel wrote:

In D65835#1623772, @ABataev wrote:

In D65835#1623756, @kkwli0 wrote:

In D65835#1622042, @ABataev wrote:

I want to be sure we're on the same page: For OpenMP 5.0, should we allow is_device_ptr with the private clauses?

Yes, since it is allowed by the standard.

Umm ... I probably missed some earlier discussions! What would be the behavior of the following code?

p = omp_target_alloc(...);
#pragma omp target private(p) is_device_ptr(p)
  p[...] = ...;   // crash or not?

It must crush, I assume. The main problem is that this construct is allowed by the standard.

Yep. We should add a warning message for it.

Upon further reflection, this is not clearly allowed by the standard. My experience is that, when reading standards, sometimes things are disallowed by contradiction (i.e., the standard does not define some behavior, and what the standard does say that's relevant is self contradictory). In this case, 2.19.3 says that list items which are privatized (and which are used) undergo replacement (with new items created as specified) while 2.12.5 says that "The is_device_ptr clause is used to indicate that a list item is a device pointer already in the device data environment and that it should be used directly." A given list item cannot simultaneously be "used directly" (2.12.5) and also undergo replacement: "Inside the construct, all references to the original list item are replaced by references to a new list item received by the task or SIMD lane" (2.19.3). Thus, it may be disallowed.

I think, this combination is still allowed. I assume that

#Pragma omp target parallel is_device_ptr(a) <dsa_clause>(a)

is the same as

#pragma omp target is_device_ptr(a)
#pragma omp parallel <dsa_clause>(a)

i.e. datasharing clauses are applied to inner sub-regions, not the target region itself.

With the parallel, that makes sense to me. In that case, however, I'd imagine that the privitization works as normal and the code wouldn't crash. Right?

[ABataev]

In D65835#1623939, @hfinkel wrote:

In D65835#1623903, @ABataev wrote:

In D65835#1623883, @hfinkel wrote:

In D65835#1623814, @hfinkel wrote:

In D65835#1623772, @ABataev wrote:

In D65835#1623756, @kkwli0 wrote:

In D65835#1622042, @ABataev wrote:

I want to be sure we're on the same page: For OpenMP 5.0, should we allow is_device_ptr with the private clauses?

Yes, since it is allowed by the standard.

Umm ... I probably missed some earlier discussions! What would be the behavior of the following code?

p = omp_target_alloc(...);
#pragma omp target private(p) is_device_ptr(p)
  p[...] = ...;   // crash or not?

It must crush, I assume. The main problem is that this construct is allowed by the standard.

Yep. We should add a warning message for it.

Upon further reflection, this is not clearly allowed by the standard. My experience is that, when reading standards, sometimes things are disallowed by contradiction (i.e., the standard does not define some behavior, and what the standard does say that's relevant is self contradictory). In this case, 2.19.3 says that list items which are privatized (and which are used) undergo replacement (with new items created as specified) while 2.12.5 says that "The is_device_ptr clause is used to indicate that a list item is a device pointer already in the device data environment and that it should be used directly." A given list item cannot simultaneously be "used directly" (2.12.5) and also undergo replacement: "Inside the construct, all references to the original list item are replaced by references to a new list item received by the task or SIMD lane" (2.19.3). Thus, it may be disallowed.

I think, this combination is still allowed. I assume that

#Pragma omp target parallel is_device_ptr(a) <dsa_clause>(a)

is the same as

#pragma omp target is_device_ptr(a)
#pragma omp parallel <dsa_clause>(a)

i.e. datasharing clauses are applied to inner sub-regions, not the target region itself.

With the parallel, that makes sense to me. In that case, however, I'd imagine that the privitization works as normal and the code wouldn't crash. Right?

It leads to the use of uninitialized memory. We create private non-initialized copy of the pointer, not the copy of the data it addresses. It shall work with firstprivate clause.

[jdenny]

I made the following changes, as suggested:

• Add back restriction for OpenMP < 5.0.
• Remove is_device_ptr changes.

Alexey, you said:

Plus, these changes are not enough to support this new feature from OpenMP 5.0. There definitely must be some changes in the codegen. If the variable is mapped in the target construct, we should not generate a code for the private clause of this variable on the target construct, since, in this case, private clauses are applied for the inner subdirectives, which are the part of the combined directive, but not the target part of the directive.

I haven't spent enough time exploring codegen here in the past, so I'm starting out by investigating the codegen output for various examples. I'm looking for differences between a combined target teams and an equivalent separate target and teams. So far, I see nothing but superficial differences. Do you have an example where there's an undesirable difference?

Thanks.

[ABataev]

In D65835#1624471, @jdenny wrote:

I made the following changes, as suggested:

• Add back restriction for OpenMP < 5.0.
• Remove is_device_ptr changes.

Alexey, you said:

Plus, these changes are not enough to support this new feature from OpenMP 5.0. There definitely must be some changes in the codegen. If the variable is mapped in the target construct, we should not generate a code for the private clause of this variable on the target construct, since, in this case, private clauses are applied for the inner subdirectives, which are the part of the combined directive, but not the target part of the directive.

I haven't spent enough time exploring codegen here in the past, so I'm starting out by investigating the codegen output for various examples. I'm looking for differences between a combined target teams and an equivalent separate target and teams. So far, I see nothing but superficial differences. Do you have an example where there's an undesirable difference?

Thanks.

Try something like target parallel firstprivate (a) map(a). Currently it will create a firstprivate copy of the variable a in target context thought it is not required at all. It may lead to increased register pressure and performance degradation.

[jdenny]

In D65835#1624477, @ABataev wrote:

Try something like target parallel firstprivate (a) map(a). Currently it will create a firstprivate copy of the variable a in target context thought it is not required at all. It may lead to increased register pressure and performance degradation.

Thanks. The only combination that appears to be affected is firstprivate and map for scalar types. I had only tried arrays and structs earlier, but they're not affected. If I had looked a little more closely at the test case this patch already introduced, I would have noticed that int is affected. The problematic analysis is in sema, where there was an apparent assumption that firstprivate wouldn't appear with map due to the previous restriction. This update fixes that.

For my previous update, I meant to point out that I introduced a fixme into the test suite. See the phabricator summary for details.

[ABataev]

In D65835#1624610, @jdenny wrote:

In D65835#1624477, @ABataev wrote:

Try something like target parallel firstprivate (a) map(a). Currently it will create a firstprivate copy of the variable a in target context thought it is not required at all. It may lead to increased register pressure and performance degradation.

Thanks. The only combination that appears to be affected is firstprivate and map for scalar types. I had only tried arrays and structs earlier, but they're not affected. If I had looked a little more closely at the test case this patch already introduced, I would have noticed that int is affected. The problematic analysis is in sema, where there was an apparent assumption that firstprivate wouldn't appear with map due to the previous restriction. This update fixes that.

For my previous update, I meant to point out that I introduced a fixme into the test suite. See the phabricator summary for details.

This is wrong. It affects all possible combinations and not only fof scalar types, all of them are affected.

[jdenny]

In D65835#1624615, @ABataev wrote:

This is wrong. It affects all possible combinations and not only fof scalar types, all of them are affected.

Are you saying the patch isn't sufficient because other types need to be fixed too? Can you give an example?

[ABataev]

In D65835#1624616, @jdenny wrote:

In D65835#1624615, @ABataev wrote:

This is wrong. It affects all possible combinations and not only fof scalar types, all of them are affected.

Are you saying the patch isn't sufficient because other types need to be fixed too? Can you give an example?

Target teams private map will produce extra private in target context, other constructs either.

[jdenny]

In D65835#1624617, @ABataev wrote:

Target teams private map will produce extra private in target context, other constructs either.

Here's what I tried:

int a;
#pragma omp target teams private(a) map(a)
  ;

The same code is generated as for:

int a;
#pragma omp target map(a)
#pragma omp teams private(a)
  ;

I haven't debugged it yet, but it seems orthogonal to whether you have a combined directive, which is what this patch is about.

[ABataev]

In D65835#1624619, @jdenny wrote:

In D65835#1624617, @ABataev wrote:

Target teams private map will produce extra private in target context, other constructs either.

Here's what I tried:

int a;
#pragma omp target teams private(a) map(a)
  ;

The same code is generated as for:

int a;
#pragma omp target map(a)
#pragma omp teams private(a)
  ;

I haven't debugged it yet, but it seems orthogonal to whether you have a combined directive, which is what this patch is about.

Did you check the mapping flags, generated during host codegen? They must be the same. With private clause it may generate just map(alloc) instead of map(tofrom).

[jdenny]

In D65835#1624624, @ABataev wrote:

In D65835#1624619, @jdenny wrote:

In D65835#1624617, @ABataev wrote:

Target teams private map will produce extra private in target context, other constructs either.

Here's what I tried:

int a;
#pragma omp target teams private(a) map(a)
  ;

The same code is generated as for:

int a;
#pragma omp target map(a)
#pragma omp teams private(a)
  ;

I haven't debugged it yet, but it seems orthogonal to whether you have a combined directive, which is what this patch is about.

Did you check the mapping flags, generated during host codegen? They must be the same. With private clause it may generate just map(alloc) instead of map(tofrom).

I diffed the .ll files for combined vs. separate constructs. The only difference is the file ID. @.offload_maptypes isn't generated in either (but it is if I replace private with firstprivate).

[ABataev]

In D65835#1624629, @jdenny wrote:

In D65835#1624624, @ABataev wrote:

In D65835#1624619, @jdenny wrote:

In D65835#1624617, @ABataev wrote:

Target teams private map will produce extra private in target context, other constructs either.

Here's what I tried:

int a;
#pragma omp target teams private(a) map(a)
  ;

The same code is generated as for:

int a;
#pragma omp target map(a)
#pragma omp teams private(a)
  ;

I haven't debugged it yet, but it seems orthogonal to whether you have a combined directive, which is what this patch is about.

Did you check the mapping flags, generated during host codegen? They must be the same. With private clause it may generate just map(alloc) instead of map(tofrom).

I diffed the .ll files for combined vs. separate constructs. The only difference is the file ID. @.offload_maptypes isn't generated in either (but it is if I replace private with firstprivate).

Maptypes array must be generated in all cases, check the host codegen.
Also, test codegen with the different kinds of maptypes, not only to from, but also alloc, to, from, etc. Yoh will see the difference in many cases and 2ith many kinds of types, not only scalars.

[jdenny]

In D65835#1624659, @ABataev wrote:

In D65835#1624629, @jdenny wrote:

In D65835#1624624, @ABataev wrote:

In D65835#1624619, @jdenny wrote:

In D65835#1624617, @ABataev wrote:

Target teams private map will produce extra private in target context, other constructs either.

Here's what I tried:

int a;
#pragma omp target teams private(a) map(a)
  ;

The same code is generated as for:

int a;
#pragma omp target map(a)
#pragma omp teams private(a)
  ;

I haven't debugged it yet, but it seems orthogonal to whether you have a combined directive, which is what this patch is about.

Did you check the mapping flags, generated during host codegen? They must be the same. With private clause it may generate just map(alloc) instead of map(tofrom).

I diffed the .ll files for combined vs. separate constructs. The only difference is the file ID. @.offload_maptypes isn't generated in either (but it is if I replace private with firstprivate).

Maptypes array must be generated in all cases, check the host codegen.
Also, test codegen with the different kinds of maptypes, not only to from, but also alloc, to, from, etc. Yoh will see the difference in many cases and 2ith many kinds of types, not only scalars.

I'll work on another patch for the bugs where directives are separate, and I'll return to this patch if more fixes are then needed to support combined directives consistently. Thanks.

[jdenny]

Rebase. Add more tests, as requested at D66247#1630441.

[ABataev]

Why there are the changes from the another patch?

[jdenny]

I think you're referring to stuff like OpenMPCaptureLevel in ScopeInfo.h. I wrote those changes for this patch first, as can be seen in phabricator history. I needed them for the other patch too, and I thought we were going to end up applying that patch first, so I was going to strip them out here. Of course, as I understand it, the plan now is to apply only this patch.

[jdenny]

Ping.

[ABataev]

In D65835#1630732, @jdenny wrote:

I think you're referring to stuff like OpenMPCaptureLevel in ScopeInfo.h. I wrote those changes for this patch first, as can be seen in phabricator history. I needed them for the other patch too, and I thought we were going to end up applying that patch first, so I was going to strip them out here. Of course, as I understand it, the plan now is to apply only this patch.

These changes required for this functionality, please remove.

[jdenny]

As requested, backed out the changes related to firstprivate scalars, and updated tests.

[ABataev]

LG

[jdenny]

I want to be sure we're on the same page. Due to the changes I just backed out, the following two examples now generate different code:

int a = 0;
#pragma omp target map(a)
#pragma omp teams firstprivate(a)
  ;

int a = 0;
#pragma omp target teams firstprivate(a) map(a)
  ;

The difference is whether a is passed by reference (the first case) or value (the second case) to the offloading function.

Is that fine for you?

[ABataev]

In D65835#1639584, @jdenny wrote:

I want to be sure we're on the same page. Due to the changes I just backed out, the following two examples now generate different code:

int a = 0;
#pragma omp target map(a)
#pragma omp teams firstprivate(a)
  ;

int a = 0;
#pragma omp target teams firstprivate(a) map(a)
  ;

The difference is whether a is passed by reference (the first case) or value (the second case) to the offloading function.

Is that fine for you?

No, this is what I warned about. We shall have the same codegen just like in the first case, the value must be passed by reference and mapped as tofrom.

[ABataev]

Need to fix the mapping

[jdenny]

In D65835#1639585, @ABataev wrote:

In D65835#1639584, @jdenny wrote:

I want to be sure we're on the same page. Due to the changes I just backed out, the following two examples now generate different code:

int a = 0;
#pragma omp target map(a)
#pragma omp teams firstprivate(a)
  ;

int a = 0;
#pragma omp target teams firstprivate(a) map(a)
  ;

The difference is whether a is passed by reference (the first case) or value (the second case) to the offloading function.

Is that fine for you?

No, this is what I warned about. We shall have the same codegen just like in the first case, the value must be passed by reference and mapped as tofrom.

If I add back those changes I just backed out, we get the same codegen. Is that what you want?

[ABataev]

In D65835#1639593, @jdenny wrote:

In D65835#1639585, @ABataev wrote:

In D65835#1639584, @jdenny wrote:

I want to be sure we're on the same page. Due to the changes I just backed out, the following two examples now generate different code:

int a = 0;
#pragma omp target map(a)
#pragma omp teams firstprivate(a)
  ;

int a = 0;
#pragma omp target teams firstprivate(a) map(a)
  ;

The difference is whether a is passed by reference (the first case) or value (the second case) to the offloading function.

Is that fine for you?

No, this is what I warned about. We shall have the same codegen just like in the first case, the value must be passed by reference and mapped as tofrom.

If I add back those changes I just backed out, we get the same codegen. Is that what you want?

Those 2 cases must result in the same codegen. But I rather doubt we need your previous changes. Check Sema::isOpenMPCapturedByRef instead, required functionality must be handled in this function.

[jdenny]

In D65835#1639612, @ABataev wrote:

In D65835#1639593, @jdenny wrote:

In D65835#1639585, @ABataev wrote:

In D65835#1639584, @jdenny wrote:

I want to be sure we're on the same page. Due to the changes I just backed out, the following two examples now generate different code:

int a = 0;
#pragma omp target map(a)
#pragma omp teams firstprivate(a)
  ;

int a = 0;
#pragma omp target teams firstprivate(a) map(a)
  ;

The difference is whether a is passed by reference (the first case) or value (the second case) to the offloading function.

Is that fine for you?

No, this is what I warned about. We shall have the same codegen just like in the first case, the value must be passed by reference and mapped as tofrom.

If I add back those changes I just backed out, we get the same codegen. Is that what you want?

Those 2 cases must result in the same codegen. But I rather doubt we need your previous changes. Check Sema::isOpenMPCapturedByRef instead, required functionality must be handled in this function.

That's the focus of my previous changes. The rest just supports the changes there.

[ABataev]

In D65835#1639617, @jdenny wrote:

In D65835#1639612, @ABataev wrote:

In D65835#1639593, @jdenny wrote:

In D65835#1639585, @ABataev wrote:

In D65835#1639584, @jdenny wrote:

I want to be sure we're on the same page. Due to the changes I just backed out, the following two examples now generate different code:

int a = 0;
#pragma omp target map(a)
#pragma omp teams firstprivate(a)
  ;

int a = 0;
#pragma omp target teams firstprivate(a) map(a)
  ;

The difference is whether a is passed by reference (the first case) or value (the second case) to the offloading function.

Is that fine for you?

No, this is what I warned about. We shall have the same codegen just like in the first case, the value must be passed by reference and mapped as tofrom.

If I add back those changes I just backed out, we get the same codegen. Is that what you want?

Those 2 cases must result in the same codegen. But I rather doubt we need your previous changes. Check Sema::isOpenMPCapturedByRef instead, required functionality must be handled in this function.

That's the focus of my previous changes. The rest just supports the changes there.

We don't need this new level counter to correctly handle this situation. Just check for the combined target directive in Sema::isOpenMPCapturedByRef and return true if it is mapped as to from or just from. The change must be very simple.

[jdenny]

In D65835#1639622, @ABataev wrote:

We don't need this new level counter to correctly handle this situation. Just check for the combined target directive in Sema::isOpenMPCapturedByRef and return true if it is mapped as to from or just from. The change must be very simple.

I don't see a way to do that without also affecting capturing on the teams region in the case of a combined target teams construct.

If it's simple, it might save time for both of us if you post the change you're thinking of here.

[ABataev]

In D65835#1639685, @jdenny wrote:

In D65835#1639622, @ABataev wrote:

We don't need this new level counter to correctly handle this situation. Just check for the combined target directive in Sema::isOpenMPCapturedByRef and return true if it is mapped as to from or just from. The change must be very simple.

I don't see a way to do that without also affecting capturing on the teams region in the case of a combined target teams construct.

If it's simple, it might save time for both of us if you post the change you're thinking of here.

Chwck 2 ifs: if (IsVariableUsedInMapClause) and the second if (IsByRef && Ty.getNonReferenceType()->isScalarType()). If the variable is mapped, IsByRef is set to true, but later it may be changed in the second if, when we check for the use in the firstprivate clause. We need to add a check, that the IsByRef must be set to false onlly if IsVariableUsedInMapClause is set to false too for combined target constructs.

[jdenny]

In D65835#1639700, @ABataev wrote:

Chwck 2 ifs: if (IsVariableUsedInMapClause) and the second if (IsByRef && Ty.getNonReferenceType()->isScalarType()). If the variable is mapped, IsByRef is set to true, but later it may be changed in the second if, when we check for the use in the firstprivate clause. We need to add a check, that the IsByRef must be set to false onlly if IsVariableUsedInMapClause is set to false too for combined target constructs.

I've tried that. It results in capturing by reference on both the target region and the teams region in the case of a combined target teams construct. Thus, we end up with different codegen than when the directives are separate, where capturing by reference is set only on the target region.

The changes I just backed out handle this correctly because they distinguish between the target and teams regions.

[ABataev]

In D65835#1639739, @jdenny wrote:

In D65835#1639700, @ABataev wrote:

Chwck 2 ifs: if (IsVariableUsedInMapClause) and the second if (IsByRef && Ty.getNonReferenceType()->isScalarType()). If the variable is mapped, IsByRef is set to true, but later it may be changed in the second if, when we check for the use in the firstprivate clause. We need to add a check, that the IsByRef must be set to false onlly if IsVariableUsedInMapClause is set to false too for combined target constructs.

I've tried that. It results in capturing by reference on both the target region and the teams region in the case of a combined target teams construct. Thus, we end up with different codegen than when the directives are separate, where capturing by reference is set only on the target region.

The changes I just backed out handle this correctly because they distinguish between the target and teams regions.

Ahh, I see, in this case, the codegen for the inner subregion is changed. Ok, I see now. Then yes, we need this change with subregion counter.
In this case, I would recommend trying to fix https://bugs.llvm.org/show_bug.cgi?id=40253 at first, as it relates to the same problem, I think. Better to fix the problem separately and then extend it for new feature rather than add 2 new features in one patch.

[jdenny]

Restore previous version of patch, and rebase.

I tried, and this patch is not sufficient to fix PR40253. If there are indeed common changes, it seems it's just as well to put them here.

[ABataev]

clang/include/clang/Sema/ScopeInfo.h

763 ↗

(On Diff #216444)

I would add a parameter for OpenMPCaptureLevel rather than use default value.

clang/include/clang/Sema/Sema.h

9026 ↗

(On Diff #216444)

Do not use default value here, just set it to 0 in the call in SemaOpenMP.cpp. We use it there tat target level always and thus CaptureLevel is always 0 there.

clang/lib/Sema/Sema.cpp

2108 ↗

(On Diff #216444)

I would suggest adding a new parameter for CapturedLevel and the same parameter to ActOnCapturedRegionStart function. And pass the value for CapturedLevel explicitly at ActOnCapturedRegionStart function call in SemaOpenMP.cpp rather than rely on this solution.
Also, for each separate OpenMP directive, this counter must start from 0. Currently, I think, it may lead to crash if we have several directives. Something like:

#pragma omp parallel
#pragma omp target teams
...

I think it will work incorrectly for this example.

2109 ↗

(On Diff #216444)

CapturedRegionScopeInfo *->auto *

[jdenny]

clang/lib/Sema/Sema.cpp

2108 ↗

(On Diff #216444)

I think the check that OpenMPLevel is the same prevents the problem you mention.

[ABataev]

clang/lib/Sema/Sema.cpp

2108 ↗

(On Diff #216444)

Ah, yes, missed this. Anyway, better to specify this explicitly as function parameters, I think, rather than trying to evaluate it here this way. We eventually may change the handling of captured regions in OpenMP, better to have everything in SemaOpenMP.cpp

[ABataev]

clang/lib/Sema/Sema.cpp

2108 ↗

(On Diff #216444)

BTW, here you can use the parameter with the default value 0.

[jdenny]

clang/lib/Sema/Sema.cpp

2108 ↗

(On Diff #216444)

OK.

CaptureLevel or CapturedLevel? Do you care?

[ABataev]

clang/lib/Sema/Sema.cpp

2108 ↗

(On Diff #216444)

Up to you

[jdenny]

Make suggested changes for passing around the capture level.

[ABataev]

clang/include/clang/Sema/Sema.h

1419 ↗

(On Diff #216469)

Better to use OpenMPCaptureLevel since this param is OpenMP specific.

1419 ↗

(On Diff #216469)

Add default value for the parameter here too.

3977 ↗

(On Diff #216469)

Same, add prefix OpenMP

9027 ↗

(On Diff #216469)

Same, OpenMPCaptureLevel here and in other places.

clang/lib/Sema/SemaExpr.cpp

17748 ↗

(On Diff #216469)

Restore original code here

clang/lib/Sema/SemaOpenMP.cpp

1350 ↗

(On Diff #216469)

Use /*param=*/ format, please.

2973 ↗

(On Diff #216469)

Same for param comment here and other places.

clang/lib/Sema/SemaStmt.cpp

4324 ↗

(On Diff #216469)

Use default value of the parameter, no need to pass 0 here

[jdenny]

Make suggested changes to default arguments, comments on literals, and parameter names.

[jdenny]

clang/lib/Sema/SemaExpr.cpp

17748 ↗

(On Diff #216469)

OK, I did. What's the reason for no newline at the end of the file?

[ABataev]

clang/lib/Sema/SemaExpr.cpp

17748 ↗

(On Diff #216469)

Changes not related to the patch, like formatting etc., better to commit in separate patches.

[ABataev]

LG

clang/lib/Sema/SemaExpr.cpp

17749 ↗

(On Diff #216502)

Still marked as changed code, better to restore it completely.

[jdenny]

Thanks.

clang/lib/Sema/SemaExpr.cpp

17749 ↗

(On Diff #216502)

There's no change here. I checked the diff I uploaded to be sure.

17748 ↗

(On Diff #216469)

Ah.

cfe/trunk/include/clang/Sema/ScopeInfo.h

public:
 
   /// The implicit parameter for the captured variables.
   ImplicitParamDecl *ContextParam;
 
   /// The kind of captured region.
   unsigned short CapRegionKind;
 
   unsigned short OpenMPLevel;
+  unsigned short OpenMPCaptureLevel;
 
   CapturedRegionScopeInfo(DiagnosticsEngine &Diag, Scope *S, CapturedDecl *CD,
                           RecordDecl *RD, ImplicitParamDecl *Context,
-                          CapturedRegionKind K, unsigned OpenMPLevel)
+                          CapturedRegionKind K, unsigned OpenMPLevel,
+                          unsigned OpenMPCaptureLevel)
       : CapturingScopeInfo(Diag, ImpCap_CapturedRegion),
         TheCapturedDecl(CD), TheRecordDecl(RD), TheScope(S),
-        ContextParam(Context), CapRegionKind(K), OpenMPLevel(OpenMPLevel) {
+        ContextParam(Context), CapRegionKind(K), OpenMPLevel(OpenMPLevel),
+        OpenMPCaptureLevel(OpenMPCaptureLevel) {
     Kind = SK_CapturedRegion;
   }
 
   ~CapturedRegionScopeInfo() override;
 
   /// A descriptive name for the kind of captured region this is.
   StringRef getRegionName() const {
     switch (CapRegionKind) {

cfe/trunk/include/clang/Sema/Sema.h

public:
   sema::LambdaScopeInfo *PushLambdaScope();
 
   /// This is used to inform Sema what the current TemplateParameterDepth
   /// is during Parsing.  Currently it is used to pass on the depth
   /// when parsing generic lambda 'auto' parameters.
   void RecordParsingTemplateParameterDepth(unsigned Depth);
 
   void PushCapturedRegionScope(Scope *RegionScope, CapturedDecl *CD,
-                               RecordDecl *RD,
-                               CapturedRegionKind K);
+                               RecordDecl *RD, CapturedRegionKind K,
+                               unsigned OpenMPCaptureLevel = 0);
 
   /// Custom deleter to allow FunctionScopeInfos to be kept alive for a short
   /// time after they've been popped.
   class PoppedFunctionScopeDeleter {
     Sema *Self;
 
   public:
     explicit PoppedFunctionScopeDeleter(Sema *Self) : Self(Self) {}
public:
   StmtResult ActOnContinueStmt(SourceLocation ContinueLoc, Scope *CurScope);
   StmtResult ActOnBreakStmt(SourceLocation BreakLoc, Scope *CurScope);
 
   void ActOnCapturedRegionStart(SourceLocation Loc, Scope *CurScope,
                                 CapturedRegionKind Kind, unsigned NumParams);
   typedef std::pair<StringRef, QualType> CapturedParamNameType;
   void ActOnCapturedRegionStart(SourceLocation Loc, Scope *CurScope,
                                 CapturedRegionKind Kind,
-                                ArrayRef<CapturedParamNameType> Params);
+                                ArrayRef<CapturedParamNameType> Params,
+                                unsigned OpenMPCaptureLevel = 0);
   StmtResult ActOnCapturedRegionEnd(Stmt *S);
   void ActOnCapturedRegionError();
   RecordDecl *CreateCapturedStmtRecordDecl(CapturedDecl *&CD,
                                            SourceLocation Loc,
                                            unsigned NumParams);
 
   enum CopyElisionSemanticsKind {
     CES_Strict = 0,
public:
   /// 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
   /// is performed.
-  bool isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level) const;
+  /// \param OpenMPCaptureLevel Capture level within an OpenMP construct.
+  bool isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level,
+                             unsigned OpenMPCaptureLevel) const;
 
   /// Check if the specified variable is used in one of the private
   /// clauses (private, firstprivate, lastprivate, reduction etc.) in OpenMP
   /// constructs.
   VarDecl *isOpenMPCapturedDecl(ValueDecl *D, bool CheckScopeInfo = false,
                                 unsigned StopAt = 0);
   ExprResult getOpenMPCapturedExpr(VarDecl *Capture, ExprValueKind VK,
                                    ExprObjectKind OK, SourceLocation Loc);

cfe/trunk/lib/Sema/Sema.cpp

 
 IdentifierInfo *Sema::getFloat128Identifier() const {
   if (!Ident___float128)
     Ident___float128 = &Context.Idents.get("__float128");
   return Ident___float128;
 }
 
 void Sema::PushCapturedRegionScope(Scope *S, CapturedDecl *CD, RecordDecl *RD,
-                                   CapturedRegionKind K) {
-  CapturingScopeInfo *CSI = new CapturedRegionScopeInfo(
+                                   CapturedRegionKind K,
+                                   unsigned OpenMPCaptureLevel) {
+  auto *CSI = new CapturedRegionScopeInfo(
       getDiagnostics(), S, CD, RD, CD->getContextParam(), K,
-      (getLangOpts().OpenMP && K == CR_OpenMP) ? getOpenMPNestingLevel() : 0);
+      (getLangOpts().OpenMP && K == CR_OpenMP) ? getOpenMPNestingLevel() : 0,
+      OpenMPCaptureLevel);
   CSI->ReturnType = Context.VoidTy;
   FunctionScopes.push_back(CSI);
 }
 
 CapturedRegionScopeInfo *Sema::getCurCapturedRegion() {
   if (FunctionScopes.empty())
     return nullptr;
 

cfe/trunk/lib/Sema/SemaExpr.cpp

static bool captureInCapturedRegion(CapturedRegionScopeInfo *RSI,
   if (S.getLangOpts().OpenMP && RSI->CapRegionKind == CR_OpenMP) {
     if (S.isOpenMPCapturedDecl(Var)) {
       bool HasConst = DeclRefType.isConstQualified();
       DeclRefType = DeclRefType.getUnqualifiedType();
       // Don't lose diagnostics about assignments to const.
       if (HasConst)
         DeclRefType.addConst();
     }
-    ByRef = S.isOpenMPCapturedByRef(Var, RSI->OpenMPLevel);
+    ByRef = S.isOpenMPCapturedByRef(Var, RSI->OpenMPLevel,
+                                    RSI->OpenMPCaptureLevel);
   }
 
   if (ByRef)
     CaptureType = S.Context.getLValueReferenceType(DeclRefType);
   else
     CaptureType = DeclRefType;
 
   // Actually capture the variable.
ExprResult Sema::ActOnObjCAvailabilityCheckExpr(
   if (getCurFunctionOrMethodDecl())
     getEnclosingFunction()->HasPotentialAvailabilityViolations = true;
   else if (getCurBlock() || getCurLambda())
     getCurFunction()->HasPotentialAvailabilityViolations = true;
 
   return new (Context)
       ObjCAvailabilityCheckExpr(Version, AtLoc, RParen, Context.BoolTy);
 }
 No newline at end of file

cfe/trunk/lib/Sema/SemaOpenMP.cpp

OpenMPDirectiveKind getCurrentDirective() const {
     const SharingMapTy *Top = getTopOfStackOrNull();
     return Top ? Top->Directive : OMPD_unknown;
   }
   /// Returns directive kind at specified level.
   OpenMPDirectiveKind getDirective(unsigned Level) const {
     assert(!isStackEmpty() && "No directive at specified level.");
     return getStackElemAtLevel(Level).Directive;
   }
+  /// Returns the capture region at the specified level.
+  OpenMPDirectiveKind getCaptureRegion(unsigned Level,
+                                       unsigned OpenMPCaptureLevel) const {
+    SmallVector<OpenMPDirectiveKind, 4> CaptureRegions;
+    getOpenMPCaptureRegions(CaptureRegions, getDirective(Level));
+    return CaptureRegions[OpenMPCaptureLevel];
+  }
   /// Returns parent directive.
   OpenMPDirectiveKind getParentDirective() const {
     const SharingMapTy *Parent = getSecondOnStackOrNull();
     return Parent ? Parent->Directive : OMPD_unknown;
   }
 
   /// Add requires decl to internal vector
   void addRequiresDecl(OMPRequiresDecl *RD) {
if (IterTarget != end()) {
         if (DSAIter != IterTarget->SharingMap.end() &&
             isOpenMPPrivate(DSAIter->getSecond().Attributes)) {
           DVar.RefExpr = DSAIter->getSecond().RefExpr.getPointer();
           DVar.CKind = OMPC_threadprivate;
           return DVar;
         }
         const_iterator End = end();
         if (!SemaRef.isOpenMPCapturedByRef(
-                D, std::distance(ParentIterTarget, End))) {
+                D, std::distance(ParentIterTarget, End),
+                /*OpenMPCaptureLevel=*/0)) {
           DVar.RefExpr =
               buildDeclRefExpr(SemaRef, VD, D->getType().getNonReferenceType(),
                                IterTarget->ConstructLoc);
           DVar.CKind = OMPC_threadprivate;
           return DVar;
         }
       }
     }
if ((Ty->isFloat16Type() && !Context.getTargetInfo().hasFloat16Type()) ||
        !Context.getTargetInfo().hasFloat128Type()) ||
       (Ty->isIntegerType() && Context.getTypeSize(Ty) == 128 &&
        !Context.getTargetInfo().hasInt128Type()))
     targetDiag(E->getExprLoc(), diag::err_omp_unsupported_type)
         << static_cast<unsigned>(Context.getTypeSize(Ty)) << Ty
         << Context.getTargetInfo().getTriple().str() << E->getSourceRange();
 }
 
-bool Sema::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level) const {
+bool Sema::isOpenMPCapturedByRef(const ValueDecl *D, unsigned Level,
+                                 unsigned OpenMPCaptureLevel) const {
   assert(LangOpts.OpenMP && "OpenMP is not allowed");
 
   ASTContext &Ctx = getASTContext();
   bool IsByRef = true;
 
   // Find the directive that is associated with the provided scope.
   D = cast<ValueDecl>(D->getCanonicalDecl());
   QualType Ty = D->getType();
 
+  bool IsVariableUsedInMapClause = false;
   if (DSAStack->hasExplicitDirective(isOpenMPTargetExecutionDirective, Level)) {
     // This table summarizes how a given variable should be passed to the device
     // given its type and the clauses where it appears. This table is based on
     // the description in OpenMP 4.5 [2.10.4, target Construct] and
     // OpenMP 4.5 [2.15.5, Data-mapping Attribute Rules and Clauses].
     //
     // =========================================================================
     // | type |  defaultmap   | pvt | first | is_device_ptr |    map   | res.  |
if (DSAStack->hasExplicitDirective(isOpenMPTargetExecutionDirective, Level)) {
     //    device instead of the value passed to it by the compiler.
 
     if (Ty->isReferenceType())
       Ty = Ty->castAs<ReferenceType>()->getPointeeType();
 
     // Locate map clauses and see if the variable being captured is referred to
     // in any of those clauses. Here we only care about variables, not fields,
     // because fields are part of aggregates.
-    bool IsVariableUsedInMapClause = false;
     bool IsVariableAssociatedWithSection = false;
 
     DSAStack->checkMappableExprComponentListsForDeclAtLevel(
         D, Level,
         [&IsVariableUsedInMapClause, &IsVariableAssociatedWithSection, D](
             OMPClauseMappableExprCommon::MappableExprComponentListRef
                 MapExprComponents,
             OpenMPClauseKind WhereFoundClauseKind) {
if (IsVariableUsedInMapClause) {
           DSAStack->getDefaultDMAAtLevel(Level) == DMA_tofrom_scalar ||
           DSAStack->hasExplicitDSA(
               D, [](OpenMPClauseKind K) { return K == OMPC_reduction; }, Level);
     }
   }
 
   if (IsByRef && Ty.getNonReferenceType()->isScalarType()) {
     IsByRef =
+        ((IsVariableUsedInMapClause &&
+          DSAStack->getCaptureRegion(Level, OpenMPCaptureLevel) ==
+              OMPD_target) ||
          !DSAStack->hasExplicitDSA(
              D,
              [](OpenMPClauseKind K) -> bool { return K == OMPC_firstprivate; },
-            Level, /*NotLastprivate=*/true) &&
+             Level, /*NotLastprivate=*/true)) &&
         // If the variable is artificial and must be captured by value - try to
         // capture by value.
         !(isa<OMPCapturedExprDecl>(D) && !D->hasAttr<OMPCaptureNoInitAttr>() &&
           !cast<OMPCapturedExprDecl>(D)->getInit()->isGLValue());
   }
 
   // When passing data by copy, we need to make sure it fits the uintptr size
   // and alignment, because the runtime library only deals with uintptr types.
Sema::CapturedParamNameType Params[] = {
         std::make_pair(".privates.", VoidPtrTy),
         std::make_pair(
             ".copy_fn.",
             Context.getPointerType(CopyFnType).withConst().withRestrict()),
         std::make_pair(".task_t.", Context.VoidPtrTy.withConst()),
         std::make_pair(StringRef(), QualType()) // __context with shared vars
     };
     ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
-                             Params);
+                             Params, /*OpenMPCaptureLevel=*/0);
     // Mark this captured region as inlined, because we don't use outlined
     // function directly.
     getCurCapturedRegion()->TheCapturedDecl->addAttr(
         AlwaysInlineAttr::CreateImplicit(
             Context, AlwaysInlineAttr::Keyword_forceinline));
     Sema::CapturedParamNameType ParamsTarget[] = {
         std::make_pair(StringRef(), QualType()) // __context with shared vars
     };
     // Start a captured region for 'target' with no implicit parameters.
     ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
-                             ParamsTarget);
+                             ParamsTarget, /*OpenMPCaptureLevel=*/1);
     Sema::CapturedParamNameType ParamsTeamsOrParallel[] = {
         std::make_pair(".global_tid.", KmpInt32PtrTy),
         std::make_pair(".bound_tid.", KmpInt32PtrTy),
         std::make_pair(StringRef(), QualType()) // __context with shared vars
     };
     // Start a captured region for 'teams' or 'parallel'.  Both regions have
     // the same implicit parameters.
     ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
-                             ParamsTeamsOrParallel);
+                             ParamsTeamsOrParallel, /*OpenMPCaptureLevel=*/2);
     break;
   }
   case OMPD_target:
   case OMPD_target_simd: {
     QualType KmpInt32Ty = Context.getIntTypeForBitwidth(32, 1).withConst();
     QualType VoidPtrTy = Context.VoidPtrTy.withConst().withRestrict();
     QualType KmpInt32PtrTy =
         Context.getPointerType(KmpInt32Ty).withConst().withRestrict();
     QualType Args[] = {VoidPtrTy};
     FunctionProtoType::ExtProtoInfo EPI;
     EPI.Variadic = true;
     QualType CopyFnType = Context.getFunctionType(Context.VoidTy, Args, EPI);
     Sema::CapturedParamNameType Params[] = {
         std::make_pair(".global_tid.", KmpInt32Ty),
         std::make_pair(".part_id.", KmpInt32PtrTy),
         std::make_pair(".privates.", VoidPtrTy),
         std::make_pair(
             ".copy_fn.",
             Context.getPointerType(CopyFnType).withConst().withRestrict()),
         std::make_pair(".task_t.", Context.VoidPtrTy.withConst()),
         std::make_pair(StringRef(), QualType()) // __context with shared vars
     };
     ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
-                             Params);
+                             Params, /*OpenMPCaptureLevel=*/0);
     // Mark this captured region as inlined, because we don't use outlined
     // function directly.
     getCurCapturedRegion()->TheCapturedDecl->addAttr(
         AlwaysInlineAttr::CreateImplicit(
             Context, AlwaysInlineAttr::Keyword_forceinline));
     ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
-                             std::make_pair(StringRef(), QualType()));
+                             std::make_pair(StringRef(), QualType()),
+                             /*OpenMPCaptureLevel=*/1);
     break;
   }
   case OMPD_simd:
   case OMPD_for:
   case OMPD_for_simd:
   case OMPD_sections:
   case OMPD_section:
   case OMPD_single:
Sema::CapturedParamNameType Params[] = {
         std::make_pair(".privates.", VoidPtrTy),
         std::make_pair(
             ".copy_fn.",
             Context.getPointerType(CopyFnType).withConst().withRestrict()),
         std::make_pair(".task_t.", Context.VoidPtrTy.withConst()),
         std::make_pair(StringRef(), QualType()) // __context with shared vars
     };
     ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
-                             Params);
+                             Params, /*OpenMPCaptureLevel=*/0);
     // Mark this captured region as inlined, because we don't use outlined
     // function directly.
     getCurCapturedRegion()->TheCapturedDecl->addAttr(
         AlwaysInlineAttr::CreateImplicit(
             Context, AlwaysInlineAttr::Keyword_forceinline));
     Sema::CapturedParamNameType ParamsTarget[] = {
         std::make_pair(StringRef(), QualType()) // __context with shared vars
     };
     // Start a captured region for 'target' with no implicit parameters.
     ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
-                             ParamsTarget);
+                             ParamsTarget, /*OpenMPCaptureLevel=*/1);
 
     Sema::CapturedParamNameType ParamsTeams[] = {
         std::make_pair(".global_tid.", KmpInt32PtrTy),
         std::make_pair(".bound_tid.", KmpInt32PtrTy),
         std::make_pair(StringRef(), QualType()) // __context with shared vars
     };
     // Start a captured region for 'target' with no implicit parameters.
     ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
-                             ParamsTeams);
+                             ParamsTeams, /*OpenMPCaptureLevel=*/2);
 
     Sema::CapturedParamNameType ParamsParallel[] = {
         std::make_pair(".global_tid.", KmpInt32PtrTy),
         std::make_pair(".bound_tid.", KmpInt32PtrTy),
         std::make_pair(".previous.lb.", Context.getSizeType().withConst()),
         std::make_pair(".previous.ub.", Context.getSizeType().withConst()),
         std::make_pair(StringRef(), QualType()) // __context with shared vars
     };
     // Start a captured region for 'teams' or 'parallel'.  Both regions have
     // the same implicit parameters.
     ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
-                             ParamsParallel);
+                             ParamsParallel, /*OpenMPCaptureLevel=*/3);
     break;
   }
 
   case OMPD_teams_distribute_parallel_for:
   case OMPD_teams_distribute_parallel_for_simd: {
     QualType KmpInt32Ty = Context.getIntTypeForBitwidth(32, 1).withConst();
     QualType KmpInt32PtrTy =
         Context.getPointerType(KmpInt32Ty).withConst().withRestrict();
 
     Sema::CapturedParamNameType ParamsTeams[] = {
         std::make_pair(".global_tid.", KmpInt32PtrTy),
         std::make_pair(".bound_tid.", KmpInt32PtrTy),
         std::make_pair(StringRef(), QualType()) // __context with shared vars
     };
     // Start a captured region for 'target' with no implicit parameters.
     ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
-                             ParamsTeams);
+                             ParamsTeams, /*OpenMPCaptureLevel=*/0);
 
     Sema::CapturedParamNameType ParamsParallel[] = {
         std::make_pair(".global_tid.", KmpInt32PtrTy),
         std::make_pair(".bound_tid.", KmpInt32PtrTy),
         std::make_pair(".previous.lb.", Context.getSizeType().withConst()),
         std::make_pair(".previous.ub.", Context.getSizeType().withConst()),
         std::make_pair(StringRef(), QualType()) // __context with shared vars
     };
     // Start a captured region for 'teams' or 'parallel'.  Both regions have
     // the same implicit parameters.
     ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
-                             ParamsParallel);
+                             ParamsParallel, /*OpenMPCaptureLevel=*/1);
     break;
   }
   case OMPD_target_update:
   case OMPD_target_enter_data:
   case OMPD_target_exit_data: {
     QualType KmpInt32Ty = Context.getIntTypeForBitwidth(32, 1).withConst();
     QualType VoidPtrTy = Context.VoidPtrTy.withConst().withRestrict();
     QualType KmpInt32PtrTy =
if (!Type->isAnyPointerType() && Type->isVariablyModifiedType() &&
            IsDecl ? diag::note_previous_decl : diag::note_defined_here)
           << D;
       continue;
     }
 
     // OpenMP 4.5 [2.15.5.1, Restrictions, p.3]
     // A list item cannot appear in both a map clause and a data-sharing
     // attribute clause on the same construct
-    if (isOpenMPTargetExecutionDirective(CurrDir)) {
+    //
+    // OpenMP 5.0 [2.19.7.1, Restrictions, p.7]
+    // A list item cannot appear in both a map clause and a data-sharing
+    // attribute clause on the same construct unless the construct is a
+    // combined construct.
+    if ((LangOpts.OpenMP <= 45 && isOpenMPTargetExecutionDirective(CurrDir)) ||
+        CurrDir == OMPD_target) {
       OpenMPClauseKind ConflictKind;
       if (DSAStack->checkMappableExprComponentListsForDecl(
               VD, /*CurrentRegionOnly=*/true,
               [&](OMPClauseMappableExprCommon::MappableExprComponentListRef,
                   OpenMPClauseKind WhereFoundClauseKind) -> bool {
                 ConflictKind = WhereFoundClauseKind;
                 return true;
               })) {
if (!IsImplicitClause) {
           reportOriginalDsa(*this, DSAStack, D, DVar);
           continue;
         }
       }
 
       // OpenMP 4.5 [2.15.5.1, Restrictions, p.3]
       // A list item cannot appear in both a map clause and a data-sharing
       // attribute clause on the same construct
-      if (isOpenMPTargetExecutionDirective(CurrDir)) {
+      //
+      // OpenMP 5.0 [2.19.7.1, Restrictions, p.7]
+      // A list item cannot appear in both a map clause and a data-sharing
+      // attribute clause on the same construct unless the construct is a
+      // combined construct.
+      if ((LangOpts.OpenMP <= 45 &&
+           isOpenMPTargetExecutionDirective(CurrDir)) ||
+          CurrDir == OMPD_target) {
         OpenMPClauseKind ConflictKind;
         if (DSAStack->checkMappableExprComponentListsForDecl(
                 VD, /*CurrentRegionOnly=*/true,
                 [&ConflictKind](
                     OMPClauseMappableExprCommon::MappableExprComponentListRef,
                     OpenMPClauseKind WhereFoundClauseKind) {
                   ConflictKind = WhereFoundClauseKind;
                   return true;
if (CKind == OMPC_map) {
             << getOpenMPSimpleClauseTypeName(OMPC_map, MapType)
             << getOpenMPDirectiveName(DKind);
         continue;
       }
 
       // OpenMP 4.5 [2.15.5.1, Restrictions, p.3]
       // A list item cannot appear in both a map clause and a data-sharing
       // attribute clause on the same construct
-      if (VD && isOpenMPTargetExecutionDirective(DKind)) {
+      //
+      // OpenMP 5.0 [2.19.7.1, Restrictions, p.7]
+      // A list item cannot appear in both a map clause and a data-sharing
+      // attribute clause on the same construct unless the construct is a
+      // combined construct.
+      if (VD && ((SemaRef.LangOpts.OpenMP <= 45 &&
+                  isOpenMPTargetExecutionDirective(DKind)) ||
+                 DKind == OMPD_target)) {
         DSAStackTy::DSAVarData DVar = DSAS->getTopDSA(VD, /*FromParent=*/false);
         if (isOpenMPPrivate(DVar.CKind)) {
           SemaRef.Diag(ELoc, diag::err_omp_variable_in_given_clause_and_dsa)
               << getOpenMPClauseName(DVar.CKind)
               << getOpenMPClauseName(OMPC_map)
               << getOpenMPDirectiveName(DSAS->getCurrentDirective());
           reportOriginalDsa(SemaRef, DSAS, CurDeclaration, DVar);
           continue;

cfe/trunk/lib/Sema/SemaStmt.cpp

else
     CurContext = CD;
 
   PushExpressionEvaluationContext(
       ExpressionEvaluationContext::PotentiallyEvaluated);
 }
 
 void Sema::ActOnCapturedRegionStart(SourceLocation Loc, Scope *CurScope,
                                     CapturedRegionKind Kind,
-                                    ArrayRef<CapturedParamNameType> Params) {
+                                    ArrayRef<CapturedParamNameType> Params,
+                                    unsigned OpenMPCaptureLevel) {
   CapturedDecl *CD = nullptr;
   RecordDecl *RD = CreateCapturedStmtRecordDecl(CD, Loc, Params.size());
 
   // Build the context parameter
   DeclContext *DC = CapturedDecl::castToDeclContext(CD);
   bool ContextIsFound = false;
   unsigned ParamNum = 0;
   for (ArrayRef<CapturedParamNameType>::iterator I = Params.begin(),
if (!ContextIsFound) {
     QualType ParamType = Context.getPointerType(Context.getTagDeclType(RD));
     auto *Param =
         ImplicitParamDecl::Create(Context, DC, Loc, ParamName, ParamType,
                                   ImplicitParamDecl::CapturedContext);
     DC->addDecl(Param);
     CD->setContextParam(ParamNum, Param);
   }
   // Enter the capturing scope for this captured region.
-  PushCapturedRegionScope(CurScope, CD, RD, Kind);
+  PushCapturedRegionScope(CurScope, CD, RD, Kind, OpenMPCaptureLevel);
 
   if (CurScope)
     PushDeclContext(CurScope, CD);
   else
     CurContext = CD;
 
   PushExpressionEvaluationContext(
       ExpressionEvaluationContext::PotentiallyEvaluated);

cfe/trunk/test/OpenMP/target_teams_distribute_parallel_for_lastprivate_messages.cpp

-// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized
 
-// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized
 
 typedef void **omp_allocator_handle_t;
 extern const omp_allocator_handle_t omp_default_mem_alloc;
 extern const omp_allocator_handle_t omp_large_cap_mem_alloc;
 extern const omp_allocator_handle_t omp_const_mem_alloc;
 extern const omp_allocator_handle_t omp_high_bw_mem_alloc;
 extern const omp_allocator_handle_t omp_low_lat_mem_alloc;
 extern const omp_allocator_handle_t omp_cgroup_mem_alloc;

 // expected-error@+1 {{lastprivate variable cannot be firstprivate}} expected-note@+1 {{defined as lastprivate}}
 #pragma omp target teams distribute parallel for lastprivate(n) firstprivate(n) // expected-error {{calling a private constructor of class 'S6'}}
   for (i = 0; i < argc; ++i) foo();
 
   static int si;
 #pragma omp target teams distribute parallel for lastprivate(si) // OK
   for (i = 0; i < argc; ++i) si = i + 1;
 
-#pragma omp target teams distribute parallel for lastprivate(k) map(k) // expected-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute parallel for' directive}} expected-note {{defined as lastprivate}}
+#pragma omp target teams distribute parallel for lastprivate(k) map(k) // le45-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute parallel for' directive}} le45-note {{defined as lastprivate}}
   for (i = 0; i < argc; ++i) foo();
 
   return foomain<S4, S5>(argc, argv); // expected-note {{in instantiation of function template specialization 'foomain<S4, S5>' requested here}}
 }

cfe/trunk/test/OpenMP/target_teams_distribute_parallel_for_simd_lastprivate_messages.cpp

-// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized
 
-// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized
 
 typedef void **omp_allocator_handle_t;
 extern const omp_allocator_handle_t omp_default_mem_alloc;
 extern const omp_allocator_handle_t omp_large_cap_mem_alloc;
 extern const omp_allocator_handle_t omp_const_mem_alloc;
 extern const omp_allocator_handle_t omp_high_bw_mem_alloc;
 extern const omp_allocator_handle_t omp_low_lat_mem_alloc;
 extern const omp_allocator_handle_t omp_cgroup_mem_alloc;

 // expected-error@+1 {{lastprivate variable cannot be firstprivate}} expected-note@+1 {{defined as lastprivate}}
 #pragma omp target teams distribute parallel for simd lastprivate(n) firstprivate(n) // expected-error {{calling a private constructor of class 'S6'}}
   for (i = 0; i < argc; ++i) foo();
 
   static int si;
 #pragma omp target teams distribute parallel for simd lastprivate(si) // OK
   for (i = 0; i < argc; ++i) si = i + 1;
 
-#pragma omp target teams distribute parallel for simd lastprivate(k) map(k) // expected-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute parallel for simd' directive}} expected-note {{defined as lastprivate}}
+#pragma omp target teams distribute parallel for simd lastprivate(k) map(k) // le45-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute parallel for simd' directive}} le45-note {{defined as lastprivate}}
   for (i = 0; i < argc; ++i) foo();
   return foomain<S4, S5>(argc, argv); // expected-note {{in instantiation of function template specialization 'foomain<S4, S5>' requested here}}
 }

cfe/trunk/test/OpenMP/target_teams_distribute_simd_firstprivate_messages.cpp

-// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized
 
-// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized
 
 typedef void **omp_allocator_handle_t;
 extern const omp_allocator_handle_t omp_default_mem_alloc;
 extern const omp_allocator_handle_t omp_large_cap_mem_alloc;
 extern const omp_allocator_handle_t omp_const_mem_alloc;
 extern const omp_allocator_handle_t omp_high_bw_mem_alloc;
 extern const omp_allocator_handle_t omp_low_lat_mem_alloc;
 extern const omp_allocator_handle_t omp_cgroup_mem_alloc;
#pragma omp target teams distribute simd firstprivate(S2::S2s)
   for (i = 0; i < argc; ++i) foo();
 
 #pragma omp target teams distribute simd firstprivate(S2::S2sc)
   for (i = 0; i < argc; ++i) foo();
 
 #pragma omp target teams distribute simd firstprivate(h) // expected-error {{threadprivate or thread local variable cannot be firstprivate}}
   for (i = 0; i < argc; ++i) foo();
 
-#pragma omp target teams distribute simd private(i), firstprivate(i) // expected-error {{private variable cannot be firstprivate}} expected-note 2 {{defined as private}}
-  for (i = 0; i < argc; ++i) foo(); // expected-error {{loop iteration variable in the associated loop of 'omp target teams distribute simd' directive may not be private, predetermined as linear}}
+#pragma omp target teams distribute simd private(i), firstprivate(i) // expected-error {{private variable cannot be firstprivate}} expected-note {{defined as private}}
+  // le45-note@-1 {{defined as private}}
+  // FIXME: Error is dropped for OpenMP 5.0.  Probably shouldn't be.
+  for (i = 0; i < argc; ++i) foo(); // le45-error {{loop iteration variable in the associated loop of 'omp target teams distribute simd' directive may not be private, predetermined as linear}}
 
 #pragma omp target teams distribute simd firstprivate(i)
   for (j = 0; j < argc; ++j) foo();
 
 #pragma omp target teams distribute simd firstprivate(i) // expected-note {{defined as firstprivate}}
   for (i = 0; i < argc; ++i) foo(); // expected-error {{loop iteration variable in the associated loop of 'omp target teams distribute simd' directive may not be firstprivate, predetermined as linear}}
 
 #pragma omp target teams distribute simd firstprivate(j)
   for (i = 0; i < argc; ++i) foo();
 
 // expected-error@+1 {{lastprivate variable cannot be firstprivate}} expected-note@+1 {{defined as lastprivate}}
 #pragma omp target teams distribute simd lastprivate(argc), firstprivate(argc)
   for (i = 0; i < argc; ++i) foo();
 
-#pragma omp target teams distribute simd firstprivate(argc) map(argc) // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} expected-note {{defined as firstprivate}}
+#pragma omp target teams distribute simd firstprivate(argc) map(argc) // le45-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} le45-note {{defined as firstprivate}}
   for (i = 0; i < argc; ++i) foo();
 
   return 0;
 }

cfe/trunk/test/OpenMP/target_teams_distribute_simd_lastprivate_messages.cpp

-// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized
 
-// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized
 
 typedef void **omp_allocator_handle_t;
 extern const omp_allocator_handle_t omp_default_mem_alloc;
 extern const omp_allocator_handle_t omp_large_cap_mem_alloc;
 extern const omp_allocator_handle_t omp_const_mem_alloc;
 extern const omp_allocator_handle_t omp_high_bw_mem_alloc;
 extern const omp_allocator_handle_t omp_low_lat_mem_alloc;
 extern const omp_allocator_handle_t omp_cgroup_mem_alloc;

 // expected-error@+1 {{lastprivate variable cannot be firstprivate}} expected-note@+1 {{defined as lastprivate}}
 #pragma omp target teams distribute simd lastprivate(n) firstprivate(n) // expected-error {{calling a private constructor of class 'S6'}}
   for (i = 0; i < argc; ++i) foo();
 
   static int si;
 #pragma omp target teams distribute simd lastprivate(si) // OK
   for (i = 0; i < argc; ++i) si = i + 1;
 
-#pragma omp target teams distribute simd lastprivate(k) map(k) // expected-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} expected-note {{defined as lastprivate}}
+#pragma omp target teams distribute simd lastprivate(k) map(k) // le45-error {{lastprivate variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} le45-note {{defined as lastprivate}}
   for (i = 0; i < argc; ++i) foo();
 
   return foomain<S4, S5>(argc, argv); // expected-note {{in instantiation of function template specialization 'foomain<S4, S5>' requested here}}
 }

cfe/trunk/test/OpenMP/target_teams_distribute_simd_private_messages.cpp

-// RUN: %clang_cc1 -verify -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp %s -Wuninitialized
 
-// RUN: %clang_cc1 -verify -fopenmp-simd %s -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd %s -Wuninitialized
 
 typedef void **omp_allocator_handle_t;
 extern const omp_allocator_handle_t omp_default_mem_alloc;
 extern const omp_allocator_handle_t omp_large_cap_mem_alloc;
 extern const omp_allocator_handle_t omp_const_mem_alloc;
 extern const omp_allocator_handle_t omp_high_bw_mem_alloc;
 extern const omp_allocator_handle_t omp_low_lat_mem_alloc;
 extern const omp_allocator_handle_t omp_cgroup_mem_alloc;
for (int k = 0; k < 10; ++k) {
     #pragma omp target teams distribute simd private(i)
     for (int x = 0; x < 10; ++x) foo();
   }
 
 #pragma omp target teams distribute simd firstprivate(i)
   for (int k = 0; k < 10; ++k) {
   }
 
-#pragma omp target teams distribute simd private(j) map(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} expected-note {{defined as private}}
+#pragma omp target teams distribute simd private(j) map(j) // le45-error {{private variable cannot be in a map clause in '#pragma omp target teams distribute simd' directive}} le45-note {{defined as private}}
   for (int k = 0; k < argc; ++k) ++k;
 
   return 0;
 }

cfe/trunk/test/OpenMP/target_teams_map_codegen.cpp

+// Test host codegen.
+// RUN: %clang_cc1 -DHAS_INT128 -verify -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s -check-prefixes=CHECK,HOST,INT128,HOST-INT128
+// RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefixes=CHECK,HOST,INT128,HOST-INT128
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s -check-prefixes=CHECK,HOST
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefixes=CHECK,HOST
+
+// Test target codegen - host bc file has to be created first.
+// RUN: %clang_cc1 -DHAS_INT128 -verify -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -DHAS_INT128 -verify -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s -check-prefixes=CHECK,INT128
+// RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t %s
+// RUN: %clang_cc1 -DHAS_INT128 -fopenmp -fopenmp-version=50 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s -check-prefixes=CHECK,INT128
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+// HOST: @[[MAPTYPES_FIRSTPRIVATE:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35]
+// HOST: @[[MAPTYPES_REDUCTION:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35]
+// HOST: @[[MAPTYPES_FROM:.offload_maptypes[0-9.]*]] = private {{.*}}constant [1 x i64] [i64 34]
+// HOST: @[[MAPTYPES_TO:.offload_maptypes[0-9.]*]] = private {{.*}}constant [1 x i64] [i64 33]
+// HOST: @[[MAPTYPES_ALLOC:.offload_maptypes[0-9.]*]] = private {{.*}}constant [1 x i64] [i64 32]
+// HOST: @[[MAPTYPES_ARRAY_R0:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35]
+// HOST: @[[MAPTYPES_ARRAY_R1:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 33, i64 33]
+// HOST-INT128: @[[MAPTYPES_INT128_R0:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 35, i64 35]
+// HOST-INT128: @[[MAPTYPES_INT128_R1:.offload_maptypes[0-9.]*]] = private {{.*}}constant [2 x i64] [i64 34, i64 34]
+//
+// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_PRIVATE:__omp_offloading_[^"\\]*mapWithPrivate[^"\\]*]]\00"
+// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_FIRSTPRIVATE:__omp_offloading_[^"\\]*mapWithFirstprivate[^"\\]*]]\00"
+// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_REDUCTION:__omp_offloading_[^"\\]*mapWithReduction[^"\\]*]]\00"
+// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_FROM:__omp_offloading_[^"\\]*mapFrom[^"\\]*]]\00"
+// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_TO:__omp_offloading_[^"\\]*mapTo[^"\\]*]]\00"
+// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_ALLOC:__omp_offloading_[^"\\]*mapAlloc[^"\\]*]]\00"
+// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_ARRAY_R0:__omp_offloading_[^"\\]*mapArray[^"\\]*]]\00"
+// CHECK: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_ARRAY_R1:__omp_offloading_[^"\\]*mapArray[^"\\]*]]\00"
+// INT128: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_INT128_R0:__omp_offloading_[^"\\]*mapInt128[^"\\]*]]\00"
+// INT128: @.omp_offloading.entry_name{{[0-9.]*}} = {{.*}} c"[[OFFLOAD_INT128_R1:__omp_offloading_[^"\\]*mapInt128[^"\\]*]]\00"
+
+// HOST: define {{.*}}mapWithPrivate
+// HOST: call {{.*}} @.[[OFFLOAD_PRIVATE]].region_id
+// HOST-NOT: offload_maptypes
+// HOST-SAME: {{$}}
+//
+// CHECK: define {{.*}} void @[[OFFLOAD_PRIVATE]]()
+// CHECK: call void ({{.*}}@[[OUTLINE_PRIVATE:.omp_outlined.[.0-9]*]]
+//
+// CHECK: define {{.*}} void @[[OUTLINE_PRIVATE]]({{.*}} %.global_tid., {{.*}} %.bound_tid.)
+void mapWithPrivate() {
+  int x, y;
+  #pragma omp target teams private(x) map(x,y) private(y)
+    ;
+}
+
+// HOST: define {{.*}}mapWithFirstprivate
+// HOST: call {{.*}} @.[[OFFLOAD_FIRSTPRIVATE]].region_id{{.*}} @[[MAPTYPES_FIRSTPRIVATE]]
+//
+// CHECK: define {{.*}} void @[[OFFLOAD_FIRSTPRIVATE]](i{{[0-9]*}}* {{[^,]*}}%x, i{{[0-9]*}}* {{[^,]*}}%y)
+// CHECK: call void ({{.*}}@[[OUTLINE_FIRSTPRIVATE:.omp_outlined.[.0-9]*]]
+//
+// CHECK: define {{.*}} void @[[OUTLINE_FIRSTPRIVATE]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x, i{{[0-9]*}} %y)
+void mapWithFirstprivate() {
+  int x, y;
+  #pragma omp target teams firstprivate(x) map(x,y) firstprivate(y)
+    ;
+}
+
+// HOST: define {{.*}}mapWithReduction
+// HOST: call {{.*}} @.[[OFFLOAD_REDUCTION]].region_id{{.*}} @[[MAPTYPES_REDUCTION]]
+//
+// CHECK: define {{.*}} void @[[OFFLOAD_REDUCTION]](i{{[0-9]*}}* {{[^,]*}}%x, i{{[0-9]*}}* {{[^,]*}}%y)
+// CHECK: call void ({{.*}}@[[OUTLINE_REDUCTION:.omp_outlined.[.0-9]*]]
+//
+// CHECK: define {{.*}} void @[[OUTLINE_REDUCTION]]({{[^)]*}} i{{[0-9]*}}* {{[^,]*}}%x, i{{[0-9]*}}* {{[^,]*}}%y)
+// CHECK: %.omp.reduction.red_list = alloca [2 x i8*]
+void mapWithReduction() {
+  int x, y;
+  #pragma omp target teams reduction(+:x) map(x,y) reduction(+:y)
+    ;
+}
+
+// HOST: define {{.*}}mapFrom
+// HOST: call {{.*}} @.[[OFFLOAD_FROM]].region_id{{.*}} @[[MAPTYPES_FROM]]
+//
+// CHECK: define {{.*}} void @[[OFFLOAD_FROM]](i{{[0-9]*}}* {{[^,]*}}%x)
+// CHECK: call void ({{.*}}@[[OUTLINE_FROM:.omp_outlined.[.0-9]*]]
+//
+// CHECK: define {{.*}} void @[[OUTLINE_FROM]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x)
+void mapFrom() {
+  int x;
+  #pragma omp target teams firstprivate(x) map(from:x)
+    ;
+}
+
+// HOST: define {{.*}}mapTo
+// HOST: call {{.*}} @.[[OFFLOAD_TO]].region_id{{.*}} @[[MAPTYPES_TO]]
+//
+// CHECK: define {{.*}} void @[[OFFLOAD_TO]](i{{[0-9]*}}* {{[^,]*}}%x)
+// CHECK: call void ({{.*}}@[[OUTLINE_TO:.omp_outlined.[.0-9]*]]
+//
+// CHECK: define {{.*}} void @[[OUTLINE_TO]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x)
+void mapTo() {
+  int x;
+  #pragma omp target teams firstprivate(x) map(to:x)
+    ;
+}
+
+// HOST: define {{.*}}mapAlloc
+// HOST: call {{.*}} @.[[OFFLOAD_ALLOC]].region_id{{.*}} @[[MAPTYPES_ALLOC]]
+//
+// CHECK: define {{.*}} void @[[OFFLOAD_ALLOC]](i{{[0-9]*}}* {{[^,]*}}%x)
+// CHECK: call void ({{.*}}@[[OUTLINE_ALLOC:.omp_outlined.[.0-9]*]]
+//
+// CHECK: define {{.*}} void @[[OUTLINE_ALLOC]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i{{[0-9]*}} %x)
+void mapAlloc() {
+  int x;
+  #pragma omp target teams firstprivate(x) map(alloc:x)
+    ;
+}
+
+// HOST: define {{.*}}mapArray
+// HOST: call {{.*}} @.[[OFFLOAD_ARRAY_R0]].region_id{{.*}} @[[MAPTYPES_ARRAY_R0]]
+// HOST: call {{.*}} @.[[OFFLOAD_ARRAY_R1]].region_id{{.*}} @[[MAPTYPES_ARRAY_R1]]
+//
+// CHECK: define {{.*}} void @[[OFFLOAD_ARRAY_R0]]([88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z)
+// CHECK: call void ({{.*}}@[[OUTLINE_ARRAY_R0:.omp_outlined.[.0-9]*]]
+//
+// CHECK: define {{.*}} void @[[OUTLINE_ARRAY_R0]]({{.*}} %.global_tid., {{.*}} %.bound_tid., [88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z)
+// CHECK: %.omp.reduction.red_list = alloca [1 x i8*]
+//
+// CHECK: define {{.*}} void @[[OFFLOAD_ARRAY_R1]]([88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z)
+// CHECK: call void ({{.*}}@[[OUTLINE_ARRAY_R1:.omp_outlined.[.0-9]*]]
+//
+// CHECK: define {{.*}} void @[[OUTLINE_ARRAY_R1]]({{.*}} %.global_tid., {{.*}} %.bound_tid., [88 x i{{[0-9]*}}]* {{[^,]*}}%y, [99 x i{{[0-9]*}}]* {{[^,]*}}%z)
+// CHECK: %.omp.reduction.red_list = alloca [1 x i8*]
+void mapArray() {
+  int x[77], y[88], z[99];
+  #pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(x,y,z)
+    ;
+  #pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(to:x,y,z)
+    ;
+}
+
+# if HAS_INT128
+// HOST-INT128: define {{.*}}mapInt128
+// HOST-INT128: call {{.*}} @.[[OFFLOAD_INT128_R0]].region_id{{.*}} @[[MAPTYPES_INT128_R0]]
+// HOST-INT128: call {{.*}} @.[[OFFLOAD_INT128_R1]].region_id{{.*}} @[[MAPTYPES_INT128_R1]]
+//
+// INT128: define {{.*}} void @[[OFFLOAD_INT128_R0]](i128* {{[^,]*}}%y, i128* {{[^,]*}}%z)
+// INT128: call void ({{.*}}@[[OUTLINE_INT128_R0:.omp_outlined.[.0-9]*]]
+//
+// INT128: define {{.*}} void @[[OUTLINE_INT128_R0]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i128* {{[^,]*}}%y, i128* {{[^,]*}}%z)
+// INT128: %.omp.reduction.red_list = alloca [1 x i8*]
+//
+// INT128: define {{.*}} void @[[OFFLOAD_INT128_R1]](i128* {{[^,]*}}%y, i128* {{[^,]*}}%z)
+// INT128: call void ({{.*}}@[[OUTLINE_INT128_R1:.omp_outlined.[.0-9]*]]
+//
+// INT128: define {{.*}} void @[[OUTLINE_INT128_R1]]({{.*}} %.global_tid., {{.*}} %.bound_tid., i128* {{[^,]*}}%y, i128* {{[^,]*}}%z)
+// INT128: %.omp.reduction.red_list = alloca [1 x i8*]
+void mapInt128() {
+  __int128 x, y, z;
+  #pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(x,y,z)
+    ;
+  #pragma omp target teams private(x) firstprivate(y) reduction(+:z) map(from:x,y,z)
+    ;
+}
+# endif
+#endif

cfe/trunk/test/OpenMP/target_teams_map_messages.cpp

-// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=40 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-version=45 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected -fopenmp-version=50 -fopenmp -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
 
-// RUN: %clang_cc1 -verify -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
-// RUN: %clang_cc1 -DCCODE -verify -fopenmp -ferror-limit 200 -x c %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -verify=expected,le45 -fopenmp-simd -ferror-limit 200 %s -Wno-openmp-target -Wuninitialized
+// RUN: %clang_cc1 -DCCODE -verify=expected,le45 -fopenmp -ferror-limit 200 -x c %s -Wno-openmp-target -Wuninitialized
 #ifdef CCODE
 void foo(int arg) {
   const int n = 0;
 
   double marr[10][10][10];
 
   #pragma omp target teams map(marr[2][0:2][0:2]) // expected-error {{array section does not specify contiguous storage}}
   {}

 
 #pragma omp target data map(always, tofrom: x)
 #pragma omp target data map(always: x) // expected-error {{missing map type}}
 #pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always', 'close', or 'mapper'}} expected-error {{missing map type}}
 #pragma omp target data map(always, tofrom: always, tofrom, x)
 #pragma omp target teams map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
   foo();
 
-#pragma omp target teams private(j) map(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target teams' directive}}  expected-note {{defined as private}}
+#pragma omp target teams private(j) map(j) // le45-error {{private variable cannot be in a map clause in '#pragma omp target teams' directive}}  le45-note {{defined as private}}
   {}
 
-#pragma omp target teams firstprivate(j) map(j)  // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams' directive}} expected-note {{defined as firstprivate}}
+#pragma omp target teams firstprivate(j) map(j) // le45-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams' directive}} le45-note {{defined as firstprivate}}
   {}
 
 #pragma omp target teams map(m)
   {}
 
   return tmain<int, 3>(argc)+tmain<from, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
 }
 #endif